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

Question

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

6

在一个 .cu 文件中，我已经尝试在全局范围内（即不在函数中）执行以下操作：

__device__ static const double cdInf = HUGE_VAL / 4;

出现了nvcc错误：

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

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

注意1：由于美观原因以及实际情况中表达式更加复杂并涉及内部数据类型（不仅仅是double），所以#define不可行。因此，在每个CUDA线程中每次调用构造函数将会太昂贵。

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

- Serge Rogatch

可能是使用CUDA常量的重复问题。 - m.s.

@m.s.，这个问题将范围限制在__constant__和#define上。难道没有一种方法可以像我的代码片段所示那样在设备上使用C++常量吗？ - Serge Rogatch

@SergeRogatch：你实际上需要在设备上以可访问的变量形式，在运行时具有地址和其他一些东西吗？因为除了#DEFINE和const之外，还有其他替代方案。如果你的回答是“不需要”，那么我会在回答中提供几个替代方案... - einpoklum

@einpoklum，我只需要在设备上使用它，但不需要它的地址。如果可以将其编译时常量化，那么在设备上它就会变成一个立即值，我会很高兴的。 - Serge Rogatch

3个回答

2

为了让您展示的代码能够编译并按预期工作，您需要在运行时而非编译时初始化变量。为此，请添加一个主机端调用 cudaMemcpyToSymbol，类似于以下内容：

__device__ double cdInf;

// ...

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

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

- talonmies

实际上，我已经在使用这个了。除了性能问题之外，还存在可维护性问题，即常量必须在定义它的地方初始化。因此，在C++中，我需要3个位置：1）在类中声明2）在.cpp文件中定义3）cudaMemcpyToSymbol初始化。前两个是不可避免的（除了int），但我会尝试摆脱第三个。 - Serge Rogatch

1

如果您想要一个纯编译时常量，并且正在使用g++，则可以只使用“const”而不需要任何cuda说明符，并在设备代码中将该常量作为立即常量值访问。 - talonmies

我正在使用MSVC++2013：它不支持constexpr，但CUDA 8RC仍然不支持MSVC++2015。我将尝试“只是const”选项，但我猜编译器会对主机const（除了整数）给出像“无法从设备访问主机变量”的错误。 - Serge Rogatch

1

我认为它不能在Visual C++上工作。但是它可以在g++上工作。 - talonmies

是的，还有另一个原因：常量值类型的构造函数是设备函数。虽然有一对主机类型，但它们是不同的类型。 - Serge Rogatch

0

为了初始化它，你必须使用cudaMemcpyToSymbol。它不是一个编译时常量，而是存储在设备的常量内存中，并且相较于全局内存具有一些优势。来自CUDA博客的原文如下：

对于一个半warp的所有线程来说，如果所有线程都读取同一个地址，那么从常量缓存器中读取的速度就像从寄存器中读取一样快。但是，半warp内的线程对于不同的地址访问是串行化的，因此成本与半warp内所有线程读取的不同地址数量呈线性关系。

你不需要使用const，也不能使用它。它不是C++常量，因为你需要通过cudaMemcpyToSymbol来修改它。所以从C++的角度来看，它并不是一个"真正"的常量。但是在设备内核中，它表现得像一个常量，因为你只能通过cudaMemcpyToSymbol来修改它，而这个函数只能从主机端调用。

- curious_amateur

就C++而言，这样的变量不会是“const”：如果我尝试使用“const”和初始化程序来定义它，我会得到问题中提到的错误。如果我使用“const”而没有初始化程序来定义它，我也会得到一个需要初始化程序的错误。所以这是编译器要求的一个循环。 - Serge Rogatch

这里对常量内存的所有引用都是无关紧要的。 - talonmies

注意2是关于性能的，我尝试澄清性能疑问。 - curious_amateur

网页内容由stack overflow 提供, 点击上面的

可以查看英文原文，
原文链接

- einpoklum · Accepted Answer

使用 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，请参阅此处了解更多详细信息（感谢@harrism）。