__device__ __constant__ const

Question

有什么区别,在CUDA程序中定义设备常量的最佳方法是什么？在C++,主机/设备程序中,如果我想将常量定义为设备常量内存,我可以这样做

__device__ __constant__ float a = 5;
__constant__ float a = 5;

问题1.在设备2.x和CUDA 4上,它是否相同,

__device__ const float a = 5;

问题2.为什么在PyCUDA SourceModule("""......""")中,它只编译设备代码,甚至以下工作？

const float a = 5;

Answer 1

在CUDA中__constant__是一个变量类型限定符,它指示声明的变量将存储在设备常量存储器中.引用CUDA编程指南的B 2.2节

__constant__可选择与之一起使用的限定符__device__声明一个变量:

• 驻留在恒定的存储空间中,
• 有应用程序的生命周期,
• 可以从网格中的所有线程访问,也可以从主机通过运行时库(cudaGetSymbolAddress()/ cudaGetSymbolSize()/ cudaMemcpyToSymbol()/ cudaMemcpyFromSymbol()用于运行时API和cuModuleGetGlobal()驱动程序API)访问.

在CUDA,常量内存是专门的,静态的,全局存储器通过高速缓存访问的区域(有专门的一套其目的PTX加载指令),这是统一和只读在运行的内核的所有线程.但是,通过使用上面引用的主机端API,可以在运行时修改常量内存的内容.这与使用声明向编译器声明变量不同const,后者是在声明范围内向变量添加只读特性.两者完全不是一回事.