具有 CUDA 内核的动态数据的 C 结构？

Posted 2023-02-21

【中文标题】具有 CUDA 内核的动态数据的 C 结构？

【英文标题】：C structures with dynamic data with CUDA kernels?

【发布时间】：2012-05-14 10:34:51

【问题描述】：

假设我有一个数据结构：

struct MyBigData 
    float * dataArray;
    float * targetArray;
    float * nodes;
    float * dataDataData;

我希望能够在一些不同的 CUDA 内核中传递这个结构。我不想将多个数组作为参数传递，所以我可以只传递结构并完成它吗？我知道内核支持 C 结构，但是 C 结构中的动态内存呢？

看来我只是为了在 CUDA 卡上制作结构：

MyBigData * mbd = (MyBigData *) cudaMalloc( sizeof(MyBigData) );

但是结构中的数组的动态内存呢？下面这一行编译但有一个运行时错误：

mbd->dataArray = (float *) cudaMalloc( 10 * sizeof(float) );

这是因为 cudaMalloc() 在 CPU 上运行，它无法读取 mdb->dataArray 来设置指针等于新的内存地址。所以存在运行时错误。但是，这可以编译并运行，但似乎不是我想要的：

MyBigData * mbd = (MyBigData *) malloc( sizeof(myBigData) );
mbd->dataArray = (float *) cudaMalloc( 10 * sizeof(float) );

因为现在，虽然这是有效的，但现在 mbd 驻留在主系统内存上，并且浮点指针指向在 CUDA 设备上分配的内存。所以我不能只传递一个指向 MyBigData 结构的指针，我必须将结构中的每个变量单独传递给内核。不干净。我想要的是：

someKernel<<<1,1>>>(mbd);

不是：

someKernel<<<1,1>>>(mbd->dataArray, mbd->targetArray, mbd->nodes, mbd->dataDataData);

所以我在想，cudaMemcpy() 怎么样？我在想这个：

MyBigData *d_mbd = cudaMemcpy( (void*) &d_mbd, (void*) mbd, SOMESIZE, CudaHostToDevice);

那么我应该为 SOMESIZE 放什么？我不能使用 sizeof(MyBigData)，因为这将包括浮点指针的大小，而不是数组的实际大小。其次，cudaMemcpy() 是否足够聪明，可以深入挖掘复杂数据结构的子对象？我认为不会。

那么，CUDA卡上不可能有包含动态内存的结构吗？或者我错过了什么。简单的方法是让 CUDA 内核分配一些内存，但不能从 CUDA 内核调用 cudaMalloc()。

想法？

更新 5 月 7 日： 我写了这段代码，它编译了，但它告诉我所有的值都是零。我认为我正在正确地创建对象并使用 CUDA 内核正确填充值。这些值只是线程 ID。我怀疑我没有正确打印这些值。想法？谢谢！

MyBigData* generateData(const int size) 
    MyBigData *mbd_host, *mbd_cuda;
    mbd_host = (MyBigData *) malloc( sizeof(MyBigData) );
    cudaMalloc( (void**) &mbd_host->dataArray, size * sizeof(float) );
    cudaMalloc( (void**) &mbd_host->targetArray, size * sizeof(float) );
    cudaMalloc( (void**) &mbd_host->nodes, size * sizeof(float) );
    cudaMalloc( (void**) &mbd_host->dataDataData, size * sizeof(float) );
    cudaMalloc( (void**) &mbd_cuda, sizeof(MyBigData) );
    cudaMemcpy( mbd_cuda, mbd_host, sizeof(mbd_host), cudaMemcpyHostToDevice );
    free(mbd_host);
    return mbd_cuda;


