필사 모드: GPU 컴파일러와 프레임워크 지형 — 그래프를 받아 커널을 만드는 하나의 문제, 층마다 다른 답

• 들어가며 — 프로파일러가 가리킨 커널을 아무도 손으로 쓰지 않았습니다
• 공통 문제 하나 — 그래프를 받아 커널을 만든다
• 바닥층 — NVCC, PTX, SASS, 그리고 LLVM의 GPU 백엔드
• MLIR과 그 위에 올라간 것 — IR을 만드는 틀, 그리고 Triton
  • 그 위에 올라간 것 — Triton이 커널을 만드는 경로
• torch.compile과 Inductor가 실제로 하는 일
• XLA, IREE, TVM — 그래프를 통째로 받는 쪽
• 어느 층을 언제 알아야 하는가
• 마치며 — 층은 성능이 아니라 결정권으로 나뉩니다
• 참고 자료

들어가며 — 프로파일러가 가리킨 커널을 아무도 손으로 쓰지 않았습니다

Nsight Compute로 학습 스텝을 잡아 보면 시간을 가장 많이 쓰는 커널 이름이 나옵니다. 그런데 그 이름이 triton_poi_fused_add_mul_native_layer_norm_7 같은 것이면 곤란해집니다. 저장소를 검색해도 없습니다. 아무도 쓰지 않았기 때문입니다. 기계가 만들었습니다.

만든 것이 무엇인지 알아내려면 이런 질문에 답할 수 있어야 합니다. 그 커널은 어느 단계에서 생성됐습니까. 파이토치가 만들었습니까, Triton이 만들었습니까, 아니면 그 아래의 LLVM이 만들었습니까. 이름을 바꾸고 싶다면, 아니 코드를 고치고 싶다면 어느 층에 손을 대야 합니까.

이 질문이 어려운 이유는 층이 많기 때문이 아니라, 층마다 이름이 다르고 소속 조직이 다르고 문서가 따로 있어서 서로 무관해 보이기 때문입니다. NVCC와 LLVM과 MLIR과 Triton과 Inductor와 XLA와 IREE와 TVM은 각자 다른 세계의 물건처럼 보입니다.

실제로는 전부 같은 문제를 풉니다. 이 글은 그 문제가 무엇인지 먼저 말하고, 각 층이 그 문제의 어느 부분을 잘라 갔는지 정리합니다. 이 시리즈의 앞선 글들이 커널을 직접 고치는 일과 세 가지 커널 작성 층위, 추론 엔진 튜닝, 두 벤더의 스택 차이를 다뤘다면, 이 글은 그것들이 놓인 좌표계입니다.

버전과 사실 관계는 2026년 8월 2일에 각 프로젝트의 공식 문서와 릴리스 노트에서 직접 확인했습니다. 이 영역은 6개월이면 지형이 바뀌므로, 확인한 것과 확인하지 못한 것을 본문에서 구분해 적었습니다.

공통 문제 하나 — 그래프를 받아 커널을 만든다

이름이 무엇이든 이 글에 나오는 도구들은 입력과 출력이 같습니다.

입력은 연산 그래프입니다. 행렬곱, 덧셈, 활성함수, 정규화 같은 노드가 텐서로 연결된 것입니다. 출력은 GPU에서 실제로 실행되는 기계어입니다. 그 사이에서 해야 할 일도 놀라울 만큼 비슷합니다.

• 퓨전 결정입니다. 어떤 노드들을 하나의 커널로 합칠 것인가. 합치면 중간 텐서를 전역 메모리에 쓰지 않아도 되므로 메모리 대역폭을 아낍니다. XLA 문서는 퓨전을 "XLA의 단일 최중요 최적화"라고 부릅니다.
• 레이아웃 결정입니다. 텐서를 메모리에 어떤 순서로 놓을 것인가. 같은 연산이라도 레이아웃에 따라 코얼레싱이 되기도 하고 안 되기도 합니다.
• 타일링 결정입니다. 큰 행렬을 어떤 크기의 조각으로 잘라 각 스레드블록에 나눠 줄 것인가. 조각이 크면 공유 메모리와 레지스터가 모자라고, 작으면 재사용이 줄어듭니다.
• 라이브러리 호출과 코드 생성 사이의 선택입니다. 이 연산은 cuBLAS를 부르는 편이 빠른가, 직접 커널을 찍는 편이 빠른가.
• 로워링입니다. 결정한 것을 실제 명령어로 내리는 단계입니다.

