访问不同 CUDA 内核中的类成员

Posted 2023-03-23

【中文标题】访问不同 CUDA 内核中的类成员

【英文标题】：Accessing Class Member in different CUDA kernels

【发布时间】：2018-05-29 21:00:36

【问题描述】：

我有一个纯 GPU 类 T，我想在 GPU 上创建它，但在 CPU 上有一个对它的引用，因此我可以将链接作为参数发送到不同的 CUDA 内核。

class T

public:
    int v;
public:
    __device__ T()  v = 10; 
    __device__ ~T() 
    __device__ int compute()  return v; 
;

这是我用来创建类实例和调用compute() 函数的内核。

__global__ void kernel(T* obj, int* out)

    if(blockIdx.x * blockDim.x + threadIdx.x == 0) 
        out[0] = obj->compute(); // no kernel error, but it returns garbage
    


__global__ void cudaAllocateGPUObj(T* obj)

    if(blockIdx.x * blockDim.x + threadIdx.x == 0) 
        obj = new T;
        // if I call `out[0] = obj->compute();` here, everything works fine
    

主函数只是为T* 类型的指针分配内存，该指针稍后用作cudaAllocateGPUObj 的参数。

int main()

    int cpu, *gpu;
    cudaMalloc((void**)&gpu, sizeof(int));
    T* obj;
    cudaMalloc((void**)&obj, sizeof(T*));
    cudaAllocateGPUObj<<<1,1>>>(obj);
    kernel<<<1,1>>>(obj, gpu);
    cudaMemcpy(&cpu, gpu, sizeof(int), cudaMemcpyDeviceToHost);
    cudaDeviceSynchronize();
    printf("cudaMemcpy\nresult: %d\n", cpu);
    return 0;

这段代码的问题（在代码中的cmets中指定）是，当我在cudaAllocateGPUObj内核中调用out[0] = obj->compute();并将获得的值传输到CPU时，一切都是正确的。但是如果我想在另一个内核中获取成员值，它就会变成垃圾，但是如果我将返回值从v 变量更改为常量，一切正常。

你能告诉我这段代码有什么问题吗？

【参考方案1】：

当您将参数传递给 CUDA 内核时，它是一种按值传递的机制。你已经从一个指向对象的指针开始：

T* obj;

然后，不是为对象分配存储空间，而是为另一个指针分配存储空间：

cudaMalloc((void**)&obj, sizeof(T*));

所以我们在这里走错了路。 （此时这是一个逻辑 C 编程错误。）接下来，在分配内核中，obj 参数（现在指向 GPU 内存空间中的某个位置）按值传递：

__global__ void cudaAllocateGPUObj(T* obj)
                                      ^^^ pass-by-value: local copy is made

现在，当你这样做时：

        obj = new T;

您创建一个新指针，并用该新指针覆盖obj 的本地副本。所以这当然可以在本地工作，但是调用环境中obj 的副本不会用那个新指针更新。

解决此问题的一种可能方法是创建适当的指针对指针方法：

$ cat t5.cu
#include <stdio.h>

class T

public:
    int v;
public:
    __device__ T()  v = 10; 
    __device__ ~T() 
    __device__ int compute()  return v; 
;

__global__ void kernel(T** obj, int* out)

    if(blockIdx.x * blockDim.x + threadIdx.x == 0) 
        out[0] = (*obj)->compute(); 
    


__global__ void cudaAllocateGPUObj(T** obj)

    if(blockIdx.x * blockDim.x + threadIdx.x == 0) 
        *obj = new T;
    


int main()

    int cpu, *gpu;
    cudaMalloc((void**)&gpu, sizeof(int));
    T** obj;
    cudaMalloc(&obj, sizeof(T*));
    cudaAllocateGPUObj<<<1,1>>>(obj);
    kernel<<<1,1>>>(obj, gpu);
    cudaMemcpy(&cpu, gpu, sizeof(int), cudaMemcpyDeviceToHost);
    cudaDeviceSynchronize();
    printf("cudaMemcpy\nresult: %d\n", cpu);
    return 0;


$ nvcc -arch=sm_35 -o t5 t5.cu
$ cuda-memcheck ./t5
========= CUDA-MEMCHECK
cudaMemcpy
result: 10
========= ERROR SUMMARY: 0 errors
$