如何解决NSight Compute - 预计银行冲突但未检测到任何
我试图检测矩阵转置内核的共享内存库冲突。第一个内核进行矩阵转置没有填充,因此应该有bank冲突,而第二个内核使用padding,应该没有bank冲突。
但是,在内存工作负载部分使用 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。
真的没有银行冲突,还是完全是其他原因造成的?我可以使用其他部分的哪些其他信息来确定是否可能存在银行冲突?
解决方法
对于 32x32 的块尺寸,我不希望任何一个内核都表现出银行冲突。银行冲突包含在 many resources 中,包括 cuda
标签上的 many questions,因此我将简要总结一下。
当同一 warp 中的两个或多个线程(并且在同一条指令期间)执行共享加载或共享存储时,会出现 Bank 冲突,其中这两个线程引用的位置在同一 bank 但不相同位置。
一个 bank 可以粗略地描述为共享内存中的一列,当共享内存被认为是一个 2D 数组,宽度为 32 个 bank 乘以每个 bank 32 位的数量,即宽度为 128 字节。
>这些定义应提供相当完整的理解并涵盖大多数感兴趣的情况。我们可以从中得出一个观察结果,即对于全局内存合并加载/存储很好地工作的相同访问模式(相邻线程访问内存中的相邻元素)也可以很好地避免组冲突。 (这不是唯一适用于共享内存的模式,但它是一种规范模式。)
转向你的代码,然后:
-
您已经(正确地)指出您不希望在第一个代码中出现共享银行冲突。该代码中的共享负载:
= tile[local_y][local_x];
将
threadIdx.x
(或包含threadIdx.x
且没有任何乘法因子的索引)作为最后一个下标,这是CUDA 中用于“nice”访问的规范模式.它表示相邻线程将从内存中的相邻位置读取。这对全局内存和共享内存都适用。对于共享商店:
tile[local_x][local_y] =
乍一看,这似乎是跨经线的“列式”访问,对 CUDA(无论是全局还是共享)来说是典型的错误,但您正在使用 shared memory offset-the-columns-by-1 trick:
__shared__ value_type tile[BLOCK_WIDTH][BLOCK_WIDTH+1]; ^^
这样的情况也得到处理/排序。对于 32x32 块配置(每个经线中的所有 32 个线程将具有单调递增的
threadIdx.x
和 constantthreadIdx.y
),此处预计不会发生 bank 冲突。 -
对于第二个代码,只有一种索引模式用于共享存储和共享加载:
input_tile[shared_index]= = input_tile[shared_index];
即:
auto shared_index = threadIdx.y*blockDim.x+threadIdx.x;
因此,要回答这种情况下的银行冲突问题,我们只需要研究一种访问模式。让我们看看我们是否可以走同样的捷径。索引模式是否包含
threadIdx.x
且没有乘法因子(在最后一个下标中)? 是的。因此,warp 中的相邻线程将访问内存中的相邻位置,这是一种典型的良好模式,即没有 bank 冲突。
版权声明:本文内容由互联网用户自发贡献,该文观点与技术仅代表作者本人。本站仅提供信息存储空间服务,不拥有所有权,不承担相关法律责任。如发现本站有涉嫌侵权/违法违规的内容, 请发送邮件至 dio@foxmail.com 举报,一经查实,本站将立刻删除。