Single __device__ variable in CUDA kernel


I’ve declared a group of single values (i.e device_ int val) in the .cu file. In my main function, I was able to assign their value just before the kernel launch. But, when I enter in the kernel, those variables have 0 as value.

Do you know how to assign value to some device single variable?



The proper way to assign a value to constant or device variables from host code is to use the cudaMemcpyToSymbol API. There are many many examples of this on the web already if you search for them.

In device code, you can read or write a device variable directly from you device code, just as if it were a “global” variable.

In device code, variables marked with constant can only be read, not written to.

Hello Robert,

Great ! It’s working now ! One more question, if I need to assign a single variable as kernel arguments, should I use the same procedure (cudaMemcpyToSymbol)?


I tried to fix values for a kernel arguments but they are still in zero when debugging the kernel launch. I followed the but it does not help.

Kernel arguments are passed in the normal fashion used by C++ to pass function arguments. CUDA ships with plenty of example apps, some of them very simple and basic, that demonstrate how kernels are invoked. Have you had a chance to look at them?

In general, a kernel invocation looks like a normal function call, but with launch configuration data added within triple angular brackets.

Hi njuffa,

Yes, it should be simple as the C++ standard. I just did it but when assigning 1.0 to a double argument, it remains 0 when debugging the kernel itself. It’s very strange for me as I invoke the kernel with a specific value.

Are you using a debug build when debugging? Can you show the exact command line used to invoke nvcc when you build your code?

What happens when you try it with the trivial example below? What does it print? What do you see in the debugger when you examine variable ‘a’?

#include <stdio.h>
#include <stdlib.h>

__global__ void kernel (double a)
    a = a + a;
    printf ("kernel: a=%23.16e\n", a);

int main (void)
    return EXIT_SUCCESS;

Hello njuffa,

Your kernel is accepting the entered argument value. Also, my own kernel accepts it if only you are running on CUDA Debugger (Legacy), so it’s the Next-Gen where it does not work. And I notice that this Next-Gen (Host/Device debugging) is using huge ressource and may lead to BSOD on my system.

I think that I will stay with Legacy now.

Thanks for your help !