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.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,我的问题似乎不合标准。我应该删除它吗?以上是关于NSight Compute - 预期银行冲突但未检测到任何的主要内容,如果未能解决你的问题,请参考以下文章
Nsight Compute 说:“此设备不支持分析” - 为啥?