CODEWITHSUNDEEP
cuda
gpuguy
gpuguy

Reputation: 4585

In cuda which thread is responsible for shared memory allocation requested inside a kernel?

I understand instructions inside a kernel is executed by all the threads. Let us consider the following case:

_global__ void staticReverse(int *d, int n)
{
  __shared__ int s[64];
  int t = threadIdx.x;
  int tr = n-t-1;
  s[t] = d[t];
  __syncthreads();
  d[t] = s[tr];
}

Basically this code will be running at different cores as thread.
Now there is a shared memory allocation. Since it(shared memory allocation) will be encountered by all the threads, will it be allocated by all the threads? (Logically not.) But I am sure at least one thread must allocate it. I want to know which thread do it ? Kindly help me understand where my understanding is wrong.

Upvotes: 1

Views: 175

Answers (1)

Robert Crovella
Robert Crovella

Reputation: 152174

No threads are responsible for this allocation - the threads do not run any SASS code that is involved in allocating this memory.

The same statement is true if you use dynamic (extern) shared allocation - no threads are responsible - meaning the threads do not run any SASS code that is involved in allocating this memory. There are no function calls or other mechanisms involved.

The memory is already allocated, and a pointer to it is already established, by the time the thread SASS code (i.e. the kernel) begins executing.

There is a wrinkle to be aware of. If the shared memory declaration involves a constructor, then the constructor will be run on all threads. This can be confusing behavior.

Upvotes: 2

Related Questions