二级缓存的内存操作是否明显快于 NVIDIA GPU 的全局内存？

Question

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

谢谢！

Answer 1

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

发现它

从带宽的角度来看，据我所知，这个数字也不是 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 倍。

二级缓存的内存操作是否明显快于 NVIDIA GPU 的全局内存？

Is memory operation for L2 cache significantly faster than global memory for NVIDIA GPU?

cuda

gpu

nvidia