gpumode · 강의 아카이브
《GPU Mode》 L057 2024 High priority transcript · failed

CuTe — CUTLASS Tensor의 layout algebra

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 으로 재구성됐다.

CuTe CUTLASS 3.x layout algebra hierarchical mode shape · stride compose · divide · coalesce MMA atom Hopper TMA
C
Speaker
Cris Cecka
NVIDIA · CUTLASS · CuTe 핵심 디자이너
강의 번호
L057
스피커
Cris Cecka
Transcript
failed · 본 노트는 재구성
학습 우선순위
High · Hopper-시대 substrate
§ 01강의가 풀려는 문제· why this lecture exists

“같은 tensor 가 다섯 가지 모양으로 보여야 하는” GPU kernel 의 자리

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. 이 다섯 사이의 변환을 손으로 짜면 인덱스 계산만으로 코드의 절반.

강의가 풀려는 두 질문.

  1. layout 자체를 1급 객체로 다루는 algebra 가 있다면 — 같은 tensor 의 여러 view 사이 변환을 어떻게 generic 하게 표현할까?
  2. 그 algebra 가 GPU kernel 작성을 어떻게 단순화하는가 — 특히 Hopper 의 새 hardware (TMA, async, MMA) 와 결합될 때.

본 노트는 transcript 실패로 — CUTLASS GitHub repo (github.com/NVIDIA/cutlass), CuTe 의 media/docs/cute/ 문서, Cris Cecka 의 GTC talk, Pradeep Ramani 의 CuTe blog 시리즈로 재구성.

강의의 frame

CuTe 의 본질은 “layout 이 함수다” 라는 통찰. layout 은 (logical multi-dim coord) → (linear index) 의 함수. 그러면 두 layout 의 합성, 역, divide 같은 algebra 가 자연스럽게 정의된다. 이 algebra 위에서 — 같은 tensor 의 여러 view 가 함수 합성으로 표현된다.

“CuTe 는 layout 을 첫 번째 객체로 승격한 algebra 다 — 그리고 그 algebra 가 GPU kernel 의 거의 모든 인덱스 연산을 generic 하게 만든다.” 학습 노트 · 재구성
§ 02layout 의 의미· shape × stride

layout = (shape, stride) — 그리고 그 자체로 함수

CuTe 의 layout 은 (shape, stride) pair — shape 은 multi-dim 의 크기, stride 는 한 차원으로 한 칸 움직일 때 linear index 가 얼마나 변하는지. 같은 데이터를 다른 shape/stride 로 보면 다른 layout — 같은 메모리, 다른 view.

layout Layout<Shape, Stride> = function (logical coord) → (linear index). shape 은 그 logical 공간의 크기, stride 는 매 차원의 곱셈 계수.
예 1 — row-major Shape = (4, 3), Stride = (3, 1). idx(i, j) = i*3 + j*1. logical (1,2) → linear 5.
예 2 — column-major Shape = (4, 3), Stride = (1, 4). idx(i, j) = i*1 + j*4. logical (1,2) → linear 9.
같은 데이터, 다른 layout 12 개 element 의 같은 메모리. row-major 와 column-major 는 다른 layout — 같은 logical (i,j) 가 다른 linear index 로 매핑.
layout 자체를 변수로 CuTe 에서는 layout 을 변수에 담고 함수에 전달 가능. 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 이 두 모드를 혼합 가능.

§ 03stride 와 shape 의 결합· multi-dim 매핑

stride 가 음수, 0, 큰 값일 때 — 어떤 view 가 만들어지는가

stride 의 의미가 “한 차원으로 한 칸 움직일 때 메모리 offset” 이라면 — 다양한 stride 가 다양한 view 를 만든다.

  • stride = 1 — 그 차원이 contiguous. 가장 빠른 차원.
  • stride = N — 그 차원이 N step 마다 다음 element. row-major matrix 의 row 차원.
  • stride = 0 — 그 차원이 broadcast. 같은 메모리 element 가 여러 logical 위치에 보임.
  • stride 음수 — 거꾸로. flip 같은 view.
  • stride = M*N — outer 차원. matrix 의 다음 “sheet”.

