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

Posted 2023-04-15

【中文标题】从同一位置读取时 CUDA 中的银行冲突

【英文标题】：Bank conflict in CUDA when reading from the same location

【发布时间】：2015-09-23 22:12:54

【问题描述】：

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

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

【讨论】：

感谢您的精彩解释！