층이 여럿 생긴 이유는 이 결정들의 성격이 서로 다르기 때문입니다. 퓨전은 그래프 전체를 봐야 판단할 수 있고, 레지스터 할당은 명령어 수준에서만 판단할 수 있습니다. 하나의 추상화로 둘 다 잘하기는 어렵습니다. 그래서 위쪽 층은 그래프를 보고, 아래쪽 층은 명령어를 보고, 그 사이에 중간층이 끼어 있습니다.

정리하면 이렇습니다.

이 표를 머리에 두고 아래를 읽으면 각 도구의 위치가 정해집니다.

바닥층 — NVCC, PTX, SASS, 그리고 LLVM의 GPU 백엔드

가장 아래부터 봅니다. 여기가 유일하게 20년째 구조가 크게 안 바뀐 곳입니다.

.cu 파일을 nvcc에 넣으면 호스트 코드와 디바이스 코드가 갈라집니다. 디바이스 코드는 두 가지 형태로 나옵니다. PTX와 CUBIN입니다. CUDA 13.3 컴파일러 문서가 설명하는 절차는 이렇습니다. 디바이스 함수가 PTX 또는 CUBIN으로 컴파일되고, 그것들이 fatbinary 컨테이너에 담기고, 원본 코드는 그 fatbinary를 품은 표준 C++로 변환되어 호스트 컴파일러에 넘어갑니다. 실행 시점에 "임베드된 fatbinary를 CUDA 런타임이 검사"해서 지금 GPU에 맞는 이미지를 고릅니다.

여기서 이해해야 할 개념이 가상 아키텍처와 실 아키텍처의 구분입니다.

# compute_90 = 가상 아키텍처, PTX를 생성할 명령어 집합 세대
# sm_90      = 실 아키텍처, 실제 칩용 바이너리
nvcc -gencode arch=compute_90,code=sm_90 \
     -gencode arch=compute_90,code=compute_90 \
     kernel.cu -o kernel

# 앞줄은 Hopper용 SASS를, 뒷줄은 PTX 자체를 fatbinary에 넣습니다.
# 뒷줄이 있어야 미래 세대 GPU에서 드라이버가 JIT으로 살려 냅니다.

문서의 표현대로, PTX를 함께 넣어 두면 "현재 GPU에 맞는 바이너리 로드 이미지가 없을 때 임베드된 PTX 코드를 CUDA 런타임이 동적으로 컴파일"합니다. 이것이 전방 호환성의 전부입니다. 그리고 이것이 PTX가 사라지지 않는 이유입니다. PTX는 사람이 읽는 어셈블리처럼 생겼지만 실제로는 배포 포맷이자 안정적인 계약입니다. SASS는 세대마다 인코딩이 바뀌고 문서화되지 않지만 PTX는 문서화되어 있고 하위 호환을 지킵니다.

실무에서 이 구분이 물리는 지점이 있습니다. PyTorch 2.13 릴리스 노트를 보면 CUDA 13 빌드에서 ptxas가 더 이상 바이너리에 번들되지 않습니다. PTX를 SASS로 바꾸는 도구가 휠 안에 없다는 뜻이고, PTX JIT 경로에 의존하는 환경이나 ptxas를 직접 호출하는 커스텀 빌드 스크립트가 갑자기 깨질 수 있습니다. 층을 몰라도 되는 것과 층이 없어도 되는 것은 다릅니다.

