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

Question

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

谢谢！

Answer 1

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

从带宽的角度来看，AFAIK这个数字也没有由NVIDIA发布(用于二级缓存(，但通过一个相当简单的复制内核测试用例应该很容易发现它。我们可以通过确保拷贝内核使用的拷贝占用空间远大于已发布的二级缓存大小(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的数字是"0"；在球场上"；。如果我们用它来判断二级缓存的测量值是大致的，那么我们可能会猜测，在某些情况下，二级缓存可能比全局内存快3-4倍。