Как передать значения constexpr функциям на стороне устройства CUDA с использованием ссылок const?

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 - функции первого типа (foo1). Если это на стороне хоста - нет проблем. constexpr переменные имеют адреса, и компилятор позаботится обо мне и сделает что-то разумное.

Но - то же самое не относится к стороне устройства. Если мы скомпилируем:

constexpr const int c { 123 };

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

Первая функция будет работать нормально, но вторая не сможет скомпилировать (GodBolt).

Я не могу предоставить оба функции, так как компилятор не сможет выбирать между ними (часто / всегда). И я не хочу просто передавать значения, потому что я хочу избежать копий больших T; или потому что я обязан предоставить foo1() с помощью какого-то формального ограничения.

Что я тогда могу сделать?

Я также упомяну, что хочу написать один и тот же код на устройстве и на стороне хоста.

Answer 1

В настоящее время я использую следующий уродливый обходной путь:

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

, который похож на то, что предлагал @Artyer, но также будет работать в шаблонном коде, так как вам не нужно указывать тип c, например:

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

Это также имеет дополнительное преимущество в том, что вам не нужно знать тип c. Конечно, это также может быть сделано в коде на стороне хоста.

Однако - мне это очень не нравится! Читатели не поймут, зачем это нужно, и запутаются.

Answer 2

Вы можете явно скопировать его, чтобы он не занимал адрес чего-то несуществующего:

return foo(int{cci});

Таким образом, вместо него берется адрес нового значения r. Это действительно делает код другим на стороне устройства.

Вы также можете предоставить две перегрузки:

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.