PTX에서 SASS로 가는 마지막 구간은 ptxas가 담당합니다. 이 단계에서 레지스터 할당과 명령어 스케줄링이 일어나고, 점유율을 좌우하는 레지스터 개수도 여기서 정해집니다. 커널 소스를 안 고쳤는데 CUDA 툴킷 마이너 버전만 올렸더니 성능이 변하는 현상은 거의 항상 ptxas의 스케줄링이 달라졌기 때문입니다. 아래 도구들로 확인할 수 있습니다.

# 실제로 생성된 SASS를 봅니다
cuobjdump -sass ./kernel | head -40

# 레지스터와 공유 메모리 사용량 (점유율 계산의 입력값)
nvcc -Xptxas -v -arch=sm_90 -c kernel.cu
# ptxas info : Used 64 registers, 8192 bytes smem, 384 bytes cmem[0]

# PTX만 뽑아 보기
nvcc -ptx -arch=compute_90 kernel.cu -o kernel.ptx

그런데 CUDA C++만 GPU 기계어를 만드는 것은 아닙니다. LLVM에는 NVPTX라는 백엔드가 있어서 LLVM IR을 PTX로 내립니다. Triton도, XLA도, 그 외 거의 모든 코드 생성기가 마지막에는 이 백엔드를 거칩니다. 즉 PTX는 CUDA의 산출물이면서 동시에 LLVM의 산출물입니다.

AMD 쪽도 구조는 같고 이름만 다릅니다. LLVM의 AMDGPU 백엔드가 LLVM IR을 AMDGCN 어셈블리로 내리고, 그것이 코드 오브젝트가 됩니다. 다만 AMD는 업스트림 LLVM을 그대로 쓰지 않고 포크를 유지합니다. ROCm 문서는 자사 컴파일러를 "llvm/llvm-project의 포크"라고 명시하고, amdclang++를 기본 컴파일러로, hipcc를 "clang 또는 nvcc를 호출하고 적절한 include와 library 옵션을 넘기는" 드라이버로 설명합니다. 조회 시점 문서의 llvm-project 버전은 22.0.0입니다.

포크를 유지하는 것이 나쁜 신호는 아닙니다. 새 칩 지원과 기능이 업스트림 릴리스 주기보다 빨리 필요하기 때문이고, AMD 변경분은 상당수가 업스트림에 올라갑니다. 다만 실무적으로는 의미가 있습니다. AMD GPU를 대상으로 하는 서드파티 컴파일러 프로젝트는 자기 LLVM 버전과 ROCm의 LLVM 버전 사이에서 조율을 해야 하고, 이것이 두 벤더 스택을 비교한 글에서 다룬 이식 마찰의 한 갈래입니다.

MLIR과 그 위에 올라간 것 — IR을 만드는 틀, 그리고 Triton

여기서부터 지형이 달라집니다.

LLVM IR은 CPU를 위해 설계됐습니다. 스칼라와 벡터, 기본 블록, SSA. 텐서라는 개념이 없고 루프 중첩이라는 개념도 없습니다. 행렬곱을 LLVM IR로 표현하면 이미 루프가 풀린 저수준 코드가 되어 있어서, "이 루프를 타일링하자" 같은 결정을 내릴 정보가 남아 있지 않습니다.

그래서 2010년대 후반의 모든 ML 컴파일러가 같은 일을 반복했습니다. 텐서 수준 IR을 새로 정의하고, 패스 인프라를 새로 짜고, 파서와 검증기와 출력기를 새로 만들고, 마지막에 LLVM IR로 내리는 코드를 새로 썼습니다. XLA도, TVM도, 각 하드웨어 벤더도 각자 했습니다.

MLIR의 제안은 그 반복을 없애자는 것입니다. IR 하나를 더 정의하는 대신, IR을 정의하는 인프라를 제공합니다. 연산과 타입의 묶음을 다이얼렉트라고 부르고, 프로젝트마다 자기 다이얼렉트를 정의하되 패스 관리, 검증, 직렬화, 로워링 프레임워크는 공유합니다.

