CUDA constant variables not loading properly in kernels

Accelerated Computing CUDA CUDA Programming and Performance
1

Hi, i’ve run into a very very weird issue.

First of all my app is quite large and split into quite some files. I’m mentioning this in case this could be related to my prob. All my kernels are in one place cu file, all my constant variables are located in a .h file which is included in my main and my kernel cu files, and of course one main cu file that invokes the kernels. I’m importing the kernels to the main app using an .cuh file.

Constant variables are just defined in the header file using the following code:

//Define constants for CUDA
__device__ __constant__ double d_invertn;
__device__ __constant__ double d_dt;
__device__ __constant__ double d_epsilon2;
__device__ __constant__ double d_pi;
__device__ __constant__ double d_Deff;
__device__ __constant__ double d_curv;
__device__ __constant__ double d_phis;

They are afterwards initialized in the main code file and are passed on the device using cudaMemcpyToSymbol().

The problem is that for some reason, all the constant variables within the kernels are 0.0…
I’ve even tried to get the constant variables back from the device using cudaMemcpyFromSymbol() function and it returns the correct value.

Do i need to define or setup the constants otherwise?

I have no idea where to start looking, and i would really appreciate some advice

thanks in advance

EDIT: I tried to assemble a sample program, putting everything into one file and this worked fine. Could this be happening because of my multifile approach?
I’m running on a GTX750 btw.

2

It is not entirely clear as to what exactly you are doing. Could you post a minimal, buildable, runnable code that reproduces the problem?

constant data has file scope (i.e. “implied static storage”, see Programming Guide, section E.2.3.1). So the calls to cudaMemcpyToSymbol() have to reside in the same file where the constant data is defined.

3

As i said i did try to do a minimal example but i moved everything on one file and it worked smoothly. I’ll try to test this out by keeping the multiple files scheme and post it asap.

Btw i think what you said sort of helps. I am heading the same direction right now. I moved the constant variables declaration in the kernels cu file, and i wrote a small wrapper function to handle all the necessary cudaMemcpyToSymbol calls. But this doesn’t seem to work as well for some reason…

I’ve declared the following function in the kernel’s file:

void setConstants(double *dt, double *pi, double *epsilon2,\
                           double *Deff, double *curv, double *phis){
    
    cudaError_t err;
    err = cudaMemcpyToSymbol(d_dt, dt, sizeof(double));
    checkCudaError(err);
    cudaMemcpyToSymbol(d_pi, pi, sizeof(double));
    cudaMemcpyToSymbol(d_epsilon2, epsilon2, sizeof(double));
    cudaMemcpyToSymbol(d_Deff, Deff, sizeof(double));
    cudaMemcpyToSymbol(d_curv, curv, sizeof(double));
    cudaMemcpyToSymbol(d_phis, phis, sizeof(double));
}

i’m calling this from the main code using:

setConstants(&dt, &pi, &epsilon2, &Deff, &curv, &phis);

but again when i’m trying to read the constants they are still zero…