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

Question

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

谢谢你！

Answer 1

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

从带宽的角度来看，据我所知，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

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