핵심은 다이얼렉트가 공존한다는 점입니다. 하나의 모듈 안에 추상도가 다른 연산이 섞여 있을 수 있고, 로워링은 그 섞임을 점진적으로 해소하는 과정이 됩니다.

// 같은 함수 안에서 추상도가 다른 두 다이얼렉트가 공존합니다.
// linalg는 "행렬곱"이라는 의미를 유지하고,
// scf/arith는 이미 루프와 스칼라 연산으로 내려와 있습니다.
func.func @mm(%A: tensor<128x256xf32>, %B: tensor<256x64xf32>)
    -> tensor<128x64xf32> {
  %init = tensor.empty() : tensor<128x64xf32>
  %C = linalg.matmul
      ins(%A, %B : tensor<128x256xf32>, tensor<256x64xf32>)
      outs(%init : tensor<128x64xf32>) -> tensor<128x64xf32>
  return %C : tensor<128x64xf32>
}

linalg.matmul 한 줄은 "이것은 행렬곱이다"라는 정보를 그대로 담고 있습니다. 타일링 패스가 이 정보를 보고 루프를 자를 수 있습니다. 같은 것을 LLVM IR로 먼저 내려 버리면 그 정보가 사라져서, 이후에는 루프 재구성이 훨씬 어려워집니다. 정보를 늦게 버릴수록 좋은 결정을 내릴 수 있다 — 이것이 MLIR이 존재하는 이유를 한 문장으로 줄인 것입니다.

MLIR은 스스로 무엇도 컴파일하지 않습니다. 위에 무엇을 올리느냐가 전부입니다. 그리고 실제로 많이 올라갔습니다. Triton의 IR도, IREE의 IR도, XLA의 일부 구성 요소도, 여러 하드웨어 벤더의 사내 컴파일러도 MLIR 위에 있습니다. 이름이 다른 프로젝트들이 같은 기반을 공유하고 있다는 사실은 이 지형을 이해하는 데 중요합니다.

그 위에 올라간 것 — Triton이 커널을 만드는 경로

Triton은 커널 작성 언어이면서 동시에 MLIR 기반 컴파일러입니다. 세 가지 커널 작성 층위를 비교한 글에서 언어로서의 Triton을 다뤘으니, 여기서는 컴파일러로서의 위치를 봅니다.

파이프라인이 명확하게 정의되어 있습니다. 파이토치 팀이 2026년 7월에 공개한 Triton 플러그인 확장 문서가 단계를 이렇게 적고 있습니다.

파이썬 커널 함수
      ↓  (AST 순회)
TTIR    — Triton IR. 타일 단위 연산. 하드웨어 무관.
      ↓  (레이아웃 결정, 스레드 매핑)
TTGIR   — TritonGPU IR. 워프 배치와 공유 메모리가 결정됨.
      ↓
LLVM IR
      ↓
PTX (NVIDIA)  또는  AMDGCN (AMD)

이 네 단계가 이 글의 앞부분에서 정리한 "층" 구조를 압축해 보여 줍니다. TTIR은 커널 층, TTGIR은 커널과 명령어의 경계, LLVM IR 아래는 명령어 층입니다. 사용자는 TTIR 수준으로 코드를 쓰고 나머지는 컴파일러에 맡깁니다. 이것이 Triton의 거래 조건입니다.

지원 백엔드는 저장소 README 기준으로 NVIDIA(Compute Capability 8.0 이상)와 AMD(ROCm 6.2 이상)이고, CPU는 "개발 중"으로 표시되어 있습니다. CPU 백엔드는 여러 갈래로 시도되고 있지만 조회 시점 기준으로 실험 단계이며, 실무에서 기대할 대상은 아직 아닙니다.

