L2 缓存的内存操作是否比 NVIDIA GPU 的全局内存快得多？答案

【问题标题】：Is memory operation for L2 cache significantly faster than global memory for NVIDIA GPU?L2 缓存的内存操作是否比 NVIDIA GPU 的全局内存快得多？
【发布时间】：2021-06-29 10:48:36
【问题描述】：

现代 GPU 架构同时具有 L1 缓存和 L2 缓存。众所周知，一级缓存比全局内存快得多。但是，在 CUDA 文档中，L2 缓存的速度不太清楚。我查阅了CUDA文档，但只能发现全局内存操作的延迟大约是300-500个周期，而L1缓存操作只需要大约30个周期。谁能给出二级缓存的速度？这些信息可能非常有用，因为如果 L2 缓存与全局内存相比速度不是很快，编程将不会专注于优化 L2 缓存的使用。如果不同架构速度不同，我只想关注最新的架构，比如NVIDIA Titan RTX 3090（Compute Capability 8.6）或者NVIDIA Telsa V100（Compute Capability 7.0）。

谢谢！

【问题讨论】：

标签： cuda gpu nvidia

【解决方案1】：

在讨论 GPU 内存时，至少有两个常用的品质因数：延迟和带宽。从延迟的角度来看，这个数字不是由 NVIDIA 发布的（据我所知），通常的做法是仔细发现它microbenchmarking。

从带宽的角度来看，AFAIK 这个数字也不是由 NVIDIA 发布的（用于 L2 缓存），但通过一个相当简单的复制内核测试用例应该很容易发现它。我们可以简单地通过确保我们的副本内核使用的副本占用空间远大于已发布的 L2 缓存大小（V100 为 6MB）来估计全局内存的带宽，而我们可以通过保持我们的副本占用空间小于那个。

这样的代码 (IMO) 编写起来相当简单：

$ cat t44.cu
template <typename T>

__global__ void k(volatile T * __restrict__ d1, volatile T * __restrict__ d2, const int loops, const int ds){

  for (int i = 0; i < loops; i++)
    for (int j = threadIdx.x+blockDim.x*blockIdx.x; j < ds; j += gridDim.x*blockDim.x)
      if (i&1) d1[j] = d2[j];
      else d2[j] = d1[j];
}
const int dsize = 1048576*128;
const int iter = 64;
int main(){

  int *d;
  cudaMalloc(&d, dsize);
  // case 1: 32MB copy, should exceed L2 cache on V100
  int csize = 1048576*8;
  k<<<80*2, 1024>>>(d, d+csize, iter, csize);
  // case 2: 2MB copy, should fit in L2 cache on V100
  csize = 1048576/2;
  k<<<80*2, 1024>>>(d, d+csize, iter, csize);
  cudaDeviceSynchronize();
}

$ nvcc -o t44 t44.cu
$ nvprof ./t44
==53310== NVPROF is profiling process 53310, command: ./t44
==53310== Profiling application: ./t44
==53310== Profiling result:
            Type  Time(%)      Time     Calls       Avg       Min       Max  Name
 GPU activities:  100.00%  6.9032ms         2  3.4516ms  123.39us  6.7798ms  void k<int>(int volatile *, int volatile *, int, int)
      API calls:   89.47%  263.86ms         1  263.86ms  263.86ms  263.86ms  cudaMalloc
                    4.45%  13.111ms         8  1.6388ms  942.75us  2.2322ms  cuDeviceTotalMem
                    3.37%  9.9523ms       808  12.317us     186ns  725.86us  cuDeviceGetAttribute
                    2.34%  6.9006ms         1  6.9006ms  6.9006ms  6.9006ms  cudaDeviceSynchronize
                    0.33%  985.49us         8  123.19us  85.864us  180.73us  cuDeviceGetName
                    0.01%  42.668us         8  5.3330us  1.8710us  22.553us  cuDeviceGetPCIBusId
                    0.01%  34.281us         2  17.140us  6.2880us  27.993us  cudaLaunchKernel
                    0.00%  8.0290us        16     501ns     256ns  1.7980us  cuDeviceGet
                    0.00%  3.4000us         8     425ns     217ns     876ns  cuDeviceGetUuid
                    0.00%  3.3970us         3  1.1320us     652ns  2.0020us  cuDeviceGetCount
$

根据分析器输出，我们可以将全局内存带宽估计为：

2*64*32MB/6.78ms = 604GB/s

我们可以将 L2 带宽估计为：

2*64*2MB/123us   = 2.08TB/s

这两个都是粗略的测量（我在这里没有做仔细的基准测试），但是这个 V100 GPU 上的bandwidthTest 报告的设备内存带宽约为 700GB/s，所以我相信 600GB/s 的数字是“在球场”。如果我们用它来判断 L2 缓存的测量结果在大致范围内，那么我们可能会猜测 L2 缓存在某些情况下可能比全局内存快约 3-4 倍。

【讨论】：

据NVIDIA 报道：“A100 L2 缓存采用新的分区交叉开关结构，提供的 L2 缓存读取带宽是 V100 的 2.3 倍。”
Zhe Jia、Marco Maggioni、Benjamin Staiger、Daniele P. Scarpazza，“通过微基准测试剖析 NVIDIA Volta GPU 架构”，Citadel 技术报告，2018 年 4 月，states NVIDIA 的 L2 缓存负载带宽V100 为 2155 GB/秒。