CUDA and shared variables among different global functions

Question

Just a few doubts about CUDA. Perhaps they may seem silly questions; I apologize for it.

If I declare a variable on the GPU (e.g., an array alpha with N elements, cudaMalloc((void**)&alpha, N * sizeof(double))) and allocate its value in a global function without freeing its memory, this variable should be available for other successive global functions, right?

Besides, is it possible (or advisable) to compute a scalar variable on the GPU and make it shared among several global functions on the GPU, or is it better to pass it every time as an argument from the CPU?

Thank you for your attention.

Thank you for all your answers. I am not able to compute all my project on the GPU as I have a recursive loop. Hence, my aim would be to compute some constant arrays on the GPU outside this loop, without deallocating the memory used, and some constant scalars on the CPU. Then, inside the recursive loop, I will use global functions as long as I can. — Pippo
– Pippo, Commented Dec 29, 2012 at 20:35

agrippa · Accepted Answer · 2012-12-29 20:06:54Z

1

Yes, if you write values into allocated global memory those values are persistent until you free that memory, even across kernel invocations.

As for accessing scalar variables (which are constants), the better approach would be to pass it as a parameter to the global kernel launch rather than putting it in global memory and reading it from there. Global memory accesses are expensive, and this avoids loading that scalar from global memory every time you need to read it.

answered Dec 29, 2012 at 20:06

agrippa

3795 silver badges19 bronze badges

Sign up to request clarification or add additional context in comments.

1 Comment

Roger Dahl Over a year ago

Kernel parameters are also stored in global memory, but they get read through the constant cache. If the parameter is a pointer to an array, the array itself can also be read through the constant cache by using the const qualifier in the kernel parameter list.

saeedn · Accepted Answer · 2012-12-29 20:09:32Z

1

If I get your question right, you're allocating an array, fill the array in a global kernel function on GPU and then process the values of that array in another kernel call.

As long as you do not free the allocated array, its values remain on global memory. So you can do that and process the same array without copying it back to CPU. Dividing jobs between multiple kernel call may come handy when you have execution time limit or one of kernel functions are in a library. But in most other cases, it seems better do all jobs in one function call.

It also seems better to pass the scalar value as an argument, because reading it from global memory has a very higher overhead.

answered Dec 29, 2012 at 20:09

saeedn

3,3881 gold badge24 silver badges20 bronze badges

Comments

Roger Dahl · Accepted Answer · 2012-12-29 20:15:12Z

If I declare a variable on the GPU (e.g., an array alpha with N elements, cudaMalloc((void**)&alpha, N * sizeof(double))) and allocate its value in a global function without freeing its memory, this variable should be available for other successive global functions, right?

You cannot call cudaMalloc() from a global function (kernel). It's a host function. You can use malloc() and new in kernels, but that can be inefficient.

You can use the same array in multiple kernels, for instance, you might perform multiple calculation steps with different kernels.

Besides, is it possible (or advisable) to compute a scalar variable on the GPU and make it shared among several global functions on the GPU, or is it better to pass it every time as an argument from the CPU?

If you pass a constant as an argument to a kernel, it is very efficiently shared among all the threads. So, it will typically be much more efficient to calculate parameters on the CPU and pass them to the kernel.

If there is a lot of parallel calculation that goes into creating the scalar, then it would be best to calculate it with a separate kernel, pass it back to the host and then pass it to the next kernel as an argument. Alternatives would just increase the complexity of your code without any performance benefit.

If there is little computation required for the scalar, then it doesn't make sense to calculate it with a kernel. Also, remember that there is no guarantee as to which order the blocks are launched in the kernel, so you would have to create a separate code path in the kernel to set up the scalar and then do costly thread index testing and synchronization to calculate the scalar and make it available to all the threads.

Thank you. About the use of cudaMalloc(), I didn't explain it well in my question, but I meant what you say (i.e., allocate the memory space on the GPU in the main function, and not in a global one).

Collectives™ on Stack Overflow

CUDA and shared variables among different global functions

3 Answers 3

1 Comment

Comments

1 Comment

Your Answer

Linked

Hot Network Questions

Collectives™ on Stack Overflow

3 Answers 3

1 Comment

Comments

1 Comment

Your Answer

Sign up or log in

Post as a guest

Linked

Related