필사 모드: Triton Gluon — 컴파일러가 숨기던 레이아웃을 손으로 쓰는 언어

• 들어가며 — Triton이 막히는 지점은 "느려서"가 아니다
• Gluon이란 무엇인가 — 같은 컴파일러 스택, 뒤집힌 계약
• 첫 커널은 실망스럽다 — 666 GB/s
• 레이아웃 — Gluon이 당신에게 넘기는 것
• 숫자로 보는 레이아웃
• 레이아웃이 틀리면 성능이 무너진다
• PyTorch를 이기는 지점, 그리고 그 대가
• 릴리스 타임라인 — 어디까지 왔나
• NVIDIA 전용이 아니다 — AMD gfx950과 MI355
• 레이아웃 밑바닥 — 선형 레이아웃
• 정직한 트레이드오프 — 언제 쓰지 말아야 하나
• 마치며
• 참고 자료

들어가며 — Triton이 막히는 지점은 "느려서"가 아니다

Triton의 매력은 분명합니다. Python으로 타일 단위 커널을 쓰면 컴파일러가 레이아웃, 메모리 할당, 데이터 이동, 비동기화를 알아서 처리해 줍니다. 대부분의 경우 이건 훌륭한 거래입니다 — CUDA C++로 내려가지 않고도 쓸 만한 성능이 나오니까요.

문제는 그 거래가 깨지는 순간입니다. Gluon 튜토리얼의 도입부가 이 상황을 정확히 서술합니다. Triton 컴파일러는 넓은 범위의 커널에 대해 효율적인 코드를 잘 생성하지만, 손으로 튜닝한 저수준 코드에 질 수 있습니다. 그리고 그런 일이 벌어졌을 때 — 튜토리얼의 표현으로는 "모든 디테일이 감춰져 있기 때문에 사용자가 성능을 유의미하게 개선하기 위해 할 수 있는 일이 거의 없습니다".

이게 핵심입니다. Triton의 한계는 "느리다"가 아니라 막혔을 때 빠져나갈 문이 없다는 것입니다. 추상화가 잘 동작할 때는 축복이고, 어긋나는 순간에는 벽입니다. 지금까지 그 벽 앞에서의 선택지는 하나였습니다 — CUDA C++로 처음부터 다시 쓰기.

Gluon은 세 번째 선택지를 만들려는 시도입니다.

Gluon이란 무엇인가 — 같은 컴파일러 스택, 뒤집힌 계약

공식 문서의 정의는 짧습니다. Gluon은 Triton의 하위 레벨 GPU 프로그래밍 모델이고, "레이아웃, 공유 메모리, 워프 특수화, 타깃별 기능을 직접 노출해서 고급 커널이 편의성을 제어권과 맞바꿀 수 있게" 합니다.

중요한 건 Gluon이 별개의 프로젝트가 아니라는 점입니다. 튜토리얼이 밝히듯 Gluon과 Triton은 같은 컴파일러 스택을 씁니다. 둘 다 타일 기반 SPMD 프로그래밍 모델을 구현하고, 둘 다 같은 프론트엔드와 JIT 인프라를 공유하는 Python DSL입니다. 호스트 쪽 코드도 거의 같습니다 — 커널을 부르는 방식, 그리드를 지정하는 방식, 심지어 triton.autotune까지 그대로 씁니다.

달라지는 건 디바이스 코드를 쓰는 방식입니다. 튜토리얼의 정리를 그대로 옮기면, Gluon은 디바이스 코드를 어떻게 쓰는지를 바꾸고, 호스트 코드는 커널에 하이퍼파라미터가 늘어나는 정도로만 바뀝니다.

가장 단순한 Gluon 커널은 Triton과 구별이 안 갑니다.

import torch
import triton
from triton.experimental import gluon
from triton.experimental.gluon import language as gl


@gluon.jit
def copy_scalar_kernel(in_ptr, out_ptr):
    value = gl.load(in_ptr)
    gl.store(out_ptr, value)

@triton.jit 대신 @gluon.jit, tl 대신 gl. 여기까지는 이름만 바뀐 것처럼 보입니다.

그리고 저 import 경로를 눈여겨보십시오 — triton.experimental입니다. 이 글의 마지막에 다시 돌아올 이야기입니다.

