我有一个CUDA内核，其中每个线程都从全局内存中读取相同的值 . 所以类似于：
__global__ void my_kernel(const float4 * key_pts)
{
if (key_pts[blockIdx.x] < 0 return;
}
内核配置如下：
dim3 blocks(16, 16);
dim3 grid(2000);
my_kernel<<<grid, blocks, 0, stream>>>(key_pts);
我的问题是，这是否会导致某种银行冲突或CUDA中的次优访问 . 我必须承认我还没有详细了解这个问题 .
我想我可以做以下事情，以防我们有次优访问：
__global__ void my_kernel(const float4 * key_pts)
{
__shared__ float x;
if (threadIdx.x == 0 && threadIdx.y == 0)
x = key_pts[blockIdx.x];
__syncthreads();
if (x < 0) return;
}
虽然做了一些时间安排，但我认为这两者之间没有任何区别，但到目前为止我的测试数据有限 .
1 回答
银行冲突适用于shared memory，而非全局记忆 .
由于所有线程都需要（最终）相同的值来做出决定，因此不会对全局内存产生次优访问，因为有一个广播机制，以便同一个warp中的所有线程，从global memory请求相同的位置/值，将检索没有任何序列化或开销 . warp中的所有线程可以同时进行维护：
此外，假设您的GPU具有缓存（cc2.0或更高版本），则遇到此问题的第一个warp从全局内存中检索到的值可能会在缓存中用于此后发生的warp .
我不希望这两种情况之间有太大的性能差异 .