Helion: 고성능 이식 가능한 ML 커널을 위한 고수준 DSL | 파이토치 한국 사용자 모임

bot · 5월 24, 2026, 12:24오전

Helion 소개 / Introduction to Helion

현대 머신러닝에서 고성능 연산에 대한 수요는 커스텀 커널의 급증으로 이어졌습니다. 이러한 커널은 뛰어난 성능을 제공할 수 있지만, 종종 저수준(low-level)의 하드웨어 종속적인 언어로 작성됩니다. 이는 장기적인 유지보수 부담을 만들어냅니다: 하나의 하드웨어 아키텍처에 맞춰 세심하게 최적화된 커널은 금세 기술 부채(technical debt)가 되어, 다른 하드웨어로 이식하기 어렵고 비용이 많이 듭니다. 이러한 문제는 개발과 혁신을 저해하며, 개발자들에게 생산성과 성능 사이에서 양자택일을 강요합니다.

In modern machine learning, the demand for high-performance computation has led to a proliferation of custom kernels. While these kernels can deliver impressive performance, they are often written in low-level, hardware-specific languages. This creates a long-term maintenance burden: A kernel meticulously optimized for one hardware architecture quickly becomes technical debt, difficult and costly to port to another. This challenge hinders development and innovation and forces developers to choose between productivity and performance.

Helion은 고수준 Python 내장 도메인 특화 언어(DSL)를 자동 튜닝된 Triton 코드로 컴파일하여 이러한 갈등을 해결합니다. Helion은 PyTorch의 사용자 친화적인 단순함과 저수준 언어의 성능을 연결하는 새로운 추상화 계층을 구축합니다. Tensor 인덱싱, 메모리 관리, 하드웨어별 튜닝과 같은 번거롭고 오류가 발생하기 쉬운 작업을 자동화함으로써, Helion은 개발자가 하드웨어별 구현 세부 사항 대신 알고리즘 로직에 집중할 수 있도록 합니다. Helion은 친숙한 PyTorch 중심 문법과 최적의 커널 구성을 자동으로 탐색하는 강력한 오토튜닝 엔진을 결합하여 이러한 균형을 달성합니다. 그 결과, 개발 노력을 대폭 줄이면서도 하드웨어 아키텍처 간 성능 이식성(performance portability)을 제공하는 시스템이 탄생했습니다.

Helion resolves this conflict by compiling a high-level Python-embedded domain-specific language (DSL) into automatically tuned Triton code. It establishes a new layer of abstraction that bridges the user-friendly simplicity of PyTorch with the performance of a lower level language. By automating tedious and error-prone tasks like tensor indexing, memory management, and hardware-specific tuning, Helion empowers developers to focus on algorithmic logic rather than hardware-specific implementation details. Helion achieves this balance by pairing a familiar, PyTorch-centric syntax with a powerful autotuning engine that automates the complex search for optimal kernel configurations. This results in a system that delivers performance portability across hardware architectures while drastically reducing development effort.

새로운 DSL의 필요성 / Motivation for a New DSL

커널 개발에 적합한 추상화 수준을 선택하는 것은 성능, 유지보수성, 개발 속도에 직접적인 영향을 미치는 전략적 결정입니다. 현재의 프로그래밍 언어와 추상화는 개발자들에게 저수준 제어와 고수준 생산성 사이의 잘못된 이분법을 강요합니다. 양쪽 모두 장단점이 있음에도 불구하고 말입니다.

Choosing the right abstraction level for kernel development is a strategic decision that directly impacts performance, maintainability, and the developer velocity. The current programming languages and abstractions force developers into a false dichotomy between low-level control and high-level productivity. Both ends of the spectrum come with advantages and drawbacks.

