Helion：用自动调优改进 vLLM 推理

简而言之TL;DR

我们将 Helion 集成到 vLLM 的线性后端，以探索经过自动调优的高级内核 DSL 如何在降低内核实现复杂度的同时提升 LLM 推理性能。单个 Helion 通用矩阵乘法（GEMM）实现即可覆盖多种算法变体，包括 Standard GEMM、Split-K 和 Swap-AB，并通过针对每种形状的自动调优自动选出最佳变体和配置。We integrated Helion into vLLM’s linear backend to explore how an autotuned, high-level kernel DSL can improve LLM inference performance while reducing kernel implementation complexity. A single Helion general matrix multiplication (GEMM) implementation can cover multiple algorithmic variants, including Standard GEMM, Split-K, and Swap-AB, with the best variant and config automatically selected through per-shape autotuning.

在 NVIDIA Hopper GPU 上，Helion 线性后端结合了按形状调优与混合调度，在所评估的各个模型中均优于 vLLM 默认的 CUTLASS 和 DeepGEMM 后端，带来稳定的端到端性能提升，并使部分工作负载的吞吐量提升超过 10%。On NVIDIA Hopper GPUs, the Helion linear backend combines per-shape tuning with hybrid dispatch to outperform the vLLM default CUTLASS and DeepGEMM backends across the evaluated models, delivering consistent end-to-end performance gains and more than 10% throughput improvement for some workloads.

vLLM 和 Helion 简介Brief Background on vLLM and Helion

vLLM 是一个用于大型语言模型（LLM）的高性能推理与服务框架。对于量化线性层（如 FP8、INT8、INT4 和 NVFP4），vLLM 提供了专用的线性后端，可集成 CUTLASS、DeepGEMM 和 FlashInfer 等库中针对硬件优化的内核实现。这些后端使 vLLM 能够针对不同的量化格式和硬件平台使用优化内核。vLLM is a high-performance inference and serving framework for large language models (LLMs). For quantized linear layers, such as FP8, INT8, INT4, and NVFP4, vLLM provides specialized linear backends that integrate hardware-optimized kernel implementations from libraries such as CUTLASS, DeepGEMM, and FlashInfer. These backends allow vLLM to use optimized kernels for different quantization formats and hardware platforms.

Helion 是一种原生支持 PyTorch、与硬件无关的内核 DSL，旨在使用 tile 编程模型编写高性能内核。它的高级抽象让开发者能够用简洁的 Python 风格代码一次性表达内核，同时由 Helion 针对不同工作负载和硬件平台生成专门优化的代码。Helion 不要求为每种工作负载和硬件目标手动特化内核，而是依靠提前（AOT）自动调优器系统地探索广阔的搜索空间，范围涵盖底层内存布局和内核调度，以及高层算法选择，并选出性能最佳的配置。高级抽象与系统化自动调优相结合，使单一内核实现能够针对多种工作负载和硬件进行优化。Helion 还支持由 LLM 引导的搜索，以进一步提升调优效率。Helion is a PyTorch-native hardware agnostic kernel DSL designed for writing high-performance kernels using a tile-programming model. Its high-level abstraction allows developers to express a kernel once in concise Pythonic code, while Helion generates specialized code for different workloads on different hardware platforms. Rather than manually specializing kernels for each workload and hardware target, Helion relies on ahead-of-time (AOT) autotuner to systematically explore a broad search space, from low-level memory layout and kernel scheduling to high-level algorithmic choices, and select the best-performing config. This combination of high-level abstraction and systematic autotuning enables a single kernel implementation to be optimized across diverse workloads and hardware. Helion also supports LLM-guided search to further improve tuning efficiency.

Helion 的采用机遇与挑战Helion Adoption Opportunities and Challenges

本节总结了在 vLLM 等推理引擎中采用 Helion 内核的主要机遇与挑战，并从宏观层面概述了 Helion 的价值主张及其权衡。This section summarizes the key opportunities and challenges of adopting Helion kernels in inference engines such as vLLM, providing a high-level overview of Helion’s value proposition and tradeoffs.

机遇Opportunities

性能：我们此前的工作表明，通过精细调优，Helion 有望在各种工作负载模式和硬件平台上实现推理内核的最先进性能。Performance: Our previous work demonstrated Helion’s potential to achieve SOTA performance for inference kernels across diverse workload patterns and hardware platforms through fine-grained tuning. 

