当GPU读取内存时会发生什么

# GPU 读取内存时会发生什么？ | Doubleword 来源：https://blog.doubleword.ai/what-happens-when-a-gpu-reads-memory 我们之前的文章（https://fergusfinn.com/blog/what-happens-when-you-run-a-gpu-kernel/）跟踪了一个向量加法内核——`c[i] = a[i] + b[i]`，每个线程处理一个浮点数——从`nvcc`编译到线程束（warp）执行的全过程。我们详细讨论了内核的启动过程，但也留下了许多未解之谜。这次，我们将弥补之前的不足，追踪关键SASS指令（一次全局加载）在硬件中的路径——具体来说，由于硬件就在我们桌子下，这里指的是RTX 4090（https://www.nvidia.com/en-gb/geforce/graphics-cards/40-series/rtx-4090/）。 我们进行此类逆向工程是出于性能优化的目的（至少理论上如此），详情可参考《Citadel微基准测试》（https://arxiv.org/pdf/1804.06826）论文中“为何这些细节重要”部分。对于更多生产级相关GPU的类似分析，请持续关注本专栏。NVIDIA并未详细记录这条路径的大部分细节，至少未达到我们期望的深度，因此我们将通过在实际硬件上运行计时实验来确定其机制。 我们研究的CUDA内核函数体包含以下两行代码： ```c __global__ void vadd(const float* a, const float* b, float* c, int n) { int i = blockIdx.x * blockDim.x + threadIdx.x; if (i < n) c[i] = a[i] + b[i]; } ``` 如果你查看编译后的SASS指令，会看到驱动这些代码的底层指令： ```assembly /*0080*/ IMAD.WIDE R4, R6, R7, c[0x0][0x168] ; // &b[i] /*00a0*/ LDG.E R4, [R4.64] ; // b[i] ``` 这些指令用于将向量`b`（我们以`b`为例，但对`a`的指令相同）的元素从全局内存加载到寄存器中，以便与`a`的元素相加完成内核计算。一条`LDG.E`指令要求32个通道各读取4字节数据。为满足此请求，需要：4个32字节扇区、1个缓存行、1次地址转换、1次交叉开关传输、36个L2切片之一，以及当所有缓存均未命中时，在DRAM芯片上执行1次行激活和4次列读取。我们将尝试追踪这条指令在硬件中的完整往返路径。 背景设定：我们的线程束驻留在SM的四个**子分区**之一中，与其他11个常驻线程束共处。每个周期，子分区的调度器会挑选一个就绪的线程束，并一次性向32个通道发出其下一条指令。我们的线程束两次赢得调度：第一次是执行`IMAD.WIDE`，几周期后（地址已存入`R4`和`R5`）执行`LDG`。我们的故事从`LDG`指令开始。 ## 从线程束到L1缓存（https://blog.doubleword.ai/what-happens-when-a-gpu-reads-memory#from-the-warp-to-the-l1-cache） 让我们从指令本身开始。`LDG.E R4, [R4.64]`是一条全局加载指令，从寄存器`R4`和`R5`（因标注`.64`而出现`R5`；寄存器为32位宽）存储的64位*地址*中读取32位数据，并将结果存入寄存器`R4`。要加载数据本身，我们首先需要从这些寄存器中获取地址。寄存器文件的一行包含所有32个通道的`R4`值（读取操作首先在**操作数收集器**中暂存；暂存机制用于处理源操作数共享寄存器文件组（bank）的指令，因为每个组每周期只能服务一次读取。共有两个组，由寄存器编号的最低位决定，因此相邻寄存器对总是跨两个组。）另一行包含`R5`值。线程束读取这两个条目，得到32个不同的64位地址（每个通道一个地址），总计读取256字节。 **寄存器检索的开销** 地址读取最多增加1个周期。共享内存加载操作若从寄存器获取地址，从指令发出到首次使用需要24个周期；若地址为立即数则需23个周期。（`LDG`指令不支持立即数地址。）所有地址解析完毕后，指令被发送到**加载/存储单元**（LSU）。LSU接收指令及其操作数地址，执行必要的地址运算（此单元可添加立即数偏移量，但`[R4.64]`无偏移量需添加），并处理作用域加载（`LDG`直接指定全局内存窗口），随后将操作码（二进制形式的“加载这些地址”）、32位有效通道掩码、计算得到的地址以及结果目标寄存器编号发送出去。 下一个目标是**合并器**。每个通道的`LDG.E`指令请求4字节数据，但我们的下一个目标——**L1缓存**——以32字节**扇区**为单位进行寻址。合并器的任务是确定满足这4字节请求所需的最少L1扇区数量。合并器判断应发出4个连续的扇区请求，以覆盖线程束请求的128字节数据¹。 ### 进入L1缓存（https://blog.doubleword.ai/what-happens-when-a-gpu-reads-memory#entering-the-l1-cache） 四个连续32字节扇区的请求被发送到**L1缓存**。L1缓存的组织单元更为粗粒度：128字节的**缓存行**。我们的4个连续扇区对应单个缓存行的4个部分，因此L1会为该缓存行发出一次请求。 首先，我们必须确定该缓存行是否已存在于缓存中。缓存被分为称为**组**（set）的槽位组（技术上，4090的L1缓存是4路组相联的。缓存介于全相联（任何缓存行可存储在缓存中任意位置）和“直接映射”（每行只能存储在一个位置）之间。），一行的地址决定其所属的组。该显卡上一个组包含4个槽位²，每个槽位携带一个**标签**（tag）标识其中的缓存行。查找操作会将所有四个标签与所需行的标签进行比较。使用的地址是程序中使用的**虚拟**地址³（大概是为了避免L1命中时的转换开销）。缓存行落入的组由其虚拟地址通过哈希方案生成⁴（这是一种复杂的奇偶校验方案（见附录），而非简单的比特切片，以防止2的幂次跨步访问（如矩阵、张量等的列）反复命中相同组导致颠簸。），该方案可被逆向工程⁵。 如果四个标签中有一个匹配，且所需扇区在该槽位中，数据将被读出，加载完成⁶。由于我们首次加载所有数据，请求未命中，必须进一步深入内存系统。 **L1命中的开销** L1命中约需15.4纳秒（40个周期）。该数据来自一个线程通过L1常驻行的随机排列跟踪依赖链的实验⁷。 ## 寻找L2：地址转换（https://blog.doubleword.ai/what-happens-when-a-gpu-reads-memory#looking-for-l2-translation） **虚拟内存**在程序命名的地址与硬件存储数据的地址之间引入了一层间接寻址。程序获得自己的连续空间，硬件则将该空间跨物理页面随意布局。**转换**（Translation）是两者之间的映射。我们刚讨论的L1使用**虚拟寻址**，因此无需关心转换。超过此点后，我们必须开始使用硬件的语言——L1未命中的请求在离开SM前必须经过转换⁸。 物理地址与虚拟地址的实际映射关系在驱动程序分配内存时建立：当分配`b`时，驱动程序为其选择物理（2MiB）页面，并将页表写入VRAM以记录分配⁹。转换单元根据这些表将虚拟地址转换为物理地址。SM在其TLB（转换后备缓冲器）中保存最近十六次转换结果，供所有线程束共享¹⁰。首次加载必然在TLB中未命中。 **转换开销** 我们在任何探测中都未观察到TLB命中的开销。未命中成本约为4.4纳秒（11个周期）。无论页面大小或来自哪个SM，相同的重填成本误差在0.1纳秒以内，表明下一级转换缓存是全局共享且开销极低的。 转换完成后，每个128字节缓存行发出一个请求：包含该行的**物理地址**，以及所需扇区的掩码。我们的请求是一个单独的请求，标记了所有四个扇区¹¹。请求通过SM离开，经由**交叉开关**（crossbar）到达**L2缓存**。 ## 迷失在L2中（https://blog.doubleword.ai/what-happens-when-a-gpu-reads-memory#lost-in-l2） 请求通过交叉开关到达36个2MiB L2切片之一，选择由其物理地址的一个相对复杂的函数决定¹²。任何SM可命中任何切片。所有切片可并行服务，因此总带宽是单个切片的36倍。 在切片内部，结构与L1类似。每个切片包含1024个**组**。缓存行所属的组由其物理地址的哈希决定。每个组现在包含16个槽位：切片是**16路组相联**的¹³。缓存行大小为128字节，与L1相同。该行不在L2中，因为我们之前从未获取过它（这可能是艺术化处理：通过PCIe从主机内存加载`b`向量可能已缓存在L2中，但那样我们就无法继续深入到DRAM了！）。 每个切片回退到12个内存控制器之一——每个控制器对应3个切片。每个内存控制器负责与单个**GDDR6X**（https://en.wikipedia.org/wiki/GDDR_SDRAM）DRAM芯片通信¹⁴。我们的请求被移交到该控制器。 **L2命中的开销** L2命中约需127纳秒（330个周期）。每个SM每周期可向交叉开关发送最多两个缓存行请求，36个切片独立服务。出口端口计数器为`l1tex__m_l1tex2xbar_req_cycles_active`。 ## 在DRAM中找到（https://blog.doubleword.ai/what-happens-when-a-gpu-reads-memory#found-in-dram） 内存控制器的任务是从其2GiB DRAM芯片加载数据。它通过总线向DRAM发出命令来实现。DRAM被分为两个独立总线，由控制器独立驱动，称为**通道**（channel）。每个通道上有16个**bank**：二维存储单元阵列。一个bank包含65,536个**行**。硬件一次可打开一行（**激活**，开销大），然后从该行返回任意32字节**列**（**读取**，行打开期间开销小）。 GDDR6X结构： * 通道0 · 1 GiB * 通道1 · 1 GiB * 一个bank · 65,536行，每行1 KiB * 一行 · 1 KiB · 32列，每列32字节 地址被最终分解以匹配此内存结构。它选择通道、bank、行和列。我们的四个扇区是同一行的四个列¹⁵。因此，为服务我们的加载请求，内存控制器必须先发送一个**激活**命令，然后发送四个**读取**命令¹⁶。 DRAM芯片如何响应这些命令？每个DRAM单元是一个电容器加一个晶体管。一行的晶体管共享一条**字线**（wordline），连接至其栅极。每个晶体管位于其电容器和**位线**（bitline）之间，位线沿列延伸，提供从每个单元（与其他行单元共享）到**感应放大器**（sense amplifiers）的路径。比特存储在电容器的电荷状态中。电容器持续漏电，因此芯片必须不时暂停每个bank以补充电荷。 DRAM结构：点击一行可模拟行解码器，从电容器释放电荷到位线并存入行缓冲器。 位线 | 行解码器 | 字线 激活命令触发**行解码器**驱动该行的字线，打开该行的晶体管，并将该行（仅该行）电容器中的电荷通过位线驱入感应放大器。感应放大器将电荷放大为全摆幅比特并保持，供控制器读取。当发出读取命令时，其列地址选择256个这样的行比特。 从感应放大器读取一次可获得大量比特，但我们需要将它们序列化到驱动数据回传总线的引脚上。每个通道有16个数据引脚。我们读取的256比特通过这些引脚作为PAM4（https://fergusfinn.com/blog/nvlink-scale-up/#fast-or-far）（GDDR6X是GDDR6标准，增加了PAM4信令）符号传输：每个符号是四种电压电平之一，携带两个比特，因此16个引脚传输16比特——即8个符号。时钟通过共享线路发送，以便控制器在正确的边沿采样。 ## 归途（https://blog.doubleword.ai/what-happens-when-a-gpu-reads-memory#the-way-back） 这些PAM4突发数据在内存控制器中被反序列化，并写入L2切片的缓存行。数据通过交叉开关返回其SM，并填充L1槽位。它们与离开时留下的记录汇合，其字节被写入寄存器`R4`的所有通道中。当加载指令发出时，已设置一个依赖屏障，寄存器写入操作清除了该屏障。线程束再次变为就绪状态，在调度器的下一个周期赢得仲裁。它发出的指令是等待`b[i]`结果的加法操作。 整个往返过程——L1、TLB、交叉开关、L2、控制器并返回——耗时约255纳秒（660个周期）。在此期间，我们的线程束一直停在其屏障上。然而芯片的其余部分并未闲置。子分区为其他11个线程束发出了相同的加载指令，该SM的其他部分为另外36个线程束服务，其他SM则为剩余的6096个线程束服务。结果是一片加载的喧嚣，任何单个加载的延迟都淹没在噪声中。 以下是模拟结果：仅模拟`vadd`内核中对应加载`b`的指令执行，时间成比例。每个SM仅加载其在真实内核中加载的地址：这些地址在正确的L1组中亮起（并未命中），然后通过交叉开关路由到正确的L2切片，在那里未命中，经过正确争用的内存控制器落到模拟的DRAM bank，然后通过L2、L1返回，并将其结果写入正确的寄存器。 SM（128个），每个L1组一个像素 交叉开关 L2（36切片，每个控制器3个），每个组一个像素 内存控制器（电平表示瞬时吞吐量） GDDR6X，12个芯片，32个bank 激活 | 行打开 | 预充电 | 刷新 | 进行中0 | 已完成0 | 激活0 | 刷新0 | GB/s 0 ## 附录：探测方法（https://blog.doubleword.ai/what-happens-when-a-gpu-reads-memory#appendix-the-probes） ### 设置（https://blog.doubleword.ai/what-happens-when-a-gpu-reads-memory#setup） 所有测量均在一台RTX 4090（`sm_89`）上进行，核心时钟锁定在2.6 GHz。周期数通过该频率下的纳秒测量值换算得出。两个主要测量工具如下：