CUDA shared object between threads

Question

I am totally new to CUDA. I want to create one object on the device, and access its member from different threads. I use nvcc -arch=sm_20 (on Tesla M2090), and if I run my code I get an 'unspecified launch failure'. Here is my code:

#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;
}

You should check for errors after each kernel call and not at the end. This also applies to any CUDA API call. — pQB
Can you try running your application with cuda-memcheck and post the result ? — Vyas

SqrtPi SqrtPi · Accepted Answer · 2013-03-22T22:30:25

You get the 'unspecified launch failure' during your first kernel code because 'a' is a pointer stored in the host, but you want to give it a value from a device function. If you want to allocate the object on the device, than you first have to allocate a pointer on the device and than you can read and write it form device (kernel) code, but be careful because it will require double indirection.

Your code should looks like something like this (the rest of the functions should be modified similarly):

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

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

PS.: pQB is absolutely right about that, you should do an error check after each kernel code to detect the errors as soon as possible (and currently for finding the exact location of the error in your code)

CUDA shared object between threads

1 Answers