Cannot define CUDA constants using other variables

Accelerated Computing CUDA CUDA Programming and Performance
1

Hello! I have some strange problems with definition of constants in CUDA. I posted this in another thread recently, but now the problem is more isolated so I thought a new thread would be best.

I have a single main.cu file, and this compiles fine:

__device__ __constant__ float d_some_constant = 1.0;

__global__ void kernel(float* time)

{

        time[threadIdx.x] += d_some_constant;

}

int main()

{

	return 0;

}

However, the following does NOT work:

const float some_constant = 1.0;

__device__ __constant__ float d_some_constant = some_constant;

__global__ void kernel(float* time)

{

        time[threadIdx.x] += d_some_constant;

}

int main()

{

	return 0;

}

The build log looks like this:

Build Log

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

Command Lines

Creating temporary file “d:\Kod\Projects\cuda_template\cuda_template\x64\Debug\BAT00000E19685500.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\BAT00000E19685500.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_0000074c_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)

Any ideas? I have friends that can compile the same code without any problems, and I have not found any other strange behaviour in my CUDA compiling yet!

2

I think the problem is : some_constant is a const value in host memory, but the d_some_constant is on device. So , it might be not proper if you directly use them like this.

3

If so, is there any other way to do this? I want to aviod declare constants that both the device and host use at two places. Also, the code have compiled just fine for other people.

4

I am not sure that my explanation is right. So, you can declare constants in both sides firstly to check whether what I said is right. If so, the try to find fine solution.

5

I’m currently trying something like this, but I’m a bit unsure on how to handle the pointers.

float some_constant = 1.0;

__device__ __constant__ float d_some_constant;

cudaMemcpyToSymbol(d_some_constant, &some_constant, sizeof(some_constant));
6

Alright, I found that this works:

float some_constant = 1.0;

__device__ __constant__ float d_some_constant;

__global__ void kernel(float* time)

{

        time[threadIdx.x] += d_some_constant;

}

void copyData()

{

    cudaMemcpyToSymbol("d_some_constant", &some_constant, sizeof(float));

}

int main()

{

	copyData();

	return 0;

}
7

Yes, I just want to tell your this function