一次搜索，双重空间：在 PyTorch Helion 中使用 NVIDIA CompileIQ 控制编译器

精选项目Featured projects

Helion 去年发布时，旨在让开发者能够使用 PyTorch 原生的 Python 编写 GPU 内核，并由其强大的自动调优器完成过去需要手写 Triton 才能完成的工作，同时屏蔽硬件特定的实现细节。Helion 的自动调优器会探索数千种内核配置，通常能达到专家调优代码性能的几个百分点以内，而且往往还能超越它。With Helion’s release last year, it aimed at enabling developers to write GPU kernels in PyTorch-native Python and let its powerful autotuner do the work that used to require hand-written Triton and abstract away hardware-specific implementation details. Helion’s autotuner explores thousands of kernel configurations and routinely lands within a few percent of expert-tuned code, and oftentimes exceeding it.

现在需要回答的问题是：如果一个内核已经经过自动调优，并且开发者已经尽其所能将其优化到最佳状态，还有什么可做的吗？The question that now needs to be answered is this: If a kernel has already been autotuned and optimized to the best of a developer’s ability, is there anything left?

确实如此。它存在于内核之下，在编译器中。There is. It lives below the kernel, in the compiler.

生态系统中的每个内核自动调优器，无论是 Helion 的、Triton 的，还是 torch.compile 内部的，搜索的都是同一类东西：瓦片形状、warp 数量、流水线阶段、内存访问策略。这些决策全都关乎每个线程块要执行哪些工作。在此之下，还有另一组决策，决定这些工作如何转化为由 NVIDIA 的 PTX 汇编器 ptxas 生成的机器汇编指令。寄存器分配策略、细粒度指令调度、代码生成启发式等设置，如今发布的所有内核都仍采用默认值——不是因为默认值最优，而是因为此前根本没有人搜索并优化这些设置！Every kernel autotuner in the ecosystem whether it is Helion’s, Triton’s, the one inside torch.compile, they all search the same thing: tile shapes, warp counts, pipeline stages, memory access strategy. All of those are decisions about what work each thread block does. Underneath them sits a second set of decisions about how that work becomes machine assembly instructions generated by ptxas, NVIDIA’s PTX assembler. Things like register allocation policy, fine-grained instruction scheduling, code-generation heuristics are left at their defaults for all kernels shipping today – not because the defaults are optimal, but because nothing was searching and optimizing them!

Helion 现在支持这一点，而本文将介绍这个第二个特定于编译器的搜索空间包含什么，以及为什么将它与内核配置搜索空间联合搜索，能进一步提升那些此前已被认定为最优的内核的性能。Helion now enables this, and this post is about what that second compiler-specific space contains, why searching it jointly with the kernel configuration space shows further improvements on kernels that were already deemed to be optimal.

搜索由 NVIDIA 的 CompileIQ 执行，它会像自动调优器优化分块大小那样优化 ptxas 控制项，并将结果打包为可移植的高级控制文件（ACF）。大多数用户不会亲自运行这项搜索：CompileIQ 会发布预调优的 ACF 目录，Helion 的自动调优器现在可以将其与已有探索的内核配置一并评估。The searching is done by NVIDIA’s CompileIQ, which optimizes ptxas controls the way an autotuner optimizes tile sizes and packages the result as a portable Advanced Control File (ACF). Most users won’t run that search themselves: CompileIQ publishes pre-tuned ACF catalogs that Helion’s autotuner can now evaluate alongside the kernel configurations it already explores.

从高度优化的内核入手，并优化最后一公里Starting off with highly-optimized kernels & optimizing the last-mile

与什么基准进行比较至关重要。如果起点是一个朴素内核，那么声称“我们让内核变快了”很容易；但对于现实世界的工作负载来说，这种说法几乎毫无意义。It matters a great deal what baseline one compares against. A “we made kernels faster” claim is easy if the starting point is a naive kernel, and close to meaningless for real-world workloads.

我们这里使用的核函数来自 Helion 仓库中的 pretuned_kernels/，大约有 27 个，覆盖了真实模型代码中常见的张量形状：FP8 和 NVFP4 量化、RMSNorm 变体、RoPE、分组 GEMM、因果深度卷积、Gated DeltaNet 解码，以及 Qwen3 解码层等融合模型区域。The kernels we use here come from the pretuned_kernels/ in the Helion repository, roughly 27 kernels covering the shapes that show up in real model code: FP8 and NVFP4 quantization, RMSNorm variants, RoPE, grouped GEMM, causal depthwise convolution, Gated DeltaNet decode, and fused model regions like a Qwen3 decode layer.

需要注意的重要一点是，这些并不是在运行时即时进行自动调优的。每种架构都会随附预先生成的启发式文件：H100 对应 _helion_aot_rms_norm_cuda_sm90.py，B200 对应 _helion_aot_rms_norm_cuda_sm100.py，等等；Helion 会在运行时选择相应文件。自动调优搜索早已离线完成，搜索结果也已预先固化。运行此程序时，用户无需承担搜索成本，就能获得经过调优的性能。The important property to note here is that these are not autotuned on the fly. Each one ships with ahead-of-time heuristic files per architecture: _helion_aot_rms_norm_cuda_sm90.py for H100, _helion_aot_rms_norm_cuda_sm100.py for B200 etc. that Helion selects at runtime. The autotuning search already happened offline, and its result is baked in. When this program is run, the user gets tuned performance with no search cost.

