Could shared variable be accessed by multiples launches of same kernel?

Question

I have a doubt about the shared variables in kernel. Assuming I launch the same kernel several times like this:

if (index>i && index<n){
  K<<<1,n>>>(i, n, n);
}

If I declare inside kernel "K" a extern __shared__ variable like this:

__global__ K(int i, int n){
  ...
  extern __shared__ int test[];
  ...
}

Do the all threads of all kernels inside block can be access it? Or the shared variable is exclusive to a kernel?

The "all kernels inside block" part makes no sense. Kernel launch consists of one or multiple thread blocks, not the other way around. Shared variables are accessible to all threads of a thread block, and are not accessible otherwise. — void_ptr, Sep 01 '15 at 23:52
shared memory is also only in scope for a particular kernel call. it does not persist between kernels. — Robert Crovella, Sep 02 '15 at 00:28
Ok! So, the shared variables are exclusive of each kernel, all right? Since I launch several times the same kernel, the threads of a kernel will access only your "test" shared variable, correct? Thank you very much, guys! — Lucas Torquato, Sep 02 '15 at 15:15

score 0 · Answer 1 · edited Sep 02 '15 at 07:42

There are three memory types in cuda: Per-thread local memory, Per-block shared memory and global memory.
__shared__ variables are stored in per-block memory. All threads in a same block can access the same __shared__ var.
If the variable need to live between kernel calls, there is a __device__ qualifier, which indicates the var should be stored in global space and have lifetime of the application.

1 Answers1