multi-dim shape 의 layout 은 — 차원별 (shape, stride) 의 tuple. 함수는 모든 차원의 contribution 의 합 → linear index.

FIG · 같은 4×3 데이터의 두 layoutrow vs column
0
1
2
·
·
·
·
·
3
4
5
·
·
·
·
·
6
7
8
·
·
·
·
·
9
10
11
·
·
·
·
·
row-major (4, 3) : (3, 1). 각 row 가 contiguous. logical (i,j) → idx = 3i + j.
stride 0 의 broadcast

흥미로운 사용 — stride 0 은 broadcast. logical 상으로 “row 가 4 개” 처럼 보이지만, 메모리는 row 1 줄만. 같은 데이터를 4번 “읽는” 효과. (4, 3) : (0, 1) — 첫 차원이 broadcast. 이걸 활용하면 — bias addition, position encoding, mask 적용 같은 패턴을 layout 안에서 표현 가능. 실제 메모리 절약.

§ 04hierarchical mode· 중첩된 layout

shape 안의 shape — block 과 thread 가 같은 framework 에서

CuTe 의 가장 큰 디자인 — hierarchical mode. shape 가 그냥 (M, N) 이 아니라 — 같은 차원이 ((tile_M, num_tiles_M), N) 같이 중첩. 이게 GPU 의 hierarchy (thread → warp → block → grid) 와 정확히 맞물린다.

예시 — matmul 의 M 차원 분해.

  • 전체 M = 4096
  • 각 thread block 이 BM = 128 줄 처리 → 32 blocks
  • 각 block 안에서 각 warp 가 wm = 64 줄 → 2 warps
  • 각 warp 안에서 각 thread 가 tm = 4 줄 → 16 threads/dim

이 hierarchy 를 CuTe 의 hierarchical shape 로 표현 — Shape = (4, 16, 2, 32) : (1, 4, 64, 128). 각 mode 가 thread/warp/block/grid 의 분해. 한 layout 에 GPU 의 4-level hierarchy 가 그대로 박힘.

왜 hierarchical 이 자연스러운가

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 를 표현 가능.

“hierarchical shape 가 GPU 의 4-level 분해를 한 layout 안에 모은다 — kernel 코드의 indexing 이 한 줄로 줄어든다.” 학습 노트 · 재구성
§ 05layout 연산· compose · divide · coalesce

layout 위의 algebra — 함수 합성으로 새 view 만들기

CuTe 의 핵심 — layout 위에 정의된 연산들. compose, divide, coalesce. 이 셋이 복잡한 layout 변환을 generic 하게 표현한다.

compose 두 layout 의 합성 — A ∘ B 는 “A 의 출력을 B 가 다시 indexing”. 같은 데이터의 두 view 를 결합. 예: row-major data 를 swizzled view 로 보고 싶으면 row-major ∘ swizzle.
divide shape 를 nested 로 분해(M=4096) → ((BM=128), (M/BM=32)). block-tile 분해의 표준 연산. matmul 의 outer/inner 루프 분해 자리.
coalesce nested layout 의 단순화((4, 8)) 같이 nested 인데 stride 가 deterministic 하면 (32) 로 단순화. 컴파일 시점 최적화.
tile 큰 layout 에서 하나의 tile 추출local_tile(layout, tile_shape, coord). 한 thread block 의 입장에서 자기 tile 만 보기.
partition tile 안에서 thread 별 분할 — 같은 tile 을 N 개 thread 가 나누어 읽기. local_partition(tile, thr_layout, thr_idx).
inverse layout 의 역 — (linear idx) → (logical coord). drop 된 차원 처리 등 미세한 디테일 있음.
divide 와 partition 의 사용 패턴

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 에서 동작.

§ 06GPU thread mapping· layout 으로 thread → data

같은 algebra 가 thread 의 indexing 도 다룬다

CuTe 의 “layout 은 함수” 이라는 통찰의 가장 큰 응용 — thread layout 도 같은 객체. (thread idx) → (어떤 element 를 다룰지) 의 매핑이 layout 함수.

