我们上一篇文章追踪了一个向量加法内核从 nvcc 到 warp 的过程。这一次,我们跟随一条关键的 SASS 指令(全局加载)穿越 RTX4090 的硬件。所研究的 CUDA 内核只有两行:__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 显示:/*0080*/ IMAD.WIDE R4, R6, R7, c[0x0][0x168] ; // &b[i] /*00a0*/ LDG.E R4, [R4.64] ; // b[i]。这些指令将向量 b 的元素从全局内存加载到寄存器中。
一条 LDG.E 在 32 个通道中每个通道请求四个字节。为它提供服务需要四个 32 字节的扇区、一条缓存行、一次地址转换、一次交叉开关穿越、三十六个 L2 切片之一,以及当处处未命中时,DRAM 芯片上的一次激活和四次列读取。我们的 warp 位于 SM 四个子分区之一上,与其他十一个常驻 warp 并存。该指令发射到加载/存储单元(LSU),后者将操作码、活动通道掩码、计算出的地址和寄存器编号发送给合并器。
合并器确定 warp 请求的 128 字节需要 4 个连续的扇区请求。进入 L1 缓存后,请求被发送到虚拟寻址的 4 路组相联缓存。查找在四个槽位的组中比较标签。由于这是首次加载,它未命中并继续向下。L1 命中大约在 15.4 纳秒内返回——40 个周期。虚拟内存需要在离开 SM 之前进行转换。SM 的 TLB 保存着十六个最近的转换。未命中的开销约为 4.4 纳秒——十一个周期。
转换之后,每个 128 字节行的一个请求穿过交叉开关,前往由其物理地址的函数选定的 36 个 2 MiB L2 切片之一。所有切片可以并行服务,使总带宽达到单个切片的 36 倍。
来源: Hacker News · 由HeadlinesBriefing整理摘要