CUDA/Gluon/TLX: CUDA와 같은 언어로 커널을 직접 작성하면 매우 넓은 범위를 제어할 수 있지만, 높은 성능을 달성하기 위해 상당한 노력이 필요할 수 있습니다. 이러한 커널은 특정 하드웨어에 고도로 특화되어 있어 새로운 아키텍처에 적응하기 어려울 수 있습니다.
Triton: Triton은 큰 진전이지만, 여전히 상당한 수작업이 필요합니다. 개발자가 직접 Tensor 인덱싱을 관리하고, 오토튜닝을 위한 탐색 공간을 정의하고, 커널 인자를 관리해야 하며, 최적화 전략을 변경하려면 종종 코드를 크게 재작성해야 합니다.
PyTorch: PyTorch와 torch.compile 같은 프레임워크는 뛰어난 사용 편의성을 제공하지만, 세밀한 제어가 제한적입니다. 정확한 퓨전(fusion) 전략을 지정해야 하는 사용자에게는 고수준 추상화가 너무 제한적일 수 있습니다.

CUDA/Gluon/TLX: Writing kernels directly in languages like CUDA offers maximum control but can require significant effort to achieve high performance. These kernels are highly specialized to specific hardware and can be difficult to adapt to new architectures.

Triton: While Triton represents a major step forward, it still requires significant manual effort. Developers are responsible for explicitly managing tensor indexing, defining search spaces for autotuning, managing kernel arguments, and changing the optimization strategy can often require significant code rewrites.

PyTorch: While frameworks like PyTorch and torch.compile offer exceptional ease of use, they provide limited fine-grained control. Users who need to specify exact fusion strategies often find the high-level abstraction too restrictive.

Helion 프로그래밍 모델: “타일이 있는 PyTorch” / Helion Programming Model: “PyTorch with Tiles”

Helion의 프로그래밍 모델은 보일러플레이트를 최소화하고 개발자의 기존 PyTorch 지식을 활용하는 것을 목표로 합니다. 이 설계 철학은 친숙하고 직관적인 문법을 제공하여 정확하고 효율적인 커널 작성을 가속화하며, “타일이 있는 PyTorch(PyTorch with Tiles)”로 설명할 수 있습니다.

The goal of Helion’s programming model is to minimize boilerplate and leverage developers’ existing knowledge of PyTorch. This design philosophy accelerates the creation of correct, efficient kernels by providing a familiar and intuitive syntax, which can be described as “PyTorch with Tiles”.

아래의 행렬 곱셈 예시와 같이, 일반적인 Helion 커널은 함께 동작하는 두 가지 구분된 부분으로 구성됩니다:

A typical Helion kernel, such as the matrix multiplication example below, is composed of two distinct parts that work in concert:

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

호스트 코드(Host Code): 가장 바깥쪽 hl.tile for 루프 밖의 코드는 표준 PyTorch 코드입니다. 주로 출력 Tensor 할당이나 shape 계산과 같은 설정 작업에 사용됩니다. Helion은 이러한 값을 디바이스 코드로 자동 전달하여 수동 인자 관리의 필요성을 없앱니다.
디바이스 코드(Device Code): 가장 바깥쪽 hl.tile for 루프 안의 코드가 커널의 핵심입니다. 이 부분은 GPU에서 병렬로 실행되는 단일 고성능 Triton 커널로 컴파일됩니다.

Host Code: The code outside the outermost hl.tile for loop is standard PyTorch code. It is primarily used for setup tasks such as allocating output tensors and computing shapes. Helion automatically handles the passing of these values to the device code, eliminating the need for manual argument management.

Device Code: The code inside the outermost hl.tile for loop is the core of the kernel. This section is compiled into a single, high-performance Triton kernel that executes in parallel on a GPU.

핵심 언어 구성 요소인 hl.tile은 커널의 반복 공간(iteration space)을 타일로 분할합니다. 프로그래머는 반복 공간을 타일링하도록 지정하기만 하면 되고, 타일 크기, 반복 순서, 메모리 레이아웃 최적화 같은 구체적인 구현 세부 사항은 Helion의 오토튜너가 대상 하드웨어에 최적인 구성을 체계적으로 탐색하여 처리합니다.

The core language construct, hl.tile, subdivides the kernel’s iteration space into tiles. The programmer only specifies to tile the iteration space, and the specific implementation details, such as tile sizes, iteration order, and memory layout optimizations, are handled by Helion’s autotuner, which systematically explores the optimal configuration for the target hardware.

