AMD ROCm：4波交错优化FP8 GEMM内核，为Instinct MI355X（CDNA 4）每波提供512个向量寄存器
AMD ROCm博客介绍了针对搭载CDNA 4架构的Instinct MI355X加速器的FP8 GEMM（混合精度矩阵乘法）内核的全新4波交错方法。该优化将布局从8波乒乓方案改为每波拥有完整512个VGPR寄存器的4波设计，通过基于XOR的swizzle消除LDS内存冲突，并通过精确交错MFMA指令与数据加载来隐藏内存延迟。
本文由人工智能基于一手来源生成。
AMD ROCm团队发布了针对采用CDNA 4架构的AMD Instinct MI355X图形处理器上FP8 GEMM（通用矩阵乘法）内核的全新优化策略详细技术说明。该优化超越了标准的8波乒乓方案，引入4波交错设计，大幅提升每波的寄存器预算，并在硬件层面消除内存冲突。
为何4波方案在FP8 GEMM上优于8波方案？
关键差异在于VGPR（向量通用寄存器）的分配。在8波设计中，「每个SIMD单元两个波」意味着每个波必须共享寄存器空间，限制了可处理的输出块大小。4波方案每个SIMD单元分配一个波，赋予其访问完整512个VGPR寄存器的权限。
这直接影响可处理的分块大小：4波方案可使用128×128元素的输出分块，而8波设计仅为64×128。整个内核的全局分块为256×256×128，四个波中的每个使用每块内的16×16 MFMA（矩阵融合乘加）操作处理四个64×64元素块。
4波交错如何在MI355X上消除内存延迟？
GPU矩阵乘法内核的核心挑战是隐藏内存延迟——GPU在等待从内存加载数据时必须保持MFMA计算忙碌。4波方案通过在子分块粒度级别精确交错指令来解决这一问题，无需依赖硬件调度器。
实现使用__builtin_amdgcn_sched_barrier(0)强制执行指令顺序——禁止编译器将所有内存操作分组到计算操作之后。每次interleaved_block调用将16条MFMA指令与8次LDS到寄存器的加载（ds_read_b128）重叠，实现一致的覆盖，无需乒乓同步的复杂性。
基于XOR的LDS swizzle如何消除存储体冲突？
MI355X的LDS（本地数据共享）内存组织为64个存储体，每个4字节，周期为256字节。问题具有结构性：16×16×128 MFMA指令天然按列访问矩阵（column-major），在标准布局下会导致8倍存储体冲突——八次访问同时等待同一存储体。
AMD通过偏移量的数学变换解决了这个问题：
row_bits = (offset % (16 * 128) >> 7) / 2
mask = row_bits << 4
swizzled_offset = offset ^ maskXOR操作重映射地址，使所有四个阶段的ds_read_b128指令均无任何存储体冲突。结果是获得LDS内存的完整带宽，而非8倍冲突情况下有效的12.5%。
硬件平台与芯片感知优化
所有优化针对采用CDNA 4架构的AMD Instinct MI355X。MI355X具有8芯片（XCD）拓扑——每个XCD拥有自己的L2缓存，末级缓存在XCD之间共享。ROCm团队添加了芯片感知网格swizzle，使用类莫顿曲线映射重新映射线程块到线程的分配，最大化单个XCD内L2的利用率以及XCD之间的末级缓存利用率。
基准测试使用了1000次热身和1000次基准迭代，在正态分布数据和轮换缓冲区上进行。ROCm团队表示「在不同ROCm版本间性能一致」，无需手动pragma调优——这与对编译器变化敏感的8波变体形成对比。
常见问题
- AMD CDNA 4的4波交错FP8 GEMM优化是什么？
- 这是针对AMD Instinct MI355X上FP8精度矩阵乘法（GEMM）的GPU内核优化。与共享寄存器空间的8波设计不同，它使用4波，每波拥有完整的512个VGPR寄存器，从而支持更大的输出块（128×128对比64×128）并更好地隐藏内存延迟。
- AMD为何在这些GEMM内核中使用基于XOR的swizzle？
- MI355X的LDS内存有64个存储体，而16×16×128 MFMA指令天然按列访问（column-major），会导致8倍存储体冲突。XOR swizzle重映射偏移量可消除冲突，确保四个阶段的ds_read_b128访问均无冲突。
- 这些优化针对哪款GPU硬件？
- 优化针对采用CDNA 4架构、具有8芯片（XCD）拓扑的AMD Instinct MI355X。除内核优化外，还包括芯片感知网格swizzle，通过类莫顿曲线映射最大化每个XCD的L2缓存利用率以及XCD间的末级缓存利用率。