understanding the shared memory what does "static" mean for shared memory?

Accelerated Computing CUDA CUDA Programming and Performance
1

From the programming guide, “shared and constant variables have implied static storage”. my understandings from C, “static” means the memory persist across multiple function calls.

If I declare a shared array, initialize it with one global kernel, and run another global kernel to read/write to this memory, will the values of the shared memory consistant across muliple calls?

here is an example to illustrate my question (taken from this thread)

__shared__ uint g_seeds[N];

__global__ void mt_rand_init(void) {

  // initialize g_seeds[idx] here ...

}

__global__ void genrand(float2 result[]){

  // take g_seeds[] values, do some math 

  // and save back to g_seeds[]

}

main(){

  ...

  mt_rand_init<<<gridsize,blocksize>>>();

  for(i=0;i<10;i++){

	   genrand<<<gridsize,blocksize>>>(res);

  }

  ...

}

if I initialize g_seeds in mt_rand_init(), shall I expect to see the same values in genrand()? and across the multiple calls of genrand()?

what would happen if gridsize/blocksize are different for mt_rand_init and genrand?

2

In normal C code, static variables have the lifetime of the entire program, instead of having the lifetime of individual functions.

But shared memory only has the lifetime of a single block, and when a block finishes, subsequent blocks that execute on the MP may or may not be able to sensibly use the values. So I don’t think it’s reasonable to expect the values to persist across multiple kernels. But within a single block of a single kernel, they will survive across multiple device function calls.

3

“implied static storage” merely means that those values exist only within the compilation unit in which they are defined. Thus, you cannot declare a global device variable in a header file and use it in multiple .cu files unless you #include them all together and only run nvcc once to compile.