系统化调优框架：与那些通过迭代生成、分析性能并改进内核实现的开放式智能体方法相比，Helion 将内核调优建模为一个在定义明确的搜索空间中进行的结构化数值优化问题。这使调优过程更加稳健可靠，同时保持了具有竞争力的调优效率，且无需依靠性能分析来引导搜索。Helion 还支持由 LLM 引导的搜索，将 LLM 的推理能力与系统化数值优化相结合，从而高效地探索搜索空间。Systematic Tuning Framework: Compared with more open-ended agentic approaches that iteratively generate, profile, and refine kernel implementations, Helion formulates kernel tuning as a structured numerical optimization problem over a well-defined search space. This makes the tuning process more robust and reliable while remaining competitive in tuning efficiency, without relying on profiling to guide the search. Helion further supports LLM-guided search, combining the reasoning capabilities of LLMs with systematic numerical optimization to efficiently navigate the search space.

可移植性与抽象：Helion 不仅是一种可移植的领域特定语言（DSL），还支持在维护单一内核实现的同时，针对不同的工作负载模式和硬件平台进行优化。Portability and Abstraction: Helion is not only a portable DSL. It is possible to maintain a single kernel implementation while optimizing it for different workload patterns and hardware platforms. 

客户端内核优化：vLLM 等推理引擎中的默认内核通常针对通用工作负载和常用模型进行优化。借助 Helion，用户可以针对特定部署进一步调优内核性能，无需具备专业的内核知识，也无需投入大量工程工作。Client-Side Kernel Optimization: Default kernels in inference engines such as vLLM are typically optimized for general workloads and commonly used models. With Helion, users can further tune kernel performance for their specific deployments without requiring specialized kernel expertise or significant engineering effort.

挑战Challenges

Helion 在集成到推理引擎方面仍面临一些挑战，Helion 团队正在持续开展工作，以解决和缓解这些问题。Helion still has some challenges integrating into inference engines, with ongoing work from the Helion team to address and mitigate them.

预先进行的内核调优开销：Helion 会自动化并系统化内核调优，但精细调优仍可能耗费数小时。这种开销主要源于调优策略的粒度，而非 Helion 本身。使用其他内核领域特定语言实现同等程度的工作负载专属优化，也会产生相当甚至更高的调优成本。Ahead-of-time kernel tuning overhead: Helion automates and systematizes kernel tuning, but fine-grained tuning can still take hours. This overhead comes primarily from the granularity of the tuning strategy rather than from Helion itself. Achieving the same level of workload-specific specialization with other kernel DSLs would incur comparable—or potentially greater—tuning costs.

启动时开销：vLLM 启动期间的 CUDA Graph 捕获会触发 Helion JIT 编译，从而增加冷启动延迟。通过缓存已编译的产物，可在热启动时大幅消除此项开销。Startup-time overhead: CUDA Graph capture during vLLM startup triggers Helion JIT compilation, increasing cold-start latency. This overhead can be largely eliminated on warm starts by caching compiled artifacts.

推理运行时开销：在 CUDA Graph 捕获范围之外，推理期间 Helion 内核的分派和启动可能会带来额外的 CPU 开销，从而抵消细粒度调优带来的性能提升。实际应用中，在 CUDA Graph 下执行 Helion 效果最佳，可避免这种 CPU 开销。Inference-runtime overhead: Outside the CUDA Graph capture range, Helion kernel dispatch and launch during inference can introduce additional CPU overhead, potentially offsetting the performance gains from fine-grained tuning. In practice, Helion is most effective when executed under CUDA Graphs to avoid this CPU overhead.

维护开销：为热门模型提供预调优配置，会带来持续的上游维护负担。大型配置文件难以维护，也无法通过单元测试和 CI 进行全面验证。Maintenance overhead: Shipping pre-tuned configs for popular models creates an ongoing upstream maintenance burden. Large config files are difficult to maintain and impractical to validate exhaustively through unit tests and CI. 

权衡三角形The Tradeoff Triangle

使用 Helion 进行精细化内核调优，需要在性能、易用性和可维护性之间权衡。这种权衡并非 Helion 独有，而是特定工作负载内核优化固有的特点：优化其中两项，往往会以牺牲第三项为代价。Fine-grained kernel tuning with Helion presents a tradeoff among performance, usability, and maintainability. This tradeoff is not specific to Helion, but is inherent to workload-specific kernel optimization: optimizing for two often comes at the expense of the third.

图 1：性能、易用性与可维护性之间的权衡三角形Fig. 1: The Performance-Usability-Maintainability tradeoff triangle