첫 커널은 실망스럽다 — 666 GB/s

튜토리얼은 정직하게도 실패에서 시작합니다. 스칼라를 하나씩 복사하는 memcpy 커널을 쓰고, XBLOCK을 오토튠하고, GB200에서 8 GB를 복사해 봅니다. 튜토리얼 소스에 주석으로 박혀 있는 결과는 이렇습니다.

Time:        24.00 ms
Throughput: 666.24 GB/s

오토튜너가 고른 최적 XBLOCK은 2048이었고, 결과는 약 666 GB/s. 튜토리얼은 이 숫자에 대해 "GPU의 8 TB/s 피크 대역폭에 한참 못 미친다"고 적습니다. 대략 피크의 8% 수준입니다.

이유는 커널을 보면 바로 보입니다.

@gluon.jit
def memcpy_kernel(in_ptr, out_ptr, xnumel, XBLOCK: gl.constexpr):
    pid = gl.program_id(0)
    start = pid * XBLOCK
    end = min(start + XBLOCK, xnumel)
    for i in range(start, end):
        value = gl.load(in_ptr + i)
        gl.store(out_ptr + i, value)

각 프로그램(GPU의 CTA, 즉 스레드 블록)이 한 번에 원소 하나씩 복사합니다. GPU의 병렬성을 전혀 쓰지 못하는 코드입니다. 여러 원소를 한꺼번에 옮기려면 스칼라가 아니라 타일을 로드해야 하고, 타일을 로드하려면 — 레이아웃을 골라야 합니다.

이 지점이 Gluon이 Triton과 갈라지는 곳입니다. Triton이었다면 컴파일러가 알아서 골랐을 것을, Gluon에서는 당신이 씁니다.

레이아웃 — Gluon이 당신에게 넘기는 것

레이아웃은 텐서의 원소들이 스레드 블록 안의 스레드들에게 어떻게 분배되는지를 규정합니다. 분배는 GPU의 계층 구조를 따라갑니다 — 스레드 블록, 워프, 레인, 그리고 각 레인의 개별 레지스터.

가장 흔한 종류가 BlockedLayout입니다.

gl.BlockedLayout(
    size_per_thread=[2, 4],
    threads_per_warp=[16, 2],
    warps_per_cta=[2, 2],
    order=[1, 0],
)

읽는 법은 이렇습니다.

• size_per_thread=[2, 4] — 각 스레드가 연속된 2x4 서브타일을 소유하고, 그건 그 스레드의 레지스터에 담깁니다.
• threads_per_warp=[16, 2] — 스레드 타일들을 이렇게 쌓아 워프 타일을 만듭니다. 곱이 워프당 스레드 수와 맞아야 합니다(NVIDIA 하드웨어에서 32).
• warps_per_cta=[2, 2] — 워프 타일들을 이렇게 쌓아 블록 전체를 만듭니다.
• order=[1, 0] — 어느 차원부터 타일링할지. 행 우선(row-major)입니다.

셋을 원소별로 곱하면 블록 모양이 나옵니다 — 여기서는 [64, 16].

order가 왜 중요한지는 스레드 하나의 타일을 그려 보면 즉시 드러납니다. order=[1, 0]이면 레지스터가 안쪽 차원을 따라 증가합니다.

[[T:0, T:1, T:2, T:3],
 [T:4, T:5, T:6, T:7]]

order=[0, 1]이면 같은 2x4 타일이 이렇게 바뀝니다.

[[T:0, T:2, T:4, T:6],
 [T:1, T:3, T:5, T:7]]

같은 원소, 같은 스레드, 다른 배치. 그리고 뒤에서 보겠지만 이 차이 하나가 대역폭을 몇 배로 갈라놓습니다.

레이아웃을 명시하는 일이 부수적으로 주는 것도 있습니다 — 레지스터 예산이 계산 가능해집니다. 튜토리얼의 예를 따라가면, 위 레이아웃(블록 [64, 16])으로 128x128 f32 텐서를 다루면 블록을 [2, 8]로 타일링하게 되고, 스레드당 블록마다 8개, 텐서 전체로는 8 * 16 = 128개의 레지스터를 씁니다. Triton에서는 컴파일러 뒤에 숨어 있던 숫자입니다.

