如何像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不仅仅是出于美学原因,而且因为在实践中表达式更复杂并且涉及内部数据类型,而不仅仅是双重的,这是不可能的.因此,每次在每个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看到更多细节(感谢@harrism).