更高的性能通常需要更精细地调整配置。这要么会增加客户端自动调优的开销，要么要求维护者在上游提供并维护更多预调优配置。因此，目标是针对目标用例找到恰当的平衡。Higher performance generally requires more fine-grained config tuning. This either increases client-side autotuning overhead or requires maintainers to provide and maintain more pre-tuned configs upstream. The goal, therefore, is to strike the right balance for the target use case. 

Helion 线性后端Helion Linear Backend 

范围Scope

这项工作为 vLLM 添加了 Helion 线性后端，并重点支持通过 Helion 的 Triton 后端在 NVIDIA Hopper GPU 上运行。FP8 和 INT8 是 Hopper 上实现高效推理的主要量化格式，因此我们针对以下格式的量化 GEMM：This work adds a Helion linear backend to vLLM and focuses on NVIDIA Hopper GPUs using Helion’s Triton backend. FP8 and INT8 are the primary quantization formats for efficient inference on Hopper, so we target quantized GEMM with the following formats:

• FP8_Dynamic：激活值按 token 缩放，权重按通道缩放。
  FP8_Dynamic: FP8 per-token activation and per-channel weight scaling.
• W8A8_INT8：激活值按 token 缩放、权重按通道缩放，均采用 INT8。
  W8A8_INT8: INT8 per-token activation and per-channel weight scaling.
• Block_FP8：采用 1×128 激活缩放和 128×128 权重缩放的 FP8。
  Block_FP8: FP8 with 1×128 activation scaling and 128×128 weight scaling. 

Helion 的 CuteDSL 后端初步结果表明，其在 NVIDIA Blackwell GPU 上的 GEMM 性能具有竞争力。随着 CuteDSL 后端逐渐成熟，这项工作可以进一步扩展，以支持 Blackwell。Initial results with Helion’s CuteDSL backend show competitive GEMM performance on NVIDIA Blackwell GPUs. As the CuteDSL backend matures, this work can be extended to support Blackwell.

内核实现Kernel Implementation

多种 GEMM 算法变体可以针对特定输入形状提升性能。具体而言：Several GEMM algorithm variants can improve performance for specific input shapes. In particular:

• Split-K：将 K 维度划分到多个线程块中，以便在 M 和/或 N 维度过小时提高并行度。
  Split-K: Partitions the K dimension across multiple thread blocks to increase parallelism when the M and/or N dimensions are too small.
• Swap-AB：将 A@B 重写为 (B.T@A.T).T，以提升 M 维度较小时的性能，同时支持更有利的分块策略并提高 GPU 利用率。
  Swap-AB: Rewrites A@B as (B.T@A.T).T to improve performance for shapes with a small M dimension by enabling more favorable tiling and better GPU utilization.

传统上，内核作者需要实现多个 GEMM 变体，将标准 GEMM 与这些专用变体进行基准测试，并制定启发式规则，根据输入形状决定调度到哪种实现。例如，vLLM 当前的 Block_FP8 线性后端在 Hopper 上遇到 M < 32 时会调度到 Swap-AB 变体。这类启发式规则通常旨在保证整体表现良好，但未必能针对每种模型或工作负载达到最优效果。Traditionally, kernel authors need to implement multiple GEMM variants, benchmark the standard GEMM against these specialized variants, and develop heuristics to determine which implementation to dispatch to based on the input shape. For example, vLLM’s current Block_FP8 linear backend dispatches to the Swap-AB variant when M < 32 on Hopper. Such heuristics are typically designed to perform well in general, but may not be optimal for every model or workload. 

使用 Helion，一个统一的 GEMM 实现即可涵盖三种变体：Standard、Split-K 和 Swap-AB。无需分别实现多个内核，也无需手动设计调度启发式规则；只需将算法选项设为可调参数。随后，AOT 自动调优器便可针对每种工作负载，自动选择最佳变体及底层内核配置。With Helion, a single unified GEMM implementation can cover all three variants—Standard, Split-K, and Swap-AB. Instead of implementing separate kernels and manually designing dispatch heuristics, the algorithmic choices are exposed as tunable parameters. The AOT autotuner can then automatically select the best variant along with the low-level kernel config for each workload.

为简化说明，下面我们使用一个基础矩阵乘法内核来阐释这种方法。本工作中使用的量化 GEMM 内核遵循相同的结构，只是额外加入了量化和缩放逻辑。For simplicity, we use a basic matrix multiplication kernel below to illustrate the approach. The quantized GEMM kernels used in this work follow the same structure, with additional quantization and scaling logic.