data layout tile 의 shape 과 stride. (M_tile, K_tile) : (stride_m, stride_k).
thread layout block 안의 thread 분배. (4, 8) : (8, 1) 같이. 같은 객체 type — 그냥 의미가 다름.
partition 두 layout 의 결합 — local_partition(tile, thr_layout, thr_idx). 같은 algebra 로.
결과 thread 가 다룰 element 들의 fragment view. compile-time 에 size 결정.
왜 thread 도 layout 인가

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 에 빠르게 적응하는 이유의 큰 부분.

“thread idx 는 단순한 정수가 아니다 — 그것도 layout 의 입력. 그 통찰이 GPU kernel 작성을 수학적 객체 위로 옮긴다.” 학습 노트 · 재구성
§ 07Tensor 위 연산· MMA · copy atom

layout 이 hardware instruction 의 모양에 맞아야 한다

CuTe 의 두 번째 핵심 — hardware instruction 자체를 “atom” 으로 추상. NVIDIA 의 MMA (Matrix Multiply Accumulate) 명령어, TMA (Tensor Memory Accelerator) copy, async copy 등이 모두 atom. 각 atom 은 자기만의 layout 을 강제.

MMA atom

  • SM_75: 16×8×8 fp16
  • SM_80: 16×8×16 fp16, 16×16×16 bf16
  • SM_90 (Hopper): 64×N×16 등 — wgmma async
  • 각 atom 은 input/output layout 을 강제
  • register fragment 의 layout 이 그 atom 에 맞아야 함
  • 코드는 atom 을 변수로 받음

copy atom

  • cp.async (SM_80+) — async global → shared
  • TMA (SM_90+) — descriptor 기반 비동기 copy
  • 각 copy atom 의 throughput / alignment 요구가 다름
  • layout 이 alignment 와 vector width 에 맞아야 함
  • 코드는 atom 만 바꾸면 다른 hardware 에서 컴파일
atom 의 의미 — generic kernel

같은 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 으로부터 자동 생성.

FA3 와 CuTe

Flash Attention 3 (Tri Dao, 2024) 가 CUTLASS 3.x 위에 구축됐음. 즉 CuTe 가 base. Hopper 의 TMA + wgmma + async pipeline 을 모두 CuTe 의 layout algebra 위에서 표현. L050 의 attention 학습 시퀀스의 마지막 자리.

§ 08CUTLASS 위 사용· SOTA GEMM 의 substrate

CUTLASS 2.x → 3.x 의 변화 — manual 에서 layout-driven 으로

CUTLASS 2.x 는 — C++ template 의 deep nesting + 거의 manual 한 indexing. 새 architecture 마다 코드 거의 새로. CUTLASS 3.x 는 CuTe 위에서 — 같은 kernel 이 layout 변경만으로 새 hardware 적응. 실용적 차이가 큼.

CUTLASS 2.x deep template, manual thread idx, architecture-specific tiling. 새 GPU 마다 거의 새 코드. SM_75 → SM_80 → SM_90 마다 별도 path.
CUTLASS 3.x CuTe 위. layout 객체로 모든 indexing. 새 hardware 는 새 atom + 새 layout. kernel 코드는 generic. 같은 source 가 multi-arch.
production 사용 cuBLAS 의 일부 GEMM, FlashAttention 3, vLLM 의 일부 kernel — CUTLASS 3.x 위에. 점점 늘고 있음.
학습 곡선의 현실

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 코드를 한 줄씩 읽기.

“CuTe 는 GPU kernel 의 lingua franca 가 되어가는 중. Hopper 시대 의 SOTA kernel 거의 모두가 CuTe 위에서 짜인다.” 학습 노트 · 재구성
§ 09학습 자료· examples · papers

CuTe 를 처음 배우는 사람을 위한 자료 목록

