CUDA恒定存储器应该是均匀访问的吗？

Question

我的CUDA应用程序具有小于8KB的恒定内存.由于它都将被缓存,我是否需要担心每个线程访问相同的地址以进行优化？

如果是,我如何确保所有线程同时访问同一地址？

Answer 1

由于它都将被缓存,我是否需要担心每个线程访问相同的地址以进行优化？

是.该缓存本身只能充当了每个周期32位字.

如果是,我如何确保所有线程同时访问同一地址？

确保任何类型的索引或解决您使用引用在常量存储区域的元件不依赖于任何一个内置的线程变量,例如threadIdx.x,threadIdx.y或threadIdx.z.请注意,实际要求不如此严格.只要索引计算给定warp中每个线程的相同数字,就可以实现必要的目标.这里有一些例子:

__constant__ int data[1024];
...
// assume 1D threadblock
int idx = threadIdx.x;
int bidx = blockIdx.x;
int a = data[idx];      // bad - every thread accesses a different element
int b = data[12];       // ok  - every thread accesses the same element
int c = data[b];        // ok  - b is a constant w.r.t threads
int d = data[b + idx];  // bad
int e = data[b + bidx]; // ok
int f = data[idx/32];   // ok - the same element is being accessed per warp