Optimizing GPU Kernels with Helion and NVIDIA CompileIQ
Achieving deeper performance gains by jointly tuning kernel configurations and compiler assembly controls.
1Beyond Standard Kernel Autotuning
Helion’s autotuner traditionally optimizes tile shapes, warp counts, and memory strategies to improve GPU kernel performance. While these configurations often match expert-tuned code, they leave lower-level compiler decisions untouched. By targeting NVIDIA’s ptxas assembler controls, developers can now optimize register allocation and instruction scheduling, unlocking performance headroom previously considered inaccessible to standard autotuners.
2The Role of Compiler Controls
Standard kernel autotuners stop at Triton configurations, leaving ptxas settings at default values. These defaults are rarely optimal for every specific workload. Helion now exposes this second layer of control, allowing users to influence how machine assembly instructions are generated. This interface, introduced in CUDA 13.3, enables fine-grained hardware-level optimization.
3Joint Optimization Strategy
Searching kernel configurations and compiler controls separately leads to suboptimal local results because these dimensions are tightly coupled. Helion’s joint search method evaluates both spaces simultaneously. This ensures that register pressure and shared-memory footprints are balanced against specific ptxas settings, preventing the performance bottlenecks that occur when one layer is frozen.
4Advanced Control Files (ACF)
Compiler settings are packaged into portable, versionable binary artifacts called Advanced Control Files. Users can pin a known-good ACF to a kernel configuration for consistent production performance. Helion allows passing a list of ACFs to its autotuner, which evaluates them alongside tile sizes to find the best overall execution time.
5NVIDIA CompileIQ Integration
CompileIQ automates the generation of ACFs by treating compiler control selection as a hyperparameter optimization problem. It uses an evolutionary search algorithm to propose, compile, and benchmark candidate control sets. This process identifies the most efficient assembly configurations for specific hardware architectures, packaging the successful results into reusable Booster Packs.
6Booster Packs for Users
Most users do not need to run the search themselves. CompileIQ publishes pre-tuned Booster Packs containing 24 ACFs validated against CUDA 13.4 and B200 hardware. Developers simply download these catalogs, point Helion’s autotuner to the files, and perform a single offline search to pin the best-performing configuration for their specific workload.
7Experimental Methodology
Researchers tested this approach on a GB200 system using sm_100 compute capability. The study focused on the nvfp4_gemv_fp4in_kernel, which already utilized a checked-in AOT configuration. By applying CompileIQ to this baseline, the team sought to determine if additional performance could be extracted specifically at the ptxas compilation stage during the decode process.
8Performance Benchmarking Results
The evolutionary search spanned eight generations with a pool size of 15 candidates. By coupling the scoring routine to the PTXAS search space, the team reduced execution time from 0.0971 ms to 0.0928 ms. This represents a 1.047x speedup, demonstrating that even highly optimized kernels retain significant room for last-mile performance improvements.
9Correctness and Validation
Every configuration candidate undergoes strict numerical validation against a PyTorch reference implementation before benchmarking. Only configurations passing this correctness gate are timed using triton.testing.do_bench. Statistical filtering, including Welch’s t-test with p < 0.05 and Cohen’s d, ensures that reported speedups are genuine and not artifacts of measurement noise or variance.
10Search Cost Considerations
Joint searching increases the number of compile-and-measure trials, making it unsuitable for default use in every codebase. It is designed as an offline, one-time optimization step for critical kernels. Once the optimal ACF is identified and pinned, the runtime performance remains high with zero additional search cost during production deployment.
11Handling Custom Workloads
For workloads not covered by existing Booster Packs, developers can install the CompileIQ package directly. The process involves defining an objective function, typically a few lines of code around a benchmark, and initiating the search. This path is fully documented, allowing teams to extend the optimization benefits to proprietary or niche model architectures.
12Future of Kernel Optimization
The integration of Helion and CompileIQ represents a shift toward holistic compiler-aware development. By automating the exploration of both high-level kernel parameters and low-level assembly controls, developers can achieve hardware-specific performance without deep compiler expertise. This methodology sets a new standard for optimizing the last mile of GPU-accelerated deep learning workloads.
Original article: PyTorch Blog — Optimizing GPU Kernels with Helion and NVIDIA CompileIQ