Every kernel autotuner I've used searches the same kind of thing. Tile shapes, warp counts, pipeline stages, whether to use TMA descriptors or plain pointer math. Those are all choices about what work each thread block does. Once the winning config is picked, the generated PTX goes to ptxas and gets compiled at default settings, the same as every other kernel on the planet.

A PyTorch blog post published Oct 19, right before PyTorch Conference in San Jose, shows that layer is now searchable from inside Helion. The search itself comes from NVIDIA's CompileIQ, and the result is a small binary file you commit next to your kernel. The numbers are modest, and the authors say so up front, which is the main reason I trust them.

What is actually being tuned

Starting with CUDA 13.3, ptxas and nvcc accept an Advanced Control File through --apply-controls. An ACF sets internal compiler decisions that no public flag exposes: register allocation policy, fine-grained instruction scheduling, loop transformation thresholds and other code generation heuristics. NVIDIA's CompileIQ launch post from May describes it as a search over those internals, with the internals kept opaque. You never see what the knobs are called. You get a blob that produces a different binary.

CompileIQ treats the choice as a hyperparameter search. It pulls a curated search space for your compiler and toolkit version, then runs an evolutionary loop: propose a population of control sets, compile each one, check correctness, time the survivors, breed the winners, repeat. You write the objective function. In most cases that's a few lines around triton.testing.do_bench.

It installs with pip install compileiq. It's free to use but ships under the NVIDIA Software License Agreement, which is a proprietary EULA. The ACFs it produces are freely redistributable, so you can commit them to an open repo. Version 1.1.0 landed on PyPI and GitHub on Oct 16. The Helion experiments in the post used 1.0.3.

Why the search has to be joint

This is the part of the post I'd want every kernel engineer to read. Helion now exposes ACFs in two ways.

  • advanced_controls_file= inside a helion.Config pins one known-good ACF to a config you've already chosen.
  • autotune_search_acf=[...] hands a list of ACFs to Helion's own autotuner, which treats them as one more dimension next to block_sizes, num_warps, num_stages, indexing and the rest.

The second mode is the interesting one, and the reasoning holds up. If you autotune the kernel first and then search compiler controls, you're searching conditioned on a config that was chosen under default controls. Block size and stage count set register pressure and shared memory footprint. Register allocation policy decides whether that pressure turns into spills. A control set that would win at block_sizes=[64, 64] never gets tried there, because you froze [128, 128] before the compiler search started. Doing it the other way round has the same problem. Searching both at once is the only order that avoids it.

Two details make the joint mode safe to try. Helion always appends the empty string to your ACF list, meaning "compiler defaults," so the baseline is always a candidate and an ACF only wins if it measures faster. The worst case is that you burn autotuning time and get your old kernel back. The cost is real, though: every ACF multiplies compile-and-measure trials. The intended workflow is to run it once, offline, on a kernel that matters, then pin the winner the way you already pin autotuned configs.

Most people won't run CompileIQ at all. NVIDIA publishes prebuilt ACF catalogs called Booster Packs. The current helion-booster-pack has 24 ACFs tuned across a range of Helion kernels, validated against ptxas from CUDA 13.4 and targeting B200. Download it, point autotune_search_acf at it, autotune once, and keep what wins.

The baseline is already tuned

What makes the results worth reading is the baseline. The post doesn't start from naive kernels. It starts from Helion's pretuned_kernels/, about 27 kernels that cover shapes real models hit: FP8 and NVFP4 quantization, RMSNorm variants, RoPE, grouped GEMM, causal depthwise convolution, Gated DeltaNet decode, and fused regions like a Qwen3 decode layer. Each one ships ahead-of-time config files per architecture, such as _helion_aot_rms_norm_cuda_sm90.py for H100 and _sm100.py for B200. The search over the kernel space has already been done by the people who wrote them.