def matmul(
    out: torch.Tensor,  # [M, N]
    a: torch.Tensor,    # [M, K]
    b: torch.Tensor,    # [K, N]
) -> None:
    M, K = a.shape
    N = b.shape[1]
    hl.specialize(K)
    hl.specialize(N)

    out_dtype = out.dtype
    acc_dtype = torch.float32

    split_k = hl.register_tunable(
        "split_k", PowerOfTwoFragment(1, 256)
    )
    k_block_size = helion.next_power_of_2(helion.cdiv(K, split_k))
    if split_k > 1:
        out.zero_()

    swap_ab = hl.register_tunable(
        "swap_ab", BooleanFragment()
    )

    for tile_m, tile_n, outer_k in hl.tile(
        [M, N, K],
        block_size=[None, None, k_block_size]
    ):
        acc = hl.zeros([tile_m, tile_n], acc_dtype)
        acc_t = acc.t()
        for tile_k in hl.tile(outer_k.begin, outer_k.end):
            if swap_ab:
                a_blk = hl.load(a, [tile_m.index[None, :], tile_k.index[:, None]])
                b_blk = hl.load(b, [tile_k.index[None, :], tile_n.index[:, None]])
                acc_t = hl.dot(
                    b_blk,
                    a_blk,
                    acc=acc_t,
                    out_dtype=acc_dtype,
                )
            else:
                acc = hl.dot(
                    a[tile_m, tile_k],
                    b[tile_k, tile_n],
                    acc=acc,
                    out_dtype=acc_dtype,
                )

        if swap_ab:
            out_blk = acc_t.t().to(out_dtype)
        else:
            out_blk = acc.to(out_dtype)

        if split_k == 1:
            out[tile_m, tile_n] = out_blk
        else:
            hl.atomic_add(out, [tile_m, tile_n], out_blk)

此实现将两种算法选择都作为可调参数公开：split_k 被限制为不超过 256 的 2 的幂值，而 swap_ab 是一个布尔参数。自动调优器会探索受约束的搜索空间，对不同组合进行基准测试，并为每种工作负载选择最佳变体。This implementation exposes both algorithmic choices as tunable parameters: split_k is constrained to power-of-two values up to 256, while swap_ab is a Boolean parameter. The autotuner explores the constrained search space, benchmarks different combinations, and selects the best variant for each workload.

混合调度Hybrid Dispatch

图 2：Helion 线性后端混合调度策略，基于运行时 num_tokens 和 CUDA Graph 覆盖率。Fig. 2: Helion linear backend hybrid dispatch strategy based on runtime num_tokens and CUDA Graph coverage.

我们采用一种基于运行时 num_tokens 和 CUDA Graph 覆盖范围的混合分发策略。对于不超过 max_helion_size 的小尺寸，线性后端会在 CUDA Graph 重放期间分发到 Helion。对于超过 max_helion_size 的更大尺寸，则回退到默认内核（CUTLASS/DeepGEMM）。We adopt a hybrid dispatch strategy based on runtime num_tokens and CUDA Graph coverage. For small shapes up to max_helion_size, the linear backend dispatches to Helion under CUDA Graph replay. For larger shapes beyond max_helion_size, it falls back to the default kernel (CUTLASS/DeepGEMM).

这种混合调度策略具有三项实际优势：This hybrid dispatch strategy provides three practical benefits:

• 消除 Helion 运行时开销。Helion 内核仅通过 CUDA Graph 重放执行，从而避免内核分派和启动带来的额外 CPU 开销。
  Eliminates Helion runtime overhead. Helion kernels execute only through CUDA Graph replay, avoiding the additional CPU overhead from kernel dispatch and launch.
• 降低内核调优开销。细粒度 Helion 调优仅针对在解码工作负载中占主导地位的小 num_tokens 范围，从而显著减少需要调优的输入形状数量，同时保留实现显著端到端性能提升的机会。
  Reduces kernel tuning overhead. Fine-grained Helion tuning is limited to the small num_tokens range that dominates decoding workloads, significantly reducing the number of input shapes that need to be tuned while preserving the opportunity for meaningful end-to-end performance gains.
• 降低配置维护开销。数量更少的调优配置让预调优配置更便于验证、发布和长期维护。
  Reduces config maintenance overhead. The smaller set of tuned configs makes pre-tuned configs more practical to validate, ship, and maintain over time.

内核自动调优Kernel Autotuning

我们使用 vLLM 中提供的实用脚本对 Helion 线性内核进行自动调优：We autotune the Helion linear kernels using the utility script available in vLLM:

