Triton 自动调优上手让 GPU 内核自己挑最快的那套参数【免费下载链接】tritonDevelopment repository for the Triton language and compiler项目地址: https://gitcode.com/GitHub_Trending/tri/triton写过 GPU 内核的人都遇到过这种场面周末花了一晚上在BLOCK_SIZE_M、BLOCK_SIZE_N、num_warps之间来回试调了三个晚上最后测出来还是比同事随手写的配置慢 8%。问题不在你不细心而是参数组合是指数级的——块大小、warp 数、流水线级数随便两三个维度就够试到天荒地老。这篇讲的就是 Triton 的自动调优autotune机制把「列出候选配置 → 逐个实测 → 记住最快那个」这件事交给框架做你只负责写内核和列配置。适合已经能用triton.jit写出内核、但还在手动网格搜索参数的人看。autotune 到底在替你干什么把triton.autotune想象成一个专职测试员你给它一张配置清单每次内核被调用时它检查这批输入比如矩阵的 M、N、K是不是见过。没见过的它把清单里的每个配置都真跑一遍、计时然后只把最快的那份留下来存档下次同样尺寸的输入再来直接命中存档不再重复测试。它内部并不神秘每个候选配置对应一次真实 benchmark带 warmup选最小耗时者。配置里编译不过、显存不够的资源类错误会被记成无穷大并跳过不会让整个调优崩掉。源码在python/triton/runtime/autotuner.py想看行为细节可以直接翻。最小上手示例三步给内核加上自动调优以矩阵乘法为例完整写法在官方教程 python/tutorials/03-matrix-multiplication.py 里。第 1 步把「可调的东西」写成配置清单。triton.Config的kwargs是内核里的编译期常量如块大小num_warps、num_stages是编译选项configs [ triton.Config({BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 64, BLOCK_SIZE_K: 32}, num_warps4, num_stages3), triton.Config({BLOCK_SIZE_M: 128, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64}, num_warps8, num_stages4), triton.Config({BLOCK_SIZE_M: 256, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64}, num_warps8, num_stages3), ]这一步的意义把「人凭直觉拍脑袋」变成「让数据说话」清单可以粗后面有剪枝兜底。第 2 步加装饰器指定key。key里的参数值一变就会重新触发一轮完整测试triton.autotune(configsconfigs, key[M, N, K]) triton.jit def matmul_kernel(a_ptr, b_ptr, c_ptr, M, N, K, stride_am, stride_ak, stride_bk, stride_bn, stride_cm, stride_cn, BLOCK_SIZE_M: tl.constexpr, BLOCK_SIZE_N: tl.constexpr, BLOCK_SIZE_K: tl.constexpr): pid tl.program_id(0) # ... 内核主体这一步的意义调优结果是按key值分桶缓存的key定得好缓存命中率才高。第 3 步正常调用即可。第一次调用某个新尺寸的输入会同步跑完一轮 benchmark之后同一尺寸直接走缓存C triton.testing.do_bench(lambda: matmul(A, B, M, N, K, *strides)) # 首测包含调优耗时调优过程想「看得见」设置环境变量TRITON_PRINT_AUTOTUNING1框架会打印每个 kernel 的调优耗时和选中配置。新手最常踩的三个坑现象同一个内核换个小尺寸输入启动又卡了几秒。原因key里的值变了缓存分桶就变了触发新一轮全量测试。如果key写得太细比如把 stride 也放进去几乎每次调用都在重新调优。 怎么办只放真正影响性能的量通常是问题规模其余参数交给缓存机制去重。现象首次调用慢到怀疑挂了。原因configs 列了几十条每条都要完整 benchmark串行跑下来自然长。 怎么办先用prune_configs_by剪枝——early_config_prune可以按输入条件直接砍掉明显不合理的配置比如小矩阵用超大块perf_modeltop_k可以只实测预测最好的前几个。剪到只剩一条时框架会跳过 benchmark 直接用。现象调优结果不可信输出张量值被写坏/累加。原因benchmark 阶段内核会被反复执行凡是会「原地累加」输出的内核第二轮起就是在脏数据上算。 怎么办用reset_to_zero把输出张量在每个配置跑之前清零需要更精细控制的还有restore_value和pre_hook/post_hook具体参数名以官方文档为准。收益与边界一句话讲清楚手动网格搜索triton.autotune覆盖配置数看心情通常个位数清单里全跑可剪枝换输入尺寸全部重来只测新尺寸旧结果复用你的时间数天起写清 configs 和 key 即可边界也老实说autotune 只在清单里选不会发明你没写的配置它优化的是「这组参数下的实测耗时」不做架构级改写。内核本身的算法质量还是你自己负责。下一步别停在本文的代码片段上打开仓库里的 python/tutorials/03-matrix-multiplication.py把其中get_autotune_config()生成的配置清单逐条读一遍然后在自己的内核上复制这套装饰器写法参数语义拿不准时查 docs/python-api/triton.rst 中autotune一节。本文基于仓库当前版本的python/triton/runtime/autotuner.py与官方 tutorial 编写不同版本间 API如 benchmark 相关参数可能有差异请以对应版本的官方文档为准。【免费下载链接】tritonDevelopment repository for the Triton language and compiler项目地址: https://gitcode.com/GitHub_Trending/tri/triton创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考 Triton 自动调优上手：让 GPU 内核自己挑最快的那套参数
Triton 自动调优上手让 GPU 内核自己挑最快的那套参数【免费下载链接】tritonDevelopment repository for the Triton language and compiler项目地址: https://gitcode.com/GitHub_Trending/tri/triton写过 GPU 内核的人都遇到过这种场面周末花了一晚上在BLOCK_SIZE_M、BLOCK_SIZE_N、num_warps之间来回试调了三个晚上最后测出来还是比同事随手写的配置慢 8%。问题不在你不细心而是参数组合是指数级的——块大小、warp 数、流水线级数随便两三个维度就够试到天荒地老。这篇讲的就是 Triton 的自动调优autotune机制把「列出候选配置 → 逐个实测 → 记住最快那个」这件事交给框架做你只负责写内核和列配置。适合已经能用triton.jit写出内核、但还在手动网格搜索参数的人看。autotune 到底在替你干什么把triton.autotune想象成一个专职测试员你给它一张配置清单每次内核被调用时它检查这批输入比如矩阵的 M、N、K是不是见过。没见过的它把清单里的每个配置都真跑一遍、计时然后只把最快的那份留下来存档下次同样尺寸的输入再来直接命中存档不再重复测试。它内部并不神秘每个候选配置对应一次真实 benchmark带 warmup选最小耗时者。配置里编译不过、显存不够的资源类错误会被记成无穷大并跳过不会让整个调优崩掉。源码在python/triton/runtime/autotuner.py想看行为细节可以直接翻。最小上手示例三步给内核加上自动调优以矩阵乘法为例完整写法在官方教程 python/tutorials/03-matrix-multiplication.py 里。第 1 步把「可调的东西」写成配置清单。triton.Config的kwargs是内核里的编译期常量如块大小num_warps、num_stages是编译选项configs [ triton.Config({BLOCK_SIZE_M: 64, BLOCK_SIZE_N: 64, BLOCK_SIZE_K: 32}, num_warps4, num_stages3), triton.Config({BLOCK_SIZE_M: 128, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64}, num_warps8, num_stages4), triton.Config({BLOCK_SIZE_M: 256, BLOCK_SIZE_N: 128, BLOCK_SIZE_K: 64}, num_warps8, num_stages3), ]这一步的意义把「人凭直觉拍脑袋」变成「让数据说话」清单可以粗后面有剪枝兜底。第 2 步加装饰器指定key。key里的参数值一变就会重新触发一轮完整测试triton.autotune(configsconfigs, key[M, N, K]) triton.jit def matmul_kernel(a_ptr, b_ptr, c_ptr, M, N, K, stride_am, stride_ak, stride_bk, stride_bn, stride_cm, stride_cn, BLOCK_SIZE_M: tl.constexpr, BLOCK_SIZE_N: tl.constexpr, BLOCK_SIZE_K: tl.constexpr): pid tl.program_id(0) # ... 内核主体这一步的意义调优结果是按key值分桶缓存的key定得好缓存命中率才高。第 3 步正常调用即可。第一次调用某个新尺寸的输入会同步跑完一轮 benchmark之后同一尺寸直接走缓存C triton.testing.do_bench(lambda: matmul(A, B, M, N, K, *strides)) # 首测包含调优耗时调优过程想「看得见」设置环境变量TRITON_PRINT_AUTOTUNING1框架会打印每个 kernel 的调优耗时和选中配置。新手最常踩的三个坑现象同一个内核换个小尺寸输入启动又卡了几秒。原因key里的值变了缓存分桶就变了触发新一轮全量测试。如果key写得太细比如把 stride 也放进去几乎每次调用都在重新调优。 怎么办只放真正影响性能的量通常是问题规模其余参数交给缓存机制去重。现象首次调用慢到怀疑挂了。原因configs 列了几十条每条都要完整 benchmark串行跑下来自然长。 怎么办先用prune_configs_by剪枝——early_config_prune可以按输入条件直接砍掉明显不合理的配置比如小矩阵用超大块perf_modeltop_k可以只实测预测最好的前几个。剪到只剩一条时框架会跳过 benchmark 直接用。现象调优结果不可信输出张量值被写坏/累加。原因benchmark 阶段内核会被反复执行凡是会「原地累加」输出的内核第二轮起就是在脏数据上算。 怎么办用reset_to_zero把输出张量在每个配置跑之前清零需要更精细控制的还有restore_value和pre_hook/post_hook具体参数名以官方文档为准。收益与边界一句话讲清楚手动网格搜索triton.autotune覆盖配置数看心情通常个位数清单里全跑可剪枝换输入尺寸全部重来只测新尺寸旧结果复用你的时间数天起写清 configs 和 key 即可边界也老实说autotune 只在清单里选不会发明你没写的配置它优化的是「这组参数下的实测耗时」不做架构级改写。内核本身的算法质量还是你自己负责。下一步别停在本文的代码片段上打开仓库里的 python/tutorials/03-matrix-multiplication.py把其中get_autotune_config()生成的配置清单逐条读一遍然后在自己的内核上复制这套装饰器写法参数语义拿不准时查 docs/python-api/triton.rst 中autotune一节。本文基于仓库当前版本的python/triton/runtime/autotuner.py与官方 tutorial 编写不同版本间 API如 benchmark 相关参数可能有差异请以对应版本的官方文档为准。【免费下载链接】tritonDevelopment repository for the Triton language and compiler项目地址: https://gitcode.com/GitHub_Trending/tri/triton创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考