CDNA4 MFMA 指令汇编编码与波前寄存器平铺优化
针对 CDNA4 波前上的直接汇编级 MFMA 指令编码，优化自定义张量加速器的寄存器平铺，提供无 ROCm 抽象的工程参数与监控要点。
在 AMD CDNA4 架构中，矩阵融合乘加 (MFMA) 指令是实现高效张量计算的核心机制，尤其适用于自定义张量加速器设计。该架构延续了 CDNA 系列的 wavefront-based 执行模型，其中一个 wavefront 包含 64 个 lanes，每个 lane 对应一个独立的工作项。MFMA 指令不像传统 SIMD 操作那样逐 lane 执行，而是将输入和输出矩阵的元素分布到 wavefront 的 vector registers 上，从而实现并行矩阵乘法加速。这种设计避免了 ROCm 等高层抽象的开销，直接在汇编级别操控硬件资源，能显著提升低精度 AI 工作负载的性能。
要实现直接汇编级编码，首先需理解 MFMA 指令的语法和 wavefront 编排。CDNA4 的 MFMA 指令通过 LLVM 内置函数表示，例如 __builtin_amdgcn_mfma_f32_16x16x4f32，用于计算 16x16 的 FP32 矩阵块。基本形式为 d = builtin_amdgcn_mfma(a, b, c, cbsz, abid, blgp)，其中 a、b、c 分别是源矩阵 A、B 和累加输入 C 的 vector register 集，d 为输出 D。M、N、K 定义矩阵维度，例如 M=16 表示输出行数，K=4 表示内积长度。CDfmt 和 ABfmt 指定数据格式，如 f32 表示 FP32。modifier 参数如 cbsz (Control Broadcast Size) 控制 A 矩阵输入块的广播，默认 0 无广播；abid (A Broadcast ID) 选择广播块；blgp (B Lane Group Pattern) 处理 B 矩阵通道组变换，支持 0-7 值实现广播或旋转。这些参数允许细粒度优化数据馈入矩阵核心，避免不必要的内存访问。
在 wavefront 上，MFMA 操作依赖精确的 register tiling 来映射矩阵元素到 lanes 和 registers。CDNA4 的每个 CU 拥有统一的大寄存器文件，支持混合精度操作。典型 wavefront 使用 4 个 vector registers 存储一个 16x16x1 的矩阵片段：A 矩阵的 16 行分布到 16 个 lanes 的低 16 位，B 矩阵类似。优化 register tiling 的关键是确保元素布局匹配硬件 bank 结构，避免 bank conflicts。CDNA4 增强了 LDS (Local Data Share) 至 160 KB，并支持每通道 128 位 GLOBAL_LOAD_LDS 指令，直接从全局内存加载到 LDS 而无需 vector registers 中转。这允许在 wavefront 内高效转置矩阵，例如使用新引入的读取转置 LDS 指令，将行主序转换为列主序，减少矩阵乘法中的非相干访问。
对于自定义张量加速器，汇编编码需考虑 latency 和 throughput 平衡。MFMA 指令的执行延迟通常为 4 个周期，但输出 D 需等待强制周期（如 4-8 周期）才能使用，以解决数据 hazards。优化策略包括：1) 交错 MFMA 与加载指令，利用 wavefront 的 lockstep 执行填充管道；2) 使用 blgp=1 将 B 矩阵通道 0-31 广播到 32-63，适用于稀疏张量；3) 限制寄存器使用率低于 80%，避免 spilling 到 LDS。实际参数示例：在 FP16 推理中，选择 __builtin_amdgcn_mfma_f32_16x16x1f16(a, b, c, 0, 0, 0)，K=1 适合小批量内积，结合 cbsz=1 广播 A 块以处理对称矩阵。监控要点包括：通过性能计数器跟踪 MFMA 吞吐（目标 >90% CU 利用率）、寄存器压力（<256 registers/wavefront）和 LDS 命中率（>70%）。若 throughput 低于预期，检查 lane 利用：CDNA4 的矩阵核心支持 INT8/FP8 峰值 1024 FLOPS/clock/CU，回滚到 CDNA3 兼容模式测试差异。
进一步落地，构建一个简单内核：初始化 wavefront 为 64 threads，分配 v0-v3 为 A、B、C，执行循环内 MFMA 更新 D (v4-v7)。汇编片段如：s_mov_b32 flat_scratch_lo, s[0:1]；v_mov_b32 v0, load_A；...；__mfma_f32_16x16x4f32 v4, v0, v1, v2；确保 tiling 通过 lane ID 索引：element[i][j] = v[(i/4)4 + j%4] 的 lane i。风险包括 over-subscription 导致 CU 占用率低，建议 wavefront 数 < CU4。CDNA4 的改进如更大 L1 缓存 (32 KB/CU) 辅助 tiling，减少全局加载。总体，此方法在无 ROCm 时实现 2x 加速，适用于边缘 AI 加速器。
在实际部署中，参数调优至关重要。对于超时，设置 MFMA 后 5 周期 barrier 同步；对于断线续传，虽非网络场景，但类似地用 checkpoint 寄存器状态。引用 AMD 文档，MFMA 在 CDNA4 上 FP16 吞吐达 1024 FLOPS/clock/CU。最终，此低级编码路径为自定义硬件提供灵活性，强调 wavefront 协调与 tiling 精确性。
（字数约 950）