How do I pass constexpr values to CUDA device-side functions taking const references?

user_1593077 0
Viewed 223

Consider the following code:

template <typename T> __host__ __device__ int foo1(const T& x);
template <typename T> __host__ __device__ int foo2(T x);

These two functions correspond to two common ways to pass "in"-parameters rather than "out" or "in/out" parameters. The second one is simpler, in that no references or addresses are involved; but the first one ensures no copying of more complex types, so it is often preferred.

My problem is with passing constexpr values - to the first kind of function (foo1). If it's on the host side - no problem. constexpr variables have addresses, and the compiler will take care of me and do something reasonable.

But - the same is not true for the device side. If we compile:

constexpr const int c { 123 };

__host__   int bar() { return foo1(c); }
__device__ int baz() { return foo1(c); }

The first function will compiler fine, but the second one will fail to compile (GodBolt).

I can't provide both functions, since the compiler won't be able to decide between them (often/always). And I don't want to just pass values, because I do want to avoid copies of large T's; or because I'm required to provide foo1() by some formal constraint.

What can I do, then?

I'll also mention I'd want to be able to write the same code on both the device and the host side.

const-reference constexpr cuda c++11 c++
2 Answers

You can explicitly copy it so it doesn't take the address of something that doesn't exist:

return foo(int{cci});

So the address of the new rvalue is taken instead. This does make the code different on the device side though.

You could also provide two overloads:

template <typename T> __host__ __device__ std::enable_if_t<!std::is_trivial_v<T>, int> foo(const T& x);
template <typename T> __host__ __device__ std::enable_if_t<std::is_trivial_v<T>, int> foo(T x);

so the copy is done for you for trivial types like int.

user_57546560

Currently, I use the following ugly workaround:

__device__ int baz() { return foo1(decltype(c){c}); }

which is similar to what @Artyer suggested, but will also work in templated code, since you don't need to specify the type of c, e.g.:

template <typename T>
__device__ int quux() { return function_taking_const_ref(decltype(c){c}); }

It also has the added benefit of not having to know the type of c. This can also be done in host-side code of course.

However - I really dislike it! Readers will not understand why it's needed, and will get somewhat confused.

user_15930770
Related