HELION_AUTOTUNER=LLMSeededLFBOTreeSearch \
HELION_BENCHMARK_CUDAGRAPH =1 \
python scripts/autotune_helion_kernels.py \
  --kernels scaled_mm block_scaled_mm \
  --autotune-effort "full"

以下各节介绍了本研究中采用的关键调优策略和配置。The following sections describe the key tuning strategies and setups used in this work.

单个形状的配置调优Per-shape Config Tuning

为了与 CUTLASS、DeepGEMM 和 FlashInfer 等高度优化的 GEMM 库竞争，我们会针对 Helion 调度范围内的每种输入形状，单独调优 Helion 内核。To compete with highly optimized GEMM libraries such as CUTLASS, DeepGEMM, and FlashInfer, we tune the Helion kernel individually for each input shape within the Helion dispatch range.

如混合调度策略中所述，本项工作将 max_helion_size 设置为 32。vLLM 会为以下 num_tokens 值捕获 CUDA Graph：As described in the hybrid dispatch strategy, we set max_helion_size = 32 for this work. vLLM captures CUDA Graphs for the following num_tokens values:

[1, 2, 4] + range(8, 256, 8) + range(256, max_graph_size + 1, 16)

因此，当 max_helion_size = 32 时，Helion 内核会针对以下内容进行自动调优：With max_helion_size = 32, the Helion kernels are therefore autotuned for:

num_tokens = [1, 2, 4, 8, 16, 24, 32]

每个 num_tokens 值都针对模型使用的相应 GEMM 形状进行了单独调优。Each num_tokens value is tuned individually for the corresponding GEMM shapes used by the model.

为自动调优器基准测试启用 CUDA GraphEnable CUDA Graph for Autotuner Benchmarking

由于内核分派和启动会带来额外的 CPU 开销，Helion 内核仅在 CUDA Graph 执行模式下使用。因此，启用 CUDA Graph 进行基准测试，可以让自动调优器在更接近实际推理执行的条件下评估配置。Due to the additional CPU overhead from kernel dispatch and launch, Helion kernels are used only under CUDA Graph execution. Benchmarking with CUDA Graph enabled therefore allows the autotuner to evaluate configs under conditions that more closely match actual inference execution.

Helion 通过 HELION_BENCHMARK_CUDAGRAPH 环境变量提供此功能，并在本工作中启用。Helion exposes this feature through the HELION_BENCHMARK_CUDAGRAPHenvironment variable and is turned on for this work.

使用 LLM 引导搜索Use LLM-Guided Search

本工作中的内核配置由 LLMSeededLFBOTreeSearch 自动调优器生成。该调优器使用 LLM 找出一组有潜力的配置候选项，作为后续 LFBOTreeSearch 的种子。本工作使用 Claude Opus 4.8。The kernel configs in this work are generated using the LLMSeededLFBOTreeSearch autotuner, which uses an LLM to identify a set of promising config candidates as seeds for the subsequent LFBOTreeSearch. For this work, we use Claude Opus 4.8.

从质量更高的候选项开始搜索，有助于自动调优器找到性能更优的配置。它还可能通过将搜索引向搜索空间中更有希望的区域，缩短整体自动调优时间。Starting the search from higher-quality candidates helps the autotuner discover better-performing configs. It may also reduce overall autotuning time by directing the search toward more promising regions of the search space.

性能评估Performance Evaluation

我们在多种稠密模型和量化格式上，从内核和端到端服务两个层面评估 Helion 线性后端。为了解性能如何随模型规模变化，我们对 Qwen3-1.7B、Qwen3-4B、Qwen3-8B、Qwen3-14B 和 Qwen3-32B 进行了基准测试。此外，我们还纳入了最新的 Qwen3.8-27B 模型，以评估该后端在较新模型架构上的表现。We evaluate the Helion linear backend at both the kernel and end-to-end serving levels across a range of dense models and quantization formats. To understand how performance scales with model size, we benchmark Qwen3-1.7B, Qwen3-4B, Qwen3-8B, Qwen3-14B, and Qwen3-32B. We additionally include the latest Qwen3.8-27B model to evaluate the backend on a more recent model architecture.

对于每个模型，我们都会在 NVIDIA Hopper GPU 上评估三种常用的 8 位量化格式。例如，对于 Qwen3.8-27B，我们会进行以下基准测试：For each model, we evaluate three commonly used 8-bit quantization formats on NVIDIA Hopper GPUs. For example, for Qwen3.8-27B, we benchmark: 

