Learning CUTLASS the hard way

CUTLASS 의 공식 자료는 양도 많고 깊이도 있다. 하지만 처음 들어온 사람은 — example 을 다 돌려보고 docs 를 다 읽어도 — “내가 짠 GEMM 한 줄이 동작하지 않을 때 어디서부터 무엇을 봐야 하는가” 가 안 잡힌다. 이 강의가 풀려는 문제가 정확히 이 자리.

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

1. CUTLASS 의 추상 hierarchy 를 어떻게 머리에 넣는가 — 한 번에 다 이해하지 않고, 위에서 아래로 한 layer 씩.
2. 어떤 자료를 어떤 순서로 읽으면 학습 곡선이 가장 가파르지 않은가 — 시간을 가장 적게 쓰는 path.

발표자의 자세는 practitioner-honest. 이 강의는 “CUTLASS 가 쉽다” 는 광고가 아니라 — “CUTLASS 가 어렵다는 사실 자체를 우회하지 않고, 그 어려움을 단계로 나누어 학습 가능한 chunk 로 만든다”. L029 Triton Internals 와 같은 자세의 follow-up.

CUTLASS 가 어려운 이유를 분해하면 — 서로 직교적인 두 차원의 어려움이 같은 코드 라인에 합쳐져 있다는 사실이 드러난다. 이 둘을 분리해서 따로따로 정복하는 게 강의의 첫 학습 전략.

그래서 학습 곡선은 — 머리 속에서는 한 곡선이지만 — 실제로는 두 곡선의 합이다. 한 쪽을 평평하게 만들고 다른 쪽을 올라가는 게 “the hard way” 의 진짜 의미.

CUTLASS 3.x (CuTe-based) 의 한 GEMM kernel 의 클래스 hierarchy 를 layer 로 풀면 다음 5 단. 위에서 아래로 — 각 layer 가 무엇을 결정하고 무엇을 다음 layer 에 위임하는지.

두 번째 달 — Level 2 (mainloop) 만. epilogue 는 그대로 두고 mainloop 안의 stages 와 tile size 를 바꿔본다. “성능이 어디서 만들어지는가” 의 첫 답이 여기서 나온다.

세 번째 달 — Level 3 (epilogue). bias / activation / scaling 의 fused 패턴을 직접 짜본다. 이 단계에서 처음으로 CuTe 의 layout을 만지게 된다 — accumulator → output tensor 의 mapping.

CUTLASS 2.x 에서는 layout 이 iterator pattern 으로 표현됐다. 3.x 부터는 layout 이 일급 객체 — Layout = (Shape, Stride). 이 한 줄을 정확히 이해하는 게 모든 mainloop / epilogue 코드의 출발점.

// 가장 단순한 row-major 4×2
using ShapeA  = Shape<_4, _2>;
using StrideA = Stride<_2, _1>;
using LayoutA = Layout<ShapeA, StrideA>;

// 좌표 (i,j) → index = 2*i + 1*j

// hierarchical — block 안에 sub-block
using Hier = Layout<Shape<_2, Shape<_2,_2>>,
                     Stride<_4, Stride<_2,_1>>>;

// 좌표 (i, (j0, j1)) → index = 4*i + 2*j0 + 1*j1

layout algebra — 자주 쓰는 5 연산

강의의 본론. 강의가 한 번에 정확히 따라간다고 약속한 “한 case” — Hopper (sm_90) FP16 GEMM. 이 한 케이스의 mainloop step-by-step 을 따라가면 다른 케이스 (FP8, sparse, FP4) 가 변형으로 보인다.

강의의 응용 파트. FlashAttention-3 의 reference 구현이 사실상 CUTLASS 위에 쓰여 있다. GEMM 의 mainloop 와 attention 의 mainloop 의 공통점/차이점을 정리하면 — 한 번 GEMM 을 이해하고 attention 으로 옮길 때 어디만 새로 봐야 하는지가 명확해진다.

FA3 가 H100 위에서 WGMMA + TMA + warp specialization 의 모든 기능을 한꺼번에 동원한다. 이 한 example 을 따라가면 — Hopper 시대의 모든 “fast kernel” 이 같은 패턴이라는 사실이 잡힌다. L086 — FlashAttention 3 와 L057 — CuTe 가 같은 자리의 깊이 보충.

