The use of host-side consexpr constants in device code is limited to
- scalars (other than
long double)
- scalar elements of aggregates used inside
constexpr functions
In particular, it's not possible to use constexpr scalars by pointer or reference (e.g. std::min() takes arguments by reference), or pass constexpr arrays as pointers, or access elements of constexpr arrays outside of constexpr functions.
The workaround that we found, is to write portable code that declare two copies of the constants, one on the host and one on the device:
#ifdef __CUDA_ARCH__
__device__
#endif
constexpr uint32_t values[N] = { ... };We should probably abstract it with something like
#ifdef __CUDA_ARCH__
#define HOST_DEVICE_CONSTANT __device__ constexpr
#else
#define HOST_DEVICE_CONSTANT constexpr
#endif
to be used as
HOST_DEVICE_CONSTANT uint32_t values[N] = { ... };Suggestions for a better name are welcome :-)
The use of host-side
consexprconstants in device code is limited tolong double)constexprfunctionsIn particular, it's not possible to use
constexprscalars by pointer or reference (e.g.std::min()takes arguments by reference), or passconstexprarrays as pointers, or access elements ofconstexprarrays outside ofconstexprfunctions.The workaround that we found, is to write portable code that declare two copies of the constants, one on the host and one on the device:
We should probably abstract it with something like
to be used as
HOST_DEVICE_CONSTANT uint32_t values[N] = { ... };Suggestions for a better name are welcome :-)