I am working on reorganizing some of my previous CUDA work. I remember seeing at one point that any kernel that uses a particular constant array must be located in the same file where the array is declared. For me, almost every kernel accesses my constant arrays, so I have one giant file with many different kernels.
Am I correct on this limitation? If not, how can I reference constant variables declared in one file from a different file?
Thanks!
You can put it in a .cuh (header file) and just have the .cu include it like a regular C header file.
eyal
There’s no linker, so doing what eyal suggested may fail miserably. Putting kernels in .cuh files and including those into a single compilation unit is probably the safest way to split a CUDA source base that accesses common global variables into multiple files.
No you can’t. The programming guide (B.2.5 in version 2.3) stats that constant variables have implied static storage, which means they are visible only within the scope of the same file.
Eh
I’ve just started to play with 3.0. A working code in 2.3 fails at runtime with this: “invalid device symbol[13]”
It seems that it has something to do with the constant I’ve defined in a .cuh - could this be because of what you’ve suggested here, Tim?
Can you please specify a more detailed example of what you’ve meant in your answer??
[edit] - I’ve moved the constant definition to the .cu file and I still get the same error. I put a check error code before the cudaMemcpyToSymbol function and
everything is fine. I then do:
int constHostIntParams[ PROJECT_INT_PARAMETERS_COUNT ];
constHostIntParams[ SYMETRICAL_TRACE_SELECTION ] = gGPUEngineSettings.m_useSymetricalTraceSelection;
....
// --> Check here for error - everything fine....
cudaMemcpyToSymbol( ProjectIntParams, &constHostIntParams[0], PROJECT_INT_PARAMETERS_COUNT * sizeof( int ), 0 );
// ---> Check here for error - invalid device symbol[13]
thanks
eyal