Extern pointer shared memory

Accelerated Computing CUDA CUDA Programming and Performance
1

I’m trying use pointer to a struct in shared memory but I don’t know how to setting memory size for each block since their size differ between blocks, something like:

struct reg{
   int *val;
   int sizeval;
};

__shared__ reg data;

__global__ void kernel1(){
   //initialize data for each block
}

__global__ void kernel2(){
   //use same data pointer here but when I try use *val gets memory error
}

int main(){
   kernel1<<< nblocks, nthreads >>>(); 
   kernel2<<< nblocks, nthreads >>>(); //here I need the size of shared memory to data
                                       //but how can pass size per block?
}

There is a way to pass different size to shared memory or allocate data as static in kernel1() and use same variable in kernel2()?

2

Shared memory is not persistent between kernels (not even between blocks), so you can’t do this with shared memory.
Use normal device memory instead, or merge the two kernels into a single one.

3

If is not persistent then why have extern shared memory (Programming guide B.2.3)?

4

Read the documentation you cited and you’ll find that “extern” on shared memory variables does something completely different than you expect - it defers the decision over their size until runtime.

Admittedly that is abusing this keyword that would otherwise have no sensible meaning in the context of shared memory.
Have you read introductory section 2 about the CUDA programming model?