大数跨境
0
0

Day 0 开发者指南:hipBLASLt 离线 GEMM 调优脚本

Day 0 开发者指南:hipBLASLt 离线 GEMM 调优脚本 AMD开发者中心
2025-11-21
1

原文作者:Han Lin, Spandan Tiwari, Xinjun Niu, Chao Li, Wei Luo, Junyan Yang, Carson Liao, Chunhung Wang, Ashish Sirasao.

本文聚焦于如何借助QuickTune 脚本优化真实模型的性能,并以在 AMD GPU 上对 Qwen 模型进行离线 GEMM 调优为例QuickTune 由 AMD Quark 团队开发,可在极小时间开销下显著提升 GEMM 性能。它是一个基于 hipBLASLt 的高级离线 GEMM 调优工具;相比手动使用 hipblaslt-bench 对模型进行调优,QuickTune 支持“一键调优”,显著减少手动配置与执行成本。

关于如何使用hipblaslt-bench 进行离线调优的基础内容,可参考前一篇博客 hipblaslt offline tuning part 2 [1]。那篇文章介绍了如何在单个 GEMM 上使用基础调优工具 hipblaslt-bench。本文在此基础上,介绍一个基于 hipblaslt-bench 的高级调优脚本,用于对模型中大量 GEMM 进行批量调优。

背景

矩阵–矩阵乘法(GEMM,General Matrix–Matrix Multiplication) 是 AI 工作负载中最核心的运算之一在训练与推理中,GEMM 通常占据主要算时。因此,即便是 GEMM 性能的小幅提升,也能显著改善端到端表现。围绕矩阵乘法的优化,业界已进行了大量研究。

GEMM 调优的目标是在给定的生产场景下、在吞吐、时延、goodput(有效吞吐)等指标上,针对特定矩阵尺寸、数据类型和硬件架构,找出表现最佳的实现。现代计算库在 CPU 和 GPU 上通常提供多种 GEMM 算法和内核实现,它们在以下方面可能不同:

  • 分块与tile 大小

  • 内存布局

  • 线程配置

  • 指令和流水线的使用方式


常见的GEMM 调优方式包括:

  • 离线调优(Offline Tuning):部署前对不同方案进行基准测试并选择最优方案

  • 在线调优(Online Tuning):运行时动态调优,根据实际输入特征选择最优方案


与在线调优相比,离线调优的优势包括:

  • 性能更高:可预先穷举/大空间搜索,更易找到最优配置

  • 运行时零调优开销:运行时直接使用预调优结果

  • 性能稳定可预测:结果在多次运行间具有确定性


因此,离线GEMM 调优近年来被广泛用于性能敏感场景。在这套生态中,核心组件是 hipBLASLt,它通过 hipblaslt-bench 支持离线预调优。诸如 PyTorch ROCm 等框架会以 TunableOp 的形式,将这套能力透明地暴露给用户,让开发者与终端用户无需手动选择内核即可获得优化后的 GEMM 性能。近期的 AITER 库中,AMD 还提供了 gradlib工具,用于在当前硬件上自动选出性能更优的 GEMM 参数。

Offline Tuning 离线调优

离线GEMM 调优是在高性能计算和深度学习框架中常用的一种性能优化技术。其目标是在已有的解空间中选择性能最佳的内核方案,并将最佳方案的索引写入调优结果文件。运行时可依据该索引直接调度对应内核,无需再次调优。流程包括:

1.Benchmark:在目标GEMM 形状和数据类型上,对候选解进行基准测试。

2.Profile:采集每个解的执行时间或吞吐(可用合成或真实数据)。

3.Select:选出性能最优解。

4.Save:将结果保存到缓存或文件。

5.Load:运行时直接加载最优解。

离线调优避免了运行期的调优开销,并保证性能的可预测性,尤其适合大规模部署或实时系统。

hipblaslt-bench

hipBLASLt 提供命令行基准工具 hipblaslt-bench,用于在 AMD GPU 上基于 hipBLASLt API 测量 GEMM 性能并验证正确性hipblaslt-bench 是 hipBLASLt 离线 GEMM 调优的核心组件。近期,hipblaslt-bench 新增了对 swizzle GEMM 内核的支持,该内核针对性能做了特别设计,可显著提升表现。

Swizzle GEMM

Swizzle GEMM 指在矩阵乘计算中,通过对输入或输出张量的内存布局进行刻意重排(swizzling) 来提升硬件效率Swizzling 一般指对数据进行结构化重排或置换,以提升:

  • 内存合并访问

  • Cache 局部性

  • 线程级并行

  • 计算单元间负载均衡