• AzatAI/Qwen3.8-27B-FP8-dynamic (FP8_Dynamic)
• Freaksterz/Qwen3.8-27B-SmoothQuant-W8A8-INT8 (W8A8_INT8)
• Qwen/Qwen3.8-27B-FP8 (Block_FP8) 

所有基准测试均在 NVIDIA H100 80GB HBM3 GPU 上进行。All benchmarks are performed on an NVIDIA H100 80GB HBM3 GPU.

内核级评估Kernel-Level Evaluation

我们首先单独评估 Helion 量化 GEMM 内核，以了解细粒度调优在不受推理栈其余部分影响的情况下带来的性能收益。我们将每个 Helion 内核与 Hopper 上默认 vLLM 线性后端使用的对应内核进行比较：We first evaluate the Helion quantized GEMM kernels in isolation to understand the performance benefit of fine-grained tuning independent of the rest of the inference stack. Each Helion kernel is compared against the corresponding kernel used by the default vLLM linear backend on Hopper:

• FP8_Dynamic: Helion vs. CUTLASS
• W8A8_INT8: Helion vs. CUTLASS
• Block_FP8: Helion vs. FlashInfer/DeepGEMM

图 3：对于端到端基准测试模型所用的所有输入形状，Helion 量化 GEMM 内核相较于默认 vLLM 内核库的加速比分布。菱形表示几何平均加速比。Fig. 3: Helion quantized GEMM kernel speedup distribution over the default vLLM kernel libraries across all input shapes used by the end-to-end benchmarked models. Diamonds indicate the geometric mean speedup.

与 CUTLASS 相比，Helion 在 FP8_Dynamic 上实现了 1.110× 的几何平均加速比，在 W8A8_INT8 上实现了 1.178× 的几何平均加速比；与 FlashInfer 相比，Helion 在 Block_FP8 上实现了 1.149× 的几何平均加速比；与 DeepGEMM 相比，Helion 在 Block_FP8 上实现了 1.177× 的几何平均加速比。该分布还表明，不同输入形状下的性能各不相同，这凸显了精细调优的价值，而非依赖单一内核配置或通用调度启发式方法。Helion achieves geometric mean speedups of 1.110× for FP8_Dynamic over CUTLASS, 1.178× for W8A8_INT8 over CUTLASS, 1.149× for Block_FP8 over FlashInfer, and 1.177× for Block_FP8 over DeepGEMM. The distribution also shows that performance varies across individual input shapes, highlighting the value of fine-grained tuning rather than relying on a single kernel config or generic dispatch heuristic.

端到端评估End-to-End Evaluation

接下来，我们评估内核层面的改进能否转化为端到端的服务性能提升。We next evaluate whether the kernel-level improvements translate into end-to-end serving performance.

服务器设置Server Setup

我们使用以下命令启动 vLLM 服务器：We use the following command to start the vLLM server:

vllm serve \
    --model "$MODEL" \
    --max-num-seqs 32 \
    --tensor-parallel-size 1 \
    --no-enable-prefix-caching \
    --linear-backend helion

相关选项如下：The relevant options are:

• --max-num-seqs 32：Helion 线性后端当前采用 32 个 num_tokens 的混合调度阈值。因此，我们将评估重点放在不超过 32 的批量大小上，此时 Helion 内核处于启用状态，可以直接影响端到端性能。
  --max-num-seqs 32: The Helion linear backend currently uses a hybrid dispatch threshold of 32 num_tokens. We therefore focus the evaluation on batch sizes up to 32, where Helion kernels are active and can directly affect end-to-end performance.
• --no-enable-prefix-caching：有意禁用前缀缓存，以避免其影响基准测试结果。
  --no-enable-prefix-caching: Prefix caching is intentionally disabled to avoid its impact on benchmark results.
• --linear-backend helion：启用 Helion 线性后端。省略此标志时，将使用默认的线性后端。
  --linear-backend helion: Enables the Helion linear backend. Omitting this flag uses the default linear backend.

基准测试设置Benchmark Setup

我们使用 ShareGPT 数据集运行端到端服务基准测试，配置如下：We run end-to-end serving benchmarks using the ShareGPT dataset with:

vllm bench serve \
    --backend vllm \
    --model "${MODEL}" \
    --endpoint /v1/completions \
    --dataset-name sharegpt \
    --dataset-path "${DATASET}" \
    --max-concurrency "${BATCH_SIZE}" \
    --num-warmups "${NUM_WARMUPS}" \
    --num-prompts "${PROMPTS}" \
    --ignore-eos

