CODEWITHSUNDEEP
multithreadingcuda
Edd
Edd

Reputation: 107

Local memory for each CUDA thread

I have a simple program below. My question is that where is "temp" actually stored? is it in global or local memory? I need array temp for each idx so that every thread has individual array temp. In this case, it is working properly. But in my actual program, when I tried to fill temp[0] from test2 it made the program stopped. Suppose we have 1024 threads then it only run the kernel around 200 threads. So, I am wondering whether temp is shared or not. If yes, maybe there is a collision there. I also did not get any error messsage. Please someone explain about this.

__device__ void test2(int temp[], int idx) {
    temp[0] = idx;
    printf("%d ", temp[0]);
}

__global__ void test() {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;

    int *temp = (int *) malloc(100 * sizeof (int));

    test2(temp, idx);
}

int main() {
    test << <1, 1024 >> >();
    return 0;
}

Upvotes: 1

Views: 1105

Answers (1)

Robert Crovella
Robert Crovella

Reputation: 151799

My question is that where is "temp" actually stored?

The allocation for temp is stored in a place called the device heap. It is a form of global memory. However the temp variable itself (i.e. the pointer value) is in local memory - not shared or visible to other threads.

I need array temp for each idx so that every thread has individual array temp.

You will get that, subject to caveats below. Each thread will have its own individual array, referenced by its local variable temp. Each thread will have a separate allocation for storage on the device heap.

People commonly have problems with in-kernel new or malloc. One of the main reasons is that the device heap is initially limited to 8MB, across all of your device heap allocations. So if enough threads do a new or malloc of enough allocation requests, you will run out of space.

When you run out of space, the API way to signal that is to return a zero pointer value for the allocation (a NULL pointer). If you then attempt to use this NULL pointer, you will have trouble.

For debugging purposes (i.e. to prove this is happening), test the pointer for NULL (i.e. == 0) before using it. If it is NULL, don't use it (perhaps print an error message instead).

You can read more about this in the documentation or in many questions here on the SO cuda tag. If you read any of these sources, you will discover that you can increase the size of the device heap.

Upvotes: 3

Related Questions