您的问题在概念上与我在 StackOverflow 上的第一个问题类似:移动 (BS_X+1)(BS_Y+1) BS_X 的全局内存矩阵BS_Y 线程 https://stackoverflow.com/questions/13771538/moving-a-bs-x1bs-y1-global-memory-matrix-by-bs-xbs-y-threads.
您面临以下问题:每个线程块的大小TILE_WIDTHxTILE_WIDTH
应该填充大小的共享内存区域(TILE_WIDTH + Mask_width - 1)x(TILE_WIDTH + Mask_width - 1)
.
4)一般来说,有两个载荷的直观解释是什么?
由于共享内存区域(TILE_WIDTH + Mask_width - 1)x(TILE_WIDTH + Mask_width - 1)
大于块大小TILE_WIDTHxTILE_WIDTH
并假设它小于2xTILE_WIDTHxTILE_WIDTH
,那么每个线程最多应将两个元素从全局内存移动到共享内存。这就是为什么要进行两阶段加载的原因。
1) The destY
and destX
指数。将输出索引除以输入图块宽度意味着什么?
这涉及指定加载的第一个加载阶段TILE_WIDTHxTILE_WIDTH
来自全局内存的元素并填充共享内存区域的最上部。
所以,操作
dest = threadIdx.y * TILE_WIDTH + threadIdx.x;
展平通用线程的 2D 坐标,同时
destX = dest % w;
destY = dest / w;
进行逆操作,因为它计算通用线程相对于共享内存区域的 2D 坐标。
2) The srcY
add srcX
指数。为什么destY
and destX
指数参与srcY
add srcX
index?
srcY = blockIdx.y * TILE_WIDTH + destY - Mask_radius;
srcX = blockIdx.x * TILE_WIDTH + destX - Mask_radius;
(blockIdx.x * TILE_WIDTH, blockIdx.y * TILE_WIDTH)
如果块大小和共享内存大小相同,则为全局内存位置的坐标。由于您还从相邻图块“借用”内存值,因此您必须将上述坐标移动(destX - Mask_radius, destY - Mask_radius)
.
3)为什么在第二次加载时我们使用偏移量TILE_WIDTH * TILE_WIDTH?
你有这个偏移量是因为在第一个记忆阶段你已经填满了“第一个”TILE_WIDTHxTILE_WIDTH
共享内存的位置。
EDIT
下图说明了压扁螺纹索引之间的对应关系dest
和共享内存位置。在图中,蓝色框代表通用图块的元素,而红色框代表相邻图块的元素。蓝色和红色框的并集对应于整体共享内存位置。如您所见,所有256
线程块的线程参与填充绿线上方共享内存的上部部分,而只有145
参与填充绿线以下共享内存的下部。现在你也应该明白了TILE_WIDTH x TILE_WIDTH
offset.
请注意,您最多有2
由于参数的特定选择,每个线程的内存负载。例如,如果您有TILE_WIDTH = 8
,那么线程块中的线程数为64
,而共享内存大小为12x12=144
,这意味着每个线程负责执行at least 2
共享内存写入144/64=2.25
.