I wonder whether we can free shared memory within a kernel…
Given that you cannot allocate shared memory within a kernel, what would this be used for?
Good point …
i have a 3-pass algorithm and each pass includes 4 shared mem buffers.
After each pass i syncthreads and load data from global memory to the shared memory buffers.
I have a feeling that there is a performance hit when i rewrite on the shared memory, probably due to optimizations that the compiler does in the preceding passes
(ie. unloading shared mem in registers etc.)
So i would like to try sth like
1st pass
declare shared mem
load from global to shared
do sth with shared data
delete shared
syncthreads
2nd pass
declare shared mem
load from global to shared
do sth with shared data
delete shared
syncthreads
3d pass
declare shared mem
load from global to shared
do sth with shared data
delete shared
syncthreads
__global__ void kernel(...)
{
{
// 1st pass
__shared__ ...
...
__syncthreads();
}
{
// 2nd pass
__shared__ ...
...
__syncthreads();
}
{
// 3rd pass
__shared__ ...
...
__syncthreads();
}
}
IIRC the compiler will reuse the same shared memory for the three declarations. If it doesn’t, or if you have a more complicated kernel structure which cannot be separated into consecutive passes, you can also use a union to overlap any variables in (shared) memory.