内核上的 CUDA 非法内存访问

Question

我正在尝试在 CUDA 上实现一个基本的设备数组类型，作为练习。作为设计目标，它应该模仿 std::array 接口。在实现operator+ 时，我遇到了非法内存访问错误，我无法解释原因。 这是代码。

#include <iostream>
#include <array>

enum class memcpy_t {
    host_to_host,
    host_to_device,
    device_to_host,
    device_to_device
};

bool check_cuda_err() {
    cudaError_t err = cudaGetLastError();
    if(err == cudaSuccess) {
        return true;
    }
    else {
        std::cerr << "Cuda Error: " << cudaGetErrorString(err) << "\n" << std::flush;
        return false;
    }
}

template <typename T, std::size_t N>
struct cuda_allocator {
    using pointer = T*;

    static void allocate(T *&dev_mem) {
        cudaMalloc(&dev_mem, N * sizeof(T));
    }

    static void deallocate(T *dev_mem) {
        cudaFree(dev_mem);
    }

    template <memcpy_t ct>
    static void copy (T *dst, T *src) {
        switch(ct) {
        case memcpy_t::host_to_host:
            cudaMemcpy(dst, src, N * sizeof(T), cudaMemcpyHostToHost);
            break;
        case memcpy_t::host_to_device:
            cudaMemcpy(dst, src, N * sizeof(T), cudaMemcpyHostToDevice);
            break;
        case memcpy_t::device_to_host:
            cudaMemcpy(dst, src, N * sizeof(T), cudaMemcpyDeviceToHost);
            break;
        case memcpy_t::device_to_device:
            cudaMemcpy(dst, src, N * sizeof(T), cudaMemcpyDeviceToDevice);
            break;
        default:
            break;
        }
    }
};

template <typename T, std::size_t N>
struct gpu_array {
    using allocator = cuda_allocator<T, N>;
    using pointer = typename allocator::pointer;
    using value_type = T;
    using iterator = T*;
    using const_iterator = T const*;

    gpu_array() {
       allocator::allocate(data);
    }

    gpu_array(std::array<T, N> host_arr) {
        allocator::allocate(data);
        allocator::template copy<memcpy_t::host_to_device>(data, host_arr.begin());
    }

    gpu_array& operator=(gpu_array const& o) {
        //allocator::allocate(data);
        allocator::template copy<memcpy_t::device_to_device>(data, o.begin());
    }

    operator std::array<T, N>() {
        std::array<T, N> res;
        allocator::template copy<memcpy_t::device_to_host>(res.begin(), data);
        return res;
    }

    ~gpu_array() {
        allocator::deallocate(data);
    }

    __device__ iterator begin() { return data; }
    __device__ iterator end() { return data + N; }
    __device__ const_iterator begin() const { return data; }
    __device__ const_iterator end() const { return data + N; }

private:
    T* data;
};

template <typename T, std::size_t N>
__global__ void add_kernel(gpu_array<T,N> **r,
                           gpu_array<T,N> const* a1,
                           gpu_array<T,N> const* a2) {
    int i = blockIdx.x*blockDim.x + threadIdx.x;
    printf("Index: %d\n", i);
    (*r)->begin()[i] = a1->begin()[i] + a2->begin()[i];
}

template <typename T, std::size_t N>
gpu_array<T, N> operator+(gpu_array<T,N> const&a1,
                          gpu_array<T,N> const&a2)
{
    gpu_array<T, N> *res = new gpu_array<T, N>;
    add_kernel<<<(N+3)/4, 4>>>(&res, &a1, &a2);
    cudaDeviceSynchronize();
    check_cuda_err();
    // ignore memory leak for now
    return *res;
}
const int N = 1<<3;

int main() {
    std::array<float, N> x,y;

    for (int i = 0; i < N; i++) {
        x[i] = 1.0f;
        y[i] = 2.0f;
    } 

    gpu_array<float, N> dx{x};
    gpu_array<float, N> dy{y};
    check_cuda_err(); // shows no error for memcpy
    std::array<float, N> res = dx + dy;

    for(const auto& elem : res) {
        std::cout << elem << ", ";
    }
}

我正在创建一个大小为 8 的数组来测试。如您所见，cuda_check_err() 从主机阵列初始化 gpu_array 后没有显示错误。我猜复制数据可以正常工作。但是在内核中，当我索引设备数组时，我收到illegal memory access 错误。这是输出：

索引：0

索引：1

索引：2

索引：3

索引：4

索引：5

索引：6

索引：7

Cuda 错误：遇到非法内存访问

9.45143e-39, 0, 6.39436e-39, 0, 0, 0, 0, 0,

如您所见，我已经为每个线程打印了计算索引，似乎没有什么超出范围。那么，什么可能导致这种非法内存访问错误呢？顺便说一句，cuda-memchecksays：

无效的全局读取大小为 8

以后

地址 0x7fff9f4c6ec0 超出范围

但我已经打印了索引，不知道为什么它超出了范围。

Answer 1

我们在这个问题中看到了两个版本的代码，不幸的是，两个版本都有相同问题的不同版本。

第一次使用引用作为内核的参数：

template <typename T, std::size_t N>
 __global__ void add_kernel(gpu_array<T,N> &r,
                       gpu_array<T,N> const&a1,
                       gpu_array<T,N> const&a2) {
    int i = blockIdx.x*blockDim.x + threadIdx.x;
    printf("Index: %d\n", i);
    r.begin()[i] = a1.begin()[i] + a2.begin()[i];
}

template <typename T, std::size_t N>
gpu_array<T, N> operator+(gpu_array<T,N> const&a1,
                      gpu_array<T,N> const&a2)
{
    gpu_array<T, N> res;
    add_kernel<<<(N+3)/4, 4>>>(res, a1, a2);
    cudaDeviceSynchronize();
    check_cuda_err();
    return res;
 }

虽然这很简洁，并且在 CUDA 内核代码中完全支持引用，但从主机通过引用传递内核参数最终会将主机地址作为设备中的参数，因为 CUDA 工具链，就像我所使用的所有其他 C++ 编译器一样意识到，使用指针实现引用。结果是非法地址的内核运行时错误。

第二个使用指针间接而不是引用，并最终将主机指针传递给 GPU，这与第一个版本几乎相同：

template <typename T, std::size_t N>
__global__ void add_kernel(gpu_array<T,N> **r,
                           gpu_array<T,N> const* a1,
                           gpu_array<T,N> const* a2) {
    int i = blockIdx.x*blockDim.x + threadIdx.x;
    printf("Index: %d\n", i);
    (*r)->begin()[i] = a1->begin()[i] + a2->begin()[i];
}

template <typename T, std::size_t N>
gpu_array<T, N> operator+(gpu_array<T,N> const&a1,
                          gpu_array<T,N> const&a2)
{
    gpu_array<T, N> *res = new gpu_array<T, N>;
    add_kernel<<<(N+3)/4, 4>>>(&res, &a1, &a2); 
    cudaDeviceSynchronize();
    check_cuda_err();
    // ignore memory leak for now
    return *res;
}

将此结构直接传递给设备内核的唯一安全实现是使用按值传递。然而，这意味着副本将超出范围并触发销毁，这将释放支持数组的内存并导致不同类型的意外错误。