这就是我们的基准：内核作者和 Helion 编译器开发人员已经在该内核运行的硬件上，全面探索过它的配置空间。如果在此基础上还有提升空间，那就是真正的提升空间，也是本文聚焦的最后一公里优化。That is our baseline: a kernel whose configuration space has already been swept, on the hardware it’s running on, by the kernel authors and the Helion compiler developers. If there’s headroom below that, it’s real headroom and that’s the last-mile focus of this post.

Helion 的自动调优器会搜索什么，以及搜索止于何处What Helion’s autotuner searches, and where it stops

Helion 内核就是普通的 PyTorch，只多了一个新概念——分块（tiles）：A Helion kernel is ordinary PyTorch with one new idea – tiles:

import torch, helion, helion.language as hl

@helion.kernel()
def softmax(x: torch.Tensor) -> torch.Tensor:
    n, _m = x.size()
    out = torch.empty_like(x)
    for tile_n in hl.tile(n):
        out[tile_n, :] = torch.nn.functional.softmax(x[tile_n, :], dim=1)
    return out

用户从未指定 tile 的大小。这正是关键所在：hl.tile 声明迭代空间采用分块方式，剩下的交给自动调优器决定。单个内核定义会隐式生成一个包含数千种 Triton 配置的搜索空间，涵盖 block_sizes、loop_orders、l2_grouping、num_warps、num_stages、索引方式（指针算术、块指针或 TMA 描述符）、pid_type、reduction_loops 等。Helion 通过 TorchInductor 将代码降级为 Triton，为每个候选配置编译并计时，最后保留性能最佳的配置。The user never says how big a tile is. That’s the point: hl.tile declares that the iteration space is tiled, and the autotuner decides the rest. A single kernel definition implicitly generates a search space of thousands of Triton configurations that span block_sizes, loop_orders, l2_grouping, num_warps, num_stages, indexing (pointer arithmetic vs. block pointers vs. TMA descriptors), pid_type, reduction_loops, and more. Helion lowers through TorchInductor to Triton, compiles each candidate, times it, and keeps the winner.

随后，胜出的 Triton 内核会交给 ptxas 处理，从这一步开始的所有操作都使用 PTXAS 的默认设置。内核自动调优器的搜索空间由内核语言能够表达的内容决定，而没有任何内核语言能够表达“以不同方式分配寄存器”。这个编译器控制接口是独立的，自 CUDA 13.3 起由 ptxas 本身提供。Then the winning Triton kernel goes to ptxas, and everything from that point down happens at PTXAS’ default settings. The kernel autotuner’s search space is defined by what the kernel language can express, and no kernel language expresses “allocate registers differently”. This compiler control surface is a separate interface, exposed by ptxas itself starting from CUDA 13.3.

必须指出，这项探索确实非常困难。各种编译器控制选项之间存在复杂的相互作用：寄存器分配的变化会影响占用率，进而改变某项调度决策究竟能带来帮助还是造成损害。正确的设置取决于具体工作负载，并没有可供公布的通用答案。而且，投入单位人工所获得的收益很低，这正是适合交由搜索算法处理的问题特征。It is imperative to mention that this exploration is genuinely hard. There are complex interactions between various compiler controls: a change to register allocation shifts occupancy, which changes whether a scheduling decision helps or hurts. The right setting is workload-specific, and there’s no universal answer to publish. And the payoff per unit of human effort is poor, which is exactly the profile of a problem one hands off to a search algorithm.

联合搜索：Helion 自动调优器与 PTXAS 的高级控制The joint search: Helion Autotuner and PTXAS’ Advanced Controls

Helion 通过两种设置提供对底层编译器控制空间的访问。控制项打包在高级控制文件（Advanced Control File，简称 ACF，通常以 .acf 结尾）中。该文件是一种在编译时应用的二进制制品，便于移植和版本管理，用户可以将其与内核一同发布。Helion exposes the underlying compiler control space through two settings. Controls are packaged in an Advanced Control File (ACF, typically ending with .acf) which is a binary artifact applied at compile time, portable and versionable, that users can ship alongside their kernel.

第一个设置会将一个已知可用的 ACF 固定为用户已选择的配置的一部分：The first setting pins a known-good ACF as part of a config the user has already chosen:

softmax_tuned = helion.kernel(
    config=helion.Config(
        advanced_controls_file="h100_example1.acf",
        block_sizes=[4],
        num_warps=8,
        num_stages=2,
        reduction_loops=[512],
        indexing=["pointer", "pointer", "tensor_descriptor",
        "tensor_descriptor", "pointer", "pointer"],
    ),
    static_shapes=True,
)(_softmax)

第二种情况更有意思。它将一组 ACFs 交给 Helion 自带的自动调优器，让它们像其他维度一样参与搜索，与分块大小和 warp 数量并列：The second is a more interesting case. It hands a list of ACFs to Helion’s own autotuner and lets them be searched as just another dimension, alongside tile sizes and warp counts:

softmax_search = helion.kernel(
    autotune_search_acf=[
        "acfs/helion-booster-pack/candidate_03.acf",
        "acfs/helion-booster-pack/candidate_11.acf",
        "acfs/helion-booster-pack/candidate_17.acf",
    ],
    autotune_compile_timeout=120,
    static_shapes=True,
)(_softmax)

这种模式下有两点值得指出：Two things are worth pointing out in this mode:

• 空字符串会自动为你追加。Helion 会自动将“”（表示“无 ACF，使用编译器默认设置”）添加到提供的列表中。因此，默认配置始终是搜索候选项，自动调优器只有在 ACF 的实测耗时确实更优时才会选择它。开启此功能最坏的结果，就是花掉一些自动调优预算，最后得到的还是你原本的内核。正是这一点，让人很想在自己关心的内核上试一试。
  The empty string is appended for you. Helion automatically adds “”, meaning “no ACF, compiler defaults”, to the list that is supplied. The default configuration is therefore always a candidate in the search, and the autotuner will only select an ACF if it actually won on measured time. The worst case of turning this on is that you spend some autotuning budget and get back the kernel you already had. That property is what makes it tempting to try on a kernel you care about.
• 顺序调优只能找到局部最优；这种联合方法则不会。这正是应该在一个循环中完成，而不是分成两个循环的实质性理由。假设你先自动调优内核，得到 block_sizes=[128, 128] 和 num_stages=4，然后再去搜索编译器控制项。此时，你搜索的是一个受内核配置制约的控制空间，而这个内核配置本身是在默认控制项下选定的。但这些维度会通过硬件相互耦合：块大小和阶段数决定了寄存器压力与共享内存占用，而寄存器分配策略恰恰决定了这种压力是否会导致溢出。对于 block_sizes=[64, 64] 来说可能表现出色的一组控制项，根本不会在该配置下接受评估，因为在编译器搜索开始之前，内核配置就已经固定了。反过来先搜索也有同样的问题。只有同时搜索两者，才能避免这个问题；这正是 autotune_search_acf 位于 Helion 自动调优器内部，而不是作为后处理步骤的原因。
  Sequential tuning finds a local optimum; this joint method doesn’t. This is the substantive argument for doing it in one loop rather than two. Suppose you autotune the kernel first, get block_sizes=[128, 128] with num_stages=4, and then go searching for compiler controls. You are now searching the control space conditioned on a kernel config that was itself chosen under default controls. But those dimensions are coupled through the hardware: block size and stage count determine register pressure and shared-memory footprint, and register allocation policy is precisely what decides whether that pressure causes spills. A control set that would shine at block_sizes=[64, 64] never gets evaluated there, because the kernel config was frozen before the compiler search began.The reverse ordering has the same problem. Searching both together is the only ordering that doesn’t have it, which is why autotune_search_acf lives inside Helion’s autotuner rather than being a post-processing step.

代价是搜索时间。每个 ACF 都会使编译和测量的试验次数成倍增加，因此不应默认对代码库中的每个内核都启用它。应针对重要的内核离线运行一次，然后通过 advanced_controls_file 固定结果，方式与如今固定自动调优配置完全相同。The cost is search time. Each ACF multiplies the number of compile-and-measure trials, so this is not something to leave on by default for every kernel in a codebase. It is something to run once, offline, on a kernel that matters, and then pin the result with advanced_controls_file, exactly the way autotuned configs are pinned today.

这些高级控制文件是从哪里来的？Where do these Advanced Control Files come from

既然我们已经了解了将 Helion 参数与 PTXAS 的高级控制文件（ACF）联合自动调优的好处，一个自然而然的问题便是：这些 ACF 要如何生成？这正是 NVIDIA CompileIQ 发挥作用的地方。Now that we’ve understood the benefits of joint autotuning of Helion parameters with PTXAS’ advanced control files (ACF), the natural question one may have is how can these ACFs be generated. That is where NVIDIA’s CompileIQ comes into the picture.

CompileIQ 将编译器控制项的选择视为一个超参数优化问题。它会针对目标编译器和工具包版本提取一个精心整理的搜索空间，然后运行进化搜索：提出一组候选控制项配置，对每个配置进行编译，根据参考实现检查正确性，为通过检查的配置计时，再对优胜者进行繁衍并重复这一过程。该搜索过程的输出是一个 ACF 文件。搜索算法本身没有任何 GPU 专属之处；GPU 专属的部分仅存在于用户编写的目标函数中，通常就是围绕 `triton.testing.do_bench` 写几行代码。CompileIQ treats compiler-control selection as a hyperparameter optimization problem. It pulls a curated search space for the target compiler and toolkit version, then runs an evolutionary search: propose a population of candidate control sets, compile each one, check correctness against a reference, time the survivors, breed the winners, repeat. The output of this search process is an ACF file. There is nothing GPU-specific about the search algorithm itself, and the only GPU-specific parts are what is entirely in the objective function that the user writes, which is usually a few lines around triton.testing.do_bench.

