时间戳

首页 > 解决方案 > 如何避免矩阵乘法 CUDA 内核中的未合并访问？

parallel-processing - 如何避免矩阵乘法 CUDA 内核中的未合并访问？

问题描述

我正在通过《大规模并行处理器编程》这本书学习 CUDA 。第 5 章的一个练习题让我很困惑：

对于 BLOCK_SIZE 的可能值范围之外的平铺矩阵乘法，对于 BLOCK_SIZE 的哪些值，内核将完全避免对全局内存的未合并访问？（你只需要考虑方块）

据我了解，BLOCK_SIZE 对内存合并几乎没有作用。只要单个warp中的线程访问连续的元素，我们就会有一个合并的访问。我无法弄清楚内核在哪里对全局内存进行了未合并的访问。你们有什么提示吗？

这是内核的源代码：

#define COMMON_WIDTH 512
#define ROW_LEFT 500 
#define COL_RIGHT 250
#define K 1000
#define TILE_WIDTH 32
__device__ int D_ROW_LEFT = ROW_LEFT;
__device__ int D_COL_RIGHT = COL_RIGHT;
__device__ int D_K = K;
.....
__global__
void MatrixMatrixMultTiled(float *matrixLeft, float *matrixRight, float *output){
    __shared__  float sMatrixLeft[TILE_WIDTH][TILE_WIDTH];
    __shared__  float sMatrixRight[TILE_WIDTH][TILE_WIDTH];  
   int bx = blockIdx.x; int by = blockIdx.y;
   int tx = threadIdx.x; int ty = threadIdx.y;
   int col = bx * TILE_WIDTH + tx;
   int row = by * TILE_WIDTH + ty;
   float value = 0;
   for (int i = 0; i < ceil(D_K/(float)TILE_WIDTH); ++i){
       if (row < D_ROW_LEFT && row * D_K + i * TILE_WIDTH  +tx < D_K){
        sMatrixLeft[ty][tx]  = matrixLeft[row * D_K + i * TILE_WIDTH  +tx];
       }
       if (col < D_COL_RIGHT && (ty + i * TILE_WIDTH) * D_COL_RIGHT  + col < D_K ){
        sMatrixRight[ty][tx] = matrixRight[(ty + i * TILE_WIDTH) * D_COL_RIGHT  + col];
       }
       __syncthreads();
       for (int j = 0; j < TILE_WIDTH; j++){
           value += sMatrixLeft[ty][j] * sMatrixRight[j][tx]; 
       }
       __syncthreads();
   }
   if (row < D_ROW_LEFT && col < D_COL_RIGHT ){
        output[row * D_COL_RIGHT + col] = value;
       }
}

标签： parallel-processingcuda

解决方案

您的问题不完整，因为您发布的代码没有提及BLOCK_SIZE，这肯定至少与书中提出的问题非常相关。更一般地，在没有启动配置的情况下提出内核的问题通常是不完整的，因为启动配置通常与内核的正确性和行为相关。

目前我还没有重读这本书的这一部分。但是，我假设内核启动配置包含类似于以下内容的块维度：（您的问题中没有此信息，但在我看来，应该包含一个明智的问题）

dim3 dimBlock(BLOCK_SIZE, BLOCK_SIZE);
dim3 dimGrid(...,...);

我将假设内核启动是由以下内容给出的：

MatrixMatrixMultTiled<<<dimGrid, dimBlock>>>(...);

您的声明：“只要单个经线中的线程访问连续元素，我们就会有合并访问。” 是一个合理的工作定义。BLOCK_SIZE考虑到上述假设以弥补您不完整问题中的空白，让我们证明这对于 的某些选择是违反的。

合并访问是一个仅适用于全局内存访问的术语。因此，我们将忽略对共享内存的访问。在本次讨论中，我们还将忽略对__device__变量的访问，例如D_ROW_LEFT. （对这些变量的访问似乎是统一的。我们可以质疑这是否构成合并访问。我的主张是它确实构成了合并访问，但我们不需要在这里解包。）因此我们只剩下 3 个“访问“点：

matrixLeft[row * D_K + i * TILE_WIDTH  +tx];
matrixRight[(ty + i * TILE_WIDTH) * D_COL_RIGHT  + col];
output[row * D_COL_RIGHT + col]

现在，举个例子，假设BLOCK_SIZE是 16。上述任何访问点是否会违反您的声明“单个经线中的线程访问连续元素”？

让我们从块 (0,0) 开始。因此row等于threadIdx.y和col等于threadIdx.x。让我们考虑该块中的第一个扭曲。因此，该 warp 中的前 16 个线程的值将为threadIdx.y0，并且它们的threadIdx.x值将从 0..15 增加。同样，该经线中的第二个 16 线程将具有threadIdx.y值 1，并且它们的threadIdx.x值将从 0..15 增加。

现在让我们计算为上面的第一个访问点生成的实际索引，跨越 warp。假设我们在第一次循环迭代中，所以i是零。因此：

matrixLeft[row * D_K + i * TILE_WIDTH  +tx];

减少为：

matrixLeft[threadIdx.y * D_K + threadIdx.x];

D_K这里只是K变量的设备副本，即 1000。现在让我们在我们选择的块 (0,0) 中跨我们选择的扭曲 (0) 评估上面的缩减索引表达式：

warp lane:    0  1  2  3  4  5  6  .. 15     16   17   18 .. 31
threadIdx.x   0  1  2  3  4  5  6     15      0    1    2    15
threadIdx.y   0  0  0  0  0  0  0      0      1    1    1     1
index:        0  1  2  3  4  5  6     15   1000 1001 1002  1015

因此，此处生成的索引模式显示了经线中第 16 和第 17 线程之间的不连续性，并且访问模式不符合您之前陈述的条件：

“单个经线内的线程访问连续元素”

在这种情况下，我们没有合并访问权限（至少对于float数量而言）。

推荐阅读