The case study is nvfp4_gemv_fp4in_kernel, a W4A4 decode GEMV (NVFP4 weights times NVFP4 activations) at the N=8192, K=28672 down_proj shape. Its checked-in config is block_sizes [4, 256], num_warps 2, num_stages 3. The setup was a GB200 (sm_100) system, ptxas 13.3.73, CompileIQ 1.0.3 and Triton 3.8.0. Clocks were locked, every candidate was checked against the PyTorch reference, and each gain was re-measured and filtered with a Welch's t-test (p < 0.05) and Cohen's d (d > 0.2).

Here's what they found:

  • Runtime as the objective (8 generations, pool of 15): 0.0971 ms down to 0.0928 ms in the search loop. Two independent re-runs landed at 1.0485x and 1.0478x, with d around 2.5.
  • Occupancy as the objective: Nsight Compute showed the kernel was register-limited, with 25.0% theoretical and 21.63% achieved occupancy, and NCU estimated 21% headroom. Searching on achieved occupancy (4 generations, pool of 6) moved it only from 21.60% to 21.76%. That still turned into a re-measured runtime gain of about 1.75 to 2.0%.
  • Multi-objective, runtime plus L1TEX throughput: the search returned one candidate that beat every other evaluated candidate on both axes. It was 1.0689x faster, and L1TEX throughput went from 81.51% to 85.96%.
  • One ACF across two shapes: a joint search over down_proj and N=15360, K=5120 produced a four-point Pareto front. The chosen candidate validated at 1.05x on the first shape and 1.01x on the second.

The occupancy result is the one I keep coming back to. NCU's cost model pointed at a 21% opportunity, and a direct search on that metric barely moved it, yet runtime still improved. Proxy metrics guide you, but the clock is what counts. It's also a fair reminder that the knobs CompileIQ can reach don't add up to "pick fewer registers."

How big this is, honestly

The post is clear about scale: on kernels that are already well tuned, compiler-control tuning has shown roughly 2 to 7% runtime gains, and less-optimized kernels have seen double digits. NVIDIA's May post says Meta has seen up to 15% on TritonBench and Helion kernels, citing a GTC talk. That's NVIDIA reporting a partner's result, and I haven't reproduced it.

Five percent on one GEMV isn't a headline. Five percent on the handful of GEMV, GEMM and attention kernels that dominate every decode step of a large serving fleet is real capacity, and it costs nothing at runtime because the change is baked into the binary. NVIDIA's own estimate is that linear-layer GEMMs make up about 70% of LLM inference FLOPs and attention variants about another 25%. That's where I'd spend the search budget.

What I'd watch before shipping ACFs

The caveats matter more than the speedups, so here they are plainly.

  1. ACFs are tied to the architecture and toolkit. An sm_90 ACF isn't valid on sm_100. The post warns that a mismatched ACF can give wrong results or runtime errors, not just slower code. So an ACF needs the same discipline as a compiled artifact: key it on GPU arch and exact CUDA toolkit version, and fail loudly on a mismatch.
  2. Keep the correctness gate in CI, not only in the search. The search drops candidates that fail the reference check, but the shapes you tuned on aren't every shape you'll serve. Put the pinned ACF through the same numeric tests as the kernel, at each shape you deploy.
  3. Mind the version skew. The Booster Pack says CUDA 13.4 and B200, while the case study ran on ptxas 13.3.73 on GB200. Neither is wrong, but you're the one who has to check that the pack manifest matches your toolchain image.
  4. You can't read what it changed. The controls are opaque by design. When a kernel regresses after a driver or toolkit bump, you'll diff binaries and benchmarks, not settings. Commit the ACF with the benchmark that justified it so someone can re-run that check later.
  5. Budget the search. Each ACF in autotune_search_acf multiplies trials. Run it offline on your top few kernels and pin the result. Don't turn it on across a whole codebase.

The broader point is the one the authors end on, and I agree with it. The compiler's control surface is a real optimization space, it can be searched, and almost nobody is searching it yet. Helion putting it in the same loop as tile sizes is what moves it from something you could do in principle to something you'll actually ship.

Sources

Anpoo Sivanadi, Staff Software Engineer