如何定义CUDA设备常量,如C++常量/常量表达式
How to define CUDA device constant like a C++ const/constexpr?
在.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
编写的变量。
使用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(。
要使显示的代码按预期编译和工作,需要在运行时初始化变量,而不是在编译时初始化。为此,向cudaMemcpyToSymbol
添加一个主机端调用,类似于:
__device__ double cdInf;
// ...
double val = HUGE_VAL / 4
cudaMemcpyToSymbol(cdInf, &val, sizeof(double));
然而,对于单个值,将其作为内核参数传递似乎更明智。编译器将自动将参数存储在所有支持的体系结构上的常量内存中,并且有一个"免费"的常量缓存广播机制,这将使在运行时访问该值的成本可以忽略不计。
要初始化它,必须使用cudaMemcpyToSymbol
。它不是编译时常数,而是存储在设备的常数存储器中,并且与全局存储器相比具有一些优势。来自CUDA博客:
对于半曲速的所有线程,从常量缓存中读取如下只要所有线程都读取相同的内容,就可以像从寄存器中读取一样快住址通过半曲速内的线程访问不同的地址是序列化的,因此成本随不同数量的线性扩展由半曲速内的所有线程读取的地址。
您不需要使用const
,也不能使用它。它不是c++常量,因为您需要通过cudaMemcpyToSymbol
修改它。因此,至少从c++的角度来看,它不是一个"真正的"常数。但它在设备内核中的行为就像一个常量,因为您只能通过只能从主机调用的cudaMemcpyToSymbol
来修改它。
- #定义c-预处理器常量..我做错了什么
- 用C++中的一个变量定义一个常量
- 什么时候在C++中返回常量引用是个好主意
- 代理对象的常量正确性
- 我想将一个对T类型的非常量左值引用绑定到一个T类型的临时值
- 通过多个头文件使用常量变量
- 在cuda线程之间共享大量常量数据
- 不能在初始值设定项列表中将非常量表达式从类型 'int' 缩小到'unsigned long long'
- 有没有什么方法可以使用一个函数中定义的常量变量,也可以由c++中同一程序中的其他函数使用
- 是默认情况下分配给char数组常量的值
- 私有类型的静态常量成员
- 类似枚举的计算常量
- 递归模板化函数不能分配给具有常量限定类型"const tt &"的变量"state"
- 为什么我可以通过引用修改常量返回
- 如何创建长度由常量参数指定的数组
- 当一个值是非常量但用常量表达式初始化时使用constexpr
- 返回常量对象引用 (getter) 和仅返回字符串有什么区别?
- 隐式常量/非常量运算符布尔
- 非常量变量只读位置的赋值
- constexpr构造函数需要常量成员函数时出现问题