文章摘要
这篇文章介绍了如何在AMD CDNA3和CDNA4架构上使用HIP内核编程矩阵核心，重点讲解了低精度数据类型(FP16/FP8/FP4)和新型矩阵核心指令。通过代码示例说明矩阵乘法运算原理、编译器内置函数及数据布局要求，帮助开发者利用矩阵核心加速AI和HPC工作负载中的矩阵运算。
文章总结
AMD CDNA3与CDNA4架构中的矩阵核心编程
核心内容概述
本文详细介绍了如何在AMD CDNA3和CDNA4架构上利用HIP内核编程矩阵核心（Matrix Cores），重点涵盖低精度数据类型（如FP16、FP8、FP4）和CDNA4新增的指数块缩放指令。通过代码示例和图示，文章提供了编程矩阵核心所需的关键知识，包括现代低精度浮点类型、矩阵核心编译器内置函数以及指令要求的数据布局。
1. 矩阵核心简介
矩阵乘法是AI和HPC工作负载的核心。AMD CDNA架构通过专用硬件矩阵核心加速矩阵融合乘加（MFMA）运算，其基本形式为D:=A*B+C。在混合精度模式下（输入矩阵使用FP16/FP8等低精度类型，输出保持FP32以保障精度），性能提升显著：
- CDNA3（如MI325X）：FP16性能达1307.4 TFLOPS（较FP32提升8倍），FP8达2614.9 TFLOPS（16倍）。
- CDNA4（如MI355X）：FP16性能2.5 PFLOPS（16倍），FP8达5 PFLOPS（32倍），并新增FP6/FP4支持（64倍提升）。
2. 低精度浮点类型详解
浮点数的二进制表示由符号位、指数位和尾数位组成（如E4M3表示4位指数+3位尾数的8位浮点）。文章对比了常见低精度类型的特性： - FP16（E5M10）：支持±65504范围，保留NaN/Infinity。 - FP8：分E4M3（如OCP标准格式）和E5M2（BF8）两种子类型，CDNA3使用FNUZ变体，CDNA4支持OCP标准。 - FP4（E2M1）：仅6位动态范围，适用于极端低精度场景。 - E8M0：专用于块缩放因子，值域为2^[-127,127]。
3. MFMA指令与性能计算
CDNA4扩展了MFMA指令集，支持：
- 新型指令：如FP6/FP4混合精度运算和块缩放MFMA（32x32x64规模）。
- 性能公式：峰值TFLOPS = 2*M*N*K * 核心数 * (时钟频率/周期数) / 1e6。例如，MI325X的FP16 MFMA（32x32x8）理论性能为1307.4 TFLOPS。
4. 编译器内置函数
通过LLVM内置函数调用MFMA指令，语法示例：
cpp
// FP16输入，FP32输出的16x16x16 MFMA
__builtin_amdgcn_mfma_f32_16x16x16f16(a, b, c, 0, 0, 0);
// CDNA4块缩放FP8指令
__builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4(a, b, c, 0, 0, 0, scale_a, 0, scale_b);
5. 实战代码示例
文章提供了5个HIP内核实现案例：
1. FP32基础MFMA：32x32x2矩阵乘法，展示线程间数据分布。
2. FP16混合精度：16x16x16运算，使用_Float16向量类型。
3. FP8高效处理：通过__hip_fp8_storage_t类型和64位长整型传参。
4. CDNA4块缩放FP8：结合E8M0缩放因子，演示32x32x64运算。
5. FP4极限优化：利用__amd_fp4x2_storage_t和位操作函数处理非连续内存数据。
关键图表说明
- 图1-4：浮点数二进制格式与转换公式。
- 图5-10：不同MFMA指令的数据布局（如线程0负责的矩阵块以红色高亮）。
- 性能对比表：清晰展示CDNA3与CDNA4在各精度下的理论算力差异。
总结与资源
本文系统性讲解了AMD矩阵核心编程的全流程，推荐进一步参考： 1. ROCm官方文档中的低精度类型指南。 2. AMD矩阵指令计算器工具。 3. OCP Microscaling Formats标准文档。
（注：原文中的数学公式、代码片段及性能表格均完整保留核心信息，次要技术细节如广播标志说明等已简化。）
评论总结
这篇评论主要围绕AMD硬件加速和GPU矩阵计算的适用性展开讨论，呈现两种对立观点：
- 支持AMD硬件加速的积极观点
- 认为AMD硬件加速能带来更多元化的选择："More diversity in this space is welcome"（"这个领域需要更多元化"）
- 对相关技术应用表示欢迎："Glad to see more articles out using AMD hardware acceleration especially for matrix math"（"很高兴看到更多使用AMD硬件加速的文章，特别是针对矩阵运算"）
- 质疑GPU矩阵计算效率的观点
- 认为GPU架构不适合矩阵乘法："GPUs are uniquely unsuited for matrix multiplication"（"GPU特别不适合矩阵乘法"）
- 用比喻说明GPU并行处理的局限性："like independently controlling one out of 32 cars on a 32 lane highway"（"就像在32车道高速路上独立控制32辆车中的一辆"）
- 指出实际硬件资源有限："you're feeding something that only exists once or twice per SM"（"你提供的资源每个流处理器只有一到两个"）