Helion:高性能内核开发新范式,代码缩减,性能暴涨2.6倍!

2025-10-22人工智能

Image

深度解析:Helion如何重塑高性能机器学习内核开发范式

在当代机器学习领域,对高性能计算的需求持续攀升,这直接推动了定制化内核(custom kernels)的广泛应用。这些定制内核虽然能够提供卓越的性能表现,但其编写往往依赖于低级别、硬件特定的编程语言。这种现状带来了长期的维护负担:一个为特定硬件架构精心优化过的内核,很快就会演变为技术债务,难以且成本高昂地移植到其他架构上。这一挑战不仅阻碍了开发和创新,也迫使开发者不得不在开发效率与性能表现之间做出艰难抉择。

Helion的出现旨在解决这一核心矛盾。它通过将一种高级别的、嵌入Python的领域特定语言(DSL)编译成可自动调优的Triton代码,从而在易用性极强的PyTorch与底层语言的卓越性能之间搭建起一个新的抽象层。Helion自动化处理了诸如张量索引、内存管理以及硬件特定调优等繁琐且易错的任务,使得开发者能够将精力集中于算法逻辑本身,而非深陷于硬件实现的细节之中。

Helion巧妙地结合了PyTorch用户熟悉的语法与强大的自动调优引擎,该引擎能够自动搜索最佳的内核配置。这一系统不仅在不同硬件架构间实现了性能的可移植性,还大幅降低了开发工作量。
Helion Logo

新一代DSL诞生的深层动因

选择正确的抽象层进行内核开发,是一项战略性决策,它直接影响着性能、可维护性以及开发者的效率。当前的编程语言和抽象工具往往迫使开发者在“底层控制”与“高级生产力”之间陷入两难境地。这两种极端情况都各有利弊:

  1. CUDA/Gluon/TLX等低级语言: 直接使用CUDA等语言编写内核能够提供最大的控制权,但要实现高性能往往需要投入巨大的精力。这些内核高度依赖于特定的硬件,并且难以适应新的架构。
  2. Triton: 尽管Triton是向前迈出的重要一步,但它仍然需要大量的手动工作。开发者需要负责显式管理张量索引,定义自动调优的搜索空间,管理内核参数,并且改变优化策略通常需要对代码进行大量的重写。
  3. PyTorch的torch.compile torch.compile虽然提供了出色的易用性,但在精细化控制方面存在局限。需要指定精确融合策略的用户,常常会发现这种高级抽象过于受限。

Helion正是为了打破这种僵局而生,旨在提供一种兼顾高效开发和极致性能的新途径。

Helion编程模型:“带瓦片(Tiles)的PyTorch”

Helion编程模型的核心目标是最大限度地减少样板代码,并充分利用开发者现有的PyTorch知识。这种设计理念通过提供熟悉且直观的语法,大大加快了正确、高效内核的创建速度,其特点可概括为“带瓦片(Tiles)的PyTorch”。一个典型的Helion内核,例如下面的矩阵乘法示例,由两个协同工作的独立部分组成:

import torch, helion, helion.language as hl

@helion.kernel()
def matmul(x: torch.Tensor, y: torch.Tensor) -> torch.Tensor:
    # --- Host Code (runs on CPU) ---
    m, k = x.size()
    k, n = y.size()
    out = torch.empty([m, n], dtype=x.dtype, device=x.device)
    # --- Device Code (compiles to a Triton kernel) ---
    for tile_m, tile_n in hl.tile([m, n]):
        acc = hl.zeros([tile_m, tile_n], dtype=torch.float32)
        for tile_k in hl.tile(k):
            acc = torch.addmm(acc, x[tile_m, tile_k], y[tile_k, tile_n])
        out[tile_m, tile_n] = acc
    return out
  1. 主机代码(Host Code): hl.tile for循环之前的代码段是标准的PyTorch代码。它主要用于设置任务,例如分配输出张量和计算形状。Helion会自动处理这些值到设备代码的传递,无需手动管理参数。
  2. 设备代码(Device Code): hl.tile for循环是内核的核心部分。此部分被编译成一个单一的高性能Triton内核,并在GPU上并行执行。hl.tile指令将内核的迭代空间细分为“瓦片”(tiles)。程序员只需指定对迭代空间进行瓦片划分,而具体的实现细节,例如瓦片大小、迭代顺序和内存布局优化,则由Helion的自动调优器处理。自动调优器系统地探索目标硬件的最佳配置。

在内核主体内,开发者可以使用标准PyTorch操作符,如torch.addmm以及其他逐点或归约操作。Helion利用TorchInductor,一个PyTorch 2的核心组件,自动将这些PyTorch调用映射到其对应的低级Triton实现。这带来了强大的可用性优势:熟悉PyTorch意味着已经掌握了Helion的大部分内容。