2026년에 이 층에서 실제로 일어난 변화 중 눈여겨볼 것은 플러그인 확장 체계입니다. 파이토치가 배포하는 pytorch-triton 3.7에 들어갔고, 파이토치 2.13은 Triton 3.7.1을 핀으로 잡고 있습니다. 이전에는 커스텀 패스나 다이얼렉트를 넣으려면 Triton을 포크해야 했습니다. 포크는 금방 업스트림에 뒤처지고, 병합 충돌이 쌓이고, 결국 새 하드웨어 지원을 못 받는 낡은 릴리스에 묶입니다. 플러그인 체계는 런타임에 공유 라이브러리를 얹어 파이프라인 임의 지점에 패스를 삽입하거나, 특정 패스를 끄거나, 단계 전체를 교체할 수 있게 합니다.

# 플러그인은 환경변수로 로드됩니다. import 순서가 중요합니다.
import os, sysconfig

dist_packages = sysconfig.get_paths()["purelib"]
os.environ["TRITON_PLUGIN_PATHS"] = os.path.join(
    dist_packages, "utlx_plugin", "libutlx.so"
)

import triton
import triton.language as tl
import utlx_plugin as tlx      # TLX 확장 연산이 여기서 들어옵니다

첫 소비자가 메타의 TLX(Triton Language Extensions)입니다. tlx.local_alloc, tlx.async_load, tlx.async_dot 같은 연산으로 공유 메모리 버퍼와 비동기 파이프라이닝을 커널 안에서 명시적으로 다룹니다. 즉 Triton이 자동으로 하던 결정 중 일부를 사람이 도로 가져오는 방향입니다. 발표된 수치로는 H100 퍼시스턴트 GEMM에서 cuBLAS 대비 2~4퍼센트 앞서고, MI350 파이프라인 GEMM에서 rocBLAS 대비 12~15퍼센트 앞선다고 되어 있습니다. 벤더 라이브러리보다 빠른 커널이 컴파일러 언어로 나온다는 것이 요점입니다.

여기서 흐름 하나가 보입니다. 추상화 층은 위로 올라가기만 하지 않습니다. Triton이 스레드를 감춰서 위로 올라갔다면, TLX는 비동기 명령어를 도로 노출해서 아래로 내려옵니다. 추상화의 높이는 성능이 아니라 어느 결정을 사람이 쥘 것인가로 정해집니다.

torch.compile과 Inductor가 실제로 하는 일

파이토치 사용자가 가장 자주 만나면서 가장 덜 이해하는 층입니다. torch.compile(model) 한 줄 뒤에서 세 단계가 순서대로 돕니다.

TorchDynamo는 파이썬 바이트코드를 가로채서 FX 그래프를 뽑아냅니다. 순수 텐서 연산은 그래프로 들어가고, 파이썬 리스트 조작이나 print 같은 것은 그래프 밖으로 밀려납니다. 그 지점이 그래프 브레이크입니다. 그래프가 여러 조각으로 잘리면 각 조각의 앞뒤에서 파이썬으로 돌아왔다 나가는 비용이 생기고, 조각이 작으면 퓨전 기회도 사라집니다. torch.compile이 기대만큼 안 빨라지는 사례의 대부분은 커널이 나빠서가 아니라 그래프가 잘려서입니다.

AOTAutograd는 그 그래프를 순전파와 역전파 두 개의 ATen 수준 그래프로 분리하고, 무엇을 저장하고 무엇을 다시 계산할지 정합니다.

TorchInductor가 실제 코드를 찍습니다. GPU 대상이면 Triton 커널을, CPU 대상이면 C++와 OpenMP를 생성합니다. 그래서 앞에서 본 triton_poi_fused_... 같은 이름이 나옵니다. 이름 자체가 정보를 담고 있습니다. poi는 pointwise, fused_ 뒤에 붙은 것들이 합쳐진 연산 목록입니다.

생성된 코드는 눈으로 볼 수 있습니다. 이것이 이 층을 이해하는 가장 빠른 길입니다.

import os
os.environ["TORCH_COMPILE_DEBUG"] = "1"        # 생성 코드를 디스크에 남깁니다
os.environ["TORCH_LOGS"] = "output_code,graph_breaks"

import torch

def block(x, w, b):
    return torch.nn.functional.gelu(x @ w + b)

