CUDA 共享内存效率为 50%？

Posted 2023-04-15

【中文标题】CUDA 共享内存效率为 50%？

【英文标题】：CUDA shared memory efficiency at 50%?

【发布时间】：2018-11-30 02:26:37

【问题描述】：

我有以下代码使用共享内存执行平铺矩阵转置以提高性能。共享内存用 1 列填充，以避免 32x32 线程块的银行冲突。

__global__ void transpose_tiled_padded(float *A, float *B, int n)

    int i_in = blockDim.x*blockIdx.x + threadIdx.x;
    int j_in = blockDim.y*blockIdx.y + threadIdx.y;
    int i_out = blockDim.x*blockIdx.y + threadIdx.x;
    int j_out = blockDim.y*blockIdx.x + threadIdx.y;

    extern __shared__ float tile[];

    // coalesced read of A rows to (padded) shared tile column (transpose)
    tile[threadIdx.y + threadIdx.x*(blockDim.y+1)] = A[i_in + j_in*n];
    __syncthreads();

    // coalesced write from (padded) shared tile column to B rows
    B[i_out + j_out*n] = tile[threadIdx.x + threadIdx.y*(blockDim.x+1)];

运行此代码，我在 NVIDIA 视觉分析器中获得了 100% 的共享内存效率，正如我所期望的那样。但是，当我使用 16x16 线程块运行它时，我只能获得 50% 的效率。这是为什么？据我所知，经线中的线程没有从具有这种布局的同一银行读取。还是我弄错了？

【参考方案1】：

是的，你错了。

考虑对 16x16 块中 warp 0 的此（读取）访问：

tile[threadIdx.x + threadIdx.y*(blockDim.x+1)];
     ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
                     "index"

以下是warp中每个线程的相关计算：

warp lane:    0  1  2  3  4  5  6  7  8  9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 23 25 26 27 28 29 30 31
threadIdx.x:  0  1  2  3  4  5  6  7  8  9 10 11 12 13 14 15  0  1  2  3  4  5  6  7  8  9 10 11 12 13 14 15
threadIdx.y:  0  0  0  0  0  0  0  0  0  0  0  0  0  0  0  0  1  1  1  1  1  1  1  1  1  1  1  1  1  1  1  1
"index":      0  1  2  3  4  5  6  7  8  9 10 11 12 13 14 15 17 18 19 20 21 22 23 24 25 26 27 28 29 30 31 32
bank:         0  1  2  3  4  5  6  7  8  9 10 11 12 13 14 15 17 18 19 20 21 22 23 24 25 26 27 28 29 30 31  0

所以我们看到，对于这个 warp，第一个和最后一个线程都从 bank 0 读取。这会导致 2-way bank 冲突、2-way 序列化和 50% 的效率。

【讨论】：

我明白了。那么是否有系统的方法来计算填充宽度以确保特定块尺寸不会发生银行冲突？

线程 x,y=[0,1] 的索引不应该是 16 吗？这意味着该曲速通道中没有银行冲突，因此，冲突必须存在于其他地方。

不，这是不正确的。应该是 17。blockDim.x 是 16。加一就是 17。研究代码。

@RobertCrovella 你是完全正确的。没注意是我的错。谢谢你的回复。