我们之前的帖子跟踪了一个向量相加内核 — c[i] = a[i] + b[i]，每个线程处理一个浮点数 — 从nvcc到底层的warp。我们详细讨论了内核的启动过程，但也遗漏了很多。这一次，我们将解决这些遗漏，并追踪关键的SASS指令（一个全局加载）在硬件中的路径 — 在这种情况下，因为它在我桌子下面，是RTX 4090。我们之所以进行这种逆向工程，至少在原则上是出于性能原因（关于为何这些细节重要的很好的理由，请参见《Citadel微基准测试》论文中的“为什么这些细节很重要”）。对于应用于更具生产相关性的GPU的相同工作，请关注这个空间。NVIDIA对这条路径的详细文档记录很少，至少没有我们希望的那么详细，因此我们将通过在硬件本身上运行时间实验来确定它。我们正在研究的CUDA内核在其函数体内有两行代码：__global__ void vadd（const float* a，const float* b，float* c，int n）{int i = blockIdx.x * blockDim.x + threadIdx.x;如果（i < n）c[i] = a[i] + b[i];} 如果你检查编译后的SASS，你会看到支持这些行的指令：/*0080*/ IMAD.WIDE R4，R6，R7，c [0x0][0x168]；// &b[i] /*00a0*/ LDG.E R4，[R4.64]；// b[i] 它们的作用是将向量b的元素从全局内存加载到寄存器中，在这里它们可以与a的元素相加以执行内核。一个LDG.E请求每个32车道4字节。服务它需要四个32字节的扇区，一个缓存行，一个地址转换，一个交叉开关，三个 L2 切片中的一个，而且当它在所有地方未命中时，需激活并在DRAM芯片上进行四次列读取。我们将尝试追踪指令通过硬件及其返回的过程。场景设定：我们的warp位于一个SM的四个子分区中的一个，旁边还有另外十一条驻留的warp。每个周期，子分区的调度器选择一个符合条件的warp，并一次性在32条通道上发出它的下一条指令。我们的warp赢了两次：一次用于IMAD.WIDE，几周期后（地址现在位于R4和R5）用于LDG。我们的故事从LDG开始。 从warp到L1缓存 让我们从指令开始。LDG.E R4，[R4.64]是从存储在寄存器R4和R5中的64位地址全局加载32位数据。R5出现是因为.64注释：寄存器的大小为32位。将结果存储在寄存器R4中。为了加载数据本身，我们首先必须从这些寄存器中获取地址。寄存器文件的一行同时为所有32个车道持有R4。读取操作首先在操作数收集器中进行分级。分级是为了共享寄存器文件的银行的指令，因为一个银行每个周期只能服务一次读取。共有两个银行，由寄存器号的低位选择，因此相邻的对总是跨越两个银行。另一个银行持有R5。warp读取两个条目，产生256字节的读取，作为32个不同的64位地址，每个车道一个地址。 寄存器读取的成本 地址读取最多增加一个周期。取自寄存器的共享内存加载从发出到首次使用需要24个周期，而地址作为直接值的相同加载需要23个周期。（LDG不能采用直接值）。在解析了所有地址后，指令发往加载/存储单元（LSU）。LSU接受指令及其操作数地址，并进行一些地址计算（如有必要）。该单元可以添加直接偏移量（[R4.64]没有附加偏移量），并且作用域加载（LDG直接命名全局窗口）。然后它发送操作码（‘加载这些地址’，以二进制形式）、活动车道的32位掩码、其计算的地址以及结果所属的寄存器号。下一个目标是合并器。 每个车道的每个LDG.E指令请求4字节，但我们的下一个目标，L1缓存，是以32字节扇区寻址的。合并器的工作是找出服务我们4字节请求所需的最小L1扇区数量。合并器确定它应该发出4个连续扇区请求，以满足warp要求的128字节的内容。 进入L1缓存 请求四个连续的32字节扇区被发往L1缓存。L1缓存的组织单位仍然更小：128字节的行。我们的4个连续扇区代表单个行的4个部分，因此已向L1发出对该缓存行的请求。首先，我们必须确定该行是否已存储在缓存中。缓存被分为称为集合的插槽组。在技术上，4090上的L1缓存是4路组相联的。缓存存储的线路在完全相联（任何缓存行可以存放在缓存的任何地方）和“直接映射”（每个缓存行只能存放在一个地方）之间的连续体上。一个行的地址决定了它属于哪个集合。这张卡上的一个集合包含四个插槽。