Memory Deallocation

Is it possible to deallocate shared memory on cuda kernels?

Thanks.

why would you need to explicitly deallocate it? just reuse it using pointer trickery or something once you’re done with the initial array/variables.

Could you please give an example?

Thanks.

No. You can’t allocate inside the kernel, so you can’t deallocate inside the kernel. You can ‘allocate’ when calling a kernel, but that gets automatically freed when the kernel returns.

For pointer trickery consult C/C++ books. Basically you get yourself a static chunk of memory (any way you can… simple array or allocated when the kernel is called) and then play with it, by keeping track of how much and what you’re using and passing pointers around.

__shared__ float* array1;

/* use array1 here */

__syncthreads();

// now that you're done with array 1, we can use a differently typed pointer to the same memory location

int* array2 = (int*) array1;

/* use array2, probably reinitializing everything first */

Just make sure to allocate the proper amount of shared memory in the kernel launch parameters if you’re going to dynamically allocate shmem like that.