주의할 함정도 같이 나옵니다. 텐서가 블록보다 작으면 브로드캐스트됩니다. 32x8 f32 텐서는 원소가 256개뿐이지만, 블록 타일링이 변하지 않으므로 각 프로그램에서 64 * 16 = 1024개의 물리 레지스터로 저장됩니다. 작은 텐서가 레지스터를 4배로 먹는 셈입니다 — 명시적 레이아웃이 없으면 알아채기 어려운 종류의 비용입니다.

숫자로 보는 레이아웃

여기서부터는 Triton 저장소의 튜토리얼 소스(python/tutorials/gluon/02-layouts.py)에 주석으로 기록된 측정값입니다. 상류 개발자들이 GB200에서 직접 재서 소스에 박아 둔 값이고, 제가 재현한 것은 아닙니다.

1D memcpy를 제대로 된 타일 로드로 다시 쓰고, XBLOCK=2048, num_warps=4로 고정한 뒤 size_per_thread(아래의 R)만 바꿔 가며 잽니다.

gl.BlockedLayout(
    size_per_thread=[R],
    threads_per_warp=[32],
    warps_per_cta=[4],
    order=[0],
)

기록된 결과입니다.

R=1   6.574 TB/s
R=2   6.476 TB/s
R=4   6.474 TB/s
R=8   6.502 TB/s
R=16  6.214 TB/s

두 가지를 읽을 수 있습니다. 첫째, 스칼라 루프의 0.666 TB/s에서 6.574 TB/s로 왔습니다 — 기록된 두 값을 그냥 나누면 약 10배입니다(제 산술이고, 상류가 그렇게 표현한 건 아닙니다). 타일 로드와 레이아웃이 하는 일의 크기가 이 정도입니다.

둘째, R을 키운다고 좋아지지 않습니다. R=1이 가장 빠르고 R=16이 가장 느립니다. 튜토리얼은 여기서 SASS를 열어 LDG.E/STG.E 명령을 직접 세어 보는 쪽으로 갑니다. 이게 Gluon의 성격을 잘 보여 줍니다 — 레이아웃을 노출한다는 건 "더 좋은 기본값"을 준다는 뜻이 아니라, 선택지와 함께 그 선택을 검증할 책임을 준다는 뜻입니다. 폭을 키우면 빨라진다는 직관은 여기서 틀립니다.

레이아웃이 틀리면 성능이 무너진다

2D로 가면 이야기가 더 날카로워집니다.

8 GB짜리 32K x 64K 텐서를 XBLOCK=1, YBLOCK=2048로, 1D에서 이겼던 R=1과 같게 동작하는 레이아웃으로 복사합니다.

layout = gl.BlockedLayout([1, 1], [1, 32], [1, 4], [1, 0])

결과는 6.260 TB/s — 1D보다 5% 느립니다. 튜토리얼은 2D 산술이 더 복잡한 것 등을 이유로 들며, 이 정도는 납득할 만하다고 넘어갑니다.

그런데 입력 텐서를 전치하면 어떻게 될까요. 같은 커널, 같은 레이아웃입니다.

0.774 TB/s

튜토리얼의 표현으로는 성능이 "폭락(craters)"합니다. 이유는 단순합니다 — 안쪽 차원이 더 이상 연속적이지 않으니 병합(coalescing)이 사라집니다. 32개 레인이 인접한 주소를 읽어야 한 번의 트랜잭션으로 묶이는데, 전치된 텐서에서는 레인마다 멀리 떨어진 주소를 짚습니다.

고치는 법도 단순합니다. 블록 크기를 바꾸고 레이아웃을 전치합니다.

layout = gl.BlockedLayout([1, 1], [32, 1], [4, 1], [0, 1])
# XBLOCK=2048, YBLOCK=1

6.590 TB/s

0.774에서 6.590으로. 커널의 알고리즘은 한 글자도 안 바뀌었고, 바뀐 건 레이아웃 서술자뿐입니다.

여기서 튜토리얼이 덧붙이는 관찰이 재밌습니다. 고친 버전(6.590 TB/s)은 1D memcpy(6.574 TB/s)보다 살짝 빠릅니다. 전치/비전치 양쪽에서 각 프로그램이 메모리에 접근하는 방식은 동일한데도 차이가 나는데, 튜토리얼은 이걸 프로그램이 GPU 어디에 스케줄되는지, 즉 데이터 지역성 탓으로 설명합니다 — TLB의 가상 주소 변환 캐싱, H100에서 L2 캐시가 서로 통신하는 파티션으로 나뉘어 있다는 점 같은 것들. 이 층위에서는 "같은 접근 패턴"조차 같은 성능을 뜻하지 않습니다.

