如何定义CUDA设备常量，如C++常量/常量表达式

Question

在.cu文件中，我在全局范围内(即不在函数中(尝试了以下操作：

__device__ static const double cdInf = HUGE_VAL / 4;

并得到nvcc错误：

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

如果可能的话，如何在设备上定义C++常量/常量表达式？

注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 %lfn", 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，请参阅此处了解更多详细信息(感谢@harrism(。

Answer 2

要使显示的代码按预期编译和工作，需要在运行时初始化变量，而不是在编译时初始化。为此，向cudaMemcpyToSymbol添加一个主机端调用，类似于：

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

然而，对于单个值，将其作为内核参数传递似乎更明智。编译器将自动将参数存储在所有支持的体系结构上的常量内存中，并且有一个"免费"的常量缓存广播机制，这将使在运行时访问该值的成本可以忽略不计。

Answer 3

要初始化它，必须使用cudaMemcpyToSymbol。它不是编译时常数，而是存储在设备的常数存储器中，并且与全局存储器相比具有一些优势。来自CUDA博客：

对于半曲速的所有线程，从常量缓存中读取如下只要所有线程都读取相同的内容，就可以像从寄存器中读取一样快住址通过半曲速内的线程访问不同的地址是序列化的，因此成本随不同数量的线性扩展由半曲速内的所有线程读取的地址。

您不需要使用const，也不能使用它。它不是c++常量，因为您需要通过cudaMemcpyToSymbol修改它。因此，至少从c++的角度来看，它不是一个"真正的"常数。但它在设备内核中的行为就像一个常量，因为您只能通过只能从主机调用的cudaMemcpyToSymbol来修改它。