x = torch.randn(1024, 512, device="cuda", dtype=torch.bfloat16)
w = torch.randn(512, 512, device="cuda", dtype=torch.bfloat16)
b = torch.randn(512, device="cuda", dtype=torch.bfloat16)

compiled = torch.compile(block)
out = compiled(x, w, b)

# 그래프 브레이크만 따로 확인하는 방법
explanation = torch._dynamo.explain(block)(x, w, b)
print("그래프 개수 :", explanation.graph_count)
print("브레이크 수 :", explanation.graph_break_count)

TORCH_LOGS="output_code"를 켜면 Inductor가 만든 Triton 소스가 표준 출력으로 나옵니다. 행렬곱은 보통 cuBLAS 호출로 남고, GELU와 바이어스 덧셈이 하나의 pointwise 커널로 합쳐져 있는 것을 볼 수 있습니다. 그 커널이 바로 앞서 프로파일러가 가리킨 이름 없는 커널의 정체입니다.

2026년에 이 층에서 생긴 중요한 변화가 하나 있습니다. 파이토치 2.13부터 Inductor가 Triton 말고 CuTeDSL 경로를 추가로 갖습니다. 엔비디아 CuTe 위에 올라간 DSL이고, 릴리스 노트에 따르면 트랜스포머의 GEMM과 RMSNorm을 겨냥해 "Triton 없이도 더 나은 품질의 행렬곱 코드를 만드는" 두 번째 고성능 경로입니다. 커널 컴파일이 스레드 풀에서 서브프로세스 풀로 옮겨진 것도 같은 릴리스인데, 파이썬 GIL이 컴파일 시간의 병목이었기 때문입니다.

같은 릴리스에 torch.compiler.set_default_backend가 들어갔습니다. torch.set_default_dtype과 같은 방식으로 프로세스 전역 기본 백엔드를 정하는 API이고, 모든 torch.compile 호출에 backend=를 붙이지 않아도 됩니다. 트리 밖 백엔드를 만드는 하드웨어 벤더에게 의미가 큰 변화입니다.

정리하면 torch.compile은 그래프 컴파일러이고, 커널 생성은 아래층에 위임합니다. 그 아래층이 Triton 하나였다가 지금은 둘이 됐습니다.

XLA, IREE, TVM — 그래프를 통째로 받는 쪽

Inductor가 파이토치에 붙어 있는 그래프 컴파일러라면, 프레임워크와 독립적으로 그래프를 받겠다는 계열이 따로 있습니다.

XLA가 가장 오래됐고 가장 많이 쓰입니다. JAX의 유일한 실행 경로이고, TPU의 유일한 실행 경로이며, 파이토치도 PyTorch/XLA로 붙습니다. 입력은 HLO이고, 프레임워크와 컴파일러 사이의 계약은 StableHLO가 맡습니다. StableHLO는 하위 5년, 상위 2년의 호환을 약속하는 연산 집합이라 저장해 둔 그래프를 몇 년 뒤 컴파일러로 여는 일이 성립합니다.

XLA:GPU의 코드 생성 방식이 이 글의 주제와 정확히 맞물립니다. 공식 문서의 표현을 그대로 옮기면 XLA:GPU는 "네이티브(LLVM 경유 PTX) 이미터와 TritonIR 이미터의 조합"을 씁니다. 세 갈래로 나뉩니다.

1. 흔한 연산은 cuBLAS, cuDNN, NCCL을 그냥 부릅니다.
2. 리덕션이나 전치처럼 패턴이 잡히는 연산은 LLVM IR을 직접 만들어 PTX로 내립니다.
3. 행렬곱이나 소프트맥스가 낀 고급 퓨전은 HLO 퓨전을 TritonIR로 변환하고, 타일 파라미터를 고른 뒤 Triton을 호출해 PTX를 받습니다.

