Accessing the same constant from multiple files?

Accelerated Computing CUDA CUDA Programming and Performance
1

Hi,
as the title says, i was trying to use the same constant definition from multiple files… after all it’s a costant, and it should be accessibile anywhere :rolleyes:

I tried defining all the constants in a .cuh file, but it did not work.
To my understainding, this way i only get identical but separate definitions in multiple modules…

There is a way to share constants between modules or i have to define (and update) a constant for each module?

Thanks External Image

2

I assume you mean const memory accessible from kernels, but this pattern will work for any other similar idea:

// constant_definitions.h

#pragma once

// define __const__ variables if you want

__const__ unsigned int constant_array[10] = {0, 1, 2, ...};

// or a preprocessor constant

#define PI 3.14

// or a const variable with global scope:

const unsigned int g_const_global_variable = 10;

The #pragma once directive is what keeps these identifiers from being multiply defined. It tells the preprocessor no matter how many times it sees this file, it should only be included once.

And then include this header from your kernel code:

// kernel.cu

#include "constant_definitions.h"

__global__ void foo(int *result)

{

  result[0] = PI * g_const_global_variable + constant_array[0];

}
3

I did the same thing but using an ifdef guard… anyway when i include the header in more than one .cu the values in the constants become undefined, and that was the problem.

After all even if you put all the constants in one single header, when it’s included they get duplicated in each single .cu or module.

So i was looking for a way to broadcast the same memory to all the symbols with the same name (but in different modules) with cudaMemcpyToSymbol…

4

I didn’t test this in CUDA, but I think this is what you need:

Just define your constant in a single .cu file and in the header file use the extern keyword. Example, if you declare:

const int x;

in the header file you should place:

extern const int x;

5

Thanks, i tried that way but it won’t work…

after examining better the CUDA manual i found this bits:

and

So, it is impossible to have a constant variable shared between two modules…

anyway there has to be a way to have two constants that point to the same memory… maybe.

6

If you dig further in the Programming Guide you will also find that the constant memory has implied static storage. This means that the variable defined at file scope

has internal linkage and thus only directly accessible in the module.

On host-side code in other modules than the one in which the constant memory was defined, you need to do run-time lookup of the variable. You do this by passing a string excactly naming the variable.

Example from my code:

in kernels.cu

[codebox]

constant int someTable[TABLE_SIZE]

…

global kernel (void) {

   // access someTable as usual

}

[/codebox]

otherFile.cpp

[codebox]

size_t arraySize;

if (cudaSuccess != cudaGetSymbolSize(&arraySize, “someTable”)) // note that we do not directly access the variable, but pass a const char *

	throw CudaInvalidSymbolError(" - Could not get someTable size");

[/codebox]