CuTe docs
cutlass/media/docs/cute/ — 0_layout, 1_tensor, 2_layout_algebra 순서. 가장 좋은 시작점.
CuTe examples
cutlass/examples/cute/ — vector add 부터 GEMM 까지 점진적 예제. 한 줄씩 따라가며 손으로.
CUTLASS GEMM examples
cutlass/examples/00_basic_gemm 부터. CUTLASS 3.x 의 GEMM 이 CuTe 위에서 어떻게 짜이는지의 reference.
Cris Cecka GTC talks
2023, 2024 GTC 의 CUTLASS / CuTe 발표. layout algebra 의 동기와 디자인.
Pradeep Ramani blog
CuTe layout algebra 의 직관적 설명. NVIDIA blog series.
FlashAttention 3 paper + code
Tri Dao 의 FA3 가 CuTe 위에서 구현. 실전 사용 사례의 reference.
Hopper architecture white paper
TMA, wgmma 의 hardware 디테일. CuTe 의 atom 이 무엇을 추상하는지의 base.
PyTorch GEMM dispatch
PyTorch 의 일부 GEMM kernel 이 CUTLASS 3.x → CuTe 위에서 동작. 어떻게 dispatch 되는지 추적 가능.
§ 10기억할 메모· key takeaways

다시 열었을 때 5분 안에 손에 잡혀야 할 것

layout 의 정의
(shape, stride) 의 pair. (logical coord) → (linear index) 의 함수.
hierarchical mode
shape 를 nested 로. ((4, 16, 2, 32)) 같은 shape 가 GPU 의 grid/block/warp/thread 분해와 일치.
core 연산
compose, divide, coalesce, tile, partition. 같은 algebra 로 모든 layout 변환.
thread layout
thread 분배도 layout 객체. data layout 과 같은 algebra 로 partition.
atom
hardware instruction (MMA, TMA) 을 추상. atom 마다 input/output layout 강제.
CUTLASS 2 → 3
manual indexing → layout-driven. 같은 kernel 이 multi-arch 에서 동작.
FA3 의 base
Flash Attention 3 가 CUTLASS 3.x = CuTe 위에서. Hopper SOTA kernel 의 substrate.
학습 진입
CUTLASS examples 의 vector add → GEMM → 자기 손으로 layout 만들고 print → FA3 코드 읽기.

손에 새기기 — 실습 시퀀스

  1. CuTe layout 만들고 printauto L = make_layout(make_shape(_4{}, _3{}), make_stride(_3{}, _1{})). print(L) 로 layout 의 모양 확인.
  2. 같은 데이터의 두 layout — row-major 와 column-major 를 둘 다 만들고 — 같은 logical (i,j) 에 대해 다른 linear idx 가 나오는지 확인.
  3. hierarchical shape — (4, 8) 의 single-mode 와 ((2, 2), 8) 의 hierarchical 두 가지로 — coalesce 후 같은 layout 인지 확인.
  4. compose · divide — 작은 layout 두 개를 compose 했을 때의 결과를 print 로 확인. divide 의 결과도 마찬가지로.
  5. local_tile · local_partition — 큰 tensor 에서 한 thread block 의 tile 을 잡고, 그 안의 한 thread 의 fragment 를 잡는 패턴.
  6. CUTLASS 의 작은 GEMM — examples/00_basic_gemm 을 빌드하고 — 코드를 한 줄씩 주석.
  7. thread layout 변경 실험 — 같은 GEMM 을 다른 thread layout (4×8 vs 8×4 vs 16×2) 으로 — 결과 동일한지, 성능 차이는 어떤지.
  8. FA3 코드의 한 forward block 정독 — Tri Dao 의 FA3 코드 (CUTLASS 위) 의 한 함수를 한 줄씩 — CuTe 의 어느 도구가 어디 쓰이는지.
§ 11다른 강의로 이어지는 길· connections

같은 자리를 다른 각도에서 다루는 강의들

§ 12열린 질문· open questions

원본 자막 실패로 비워둔 자리들

검증 메모

본 노트의 모든 코드 스니펫과 layout 식은 — CUTLASS docs 와 example 의 일반적 사용 패턴 재구성. 실제 CuTe API 는 빠르게 발전 중 — 정확한 함수 이름과 type 은 CUTLASS 의 main branch 를 직접 확인. 그리고 CuTe 의 학습은 — 개별 함수 이해보다 “layout 이 함수다” 라는 통찰의 흡수가 본질. 작은 예제를 손으로 한참 만들어보는 게 가장 빠른 길.

← Lecture 056 Kernel Benchmarking Tales Lecture 058 → 다음 강의