线程间的CUDA共享对象答案

【问题标题】：CUDA shared object between threads线程间的CUDA共享对象
【发布时间】：2013-03-23 06:30:25
【问题描述】：

我对 CUDA 完全陌生。我想在设备上创建一个对象，并从不同的线程访问它的成员。我使用 nvcc -arch=sm_20 （在 Tesla M2090 上），如果我运行我的代码，我会得到一个“未指定的启动失败”。这是我的代码：

#include <stdio.h>
#include <string>

using namespace std;

#ifdef __CUDACC__
#define CUDA_CALLABLE __host__ __device__
#else
#define CUDA_CALLABLE
#endif

class SimpleClass {
public:
    int i;
    CUDA_CALLABLE SimpleClass(){i=1;};
    CUDA_CALLABLE ~SimpleClass(){};
};

__global__ void initkernel(SimpleClass *a){
    a = new SimpleClass();
}
__global__ void delkernel(SimpleClass *a){
    delete a;
}
__global__ void kernel(SimpleClass *a){
printf("%d\n", a->i);
}

int main() {
    SimpleClass *a;
    initkernel<<<1,1>>>(a);
    cudaThreadSynchronize();
    kernel<<<1,10>>>(a);
    cudaThreadSynchronize();
    delkernel<<<1,1>>>(a);
    cudaThreadSynchronize();

    cudaError_t error = cudaGetLastError();
    string lastError = cudaGetErrorString(error);
    printf("%s\n",lastError.c_str());
    return 0;
}

【问题讨论】：

您应该在each 内核调用之后而不是最后检查错误。这也适用于任何 CUDA API 调用。
您可以尝试使用 cuda-memcheck 运行您的应用程序并发布结果吗？

标签： cuda

【解决方案1】：

您在第一个内核代码期间收到“未指定的启动失败”，因为“a”是存储在主机中的指针，但您想从设备函数中为其指定一个值。如果要在设备上分配对象，首先必须在设备上分配一个指针，然后才能从设备（内核）代码中读取和写入它，但要小心，因为它需要双重间接。

您的代码应如下所示（其余功能应类似修改）：

__global__ void initkernel(SimpleClass** a){
    *a = new SimpleClass();
}

int main() {
    SimpleClass** a;
    cudaMalloc((void**)&a, sizeof(SimpleClass**));
    initkernel<<<1,1>>>(a);
    cudaThreadSynchronize();
}

PS.：pQB 是绝对正确的，你应该在每个内核代码之后做一个错误检查，以尽快检测到错误（目前是为了在你的代码中找到错误的确切位置）

【讨论】：