GEMM优化挑战赛优秀选手技术方案解析

2026年初，摩尔学院举办了“GEMM优化挑战赛”，要求在摩尔线程 S4000 GPU 上优化一个超大规模半精度矩阵乘法：C = A × B，规模为 8192×8192×16384。选手必须使用 FP32 累加，最终输出 FP16，并保证数值误差小于 1e-2。

这项任务不仅考验对张量核心（Tensor Core）的驾驭能力，更检验对 MUSA 编程模型、数据布局、共享内存、流水线同步等底层细节的掌握程度。最终，部分优秀选手在严格正确性前提下, 逼近甚至超越官方基线的高吞吐性能。

本次博客分享四位优秀选手的技术报告，梳理出共性优化策略、各自独特的杀手锏以及对国产 GPU 生态的启示。

硬件基础：MTT S4000 的关键特性

在深入优化之前，必须理解目标硬件：

                             表1. MTT S4000 硬件规格表

规格类别	规格项	参数
产品名称	市场名称	MTT S4000
芯片	图形芯片	QY102AA-800
	GPU 架构	曲院 (Qu Yuan) 架构
	MUSA 核心数	8192
核心频率	GPU 频率	1.5 GHz
	NOC 频率	1.3-1.4 GHz
显存规格	显存容量	48 GB
	显存类型	GDDR6
	显存速率	16 Gbps
	显存位宽	384 bits
	显存带宽	768 GB/s
算力（向量）	FP32 / FP16	25 TFLOPS
算力（张量）	TF32	50 TFLOPS
	FP16 / BF16	100 TFLOPS
	INT8	200 TOPS
	FP64	支持
缓存	L2 缓存	1 MB per 4 MPs，全局 24 MB L3
线程模型	Warp 大小	128个线程
共享内存	Bank 数量	32 个
张量核心 MMA 尺寸	-	支持 16×8×8、16×16×8、16×16×16、32×32×16
渲染性能	Texture Rate	768 GTexels/s
	Pixel Rate	768 GPixels/s
多媒体	视频编码	48 路 1080p@30fps (HEVC)
	编码格式	AV1、H.264、H.265
	解码格式	AV1、H.264、H.265、VP9、AVS2
	显示输出	超高清 8K HDR
功耗	TGP	450W
互联	多卡互联	自研 MTLink 技术
虚拟化	硬件虚拟化	支持 GPU 核心和视频处理核心的硬件切分与隔离
安全	安全引擎	支持 TEE 与国密标准
软件兼容	CUDA 兼容	支持 (通过 musify 工具自动转化)

产品定位与特点

MTT S4000 是基于摩尔线程曲院 GPU 架构打造的全功能元计算卡，主要特点包括：

• 大模型训推一体：专为千亿规模大语言模型的训练、微调和推理优化，第三代 MUSA 核心配备新一代 Tensor 核心

• 全精度支持：支持 FP64/FP32/TF32/FP16/BF16/INT8 等完整计算精度

• 云原生支持：全面支持云端虚拟化与容器化 GPU 的使用和调度，兼容 K8S 生态

• 图形渲染能力：支持 OpenGL、OpenGL ES、Vulkan、DirectX 等现代图形 API

• CUDA 兼容：通过 musify 移植工具可自动转化并编译 CUDA 代码为 MUSA 代码

共性优化策略

四位选手的优化路径虽各有侧重，但都遵循了一套经典的高性能 GEMM 优化范式，并针对 MTT S4000 做了适配。

• 分块策略：找到“最稳”的平衡点

• 内存层次优化：搬得快、算得密

• 缓存局部性优化：CTA Swizzle

• 编译器与指令级优化

除了上述共性优化，每位选手还贡献了独特的工程智慧。详情请参看每位选手的具体文档。

四篇技术文档列表:

• 《MTT S4000 GEMM 优化全实现 - B矩阵预转置缓存 + 异步流水线双缓冲》- (作者: 朱振华)

  • 在第一次调用时启动 transpose_b_kernel 将列主序 B 转为行主序，并利用静态指针缓存避免重复转置。评测循环中后续调用直接复用，平均吞吐显著提升（50.6T → 53.0T）。

  • 给出了双缓冲 + 异步拷贝的完整代码宏（COMPUTE_FROM_BUF），并对 load_matrix_sync / mma_sync 做循环展开。

  • 性能最终达到 53.57 TFLOPS。

