将内核函数的参数作为 C++ 结构传递？答案

【问题标题】：passing parameters of an kernel function as C++ struct?将内核函数的参数作为 C++ 结构传递？
【发布时间】：2013-01-18 09:20:18
【问题描述】：

我想将内核函数的参数作为给定的结构传递

struct kernel_data {
    double *A;
    double *B;
    double *C;
    const int *A_dims;
    const int *B_dims;
    int C_dims[2];
};

为此，我需要使用cudaMalloc() 初始化设备指针，但我该如何初始化。这个函数的这样一个结构，包括我打算传递的这些参数。还是我必须单独通过？

【问题讨论】：

标签： c++ struct cuda parameter-passing

【解决方案1】：

您可以通过复制将此结构传递给内核。请记住，传递给内核的参数总大小在 Fermi 之前的卡上不得超过 256B，在 Fermi 上不得超过 4KB。

因此，您必须使用 cudaMalloc 来分配设备内存，而不是在结构中设置指向设备内存的指针。最后，您通过复制将结构传递给内核。

我强烈建议您避免在 C++ 代码中使用此结构。而不是

struct kernel_data {
    double *A;
    double *B;
    double *C;
    const int *A_dims;
    const int *B_dims;
    int C_dims[2];
};

你应该这样做

class DeviceData{
public:   
    DeviceData(...){//Do cudaMalloc here}
    ~DeviceData(...){//Do cudaFree here}
private:
    double *_A;
    int _dims;    
};

此类将保存设备上可用的数据，并且是异常安全的。你可以实现一个可以传递给内核的包装器

class DeviceDataWrapper{
public:
    __host__ DeviceDataWrapper(DeviceData& device):
        _A(device._A),
        _dims(device._dims)
        {}

    __forceinline__ __device__ double* data(){return _A;}
    __forceinline__ __device__ int dims()const{return _dims;}

private:
    double *_A;
    int _dims;  
}

然后这样调用内核

__global__ void myKernel(DeviceDataWrapper a, DeviceDataWrapper b, DeviceData2Wrapper c){
 //do something like a.data()[0] = 1;
}

DeviceData A,B;
DeviceData2 C;
myKernel<<< >>>(A,B,C);

【讨论】：

该结构的副本大小将大大小于 256 字节（并且该限制无论如何仅适用于计算 1.x 功能的硬件）。
这只是一个提醒，如果您构建的层次结构可能会经常达到限制。

【解决方案2】：

只需像传递任何其他参数一样将结构按值传递给内核：

struct kernel_data args;

cudaMalloc(&(args.A), sizeof(double)*.....);
cudaMalloc(&(args.B), sizeof(double)*.....);
cudaMalloc(&(args.C), sizeof(double)*.....);
cudaMalloc(&(args.A_dims), sizeof(int)*.....);
cudaMalloc(&(args.B_dims), sizeof(int)*.....);

kernel<<<....>>>(args);

参数列表大小有一个理论上的限制，从 256 字节到 4Kb，取决于您使用的硬件，如果您超过它，请将 args 结构复制到设备分配并将其作为指针传递，或将其复制到一个常量内存指针。

要从主机初始化数组，只需使用标准的cudaMemcpy 调用：

cudaMemcpy(args.A, hostA, sizeof(double)*....., cudaMemcpyHostToDevice);

等等

【讨论】：