为什么 CUDA 共享内存在平铺矩阵乘法中比全局内存慢？答案

【问题标题】：Why CUDA shared memory is slower than global memory in tiled matrix multiplication?为什么 CUDA 共享内存在平铺矩阵乘法中比全局内存慢？
【发布时间】：2018-05-18 03:15:53
【问题描述】：

我有平铺矩阵乘法代码有和没有共享内存。下面是使用全局内存的矩阵乘法：

__global__ 
void MatrixMulKernel(float* M, float* N, float* P, int Width)
{
int Row = blockIdx.y*blockDim.y + threadIdx.y;
int Col = blockIdx.x*blockDim.x + threadIdx.x;
if ((Row < Width) && (Col < Width)) {
    float Pvalue = 0;
    for (int k = 0; k < Width; ++k)
    {
        Pvalue += M[Row*Width + k] * N[k*Width + Col];
    }

    P[Row*Width + Col] = Pvalue;
}
}

下面是使用共享内存的矩阵乘法：

__global__
void MatrixMulKernel(float* d_M, float* d_N, float* d_P, int Width)
{
__shared__ float Mds[blockWidth][blockWidth];
__shared__ float Nds[blockWidth][blockWidth];
int tx = threadIdx.x; int ty = threadIdx.y;
int bx = blockIdx.x; int by = blockIdx.y;

int row = by * blockWidth + ty;
int col = bx * blockWidth + tx;
float pvalue = 0;

for (int m = 0; m < Width / blockWidth; ++m)
{
    Mds[ty][tx] = d_M[row * Width + m*blockWidth + tx];
    Nds[ty][tx] = d_N[(m*blockWidth + ty)*Width + col];
    __syncthreads();
    for (int k = 0; k < blockWidth; ++k)
    {
        pvalue += Mds[ty][k]*Nds[k][tx];
    }
    __syncthreads();
}
d_P[row*Width + col] = pvalue;
}

据我所知，使用共享内存应该更快，但在比较这两个代码时，我发现没有共享内存的代码对于 1600*1600 矩阵的运行速度要快约 2 秒。对于这种速度差异或我的代码有什么问题有什么解释吗？

我的老师使用“Programming Massively Parallel Processors”一书作为主要文本资源，这两个代码来自该书。

编辑：

内核配置：

int NumBlocks =ceil( Width / blockWidth);  // blockWidth = 16
dim3 dimGrid(NumBlocks, NumBlocks,1); // Width = 1600
dim3 dimBlock(blockWidth, blockWidth,1);
clock_t startGpuCalculation = clock();
MatrixMulKernel <<<dimGrid, dimBlock >>>(d_M, d_N, d_P, Width);
cudaThreadSynchronize();
clock_t endGpuCalculation = clock();

【问题讨论】：

你也可以添加你的内核配置（块和网格大小）吗？
我建议提供一个minimal reproducible example 来比较这两种情况。
当我在 CentOS 7、CUDA 8 和驱动程序 375.66 的 Tesla K20Xm 上运行您的代码并使用 nvprof 计时内核执行时，共享内存的执行时间大约为 36 毫秒内核和非共享内存内核的执行时间大约为 92 毫秒。 Here 是完整的成绩单。所以我无法重现共享内存内核速度较慢的观察结果。
当您edit 生成minimal reproducible example 时，您还应该提及您使用了哪些编译器优化标志。

标签： c++ memory cuda

【解决方案1】：

我在调试模式下运行项目（VS 2017 和 CUDA 9）。我在发布模式下运行代码，共享内存比全局内存快得多。我的错。

【讨论】：