Shared memory between kernels?

Does any know if its possible to use shared memory between kernels, or between functions in the same kernel?

Something like this is what i’m hoping for…

extern shared int sharedMem;

device float kernel1_sub ()
{
int* neighbors = (int*) &sharedMem[ threadIdx.x31 ];
float
ndist = (float*) &sharedMem[ threadIdx.x31 + 31];
//…
}
global float kernel1 ()
{
int
neighbors = (int*) &sharedMem[ threadIdx.x31 ];
float
ndist = (float*) &sharedMem[ threadIdx.x31 + 31];
//…
kernel1_sub (); // invoke another function from same kernel
}
global float kernel2 () // invoke another kernel with same thread/block config
{
int
neighbors = (int*) &sharedMem[ threadIdx.x31 ];
float
ndist = (float*) &sharedMem[ threadIdx.x*31 + 31];
//…
}

Ideally, the neighbors & ndist arrays would be maintained between all of these functions. Is this possible? How about just for kernel1_sub, which is part of the same kernel call?

The docs say “shared between blocks”, but does this also mean between kernels with the same block/thread access pattern? What happens to the memory between kernel calls?

Function calls: Yes
Kernel calls: No