CuTile·TileIR로 Python 커널이 SASS 되는 법
“Python으로 쓴 커널 한 줄이 GPU에서 어떤 **SASS(실제 머신 명령)** 로 바뀌는지” 궁금했던 적 있나요? NVIDIA의 `CuTile`과 그 뒤를 받치는 `TileIR`은 그 과정을 꽤 노골적으로(그리고 흥미롭게) 보여주는 최신 스택이에요.
## CuTile이 바꾼 것: 스레드가 아니라 **타일(Tile)** 중심으로 생각하기
`CuTile`은 NVIDIA가 내놓은 **타일 중심(tile-centric) 프로그래밍 모델**이에요. 개발자는 “이 일을 어떤 블록/타일로 쪼갤지” 같은 고수준 구조에 집중하고, 수백 개 스레드에 어떻게 매핑할지는 컴파일러가 맡는 구조죠. 그래서 `ct.mma(a, b, acc)`처럼 한 줄로 쓴 연산이 실제로는 여러 개의 텐서 코어 명령으로 확장될 수 있어요.
이게 중요한 이유는, 최적화의 중심이 “스레드별 인덱싱”에서 “**데이터 조각(타일)의 이동과 계산**”으로 바뀌기 때문이에요. 특히 MoE(Mixture-of-Experts)처럼 gather/scatter가 많은 모델에서는 스레드 단위로 생각하면 금방 복잡도가 폭발하거든요. CuTile은 그런 복잡도를 “타일 단위”로 압축해주는 셈입니다.
## TileIR 파이프라인 한눈에 보기: `.cutile`에서 SASS까지 “점진적 하향(lowering)”
`TileIR`은 CuTile을 실제 GPU 코드로 내리는 **MLIR 기반 컴파일러 인프라**예요. 핵심은 한 번에 PTX/SASS로 뛰지 않고, 여러 단계의 IR(중간표현)을 거치며 점점 구체화를 한다는 점입니다.
글에서 제시한 전체 흐름은 아래처럼 정리돼요(문서 기준 CUDA 13.1):
– Python(CuTile API) → `.cutile` 바이트코드
– `cuda_tile` (MLIR dialect): 아키텍처 독립적인 고수준 타일/텐서 연산
– `nv_tileaa` (MLIR dialect): 메모리 참조가 명시되는 타일 로우레벨
– `nv_tileas` (MLIR dialect): 스케줄링/async 파이프라인/TCGEN 등 아키텍처 특화
– NVVM/LLVM IR → PTX → **SASS(sm_120 등 실제 명령)**
여기서 사용자 입장에서 눈여겨볼 도구가 `tileiras`예요. 글에서도 “TileIR용 ptxas 같은 존재”라고 표현하는데, 실제로 `.cutile`을 AOT(ahead-of-time)로 `cubin`까지 뽑습니다.
“`bash
tileiras –gpu-name sm_120 MoE.cutile -o moe.cubin
“`
## Running Example(MoE 커널): `gather / load / mma`가 어떻게 변하는가
글은 MoE 커널을 예제로, 특히 아래 3개 핵심 연산이 단계별로 어떻게 바뀌는지 추적해요.
– `ct.gather(A, indices)` → 간접 로드(토큰 인덱스로 모아 읽기)
– `ct.load(B, …)` → 직접 로드(연속 타일 로드)
– `ct.mma(a, b, acc)` → 텐서 코어 MMA(행렬 곱 누적)
단계별 매핑이 재밌는데요, 요지는 “이름이 바뀌는 게 아니라 **추상화가 벗겨진다**”예요.
– `cuda_tile`에서는 `load_view_tko`, `load_ptr_tko`, `mmaf` 같이 “의미 중심”이고
– `nv_tileaa`로 내려가면 `make_memref`, `tiled_load`, `dot`처럼 “주소/형상”이 더 노출돼요.
– `nv_tileas`에선 `async.pipeline.*`, `make_tiled_tma_desc`, `tcgen05.*` 같은 “하드웨어 실행 방식”이 본격 등장합니다.
즉, “내가 타일을 이렇게 쓰고 싶다” → “그럼 글로벌/공유/텐서 메모리를 이렇게 오가며” → “이 파이프라인과 배리어로 겹쳐 돌리고”로 바뀌는 흐름이에요.
## TileIR 파이프라인 한눈에 보기: `.cutile`에서 SASS까지 “점진적 하향(lowering)”
`TileIR`은 CuTile을 실제 GPU 코드로 내리는 **MLIR 기반 컴파일러 인프라**예요. 핵심은 한 번에 PTX/SASS로 뛰지 않고, 여러 단계의 IR(중간표현)을 거치며 점점 구체화를 한다는 점입니다.
글에서 제시한 전체 흐름은 아래처럼 정리돼요(문서 기준 CUDA 13.1):
– Python(CuTile API) → `.cutile` 바이트코드
– `cuda_tile` (MLIR dialect): 아키텍처 독립적인 고수준 타일/텐서 연산
– `nv_tileaa` (MLIR dialect): 메모리 참조가 명시되는 타일 로우레벨
– `nv_tileas` (MLIR dialect): 스케줄링/async 파이프라인/TCGEN 등 아키텍처 특화
– NVVM/LLVM IR → PTX → **SASS(sm_120 등 실제 명령)**
여기서 사용자 입장에서 눈여겨볼 도구가 `tileiras`예요. 글에서도 “TileIR용 ptxas 같은 존재”라고 표현하는데, 실제로 `.cutile`을 AOT(ahead-of-time)로 `cubin`까지 뽑습니다.
“`bash
tileiras –gpu-name sm_120 MoE.cutile -o moe.cubin
“`
## Running Example(MoE 커널): `gather / load / mma`가 어떻게 변하는가
글은 MoE 커널을 예제로, 특히 아래 3개 핵심 연산이 단계별로 어떻게 바뀌는지 추적해요.
– `ct.gather(A, indices)` → 간접 로드(토큰 인덱스로 모아 읽기)
– `ct.load(B, …)` → 직접 로드(연속 타일 로드)
– `ct.mma(a, b, acc)` → 텐서 코어 MMA(행렬 곱 누적)
단계별 매핑이 재밌는데요, 요지는 “이름이 바뀌는 게 아니라 **추상화가 벗겨진다**”예요.
– `cuda_tile`에서는 `load_view_tko`, `load_ptr_tko`, `mmaf` 같이 “의미 중심”이고
– `nv_tileaa`로 내려가면 `make_memref`, `tiled_load`, `dot`처럼 “주소/형상”이 더 노출돼요.
– `nv_tileas`에선 `async.pipeline.*`, `make_tiled_tma_desc`, `tcgen05.*` 같은 “하드웨어 실행 방식”이 본격 등장합니다.
즉, “내가 타일을 이렇게 쓰고 싶다” → “그럼 글로벌/공유/텐서 메모리를 이렇게 오가며” → “이 파이프라인과 배리어로 겹쳐 돌리고”로 바뀌는 흐름이에요.
[IMAGE2:A compilation pipeline diagram from Python to MLIR dialects (cuda_tile, nv_tileaa, nv_tileas) to LLVM/NVVM to PTX and SASS, with arrows and small code/IR snippets at each stage]
## 디버깅/튜닝 포인트: `–print-before-all`과 (비공식) 환경변수들
실제로 분석하려면 `tileiras` 옵션이 꽤 유용해요. 글에서 강조하는 게 `–print-before-all`인데, 이걸 켜면 **각 패스 전 LLVM IR 덤프**를 뽑아볼 수 있어요.
“`bash
tileiras –print-before-all –gpu-name=sm_120 MoE.cutile -o moe.cubin 2>&1
“`
또 흥미로운 건 **문서화되지 않은 환경변수**들이에요(향후 변경 가능):
– `TILEIR_ALWAYS_SWIZZLE`: 강제로 swizzle 모드 적용(메모리 레이아웃 최적화 방향 고정)
– `TILEIR_PREFER_TMA_FOR_LOAD_STORE`: load/store에 TMA(Tensor Memory Accelerator) 선호
– `TILEIR_DELAY_TMA_STORE_WAIT`: store wait을 늦춰 연산-메모리 오버랩을 노리는 최적화
실전 시나리오로는, MoE처럼 메모리 이동 비중이 큰 커널에서 “TMA를 강제로 태우면 이득인가?”를 실험해볼 수 있겠죠. 덤프를 통해 `cp.async`, `mbarrier`, `tcgen05` 계열로 어떻게 내려갔는지 확인하면서요.
## 마무리: “내 커널이 왜 빠르거나 느린지”를 IR에서 설명할 수 있게 돼요
TileIR의 매력은 단순히 새 DSL이 아니라, **Python→MLIR→LLVM→SASS로 내려가는 전 과정을 한 프레임워크에서 관찰**할 수 있다는 점이에요. 특히 `nv_tileas` 단계에서 async 파이프라인/더블버퍼링/배리어가 어떻게 구성되는지 보면, 성능이 “감”이 아니라 구조로 이해되기 시작합니다.
다음에 CuTile 샘플(MoE 포함)을 한 번 직접 `tileiras`로 빌드해보고, `–print-before-all`로 IR 덤프를 떠서 **`ct.mma`가 최종적으로 어떤 `HMMA`/`TCGEN05.MMA`로 귀결되는지**를追적해보세요. 그 과정 자체가 GPU 성능 최적화의 감각을 한 단계 올려줄 거예요.” />
## 디버깅/튜닝 포인트: `–print-before-all`과 (비공식) 환경변수들
실제로 분석하려면 `tileiras` 옵션이 꽤 유용해요. 글에서 강조하는 게 `–print-before-all`인데, 이걸 켜면 **각 패스 전 LLVM IR 덤프**를 뽑아볼 수 있어요.
“`bash
tileiras –print-before-all –gpu-name=sm_120 MoE.cutile -o moe.cubin 2>&1
“`
또 흥미로운 건 **문서화되지 않은 환경변수**들이에요(향후 변경 가능):
– `TILEIR_ALWAYS_SWIZZLE`: 강제로 swizzle 모드 적용(메모리 레이아웃 최적화 방향 고정)
– `TILEIR_PREFER_TMA_FOR_LOAD_STORE`: load/store에 TMA(Tensor Memory Accelerator) 선호
– `TILEIR_DELAY_TMA_STORE_WAIT`: store wait을 늦춰 연산-메모리 오버랩을 노리는 최적화
실전 시나리오로는, MoE처럼 메모리 이동 비중이 큰 커널에서 “TMA를 강제로 태우면 이득인가?”를 실험해볼 수 있겠죠. 덤프를 통해 `cp.async`, `mbarrier`, `tcgen05` 계열로 어떻게 내려갔는지 확인하면서요.
## 마무리: “내 커널이 왜 빠르거나 느린지”를 IR에서 설명할 수 있게 돼요
TileIR의 매력은 단순히 새 DSL이 아니라, **Python→MLIR→LLVM→SASS로 내려가는 전 과정을 한 프레임워크에서 관찰**할 수 있다는 점이에요. 특히 `nv_tileas` 단계에서 async 파이프라인/더블버퍼링/배리어가 어떻게 구성되는지 보면, 성능이 “감”이 아니라 구조로 이해되기 시작합니다.
다음에 CuTile 샘플(MoE 포함)을 한 번 직접 `tileiras`로 빌드해보고, `–print-before-all`로 IR 덤프를 떠서 **`ct.mma`가 최종적으로 어떤 `HMMA`/`TCGEN05.MMA`로 귀결되는지**를追적해보세요. 그 과정 자체가 GPU 성능 최적화의 감각을 한 단계 올려줄 거예요.