不过，对大多数用户来说，额外的搜索预算并非必要，因为搜索工作已经完成。CompileIQ 会在其 GitHub releases 上发布 Booster Packs（预先调优的 ACF 目录）。当前的 helion-booster-pack 包含 24 个 ACF，针对一系列 Helion 内核进行了调优，并通过 CUDA 13.4 中的 ptxas 验证，目标平台为 B200。因此，对用户来说，实际的工作流程是：下载该包，将 Helion 的 autotune_search_acf 指向它，运行一次自动调优，然后固定胜出的配置。无需设计搜索空间，无需编写目标函数，也不需要编译器方面的专业知识！For most users however, the additional search budget is unnecessary, because the searching has been done. CompileIQ publishes Booster Packs – pre-tuned ACF catalogs – on its GitHub releases. The current helion-booster-pack contains 24 ACFs tuned across a range of Helion kernels, validated against ptxas from CUDA 13.4, targeting B200. So the realistic workflow for users is: download the pack, point Helion’s autotune_search_acf at it, autotune once, pin what wins. No search space to design, no objective function to write, no compiler expertise required!

如果你的工作负载不在该测试包的覆盖范围内，可以运行 pip install compileiq，然后自行开展有针对性的搜索。NVIDIA 的 CompileIQ 博客文章和各种 CompileIQ 示例中都介绍了这一方法。If you have a workload the pack doesn’t cover, pip install compileiq and run your own targeted search. That path is documented in NVIDIA’s CompileIQ blog post and in various CompileIQ examples.

Helion and CompileIQ: Experimental Results

本节重点介绍了在 Helion 内核上运行 CompileIQ 优化活动的多种可能方式。作为一个示例案例，我们将 CompileIQ 应用于预先调优的 nvfp4_gemv_fp4in_kernel。这是一个 Helion W4A4（NVFP4 权重 × NVFP4 激活）解码 GEMV 内核，其目标形状已有随代码一同提交的 Helion AOT 配置。问题是：在 Helion 自身的自动调优层之下，也就是 ptxas 编译阶段，CompileIQ 能否进一步挖掘性能？This section highlights many possible ways in which one can run a CompileIQ optimization campaign on Helion kernels. As a sample case study we applied CompileIQ to the pretuned nvfp4_gemv_fp4in_kernel, a Helion W4A4 (NVFP4 weight x NVFP4 activation) decode GEMV kernel that already ships a checked-in Helion AOT config for its target shape. The question: is there performance CompileIQ can extract below Helion’s own autotuning layer ,i.e., at the ptxas compilation stage?

环境与方法论Environment and Methodology

基准测试在一台配备 sm_100（Blackwell，计算能力 10.0）的 GB200 系统上进行，使用了 ptxas 13.3.73（CUDA 13.3）、CompileIQ 1.0.3、Triton 3.8.0，以及 Helion 的可编辑源代码安装版本。所有计时评估均通过 triton.testing.do_bench 执行，并锁定 GPU 时钟。每项观测到的性能提升都经过独立复测，并与 PyTorch 参考实现对比以验证数值正确性；同时通过 Welch’s t 检验（p < 0.05）和 Cohen’s d（d > 0.2）进行统计筛选。下方每项结果均附有完整的统计指标。Benchmarking was conducted on a GB200 system with sm_100 (Blackwell, compute capability 10.0), utilizing ptxas 13.3.73 (CUDA 13.3), CompileIQ 1.0.3, Triton 3.8.0, and an editable source installation of Helion. All timing evaluations were executed via triton.testing.do_bench with the GPU clock locked. Every observed performance improvement was re-measured independently, validated against the PyTorch reference implementation for numerical correctness, and statistically filtered using a Welch’s t-test (p < 0.05) alongside Cohen’s d (d > 0.2). Complete statistical metrics are provided alongside each result below.

内核描述Kernel Description

helion/pretuned_kernels/nvfp4_gemv/nvfp4_gemv.py 中的 nvfp4_gemv_fp4in_kernel，用于计算解码阶段（一个输出行）的 GEMV：out = (weight @ x) * alpha。每个 K 组都会通过 hl.load_float4_e2m1fn_x16_to_float16 解包打包的 NVFP4 权重和激活值字节，在一个 16 路静态范围内以 fp16 累加，并通过解码后的 E4M3 块缩放因子进行缩放。nvfp4_gemv_fp4in_kernel in helion/pretuned_kernels/nvfp4_gemv/nvfp4_gemv.py, computing out = (weight @ x) * alpha for a decode (one output row) GEMV. Both the packed NVFP4 weight and activation bytes are unpacked per K-group via hl.load_float4_e2m1fn_x16_to_float16, accumulated in fp16 over a 16-way static range, and scaled by a decoded E4M3 block scaler.

This kernel includes a checked-in AOT configuration specifically for the N=8192, K=28672 “down_proj” decode-GEMV shape in _helion_aot_nvfp4_gemv_cuda_sm100.py:

_CONFIG_FP4IN_DEFAULT = {
    "block_sizes": [4, 256],
    "num_warps": 2,
    "num_stages": 3,
    "range_multi_buffers": [None, True],
}

将 Runtime 用作 CompileIQ 的目标函数Using Runtime as CompileIQ Objective Function

