How to best make use of the same constants in both host and device code?

Question

Suppose I have some global constant data I use in host-side code:

const float my_array[20] = { 45.146, 54.633, 74.669, 12.734, 74.240, 100.524 };

(Note: I've kept them C-ish, no constexpr here.)

I now want to also use these in device-side code. I can't simply start using them: They are not directly accessible from the device, and trying to use them gives:

error: identifier "my_array" is undefined in device code 

What is, or what are, the idiomatic way(s) to make such constants usable on both the host and the device?

Answer 1

This approach was suggested by Mark Harris in an answer in 2012:

#define MY_ARRAY_VALUES 45.146, 54.633, 74.669, 12.734, 74.240, 100.524
__constant__ float device_side_my_array[2] = { MY_ARRAY_VALUES };
const        float host_side_my_array[2]   = { MY_ARRAY_VALUES };
#undef MY_ARRAY_VALUES

__device__ __host__ float my_array(size_t i) {
#ifdef __CUDA_ARCH__
    return device_side_my_array[i];
#else
    return host_side_my_array[i];
#endif
}

But this has some drawbacks:

• Not actually using the same constants, just constants with the same value.
• Duplication of data.
• Takes up constant memory, which is a rather limited resource.
• Seems a bit verbose (although maybe other options are even more so).

I wonder if this is what most people use in practice.

Note:

• In C++ one might use the same name, but in different sub-namespaces within the detail:: namespace.
• Doesn't use cudaMemcpyToSymbol().