hipblaslt-bench API 关键参数

完整API 参见官方 README [2] 与前文 hipblaslt offline tuning part 2 [1]。

  • -m / -n / -k:定义待调优GEMM 的问题规模 (m, n, k)

  • --lda / --ldb / --ldc / --ldd定义矩阵A/B/C/D 的 leading dimension

  • --a_type / --b_type / --c_type / --d_type定义矩阵A/B/C/D 的数据类型

  • --transA / --transB指定矩阵A/B 是否转置

  • --cold_iters / --iters预热轮数与计时迭代数

  • --algo_method搜索模式(heuristic / all / index)

  • --solution_index--algo_method index 搭配,指定解索引

  • --requested_solutionheuristic 搭配,指定候选解数量(-1 表示全解空间,1 表示仅测基线解)

  • --swizzleA启用swizzle GEMM,对矩阵 A(权重)的内存布局进行重排


QuickTune 工作流

QuickTune 是 AMD Quark 团队开发的 Python 脚本工具,用于辅助开发者与客户执行 hipBLASLt 离线调优。完整的离线调优流程包括:输入日志处理、参数设置、调优执行和结果分析。借助 QuickTune,即便缺乏专业调优经验,也能在极小时间成本下显著提升 hipBLASLt 的 GEMM 性能。整个离线流程可通过运行一个 Python 脚本一次完成,也可按步骤分别执行。

1. hipBLASLt 离线调优流程

调优前:获取 hipBLASLt 输入日志(Phase 1)

QuickTune 的输入是 hipBLASLt 日志文件。通过设置如下环境变量并运行模型,即可生成该日志:

export HIPBLASLT_LOG_MASK=32export HIPBLASLT_LOG_FILE=<path/to/hipblaslt/log>

日志会保存到HIPBLASLT_LOG_FILE 指定路径。其内容本质上是:模型运行过程中每次调用 hipBLASLt GEMM 内核时所对应的 hipblaslt-bench 命令行,这些命令行已包含离线调优所需的大部分参数。示例如下:

2. hipBLASLt 日志示例

注意日志只会包含实际调用hipBLASLt 内核的 GEMM 形状,并非模型中所有 GEMM都会被记录。

Step 1:hipBLASLt Log 去重

模型运行中,同一GEMM 形状往往会被多次调用,因此日志里会出现大量重复行。QuickTune 将:

  • 删除重复命令行

  • 并按需要调整部分参数,用于基线性能测试与不使用swizzle 的离线调优等


示例命令:

python remove_duplicate.py    --input_file example/Qwen3_hipblaslt.log    --output_path tuning_test

Step 2:一键调优

以下示例展示如何对Qwen3-32B 模型使用 hipblaslt-bench 进行 GEMM 调优:

nohup python gemm_tuning.py --input_file example/Qwen3-32B_ali_hipblaslt.log --output_path test_tuning --requested_solution 128 --swizzleA > output.log 2>&1 &

关键参数说明:

  • input_file包含待调优GEMM 的 hipBLASLt 日志文件(脚本运行时会自动去重)

  • output_path结果输出目录,用于保存调优结果文件

  • --swizzleA启用后对矩阵A 的内存布局进行重排,以减少 bank conflict,从而降低时延。若要启用带 swizzle 的离线调优,transA 需设置为 T

  • --requested_solution控制搜索空间大小(默认128)

    o 1:仅测量原始内核的基线时延

    o -1:搜索整个解空间

    o N > 1:在 N 个候选解中搜索最佳方案

  • --cold_iters / --iters预热轮数/ 计时迭代数。推荐设置:

    小:时延< 50 μs → (200, 100)

    中:100 μs ≤ 时延 < 1000 μs → (50, 20)

    o 大:时延≥ 1000 μs → (10, 2)

  • --gpu_id指定用于调优的GPU。调优对时延敏感,建议用 rocm-smi 找到一块空闲 GPU,以获得更可靠的结果

  • stablize_gpu启用后GPU 处于更一致的运行设置,整体性能略低,但硬件状态更可控(请以脚本实际参数名为准)


调优结果会保存在<output_path>/tuning.txt 。

Step 3:调优结果分析

为评估调优收益,QuickTune 会生成 CSV,包含基线与调优后解的索引及其时延,并基于此计算提升比例。由于不同 GEMM 的出现次数与时延占比不同,整体加速比不能用简单平均表示。提供的脚本可基于输入日志与调优结果 CSV 计算 hipBLASLt GEMM 的整体加速比(注意:仅针对 hipBLASLt GEMM,而非整个模型):