• 《MTT S4000 GEMM 优化中的分块、并发与尾处理实战复盘》- (作者:赵奕铭)

  • 对比 N64/N128/N256 分块，确认 N128 + K32 + 8 wave 为最平衡配置。

  • 发现主核单流 + 尾处理双流是最稳定的并发结构，性能 57.62 TFLOPS。

  • 详细记录了“失败路线”（如 16B/线程 loader 掉到 47.98T），强调实验框架与版本管理的重要性

• 《MTT S4000 张量核心 GEMM 优化的推导、编译器调优与异步边界探索》- (作者:林瑞晨)

  • 通过表格穷举不同分块下的寄存器、共享内存占用，推导出最优组合（每个 MP 1 个 block，每 block 8 个 warp，每 warp 8 次 32×32×16 MMA）。

  • 探索了 MUSA 编译器属性（mtgpu_tiny_offset, mtgpu_enable_ldma_index 等）带来的额外收益（~3 TFLOPS）。

  • 提出奇偶切换读取顺序，根据 K 轮次奇偶性决定先读 A 还是先读 B，提高缓存命中率，性能提升 3 TFLOPS。

  • 激进实验：删除第二个 __syncthreads()，在特定条件下达到 63 TFLOPS（奇思妙想）。

  • 最终稳定性能 58.5 TFLOPS。

• 《MTT S4000 FP16 GEMM 优化实战逼近 mcc 编译器的物理上限》- (作者:陈华宾)

  三段式工程复盘（开发→塌陷→重建），物理 warp 粒度错配的根本性发现：

  • 开发阶段：从 64×64 一路迭代到 256×256 mega tile + 32×32×16 WMMA，性能爬到 ~90 TFLOPS，但正确性测试只覆盖了小尺寸。

  • 塌陷：评审期新增 8k 正确性用例，发现 90T 版本 75% 元素错误。根因是:

    • 代码层根因(75% 错的物理原因):mega tile kernel 把 warp 当成 32 线程,在 MP22 上同一 128 线程物理 warp 内的 4 个 32 线程子 warp给 WMMA 传了分歧地址,硬件只采纳其中一个,3/4 输出是数据副本。

    • 测试层根因(为什么没早点暴露):dispatch 分流让所有 legacy 正确性用例都走 fallback 单 warp 内核,高性能 mega tile路径只在性能测试上跑过,而性能测试不校验输出。

  • 重建（6 天）：

    • 修复物理 warp 粒度：将 warp 从 32 线程改为 128 线程（物理 warp），同一 warp 内所有线程地址必须一致。修复后正确性通过，真实性能 22T（原 90T 是 1/4 工作量的幻觉）。

    • stride‑33 bank conflict 消除：epilogue 中 stride 从 32 改为 33（与 16 bank 互质），22T → 48T（+118%）。

    • M‑first 蛇形 CTA 排序：48T → 54T（+12%）。

    • 256×256×32 + 1024 线程：54T → 57T（+6%）。

    • postRA + ILP 调度 flag：57T → 67T（+18%）。

    • FenceSetting=mixed：67T → 73T（+9%）。

    • 最后 4 个 load‑store flag 边际提升至 73.8 TFLOPS。

  • 未走通的路：internal-regs-ra 导致正确性挂（但性能可到 79T），mma.h:761 头文件可能多分配 3 倍寄存器。

  • 结论：73.8 TFLOPS 是当前 C++ WMMA API + mcc 3.1.0 编译路径的经验上限（VLIW 165/165 满打包，C++ 双缓冲 5 种方案全败，手动 fragment 加载被硬件 bank 封死）。相比 muBLAS 78.5T 的差距定位在 5 个地址计算冗余的 VLIW word 上，需要手写汇编或工具链修复才能突破。

四位选手的方案均超越了 muBLAS 基线效率的 30% 满分线（按比赛评分标准，≥30% 效率得 70 分）。考虑到 muBLAS 在 8192×8192×16384 上的实测约 86 TFLOPS，选手们达到了 62% ~ 85% 的峰值性能，对于手工优化算子而言已是顶尖水平。

对国产 GPU 生态的启示

• 软件栈成熟度是关键：MUSA SDK 提供了 WMMA 库、异步拷贝原语、性能分析工具（msys, mthreads-gmi），已能支撑深度优化。

• 硬件特性需深度理解：128 warp 和 32×32×16 Tensor Core 是优势，但也要求分块策略重新设计，不能简单照搬 CUDA 经验。物理 warp 粒度是最容易被忽视的陷阱。

• 未来方向：warp 专用化、异步同步移除值得进一步探索。