c++ - 从 cuda 内核访问类数据成员——如何设计适当的主机/设备交互？

Question

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

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

myClass 中的方法应该如何定义，以便数组和 bool 成员从设备可见，然后数组可以复制回主机 ?我使用的是 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
$