커널 본문 내에서 개발자는 torch.addmm이나 기타 포인트와이즈(pointwise) 또는 리덕션(reduction) 연산과 같은 표준 PyTorch 연산자를 사용할 수 있습니다. Helion은 PyTorch 2의 핵심 구성 요소인 TorchInductor를 활용하여 이러한 PyTorch 호출을 해당하는 저수준 Triton 구현으로 자동 매핑합니다. 이는 강력한 사용성 이점을 제공합니다: PyTorch에 익숙하다면 Helion의 대부분을 이미 아는 것과 마찬가지입니다.

Within the kernel body, developers can use standard PyTorch operators like torch.addmm and other pointwise or reduction operations. Helion leverages TorchInductor, a core component of PyTorch 2, to automatically map these PyTorch calls to their corresponding low-level Triton implementations. This delivers a powerful usability benefit: Familiarity with PyTorch means you already know most of Helion.

Helion은 또한 커널에 인자로 전달할 수 있는 람다 함수(클로저에서 추가 인자를 캡처할 수 있는)를 통한 템플릿 기능을 제공합니다. 이 예시에서 보듯이, 이는 커스터마이즈 가능한 에필로그(epilogue)가 있는 범용 커널을 구현할 때 특히 유용합니다. 예를 들어, 람다 함수가 주변 스코프에서 정의된 Tensor를 캡처할 수 있습니다. Helion 컴파일러는 이 변수를 자동으로 감지하여 생성된 Triton 커널의 인자로 만듭니다. 이는 새로운 입력을 여러 계층의 함수 호출을 통해 전달하기 위한 상당량의 보일러플레이트 코드를 없애고, 높은 재사용성을 갖는 범용 커널의 생성을 가능하게 합니다.

Helion also includes templating capability that allows lambda functions, that may capture additional arguments in closures, to be passed in as arguments to a kernel. As shown in this example, this is particularly useful for implementing generic kernels with customizable epilogues. For instance, a lambda function can capture a tensor defined in the surrounding scope. Helion’s compiler automatically detects this variable and makes it into an argument in the generated Triton kernel. This eliminates a significant amount of boilerplate code to pass new inputs through multiple layers of function calls, enabling creation of highly reusable and generic kernels.

Helion은 커널 구현을 획기적으로 단순화합니다: Attention 커널이 Helion에서는 30줄에 불과한 반면, Triton에서는 120줄, CUDA에서는 수천 줄이 필요합니다. 이 고수준 선언적 프로그래밍 모델의 성능은 시스템 전체를 뒷받침하는 핵심 메커니즘인 오토튜닝 엔진에 의해 보장됩니다.

Helion makes kernel implementations radically simpler: the Attention kernel is just 30 lines in Helion, compared to 120 lines in Triton and thousands of lines in CUDA. This high-level, declarative programming model is made performant by the core mechanism that underpins the entire system: the autotuning engine.

Helion의 오토튜너: 암묵적 탐색 공간을 통한 최적 커널 생성 / Helion’s Autotuner: Generating Optimal Kernels via Implicit Search Spaces

Helion의 핵심 차별점은 자동화된 AOT(ahead-of-time) 오토튜닝 엔진입니다. Triton에서는 개발자가 최적화를 위한 탐색 공간을 수동으로 정의해야 합니다. 이는 테스트할 모든 구성을 명시적으로 나열해야 하는 번거로운 작업으로, 탐색 범위가 제한됩니다.

Helion’s key differentiator is its automated, ahead-of-time (AOT) autotuning engine. In Triton, developers are responsible for manually defining the search space for optimizations. This requires explicitly enumerating every configuration to be tested, a tedious process that limits the scope of exploration.

Helion은 암묵적 탐색 공간(implicit search spaces)으로 이러한 방식을 바꿉니다. 고수준 언어가 구현 선택에 대한 광범위한 다차원 탐색 공간을 자동으로 구성합니다. 예를 들어, 하나의 hl.tile 호출은 오토튜너에게 다양한 블록 크기, 루프 순서, 반복 공간을 단일 차원으로 평탄화(flatten)할지 여부를 탐색하도록 암묵적으로 지시합니다. 따라서 하나의 Helion 커널 정의가 수천 개의 Triton 구성에 매핑되어, 오토튜너가 훨씬 크고 풍부한 탐색 공간에서 더 우수한 구성을 발견할 수 있습니다.

