从 cuda 内核中访问类数据成员 - 如何设计正确的主机/设备交互？

Question

我一直在尝试将一些 cuda/C 代码转换为更面向对象的代码，但就我目前对 cuda 功能机制的理解而言，我的目标似乎并不容易实现。在这种情况下，我也找不到很好的解释。毕竟这可能是不可能的。

我有一个 myClass 类的 global 对象，其中包含一个要填充到内核中的数组。

应该如何定义 myClass 中的方法，以便数组和布尔成员对设备可见，然后可以将数组复制回主机?我使用的是cuda 7.5，我的卡的计算能力是3.5。

这是描述情况的暂定结构：

#include <cstdio>
#include <cuda.h>
#include <cuda_runtime.h>

class myClass
{
public:
        bool bool_var;    // Set from host and readable from device
        int  data_size;   // Set from host
        __device__ __host__ myClass();
        __device__ __host__ ~myClass();
        __host__ void setValues(bool iftrue, int size);
        __device__ void dosomething(int device_parameter);
        __host__ void export();

        // completely unknown methods
        __host__ void prepareDeviceObj();
        __host__ void retrieveDataToHost();
private:
        int *data; // Filled in device, shared between threads, at the end copied back to host for data output
};

__host__ __device__ myClass::myClass()
{
}

__host__ __device__ myClass::~myClass()
{
#ifdef __CUDACC__
        if(bool_var)
                cudaFree(data);
#else
        free(data);
#endif
}

__host__ void myClass::setValues(bool iftrue, int size)
{
        bool_var  = iftrue;
        data_size = size;
}

__device__ void myClass::dosomething(int idx)
{
        int toadd = idx+data_size;
        atomicAdd(&data[idx], toadd); // data should be unique among threads
}


__global__ void myKernel(myClass obj)
{
        const int idx = blockIdx.x*blockDim.x + threadIdx.x;
        if(idx < obj.data_size)
        {
                if(!obj.bool_var)
                        printf("Object is not up to any task here!");
                else
                {
                        printf("Object is ready!");
                        obj.dosomething(idx);
                }
        }
}


myClass globalInstance;

int main(int argc, char** argv)
{
        int some_number = 40;
        globalInstance.setValues(true, some_number);
        globalInstance.prepareDeviceObj();           // unknown
        myKernel<<<1,some_number>>>(globalInstance); // how to pass the object?
        globalInstance.retrieveDataToHost();         // unknown
        globalInstance.export();
        exit(EXIT_SUCCESS);
}

Answer 1

你的方法应该是可行的。当您通过值作为内核参数传递对象时（如您所指出的），实际上不需要进行太多与从主机到设备的传输相关的设置。

您需要在主机和设备上正确分配数据，并在适当的点使用cudaMemcpy类型的操作来移动数据，就像在普通的CUDA程序中一样。

在全局范围内声明对象时需要注意的一点是，建议不要在对象的构造函数或析构函数中使用 CUDA API 调用。原因已覆盖here，这里不再赘述。尽管这种处理主要关注在 main 之前启动的内核，但 CUDA 延迟初始化也会影响在 main 范围之外执行的任何 CUDA API 调用，这适用于在全局范围内实例化的对象的构造函数和析构函数。

以下是您所展示内容的充实示例。我基本上没有更改您已经编写的代码，只是为您未编写的代码添加了一些方法定义。这里显然有很多不同的可能方法。有关更多示例，您可能需要查看CUDA C++ integration sample code。

这是一个围绕您所展示内容的工作示例：

$ cat t1236.cu
#include <cstdio>

class myClass
{
public:
        bool bool_var;    // Set from host and readable from device
        int  data_size;   // Set from host
        __host__ myClass();
        __host__ ~myClass();
        __host__ void setValues(bool iftrue, int size);
        __device__ void dosomething(int device_parameter);
        __host__ void export_data();

        // completely unknown methods
        __host__ void prepareDeviceObj();
        __host__ void retrieveDataToHost();
private:
        int *data; // Filled in device, shared between threads, at the end copied back to host for data output
        int *h_data;
};

__host__ myClass::myClass()
{
}

__host__ myClass::~myClass()
{
}

__host__ void myClass::prepareDeviceObj(){
        cudaMemcpy(data, h_data, data_size*sizeof(h_data[0]), cudaMemcpyHostToDevice);
}
__host__ void myClass::retrieveDataToHost(){
        cudaMemcpy(h_data, data, data_size*sizeof(h_data[0]), cudaMemcpyDeviceToHost);
}

__host__ void myClass::setValues(bool iftrue, int size)
{
        bool_var  = iftrue;
        data_size = size;
        cudaMalloc(&data, data_size*sizeof(data[0]));
        h_data = (int *)malloc(data_size*sizeof(h_data[0]));
        memset(h_data, 0, data_size*sizeof(h_data[0]));
}

__device__ void myClass::dosomething(int idx)
{
        int toadd = idx+data_size;
        atomicAdd(&(data[idx]), toadd); // data should be unique among threads
}
__host__ void myClass::export_data(){
        for (int i = 0; i < data_size; i++) printf("%d ", h_data[i]);
        printf("\n");
        cudaFree(data);
        free(h_data);
}


__global__ void myKernel(myClass obj)
{
        const int idx = blockIdx.x*blockDim.x + threadIdx.x;
        if(idx < obj.data_size)
        {
                if(!obj.bool_var)
                        printf("Object is not up to any task here!");
                else
                {
                        //printf("Object is ready!");
                        obj.dosomething(idx);
                }
        }
}


myClass globalInstance;

int main(int argc, char** argv)
{
        int some_number = 40;
        globalInstance.setValues(true, some_number);
        globalInstance.prepareDeviceObj();
        myKernel<<<1,some_number>>>(globalInstance);
        globalInstance.retrieveDataToHost();
        globalInstance.export_data();
        exit(EXIT_SUCCESS);
}
$ nvcc -o t1236 t1236.cu
$ cuda-memcheck ./t1236
========= CUDA-MEMCHECK
40 41 42 43 44 45 46 47 48 49 50 51 52 53 54 55 56 57 58 59 60 61 62 63 64 65 66 67 68 69 70 71 72 73 74 75 76 77 78 79
========= ERROR SUMMARY: 0 errors
$

Answer 2

对我来说最有效的方法是只将常规 CUDA 函数、内核和内核启动放在 CUDA C (.cu) 文件中，然后在此之上构建一个面向对象的接口，使用 C++ 中的类 (@987654322 @) 文件。

因此，在您的类构造函数中，您调用 .cu 文件中分配和初始化内存的函数，在您的方法中，您调用启动内核的函数等。

这也可以加快开发周转速度，因为您可以经常更改您的类而无需重新编译 .cu 文件，这比编译纯 .cpp 文件要慢得多。