Use constexpr in kernel

Accelerated Computing CUDA CUDA Programming and Performance
1

My kernel has a template parameter that can be either unsigned int or unsigned long long int.
I would like to create masks that match this type, at compile time… Here is my specialized template class
for doing this. Unfortunately, I get error

error: calling a constexpr __host__ function("gen") from a __global__ function("foo") is not allowed.

template<typename T> struct MaskGen {
};
template<> struct MaskGen<uint32_t> {
    consteval uint32_t gen(uint8_t tmpl) { return uint32_t(0x01010101) * tmpl; }
};
template<> struct MaskGen<uint64_t> {
    consteval uint64_t gen(uint8_t tmpl) { return uint64_t(0x0101010101010101) * tmpl; }
};

template <typename T> __global__ void foo(void) {
    	constexpr auto ALL_EIGHTS = MaskGen<T>().gen(0x88);

}

Note: I found a simple old-school way of doing this:

#define MASK(b) ((T(-1) / 0xFF) * (b))

where T is the type.

2

Host functions cannot be called in a kernel. You could make gen __host__ __device__

2 Likes
3

You could also compile with --expt-relaxed-constexpr (if you use NVCC). That switch lets you use host-side constexpr (probably also consteval) functions in device-side code.

1 Like
4

Thanks, all. In the end, I used the simple macro, and it works fine.