Helion changes this dynamic with implicit search spaces. The high-level language automatically constructs a vast, multi-dimensional search space over implementation choices. For example, a single hl.tile call implicitly instructs the autotuner to explore different block sizes, loop orderings, and whether to flatten the iteration space into a single dimension. One Helion kernel definition thus maps to thousands of Triton configurations, allowing the autotuner to create a much larger and richer search space to discover a superior configuration.

오토튜닝 워크플로우 / The Autotuning Workflow

지정된 구성 없이 커널을 처음 실행하면, 오토튜너가 자동 탐색을 시작합니다. 이 과정은 일반적으로 약 10분이 소요되며, 차등 진화(Differential Evolution)나 패턴 탐색(Pattern Search) 같은 탐색 전략을 사용하여 주어진 입력 shape과 하드웨어에 최적화된 매개변수 집합을 찾기 위해 수천 개의 후보 Triton 커널 구성을 평가합니다. 완료 시, 오토튜너는 발견한 최적의 구성을 출력합니다:

When a kernel is run for the first time without a specified configuration, the autotuner initiates an automated search. This process, which typically takes around 10 minutes, evaluates thousands of candidate Triton kernel configurations using search strategies like Differential Evolution or Pattern Search to identify optimized sets of parameters for the given input shapes and hardware. Upon completion, the autotuner prints the single best configuration it discovered:

