"conserve" SMEM data from one kernel to another

Accelerated Computing CUDA CUDA Programming and Performance
1

Hi all,

as far as I’ve understood, in the programming guide they say that data stored in SMEM (or in any GPU memory types) during the execution of a kernel are erased once the kernel execution ends.
But if one does not free() the memory, can we use data created by kernel_1 in kernel_2?
is that possible?

Thank you so much in advance

2

For global and constant memory, yes, for shared memory (block scope) and local memory (thread scope), no.

3

thank you avidday for you reply.

So I guess I can reuse data in kernel_2 by calling the variable names allocated by cudaMalloc() before kernel_1 execution. That is, I use the same variable names as in kernell_1. right? :-)

4 
float* deviceArray = 0;

cudaMalloc((void**)&deviceArray, size*sizeof(float));

kernel1<<<...>>>(deviceArray, size);

kernel2<<<...>>>(deviceArray, size);

.

.

kernelN<<<...>>>(5.0, a, b, deviceArray, size);

cudaFree(deviceArray);

Whatever changes to deviceArray a kernel makes, subsequent kernels will see. It works like standard C functions, really. Local (“stack”) variables have the lifetime of a function and don’t propagate beyond the closing brace, if you want to externalize data from a function you need to pass it either as a return value (impossible in case of CUDA kernels) or back through the parameter list (“by reference”).

5

If you pass them by argument to both kernels, then yes. If you copy the addresses of the malloc() memory onto global or constant memory pointers before running the first kernel, then also yes. Otherwise no.

6

I remember posting something a few months ago about my annoyance that i could call the same kernel several times and the values in smem just kept adding up (the kernel didnt begin with setting the smem to any value but simply reused whatever was lying there).

This might depend on the card but the values stored in smem dont seem to be erased after a kernel has finished executing.

But trying to use this is unsupported and will probably be very problematic!

7

thanks for your answers!!

it is a good news I can reuse data in GMEM and constant mem :-)
…and it is better not to count on what smem keeps stored (if any).

Thank you!

8

…which is perfectly acceptable behaviour for undefined variables on the CPU or GPU…

9

Yes, YDD, i accept this :)