__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__是一个变量类型限定符，它指示要声明的变量将存储在设备常量内存中。引用B节2.2 of CUDA编程指南

__constant__限定符(可与__device__一起使用)声明了以下变量：

• 驻留在常量内存空间中，
• 具有应用程序的生存期，
• 可以从网格中的所有线程以及通过运行时库(运行时API的cudaGetSymbolAddress()/ cudaGetSymbolSize() / cudaMemcpyToSymbol() / cudaMemcpyFromSymbol()和驱动程序API的cuModuleGetGlobal() )访问。

在CUDA中，常量内存是通过缓存访问的专用静态全局内存区域(为此目的有一组专用的PTX加载指令)，对于运行中的内核中的所有线程来说，这些指令是统一的和只读的。但是，通过使用上面引用的主机端API，可以在运行时修改常量内存的内容。这与使用const声明向编译器声明变量不同，后者在声明范围内向变量添加只读特性。这两件事完全不一样。