세 번째가 중요합니다. XLA와 Triton은 경쟁 관계가 아니라 호출 관계입니다. XLA는 그래프 층의 결정을 하고, 커널 층의 결정은 Triton에 넘깁니다. 앞의 표에서 층을 나눈 것이 그대로 구현되어 있는 셈입니다. 마지막으로 XLA 런타임이 커널 호출과 라이브러리 호출의 나열을 RuntimeIR이라는 자체 MLIR 다이얼렉트로 옮기고, 거기서 CUDA 그래프를 추출합니다.

IREE는 같은 문제를 더 넓은 타깃 범위에서 풉니다. README의 표현으로는 "ML 모델을 통합 IR로 내려서 데이터센터의 요구까지 확장되고 모바일과 엣지의 제약까지 축소되는" MLIR 기반 컴파일러이자 런타임입니다. 저장소 주제 태그에 JAX, PyTorch, ONNX, TensorFlow가 프런트엔드로, CUDA, ROCm, Vulkan, SPIR-V가 타깃으로 올라 있습니다. 2024년 5월에 LF AI & Data 재단의 샌드박스 단계 프로젝트로 들어갔습니다.

IREE의 특징은 컴파일러와 런타임을 함께 설계했다는 점입니다. 결과물이 파이썬 프로세스가 필요한 객체가 아니라 자립 실행 모듈이라, 파이썬이 없는 임베디드 타깃까지 같은 파이프라인으로 갑니다. 데이터센터에서 IREE를 만나는 경우는 아직 드물고, AMD가 자사 스택의 일부로 밀고 있는 것이 조회 시점의 가장 큰 채택 사례입니다.

TVM은 이 지형에서 위치가 가장 많이 바뀐 프로젝트입니다. 2018년 무렵에는 유일하게 진지한 오픈소스 딥러닝 컴파일러였고, 오토튜닝으로 벤더 라이브러리를 이긴다는 주장으로 주목받았습니다. 지금은 Relax를 그래프 수준 표현으로, TensorIR을 텐서 수준 표현으로 쓰는 교차 층 설계이며, 문서가 밝히는 설계 목표는 "대부분의 변환을 파이썬에서 커스터마이즈 가능하게 만들어 ML 컴파일러를 접근 가능하게" 하는 것입니다.

솔직하게 적자면, 프런티어 LLM 학습과 서빙에서 TVM을 만날 일은 이제 거의 없습니다. 그 자리는 Inductor와 XLA가 가져갔습니다. TVM이 살아 있는 곳은 다릅니다. 벤더 라이브러리가 없는 하드웨어, 파이썬 런타임을 올릴 수 없는 엣지 타깃, 그리고 MLC LLM처럼 브라우저나 모바일에서 LLM을 돌리는 배포 경로입니다. 오토튜닝으로 튜닝된 라이브러리를 이긴다는 원래 명제는 cuBLAS와 CUTLASS가 있는 곳에서는 성립하기 어렵고, 그런 것이 없는 곳에서는 여전히 성립합니다.

한 표로 비교하면 이렇습니다.

어느 층을 언제 알아야 하는가

층이 많다고 다 알아야 하는 것은 아닙니다. 필요할 때만 내려가면 됩니다. 그 "필요할 때"의 신호가 무엇인지가 이 글에서 가장 실용적인 부분입니다.

경험상 실무자의 90퍼센트는 위 두 줄에서 끝납니다. 그래프 브레이크 하나를 없애는 것이 커널을 손으로 최적화하는 것보다 대개 더 큰 이득을 줍니다. 아래로 내려갈수록 얻는 이득은 커지지만 소요 시간이 훨씬 더 가파르게 늘어납니다.

Triton IR을 직접 보고 싶다면 덤프 방법은 이렇습니다.

# 각 단계의 IR을 파일로 남깁니다
export TRITON_KERNEL_DUMP=1
export TRITON_DUMP_DIR=/tmp/triton_ir
python train.py

