我如何将constexpr值传递给采用const引用的CUDA设备端函数？

Question

请考虑以下代码：

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

这两个函数对应于传递“输入”参数而不是“输出”或“输入/输出”参数的两种常用方法。第二个比较简单，因为不涉及引用或地址。但是第一个确保不会复制更复杂的类型，因此通常是首选。

我的问题是将constexpr值传递给第一类函数（foo1）。如果在主机端-没问题。 constexpr变量具有地址，并且编译器将照顾我并做一些合理的事情。

但是-对于设备端而言并非如此。如果我们编译：

constexpr const int c { 123 };

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

第一个函数可以编译，但是第二个函数可以fail to compile（GodBolt）。

我不能同时提供这两个功能，因为编译器无法（通常/总是）在它们之间做出决定。而且我不想只传递值，因为我想避免复制大的T。或因为某些形式上的限制我需要提供foo1()。

那我该怎么办？

我还将提到我希望能够在设备和主机端都编写相同的代码。

Answer 1

您可以显式复制它，这样它就不会使用不存在的内容的地址：

return foo(int{cci});

因此将采用新右值的地址。但这确实使代码在设备端有所不同。

您还可以提供两个重载：

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);

因此，对于琐碎的类型，例如int，将为您完成复制。

Answer 2

当前，我使用以下丑陋的解决方法：

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

与@Artyer建议的类似，但也将在模板代码中工作，因为您无需指定c的类型，例如：

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`.