Helion还具备模板化功能,允许将lambda函数作为参数传递给内核,这些lambda函数可以在闭包中捕获额外的参数。例如,如此示例所示,这对于实现带有可定制后处理逻辑(epilogues)的通用内核尤其有用。一个lambda函数可以捕获在周围作用域中定义的张量,Helion的编译器会自动检测此变量并将其作为参数纳入生成的Triton内核中。这显著减少了通过多层函数调用传递新输入的样板代码,从而能够创建高度可重用和通用的内核。

Helion极大地简化了内核实现:Attention内核在Helion中仅需30行代码,而在Triton中需要120行,在CUDA中则需要数千行。这种高级声明式编程模型之所以能实现高性能,归功于支撑整个系统的核心机制:自动调优引擎。

Helion的自动调优器:通过隐式搜索空间生成最优内核

Helion的关键差异化优势在于其自动化的、AOT(Ahead-Of-Time)自动调优引擎。在Triton中,开发者需要手动定义优化的搜索空间,这要求显式列举所有要测试的配置,这一过程既繁琐又限制了探索的范围。Helion通过引入隐式搜索空间改变了这种局面。高级语言会自动构建一个包含各种实现选择的庞大多维搜索空间。例如,一个简单的hl.tile指令就能引发一系列复杂的优化考量。

自动调优工作流程

当内核首次运行时未指定配置时,自动调优器会启动一个自动化搜索过程。这个过程通常需要大约10分钟,它会使用差分进化(Differential Evolution)或模式搜索(Pattern Search)等搜索策略,评估数千个候选的Triton内核配置,以识别针对给定输入形状和硬件的优化参数集。完成后,自动调优器会打印出它发现的最佳配置:

[ 586s ] Autotuning complete in 586 . 6s after searching 1520 configs.
One can hardcode the best config and skip autotuning with:
    @helion.kernel(config=helion.Config(block_sizes=[ 64 , 64 , 64 ], loop_orders=[[ 0 , 1 ]], l2_groupings=[ 4 ], range_unroll_factors=[ 0 , 1 ], range_warp_specializes=[None, False], range_num_stages=[ 0 , 3 ], range_multi_buffers=[None, False], range_flattens=[None, None], num_warps= 8 , num_stages= 6 , indexing= 'block_ptr' , pid_type= 'flat' ))

开发者可以将此配置复制到其源代码的@helion.kernel()装饰器中。这指示Helion在后续运行期间完全跳过搜索过程。在生产环境中,这会带来快速、确定性的编译,生成单一的、预优化过的Triton内核,以远低于手动工作的精力,实现与精心手动调优内核相当的性能。

@helion.kernel(config=helion.Config(
    block_sizes=[ 64 , 64 , 64 ],
    loop_orders=[[ 0 , 1 ]],
    l2_groupings=[ 4 ],
    range_unroll_factors=[ 0 , 1 ],
    range_warp_specializes=[ None , False],
    range_num_stages=[ 0 , 3 ],
    range_multi_buffers=[ None , False],
    range_flattens=[ None , None ],
    num_warps= 8 ,
    num_stages= 6 ,
    indexing= 'block_ptr' ,
    pid_type= 'flat'
))
def matmul(x: torch.Tensor, y: torch.Tensor) -> torch.Tensor:
    ...

即便未指定完整的配置,开发者也可以在@helion.kernel()中提供部分配置,此时Helion只会探索这些指定的配置以选择最快的实现。

配置空间

配置空间代表了Helion自动化的一系列实现选择。这个空间是Helion性能可移植性的主要来源,因为它允许单个内核定义适应不同硬件架构和输入张量尺寸的独特特性。探索这个空间是Helion优于手动编写内核的原因,因为手动编写的内核通常是针对特定条件进行调优的。自动调优器探索了从数据移动到线程映射等各个方面的广泛参数。下表详细列出了配置选项。

参数 描述
indexing 在Triton中,开发者可以在三种不同的内存访问方法之间进行选择:指针算术(pointer arithmetic)、块指针(block pointers)和张量描述符(tensor descriptors)。张量描述符利用英伟达 Hopper/Blackwell GPU 上的张量内存加速器(TMAs)。最佳选择取决于硬件架构和内存访问模式,但在这几种方法之间切换通常需要大量的代码重写。Helion通过配置pointer, block_ptr, tensor_descriptor来自动化选择。
block_sizes block_sizes 参数决定了hl.tile操作在每个维度上如何划分数据块。不同的块大小会显著影响缓存效率和并行度。
flatten_loops flatten_loops 配置控制是否将hl.tile生成的嵌套循环进行扁平化处理,这有助于在某些情况下提升指令并行度。
loop_orders, l2_grouping loop_orders 指明了迭代循环的顺序,而l2_grouping则优化了二级缓存的访问模式,这些都对内存局部性和性能有关键影响。
reduction_loops 在执行归约操作(例如,对张量维度进行sum()求和)时,一种“持久归约”(persistent reduction)方法会为单个瓦片处理整个归约维度,这对于小维度而言速度很快。然而,如果归约维度很大,这种方法可能会导致较高的寄存器压力,进而引发寄存器溢出,降低性能。作为替代方案,开发者可以编写一个循环,以较小的块迭代归约维度。Helion通过自动调优来选择最优策略,无需手动更改代码。
pid_type Helion自动化计算网格大小和程序ID(PIDs)到数据瓦片(data tiles)的映射。pid_type配置允许自动调优器探索各种映射策略,而无需任何手动代码更改:包括flat(扁平)、xyz(三维)、persistent_blocked(持久分块)和persistent_interleaved(持久交错)。
load_eviction_policy 此参数影响tl.load操作的数据逐出策略,对于GPU共享内存管理至关重要。
Triton configs: num_warps, num_stages, range_unroll_factors, range_warp_specializes, range_num_stages, range_multi_buffers, range_flattens Helion自动化探索标准的Triton可调参数,减轻了开发者手动调优的工作量。这些参数控制着线程组的数量、流水线阶段、循环展开因子、线程束专门化、多缓冲区使用和循环扁平化等,均旨在最大化GPU的吞吐量和效率。

