Gluon and Linear Layouts

Triton 의 핵심 매력 — “tile 단위 op 만 쓰면 layout / scheduling / register 할당이 자동”. 이 자동화가 L001 시점부터 GPU 커널을 짜는 자세를 바꿨다. 그러나 — Hopper / Blackwell 의 새 instruction (TMA, WGMMA, multicast) 이 들어오면서 — “heuristic 이 잡지 못하는 자리” 가 생겼다. 그 자리에서 사용자가 layout 을 직접 control 할 수 있는 — “Triton 과 같은 syntax, 그러나 layout 명시” — 가 Gluon 의 자리.

강의가 던지는 두 개의 질문.

1. Triton 의 자동 layout 결정이 부족한 자리는 정확히 어디인가 — 그리고 거기서 사용자가 무엇을 더 control 해야 하는가.
2. Linear Layouts 라는 새 추상이 — Triton 의 기존 layout 들 (blocked, shared, NVidia MMA) 을 어떻게 통합하는가 — 그리고 그것이 layout conversion 의 quadratic 복잡도를 어떻게 푸는가.

Triton 의 사용자 모델은 — “tile 의 shape 만 정하면 layout 은 컴파일러가 결정”. 이 모델이 못 잡는 자리들이 있다. 모두 새 hardware 기능과 연결.

• warp specialization — Hopper 의 producer-consumer pattern 은 warp 별로 다른 일을 한다. 한 warp 가 TMA load, 다른 warp 가 WGMMA. Triton 의 “모든 thread 가 같은 일” 모델로는 자연스러운 표현이 없음.
• TMA descriptor 의 fine control — multicast, swizzle mode, prefetch 의 상세한 결정. 자동 lowering 이 best 를 항상 못 찾음.
• shared memory 의 정확한 layout — bank conflict 회피용 swizzle 의 mode (3, 0, 3) 또는 (5, 0, 5) 의 결정. 자동 결정이 보수적이거나 잘못된 선택.
• register tile 의 partition — WGMMA 의 native shape (m64n*k16) 와 사용자 tile 의 align. 사용자가 tile 을 의도한 모양으로 못 강제하는 자리.

이 자리들은 — heuristic 이 영원히 못 잡는다 는 게 아니라 — 지금 시점에서 jaw bone 의 자동 결정이 hand-tuned 만큼 좋지 않다. 시간이 가면 일부는 자동화로 다시 옮겨갈 가능성. 단 hardware 의 진화 속도가 너무 빨라서 — “항상 자동화의 한 발 뒤” 라는 격차가 구조적으로 존재.

Gluon 의 위치를 한 줄로 — “Triton 의 layer L0 위에 명시적 layout argument 를 붙인 변형”. tl.load 의 같은 operation 이 — Triton 에서는 자동 layout, Gluon 에서는 사용자가 layout 명시.

paper 의 main thesis — “tensor layout 은 𝔽₂ (binary) 위의 행렬로 표현 가능하고, 그렇게 하면 모든 layout 변환이 행렬곱으로 통일된다”. 기존의 blocked / shared / MMA / dot operand 같은 layout 들이 특수 케이스의 행렬이 된다.

강의의 가장 구체적인 자리. Hopper 의 producer-consumer pattern 을 Triton 으로 짜는 것 vs Gluon 으로 짜는 것의 직접 대조. 의도적으로 단순화한 의사 코드 — 정확한 API 는 Triton repo 의 experimental/gluon 참조.

# 단순한 GEMM tile
@triton.jit
def gemm_tt(A, B, C, M, N, K,
            BM: tl.constexpr,
            BN: tl.constexpr,
            BK: tl.constexpr):
    pid_m = tl.program_id(0)
    pid_n = tl.program_id(1)
    a = tl.load(A_ptrs)        # layout 자동
    b = tl.load(B_ptrs)        # layout 자동
    acc = tl.zeros((BM, BN))
    for k in range(0, K, BK):
        a = tl.load(...)
        b = tl.load(...)
        acc += tl.dot(a, b)    # MMA layout 자동
    tl.store(C_ptrs, acc)

# 명시적 layout · pseudo
@gluon.jit
def gemm_gl(A, B, C, M, N, K, BM, BN, BK):
    # shared memory 의 swizzle 명시
    smem_a = gl.alloc_shared(
        (BM, BK), dtype=fp16,
        layout=SwizzledShared(3,0,3))
    smem_b = gl.alloc_shared(
        (BK, BN), dtype=fp16,
        layout=SwizzledShared(3,0,3))
    # producer warp: TMA load
    if warp_id == 0:
        gl.tma_load(A_desc, smem_a, ...)
        gl.tma_load(B_desc, smem_b, ...)
        gl.barrier_arrive(bar)
    # consumer warps: WGMMA
    else:
        gl.barrier_wait(bar)
        acc = gl.wgmma(smem_a, smem_b,
            layout=MMALayout(m64n128k16))
        gl.store(C_ptrs, acc)

