dynamic shared memory and garbage collection

Hi,

I am using dynamic shared memory, here’s some code:

__global__ kernelOne (){

  extern __shared__ int smemPoint[];

  int *smem = smemPoint;

}

__global__ kernelTwo (){

  extern __shared__ int smemPoint[];

  int *smem = smemPoint;

}

void wrapper(){

   kernelOne <<<Grids, Blocks>>>();

   kernelTwo <<<Grids, Blocks>>>();

}

I have left out code not relevant to my point.

When I go to compile this code I get:

The error occurs on the redeclaration of smemPoint in kernelTwo

I can get around this by just calling the shared memory in kernel two something else.

But I’m wondering about the garbage collection of shared memory. When is shared memory deallocated when using extern, when is it deallocated when just using the plain old shared.

Is there a way to deallocate extern shared memory in code?

Thanks

[font=“Courier New”]extern shared[/font] memory exists for as long as the block is running. If the variable name is visible outside of that scope, it might be caused by Nvidia (ab-) using a compiler mechanism that is meant for a different purpose. I’d expect you could just use a different name for the [font=“Courier New”]extern shared[/font] variable in each kernel.