从同一位置读取时 CUDA 中的银行冲突

Question

我有一个 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;
}

虽然做了一些计时，但我没有看到两者之间有任何区别，但到目前为止我的测试数据有限。

Answer 1

bank 冲突适用于 shared memory，而不是全局内存。

由于所有线程（最终）需要相同的值来做出决定，因此这不会产生对全局内存的次优访问，因为存在 广播机制 以便同一个 warp 中的所有线程，从 global memory 请求相同的 location/value，将在没有任何序列化或开销的情况下检索它。可以同时服务warp中的所有线程：

Note that threads can access any words in any order, including the same words.

此外，假设您的 GPU 有缓存（cc2.0 或更新版本），从全局内存中检索到的第一个扭曲的值很可能会在缓存中用于达到此点的后续扭曲。

我预计这两种情况之间的性能差异不大。

从同一位置读取时 CUDA 中的银行冲突

Bank conflict in CUDA when reading from the same location

cuda