将运行时间用作 CompileIQ 的目标函数，使我们能够直接优化执行速度。目标函数接收一个候选 ACF 配置，将其应用于固定的 Helion 内核并进行评估：内核运行一次，其输出会与 PyTorch 参考实现进行比对；只有通过正确性检查的配置才会使用 triton.testing.do_bench 进行基准测试。任何未通过正确性检查的配置都会返回无效分数，并在产生任何计时结果之前被移出搜索。Using runtime as the CompileIQ objective function lets us optimize execution speed directly. The objective function takes a candidate ACF configuration, applies it to our fixed Helion kernel, and evaluates it: the kernel runs once, its output is checked against the PyTorch reference implementation, and only configurations that pass correctness are then benchmarked with triton.testing.do_bench. Any configuration that fails correctness returns an invalid score and is dropped from the search before it ever contributes a timing measurement.

def objective(config) -> float:
    acf_path = ... # write `config` to a .acf, or None for baseline
    kernel = _make_kernel(acf_path)
    got = run()
    _check(got, ref, "fp4in", N, K) # correctness gate
    return do_bench(run, warmup=25, rep=100, return_mode="mean")

我们初始化一个 Search 实例，将评分例程与针对 PTXAS 编译器控制项的进化式自动调优器耦合起来，然后调用 .start()。它会在一个大小为 pool_size 的种群中，经过 generations 代迭代评估候选 ACF，并将表现最佳的控制配置繁衍至后续每一代。We initialize a Search instance that couples our scoring routine to an evolutionary autotuner over PTXAS compiler controls, then invoke .start(). It sweeps candidate ACFs across generations generations of a pool_size-sized population, breeding the top-performing control configurations forward into each subsequent iteration.

config = SearchConfiguration(problem_type="min", generations=8, pool_size=15)
tuner = Search(
    objective_function=objective,
    search_space=PtxasSearchSpace(version=cuda_version, variant="default"),
    search_config=config,
    dump_results="nvfp4_gemv_results.csv",
    worker_type=WorkerTypes.ISOLATED, # crashing ACF candidates shouldn't kill the search
)

with gpu_benchmark_mode(clock_mhz=target_clock_mhz, raise_on_failure=False):
    results = tuner.start(task_timeout=60)

使用 CompileIQ 在 PTXAS ACF 范围内执行进化搜索，针对我们在 GB200 上固定的 Helion 配置，经过 8 代、每代候选池大小为 15 的搜索后，得到了优化后的 ACF 控制集。该候选方案将执行时间从基准的 0.0971 ms 降至 0.0928 ms，根据搜索评分循环中的直接测量，实现了 1.047 倍（4.7%）的加速。Executing an evolutionary search with CompileIQ across PTXAS ACFs, spanning 8 generations with a pool size of 15 candidates over our fixed Helion configuration on GB200 yielded an optimized ACF control set. This candidate reduced execution time from a baseline of 0.0971 ms down to 0.0928 ms, delivering a 1.047x (4.7%) speedup as measured directly within the search’s scoring loop.

sample baseline run: 0.09713 ms
Best runtime (CompileIQ Search): 0.0928 ms
Speedup: 1.047x

CompileIQ 搜索结果验证Validation of CompileIQ Search Results

为确认我们的研究结果，我们在固定 GPU 时钟频率下，通过增加预热和重复迭代次数，独立重新测量每项加速效果。每种配置都要严格对照参考模型进行正确性验证，同时通过 Welch t 检验和 Cohen’s d 筛选。这构成了我们严格的统计阈值（要求 p < 0.05、cohens_d > 0.2 且 optimized.mean < baseline.mean），用于在采信数据之前验证性能提升。To confirm our findings, every speedup is re-measured independently using increased warmup and repetition iterations under fixed GPU clock frequencies. Each configuration undergoes strict correctness verification against the reference model alongside a Welch’s t-test and Cohen’s d filter. This serves as our rigorous statistical threshold (requiring p < 0.05, cohens_d > 0.2, and optimized.mean < baseline.mean) to verify performance gains before trusting the numbers.

# run 1
baseline:  n=1744 mean=0.09752ms std=0.00218
optimized: n=2018 mean=0.09301ms std=0.00133
speedup_mean:   1.0485
p_value:  0.0
cohens_d: 2.501
significant: True

# run 2 (independent re-run)
baseline:  n=1757 mean=0.09744ms std=0.00216
optimized: n=2028 mean=0.09299ms std=0.00130
speedup_mean:   1.0478
p_value:  0.0
cohens_d: 2.489
significant: True

两次运行在实验方差范围内几乎完全相同：d ≈ 2.5 表明效应量巨大，远超 0.2 的统计阈值，从而验证了 ACF 性能提升确实存在。The two executions remain virtually identical within experimental variance: with d ≈ 2.5 demonstrating a massive effect size that easily surpasses our statistical threshold of 0.2 , verifying that this ACF performance gain is genuine.

两次运行在实验方差范围内几乎完全一致，而且两者的 d 值都约为 2.5，属于极大的效应量，接近我们 d > 0.2 门槛的 12.5 倍，证实这项 ACF 性能提升确实存在，并非噪声。The two runs are virtually identical within experimental variance, and d ≈ 2.5 in both is a massive effect size nearly 12.5x past our d > 0.2 gate, confirming this ACF performance gain is genuine, not noise.