性能分析与基准测试

英伟达B200上的性能表现

AMD MI350X上的性能表现

案例研究一:超越高度优化的CuTe DSL内核

Helion实现了一个RMSNorm反向传播内核,这个内核在不到一天的时间内完成编写。在与高度手动优化的Quack内核(使用CuTe DSL编写)、torch.compile(带有最大自动调优)以及手动编写的Triton内核进行对比测试后发现,Helion的性能表现与前者持平或在许多情况下超越。

这些结果得益于Helion底层的编译器架构,该架构旨在高效支持大规模的自动调优搜索。

案例研究二:Helion与TileLang的基准测试

新媒网跨境获悉,研究还比较了Helion与TileLang中的对应实现进行了比较。在H100 GPU上,Helion展现了最高的性能,在不同配置下比TileLang加速2.12倍至2.63倍,比Triton加速1.2倍至1.85倍。

高级编译器架构

Helion的编译器架构旨在将Python函数逐步降低为通过TorchInductor优化后的Triton代码。编译流水线包含以下关键阶段:

  1. Python AST解析: 过程始于将内核的Python源代码解析为抽象语法树(AST)。
  2. 类型传播与元数据: 一个自定义的遍历(pass)会遍历AST,用类型信息和其他基本元数据标注每个节点,以创建扩展的AST。
  3. 降级到设备IR: 经过标注的树被降级为Helion的主要中间表示(IR)。设备IR是静态单赋值(SSA)形式的FX图集合,其中每个图代表程序的一个基本块。该图中的每个节点都包含一个指向Inductor IR节点的指针,用于代码生成。
  4. 编译器遍历(Compiler Passes): 一系列转换遍历(transformation passes)被应用于设备IR。这些遍历实现了关键的语义更改,例如归约滚动优化(reduction rolling optimization),它将持久归约转换为循环归约。
  5. 带配置的代码生成: 在最后阶段,代码生成器接收两个输入:转换后的设备IR和自动调优的配置。它使用这些信息生成最终的Triton代码。

一个关键的架构决策是,性能关键的配置仅在流水线的最后阶段——代码生成时才应用。这使得编译过程的大部分,从解析到IR转换,可以在自动调优搜索之前只运行一次,从而确保探索数千种配置的计算效率。

结语

Helion填补了当前机器学习内核创作领域的一个关键空白。它将熟悉且高级的类PyTorch语法与强大、提前(AOT)的自动调优引擎相结合,在开发者生产力、精细控制和性能可移植性之间取得了独特的平衡。Helion赋能开发者编写出可移植、面向未来且能够达到最先进性能的内核,而无需深入的硬件专业知识,从而为高性能机器学习内核开辟了一种新的、更高效的范式。

Helion已于2025年10月22日发布Beta版本,期待来自社区的反馈、错误报告和贡献。更多信息请参考:

  • Helion 源代码
  • Helion 在 PTC 2025 的演讲

新媒网(公号: 新媒网跨境发布),是一个专业的跨境电商、游戏、支付、贸易和广告社区平台,为百万跨境人传递最新的海外淘金精准资讯情报。

本文来源:新媒网 https://nmedialink.com/posts/helion-redefines-kernel-dev-26x-boost.html

评论(0)

暂无评论,快来抢沙发~
Helion, a new DSL, simplifies high-performance machine learning kernel development by bridging PyTorch ease-of-use and low-level language performance. It automates tuning and hardware-specific tasks, enabling portable, efficient kernels with less code, outperforming optimized kernels in some cases. Released as a Beta in October 2025.
发布于 2025-10-22
查看人数 79
人民币汇率走势
CNY
亚马逊热销榜
共 0 SKU 上次更新 NaN:NaN:NaN
类目: 切换分类
暂无数据
暂无数据
关注我们
新媒网跨境发布
本站原创内容版权归作者及NMedia共同所有,未经许可,禁止以任何形式转载。