PyTorch Helion can now tune NVIDIA compiler controls alongside GPU kernel settings

The layer below the kernel

According to the PyTorch blog post, Helion was released last year so that developers could write GPU kernels in PyTorch-native Python and leave the tuning work to its autotuner. The authors say the autotuner explores thousands of kernel configurations and routinely lands within a few percent of expert-tuned code, and often exceeds it.

The post argues that room for improvement remains even after that. Kernel autotuners, whether in Helion, Triton or torch.compile, search the same things: tile shapes, warp counts, pipeline stages and memory access strategy. A second set of decisions sits underneath, covering how the work becomes machine instructions generated by ptxas. Register allocation policy, fine-grained instruction scheduling and code-generation heuristics stay at their defaults for all shipping kernels. The authors say this is not because the defaults are optimal, but because nothing was searching them.

This compiler control surface is exposed by ptxas itself starting from CUDA 13.3. The post notes that exploring it is hard, because the controls interact: a change to register allocation shifts occupancy, which can change whether a scheduling decision helps or hurts. The right setting is workload-specific.

How the joint search works in Helion

Helion exposes the compiler controls through two settings. An ACF is a binary artifact applied at compile time, portable and versionable, which users can ship alongside their kernel.

• advanced_controls_file pins a known-good ACF inside a config the user has already chosen.
• autotune_search_acf takes a list of ACFs and lets Helion's autotuner search them as one more dimension, next to tile sizes and warp counts.

In the search mode, Helion automatically adds an empty string to the list, meaning no ACF and compiler defaults. The default is therefore always a candidate, and an ACF is only chosen if it wins on measured time. According to the post, the worst case is spending some autotuning budget and getting back the kernel you already had.

The authors argue that tuning in one loop beats tuning in two steps. Block size and stage count determine register pressure and shared-memory footprint, and register allocation policy decides whether that pressure causes spills. If the kernel config is frozen first, a control set that would work well with a different config is never evaluated. The reverse ordering has the same problem. That is why the search lives inside the autotuner rather than being a post-processing step.

The cost is search time, since each ACF multiplies the number of compile-and-measure trials. The post recommends running it once, offline, on a kernel that matters, then pinning the result with advanced_controls_file, the way autotuned configs are pinned today.

Where the ACFs come from

NVIDIA's CompileIQ generates the files. It treats compiler-control selection as a hyperparameter optimization problem, pulling a curated search space for the target compiler and toolkit version. It then runs an evolutionary search: propose a population of candidate control sets, compile each, check correctness against a reference, time the survivors, breed the winners and repeat. The output is an ACF.

The search algorithm has nothing GPU-specific in it. The GPU-specific part is the objective function the user writes, usually a few lines around triton.testing.do_bench. The post says most users will not need to run this search, because CompileIQ publishes pre-tuned ACF catalogs that Helion's autotuner can evaluate alongside its existing kernel configurations.

A deliberately hard baseline

The authors stress that the starting point matters. Their baseline is the pretuned_kernels/ directory in the Helion repository, roughly 27 kernels covering shapes from real model code: FP8 and NVFP4 quantization, RMSNorm variants, RoPE, grouped GEMM, causal depthwise convolution, Gated DeltaNet decode and fused regions such as a Qwen3 decode layer.

These kernels are not tuned on the fly. Each ships with ahead-of-time heuristic files per architecture, for example one for H100 and one for B200, which Helion selects at runtime. The search happened offline, so users get tuned performance with no search cost. Any gain beyond that, the post says, is real headroom.

What it means for developers

For developers who already rely on Helion, compiler-level tuning becomes an option inside the same autotuning workflow. Because the default stays in the candidate list, trying it on a performance-critical kernel carries little risk beyond search time. Teams that do not want to run the search can use CompileIQ's published ACF catalogs instead.