NSight Compute - 预计银行冲突但未检测到任何

问题描述

我试图检测矩阵转置内核的共享内存库冲突。第一个内核进行矩阵转置没有填充，因此应该有bank冲突，而第二个内核使用padding，应该没有bank冲突。

但是，在内存工作负载部分使用 NSight Compute 进行分析显示两个内核的存储体冲突为 0。

我将内核实现为这样的设备功能

// tiled,with padding (expecting no bank conflicts)
template <class value_type,class container_type = value_type*>
__device__
void
transpose_padded(container_type m1,container_type m2,size_t width)
{
    __shared__ value_type tile[BLOCK_WIDTH][BLOCK_WIDTH+1];
    // BLOCK_WIDTH = 32,global scope constant
    auto row = blockDim.y*blockIdx.y + threadIdx.y;
    auto col = blockDim.x*blockIdx.x + threadIdx.x;
    auto index = row * width + col;

    auto tr_row = blockDim.y * blockIdx.x + threadIdx.y;
    auto tr_col = blockDim.x * blockIdx.y + threadIdx.x;
    auto tr_index = tr_row * width + col;

    auto local_x = threadIdx.x;
    auto local_y = threadIdx.y;
    tile[local_x][local_y] = m1[index];
    __syncthreads();
    if (tr_row < width && tr_col < width)
    {
        m2[tr_index] = tile[local_y][local_x];
    }
    
    return;
}

// tiled,without padding (expecting bank conflicts)
template <class value_type,class container_type = value_type*>
__device__
void
transpose_tiled(container_type input,container_type output,size_t width)
{
    // assuming square blocks
    extern __shared__ value_type input_tile[];
    auto row = blockDim.y*blockIdx.y + threadIdx.y;
    auto col = blockDim.x*blockIdx.x + threadIdx.x;
    auto matrix_index = row*width + col;

    auto tr_row = col;
    auto tr_col = row;
    auto tr_index = tr_row*width + tr_col;
    
    // coalesced global memory access
    auto shared_index = threadIdx.y*blockDim.x+threadIdx.x;
    input_tile[shared_index]= input[matrix_index];
    __syncthreads();
    if (tr_row < width && tr_col < width)
        output[tr_index] = input_tile[shared_index];
    return;
}

我使用的输入矩阵的尺寸为 100x100。在两个内核中，块大小都是 32x32 线程。实例化的值类型为 double。

真的没有银行冲突，还是完全是其他原因造成的？我可以使用其他部分的哪些其他信息来确定是否可能存在银行冲突？

解决方法

对于 32x32 的块尺寸，我不希望任何一个内核都表现出银行冲突。银行冲突包含在 many resources 中，包括 cuda 标签上的 many questions，因此我将简要总结一下。

当同一 warp 中的两个或多个线程（并且在同一条指令期间）执行共享加载或共享存储时，会出现 Bank 冲突，其中这两个线程引用的位置在同一 bank 但不相同位置。

一个 bank 可以粗略地描述为共享内存中的一列，当共享内存被认为是一个 2D 数组，宽度为 32 个 bank 乘以每个 bank 32 位的数量，即宽度为 128 字节。

这些定义应提供相当完整的理解并涵盖大多数感兴趣的情况。我们可以从中得出一个观察结果，即对于全局内存合并加载/存储很好地工作的相同访问模式（相邻线程访问内存中的相邻元素）也可以很好地避免组冲突。（这不是唯一适用于共享内存的模式，但它是一种规范模式。）

转向你的代码，然后：

您已经（正确地）指出您不希望在第一个代码中出现共享银行冲突。该代码中的共享负载：
```
 = tile[local_y][local_x];
```
将 threadIdx.x（或包含 threadIdx.x 且没有任何乘法因子的索引）作为最后一个下标，这是CUDA 中用于“nice”访问的规范模式.它表示相邻线程将从内存中的相邻位置读取。这对全局内存和共享内存都适用。

对于共享商店：
```
tile[local_x][local_y] = 
```
乍一看，这似乎是跨经线的“列式”访问，对 CUDA（无论是全局还是共享）来说是典型的错误，但您正在使用 shared memory offset-the-columns-by-1 trick：
```
__shared__ value_type tile[BLOCK_WIDTH][BLOCK_WIDTH+1];
                                                   ^^
```
这样的情况也得到处理/排序。对于 32x32 块配置（每个经线中的所有 32 个线程将具有单调递增的 threadIdx.x 和 constant threadIdx.y），此处预计不会发生 bank 冲突。
对于第二个代码，只有一种索引模式用于共享存储和共享加载：
```
input_tile[shared_index]=
= input_tile[shared_index];
```
即：
```
auto shared_index = threadIdx.y*blockDim.x+threadIdx.x;
```
因此，要回答这种情况下的银行冲突问题，我们只需要研究一种访问模式。让我们看看我们是否可以走同样的捷径。索引模式是否包含 threadIdx.x 且没有乘法因子（在最后一个下标中）？ 是的。因此，warp 中的相邻线程将访问内存中的相邻位置，这是一种典型的良好模式，即没有 bank 冲突。