How to pass structures into CUDA device?

Question

I've been stuck on this for a while. When I pass my structures into CUDA via kernel parameters, they contain no data and everything is undefined inside of them.

Out in host global space

struct matl1
{
    static const double cond;
};

const double matl1::cond = 420.5;

Then inside of main()

matl1 * h_matl1 = (matl1*)malloc(sizeof(matl1));
matl1 * d_matl1;
cudaMalloc((void**)&d_matl1, sizeof(matl1));
cudaMemcpy(d_matl1, h_matl1, sizeof(matl1), cudaMemcpyHostToDevice);
kernel<<<1,1>>>(d_matl1,...);

Then inside of kernel()

__global__ void kernel(matl1* d_matl1,...)
{
double cond = d_matl1->cond;
}

And I get the following error:

error : identifier "matl1::cond" is undefined in device code

As a quick test, if I do the following on the host in main()

cout << h_matl1->cond << endl;

It shows me the correct output of 420.5. I am not sure why this isn't making it into the device.

Here's the rest of my output

uild started: Project: test_struct, Configuration: Debug Win32 ------

Compiling CUDA source file kernel.cu...

C:\Users\User\Documents\Visual Studio 2012\Projects\test_struct\test_struct>"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v5.5\bin\nvcc.exe" -gencode=arch=compute_35,code=\"sm_35,compute_35\" --use-local-env --cl-version 2012 -ccbin "C:\Program Files (x86)\Microsoft Visual Studio 11.0\VC\bin" -I"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v5.5\include" -I"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v5.5\include" -G --keep-dir Debug -maxrregcount=0 --machine 32 --compile -cudart static -g -DWIN32 -D_DEBUG -D_CONSOLE -D_MBCS -Xcompiler "/EHsc /W3 /nologo /Od /Zi /RTC1 /MDd " -o Debug\kernel.cu.obj "C:\Users\User\Documents\Visual Studio 2012\Projects\test_struct\test_struct\kernel.cu" 1>C:/Users/User/Documents/Visual Studio 2012/Projects/test_struct/test_struct/kernel.cu(15): error : identifier "matl1::cond" is undefined in device code

C:\Program Files (x86)\MSBuild\Microsoft.Cpp\v4.0\V110\BuildCustomizations\CUDA 5.5.targets(592,9): error MSB3721: The command ""C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v5.5\bin\nvcc.exe" -gencode=arch=compute_35,code=\"sm_35,compute_35\" --use-local-env --cl-version 2012 -ccbin "C:\Program Files (x86)\Microsoft Visual Studio 11.0\VC\bin" -I"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v5.5\include" -I"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v5.5\include" -G --keep-dir Debug -maxrregcount=0 --machine 32 --compile -cudart static -g -DWIN32 -D_DEBUG -D_CONSOLE -D_MBCS -Xcompiler "/EHsc /W3 /nologo /Od /Zi /RTC1 /MDd " -o Debug\kernel.cu.obj "C:\Users\User\Documents\Visual Studio 2012\Projects\test_struct\test_struct\kernel.cu"" exited with code 2.

========== Build: 0 succeeded, 1 failed, 0 up-to-date, 0 skipped ==========

It seens to work for me. See here I'm doing this on linux with CUDA 5.5. You'll need to provide more specifics and an actual complete example, just as I have, so we can see all the code and output of the compiler. Identify the platform (linux, windows, etc.), which version of CUDA, and provide the complete code, and the complete output from the compile command, as well as the complete compile command. — Robert Crovella

Robert Crovella Robert Crovella · Accepted Answer · 2013-11-21T03:42:55

I was able to reproduce your error if I do this:

struct matl1
{
    static const double cond;
};



__global__ void kernel(matl1* d_matl1)
{
double cond = d_matl1->cond;
printf("cond = %lf\n", cond);
}

const double matl1::cond = 420.5;

But not if I do this:

struct matl1
{
    static const double cond;
};

const double matl1::cond = 420.5;

__global__ void kernel(matl1* d_matl1)
{
double cond = d_matl1->cond;
printf("cond = %lf\n", cond);
}

You need to define the constant initializer before the kernel definition. Here is a complete example that works for me.

How to pass structures into CUDA device?

1 Answers