ls /tmp/triton_ir/*/
# *.ttir     TTIR   — 타일 단위, 하드웨어 무관
# *.ttgir    TTGIR  — 레이아웃과 워프 배치가 결정된 상태
# *.llir     LLVM IR
# *.ptx      PTX
# *.cubin    SASS를 담은 최종 바이너리

같은 커널의 다섯 파일을 나란히 열어 보는 것이 이 글 전체를 읽는 것보다 이해에 빠릅니다. 위에서 아래로 내려갈수록 무엇이 결정되고 무엇이 사라지는지가 눈에 보입니다.

마지막으로 이번 조사에서 확인하지 못한 것을 밝혀 둡니다.

• Triton의 안정 릴리스 버전과 날짜입니다. 저장소 README는 버전을 명시하지 않고, 파이토치가 핀으로 잡은 3.7.1만 확인했습니다. Triton 자체의 릴리스 번호와 파이토치가 배포하는 pytorch-triton 번호는 일치하지 않을 수 있습니다.
• IREE의 최신 릴리스 태그와 날짜입니다. README에 표시되지 않았고 별도로 확인하지 않았습니다.
• Apache TVM의 최근 릴리스 시점과 활동 수준입니다. 프로젝트 문서로 설계 방향은 확인했지만 릴리스 이력을 직접 조회하지는 않았습니다. 위의 서술은 채택 현황에 대한 판단이며 통계로 뒷받침한 것이 아닙니다.
• Helion은 파이토치가 2026년에 밀고 있는 커널 DSL이고 TPU 백엔드가 Pallas로 컴파일된다는 것까지는 확인했지만, NVIDIA와 AMD 경로에서 Triton과 어떤 관계인지는 해당 문서에서 확인하지 못했습니다.

마치며 — 층은 성능이 아니라 결정권으로 나뉩니다

이 글의 지도를 한 문장으로 접으면 이렇게 됩니다. 층이 나뉘는 기준은 속도가 아니라 누가 어떤 결정을 쥐느냐입니다.

그래프 층은 퓨전과 레이아웃을 가져갑니다. 그 대가로 커널 하나의 세부는 포기합니다. 커널 층은 타일과 파이프라이닝을 가져갑니다. 그 대가로 그래프 전체를 못 봅니다. 명령어 층은 레지스터와 스케줄링을 가져갑니다. 그 대가로 무엇을 계산하는지 모릅니다. 어떤 층도 다른 층을 대체하지 못하는 이유가 여기 있고, 30년 전 CPU 컴파일러가 겪은 층 분화와 구조가 같습니다.

그래서 새 프로젝트 이름이 나올 때 던질 질문도 정해져 있습니다. 무엇이 새로운가가 아니라, 어느 결정을 누구에게서 가져갔는가입니다. Triton은 스레드 매핑을 사람에게서 가져갔습니다. TLX는 비동기 명령어를 컴파일러에게서 도로 가져왔습니다. Inductor는 퓨전을 사용자에게서 가져갔습니다. MLIR은 IR 인프라를 각 프로젝트에게서 가져갔습니다.

이 질문 하나면 새로 나오는 프레임워크를 이 지도의 어디에 놓을지 대체로 정할 수 있습니다. 그리고 그것이, 6개월마다 바뀌는 이름을 전부 쫓아다니는 것보다 오래 쓸 수 있는 지식입니다.

참고 자료

• NVIDIA CUDA Compiler Driver NVCC 13.3 — 컴파일 궤적과 가상·실 아키텍처
• ROCm LLVM Project 문서 — amdclang, hipcc, comgr
• Triton 저장소 README — 지원 백엔드와 MLIR 기반 재작성
• Triton Plugin Extensions: TLX and Custom Compiler Passes — PyTorch 블로그, 2026-07-15
• PyTorch 2.13 릴리스 노트 — CuTeDSL 백엔드, Triton 3.7.1 핀, CUDA 13 변경
• XLA:GPU Architecture Overview — 네이티브 이미터와 TritonIR 이미터
• StableHLO — 프레임워크와 컴파일러 사이의 호환 계층
• IREE 저장소 README — MLIR 기반 컴파일러와 런타임
• Apache TVM — Relax와 TensorIR 교차 층 설계
• MLIR 프로젝트 — 다이얼렉트와 점진적 로워링