如何评估GPU显存带宽

Question

我想评估我的 GPU 的内存带宽。我编写了一个片段，将

int8_t

数组复制到另一个数组并测量时间。我设置了一个包含 64 个线程的块和一个由

(<size of my array>/<block size>/<data size a thread needs to copy>)

块组成的网格。我用

cudaEvent

和

chrono::high_resolution

来测量时间。结果显示，它只能达到几十GB/s，与Nvidia给出的4000GB/s带宽相差甚远。那么，问题是为什么我的计划不起作用以及如何解决？

我的硬件是 Nvidia H20 和 Intel Xeon 芯片。操作系统是CentOS。

我的核函数如下：

 __global__ void copyGpuMem(int8_t *d_B,int8_t *d_A,size_t oneThreadDataSize){
  int tid=threadIdx.x;
  int bid=blockIdx.x;

  size_t off=(bid*blockDim.x+tid)*oneThreadDataSize;
  
  for(int i=0;i<oneThreadDataSize;i++){
    d_B[off+i]=d_A[off+i];
  }
}

我的启动内核代码如下：

int8_t *d_A;
int8_t *d_B;
size_t blockSize=64;
const size_t oneThreadDataSize=1e3; 
const size_t oneBlockDataSize=blockSize*oneThreadDataSize;
size_t n=(1e9+oneBlockDataSize-1)/oneBlockDataSize*oneBlockDataSize;
int loop=100;
size_t gridSize=n/oneBlockDataSize;
CHECK_CUDA(cudaMalloc(d_A,n*sizeof(int8_t)));
CHECK_CUDA(cudaMalloc(d_B,n*sizeof(int8_t)));

auto start=std::chrono::high_resolution_clock::now();


for(int i=0;i<loop;i++){
  copyGpuMem<<<int(gridSize),int(blockSize)>>>(d_B,d_A,oneThreadDataSize);
}
auto end=std::chrono::high_resolution_clock::now();
double bw=(2*n*loop)/(start-end)/1e9 //unit GB/s

我的结果带宽约为 200 GB/s，但给定带宽为 4000 GB/s。我对此一无所知。

我尝试降低每个线程需要处理的数据大小。它将结果从数十 GB/s 提高到最多 200 GB/s。是的，这是事实。当我设置

oneThreadDataSize=1e6

时，结果约为10-20 GB/s。但是当我继续减小尺寸时，效果就不太好了。

Answer 1

来自评论：

您的内存访问模式未正确合并。请查看最佳实践指南中的相关部分。此外，为了获得最佳的每线程性能，请使用

int4

或其他大型对齐数据类型。一次至少一个

int

。

后续问题：

我遵循@Homer512的建议并尝试使用uint4而不是int8_t来测量带宽。我得到的结果接近 3200 GB/s，更接近官方规格。是不是因为gpu一次性把16个字节一起传输给核心？

不，当前 CUDA 硬件上的内存传输以 32 字节事务进行。然而，通过 1 字节非合并访问，您可以让系统对 1 字节使用的有效负载执行 32 字节传输。现在使用了 32 字节中的 16 个。它还可能对缓存命中率产生积极影响，因为该事务的后半部分仍然被拉入 L2 缓存，并且在下次访问发生时可能仍然存在。

我还有一个问题：L1/L2 是否会影响我的测量？

是的，这会有所帮助，但数据集的大小超出了缓存大小。这将限制其有效性。只需使用视觉配置文件Nsight Compute。它会告诉您带宽和缓存命中率。

我如何编码来实现合并内存访问。

最佳实践指南详细解释了这一点，但它相当简单：一个扭曲（32 个线程的组）充当一个。内存访问发生在 32 字节事务中。在一次内存操作中可以访问的最大类型是 16 字节的

int4

和

float

。因此，您要么需要两个线程组访问两个相邻的

int4

，要么需要 32 个线程访问单个字节，或者介于两者之间。

通常，您只需让所有 32 个线程访问一行中的 32 个元素，无论它们是什么类型。内存对齐也发挥了作用，但如果您一次处理整个数组，这通常是给定的。

这是应用于您的代码的基本模式：

__global__ void copyGpuMem(int8_t *d_B,int8_t *d_A,size_t oneThreadDataSize){
  int tid=threadIdx.x;
  int bid=blockIdx.x;

  size_t off=bid*blockDim.x+tid;
  size_t stride = blockDim.x * gridDim.x;
  
  for(size_t i=0;i<oneThreadDataSize;i++){
    d_B[off+i*stride]=d_A[off+i*stride];
  }
}

注意在每次循环迭代中，相邻线程如何访问相邻元素，偏移量是线程

tid

。当然现在有点尴尬，因为那个

oneThreadDataSize

实在是太不方便了。通常，您会像这样编写这些类型的循环：

__global__ void copyGpuMem(int8_t *d_B, int8_t *d_A, size_t fullArraySize){
  unsigned tid=threadIdx.x;
  unsigned bid=blockIdx.x;

  size_t off=bid*blockDim.x+tid;
  size_t stride = blockDim.x * gridDim.x;
  
  for(size_t i=off; i<fullArraySize; i+=stride){
    d_B[i]=d_A[i];
  }
}

注意最后一个函数参数含义的变化。

对于单个线程来说，一次一个字节的吞吐量仍然非常低。您可以将其与更大的矢量类型结合起来。更通用的

memcpy

风格的内核可能如下所示：

#include <stdint.h>

__global__ void copyGpuMem(
      int8_t *d_B, const int8_t *d_A, size_t fullArraySize) {

  size_t off = blockIdx.x*blockDim.x+threadIdx.x;
  size_t stride = blockDim.x * gridDim.x;
  uintptr_t bAddress = reinterpret_cast<uintptr_t>(d_B);
  uintptr_t aAddress = reinterpret_cast<uintptr_t>(d_A);
  if(! (bAddress % alignof(int4) || aAddress % alignof(int4))) {
    /*
     * Likely case: Addresses are 16 byte aligned for int4 access
     */
    const size_t bytes_at_once = sizeof(int4) / sizeof(int8_t);
    size_t int4ArraySize = fullArraySize / bytes_at_once;
    int4* d_B_i4 = reinterpret_cast<int4*>(d_B);
    const int4* d_A_i4 = reinterpret_cast<const int4*>(d_A);
    /* Copy 16 byte at once */
    for(size_t i=off; i<int4ArraySize; i+=stride){
      d_B_i4[i] = d_A_i4[i];
    }
    /*
     * If the full size is not divisible by 16, there is a tail.
     * We adjust the offset so that the following loop can deal with it
     */
    off += int4ArraySize * bytes_at_once;
  }
  /*
   * Either deals with unaligned arrays or with the last few entries
   * in an aligned array case
   */
  for(size_t i=off; i<fullArraySize; i+=stride){
    d_B[i] = d_A[i];
  }
}