每种工作负载均使用默认线性后端和 Helion 线性后端进行基准测试。在 Hopper 上用作基准的默认 vLLM 线性后端如下：Each workload is benchmarked with both the default and Helion linear backends. The default vLLM linear backends used as baselines on Hopper are:

• FP8_Dynamic: CutlassFP8ScaledMMLinearKernel
• W8A8_INT8: CutlassInt8ScaledMMLinearKernel
• Block_FP8: FlashInferFp8DeepGEMMDynamicBlockScaledKernel

端到端基准测试结果End-to-End Benchmark Results

下图展示了在不同模型规模、批次大小和量化格式下，Helion 线性后端相较于对应默认后端的端到端吞吐量提升。The following Figure shows the end-to-end throughput speedup of the Helion linear backend over the corresponding default backend across different model sizes, batch sizes, and quantization formats.

图 4：Helion 线性后端端到端吞吐量加速比。Fig. 4: Helion linear backend end-to-end throughput speedup.

总体而言，Helion 线性后端在所有经过评估的模型和量化格式中，均实现了稳定的端到端性能提升，部分工作负载的吞吐量提升超过 10%。Overall, the Helion linear backend delivers consistent end-to-end performance gains across the evaluated models and quantization formats, with more than 10% throughput improvement for some workloads.

迈向实用的采用模型Toward a Practical Adoption Model

本文介绍的 Helion 线性后端目前已集成到我们的 vLLM 分支中，可投入生产使用。该分支还包含自动调优工具及相关说明，帮助终端用户为其他模型和工作负载生成优化配置。要将其合并到上游，目前仍面临一个挑战：维护、发布并验证大量预调优配置所带来的开销。The Helion linear backend presented in this work is currently available in our vLLM fork and ready for production use. The fork also includes the autotuning tooling and instructions needed for end users to generate optimized configs for additional models and workloads. One remaining challenge for upstream adoption is the maintenance overhead of shipping and validating a large collection of pre-tuned configs.

对于本工作研究的 GEMM 内核等延迟关键型内核，我们正在探索一种模式：将 Helion 内核和集成框架维护在上游，并提供用于功能测试和 CI 的默认配置，同时将针对特定工作负载的自动调优交由最终用户完成。本工作使用的同一套自动化工具可让用户在部署前为目标模型和硬件生成优化配置。这种模式更注重性能和可维护性，而非开箱即用的易用性。我们认为，对于延迟关键型内核，这种权衡切实可行：性能收益足以抵偿额外的离线优化工作，尤其是对于愿意在生产部署前投入这项工作的模型服务提供商而言。For latency-critical kernels such as the GEMM kernels studied in this work, we are exploring a model in which the Helion kernels and integration framework are maintained upstream with a default config for functional testing and CI, while workload-specific autotuning is delegated to end users. The same automated tooling used in this work allows users to generate optimized configs for their target models and hardware before deployment. This model favors performance and maintainability over out-of-the-box usability. We believe this tradeoff is practical for latency-critical kernels, where the performance benefits can justify the additional offline optimization effort, particularly for model serving providers willing to make this investment before production deployment.

不过，精细调优并非 Helion 内核在实践中的唯一采用策略。对于量化、激活和归一化等较小的辅助内核，牺牲部分峰值性能，换取一个小得多且能适用于多种形状的配置集，可能更为可取。我们的初步实验表明，仅用六个配置就能显著提升这些内核的性能。这为辅助内核进入上游提供了另一条路径，同时大幅降低了调优和维护开销。Fine-grained tuning, however, is not the only practical adoption strategy for Helion kernels. For smaller auxiliary kernels, such as quantization, activation, and normalization kernels, it can be preferable to trade some peak performance for a much smaller config set that generalizes across shapes. Our initial experiments show that meaningful performance improvements can be achieved for these kernels with as few as six configs. This provides another path toward upstream adoption with substantially lower tuning and maintenance overhead for auxiliary kernels.

我们鼓励感兴趣的用户试用当前实现，并在 vLLM Helion 线性后端 RFC 中分享使用体验。来自实际部署的反馈将帮助我们评估这些方案，并推动集成以实现更广泛的采用。We encourage interested users to try the current implementation and share their experience in the ⁠vLLM Helion linear backend RFC. Feedback from real-world deployments will help us evaluate these approaches and guide the integration toward broader adoption.

未来工作Future work

