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

Posted 2023-03-23

【中文标题】NSight Compute - 预期银行冲突但未检测到任何

【英文标题】：NSight Compute - expecting bank conflicts but not detecting any

【发布时间】：2021-07-19 20:45:46

【问题描述】：

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

但是，在内存工作负载部分使用 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。

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

【问题讨论】：

银行冲突取决于执行参数和类型。一些未实例化的模板没有说明可能发生或可能不会发生的事情

我编辑了实例化和内核启动的细节。还有什么我应该补充的来改进这个问题吗？

Are there really no bank conflicts 您的任何一个内核都不应显示 32x32 块大小的存储库冲突。分析器是正确的。由于您在第二个代码中有注释expecting bank conflicts，也许您应该解释一下您的理由。

【参考方案1】：

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

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

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

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

然后转到您的代码：

您已经（正确地）表明您不希望在第一个代码上出现共享银行冲突。该代码中的共享负载：

 = tile[local_y][local_x];

具有threadIdx.x（或包含threadIdx.x 且没有任何乘法因子的索引）作为最后一个下标，这是CUDA 中用于“良好”访问的规范模式。它表示相邻线程将从内存中的相邻位置读取。这适用于全局和共享内存。

对于共享存储：

tile[local_x][local_y] = 

乍一看，这似乎是跨经线的“列式”访问，对 CUDA（无论是全局还是共享）来说都是不利的，但您使用的是 shared memory offset-the-columns-by-1 trick：

__shared__ value_type tile[BLOCK_WIDTH][BLOCK_WIDTH+1];
                                                   ^^

所以这种情况也得到处理/排序。对于 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;

因此，要回答这种情况下的银行冲突问题，我们只需要研究一种访问模式即可。让我们看看我们是否可以走同样的捷径。索引模式是否包括 threadIdx.x 没有乘法因子（在最后一个下标中）？ 是的。因此，warp 中的相邻线程将访问内存中的相邻位置，这是一种典型的良好模式，即没有银行冲突。

【讨论】：

感谢您抽出宝贵时间来很好地解释这一点。我一直在尝试重现我在其他地方从记忆中看到的银行冲突示例。我刚刚找到它并意识到他们使用了不同的访问模式。总的来说，看到反对票的数量和 cmets，我的问题似乎不合标准。我应该删除它吗？