如何像 C++ const/constexpr 一样定义 CUDA 设备常量？

Question

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

__device__ static const double cdInf = HUGE_VAL / 4;

并出现 nvcc 错误：

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

如果可能的话，如何在设备上定义 C++ const/constexpr？

注意 1：#define 是不可能的，不仅出于审美原因，而且因为在实践中表达式更复杂并且涉及内部数据类型，而不仅仅是 double。所以每次在每个 CUDA 线程中调用构造函数会太昂贵。

注意 2：我怀疑 __constant__ 的性能，因为它不是编译时常量，而更像是用 cudaMemcpyToSymbol 编写的变量。

Answer 1

要初始化它，您必须使用 cudaMemcpyToSymbol。它不是编译时常量，而是存储在设备的常量内存中，与全局内存相比具有一些优势。来自 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++ 的角度来看，它不是 "real" 常量。但它在设备内核中的行为就像一个常量，因为你只能通过 cudaMemcpyToSymbol 修改它，而它只能从主机调用。

Answer 2

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

__device__ double cdInf;

// ...

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

但是，对于单个值，将其作为内核参数传递似乎更为明智。编译器会自动将参数存储在所有支持的架构上的常量内存中，并且有一个 "free" 常量缓存广播机制，可以忽略运行时访问值的成本。

Answer 3

使用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）。

如何像 C++ const/constexpr 一样定义 CUDA 设备常量？

How to define CUDA device constant like a C++ const/constexpr?

c++

cuda

constants

compile-time-constant