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

Question

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

4

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

谢谢！

- zbh2047

1个回答

网页内容由stack overflow 提供, 点击上面的

可以查看英文原文，
原文链接

- Robert Crovella · Accepted Answer

至少有两个与GPU内存相关的性能指标：延迟和带宽。从延迟角度来看，NVIDIA没有公布这个数字（据我所知），通常的做法是通过仔细的微基准测试来发现它。

从带宽的角度来看，据我所知，NVIDIA也没有公布这个数字（对于L2缓存），但是通过一个相当简单的复制内核测试用例，应该很容易发现它。我们可以通过确保我们的复制内核使用的复制占用空间比发布的L2缓存大小要大得多来估算全局内存的带宽，而我们可以通过将复制占用空间保持较小来估算L2的带宽。

这样的代码（在我看来）相当容易编写：

$ 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倍。