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

Question

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

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

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

Answer 1

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

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

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

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

然后转到您的代码：

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

    = 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）。

2. 对于第二个代码，只有一种索引模式用于共享存储和共享加载：

   input_tile[shared_index]=
   = input_tile[shared_index];
   

   这是：

   auto shared_index = threadIdx.y*blockDim.x+threadIdx.x;
   

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