python tuning_analysis.py --input_log <input hipblaslt log> --input_csv <tuning_csv> --output_csv <analysis csv file>

分析结果示意如下:

3. 调优结果分析示例

调优后:应用调优结果

要在推理/训练中实际使用调优结果,设置如下环境变量并重新运行模型。之后,hipBLASLt 会依据调优结果自动选择相应内核。流程见图 4。

unset HIPBLASLT_TUNING_FILEexport HIPBLASLT_TUNING_OVERRIDE_FILE=<path/to/tuning/result>

4. 应用调优结果的流程

可选步骤:结果验证

你可以重复前文“调优前:获取 hipBLASLt 输入日志(Phase 1)” 的流程:在应用调优结果后,再次运行模型并收集 hipBLASLt 日志,检查日志中的 solution index 是否与调优结果文件一致;若一致,说明调优结果已成功生效。

Qwen 离线调优示例

Qwen3-32B 为例,对比调优前后的性能表现,分别关注:

  • 不启用swizzle 的离线调优

  • 启用swizzle 的离线调优


1 展示了 Qwen3-32B 中 hipBLASLt GEMM 内核的时延数据:

  • A、B、C:每个 GEMM 的 (m, n, k) 形状

  • D:离线调优前的原始时延

  • E 与 G:不启用 swizzle 时离线调优后的时延

  • F 与 H:调优前时延 / 调优后时延 的比值(提升比),数值越大表示收益越大


1. Qwen3-32B 中 hipBLASLt GEMM 内核时延

从表中可见,在不启用swizzle 的情况下,提升比大致分布在 100%~139%。其中 100% 表示调优未带来性能提升,这通常说明原始内核在当前搜索空间中已是最优解。对于已接近 roofline 的内核,即使调优,提升空间也会非常有限。

启用swizzle 的离线调优在各 GEMM 形状上均取得更高提升:平均提升比为 1.1043×(+10.43%)(不启用 swizzle)与 1.3147×(+31.47%)(启用 swizzle)。可见 swizzling 是 GEMM 调优中行之有效的优化手段。

我们还统计了端到端推理时间,以展示离线调优对Qwen3-32B 的影响:调优前为 9402.67 ms,离线调优后为 8000.73 ms,端到端整体性能提升 14.91%。

2. Qwen3-32B 端到端性能对比

为方便复现实验,部分环境信息如下:

总结

GEMM 位于大规模 AI 工作负载的核心位置,在 Qwen 等 LLM 中几乎决定训练与推理的整体性能,因此每一次优化都具有价值。通过本文可见:

  • 离线调优 是将GEMM 性能适配到具体模型形状与硬件上的一种有效方法

  • Swizzling 通过优化内存访问效率,能稳定地增强调优效果


随着模型与硬件持续演进,系统化的离线GEMM 调优将长期是充分挖掘 GPU 性能潜力的重要手段。

更多hipblaslt-bench 与基础调优方法,参见前文:hipblaslt offline tuning part 2 [1]。建议先掌握该文中的基础用法,再按本文 Step 1–3 执行一键调优与结果应用。

致谢

在本文撰写与相关工作推进过程中,我们特别感谢以下同事在技术与反馈方面给予的大力支持:

  • Ethan Yang

  • Jiangyong Ren

  • Hang Yang

  • Yilin Zhao

  • Yu Zhou

  • Yu-Chen Lin

  • Jacky Zhao

  • Thiago Fernandes Crepaldi

  • Bowen Bao


以及AMD Quark 团队的所有成员。

参考链接

[1] hipblaslt offline tuning part 2:https://rocm.blogs.amd.com/software-tools-optimization/hipblaslt-offline-tuning-part2/README.html

[2] README:https://rocm.docs.amd.com/projects/hipBLASLt/en/latest/clients.html#hipblaslt-bench

【声明】内容源于网络
0
0
AMD开发者中心
AMD开发者中心为开发者提供工具、技术和资源,助力构建AI解决方案。ROCm、Ryzen AI软件和ZenDNN,帮助您实现模型加速与部署。开发者可通过文档、SDK及教程快速上手。立即关注AMD开发者中心,开启智能未来!
内容 65
粉丝 0
AMD开发者中心 AMD开发者中心为开发者提供工具、技术和资源,助力构建AI解决方案。ROCm、Ryzen AI软件和ZenDNN,帮助您实现模型加速与部署。开发者可通过文档、SDK及教程快速上手。立即关注AMD开发者中心,开启智能未来!
总阅读49
粉丝0
内容65