针对硬件指标而非运行时间进行调优Tuning against a hardware metric instead of runtime

虽然执行速度是常见的优化目标，但 CompileIQ 并不局限于运行时性能。通过直接集成 Nsight Compute (ncu)，目标函数可以针对该工具所提供的任何硬件指标。为了确定限制内核性能的关键架构瓶颈，我们首先评估其性能剖析结果；在 GB200 系统上分析 nvfp4_gemv_fp4in_kernel 后，发现以下瓶颈：While execution speed is a common optimization goal, CompileIQ is not limited to runtime alone. Integrating directly with Nsight Compute (ncu) enables the objective function to target any hardware metric exposed by the tool. To determine the key architectural bottleneck limiting kernel performance, we start by evaluating its profile and analyzing nvfp4_gemv_fp4in_kernel on a GB200 system highlights the following bottlenecks:

Section: GPU Speed of Light Throughput
Memory Throughput               %       78.99
L1/TEX Cache Throughput         %       89.35
Compute (SM) Throughput         %       37.29

Section: Occupancy
Theoretical Occupancy                   %    25.0
Achieved Occupancy                      %    21.63

  Est. Speedup: 21.01%
  ... theoretical occupancy (25.0%) is limited by the number of required
  registers.

性能分析显示，该内核受寄存器限制。NCU 自身的成本模型认为，占用率带来了 21% 的“估计加速”空间；而占用率上限由寄存器数量决定，这正是 PTXAS 的寄存器分配控制项可以直接影响的因素。Profiling shows the kernel is register-limited. NCU’s own cost model attributes a 21% ‘Est. Speedup’ opportunity to occupancy, and this ceiling is set by register count, something PTXAS’s register-allocation controls can directly influence.

将目标从运行时间切换为占用率，只需改动两处：目标函数的评分对象，以及搜索的 problem_type：Switching the objective from runtime to occupancy only touches two things: what the objective function scores, and the search’s problem_type:

# objective(): do_bench(...) becomes an ncu metric read
return _measure_occupancy(acf_path) # ncu: sm__warps_active.avg.pct_of_peak_sustained_active

# search config: minimize runtime -> maximize occupancy
config = SearchConfiguration(problem_type="max", generations=4, pool_size=6)

在 4 代 × 池大小为 6 的范围内，直接以实际占用率本身作为搜索目标，而非运行时，找到了一种候选方案，使代理指标提高了 +0.16 个百分点（实际占用率从 21.60% 提升至 21.76%，p = 9.2e-23，d = 13.17）；经独立重新测量，这一改进对应着幅度不大但统计上相当显著的运行时提升，约为 1.75%–2.0%（p = 3.6e-116，d = 0.78）。Pointing the search at achieved occupancy itself across 4 generations × pool size 6, rather than runtime, found a candidate that moved the proxy metric by +0.16pp (21.60% → 21.76% achieved occupancy, p = 9.2e-23, d = 13.17) and, independently re-measured, translated into a modest but statistically solid runtime win of ~1.75-2.0% (p = 3.6e-116, d = 0.78).

多目标调优Multi-Objective Tuning

CompileIQ 还可以执行多目标调优：评分例程会通过返回元组而非标量浮点数，同时从多个维度评估候选项。自动调优过程不再只筛选出单个表现最佳的候选项，而是沿帕累托前沿绘制候选项分布，找出在所有目标指标上均未被支配的控制集。由此得到的统一 ACF 可针对执行运行时间和 L1TEX 缓存吞吐量进行联合优化。CompileIQ can also perform multi-objective tuning, where the scoring routine evaluates candidates along multiple dimensions simultaneously by returning a tuple rather than a scalar float. Instead of isolating a single top performer, the autotuning process maps out candidates along the Pareto frontier, identifying control sets that are not dominated across all target metrics. This yields a unified ACF optimized jointly for execution runtime alongside L1TEX cache throughput.

def objective(config) -> tuple[float, float]:
    runtime_ms = do_bench(run, warmup=25, rep=100, return_mode="mean")
    l1_throughput = _measure_l1throughput(acf_path)
    return runtime_ms, -l1_throughput

8 代 × 15 个候选的种群规模，最终在帕累托前沿上收敛到一个占主导地位的候选方案。与权衡曲线不同，这个单点按照搜索本身的评分，在两个指标上同时胜过其他所有已评估的候选方案。使用全新缓存，并采用正式发布级别的试验次数进行验证：8 generations × pool size 15 converged to a single dominant candidate on the Pareto front. Unlike a tradeoff curve, this single point beats every other evaluated candidate on both axes simultaneously in the search’s own scoring. Validating it on a fresh cache at shipping-grade trial counts:

=== RUNTIME (ms, lower is better) ===
p_value=3.516e-295  cohens_d=87.024
speedup_mean=1.0689x
=== L1TEX THROUGHPUT (%, higher is better; internally negated) ===
p_value=7.123e-19  cohens_d=3.686
baseline L1TEX THROUGHPUT mean=81.5087%
candidate L1TEX THROUGHPUT mean=85.9620%

