Newbie question: How to allocate global constants

Hi!

I’m just beginning to learn CUDA programming and have run in to the following problem. I have some global constants that the whole program shares. Part of some calculations are done with CUDA kernels, and I use OpenGL (GLUT) to visualize the results later. CUDA just returns values as this point.

Since I want the constants that CUDA use to be the same that the rest of the program use, I tried the following:

In a file called “constants.hpp”:

const float TIMESTEP = 0.001;

Then, in another file called “functions.cu” I tried this:

#include "constants.hpp"

...

...

__device__ __constant__ float D_TIMESTEP = TIMESTEP;

...

...

But the compiler returns the following error message:

“1>kernel_functions.cu(11): error: can’t generate code for non empty constructors or destructors on device”

What am I doing wrong, and is this a good idea to start with?

Thanks

Victor

you don’t need the device specifier to declare a constant in device memory. By using device it looks like nvcc is trying a treat your constant as a class or function declaration and choking on it.

I tried to remove the constant or the device specifiers, but the same error occurs anyway. I just tried

__device__ __constant__ D_TIMESTEP = 0.001;

instead of

__device__ __constant__ D_TIMESTEP = TIMESTEP;

and it worked fine! But it takes away the point, which is having the variables set at only one point in the code!

Something like this works for me with CUDA 3.2 on linux :

#include "constant.h"

__constant__ float d_timestep = timestep;

__global__ void kernel(float *time)

{

	time[threadIdx.x] += d_timestep;

}

with constant.h containing

const float timestep = 1.0e-3f;

If you can’t compile that successfully, then you might have found a compiler bug.

Thanks, I tried that and have exactly the following:

constants.h

const float timestep = 0.001;

main.cu

#include "constants.h"

__constant__ float d_timestep = 0.001;

__global__ void kernel(float *time)

{

        time[threadIdx.x] += d_timestep;

}

int main()

{

	return 0;

}

No other code. This compiles just fine, but when I change “0.001” to “timestep” the compiler gives me this:

1>------ Build started: Project: cuda_template, Configuration: Debug x64 ------

1>Compiling...

1>main.cu

1>main.cu(3): error: can't generate code for non empty constructors or destructors on device

1>1 error detected in the compilation of "C:/Users/Victor/AppData/Local/Temp/tmpxft_00001994_00000000-6_main.cpp1.ii".

1>Project : error PRJ0019: A tool returned an error code from "Compiling..."

1>Build log was saved at "file://d:\Kod\Projects\cuda_template\cuda_template\x64\Debug\BuildLog.htm"

1>cuda_template - 2 error(s), 0 warning(s)

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

This is the BuildLog.html file:

Build Log

Build started: Project: cuda_template, Configuration: Debug|x64

Command Lines

Creating temporary file “d:\Kod\Projects\cuda_template\cuda_template\x64\Debug\BAT00004311805984.bat” with contents

[

@echo off

nvcc.exe -ccbin “C:\Program Files (x86)\Microsoft Visual Studio 8\VC\bin” -c -Xcompiler "/EHsc /W3 /nologo /Wp64 /O2 /Zi /MT " -o Debug\main.obj main.cu

if errorlevel 1 goto VCReportError

goto VCEnd

:VCReportError

echo Project : error PRJ0019: A tool returned an error code from “Compiling…”

exit 1

:VCEnd

]

Creating command line “d:\Kod\Projects\cuda_template\cuda_template\x64\Debug\BAT00004311805984.bat”

Output Window

Compiling…

main.cu

main.cu(3): error: can’t generate code for non empty constructors or destructors on device

1 error detected in the compilation of “C:/Users/Victor/AppData/Local/Temp/tmpxft_00001994_00000000-6_main.cpp1.ii”.

Project : error PRJ0019: A tool returned an error code from “Compiling…”

Results

Build log was saved at “file://d:\Kod\Projects\cuda_template\cuda_template\x64\Debug\BuildLog.htm”

cuda_template - 2 error(s), 0 warning(s)

Try this - it has always worked for me.

constants.h

const float timestep = 0.001;

main.cu

#include "constants.h"

__global__ void kernel(float *time)

{

        time[threadIdx.x] += timestep;

}

The compiler is smart enough to take a const variable and put it into the appropriate memory space without you having to tell it so. I’ve always found that initializing constant memory requires a cudaMemcpyToSymbol which is just a pain to manage for numerous small constants.

Tryin that gives me:

1>main.cu(6): error: identifier “timestep” is undefined.

Seems that if I change “float” to “int” in the original and remove the decimals, it compiles just fine.

const int timestep = 1;
#include "constants.h"

__device__ __constant__ int d_timestep = timestep;

__global__ void kernel(float *time)

{

        time[threadIdx.x] += d_timestep;

}

int main()

{

	return 0;

}

Weird. It works fine for me (on Mac OS X) and I’ve been doing this for years.

I tried with an int, and then it works fine in the same way as in my test two posts up. When I change it back to float I get the compiler errors again.

I’ve also encountered another strange problem, maybe it can be good to post it as well.

When I first tried to use several blocks, using gridDim.x did not work. It compiled, but my code did not do any work on any variables except those in the first block. However, when I hard-coded the number of blocks as an integer in the kernels, it worked fine (even if I launched the kernel with the exact same integer as grid dimension and using gridDim in the kernel). I will see if I can force this error again and post it.

Could my nvcc settings be wrong in any way?