Cris Cecka (NVIDIA · CUTLASS 핵심) 가 깐 CuTe — CUTLASS 3.x 의 기반이 되는 C++ template library. tensor 의 layout 을 1급 객체로 다루는 algebra 를 도입해 — multi-dim shape, stride, hierarchical mode 를 같은 frame 에서 합치고/나누고/coalesce 한다. Hopper 시대의 SOTA GPU kernel (Flash Attention 3, GEMM, GEMV) 의 공통 substrate. 본 페이지는 transcript 가 실패해 CUTLASS GitHub, NVIDIA 공식 문서, Cris Cecka 의 GTC talk 으로 재구성됐다.
SOTA GPU GEMM 또는 attention kernel 안 — 같은 tensor 가 동시에 여러 모양으로 보여야 한다. global memory 에서는 (M, K), shared memory 에서는 swizzled (M_block, K_block), register 에서는 thread 별로 분할된 tile, MMA instruction 에서는 architecture 별로 정의된 shape — 같은 데이터에 대한 다섯 layout. 이 다섯 사이의 변환을 손으로 짜면 인덱스 계산만으로 코드의 절반.
강의가 풀려는 두 질문.
본 노트는 transcript 실패로 — CUTLASS GitHub repo (github.com/NVIDIA/cutlass), CuTe 의 media/docs/cute/ 문서, Cris Cecka 의 GTC talk, Pradeep Ramani 의 CuTe blog 시리즈로 재구성.
CuTe 의 본질은 “layout 이 함수다” 라는 통찰. layout 은 (logical multi-dim coord) → (linear index) 의 함수. 그러면 두 layout 의 합성, 역, divide 같은 algebra 가 자연스럽게 정의된다. 이 algebra 위에서 — 같은 tensor 의 여러 view 가 함수 합성으로 표현된다.
CuTe 의 layout 은 (shape, stride) pair — shape 은 multi-dim 의 크기, stride 는 한 차원으로 한 칸 움직일 때 linear index 가 얼마나 변하는지. 같은 데이터를 다른 shape/stride 로 보면 다른 layout — 같은 메모리, 다른 view.
auto L = make_layout(make_shape(_4, _3), make_stride(_3, _1));
일반적 GPU kernel 코드는 — index 계산이 row*stride_a + col*stride_b 같이 hard-coded 흩어져 있음. CuTe 는 layout 을 객체로 두고 — 한 번 선언하면 같은 layout 을 여러 자리에서 재사용. 코드의 로직과 인덱스의 로직이 분리되어 가독성과 generic 성이 동시에 잡힘.
그리고 — layout 은 compile-time 또는 runtime. shape/stride 의 일부가 compile-time const 면 — 컴파일러가 그 부분을 unroll/inline. 일부가 runtime 이면 그 부분만 일반 코드. CuTe 의 type system 이 두 모드를 혼합 가능.
stride 의 의미가 “한 차원으로 한 칸 움직일 때 메모리 offset” 이라면 — 다양한 stride 가 다양한 view 를 만든다.
multi-dim shape 의 layout 은 — 차원별 (shape, stride) 의 tuple. 함수는 모든 차원의 contribution 의 합 → linear index.
흥미로운 사용 — stride 0 은 broadcast. logical 상으로 “row 가 4 개” 처럼 보이지만, 메모리는 row 1 줄만. 같은 데이터를 4번 “읽는” 효과. (4, 3) : (0, 1) — 첫 차원이 broadcast. 이걸 활용하면 — bias addition, position encoding, mask 적용 같은 패턴을 layout 안에서 표현 가능. 실제 메모리 절약.
CuTe 의 가장 큰 디자인 — hierarchical mode. shape 가 그냥 (M, N) 이 아니라 — 같은 차원이 ((tile_M, num_tiles_M), N) 같이 중첩. 이게 GPU 의 hierarchy (thread → warp → block → grid) 와 정확히 맞물린다.
예시 — matmul 의 M 차원 분해.
이 hierarchy 를 CuTe 의 hierarchical shape 로 표현 — Shape = (4, 16, 2, 32) : (1, 4, 64, 128). 각 mode 가 thread/warp/block/grid 의 분해. 한 layout 에 GPU 의 4-level hierarchy 가 그대로 박힘.
GPU 의 분해는 본질적으로 hierarchical — grid → block → warp → thread. CuTe 가 layout 의 shape 자체에 hierarchy 를 내장하면 — “이 thread 는 어떤 element 를 읽나” 를 layout 함수 적용 한 번으로 얻을 수 있음. 손으로 indexing 안 짠다.
// hierarchical shape — CuTe
// 4096 = 4 * 16 * 2 * 32
auto shape = make_shape(_4{}, _16{}, _2{}, _32{});
auto stride = make_stride(_1{}, _4{}, _64{}, _128{});
auto layout = make_layout(shape, stride);
// 한 thread 의 입장 — (t, w, w_idx, b)
// t = thread idx within warp (0..3)
// w_idx = thread group within warp (0..15)
// w = warp within block (0..1)
// b = block idx (0..31)
auto coord = make_coord(t, w_idx, w, b);
auto linear_idx = layout(coord);
// linear_idx 가 그 thread 가 다룰 element
hierarchical layout 의 또 다른 사용 — data 의 swizzling. shared memory 의 bank conflict 를 피하기 위해 data 를 permuted layout 으로 두기. 같은 algebra 안에서 swizzled view 를 표현 가능.
CuTe 의 핵심 — layout 위에 정의된 연산들. compose, divide, coalesce. 이 셋이 복잡한 layout 변환을 generic 하게 표현한다.
(M=4096) → ((BM=128), (M/BM=32)). block-tile 분해의 표준 연산. matmul 의 outer/inner 루프 분해 자리.
((4, 8)) 같이 nested 인데 stride 가 deterministic 하면 (32) 로 단순화. 컴파일 시점 최적화.
local_tile(layout, tile_shape, coord). 한 thread block 의 입장에서 자기 tile 만 보기.
local_partition(tile, thr_layout, thr_idx).
CuTe 의 GEMM kernel 은 거의 — local_tile 로 block 의 tile 을 잡고, local_partition 으로 thread 의 fragment 를 잡는 두 단계만으로 모든 indexing 처리. 손으로 thread idx 계산 안 함. 이게 CuTe 의 가장 큰 생산성 향상 자리.
// CUTLASS 3.x GEMM 의 hot loop (의사코드)
auto A_tile = local_tile(A, tile_shape, block_coord);
auto B_tile = local_tile(B, tile_shape, block_coord);
auto A_thr = local_partition(A_tile, thr_layout, thr_idx);
auto B_thr = local_partition(B_tile, thr_layout, thr_idx);
// thr 단위 fragment register
auto rA = make_fragment_like(A_thr);
auto rB = make_fragment_like(B_thr);
auto rC = make_fragment_like(C_part);
// load + MMA
copy(A_thr, rA);
copy(B_thr, rB);
gemm(rA, rB, rC); // MMA atom
위 코드의 인상 — thread idx 가 한 번도 직접 등장하지 않는다. 모든 indexing 이 layout 객체 위의 연산 (local_tile, local_partition) 으로 표현. 같은 코드가 다른 tile 사이즈 / 다른 thread 분배에 대해 generic.
이게 CUTLASS 2.x 와 결정적 차이. 2.x 는 manual indexing 으로 매 새 architecture 마다 코드를 거의 새로 짰음. 3.x (CuTe 위) 는 — layout 만 바꾸면 같은 코드가 새 architecture 에서 동작.
CuTe 의 “layout 은 함수” 이라는 통찰의 가장 큰 응용 — thread layout 도 같은 객체. (thread idx) → (어떤 element 를 다룰지) 의 매핑이 layout 함수.
local_partition(tile, thr_layout, thr_idx). 같은 algebra 로.
thread block 안의 thread 분배 — “(M, K) 를 (4, 8) thread 로 나누면 어떻게 되나” 의 문제. 이게 그냥 layout 의 divide. 그래서 — 새 thread 분배를 시도할 때 layout 한 줄만 바꿔서 시도 가능. coalesced read, bank conflict 회피, swizzling 모두 layout 변경으로 해결.
실제 사용 — autotuning 의 base. 같은 kernel 의 여러 thread 분배 후보 (4×8, 8×4, 16×2, 2×16) 를 각각 다른 thread layout 으로 두고 — 같은 source 가 컴파일됨. CUTLASS 의 GEMM 이 새 architecture 에 빠르게 적응하는 이유의 큰 부분.
CuTe 의 두 번째 핵심 — hardware instruction 자체를 “atom” 으로 추상. NVIDIA 의 MMA (Matrix Multiply Accumulate) 명령어, TMA (Tensor Memory Accelerator) copy, async copy 등이 모두 atom. 각 atom 은 자기만의 layout 을 강제.
같은 GEMM kernel 이 — Volta MMA, Ampere MMA, Hopper wgmma 위에서 동작. atom 만 바꾸면 됨. 이게 CUTLASS 3.x 의 cross-architecture 약속의 본질. 새 hardware 가 나오면 — 새 atom 정의 + 그에 맞는 layout. kernel 코드는 그대로.
Hopper 의 새 atom — TMA (Tensor Memory Accelerator). async copy 의 super-set. 한 명령어가 큰 multi-dim tile 의 copy 를 hardware 가 자동으로. CUTLASS 3.x 의 GEMM 이 H100 에서 SOTA 인 이유의 큰 부분.
TMA 의 가장 흥미로운 점 — descriptor 기반. CPU 에서 descriptor 를 만들어 GPU 에 전달, GPU 가 그것으로 copy 자동 수행. layout 정보가 descriptor 안에 들어감. CuTe 가 그 descriptor 를 layout 으로부터 자동 생성.
Flash Attention 3 (Tri Dao, 2024) 가 CUTLASS 3.x 위에 구축됐음. 즉 CuTe 가 base. Hopper 의 TMA + wgmma + async pipeline 을 모두 CuTe 의 layout algebra 위에서 표현. L050 의 attention 학습 시퀀스의 마지막 자리.
CUTLASS 2.x 는 — C++ template 의 deep nesting + 거의 manual 한 indexing. 새 architecture 마다 코드 거의 새로. CUTLASS 3.x 는 CuTe 위에서 — 같은 kernel 이 layout 변경만으로 새 hardware 적응. 실용적 차이가 큼.
CuTe 는 — 읽고 이해하기 쉽지 않음. C++ template 의 깊이 + layout algebra 의 추상이 합쳐져 — 처음 보면 의도가 안 잡힘. 그러나 한번 잡히면 — 같은 코드가 모든 GPU 에서 동작하는 magical 한 효과. 학습 비용 크지만 ROI 도 큼.
실용적 진입 시퀀스 — (1) CUTLASS examples 의 가장 작은 GEMM 부터 정독. (2) media/docs/cute/ 의 layout tutorial 따라가기. (3) 자기 손으로 layout 객체를 만들고 print 해보기 — print(layout) 가 동작. (4) 작은 GEMM 을 CuTe 로 짜보기. (5) FA3 코드를 한 줄씩 읽기.
cutlass/media/docs/cute/ — 0_layout, 1_tensor, 2_layout_algebra 순서. 가장 좋은 시작점.cutlass/examples/cute/ — vector add 부터 GEMM 까지 점진적 예제. 한 줄씩 따라가며 손으로.cutlass/examples/00_basic_gemm 부터. CUTLASS 3.x 의 GEMM 이 CuTe 위에서 어떻게 짜이는지의 reference.auto L = make_layout(make_shape(_4{}, _3{}), make_stride(_3{}, _1{})). print(L) 로 layout 의 모양 확인.본 노트의 모든 코드 스니펫과 layout 식은 — CUTLASS docs 와 example 의 일반적 사용 패턴 재구성. 실제 CuTe API 는 빠르게 발전 중 — 정확한 함수 이름과 type 은 CUTLASS 의 main branch 를 직접 확인. 그리고 CuTe 의 학습은 — 개별 함수 이해보다 “layout 이 함수다” 라는 통찰의 흡수가 본질. 작은 예제를 손으로 한참 만들어보는 게 가장 빠른 길.