多目标调优找到了一种 ACF，实现了经验证的 6.9% 运行速度提升，同时 L1TEX 吞吐量提高了 4.5 个百分点（81.51% → 85.96%）。两项结果均具有统计显著性（p ≪ 0.05，Cohen’s d ≫ 0.2），且两项指标之间不存在此消彼长的取舍。Multi-objective tuning found an ACF delivering a validated 6.9% runtime speedup, accompanied by a simultaneous 4.5pp L1TEX throughput gain (81.51% → 85.96%). Both results are statistically significant (p ≪ 0.05, Cohen’s d ≫ 0.2), with neither metric being traded off against the other.

跨多种输入形状联合优化 ACFJoint ACF Optimization Across Multiple Input Shapes

虽然此前的评估聚焦于单一内核形状 N=8192、K=28672，但实际模型部署需要覆盖多种配置。nvfp4_gemv 实现附带了多种预先调优的 Helion AOT 配置，其中包括专为 N=15360、K=5120 设计的配置。无需为每种形状分别维护控制文件，CompileIQ 便可通过统一搜索找到一个联合 ACF，使两种形状的整体性能达到最大化。While earlier evaluations focused on a single kernel shape, N=8192, K=28672, real-world model deployment demands coverage across multiple configurations. The nvfp4_gemv implementation ships with several pre-tuned Helion AOT configurations, including a dedicated setup for N=15360, K=5120. Instead of maintaining distinct control files per shape, CompileIQ enables a unified search to discover a single, joint ACF that maximizes performance across both shapes simultaneously.

def objective(config) -> tuple[float, ...]:
    """One shared ACF, one runtime score per shape (both minimize)."""
    acf_path = ... # write `config` to a temp .acf, or None for baseline
    return tuple(_bench_shape(acf_path, shape) for shape in SHAPES)

对这两种形状分别执行种群规模为 15 的 8 代搜索后，得到了一个包含 4 个最优候选项的帕累托前沿：Executing an 8-generation search with a pool size of 15 across both shapes yielded a Pareto front containing 4 optimal candidates:

与 down_proj 的默认基准耗时 0.0972 ms 和 15360_5120 的 0.0315 ms 相比，pareto_0 脱颖而出，成为在两个目标张量维度上都始终快于基线的配置。Compared to the default benchmark times of 0.0972 ms for down_proj and 0.0315 ms for 15360_5120, pareto_0 emerges as the standout configuration that consistently outpaces the baseline across both target tensor dimensions.

为确认这些性能提升，我们对每种形状上的 pareto_0 分别进行了独立验证。两种配置均轻松通过统计验证（down_proj 的加速比为 1.05x，p 值为 1.1e-224，Cohen’s d 为 67.71；在 15360_5120 形状上的加速比为 1.01x，p 值为 6.2e-73，d 值为 4.54），远超我们的基准测试要求，证明这些提升是真实的性能改进，而非系统噪声。单个联合 ACF 在两个目标维度上同时成功超越了默认基线。To confirm these performance gains, we execute an independent validation of pareto_0 on each individual shape. Both configurations comfortably pass our statistical validation (1.05x speedup with a p-value of 1.1e-224 and Cohen’s d of 67.71 for down_proj, and 1.01x speedup with a p-value of 6.2e-73 and a d of 4.54 on the 15360_5120 shape), far exceeding our benchmark requirements and proving these gains represent genuine performance improvements rather than system noise. A single joint ACF successfully outperforms the default baseline across both target dimensions simultaneously.

此工作流在端到端模型调优中的适用性Applicability of this workflow for end-to-end model tuning

虽然这篇文章的重点是优化底层编译器，以从单个内核中榨取更多性能，但我们也应当跳出局部视角，重点说明如何利用这样的工作流，在 PyTorch 上对模型进行端到端调优。While the focus of this entire post is about optimizing the underlying compiler for extracting more performance from individual kernels, it is important to zoom out and highlight how a workflow like this could be utilized for end-to-end tuning of models on PyTorch.

第一步是对你希望优化的模型中最重要的内核进行性能分析和特征表征。随后，可以将每个内核作为 CompileIQ 优化的目标，并为其关联一个高级控制文件，以从该内核实例中提取最高性能。这些配置会在运行时通过 Helion 固定的自动调优配置应用，具体使用 Triton 的 ptxas_options API，还是直接通过 PTXAS 命令行应用，则取决于各内核实现的入口点。The first step is to profile and characterize the most important kernels that are part of the model which you wish to optimize. Then, each of these kernels can then be a target of CompileIQ optimization, with an associated Advanced Control File to extract maximum performance from that instantiation of the kernel. These get applied during runtime using Helion’s pinned autotune configurations, using Triton’s ptxas_options API or directly via the PTXAS command line depending on the entry point of each kernel’s implementation.

我们希望坦诚说明预期收益：对于已经高度优化的内核，编译器控制调优带来的运行时提升约为 2%–7%；对于顶尖模型中的关键内核而言，这仍可能意义重大。优化程度较低的内核则实现了两位数的性能提升，因此更有望显著改善端到端表现。We want to be upfront about the expected gains: On kernels that are already highly optimized, compiler-control tuning has shown runtime improvements on the order of 2-7% which may still mean a lot for critical kernels in top models. Less optimized kernels have achieved double-digit performance gains, offering greater potential for meaningful end-to-end improvements.

