无法在 cuda 内核中使用我的模板 class

Question

我以为我知道如何编写一些干净的 cuda 代码。直到我尝试制作一个简单的模板 class 并在一个简单的内核中使用它。我这几天一直在解决问题。我访问的每一个线程都让我觉得有点愚蠢。

为了检查错误，我使用了这个

这是我的 class.h:

#pragma once
template <typename T>
class MyArray
{
public:
    const int size;

    T *data;

    __host__ MyArray(int size); //gpuErrchk(cudaMalloc(&data, size * sizeof(T)));

    __device__ __host__ T GetValue(int); //return data[i]
    __device__ __host__ void SetValue(T, int); //data[i] = val;
    __device__ __host__ T& operator()(int); //return data[i];

    ~MyArray(); //gpuErrchk(cudaFree(data));
};

template class MyArray<double>;

class.cu的相关内容在评论里。如果您认为整件事是相关的，我很乐意添加它。

现在主要 class:

__global__ void test(MyArray<double> array, double *data, int size)
{
    int j = threadIdx.x;
        //array.SetValue(1, j);  //doesn't work
        //array(j) = 1;  //doesn't work
        //array.data[j] = 1; //doesn't work
        data[j] = 1;   //This does work !
        printf("Reach this code\n");
    }
}
int main(int argc, char **argv)
{
    MyArray x(20);
    test<<<1, 20>>>(x, x.data, 20);

    gpuErrchk(cudaPeekAtLastError());
    gpuErrchk(cudaDeviceSynchronize());
}

当我说 "doesn't work" 时，我的意思是程序停在那里（在到达 printf 之前）而没有输出任何错误。另外，我从 cudaDeviceSynchronize 和 cudaFree 得到以下错误：

an illegal memory access was encountered

我无法理解的是，内存管理应该没有问题，因为将数组直接发送到内核工作正常。那么为什么当我发送 class 并尝试访问 classes 数据时它不起作用？为什么当我的代码明显遇到错误时却没有收到警告或错误消息？

这是nvcc --version

的输出

nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2017 NVIDIA Corporation
Built on Fri_Nov__3_21:07:56_CDT_2017
Cuda compilation tools, release 9.1, V9.1.85

Answer 1

（编者按：这个问题的评论中有不少虚假信息，所以我整理了一个答案作为社区维基条目。）

模板 class 不能作为参数传递给内核没有特别的原因。在这样做之前需要清楚地了解一些限制：

CUDA 内核参数，出于所有意图和目的，始终按值传递。在极其有限的情况下支持通过引用传递（所讨论的参数必须存储在托管内存中）。这不适用于此处。
作为 (1) 的结果，POD 参数可以正常工作，因为它们可以简单地复制并且不依赖于任何特殊行为
类不同，因为当您按值传递 class 时，您是在隐式调用复制构造或移动构造语义。这意味着作为内核参数按值传递的 classes 必须可以简单地复制构造。作为内核启动的一部分，无法在设备上运行非平凡的复制构造函数。
CUDA 进一步要求 classes 不包含虚拟成员
虽然 <<< >>> 内核启动语法看起来像一个简单的函数调用，但事实并非如此。在您在主机代码中编写的内容与主机端工具链实际发出的内容之间存在多层抽象样板和 API 调用。这意味着您的代码和 GPU 之间存在多个复制构造操作。如果您执行类似在析构函数中调用 cudaFree 之类的操作，您应该假设它将作为函数调用序列的一部分被调用，当其中一个副本超出范围时启动内核。你不想要那个。

你没有展示 class 成员函数在这种情况下是如何实际实现的，所以除了通过原始指向内核的指针，它起作用是因为它是一个简单可复制的 POD 值，而 class 几乎肯定不是。

这是一个简单、完整的示例，展示了如何进行这项工作：

$cat classy.cu
#include <vector>
#include <iostream>

#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }
inline void gpuAssert(cudaError_t code, const char *file, int line, bool abort=true)
{
    if (code != cudaSuccess)
    {
        fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);
        if (abort) exit(code);
    }
}

template <typename T>
class MyArray
{
    public:
        int len;
        T *data;

        __device__ __host__ void SetValue(T val, int i) { data[i] = val; };
        __device__ __host__ int size() { return sizeof(T) * len; };

        __host__ void DevAlloc(int N) {
            len = N;
            gpuErrchk(cudaMalloc(&data, size()));
        };

        __host__ void DevFree() {
            gpuErrchk(cudaFree(data));
            len = -1;
        };
};

__global__ void test(MyArray<double> array, double val)
{
    int j = threadIdx.x;
    if (j < array.len)
        array.SetValue(val, j);
}

int main(int argc, char **argv)
{
    const int N = 20;
    const double val = 5432.1;

    gpuErrchk(cudaSetDevice(0));
    gpuErrchk(cudaFree(0));

    MyArray<double> x;
    x.DevAlloc(N);

    test<<<1, 32>>>(x, val);
    gpuErrchk(cudaPeekAtLastError());
    gpuErrchk(cudaDeviceSynchronize());

    std::vector<double> y(N);
    gpuErrchk(cudaMemcpy(&y[0], x.data, x.size(), cudaMemcpyDeviceToHost));
    x.DevFree();

    for(int i=0; i<N; ++i) std::cout << i << " = " << y[i] << std::endl;

    return 0;
}

编译和运行像这样：

$ nvcc -std=c++11 -arch=sm_53 -o classy classy.cu
$ cuda-memcheck ./classy
========= CUDA-MEMCHECK
0 = 5432.1
1 = 5432.1
2 = 5432.1
3 = 5432.1
4 = 5432.1
5 = 5432.1
6 = 5432.1
7 = 5432.1
8 = 5432.1
9 = 5432.1
10 = 5432.1
11 = 5432.1
12 = 5432.1
13 = 5432.1
14 = 5432.1
15 = 5432.1
16 = 5432.1
17 = 5432.1
18 = 5432.1
19 = 5432.1
========= ERROR SUMMARY: 0 errors

（CUDA 10.2/gcc 7.5 在 Jetson Nano 上）

请注意，我已经包含了用于分配和释放的主机端函数，这些函数不与构造函数和析构函数交互。否则 class 与您的设计极为相似并且具有相同的属性。

无法在 cuda 内核中使用我的模板 class

Can't use my template class in cuda kernel

c++

cuda

nvcc