CUTLASS 와 CuTe 의 자료는 너무 많고 산재. 어떤 자료를 어떤 순서로 읽으면 학습 곡선이 가장 가파르지 않은지 — 강의의 가장 실용적인 부분이자 강의 후 가장 자주 다시 보게 되는 자리.

• CuTe DSL (Python) — CUTLASS 의 layout algebra 를 Python 에서 직접. 빠른 prototyping. 단 — C++ 으로 한 번 통과한 뒤에 옮기는 게 좋다.
• CUTLASS 4 (Blackwell) — sm_100 의 새 명령들, FP4 / FP6 native, 새 TMA 변형. 큰 mainloop 변화. 학습한 H100 골격이 80% 는 옮아간다.
• QuTLASS — IST-DASLab 의 CUTLASS-기반 quantized BLAS. MXFP4 / NVFP4 의 GEMM 을 보고 싶다면. L102 Quartet v2 와 직접 연결.
• Triton 과의 비교 — 같은 GEMM 을 Triton 으로 짜고 CUTLASS 와 비교. 두 stack 의 trade-off (제어 vs 생산성) 가 손에 잡힌다. L029 와 같은 시리즈.
• custom mainloop — 자기 모델 / 자기 dtype 의 mainloop 를 직접. 이 단계가 진정한 “CUTLASS practitioner” 의 출발점.

손에 새기기 — 12주 path

1. 주 1 — CUTLASS clone, build env (CUDA 12.x, sm_90 capable GPU), 00_basic_gemm example 을 그대로 build & run. 측정만.
2. 주 2 — CuTe quickstart 와 01_layout. 좌표 → index 의 함수 식 직접 손으로. 5 개의 layout 예시를 실제 메모리 mapping 으로 그림.
3. 주 3 — 02_layout_algebra. logical_divide 와 logical_product 의 작은 예시들 (tile of 4×4 in 16×16) 직접 풀어보기.
4. 주 4 — Colfax matrix transpose tutorial. 세 단계 (naïve, SMEM, swizzle) 를 모두 build 해 측정. swizzle 이 왜 bank conflict 를 푸는지 손에 잡힐 때까지.
5. 주 5 — Colfax TMA tutorial. cp.async.bulk.tensor 의 mode (load, store, multicast) 한 번씩 실행.
6. 주 6 — CUTLASS examples/cute/tutorial 의 sgemm path. 한 mainloop 의 한 stage 의 모든 step 을 손으로 그려본다.
7. 주 7 — Hopper GEMM example. tile size 한 번, stages 한 번 변경. 측정 비교.
8. 주 8 — epilogue swap. activation (relu) 추가, bias 추가 — 같은 mainloop 위에 다른 epilogue 만.
9. 주 9 — FlashAttention example build. mainloop 두 GEMM 부분만 분리해 GEMM 처럼 본다.
10. 주 10 — online softmax 의 의미 손으로. running max + running sum 의 invariant.
11. 주 11 — NCU 로 mainloop profile. async pipeline 의 stage-level 시각으로 stall 추적.
12. 주 12 — 자기 모델의 한 GEMM 을 CUTLASS 로 짜본다. 첫 “custom” mainloop. 이 시점이 강의가 약속한 “practitioner” 진입점.

• 강의에서 직접 보여준 예제 — Kapil 이 어떤 specific kernel 을 라이브로 풀어가며 보여줬는지 자막 실패로 확인 안 됨. 본인 GitHub repo 에 동반 코드가 있을 가능성. (확인 필요)
• CUTLASS 4 / Blackwell 의 변화점이 이 강의의 어디까지 다뤄졌는지 — 발표 시점 (2025) 에 따라 다름. (확인 필요)
• CuTe DSL (Python) 에 대한 발표자의 입장 — “Python 으로 옮긴 학습 path” 가 권고되었는지 / 거부되었는지. (확인 필요)
• 사용자 정의 epilogue 의 fusion 패턴 — softmax / layernorm 같은 더 복잡한 fused epilogue 가 다뤄졌는지. (확인 필요)
• WGMMA 의 swizzle pattern — Hopper 의 swizzle 이 layout algebra 로 어떻게 표현되는지의 구체 — 강의에서 깊게 들어갔을 가능성 큼. (확인 필요)
• Triton 과의 직접 throughput 비교 — 같은 GEMM 의 Triton vs CUTLASS 측정이 발표 안에서 시연됐는지. (확인 필요)