需要注意的要点是：它可以自动找到，几乎不需要编译器方面的专业知识；它会根据你自己的正确性检查进行验证；最终生成的文件可以提交到内核文件旁边，并在每次构建时重新应用。这正是理论上存在的几个百分点提升，与实际交付的几个百分点提升之间的区别。The important points to note are that it is found automatically, and requires little to no compiler expertise, it is verified against your own correctness check, and the result is a file you can commit next to the kernel and reapply on every build. It is the difference between a few percent that exists in principle and a few percent you actually ship.

它带来的收益也恰好会在最需要的地方不断累积。大型集群中有一些高价值内核，会在推理循环中持续运行。在这种规模下，即使是在软件栈底层、针对主导解码步骤的操作提升几个百分点，也很有价值。It also compounds exactly where you’d want it to. There are high-value kernels inside inference loops that run continuously across large fleets. Even a few percent at the bottom of the stack, applied to the operations that dominate a decode step, is valuable at that scale.

注意事项Caveats

• ACF 与硬件和工具包版本相关：针对 sm_90 调优的 ACF 在 sm_100 上无效。应用不匹配的 ACF 可能导致结果错误或运行时错误，而不仅仅是代码运行变慢。加速包清单会注明其目标架构和工具包版本。这也正是搜索循环中的正确性检查不可省略的原因。
  ACFs are hardware- and toolkit-specific: An ACF tuned for sm_90 is not valid on sm_100. Applying a mismatched ACF can produce incorrect results or runtime errors, not just slower code. The booster pack manifest states its target architecture and toolkit version. This is also why the correctness check in the search loop is not optional.
• 搜索阶段的开销不容小觑：autotune_search_acf 中的每个 ACF 都会增加自动调优的试验次数。离线时将其用于重要的内核，然后固定胜出的配置。
  Search-time is a nontrivial cost: Every ACF in autotune_search_acf multiplies your autotuning trials. Use it offline on kernels that matter, then pin the winner.
• 性能提升因工作负载而异：结果取决于工作负载、形状、架构，以及内核还剩多少优化空间。对你的内核来说，唯一重要的数字就是你在自己的内核上测得的数字。
  Gains are workload-specific: The result depends on the workload, the shapes, the architecture, and how much headroom the kernel has left. The only number that matters for your kernel is the one you measure on your kernel.

现在轮到你来试试了！Your turn to try this out!

如果你从未编写过 GPU 内核：请从 pretuned_kernels/ 开始。挑选一个与你的需求相近的内核，阅读并运行它即可。这些内核都很短，而且无需等待自动调优，运行后就能立即达到调优后的速度。If you’ve never written a GPU kernel: start at pretuned_kernels/. Pick something close to what you need, read it and then run it. The kernels are short, and they run at tuned speed immediately, with no autotuning wait.

如果你已经在使用 Helion：下载 helion-booster-pack，将 autotune_search_acf=[…] 添加到你关注的一个内核中，然后对其进行自动调优。不使用 ACF 的基线方案仍会保留在搜索范围内，因此最坏的代价也就是你花费的时间。有关完整示例，请参阅 Helion 的 examples/acfs/。If you already use Helion: download the helion-booster-pack, add autotune_search_acf=[…] to one kernel you care about, and autotune it. The no-ACF baseline stays in the search, so the downside is bounded at the time you spend. See Helion’s examples/acfs/ for a worked example.

如果你以编写内核为生，这里的要点并不是工具本身，而是编译器的控制面是一个值得探索的优化空间，可以对其进行搜索，而大多数人还没有开始这么做。任何暴露可调控选项的编译器，在任何后端上，都是如此。希望我们激发了你探索这一领域的兴趣！If you work on kernels for a living: the takeaway here is not the tool itself. It is that the compiler’s control surface is a legitimate optimization space, it is searchable, and most people are not yet searching it. That is true of any compiler that exposes tunable controls, on any backend. We’re hopeful that we’ve inspired you to explore this space!

许可协议Licensing

Helion 是开源项目，由 PyTorch 基金会托管。CompileIQ 可免费使用，并根据 NVIDIA 软件许可协议（这是一项专有的最终用户许可协议）进行分发。它生成的 ACF 文件可自由分发，因此可以发布加速包，也可以将 ACF 与内核一同提交到开放仓库中。Helion is open source and hosted by the PyTorch Foundation. CompileIQ is free to use and it is distributed under the NVIDIA Software License Agreement, a proprietary EULA. The ACF files it produces are freely distributable, which is why booster packs can be published and why an ACF can be committed alongside a kernel in an open repository.

链接Links

• Helion
  • https://github.com/pytorch/helion
  • https://helionlang.com
  • https://github.com/pytorch/helion/tree/main/examples/acfs
  • https://github.com/pytorch/helion/tree/main/pretuned_kernels
• CompileIQ
  • https://github.com/NVIDIA/CompileIQ
  • https://nvidia.github.io/CompileIQ/stable/
  • https://github.com/NVIDIA/CompileIQ/releases?q=booster-packs&expanded=true
  • https://developer.nvidia.com/blog/extract-more-kernel-performance-with-nvidia-compileiq-auto-tuning