[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 커널을 생성하며, 훨씬 적은 노력으로 세심하게 수동 튜닝한 커널과 동등한 성능을 제공합니다.

The developer can copy this config into the @helion.kernel() decorator in their source code. This instructs Helion to bypass the search process entirely during subsequent runs. In a production environment, this results in fast, deterministic compilation that generates the single, pre-optimized Triton kernel, delivering performance equivalent to a meticulously hand-tuned kernel with far less effort.

@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은 해당 구성들만 탐색하여 가장 빠른 구현을 선택합니다.

The developer can also specify a list of configs in @helion.kernel(), in which case Helion will explore only those configs to choose the fastest implementation.

구성 공간 / The Configuration Space

구성 공간(configuration space)은 Helion이 자동화하는 구현 선택의 집합을 나타냅니다. 이 공간은 단일 커널 정의를 다양한 하드웨어 아키텍처와 입력 Tensor 크기의 고유한 특성에 맞게 적응시킬 수 있게 해주는, Helion 성능 이식성의 핵심 원천입니다. 이 공간을 탐색하는 것이 수동으로 작성된 커널(종종 특정 조건에 맞춰 튜닝됨)에 비해 Helion이 우위를 갖는 이유입니다.

The configuration space represents the set of implementation choices that Helion automates. This space is the primary source of Helion’s performance portability, as it allows a single kernel definition to be adapted to the unique characteristics of different hardware architectures and input tensor sizes. Exploring this space is what gives Helion its advantage over manually written kernels, which are often tuned for a specific set of conditions.

오토튜너는 데이터 이동부터 스레드 매핑까지 모든 것을 제어하는 광범위한 매개변수를 탐색합니다. 아래 표는 구성 옵션을 상세히 설명합니다.

The autotuner explores a wide range of parameters that control everything from data movement to thread mapping. The table below details the configuration options.

매개변수 / Parameter 설명 / Description indexing Triton에서 개발자는 메모리 접근을 위해 포인터 산술(pointer arithmetic), 블록 포인터(block pointer), 그리고 NVIDIA Hopper/Blackwell GPU의 Tensor Memory Accelerator(TMA)를 활용하는 Tensor 디스크립터(tensor descriptor) 등 세 가지 방식을 선택할 수 있습니다. 최적의 선택은 하드웨어 아키텍처와 메모리 접근 패턴에 따라 다르지만, 방식 간 전환에는 상당한 코드 재작성이 필요합니다.Helion은 indexing 구성 매개변수('pointer', 'block_ptr', 'tensor_descriptor')로 이 복잡성을 추상화하여, 오토튜너가 동일한 소스에서 세 가지 메모리 접근 패턴의 코드를 자동으로 생성하고 벤치마킹할 수 있게 합니다. 이는 개발을 단순화할 뿐만 아니라 코드 변경 없이 TMA 같은 최신 하드웨어 기능을 활용할 수 있도록 보장합니다. block_sizes block_sizes 매개변수는 hl.tile 루프의 각 차원에 대한 타일 크기 목록으로, 각 스레드 블록이 처리하는 데이터 양을 결정합니다. 이는 레지스터 사용량, 공유 메모리 요구사항, 병렬성에 영향을 미칩니다. flatten_loops flatten_loops 옵션은 hl.tile 루프의 다차원 타일링 공간을 단일 차원으로 평탄화하는 것을 제어하여, 개발자가 커널 코드를 재작성하지 않고도 오토튜너의 탐색 공간을 확장합니다. loop_orders, l2_grouping 데이터 지역성(locality)을 최적화하기 위해 Helion은 두 가지 구성 옵션을 제공합니다. loop_orders 매개변수는 오토튜너가 중첩 타일의 반복 순서를 치환할 수 있게 하여, Tensor 레이아웃에 따라 캐시 적중률에 영향을 줄 수 있습니다. l2_grouping 구성은 L2 캐시에서의 데이터 재사용을 개선하기 위해 스레드 블록 할당을 재정렬하는 기법인 PID 스위즐링(swizzling)을 활성화합니다. Triton에서는 이러한 변환을 위해 복잡한 루프 구조와 인덱스 계산을 재작성해야 합니다. reduction_loops 리덕션(예: Tensor 차원에 대한 sum()) 수행 시, 영속적 리덕션(persistent reduction)은 단일 타일에 대한 전체 리덕션 차원을 처리하므로 작은 차원에서는 빠릅니다. 그러나 리덕션이 크면 높은 레지스터 압력을 야기하여 레지스터 스필링(spilling)과 낮은 성능으로 이어질 수 있습니다. 대안으로 개발자가 리덕션 차원을 더 작은 청크로 반복하는 루프를 작성할 수 있습니다.두 전략 간 전환에는 보통 코드 재작성이 필요하지만, Helion은 reduction_loops 구성으로 이 선택을 자동화하여, 오토튜너가 영속적 버전과 루프 버전 모두를 벤치마킹하고 최적의 것을 선택할 수 있습니다. pid_type Helion은 그리드 크기 계산과 Program ID(PID)에서 데이터 타일로의 매핑을 자동화합니다. pid_type 구성을 통해 오토튜너가 수동 코드 변경 없이 다양한 매핑 전략을 탐색할 수 있습니다: 'flat': 단순한 1D 그리드를 사용합니다. 'xyz': PID를 다차원 그리드에 매핑합니다. 'persistent_blocked/persistent_interleaved': 스트리밍 멀티프로세서(SM)당 하나의 스레드 블록만 실행하는 영속적 커널을 생성하고, 커널 내부 루프를 사용하여 가상 PID를 반복하며 여러 타일의 작업을 처리합니다. 이 전략은 커널 실행 오버헤드 없이 컴퓨팅 리소스가 포화 상태를 유지하도록 하여 SM 활용률을 향상시킬 수 있습니다. load_eviction_policy Helion은 Triton tl.load의 eviction_policy 매개변수에 대해 오토튜닝을 수행하여 GPU L1 캐시 상주성(residency)에 영향을 미칩니다. 오토튜닝을 통해 커널의 메모리 접근 패턴에 가장 적합한 제거 힌트(eviction hint) 조합을 선택할 수 있습니다. Triton 구성:num_warps, num_stages, range_unroll_factors, range_warp_specializes, range_num_stages, range_multi_buffers, range_flattens Helion은 표준 Triton 튜닝 가능한 매개변수를 자동으로 탐색하여, 수동 튜닝에 필요한 개발자의 노력을 덜어줍니다.

성능 분석 및 벤치마크 / Performance Analysis and Benchmarks

다양한 커널과 shape에 대해 NVIDIA B200 및 AMD MI350X GPU에서 즉시 실행(eager mode) 대비 각각의 속도 향상을 측정하기 위해, Helion을 torch.compile(max-autotune 적용)과 직접 작성한 Triton 커널과 비교 벤치마킹했습니다. 대부분의 직접 작성한 Triton 커널은 Liger-Kernel 벤치마크 스위트에서 가져왔습니다.

We benchmark the performance of Helion to torch.compile(with max-autotune), and hand-written Triton to measure their respective speedups over eager mode execution across a wide variety of kernels and shapes on NVIDIA B200 and AMD MI350X GPUs. Most of the hand-written Triton kernels are from the Liger-Kernel benchmark suite.

NVIDIA B200에서의 성능 / Performance on NVIDIA B200

다음 표는 NVIDIA B200 GPU에서의 성능을 요약하며, 초록색으로 강조된 셀은 각 커널에 대해 즉시 실행 대비 가장 높은 속도 향상을 나타냅니다. 모든 벤치마크에서 Helion은 3.27배의 최고 기하 평균 속도 향상을 달성했으며, torch.compile(max-autotune 적용)은 2.7배, 직접 작성한 Triton 커널은 1.76배를 기록했습니다. 평균적으로 Helion은 torch.compile 대비 1.21배, Triton 커널 대비 1.85배의 속도 향상을 제공합니다. 특히 softmax 커널에서 Helion이 torch.compile 대비 2.28배의 속도 향상을 달성했고, jsd 커널에서는 직접 작성한 Triton 대비 6.22배의 성능 향상을 보였습니다.

The following table summarizes the performance on NVIDIA B200 GPU, with green highlighted cells indicating the highest speedup over eager mode execution for each kernel. Across all benchmarks, Helion achieves the highest geomean speedup of 3.27x, followed by torch.compile (with max-autotune) at 2.7x, and hand-written Triton kernels at 1.76x. On average, Helion delivers 1.21x speedup over torch.compile and 1.85x speedup over Triton kernels. Performance gains are especially notable for the softmax kernel, with Helion achieving 2.28x speedup over torch.compile, and for the jsd kernel, where Helion outperforms hand-written Triton by 6.22x.

AMD MI350X에서의 성능 / Performance on AMD MI350X

AMD MI350X에서의 성능도 유사한 추세를 보이며, Helion은 즉시 실행 대비 2.37배의 최고 기하 평균 속도 향상을 달성했고, torch.compile은 2.26배, Triton 커널은 1.65배를 기록했습니다. 평균적으로 Helion은 torch.compile 대비 1.05배, Triton 커널 대비 1.44배의 속도 향상을 제공합니다. 특히 int4_gemm과 jsd 커널에서 성능 향상이 두드러지며, Helion이 직접 작성한 Triton 커널 대비 각각 4.5배와 4.4배의 성능을 보였습니다.

Performance on AMD MI350X shows a similar trend, with Helion achieving the highest geomean speedup of 2.37x over eager mode execution, compared to 2.26x for torch.compile and 1.65x for Triton kernels. On average, Helion delivers 1.05x speedup over torch.compile and 1.44x speedup over Triton kernels. The performance gains are particularly pronounced for the int4_gemm and jsd kernels, with Helion outperforming hand-written Triton kernels by 4.5x and 4.4x, respectively.

사례 연구 1: 고도로 최적화된 CuTe DSL 커널을 능가하다 / Case Study 1: Outperforming Highly Optimized CuTe DSL Kernel

하루도 채 걸리지 않아 작성한 Helion의 RMSNorm 역전파 커널 구현은, CuTe DSL로 작성된 고도로 수동 최적화된 Quack 커널과 동등하거나 이를 능가하는 성능을 보여줍니다. H100 GPU에서 다양한 리덕션 차원에 걸쳐 Helion은 수동 튜닝된 커널과 일관되게 동등하거나 이를 능가하는 성능을 보여주며, 더 높은 수준의 추상화에서 개발하면서도 전문가 수준의 성능을 달성할 수 있음을 입증합니다. torch.compile(max-autotune 적용) 및 직접 작성한 Triton과도 비교했으며, 많은 경우에서 Helion이 우수한 성능을 보였습니다.

A Helion implementation of the RMSNorm backward kernel, written in less than a day, demonstrates performance on par with or exceeds a highly hand-optimized Quack kernel written in CuTe DSL. Across a range of reduction dimensions on H100 GPU, Helion consistently matches or outperforms the manually-tuned kernel, demonstrating its ability to match expert-level performance with significant productivity boost from developing at a higher level of abstraction. We also compare to torch.compile (with max-autotune) and hand-written Triton, where Helion outperformed in many cases.

이러한 결과는 대규모 오토튜닝 탐색을 효율적으로 지원하도록 설계된 Helion의 컴파일러 아키텍처에 의해 가능합니다.

These results are made possible by Helion’s underlying compiler architecture, which is designed to efficiently support large-scale search for autotuning.

사례 연구 2: TileLang과의 벤치마크 비교 / Case Study 2: Benchmarking Helion to TileLang

고성능 GPU 커널 개발을 위한 오픈 소스 DSL인 TileLang과도 Helion의 성능을 비교했습니다. Mamba-2 아키텍처의 핵심인 선택적 스캔(selective scan) 연산인 Mamba-2-chunk-scan 커널을 Helion으로 구현하여, TileLang 및 Triton 구현과 비교했습니다. H100에서 Helion은 최고 성능을 달성하여, 다양한 구성에서 TileLang 대비 2.12~2.63배, Triton 대비 1.2~1.85배의 속도 향상을 기록했습니다.

We also compare the performance of Helion to TileLang, an open-source DSL for developing high-performance GPU kernels. We implemented the Mamba-2-chunk-scan kernel — a selective scan operation central to the Mamba-2 architecture — in Helion to compare against its TileLang and Triton counterparts. On H100, Helion delivers the highest performance, achieving 2.12x–2.63x speedups over TileLang and 1.2x–1.85x over Triton across different configurations.

고수준 컴파일러 아키텍처 / High-Level Compiler Architecture

Helion의 컴파일러 아키텍처는 Python 함수를 TorchInductor를 통해 고도로 최적화된 Triton 코드로 점진적으로 저수준화(lower)하도록 설계되어 있습니다.

Helion’s compiler architecture is designed to progressively lower a Python function into highly optimized Triton code with TorchInductor.

컴파일 파이프라인은 다음의 핵심 단계를 거칩니다:

The compilation pipeline proceeds through the following key stages:

Python AST 파싱: 커널의 Python 소스 코드를 추상 구문 트리(AST, Abstract Syntax Tree)로 파싱하는 것으로 시작됩니다.
타입 전파 및 메타데이터: 커스텀 패스가 AST를 순회하면서 각 노드에 타입 정보와 기타 필수 메타데이터를 주석으로 달아 확장된 AST를 생성합니다.
디바이스 IR로 저수준화: 이 주석이 달린 트리는 Helion의 주요 중간 표현(IR, intermediate representation)으로 저수준화됩니다. 디바이스 IR은 정적 단일 할당(SSA, Static Single-Assignment) 형태의 FX 그래프 모음으로, 프로그램의 각 기본 블록을 나타내는 그래프로 구성됩니다. 이 그래프의 모든 노드는 코드 생성에 사용되는 Inductor IR 노드에 대한 포인터를 포함합니다.
컴파일러 패스: 디바이스 IR에 일련의 변환 패스가 적용됩니다. 이 패스들은 영속적 리덕션을 루프 기반으로 변환하는 리덕션 롤링(reduction rolling) 최적화와 같은 핵심 의미 변경(semantic changes)을 구현합니다.
구성 기반 코드 생성: 최종 단계에서 코드 생성기는 변환된 디바이스 IR과 오토튜닝된 구성이라는 두 가지 입력을 받아 최종 Triton 코드를 생성합니다.

Python AST Parsing: The process begins by parsing the kernel’s Python source code into an Abstract Syntax Tree (AST).

Type Propagation & Metadata: A custom pass traverses the AST, annotating each node with type information and other essential metadata to create an extended AST.

Lowering to Device IR: This annotated tree is lowered into Helion’s primary intermediate representation (IR). The Device IR is a collection of FX Graphs in Static Single-Assignment (SSA) form, with one graph representing each basic block of the program. Every node in this graph contains a pointer to an Inductor IR node, which is used for code generation.

Compiler Passes: A series of transformation passes are applied to the Device IR. These passes implement key semantic changes, such as the reduction rolling optimization, which converts a persistent reduction into a looped one.

Codegen with Configuration: In the final stage, the code generator takes two inputs: the transformed Device IR and the autotuned config. It uses these to generate the final output Triton code.

핵심적인 아키텍처 결정은 성능에 중요한 구성(config)이 코드 생성 시 파이프라인의 가장 마지막 단계에서만 적용된다는 것입니다. 이를 통해 파싱부터 IR 변환까지의 대부분의 컴파일 과정이 오토튜닝 탐색 전에 한 번만 실행되므로, 수천 개의 구성 탐색이 계산적으로 효율적이 됩니다.

A key architectural decision is that the performance-critical config is applied only at the very end of the pipeline during code generation. This allows the majority of the compilation process, from parsing to IR transformation, to be run just once before the autotuning search, making the exploration of thousands of configurations computationally efficient.

결론 / Conclusions

Helion은 머신러닝을 위한 현재의 커널 작성 영역에서 중요한 격차를 해소합니다. 친숙하고 고수준인 PyTorch 유사 문법과 강력한 사전 컴파일 오토튜닝 엔진을 결합하여, 개발자 생산성, 세밀한 제어, 성능 이식성의 독보적인 균형을 제공합니다. Helion은 깊은 하드웨어 전문 지식 없이도 최첨단 성능을 달성하는 이식 가능하고 미래 지향적인 커널을 작성할 수 있게 해주며, 고성능 머신러닝 커널을 위한 새롭고 더 생산적인 패러다임을 확립합니다.

Helion addresses a critical gap in today’s kernel authoring space for machine learning. By combining a familiar and high-level PyTorch-like syntax with a powerful, ahead-of-time autotuning engine, it provides a unique balance of developer productivity, fine-grained control, and performance portability. It empowers developers to write portable, future-proof kernels that achieve state-of-the-art performance without requiring deep hardware expertise, establishing a new and a more productive paradigm for performant machine learning kernels.

Helion은 2025년 10월 22일에 베타로 출시되었으며, 커뮤니티의 피드백, 버그 리포트 및 기여를 환영합니다. 자세한 내용은 아래를 참고하세요:

Helion is being released in Beta on Oct. 22nd, 2025, and we welcome feedback, bug reports, and contributions from the community. For more information, see:

감사의 글 / Acknowledgements

Helion은 Jason Ansel, Oguz Ulgen, Will Feng, Jongsok Choi, Markus Hoehnerbach, Manman Ren, Jie Liu, Paul Zhang, Driss Guessous, Joy Dong, Xuan Zhang, Karthick Panner Selvam, Peng Wu, Hongtao Yu, Neil Dhar, Nick Riasanovsky, Shane Nay, Alexey Loginov를 비롯하여 Meta, NVIDIA, AMD, Intel 팀의 많은 분들의 노력으로 만들어졌습니다.

Helion is the work of many hands including: Jason Ansel, Oguz Ulgen, Will Feng, Jongsok Choi, Markus Hoehnerbach, Manman Ren, Jie Liu, Paul Zhang, Driss Guessous, Joy Dong, Xuan Zhang, Karthick Panner Selvam, Peng Wu, Hongtao Yu, Neil Dhar, Nick Riasanovsky, Shane Nay, Alexey Loginov, as well as teams at Meta, NVIDIA, AMD, and Intel

참고: B200 성능 수치는 10/23에 업데이트되었습니다.

Note: B200 Performance numbers updated 10/23.

덧글 작성 및 공유를 위해 https://pytorch.kr/blog/2025/helion/에 게시된 글이 커뮤니티에도 함께 게시되었습니다.