여기서 얻을 실무 교훈: 최적 레이아웃은 커널이 아니라 전역 메모리에 놓인 텐서의 스트라이드에 달려 있습니다. 그래서 튜토리얼은 결국 이런 헬퍼로 수렴합니다.

def get_layout_for_gmem_access(tensor, num_warps):
    if len(tensor.shape) == 1:
        return gl.BlockedLayout([1], [32], [num_warps], [0])

    assert len(tensor.shape) == 2, "only 1D and 2D tensors are supported"
    assert 1 in tensor.stride(), "expected at least 1 contiguous dimension"
    if tensor.stride(1) == 1:
        return gl.BlockedLayout([1, 1], [1, 32], [1, num_warps], [1, 0])
    else:
        return gl.BlockedLayout([1, 1], [32, 1], [num_warps, 1], [0, 1])

레이아웃을 상수로 박는 게 아니라 런타임 스트라이드에서 유도합니다. Gluon 커널이 유연하면서도 빠르려면 이런 식의 코드가 필요해집니다.

PyTorch를 이기는 지점, 그리고 그 대가

Gluon 2D memcpy가 값을 하는 경우도 튜토리얼이 짚습니다. 8 GB 텐서에서 한 행씩 건너뛴 뷰(비연속 텐서)를 연속 텐서로 복사하는, PyTorch의 x.contiguous()와 같은 작업입니다.

2D memcpy:                6.258 TB/s
torch.Tensor.contiguous:  2.946 TB/s
2D memcpy (transposed):   6.398 TB/s

튜토리얼의 표현은 "이미 PyTorch 구현보다 2배 이상 빠르다"입니다. 이건 Gluon이 실제로 값을 하는 종류의 사례입니다 — 입력이나 출력이 이국적인 레이아웃을 가질 때. 다만 1D memcpy는 입력과 출력이 둘 다 연속 메모리 블록의 뷰일 때만 쓸 수 있다는 조건이 붙습니다.

그럼 입력과 출력의 레이아웃이 서로 반대면? 튜토리얼은 이 경우도 재 봅니다. 입력은 dim 1을 따라, 출력은 dim 0을 따라 연속인 32K x 32K 텐서입니다.

2D memcpy (order=[1, 0]): 0.978 TB/s
2D memcpy (order=[0, 1]): 1.674 TB/s

어느 레이아웃을 골라도 끔찍합니다. 하나를 만족시키면 다른 하나가 깨지니까요. 해법은 로드와 스토어에 서로 다른 레이아웃을 쓰고 중간에 변환하는 것입니다.

@gluon.jit
def memcpy_2d_inout_kernel(in_ptr, out_ptr, ..., layout_in: gl.constexpr, layout_out: gl.constexpr, ...):
    ...
    value = gl.load(in_ptr + in_offsets, mask=mask_in)

    # Use `gl.convert_layout` to perform layout conversions.
    value = gl.convert_layout(value, layout_out)

    gl.store(out_ptr + out_offsets, value, mask=mask_out)

2D memcpy (in/out layouts): 4.814 TB/s

0.978에서 4.814로 올랐지만 — 6.5 TB/s대에는 못 미칩니다. 그리고 튜토리얼은 그 이유를 감추지 않습니다. convert_layout은 공유 메모리를 써야 하고, 공유 메모리는 GPU에서 귀한 자원입니다. 레이아웃 변환에 공유 메모리를 쓰면 점유율(occupancy)과 최대 파이프라인 깊이가 줄어 성능에 악영향을 줍니다. 튜토리얼의 결론은 "이 경우엔 변환 비용이 불가피하고, 비효율적인 전역 메모리 접근 비용보다는 훨씬 싸다"는 것입니다.

즉 4.814 TB/s는 승리가 아니라 가장 덜 나쁜 선택입니다. 이런 식으로 비용을 눈에 보이게 만드는 게 Gluon의 성격입니다.

릴리스 타임라인 — 어디까지 왔나

