Как определить константу устройства CUDA, как C++ const/constexpr?

Question

В .CU файл я попытался следующие действия в глобальном масштабе (т.е. не в функции):

__device__ static const double cdInf = HUGE_VAL/4; 

И получил NVCC ошибку:

error : dynamic initialization is not supported for __device__, __constant__ and __shared__ variables. 

Как определить C++ сопзЬ/constexpr на устройстве, если это возможно?

ПРИМЕЧАНИЕ 1: #define не может быть рассмотрено не только по эстетическим соображениям, но и потому, что на практике выражение более сложное и включает в себя внутренний тип данных, а не только двойной. Поэтому вызов конструктора каждый раз в каждом потоке CUDA будет слишком дорогостоящим.

ПРИМЕЧАНИЕ 2. Я сомневаюсь в производительности __constant__, потому что это не константа времени компиляции, а скорее переменная, написанная с cudaMemcpyToSymbol.

Answer 1

Используйте constexpr __device__ функцию:

#include <stdio.h> 
__device__ constexpr double cdInf() { return HUGE_VAL/4; } 
__global__ void print_cdinf() { printf("in kernel, cdInf() is %lf\n", cdInf()); } 
int main() { print_cdinf<<<1, 1>>>(); return 0; } 

PTX должно быть что-то вроде:

.visible .entry print_cdinf()(

) 
{ 
     .reg .b64  %SP; 
     .reg .b64  %SPL; 
     .reg .b32  %r<2>; 
     .reg .b64  %rd<7>; 


     mov.u64   %rd6, __local_depot0; 
     cvta.local.u64 %SP, %rd6; 
     add.u64   %rd1, %SP, 0; 
     cvta.to.local.u64  %rd2, %rd1; 
     mov.u64   %rd3, 9218868437227405312; 
     st.local.u64 [%rd2], %rd3; 
     mov.u64   %rd4, $str; 
     cvta.global.u64   %rd5, %rd4; 
     // Callseq Start 0 
     { 
     .reg .b32 temp_param_reg; 
     // <end>} 
     .param .b64 param0; 
     st.param.b64 [param0+0], %rd5; 
     .param .b64 param1; 
     st.param.b64 [param1+0], %rd1; 
     .param .b32 retval0; 
     call.uni (retval0), 
     vprintf, 
     (
     param0, 
     param1 
     ); 
     ld.param.b32 %r1, [retval0+0]; 

     //{ 
     }// Callseq End 0 
     ret; 
} 

При отсутствии кода для функции constexpr. Вы также можете использовать функцию constexpr __host__, но это экспериментально в CUDA 7: использовать параметры командной строки nvcc, похоже, --expt-relaxed-constexpr и см. here для получения дополнительной информации (спасибо @harrism).

Answer 2

Для его инициализации вы должны использовать cudaMemcpyToSymbol. Это не постоянная времени компиляции, а хранится в постоянной памяти устройства и имеет некоторые преимущества перед глобальной памятью. С Blogspot CUDA:

For all threads of a half warp, reading from the constant cache is as fast as reading from a register as long as all threads read the same address. Accesses to different addresses by threads within a half warp are serialized, so cost scales linearly with the number of different addresses read by all threads within a half warp.

Вам не нужно использовать const, и вы не можете использовать его. Это не константа C++, так как вам нужно изменить ее до cudaMemcpyToSymbol. Таким образом, это не «реальная» константа, по крайней мере, с точки зрения C++. Но он ведет себя как константа внутри ядер устройства, потому что вы можете изменить его только через cudaMemcpyToSymbol, который может быть вызван только с хоста.

Answer 3

Чтобы сделать код, который вы показали, компилируется и работает должным образом, вам нужно инициализировать переменную во время выполнения, а не время компиляции. Чтобы сделать это, добавьте боковой вызов хоста cudaMemcpyToSymbol, что-то вроде:

__device__ double cdInf; 

// ... 

double val = HUGE_VAL/4 
cudaMemcpyToSymbol(cdInf, &val, sizeof(double)); 

Однако, для одного значения, передав его в качестве аргумента ядра будет казаться гораздо более разумным. Компилятор автоматически сохранит аргумент в постоянной памяти на всех поддерживаемых архитектурах, и существует «бесплатный» механизм широковещания с постоянным кешем, который должен сделать стоимость доступа к значению во время выполнения незначительным.