§ 00 · Why autotune existsPrologue
§ 00 · autotune 为何存在序言
A GPU kernel has knobs that do not change what it computes, only how fast — tile sizes, the number of warps, software-pipeline stages, register limits, cluster width. The fastest setting of those knobs is not a constant, and it is worth being precise about why, because that root cause is the reason autotuning exists at all. A GPU runs fastest when two things are balanced: enough parallel work in flight to hide memory latency (occupancy), and a tile large enough to reuse data instead of re-fetching it. Bigger tiles reuse more but need more registers and shared memory, which lowers occupancy — so the sweet spot is a trade-off, and where it lands moves with the problem. A GEMM at M=1 is memory-bound and wants a small tile; the same GEMM at M=8192 is compute-bound and wants a large one. The dtype changes how many bytes each element costs; the exact GPU changes how many registers and how much shared memory you have to spend. No formula on a datasheet folds all of that together, so you cannot read the best config off — you have to run the kernel and measure. That measure-and-pick process is autotuning, and every modern kernel DSL has some version of it.
一个 GPU kernel 有一些旋钮, 拧它们不改变算出来的结果, 只改变快慢 —— tile 大小、 warp 数、 软件流水线级数、 寄存器上限、 cluster 宽度。 这些旋钮的最优值不是常数, 而且值得说清楚为什么, 因为这个根因正是 autotune 存在的理由。 GPU 跑得最快, 要同时平衡两件事: 有足够多的并行任务在飞、 好把访存延迟藏住 (occupancy), 以及 tile 大到能复用数据、 而不是反复重取。 tile 越大复用越多, 但要吃更多寄存器和共享内存, 从而压低 occupancy —— 所以最优点是一个权衡, 而它落在哪会随问题变化。 GEMM 在 M=1 时是访存受限、 想要小 tile; 同一个 GEMM 在 M=8192 时是计算受限、 想要大 tile。 dtype 改变每个元素占几个字节; 具体那块 GPU 改变你有多少寄存器、 多少共享内存可花。 没有哪个数据手册上的公式能把这些一起折叠出来, 所以你读不出最优配置 —— 只能把 kernel 跑一遍、 测出来。 这个"测量再挑选"的过程就是 autotune, 每个现代 kernel DSL 都有它的某种版本。
This guide reads six real implementations from source — Triton (the template everyone copies), quack and CuteDSL (two DSLs that build on the idea), aiter (offline tuning at a kernel library), SGLang (an inference engine that consumes tuned configs), and FlyDSL (which has the machinery but no adoption) — and uses them to answer one question: how should FlyDSL's autotune actually be designed? The five references are evidence; the FlyDSL design and its implementation guide are the destination.
这篇 guide 从源码精读六个真实实现 —— Triton (大家都在抄的范本)、 quack 和 CuteDSL (两个在这想法上扩展的 DSL)、 aiter (kernel 库的离线 tuning)、 SGLang (一个消费已 tune 配置的推理引擎)、 以及 FlyDSL (有机制但没人用) —— 用它们回答一个问题: FlyDSL 的 autotune 到底应该怎么设计? 五个参考是论据; FlyDSL 的设计和实现 guide 是终点。
Autotuning runs the same kernel dozens of times to pick a winner. If the kernel writes its output in place or accumulates, those repeated runs corrupt the data and you measure garbage. Every correct autotuner solves this; most naive ports forget it. Hold this thought — it returns in § 03.
autotune 会把同一个 kernel 跑几十遍来挑最快的。 如果这个 kernel 原地写输出、 或者做累加, 这些重复运行会污染数据, 你测到的就是垃圾。 每个正确的 autotuner 都解决了这个问题; 大多数草率的移植都忘了它。 记住这一点 —— 它在 § 03 会回来。
§ 01 · The shape of every autotunerThe autotune loop
§ 01 · 每个 autotuner 的形状autotune 循环
Strip away the framework differences and every autotuner is the same four-step loop, wrapped in a cache. A search space of candidate configs is defined; each candidate is compiled and benchmarked; the fastest is selected; the winner is cached under a key so the next call with the same shape skips the search entirely. The interesting engineering is all in the edges: what the key contains, how you avoid the search when you can, how you keep benchmarking from corrupting state, and whether the loop runs online (at first call) or offline (ahead of time).
剥掉框架差异, 每个 autotuner 都是同一个四步循环, 外面套一层缓存。 先定义一个候选配置的搜索空间; 每个候选编译并 benchmark; 挑出最快的; 把赢家按一个 key 缓存起来, 这样下次同样形状的调用就完全跳过搜索。 有意思的工程全在边角: key 里放什么、 怎么在能省的时候省掉搜索、 怎么让 benchmark 不污染状态、 以及这个循环是在线跑 (首次调用时) 还是离线跑 (提前)。
§ 02 · The design everyone copiesTriton · the template
§ 02 · 大家都在抄的设计Triton · 范本
Triton's @triton.autotune is the design quack and FlyDSL both descend from, so it is worth reading first and reading closely. The whole thing lives in one file, python/triton/runtime/autotuner.py (511 lines). The core is Autotuner.run: build a key from the designated key-args plus their dtypes, and if it misses, benchmark the pruned configs and cache the winner.
Triton 的 @triton.autotune 是 quack 和 FlyDSL 共同继承的设计, 所以值得先读、 而且要细读。 整个东西就在一个文件里, python/triton/runtime/autotuner.py (511 行)。 核心是 Autotuner.run: 用指定的 key-arg 加上它们的 dtype 构造一个 key, 如果 miss 就 benchmark 剪枝后的配置、 缓存赢家。
:217.:217。Three things in this file separate a real autotuner from a toy, and all three are worth stealing.
这个文件里有三样东西, 把一个真正的 autotuner 和一个玩具区分开, 三样都值得偷。
@heuristics — the skip-the-search door
@heuristics — 跳过搜索的那扇门
Not every knob needs a search. If a parameter can be computed from the inputs (a block size that should just be the next power of two above N), Triton lets you derive it with @triton.heuristics instead of enumerating it. Autotuning is reserved for the genuinely uncertain knobs. This is the single biggest lever on tuning cost: the cheapest search is the one you never run.
不是每个旋钮都需要搜索。 如果一个参数能从输入算出来 (比如 block size 就该取大于 N 的下一个 2 的幂), Triton 让你用 @triton.heuristics 推导它, 而不是去枚举。 autotune 只留给真正不确定的旋钮。 这是 tuning 成本上最大的杠杆: 最便宜的搜索是你根本没跑的那次。
prune_configs_by — cut before you compile
prune_configs_by — 编译前先砍
Compiling and benchmarking is the expensive part. prune_configs_by takes a perf_model + top_k (keep only the model's predicted-best k) and an early_config_prune (drop configs that can't fit, e.g. too much shared memory) — both run before any compilation. You only pay to benchmark candidates that have a chance.
编译和 benchmark 是最贵的部分。 prune_configs_by 接受一个 perf_model + top_k (只保留模型预测最好的 k 个) 和一个 early_config_prune (砍掉跑不起来的配置, 比如 shared memory 超了) —— 两者都在任何编译之前跑。 你只为有机会的候选付 benchmark 的成本。
The cache key — what makes a config stale
cache key — 什么让一个配置过期
Triton's disk-cache key (check_disk_cache, :175) is the most complete of any reference here. It is a SHA-256 of five things, and each one answers "when should I throw the cached config away?"
Triton 的磁盘缓存 key (check_disk_cache, :175) 是这里所有参考中最完整的。 它是五样东西的 SHA-256, 每一样都回答"我什么时候该把缓存的配置扔掉?"
:175.:175。§ 03 · restore_value / reset_to_zeroThe correctness soul
§ 03 · restore_value / reset_to_zero正确性的灵魂
Here is the rule from the prologue, made concrete. To benchmark a config, Triton runs the kernel many times back-to-back and takes the median. If the kernel mutates one of its inputs — an in-place accumulate, a fused residual-add that writes back, a softmax-over-self — then run 2 sees the output of run 1 as its input. The numbers drift, the timing is measured on corrupted data, and the "fastest" config you select may simply be the one that diverged fastest. The selection is wrong in a way no unit test on the kernel itself will catch.
把序言里那条铁律说具体。 为了 benchmark 一个配置, Triton 会把 kernel 连续跑很多遍取中位数。 如果这个 kernel 会改写它的某个输入 —— 原地累加、 写回的 fused residual-add、 对自身做的 softmax —— 那么第 2 遍就会把第 1 遍的输出当成自己的输入。 数值开始漂移, 计时是在被污染的数据上测的, 你选出来的"最快"配置可能只是漂移得最快的那个。 这个选择是错的, 而且 kernel 自己的单元测试一个都抓不到。
Triton's answer is two hooks on the decorator: reset_to_zero (zero the named tensors before each rep) and restore_value (snapshot the named tensors, restore them after each rep). Both run around every benchmarked call, so each config is measured against the same clean inputs.
Triton 的解法是装饰器上的两个钩子: reset_to_zero (每遍之前把指定张量清零) 和 restore_value (把指定张量存档, 每遍之后还原)。 两者都在每次被 benchmark 的调用前后跑, 这样每个配置都是在同样干净的输入上测的。
restore_value, so autotuning any in-place kernel today would select on corrupted reruns. This is a prerequisite, not a nice-to-have — it returns in the FlyDSL guide (§ 11).restore_value 的对应物, 所以今天 autotune 任何原地 kernel 都会在被污染的重跑上做选择。 这是前提, 不是可选项 —— 它在 FlyDSL guide (§ 11) 会回来。§ 04 · Tri Dao's CuTe-based kernelsquack · two-track configs
§ 04 · Tri Dao 的 CuTe 系 kernelquack · 双轨配置
quack is Tri Dao's kernel library built on CuteDSL. Its autotuner (quack/autotuner.py, 849 lines) is the most production-hardened of the references, and unlike FlyDSL it is actually adopted — rmsnorm and gemm both wear @autotune(configs=[...]). It takes the Triton design and adds three things worth studying.
quack 是 Tri Dao 基于 CuteDSL 的 kernel 库。 它的 autotuner (quack/autotuner.py, 849 行) 是所有参考里最经得起生产考验的, 而且不像 FlyDSL, 它真的被采用了 —— rmsnorm 和 gemm 都戴着 @autotune(configs=[...])。 它拿过 Triton 的设计, 加了三样值得研究的东西。
Two-track configs — heuristic default and exhaustive search
双轨配置 — 启发式默认 和 穷举搜索
Each kernel has a *_config.py (e.g. quack/rmsnorm_config.py) that exposes both an analytical get_default(N, dtype, arch) — a hand-tuned heuristic that picks a good-enough config with zero search — and get_all_fwd_configs(), the full exhaustive space. Day-to-day runs take the heuristic; you only pay for the exhaustive search when you explicitly autotune. This is the same idea as Triton's @heuristics, made into a per-kernel convention. It is the cure for FlyDSL's current all-or-nothing.
每个 kernel 有一个 *_config.py (比如 quack/rmsnorm_config.py), 同时暴露一个解析式的 get_default(N, dtype, arch) —— 一个手调的启发式, 零搜索就挑出够好的配置 —— 和 get_all_fwd_configs(), 完整的穷举空间。 日常运行走启发式; 只有你显式 autotune 时才付穷举搜索的成本。 这和 Triton 的 @heuristics 是同一个想法, 做成了 per-kernel 的约定。 它正是 FlyDSL 当前"全有或全无"的解药。
Multi-layer cache with a source fingerprint
带源码指纹的多层缓存
The problem the cache solves: compiling a kernel is slow (seconds), but a program calls the same kernel thousands of times. You want to compile once and reuse — but "reuse" is only safe if nothing that affects the compiled code has changed. quack's caching doc (AI/caching_and_parallel_compilation.md) layers four caches, each one catching a case the layer above can't.
缓存要解决的问题: 编译一个 kernel 很慢 (秒级), 但一个程序会把同一个 kernel 调用成千上万次。 你想只编译一次、 之后复用 —— 但"复用"只有在"影响编译结果的东西都没变"时才安全。 quack 的缓存文档 (AI/caching_and_parallel_compilation.md) 叠了四层缓存, 每一层都兜住上一层兜不住的情况。
- Layer 1 — in-memory dict. Within one running process, the same kernel called twice returns the cached compiled object instantly. Fastest, but it dies when the process exits.
- 第 1 层 —— 内存 dict。 在同一个运行中的进程里, 同一个 kernel 第二次调用直接返回缓存好的编译产物, 瞬间命中。 最快, 但进程一退出就没了。
- Layer 2 — on-disk
.ofile. So the next process (a new run, a test worker) doesn't recompile from scratch, the compiled object is written to disk, keyed by(function name, its arguments). Survives across runs. - 第 2 层 —— 磁盘上的
.o文件。 为了让下一个进程 (新一次运行、 一个测试 worker) 不用从头重编, 编译产物被写到磁盘, 用(函数名, 它的参数)作 key。 能跨运行存活。 - Layer 3 — a source fingerprint on that disk key. Here's the trap the first two layers can't catch: if you edit the kernel source, the disk key
(name, args)is unchanged, so layer 2 would happily hand back a stale.ocompiled from the old code. quack prevents this by folding a fingerprint into the key — a SHA-256 of everyquack/*.pyfile plus the CUTLASS and Python versions. Change one line of kernel source, or bump CUTLASS, and the fingerprint changes, so the old cache entry is simply never found and the kernel recompiles. This is what makes the disk cache safe to keep around. - 第 3 层 —— 给磁盘 key 加一个源码指纹。 这里是前两层兜不住的坑: 如果你改了 kernel 源码, 磁盘 key
(名字, 参数)并没变, 于是第 2 层会高高兴兴地把一个用旧代码编出来的、 过期的.o还给你。 quack 的做法是往 key 里揉进一个指纹 —— 每个quack/*.py文件加上 CUTLASS、 Python 版本的 SHA-256。 改一行 kernel 源码、 或升一次 CUTLASS, 指纹就变了, 于是旧的缓存条目根本查不到、 kernel 自动重编。 正是这一层让磁盘缓存"留着也安全"。 - Layer 4 — an optional autotune-result cache. Layers 1–3 cache the compiled kernel. This one caches the autotune decision — which config won the benchmark for a given shape — so you don't re-run the whole search next time.
- 第 4 层 —— 可选的 autotune 结果缓存。 第 1–3 层缓存的是编译好的 kernel。 这一层缓存的是 autotune 的决定 —— 对某个形状, 哪个配置在 benchmark 里赢了 —— 这样下次不用把整个搜索重跑一遍。
One more detail worth stealing: when quack builds the tuning key, it doesn't store the raw stride of each tensor — it normalizes it to just {0, 1, other}. Why: a stride of 0 means a broadcast dimension and 1 means contiguous, and those are the only distinctions that change which kernel is optimal; the exact stride number (say 4096 vs 8192) doesn't. Collapsing strides this way means a contiguous tensor and a broadcast tensor of the same logical shape land on the right cache entries without exploding the key space with numbers that don't matter.
还有一个细节值得偷: quack 构造 tuning key 时, 不存每个张量的原始 stride, 而是把它规整成只剩 {0, 1, other} 三种。 为什么: stride 为 0 表示这是个广播维度, 为 1 表示连续, 而只有这两种区别会改变"哪个 kernel 最优"; 具体的 stride 数字 (比如 4096 还是 8192) 不会。 这样折叠 stride, 既能让同一逻辑形状下的连续张量和广播张量落到各自正确的缓存条目, 又不会让一堆无所谓的数字把 key 空间炸开。
Parallel precompile — kill the first-run cost
并行预编译 — 干掉首跑成本
The main pain of online autotune is the slow first call. quack's _precompile (_compile_worker.py) compiles all candidate configs in a subprocess pool to warm the .o cache before benchmarking, with FileLock for cross-process safety. This is the practical answer to the two-phase-compile idea FlyDSL floated and dropped (#266).
在线 autotune 最大的痛点是首次调用慢。 quack 的 _precompile (_compile_worker.py) 用子进程池并行编译所有候选配置、 在 benchmark 之前把 .o 缓存焐热, 用 FileLock 做跨进程安全。 这正是 FlyDSL 提过又放掉的两阶段编译想法 (#266) 的实用答案。
quack is the FlyDSL → downstream PoC target (RFC #749). Its autotuner is both the closest design analog (DSL, JIT, Triton-descended) and the most mature. If FlyDSL is to back a quack-style library, matching this shape is the bar.
quack 是 FlyDSL → 下游的 PoC 目标 (RFC #749)。 它的 autotuner 既是最接近的设计类比 (DSL、 JIT、 Triton 血统), 也是最成熟的。 如果 FlyDSL 要给一个 quack 风格的库当后端, 对齐这个形状就是门槛。
§ 05 · NVIDIA CUTLASS's Python DSLCuteDSL · autotune_jit
§ 05 · NVIDIA CUTLASS 的 Python DSLCuteDSL · autotune_jit
CuteDSL is the layer quack is built on, so it shows the autotune idea one level lower — closer to how FlyDSL itself would expose it. Its decorator autotune_jit(params_dict, update_on_change) (cutlass/cute/testing.py) takes the parameter grid directly, then on first call for a tuning key it sweeps the Cartesian product, compiles+benches each, and caches the best kernel in func._best_kernel[tuning_key].
CuteDSL 是 quack 构建其上的那一层, 所以它把 autotune 的想法展示在更低一层 —— 更接近 FlyDSL 自己会怎么暴露它。 它的装饰器 autotune_jit(params_dict, update_on_change) (cutlass/cute/testing.py) 直接接受参数网格, 然后在某个 tuning key 的首次调用时遍历笛卡尔积, 对每个编译+bench, 把最优 kernel 缓存在 func._best_kernel[tuning_key] 里。
The instructive part is the split between params_dict and update_on_change, because it answers a question every autotuner faces: which arguments should trigger a re-tune? params_dict is the search space — the knobs to sweep. update_on_change is the opposite: values that are not tuned, but whose change must still invalidate the cached winner (a dtype, a flag that alters the kernel's shape). Without that distinction you hit one of two bugs: put too much in the key and you re-tune on noise; put too little and you silently reuse a config tuned for a different situation. It is the same tension as Triton's cache key, surfaced here as an explicit knob.
有教益的是 params_dict 和 update_on_change 的分工, 因为它回答了每个 autotuner 都要面对的问题: 哪些参数该触发重新 tune? params_dict 是搜索空间 —— 要扫的旋钮。 update_on_change 相反: 那些不参与 tune、 但一旦变化就必须让缓存赢家失效的值 (一个 dtype、 一个会改变 kernel 形态的 flag)。 没有这个区分, 你会撞上两种 bug 之一: key 里放太多, 会为噪声反复重 tune; 放太少, 会静默复用一个为别的情形 tune 出来的配置。 这和 Triton cache key 是同一个张力, 在这里被暴露成一个显式旋钮。
Two CuteDSL details transfer directly to FlyDSL. First, the GEMM autotuning doc is explicit that compilation is the cost, so you cute.compile() once per config and reuse the JIT executor — exactly FlyDSL's build_*_module situation, where the build is the expensive step. Second, the implicit JIT cache (dsl_jit_caching.rst) keys on a hash of the source plus all DSL environment variables — the same "env is part of the key" lesson Triton encodes — and CUTE_DSL_LOG_AUTOTUNE makes the search observable.
两个 CuteDSL 细节可以直接搬到 FlyDSL。 第一, GEMM autotuning 文档明确说编译才是成本所在, 所以每个配置 cute.compile() 一次、 复用 JIT executor —— 这正是 FlyDSL 的 build_*_module 处境, build 是那个昂贵的步骤。 第二, 隐式 JIT 缓存 (dsl_jit_caching.rst) 的 key 是源码加上所有 DSL 环境变量的哈希 —— 和 Triton 编码的"env 是 key 的一部分"是同一课 —— 而 CUTE_DSL_LOG_AUTOTUNE 让搜索过程可观测。
§ 06 · AMD's kernel libraryaiter · offline CSV
§ 06 · AMD 的 kernel 库aiter · 离线 CSV
aiter is the opposite paradigm from the JIT autotuners: tuning is offline and checked in. A tuner reads an untuned shape list (aiter/configs/*_untuned_gemm.csv), sweeps configs for each shape, and writes the winners to *_tuned_gemm.csv committed to the repo; per-model variants live under aiter/configs/model_configs/. At serving time there is no search — the kernel reads its config from the CSV. A CI guard, check_tuned_op_regression.sh, keeps the committed configs from regressing.
aiter 是和 JIT autotuner 相反的范式: tuning 是离线的、 入库的。 一个 tuner 读取未 tune 的形状列表 (aiter/configs/*_untuned_gemm.csv), 为每个形状扫配置, 把赢家写进提交到仓库的 *_tuned_gemm.csv; per-model 变体放在 aiter/configs/model_configs/。 服务时没有搜索 —— kernel 从 CSV 读它的配置。 一个 CI 守卫 check_tuned_op_regression.sh 防止入库的配置退化。
aiter ships a FlyDSL-backed tuning path at aiter/ops/flydsl/gemm_tune/ (e.g. gemm_a16w16_tune.py imports flydsl). So the question "how does an external lib expect to tune a FlyDSL kernel?" already has a partial answer in tree — and it is the offline-CSV shape. This ties directly to RFC #749.
aiter 自带一条 FlyDSL 后端的 tuning 路径, 在 aiter/ops/flydsl/gemm_tune/ (比如 gemm_a16w16_tune.py 就 import 了 flydsl)。 所以"外部库期待怎么 tune 一个 FlyDSL kernel?"这个问题在代码里已经有了部分答案 —— 而且是离线 CSV 的形状。 这直接关联到 RFC #749。
§ 07 · Two timelinesOnline vs offline
§ 07 · 两条时间线在线 vs 离线
The JIT autotuners (Triton, quack, CuteDSL, FlyDSL) tune online — at the first call for a shape, in the user's process. aiter and SGLang tune offline — ahead of time, results committed, looked up at serving. Neither is strictly better; they answer different constraints. Online covers any shape that shows up, including ones nobody enumerated, and is friendly to iteration; it pays a slow first call. Offline pays nothing at serving and is reproducible and CI-gateable; but it only covers shapes someone tuned in advance, and every new GPU or kernel change means re-running the offline sweep.
JIT autotuner (Triton、 quack、 CuteDSL、 FlyDSL) 在线 tune —— 在某个形状首次调用时、 在用户进程里。 aiter 和 SGLang 离线 tune —— 提前算好、 结果入库、 服务时查表。 没有谁绝对更好; 它们回答不同的约束。 在线覆盖任何冒出来的形状, 包括没人枚举过的, 且对迭代友好; 代价是首次调用慢。 离线服务时零成本、 可复现、 可进 CI; 但它只覆盖有人提前 tune 过的形状, 而且每块新 GPU 或每次 kernel 改动都意味着要重跑离线 sweep。
§ 08 · How an engine consumes tuned configsSGLang · the consumer
§ 08 · 引擎怎么消费已 tune 的配置SGLang · 消费者
SGLang answers a different question from the others. It is not a kernel DSL and does not write an autotuner; it is an inference engine that consumes tuned configs. It is on this map because it is a real downstream FlyDSL consumer, and it shows exactly what the offline artifact should look like at scale — which is a direct design input for FlyDSL.
SGLang 回答的是和别人不同的问题。 它不是 kernel DSL, 也不写 autotuner; 它是一个消费已 tune 配置的推理引擎。 它在这张图上, 是因为它是 FlyDSL 真实的下游消费者, 而且它精确地展示了离线产物在规模上应该长什么样 —— 这对 FlyDSL 是直接的设计输入。
SGLang's fused-MoE configs are committed JSON whose filename is the lookup key: E=8,N=14336,device_name=NVIDIA_A100-SXM4-80GB,dtype=int8_w8a16.json, under a triton-version directory, with AMD gfx950 variants alongside the NVIDIA ones. At runtime get_moe_configs (fused_moe_triton_config.py) looks up by device + shape; on a miss it falls back to get_default_config — the heuristic-default idea again, this time as the safety net for an unseen shape.
SGLang 的 fused-MoE 配置是入库的 JSON, 文件名就是查找 key: E=8,N=14336,device_name=NVIDIA_A100-SXM4-80GB,dtype=int8_w8a16.json, 放在按 triton 版本分的目录下, NVIDIA 的旁边就是 AMD gfx950 变体。 运行时 get_moe_configs (fused_moe_triton_config.py) 按设备 + 形状查; miss 时 fallback 到 get_default_config —— 又是启发式默认的想法, 这次是给没见过的形状当安全网。
If FlyDSL kernels are to be tuned-and-consumed by an engine, the tuned-config artifact should be keyed by device + shape + dtype, discoverable as a file, and paired with a heuristic default for misses. SGLang shows the exact contract a downstream expects — the offline half of FlyDSL's design should target it.
如果 FlyDSL 的 kernel 要被一个引擎 tune-then-consume, 那么 tuned-config 产物应该按 设备 + 形状 + dtype 建 key、 可以作为文件被发现、 并配一个启发式默认兜底 miss。 SGLang 展示了下游期待的确切契约 —— FlyDSL 设计的离线那一半应该瞄准它。
Training vs inference — the cost tolerance differs
训练 vs 推理 — 成本容忍度不同
One nuance the consumer side reveals: training and inference tolerate first-run cost differently. A training run is long; a one-time online autotune at startup amortizes over millions of steps, so JIT tuning is fine. Inference serving is latency-sensitive and elastic; a slow first request per new shape is a visible regression, which is why engines like SGLang lean on committed offline configs with a default fallback. FlyDSL's design should let the same kernel be tuned online during a training/dev session and shipped as an offline config for inference.
消费侧揭示的一个微妙点: 训练和推理对首跑成本的容忍度不同。 训练跑得久; 启动时一次性的在线 autotune 会摊薄到几百万步上, 所以 JIT tuning 没问题。 推理服务对延迟敏感且弹性伸缩; 每个新形状的首个请求慢就是看得见的退化, 这正是 SGLang 这类引擎倚重入库的离线配置 + 默认兜底的原因。 FlyDSL 的设计应该让同一个 kernel 既能在训练/开发 session 里在线 tune, 又能作为离线配置发给推理。
§ 09 · What goes in the keyCache-key anatomy
§ 09 · key 里到底放什么cache key 解剖
A cache key is a contract: "two calls with the same key may reuse the same tuned config." Get it too narrow and you reuse a stale or wrong config; too wide and you never hit the cache. Every reference's key is a different answer to "what, if it changed, should force a re-tune?" — and FlyDSL's current key answers the least.
cache key 是一份契约: "两个 key 相同的调用可以复用同一份 tuned 配置。" 太窄, 你会复用过期或错误的配置; 太宽, 你永远命不中缓存。 每个参考的 key 都是对"什么东西一改就该强制重 tune?"的不同回答 —— 而 FlyDSL 当前的 key 回答得最少。
§ 10 · The destinationDesigning FlyDSL's autotune
§ 10 · 终点设计 FlyDSL 的 autotune
FlyDSL already has an autotuner — python/flydsl/autotune.py, 296 lines: a Config (constexpr params + waves_per_eu / maxnreg hints), an Autotuner (benchmark, pick fastest, disk cache under FLYDSL_AUTOTUNE_CACHE_DIR), a _make_key, a _prune, and a public @autotune. The problem is not that it is missing; it is that no kernel uses it — a repo-wide search finds only the definition and docs, no adoption, no tests. It is well-built dead code. Reading it against the five references above, here is the gap and the target design.
FlyDSL 已经有一个 autotuner —— python/flydsl/autotune.py, 296 行: 一个 Config (constexpr 参数 + waves_per_eu / maxnreg hint)、 一个 Autotuner (benchmark、 挑最快、 在 FLYDSL_AUTOTUNE_CACHE_DIR 下磁盘缓存)、 一个 _make_key、 一个 _prune、 和一个公开的 @autotune。 问题不是它缺失; 而是没有 kernel 用它 —— 全仓搜索只找到定义和文档, 没有采用、 没有测试。 它是构建良好的死代码。 把它对着上面五个参考读, 下面是差距和目标设计。
| Capability | 能力 | FlyDSL today | FlyDSL 今天 | Target (from refs) | 目标 (取自参考) |
|---|---|---|---|---|---|
| Adoption | 采用 | none | 1 kernel + offline config | ||
| Avoid-search path | 免搜路径 | none | heuristic default (quack / Triton @heuristics) | ||
| In-place correctness | 原地正确性 | reset_to_zero only | + restore_value snapshot (Triton) | ||
| Cache key | cache key | shape/dtype only | + arch + src fingerprint + compiler + env + stride | ||
| Offline artifact | 离线产物 | none | device+shape config file (SGLang-style) | ||
| Parallel precompile | 并行预编译 | none | subprocess pool (quack, optional) | ||
| CI guard | CI 守卫 | none | GPU-free unit + regression gate (aiter) |
build_*_module / launch_* structure.build_*_module / launch_* 结构上。§ 11 · A concrete pathImplementation guide
§ 11 · 一条具体路径实现 guide
Sequencing matters: each step is shippable on its own and de-risks the next. Below is the order I would take FlyDSL's autotune from dead code to adopted, mapped to the deliverables in #770.
顺序很重要: 每一步都能单独上线, 并为下一步降风险。 下面是我会采取的顺序, 把 FlyDSL 的 autotune 从死代码带到被采用, 对应 #770 里的交付物。
Step 1 · Harden the key, add restore_value
第 1 步 · 加固 key, 加上 restore_value
Before any kernel adopts autotune, fix the two correctness gaps. Extend _make_key (currently shape/dtype only, autotune.py:139) with arch, a source fingerprint, the compiler/version, and cache-invalidating env — the Plate-IV target. And add restore_value alongside the existing reset_to_zero (:163): snapshot listed tensors before the bench loop, restore after each rep. Without this, autotuning fused-add rmsnorm is incorrect.
在任何 kernel 采用 autotune 之前, 先修两个正确性缺口。 扩展 _make_key (当前只有 shape/dtype, autotune.py:139), 加上 arch、 源码指纹、 编译器/版本、 失效性 env —— 即图 IV 的目标。 并在已有的 reset_to_zero (:163) 之外加上 restore_value: 在 bench 循环前存档指定张量, 每遍之后还原。 没有这个, autotune fused-add rmsnorm 就是错的。
Step 2 · Two-track config for one kernel
第 2 步 · 给一个 kernel 做双轨配置
Pick one kernel — softmax or rmsnorm, small bounded space, good for validating the path. Give it a get_default(N, dtype, arch) heuristic and a get_all_configs() exhaustive space, quack-style. Normal callers get the heuristic for free; autotune is opt-in.
挑一个 kernel —— softmax 或 rmsnorm, 空间小而有界, 适合验证路径。 给它一个 get_default(N, dtype, arch) 启发式和一个 get_all_configs() 穷举空间, quack 风格。 普通调用者免费拿启发式; autotune 是 opt-in。
launch_softmax signature (kernels/softmax_kernel.py:244) and the existing @autotune (autotune.py:262). The default= and restore_value= args are the proposed additions.launch_softmax 签名 (kernels/softmax_kernel.py:244) 和已有的 @autotune (autotune.py:262)。 default= 和 restore_value= 是提议新增的参数。Step 3 · Offline emit + lookup (SGLang-style)
第 3 步 · 离线产出 + 查表 (SGLang 风格)
Add a mode where the online search, instead of only caching in-process, writes a config file keyed by device + shape + dtype — the SGLang contract. A serving run with tuning off looks that file up and falls back to get_default on a miss. This is what lets aiter's ops/flydsl/gemm_tune/ and SGLang consume FlyDSL kernels without paying search at serving.
加一个模式: 在线搜索不只在进程内缓存, 而是写一个按 设备 + 形状 + dtype 建 key 的配置文件 —— 即 SGLang 契约。 一个关掉 tuning 的服务运行查这个文件、 miss 时 fallback 到 get_default。 这正是让 aiter 的 ops/flydsl/gemm_tune/ 和 SGLang 能消费 FlyDSL kernel、 且服务时不付搜索成本的东西。
Step 4 · CI guard, GPU-free first
第 4 步 · CI 守卫, 先做无 GPU 的
Add GPU-free unit tests for Config serialization, cache-key construction, and pruning (these need no hardware and catch most regressions), then an opt-in GPU integration test, then a tuned-config regression gate modeled on aiter's check_tuned_op_regression.sh. Tuning must be opt-in via env/marker so normal CI never pays the search.
加无 GPU 的单元测试覆盖 Config 序列化、 cache-key 构造、 剪枝 (这些不需要硬件, 能抓住大部分回归), 然后一个 opt-in 的 GPU 集成测试, 然后一个仿 aiter check_tuned_op_regression.sh 的 tuned-config 回归门。 tuning 必须通过 env/marker opt-in, 这样正常 CI 永远不付搜索成本。
Step 5 · (optional) parallel precompile
第 5 步 · (可选) 并行预编译
If the online search proves too slow on real spaces, add quack-style subprocess precompile to warm the build cache for all candidates before benchmarking — the practical revival of the two-phase-compile idea from #266.
如果在线搜索在真实空间上太慢, 加 quack 风格的子进程预编译, 在 benchmark 前为所有候选焐热 build 缓存 —— 即 #266 两阶段编译想法的实用复活。
§ 12 · What to build firstEpilogue
§ 12 · 先造什么尾声
If you read this for one decision, it is this: do not start FlyDSL's autotune by adding configs. Start by making it correct and cheap to not use — the full cache key and restore_value (so tuning an in-place kernel is sound), and a heuristic default (so the common path pays nothing). Triton, quack, and CuteDSL all earn their autotuners by getting those two things right first; aiter and SGLang show that the payoff — committed, looked-up configs — is what an engine actually consumes. FlyDSL has the loop already. What it needs is the edges: the correctness soul, the skip-the-search door, and an offline artifact a downstream can read.
如果你为一个决定读这篇, 就是这个: 不要从加配置开始做 FlyDSL 的 autotune。 从让它正确且不用它也便宜开始 —— 完整的 cache key 和 restore_value (这样 tune 一个原地 kernel 才是可靠的)、 加一个启发式默认 (这样常见路径零成本)。 Triton、 quack、 CuteDSL 都是先把这两件事做对, 才配得上它们的 autotuner; aiter 和 SGLang 表明回报 —— 入库、 可查的配置 —— 才是一个引擎真正消费的东西。 FlyDSL 已经有那个循环了。 它需要的是边角: 正确性的灵魂、 跳过搜索的那扇门、 和一个下游能读的离线产物。
§ R · Sources readReferences
§ R · 精读来源参考资料
- Triton — runtime/autotuner.py (
Autotuner.run:217,check_disk_cache:175,restore_value/reset_to_zero,@heuristics,prune_configs_by). - quack — quack/autotuner.py, rmsnorm_config.py (two-track), caching_and_parallel_compilation.md, _compile_worker.py.
- CuteDSL — cute/testing.py (
autotune_jit:561), autotuning_gemm.rst, dsl_jit_caching.rst. - aiter —
aiter/configs/*_tuned_gemm.csv, check_tuned_op_regression.sh,aiter/ops/flydsl/gemm_tune/. - SGLang — fused_moe_triton_config.py (
get_moe_configs/get_default_config), committed JSON configs,.github/workflows/auto-tune.yml. - FlyDSL — python/flydsl/autotune.py (296 ln); issues #770 (autotune), #769 (rmsnorm bwd), #749 (downstream RFC), #266 (two-phase compile).