cuTile — NVIDIA 의 tile-based GPU programming model

2024–2025 시점 GPU programming 의 사용자 표면은 — 한 마디로 — fragmented. CUDA C++ 의 직접 작성, CUTLASS 의 template, CuTe 의 layout algebra, Triton 의 Python DSL, Helion 의 PyTorch-native, JAX 의 Pallas, OpenAI 의 ttgir. 사용자가 매번 “어떤 도구를 어디에 쓸 것인가” 의 결정에 시간을 쓴다.

이 fragmentation 안에서 NVIDIA 가 풀어야 하는 문제는 분명하다 — 새 hardware 의 새 명령을 사용자에게 가장 빨리 노출하는 길. CUTLASS 가 그 자리였지만 진입 장벽이 너무 높다. Triton 이 빠르게 받아주지만 — 그건 OpenAI 가 끄는 프로젝트. NVIDIA 가 자기 컴파일러 stack 위에서 같은 가치를 직접 제공할 수 있어야 한다는 입장.

1. cuTile 은 NVIDIA 의 자체 tile-based DSL — Python 표면, MLIR 백엔드, NVIDIA 컴파일러 stack 과 deeply integrated.
2. Triton 의 사상은 받되, NVIDIA 의 새 명령 (Hopper TMA, Blackwell tensor memory accelerator, 새 mma 변형) 을 가장 먼저 노출.
3. 같은 사용자 표면 위에서 CuTe 와 cuBLAS 까지 seamless 하게 — DSL 로 짠 커널이 cuBLAS 와 같이 잘 도는 자리.

Mehdi Amini 가 NVIDIA 에 합류하면서 들고 온 자산은 — MLIR 의 깊은 이해. cuTile 은 MLIR-first 디자인이 강조된다 — 즉, 컴파일러 인프라의 한 dialect 으로 시작해서 사용자 표면이 따라온다. Triton 도 결국 MLIR-based 가 됐지만, cuTile 은 그 사상이 더 명시적.

tile-based 추상이 무엇인지 — Triton 과 cuTile 이 공유하는 사상의 핵심을 한 번 정리. 그 위에서 cuTile 의 차이가 보인다.

이 모델이 왜 자연스러운가 — GPU 의 SIMT 와 systolic 의 본질이 모두 “같은 일을 여러 element 에 동시에” 다. tile 이 그 단위를 사용자 표현 차원에서 끌어올린 것.

같은 영역의 세 도구. 쟁점 — 사용자가 어디에 묶일 것인가. 산업의 답이 보통 “하나로 통합” 이 아니라 “여러 도구가 같이 산다” 의 형태로 끝나는 게 일반적. 각 도구의 sweet spot 이 다르다.

이 디자인의 가치는 — MLIR 의 일반적 가치와 같다. 각 lowering pass 가 독립적으로 작성, 검증, 재사용 가능. 같은 NVGPU dialect 가 cuTile 뿐 아니라 다른 사용자 표면 (linalg, OpenXLA) 에서도 lower 의 target 으로 쓰일 수 있다.

실제 API 는 NVIDIA 의 공식 release 가 답. 아래는 외부 자료 + 도메인 지식 기반의 개념적 형태. 강의에서 발표자가 라이브로 보여줬을 가능성이 높은 두 예시.

# 예시 1 — vector add (개념적)
import nvidia.cutile as cutile

@cutile.kernel
def vec_add(x: cutile.Tensor, y: cutile.Tensor,
            out: cutile.Tensor, n: int):
    pid    = cutile.program_id(0)
    offs   = pid * BLOCK + cutile.arange(0, BLOCK)
    mask   = offs < n
    a      = cutile.load(x + offs, mask=mask)
    b      = cutile.load(y + offs, mask=mask)
    cutile.store(out + offs, a + b, mask=mask)

