减少共享内存库冲突

Posted 2023-04-15

【中文标题】减少共享内存库冲突

【英文标题】：Reducing Shared Memory Bank Conflicts

【发布时间】：2021-01-19 04:21:33

【问题描述】：

Nvprof 报告说我的 sgemm 内核中有大约 2 亿个 shared_ld_bank_conflict 和一些 shared_st_bank_conflict。我尝试了填充技巧__shared__ float smem[SIZE + OFFSET];，它将存储库冲突减少到 0，但负载库冲突仍然存在。我不知道如何进一步改进它。

__global__ void sgemm(
  const float* __restrict__ A,
  const float* __restrict__ B,
  float* __restrict__ C,
  int M, int N, int K
)
  int tid = threadIdx.x;
  int gStartx = blockIdx.x * 128;
  int gStarty = blockIdx.y * 128;

  int dx = tid % 8;
  int dy = tid / 8;
  int vx = tid % 16;
  int vy = tid / 16;

  __shared__ volatile float aSM[8][128+4];
  __shared__ volatile float bSM[8][128+4];
  float aBuffer1[4];
  float bBuffer1[4];
  float aBuffer2[4];
  float bBuffer2[4];

  float cCache[8][8];
#pragma unroll
  for (int i=0; i<8; i++) 
#pragma unroll
    for (int j=0; j<8; j++)
      cCache[i][j] = 0.f;

//load first two tiles
#pragma unroll
  for (int i=0; i<4; i++)
    aBuffer1[i] = A[(gStarty + dy + i*32)*K + (dx)];
    bBuffer1[i] = B[(gStartx + dy + i*32)*K + (dx)];
  
  int nIt = (K + 8 - 1) / 8;
#pragma unroll
  for (int itr=0; itr<nIt; itr++)
    int gStartk = itr * 8;
    int is_odd = itr & 1;
    if (is_odd == 0)
#pragma unroll
      for (int i=0; i<4; i++)
        if (itr != (nIt - 1))
          // prefetch next tiles
          aBuffer2[i] = A[(gStarty + i*32 + dy)*K + (gStartk + 8 + dx)];
          bBuffer2[i] = B[(gStartx + i*32 + dy)*K + (gStartk + 8 + dx)];
        
        //move current tiles to SMEM
        aSM[dx][dy+i*32] = aBuffer1[i];
        bSM[dx][dy+i*32] = bBuffer1[i];
      
     else 
#pragma unroll
      for (int i=0; i<4; i++)
        if (itr != (nIt - 1))
          //prefetch next tiles to another buffer
          aBuffer1[i] = A[(gStarty + i*32 + dy)*K + (gStartk + 8 + dx)];
          bBuffer1[i] = B[(gStartx + i*32 + dy)*K + (gStartk + 8 + dx)];
        
        aSM[dx][dy+i*32] = aBuffer2[i];
        bSM[dx][dy+i*32] = bBuffer2[i];
      
    
    __syncthreads();

    float aCache[8][4];

#pragma unroll
    for (int p=0; p<2; p++)
#pragma unroll
      for (int ki=0; ki<8; ki++)
#pragma unroll 
        for (int mi=0; mi<4; mi++)
          aCache[ki][mi] = aSM[ki][8*vy + 4*p +mi];
        
      

#pragma unroll
      for (int ki=0; ki<8; ki++)
#pragma unroll
        for (int ni=0; ni<8; ni++)
        float b = bSM[ki][8*vx + ni];
#pragma unroll
          for (int mi=0; mi<4; mi++)
            float a = aCache[ki][mi];
            cCache[mi + 4*p][ni] = fma(a, b, cCache[mi + 4*p][ni] );
          
        
      
     
    __syncthreads();
  

#pragma unroll
  for (int i=0; i<8; i++)
    for (int j=0; j<8; j++)
      C[(gStarty + vy*8 + i)*N + (gStartx + vx*8 + j)] = cCache[i][j];
    
  

A (2048x2048) 矩阵主要是行，B (2048x2048) 是主要列，每个块有 256 个线程，每个块计算 C 的 128x128 部分，每个线程计算 8x8x8。 GPU 是 Tesla P100。

【问题讨论】：

例如，float b = bSM[ki][8*vx + ni]; 为您提供来自共享内存的银行冲突负载。考虑一个 warp 中的前 16 个线程。 vx 在这些线程中的范围为 0..15。您将其乘以 8。这将导致 4 路银行冲突。要使用共享内存有效地进行矩阵-矩阵乘法，我推荐programming guide 中给出的示例。当然，如果您对使用快速矩阵矩阵乘法很认真，请使用 CUBLAS。

@RobertCrovella 首先感谢您详细回答我的所有问题！我试图从this paper 实现采样softmax 的修改版本，如果我们说A 和B 是行向量和列向量的集合，那么A 中的每个向量都乘以向量的唯一 子集例如在 B 中，A 的第一行乘以 B 矩阵的第 2、40、800 列，A 的第二行乘以 B 的第 5、80、400 列，以此类推。我不认为它可以用 CUBLAS GEMM 来实现，这就是我尝试在 cuda 中实现它的原因。

@RobertCrovella 有哪些策略可以避免银行冲突？我应该像 this video 中建议的那样使用 float3 而不是 float 吗？

【参考方案1】：

好的，我找到了解决方案：存储到bSM时，在第二维的每32个单词之间插入一个填充词

//bSM[dx][dy+i*32] = bBuffer1[i];
bSM[dx][dy+i*33] = bBuffer1[i]; //we're skipping column 32, 65, 98, 131

阅读bSM[i][j]时，请这样阅读：bSM[i][j/32 + j]

//float b = bSM[ki][8*vx + ni];
float b = bSM[ki][(8*vx) / 32 + 8*vx + ni];
// (8*vx+ni)/32 is the same as (8*vx)/32, since vi is always less than 8

现在它在 tesla p4 上给了我 cublas gemm 55% 的性能