Gluon이 언제 어디까지 왔는지는 릴리스 노트로 추적할 수 있습니다. 아래 날짜는 GitHub 릴리스 API에서 확인한 값입니다.

읽히는 흐름이 있습니다. 2025년 중반에 뼈대가 들어왔고, 2026년 들어 멀티 CTA와 최신 하드웨어 기능(Blackwell의 tcgen05, 클러스터 배리어)으로 확장되는 중입니다. 그리고 v3.7.1이 순수 버그 수정이었다는 사실은 — 최신 안정 릴리스 기준으로 Gluon 표면이 한 달쯤 잠잠했다는 뜻이기도 합니다.

3.6.0의 Triton → Gluon 변환기는 따로 언급할 만합니다. 마이그레이션을 쉽게 해 주는 도구처럼 들리지만, 해당 PR(#8417, 2025-10-11 머지)의 설명문은 기대를 정확히 눌러 놓습니다. "이건 프로덕션용이 아니며, Triton 커널을 순진한(naive) Gluon 버전으로 변환할 수 있게 해 줄 뿐이다. 물론 그 Gluon 버전은 상당히 더 느릴 것이다."

정직한 문장입니다. 그리고 Gluon의 본질을 요약합니다 — Gluon으로 옮기는 것 자체는 아무 성능도 주지 않습니다. 성능은 옮긴 다음에 당신이 레이아웃을 제대로 고를 때 나옵니다.

NVIDIA 전용이 아니다 — AMD gfx950과 MI355

위 표에서 눈에 띄는 건 AMD 항목이 계속 나온다는 점입니다. gfx1250용 async_copy, AMD 워프 파이프라인, gfx1250용 Stream-K 커널, AMD 클러스터 배리어. Gluon은 NVIDIA 전용 탈출구로 설계되지 않았습니다.

AMD도 이걸 밀고 있습니다. 2026년 5월 22일 ROCm 블로그에 Lixun Zhang, Jason Furmanek, Peng Sun, Emad Barsoum이 쓴 "From Naive to Near-Peak: Building High-Performance GEMM Kernels with Gluon"이 올라왔습니다. gfx950(CDNA4) 기반 MI350/MI355에서 Gluon으로 GEMM 커널을 v0부터 v9까지 단계적으로 최적화하는 튜토리얼입니다.

아래 숫자는 전부 AMD의 벤더 자체 측정입니다. 블로그가 밝힌 측정 조건은 MI355 한 장, ROCm 7.0, gfx950-tutorial-v0.1 태그로 빌드한 Triton입니다.

여기서 오해하기 쉬운 지점을 분명히 해 두겠습니다. 520 → 1489 TFLOPS의 약 3배는 Gluon이 Triton을 3배 이긴다는 뜻이 아닙니다. v0도 v9도 전부 Gluon 커널입니다 — 같은 저장소(ROCm/gfx950-gluon-tutorials)의 v0_naive부터 v9_beyond_hotloop까지, 교육용으로 설계된 최적화 여정의 양 끝일 뿐입니다. 순진한 Gluon 커널이 느린 건 당연하고(앞의 변환기 PR이 말한 그대로입니다), 3배는 "레이아웃과 스케줄링을 제대로 하면 이만큼 온다"는 이야기이지 언어 간 비교가 아닙니다.

그리고 블로그 스스로 선을 긋습니다 — 이건 hipBLASLt 같은 프로덕션 라이브러리의 대체재가 아니라고 명시합니다. 제가 확인한 본문에는 hipBLASLt와의 직접 비교 수치가 없습니다. 즉 "Gluon 커널이 벤더 라이브러리를 이겼다"는 주장은 이 자료에서 나오지 않습니다.

방법론 측면에서 이 블로그가 정직한 대목도 있습니다. 주된 지표로 TFLOPS가 아니라 MFMA 효율을 쓰는데, 그 이유로 MFMA 효율이 클럭과 무관하고 실행 간 재현성이 raw TFLOPS보다 낫다는 점을 듭니다. 부스트 클럭에 따라 출렁이는 TFLOPS만 들이대는 벤더 자료보다 한 수 위의 태도입니다. (물론 MFMA 효율 99%가 곧 워크로드가 빠르다는 뜻은 아닙니다 — 핫 루프 안에서 행렬 엔진이 노는 시간이 거의 없다는 뜻이고, 그건 커널이 옳은 알고리즘을 쓰고 있을 때만 의미가 있습니다.)

레이아웃 밑바닥 — 선형 레이아웃

한 겹 더 들어가면 선형 레이아웃(linear layout)이 있습니다. 튜토리얼이 설명하듯 Gluon에는 정규(canonical) 레이아웃 표현이 없습니다 — 여러 레이아웃이 같은 원소 매핑을 나타낼 수 있습니다. 예를 들어 아래 둘은 동등합니다.

gl.BlockedLayout([1], [32], [4], [0])
gl.SliceLayout(1, gl.BlockedLayout([1, 1], [32, 1], [4, 1], [1, 0]))

동등하다는 걸 당신이 알고 있다면 gl.convert_layout(x, layout, assert_trivial=True)로 변환이 실제로 공짜인지(레지스터 재정렬 이상은 아닌지) 컴파일러에 검증시킬 수 있습니다. 명시적 언어다운 API입니다.

그리고 모든 Gluon 레이아웃은 선형 레이아웃으로 표현 가능합니다. 튜토리얼은 이걸 "가장 표현력 있고 강력한 표현"이라고 부르면서, 동시에 "상대적으로 드물고 이해하기 어려울 수 있다"고 덧붙입니다. 기반 이론은 논문으로 나와 있습니다 — Keren Zhou, Mario Lezcano, Jeff Niu, Phil Tillet, Thomas Raoux 등의 "Linear Layouts: Robust Code Generation of Efficient Tensor Computation Using F2"(arXiv:2505.23819, 2025년 5월 28일 제출, 2026년 3월 6일 개정). 레이아웃을 GF(2) 위의 이진 행렬로 모델링해서 기존 시스템의 이차 복잡도를 피한다는 접근입니다.

실무자 입장에서의 요점은 이겁니다 — 대부분의 경우 BlockedLayout과 SliceLayout으로 충분하고, 선형 레이아웃은 필요할 때까지 몰라도 됩니다. 다만 zero-cost split/join/reshape/permute가 필요해지는 순간이 오면, 그게 있는 곳이 여기입니다.

정직한 트레이드오프 — 언제 쓰지 말아야 하나

이제 정직해질 부분입니다.

아직 experimental입니다. import 경로가 triton.experimental.gluon이고, 저장소에서도 python/triton/experimental/gluon에 있습니다. 공식 Gluon 개요 문서에는 실험적 상태나 안정성에 대한 경고 문구가 따로 없지만, 경로 자체가 API 안정성을 약속하지 않는다고 말하고 있습니다. 3.4.0부터 3.7.0까지 릴리스마다 API가 계속 추가·변경돼 온 이력도 이 해석과 일치합니다. 프로덕션 커널을 여기에 올린다면 Triton 버전을 올릴 때마다 깨질 각오를 해야 합니다.

요구하는 지식의 양이 다릅니다. 튜토리얼 도입부가 직설적입니다 — Gluon 커널을 쓰려면 GPU 하드웨어와 GPU 프로그래밍의 여러 측면에 대한 더 깊은 이해가 필요합니다. 워프당 스레드가 32라는 것, 병합이 왜 안쪽 차원에 달렸는지, 공유 메모리 뱅크 충돌이 무엇인지, 레지스터 압력이 점유율에 어떻게 작용하는지를 모르면 BlockedLayout의 네 인자를 채울 수 없습니다. Triton이 이걸 감춰 준 건 게으름이 아니라 서비스였습니다.

Gluon으로 옮기는 것 자체는 공짜 성능이 아닙니다. 변환기 PR이 말한 그대로 — 순진하게 옮긴 Gluon은 상당히 더 느립니다. 그리고 앞에서 본 R 스윕이 보여 주듯, 직관적인 선택(폭을 키우기)이 오히려 나쁠 수 있습니다. 자동화된 컴파일러가 하던 일을 사람이 하기로 했다면, 사람이 그 일을 더 잘한다는 증거를 매번 벤치마크로 만들어야 합니다.

튜토리얼 전체가 memcpy입니다. 이 글에서 인용한 GB200 숫자들은 상류 개발자들이 측정해 소스에 기록한 값이지, 제가 재현한 게 아닙니다. 그리고 그 대상은 메모리 대역폭에 갇힌 memcpy입니다 — 당신의 커널이 연산에 갇혀 있다면 레이아웃 스윕에서 10배가 나오지 않습니다.

정리하면 판단 기준은 이렇습니다.

써서 값을 하는 경우

• Triton으로 이미 짰고, 프로파일링으로 컴파일러가 뽑은 코드가 하드웨어 한계에 한참 못 미친다는 걸 측정으로 확인했다.
• 병목이 레이아웃·데이터 이동·비동기화에 있고, Triton 수준에서는 손댈 손잡이가 없다.
• 최신 하드웨어 기능(Blackwell tcgen05, Hopper WGMMA, 2CTA, TMA, CDNA4의 MXFP4)을 직접 찔러야 한다.
• 커널이 오래 살고 자주 돌아서, 튜닝 시간과 버전 업 때마다의 유지 비용을 상환한다.

과잉인 경우

• 아직 Triton으로 프로파일링해 보지 않았다. (이게 대부분입니다.)
• torch.compile이 뽑아 주는 코드로 충분하다.
• 커널이 연산 바운드이고 레이아웃이 병목이 아니다.
• 벤더 라이브러리(cuBLAS, hipBLASLt, cuDNN)가 이미 그 연산을 커버한다 — AMD 블로그조차 자기 튜토리얼이 hipBLASLt의 대체재가 아니라고 못 박습니다.
• 팀에 GPU 마이크로아키텍처를 아는 사람이 없다.

마치며

Gluon은 "더 좋은 Triton"이 아닙니다. 다른 계약입니다.

Triton의 계약은 "디테일은 내가 처리할 테니 너는 알고리즘만 써라"입니다. 대부분의 경우 좋은 거래고, 앞으로도 대부분의 커널은 Triton에 머물러야 합니다. Gluon의 계약은 "디테일을 네가 처리해라, 대신 막혔을 때 벽 대신 문이 있다"입니다.

이 스펙트럼이 중요합니다. 예전에는 Triton과 CUDA C++ 사이에 절벽이 있었습니다 — 컴파일러가 못 뽑아 주면 언어와 생태계를 통째로 갈아타야 했습니다. Gluon은 그 절벽에 계단을 놓습니다. 같은 컴파일러 스택, 같은 프론트엔드, 같은 JIT, 같은 오토튜너에 머문 채로 필요한 부분만 아래로 내려갑니다. 그리고 그 계단이 NVIDIA와 AMD 양쪽에 놓이고 있다는 점 — gfx950 튜토리얼을 AMD가 직접 쓰고 릴리스 노트마다 AMD 항목이 붙는다는 점 — 이 어쩌면 더 중요한 소식입니다.

다만 순서를 지키십시오. Triton으로 짜고, 프로파일링하고, 컴파일러가 실제로 졌다는 걸 확인하고, 그다음에 Gluon을 꺼내십시오. 0.774 TB/s와 6.590 TB/s를 가른 게 알고리즘이 아니라 레이아웃 서술자 한 줄이었다는 사실은 양날입니다 — 제대로 고르면 10배를 얻지만, 잘못 고르면 10배를 잃습니다. Triton은 당신이 그 선택을 하지 않아도 되게 해 줬습니다. Gluon은 그 선택을 당신에게 돌려줍니다. 그게 선물인지 청구서인지는, 당신이 프로파일러를 열어 봤는지에 달려 있습니다.

참고 자료

• Gluon Overview — Triton 공식 문서
• Introduction to Gluon — 튜토리얼 1편
• 01-intro.py — 튜토리얼 원본 소스(GB200 memcpy 측정치 포함)
• 02-layouts.py — 레이아웃 튜토리얼 원본 소스(TB/s 측정치 포함)
• triton-lang/triton 릴리스 노트 — 3.4.0부터 3.7.1까지
• PR #8417 — 실험적 Triton to Gluon 변환기
• From Naive to Near-Peak: Building High-Performance GEMM Kernels with Gluon — ROCm 블로그(AMD 벤더 자체 측정, 2026-05-22)
• ROCm/gfx950-gluon-tutorials — AMD 블로그의 커널 소스
• Linear Layouts: Robust Code Generation of Efficient Tensor Computation Using F2 (arXiv:2505.23819)