# launch — 사용자 측에서 grid 만 깐다
vec_add[(N // BLOCK,)](x, y, out, N)

# 예시 2 — GEMM (개념적)
@cutile.kernel
def gemm(A, B, C, M, N, K):
    pid_m  = cutile.program_id(0)
    pid_n  = cutile.program_id(1)

    offs_m = pid_m * BM + cutile.arange(0, BM)
    offs_n = pid_n * BN + cutile.arange(0, BN)

    acc    = cutile.zeros((BM, BN), dtype=cutile.float32)

    for k in range(0, K, BK):
        offs_k = k + cutile.arange(0, BK)
        a = cutile.load(A[offs_m, offs_k])    # tile-level slicing
        b = cutile.load(B[offs_k, offs_n])
        acc += cutile.mma(a, b)               # NVIDIA-specific mma 명령

    cutile.store(C[offs_m, offs_n], acc.to(C.dtype))

표면에서는 — 거의 Triton 코드 같다. 차이는 — (1) cutile.mma 같은 NVIDIA-specific primitive 가 first-class, (2) launch grammar 가 NVIDIA 컴파일러와 deep integration, (3) layout / scheduling 의 hint 가 사용자 표면에서 더 정밀하게 노출 가능 — 같은 자리.

이 그림에서 cuTile 은 — “사용자 표면” 의 자리에서 새 entrant. 같은 자리에 Triton, Helion 이 있다. NVIDIA 의 입장은 — 자기 hardware 위에서 가장 잘 도는 사용자 표면을 own 한다는 것. Triton / Helion 과 경쟁이 아니라 “NVIDIA-first 사용자에게는 cuTile, cross-vendor 한 사용자에게는 Triton” 의 분담.

• NVIDIA 내부의 새 GEMM 변형 prototyping — 새 dtype (FP4/MXFP4), 새 mma 명령의 first-class 노출.
• NVIDIA 의 production AI stack — TensorRT-LLM, NIM 같은 inference engine 의 일부 hot path. 확인 필요.
• partner 협업 — 큰 customer 의 customized 커널 작성. CUTLASS C++ 보다 빠른 iteration.
• 학계 사용자 — MLIR-based 디자인이 학계 컴파일러 연구자의 자연스러운 표면.
• 새 hardware 의 시연 — Blackwell / Rubin 등의 새 명령을 보여줄 때 cuTile 코드가 가장 자연스러움.

발표 시점 (2025) 기준 cuTile 의 release 상태는 — 확인 필요. NVIDIA 내부 도구로 시작해 점차 public release 가 표준 패턴. 사용자 base 형성에 시간이 걸린다.

1. public release 와 ecosystem 형성 — repo, 문서, 예시, 튜토리얼. Triton 이 5 년 걸린 길.
2. PyTorch 통합 — torch.compile / Inductor 와의 backend 관계. Helion 같은 layer 가 cuTile 을 backend 로 쓸 가능성.
3. autotune — Triton 의 @triton.autotune 같은 자동 튜닝 인프라. Helion 의 search space 와 어떻게 합쳐질지.
4. CuTe DSL 과의 정확한 분담 — 두 도구가 NVIDIA 안에서 같이 산다. 어디서 cuTile, 어디서 CuTe 의 정확한 가이드.
5. 비-NVIDIA backend — 매우 확인 필요. NVIDIA-first 디자인이지만, MLIR 위에 있어 기술적으로는 가능. 정책의 문제.
6. 새 hardware 적응 — Rubin / 그 다음 세대의 명령을 가장 먼저 노출하는 게 cuTile 의 전략적 가치.

손에 새기기 — 학습 시퀀스

1. cuTile 공식 release 추적 — NVIDIA Developer Blog, GTC 발표, github.com/NVIDIA 의 새 repo 모니터링.
2. 같은 vector add 세 도구로 비교 — 같은 vector add 를 Triton, Helion, cuTile 로 짜고 PTX 까지 dump. 어떤 명령이 다른지 추적.
3. 같은 GEMM 비교 — 위와 같은 비교 실험을 GEMM 으로. 작은 / 큰 GEMM 모두에 대해 측정.
4. MLIR dialect 펼치기 — cuTile 의 lowering pass 를 한 번 dump 해서 본다. 같은 source 가 어떤 IR 을 거쳐 PTX 가 되는지.
5. 의사결정 사다리 — 자기 워크로드에 대해 “이 hot path 는 어떤 도구가 답인가” 의 결정 framework 작성. portable 인가, NVIDIA-only 인가, layout 통제 필요한가, autotune 자동 원하는가.

• cuTile 의 정확한 API — 노트의 코드 예시는 외부 자료 + Triton 의 일반적 형태를 기반으로 한 개념적 형태. 정확한 함수명 / namespace / decorator 는 NVIDIA 공식 release 이후 확인.
• release 상태 — public release 인지, NVIDIA 내부 도구인지. 발표 시점 (2025) 기준의 정확한 상태.
• Triton 과의 정확한 성능 비교 — 같은 GEMM / attention 의 cuTile vs Triton. NVIDIA 가 발표한 벤치마크 추적.
• PyTorch 통합 방식 — torch.compile backend 로 등장하는지, custom op 으로 register 하는지.
• Helion 과의 관계 — Helion 이 cuTile 을 backend 로 쓸 수 있는지. NVIDIA-first 와 PyTorch-native 의 합쳐짐.
• 비-NVIDIA backend 의 정책 — 기술적으로는 MLIR 위라 가능. NVIDIA 의 입장이 무엇인지.
• CuTe DSL 과의 분담 — 같은 NVIDIA 안의 두 사용자 표면이 어떻게 같이 사는지. 공식 가이드라인.