未来的工作将重点扩大 Helion 内核对各类推理工作负载和硬件平台的覆盖范围。Future work will focus on expanding Helion kernel coverage across inference workloads and hardware platforms.

硬件覆盖。Helion 团队正在持续扩展并优化对各类硬件平台的后端支持，包括面向 NVIDIA Blackwell GPU 的 CuteDSL 后端，以及针对 AMD GPU 和 TPU 的持续性能优化。随着这些后端日趋成熟，我们计划将 Helion 线性后端扩展到这些硬件目标平台，并重新评估其表现。Hardware coverage. The Helion team is continuing to expand and optimize backend support across hardware platforms, including the CuteDSL backend for NVIDIA Blackwell GPUs and ongoing performance work for AMD GPUs and TPUs. As these backends mature, we plan to extend and re-evaluate the Helion linear backend across these hardware targets.

模型覆盖范围。本工作主要针对稠密模型，在这类模型中，线性后端占推理延迟的很大一部分。对于 MoE 模型，MoE 后端则是更重要的优化目标。我们正在开发 Helion MoE 后端，将 Helion 的内核优化能力应用于这些工作负载，并扩大模型覆盖范围。Model coverage. This work primarily targets dense models, where the linear backend accounts for a significant portion of inference latency. For MoE models, the MoE backend becomes the more important optimization target. We are working on a Helion MoE backend to bring Helion’s kernel optimization capabilities to these workloads and broaden model coverage.

结论Conclusion

Helion 的高级抽象使得开发者能够编写并维护单一内核实现，同时针对不同工作负载和硬件目标进行优化。正如本工作中的量化 GEMM 内核所展示的那样，即使是 Standard GEMM、Split-K 和 Swap-AB 等算法变体，也可以统一到一个实现中，并作为可调选项提供给自动调优器。结合针对每种形状的细粒度内核调优和混合调度，Helion 线性后端在所评估的模型中，在 Hopper GPU 上的表现优于默认的 CUTLASS、DeepGEMM 和 FlashInfer 后端，实现了稳定的端到端性能提升，部分工作负载的吞吐量提升超过 10%。Helion’s high-level abstraction makes it possible to express and maintain a single kernel implementation while optimizing it across different workloads and hardware targets. As demonstrated by the quantized GEMM kernels in this work, even algorithmic variants such as Standard GEMM, Split-K, and Swap-AB can be unified into one implementation and exposed as tunable choices for the autotuner. Combined with per-shape fine-grained kernel tuning and hybrid dispatch, the Helion linear backend outperforms the default CUTLASS, DeepGEMM, and FlashInfer backends on Hopper GPUs across the evaluated models, delivering consistent end-to-end performance gains and throughput improvements exceeding 10% for some workloads.

然而，这些性能提升也伴随着取舍。针对延迟敏感型内核进行细粒度调优可以提升性能，但会增加 AOT 调优工作量；而预先调优好的配置虽能提升开箱即用体验，却需要持续维护。这体现了本文始终在探讨的性能、易用性与可维护性之间更广泛的权衡。These performance gains, however, come with tradeoffs. Fine-grained tuning for latency-critical kernels improves performance but increases AOT tuning effort, while shipping pre-tuned configs improves out-of-the-box usability at the cost of ongoing maintenance. This reflects the broader performance-usability-maintainability tradeoff discussed throughout this post.

如今，我们随 vLLM 分支一同提供预调优配置。为了推动上游采用，我们正在探索另一种模式：将 Helion 内核和集成框架维护在上游，同时把针对特定工作负载的自动调优交由最终用户完成。这种模式更侧重性能和可维护性，而非开箱即用的易用性；不过，Helion 的自动调优框架让性能敏感型部署中的额外优化步骤切实可行。Today, we ship the pre-tuned configs with our vLLM fork. For upstream adoption, we are exploring a different model: maintain the Helion kernels and integration framework upstream while delegating workload-specific autotuning to end users. This favors performance and maintainability over out-of-the-box usability, but Helion’s automated tuning framework makes the additional optimization step practical for performance-sensitive deployments.

致谢Acknowledgments

这项工作得到了红帽 OCTO 和 vLLM 团队，以及 Meta Helion 团队众多成员的支持。特别感谢我们的同事 Richard Zou 和 Jongsok Choi 在整个工作过程中提供的反馈与支持。This work was supported by many contributors across the OCTO and vLLM teams at Red Hat, as well as the Helion team at Meta. In particular, we would like to thank our colleagues: Richard Zou and Jongsok Choi for their feedback and support throughout this work.