실제 API 는 — Triton repo 의 python/triton/experimental/gluon 디렉토리에서 확인. experimental 이라는 단어가 보여주듯 — 2025 시점에서는 active development. API 의 detail 은 빠르게 변할 수 있다.

새 추상이 들어올 때 가장 흔한 문제 — “기존 코드를 다 다시 짜야 하는가”. Gluon 의 답은 아니다. Triton 코드는 그대로 도는 동시에, 일부 함수만 Gluon 으로 바꿔도 된다.

• 같은 .py 파일 — Triton 의 @triton.jit 와 Gluon 의 @gluon.jit 가 같은 파일에 공존.
• 같은 launcher — kernel 호출의 syntax 가 거의 동일. tile shape, grid, num_warps 등.
• 같은 IR 로 lowering — Triton-IR / TTGIR 까지 두 stack 이 통합. Linear Layout 추상이 두 surface 의 통합을 가능하게 함.
• profiling 도 통합 — Triton 의 TRITON_INTERPRET, ~/.triton/cache 등 모두 Gluon 에도 적용.
• 점진적 migration — 한 GEMM kernel 의 mainloop 만 Gluon 으로, epilogue 는 Triton 그대로. 충분히 가능.

2025 시점에서 Gluon 은 active development 중이지만 — 일부 production 사용처가 보이기 시작. 강의 시점에서 거론될 가능성이 큰 사례들.

실전 결정의 한 가지 trick — “이 kernel 의 maintenance window 가 얼마나 긴가”. 짧으면 (research 단계) Gluon, 길면 (production 다년간 도는 stack) CUTLASS 의 깊이 정확한 control 이 가치 있을 수 있음. 단 — Gluon 의 Linear Layouts 가 maturity 를 갖추면 이 결정이 흐려질 가능성.

• AMD support — Triton 이 AMD MI300 에 도착. Gluon 도 AMD 의 layout 결정에 같은 추상으로 적용. Linear Layouts 가 vendor-neutral.
• TMA descriptor 의 first-class API — 현재는 partial. multicast / swizzle mode / prefetch 의 fine control 이 Python 에서 자연스럽게.
• per-thread operation — Triton issue #8580 — Gluon 안에서 thread 별로 다른 operation. warp specialization 의 더 깊이.
• aggregate / OOP — issue #8781 — Gluon 안에서 type 추상 (struct, class). 큰 kernel 의 modularity.
• TTGIR layout attribute 의 노출 — issue #9789 — sizePerThread, warpsPerCTA 의 직접 접근. 사용자 control 의 더 깊이.
• Blackwell 의 새 instruction — sm_100a 의 MMA 변형, FP4 native, TMEM (tensor memory). 새 layout 추상이 추가될 가능성.

손에 새기기 — 실습 시퀀스

1. Triton GEMM baseline — Triton 으로 simple GEMM. NCU 로 측정. Hopper GPU (H100) 에서 CUTLASS reference 와 throughput gap 측정.
2. Gluon experimental setup — Triton 의 experimental/gluon 디렉토리 코드를 빌드. 단순한 vector add 부터.
3. Gluon GEMM 한 줄 — 같은 GEMM 을 Gluon 으로 변환. shared layout, MMA layout 명시. throughput 비교.
4. warp specialization — producer (TMA load) + consumer (WGMMA) 분리. mbarrier 의 sync. throughput 의 다음 step.
5. Linear Layout 손계산 — 8-element vector 의 4 layout (blocked, strided, swizzled, MMA) 을 𝔽₂ 위 3×3 행렬로 직접 적어본다.
6. conversion = matmul — 두 layout 사이의 conversion 행렬을 곱해 결과 layout 의 행렬 도출. 코드 없이 손으로.
7. 점진적 migration — 자기 Triton kernel 의 mainloop 만 Gluon 으로 옮기고 나머지는 그대로. 둘이 같은 파일에서 도는지 확인.

• Gluon 의 정확한 surface API — slides 와 강의 라이브 데모로 확인 필요. experimental/gluon 의 detail 이 빠르게 변하는 단계. (확인 필요)
• 구체적 throughput gap — Hopper 위 GEMM 의 Triton vs Gluon vs CUTLASS 의 측정 — 강의 안에서 보였을 가능성. (확인 필요)
• Linear Layout 의 행렬 표현 detail — 강의에서 어떤 행렬 예시를 보였는지. paper 본문과 강의 슬라이드의 일치 / 차이. (확인 필요)
• AMD 측 layout — AMD MI300 의 MFMA layout 이 Linear Layout 으로 어떻게 표현되는지. (확인 필요)
• future API stability — Gluon 이 stable Triton API 로 졸업하는 timeline. experimental → stable. (확인 필요)
• torch.compile 와의 통합 — torch.compile 이 Gluon 코드를 직접 lowering target 으로 쓸 수 있는지. (확인 필요)
• per-thread operation 의 모양 — issue #8580 의 진행 상태 — 강의 시점에서 어디까지. (확인 필요)