NSight Compute - 期待银行冲突但没有检测到任何

Question

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

但是，在内存工作负载部分使用 NSight Compute 进行的分析显示两个内核的 bank 冲突为 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。

是否真的没有银行冲突，或者这完全是由其他原因引起的？我可以使用其他部分的哪些其他信息来确定是否存在银行冲突？

Answer 1

对于 32x32 的块尺寸，我不希望任何一个内核都出现内存冲突。银行冲突在 cuda 标签的 many resources including many questions 中有介绍，所以我将简要总结一下。

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

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

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

转向您的代码，然后：

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

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

NSight Compute - 期待银行冲突但没有检测到任何

NSight Compute - expecting bank conflicts but not detecting any

c++

profiling

cuda

gpu

nsight