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

Question

考虑以下代码:

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

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

我的问题是通过 constexpr values - 第一种函数( foo1 )。如果它在主机端 - 没问题。 constexpr变量有地址，编译器会照顾我并做一些合理的事情。

但是 - 设备端情况并非如此。如果我们编译:

constexpr const int c { 123 };

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

第一个函数将编译正常，但第二个函数将 fail to compile (神箭)。

我不能同时提供这两个函数，因为编译器将无法在它们之间做出决定(经常/总是)。而且我不想只传递值，因为我确实想避免大型 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 .