我在常量内存中有一个数组（它是一个全局变量），并通过函数调用cudaGetSymbolAddress获取了对它的引用。当我使用这个引用来获取常量数据时，与使用全局变量相比，我的内核运行缓慢。这是什么原因呢？
__constant__ int g[2] = {1,2};
// __device__ int g[2] = {1,2};
// kernel: use by reference
__global__ void add_1( int *a, int *b, int *c, int *f )
{
int tid = blockIdx.x * blockDim.x + threadIdx.x;
c[tid] = f[0] * a[tid] + f[1] * b[tid];
}
// kernel: use global variable
__global__ void add_2( int *a, int *b, int *c, int *f )
{
int tid = blockIdx.x * blockDim.x + threadIdx.x;
c[tid] = g[0] * a[tid] + f[1] * b[tid];
}
int main()
{
......
// a,b,c are large arrays in device memory of size 40960.
int *f;
cudaGetSymbolAddress( (void **)&f, (char *)&g);
add_1 <<< 160, 256 >>> ( a, b, c, f );
......
}
这是示例代码，所有线程在warp中同时加载相同的位置。注释掉的代码是通过直接访问常量内存来实现的。
为什么不使用常量内存缓存的解释（by talonmies）
原因是缺乏常量缓存。
只有当编译器在被明确标记为处于常量状态空间的变量上发出特定的PTX指令（ld.const）时，才会发生缓存访问。编译器知道如何做到这一点的方式是当一个变量被声明为__constant__时——它是一个静态的、编译时的属性，影响着代码生成。同样的过程不能在运行时发生。
如果你传递一个指针到全局内存中，并且编译器无法确定该指针是否在常量状态空间中，那么它就不会生成正确的PTX来通过常量缓存访问该内存。结果访问会变慢。
未解决的问题
为什么即使数组g被声明为__device__变量，当引用它时代码也会变慢。通过查看PTX代码，可以发现将全局内存加载到寄存器中：
- 使用2个
ld.global.s32指令，将4个字节加载到一个寄存器中。（在使用引用的代码中） - 使用1个
ld.global.v2.s32指令，将8个字节加载到2个寄存器中。（在使用全局变量的代码中）
有什么区别，是否有任何文档参考？
__constant__声明。你如何调用内核？ - tera