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

Posted

技术标签:

【中文标题】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.xconstant threadIdx.y)。

    对于第二个代码,只有一种索引模式用于共享存储和共享加载:

    input_tile[shared_index]=
    = input_tile[shared_index];
    

    这是:

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

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

【讨论】:

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

以上是关于NSight Compute - 预期银行冲突但未检测到任何的主要内容,如果未能解决你的问题,请参考以下文章

Nsight Compute 如何确定/显示共享内存指标?

如何减少此代码中的银行冲突?

Nsight Compute 说:“此设备不支持分析” - 为啥?

如何修复“错误:预期在 'ProxyZone' 中运行,但未找到。”在摩卡测试?

Jest Vue 预期的模拟函数已被调用,但未调用

无效属性:预期的数组,但未定义