void printCudaData(MyBigData* mbd_cuda, const int size) 
    MyBigData *mbd;
    cudaMemcpy( mbd, mbd_cuda, sizeof(mbd_cuda), cudaMemcpyDeviceToHost);
    MyBigData *mbd_host = (MyBigData *) malloc( sizeof(MyBigData));
    mbd_host->dataArray = (float*) malloc(size * sizeof(float));
    mbd_host->targetArray = (float*) malloc(size * sizeof(float));
    mbd_host->nodes = (float*) malloc(size * sizeof(float));
    mbd_host->dataDataData = (float*) malloc(size * sizeof(float));

    cudaMemcpy( mbd_host->dataArray, mbd->dataArray, size * sizeof(float), cudaMemcpyDeviceToHost);
    cudaMemcpy( mbd_host->targetArray, mbd->targetArray, size * sizeof(float), cudaMemcpyDeviceToHost);
    cudaMemcpy( mbd_host->nodes, mbd->nodes, size * sizeof(float), cudaMemcpyDeviceToHost);
    cudaMemcpy( mbd_host->dataDataData, mbd->dataDataData, size * sizeof(float), cudaMemcpyDeviceToHost);

    for(int i = 0; i < size; i++) 
        printf("data[%i] = %f\n", i, mbd_host->dataArray[i]);
        printf("target[%i] = %f\n", i, mbd_host->targetArray[i]);
        printf("nodes[%i] = %f\n", i, mbd_host->nodes[i]);
        printf("data2[%i] = %f\n", i, mbd_host->dataDataData[i]);
    

    free(mbd_host->dataArray);
    free(mbd_host->targetArray);
    free(mbd_host->nodes);
    free(mbd_host->dataDataData);
    free(mbd_host);

这是我的内核和调用它的函数：

__global__ void cudaInitData(MyBigData* mbd) 
    const int threadID = threadIdx.x;
    mbd->dataArray[threadID] = threadID;
    mbd->targetArray[threadID] = threadID;
    mbd->nodes[threadID] = threadID;
    mbd->dataDataData[threadID] = threadID;


void initData(MyBigData* mbd, const int size) 
    if (mbd == NULL)
        mbd = generateData(size);

    cudaInitData<<<size,1>>>(mbd);

我的main() 电话：

MyBigData* mbd = NULL;
initData(mbd, 10);
printCudaData(mbd, 10);

【问题讨论】：

我不是 CUDA 开发人员，但听起来您所描述的内容与您所描述的方式不太可能 - 当您在两个谨慎的内存块之间共享指针时，事情只是不工作。 memcopy 系列函数需要一个连续的数据块，而这是您没有的。我很好奇的是常数 10- 如果你的数组总是长度为 10，为什么不将你的数据结构构建为 4 * ((sizeof(float*) + (10 * sizeof(float)))？

【参考方案1】：

其次，cudaMemcpy() 是否足够聪明，可以深入挖掘复杂数据结构的子对象？我认为不会。

你说得对，cudaMemcpy() 不会进行递归复制。为了实现你想要的，你应该这样做：

// Create mbd on host
MyBigData *mbd_host, *mbd;
mbd_host = (MyBigData *) malloc( sizeof(myBigData) );
// Fill it with pointers to device arrays
cudaMalloc( &mbd_host->dataArray, 10 * sizeof(float) );
// etc for other structure fields
// Create mbd on device
cudaMalloc( &mbd, sizeof(MyBigData) );
// Copy structure, filled with device addresses, to device memory
cudaMemcpy( mbd, mbd_host, sizeof(mbd), cudaMemcpyHostToDevice );
// Voila!

顺便说一句，最好不要将 MyBigData 结构存储在 __global__ 中，而是存储在设备的 __constant__ 内存中（您必须声明一个常量，而不是使用 cudaMalloc 分配 mbd并使用cudaMemcpyToSymbol 而不是最后一个cudaMemcpy)

【讨论】：

我在上面包含了一些代码。我不确定我是否正确打印出我的值，一切都只是零，但应该是 0-9，因为我用 10 个线程调用内核，并将值设置为线程 ID。我是否正确地从 GPU 检索数据以进行打印？

@RichardŻak 在printCudaData 你应该首先分配mbd。现在它只是指向无处的指针，将数据复制到它是“未定义的行为”。此外，总是检查来自cuda...函数的返回值，错误可能来自任何地方。