제3부 딥러닝 라이브러리

3.4 CUTLASS I — 계층·GEMM 구성

3.4.1 이 장에서 배우는 것 / 선수 지식

cuBLAS(3.1)·cuDNN(3.2~3.3)은 닫힌 라이브러리다 — 호출은 하지만 안을 바꾸지 못한다. 라이브러리에 없는 정밀도 조합, 라이브러리가 융합하지 않는 에필로그, 비표준 형상·레이아웃이 필요해지는 순간 한계에 부딪힌다. CUTLASS(CUDA Templates for Linear Algebra Subroutines) 는 그 GEMM·conv 커널을 C++ 템플릿으로 직접 조립하게 해 주는 오픈소스 라이브러리다. 이 장은 CUTLASS가 GEMM을 GPU 하드웨어 계층에 그대로 포갠 타일 계층으로 분해하는 구조를 익힌다. 레이아웃 대수(CuTe)와 실습은 3.5로 잇는다.

배우는 것:

• CUTLASS가 GEMM을 device → kernel → threadblock → warp → thread(MMA 명령) 의 5계층 타일로 분해하는 구조와, 그것이 왜 GPU 메모리 계층(1.3~1.8)과 1:1로 대응하는지.
• threadblock tile · warp tile · MMA tile 크기가 점유율·공유메모리·텐서코어 활용에 어떻게 묶이는지(1.6·1.11 연결).
• 소프트웨어 파이프라이닝(multistage) 으로 글로벌 로드 지연을 수학(MMA)으로 은닉하는 원리.
• device 수준 GEMM API 골격을 작성하고, CUTLASS 2.x(고정 계층) vs 3.x(CuTe·collective) 의 갈림을 안다.

선수 지식: 1.3~1.8(칩·SM·온칩 메모리 계층 — CUTLASS 타일이 이 계층에 매핑됨), 1.11~1.12(텐서코어 MMA — CUTLASS의 가장 안쪽 타일), 2.7(타일 행렬곱 — 직접 작성한 그 구조의 production판), 1.6(점유율·limiter), 3.1(cuBLAS — CUTLASS가 대체·보완하는 대상). 환경은 0.3. CUTLASS는 헤더-온리 C++ 템플릿으로 버전마다 API·디폴트가 크게 바뀐다 — 코드는 개략 골격이며 정확한 타입·이름은 설치 버전(include/cutlass/)·예제로 확인.

3.4.2 왜 CUTLASS인가 — 닫힌 라이브러리의 벽

cuBLAS는 "GEMM 전체"를 세계 최고 수준으로 돌려준다. 그런데 실무는 자주 순수 GEMM이 아니다:

• 비표준 에필로그. cuBLASLt 에필로그(3.1)는 bias·GELU·ReLU 같은 정해진 꼬리만 융합한다. "GEMM 결과에 다른 텐서를 element-wise 곱하고, 그 결과로 또 reduction" 같은 커스텀 후처리는 라이브러리에 없다 → 별도 커널 → HBM 왕복.
• 비표준 정밀도·혼합. "FP8 입력 · FP16 누적 · 출력은 INT8 양자화" 같은 조합이 라이브러리 디폴트에 없을 수 있다.
• 융합된 형상. back-to-back GEMM(첫 GEMM 결과를 HBM에 안 쓰고 바로 둘째 GEMM 입력으로), grouped GEMM(크기가 제각각인 수백 개 작은 GEMM을 한 런치로 — MoE), conv를 implicit GEMM으로.
• 연구·커널 저작. FlashAttention(3.6)·MoE·새 양자화 커널은 전부 "GEMM 골격에 커스텀 수학을 끼운" 것이다. 그 골격을 맨바닥부터 작성하면 텐서코어·파이프라이닝·뱅크 회피를 다시 발명해야 한다.

CUTLASS의 자리: cuBLAS를 못 이기는 곳에서 cuBLAS를 쓰고(순수 대형 GEMM은 cuBLAS가 보통 최고이거나 동급), 라이브러리가 표현조차 못 하는 융합·형상·정밀도가 필요할 때 CUTLASS로 내려간다. cuBLAS 자체도 내부적으로 CUTLASS 계열 커널을 다수 포함한다(같은 계층 사상). 즉 CUTLASS는 "cuBLAS가 쓰는 그 빌딩블록을 직접 만지는" 라이브러리다.

3.4.3 핵심 아이디어 — GEMM을 하드웨어 계층에 포개기

CUTLASS의 단 하나의 큰 생각: GEMM $C = A\cdot B$의 타일 분해를 GPU 실행·메모리 계층에 1:1로 포갠다. 2.7에서 직접 구현한 타일링은 사실 한 층(threadblock + 공유메모리)이었다. CUTLASS는 그것을 다섯 층으로 완성한다.

읽는 법(위→아래로 "분할한다", 아래→위로 "쌓는다"):

• Device GEMM: 전체 문제 $M\times N\times K$. 출력 $C(M\times N)$를 threadblock tile 크기(예 $128\times128$)로 격자 분할 → 각 타일 하나를 한 threadblock에 배정(그리드). K 차원은 타일 안에서 루프.
• Threadblock tile: 한 블록이 맡은 $\mathrm{BM}\times\mathrm{BN}$ 출력 조각. K를 BK씩(예 32·64) 끊어, A의 $\mathrm{BM}\times\mathrm{BK}$ · B의 $\mathrm{BK}\times\mathrm{BN}$ 조각을 공유메모리로 적재하고 그 위에서 곱-누산. 2.7의 그 구조.
• Warp tile: 블록 안의 각 warp가 $\mathrm{BM}\times\mathrm{BN}$의 더 작은 부분(예 $64\times64$)을 맡는다. 입력 조각을 레지스터로 적재해 거기서 MMA를 반복.
• MMA tile (instruction): warp가 텐서코어 명령 한 번에 처리하는 최소 행렬곱(예 m16n8k16, 1.11). 가장 안쪽 루프가 이 명령들을 배치해 warp tile을 채운다.

이 사상이 핵심인 이유: 각 계층의 타일 크기가 곧 그 계층 자원의 사용량이다. threadblock tile↑ → 공유메모리·점유율 압박(1.6). warp tile → 레지스터 수. MMA tile → 어떤 텐서코어 명령(세대·정밀도 의존, 1.11~1.12). "GEMM 튜닝"은 결국 이 세 타일 크기와 파이프라인 단계 수를 고르는 일이다.

3.4.4 데이터플로 — 공유메모리 스테이징과 소프트웨어 파이프라이닝

타일을 분할하는 것만으로는 빠르지 않다. 핵심 성능 트릭은 글로벌 로드 지연(수백 사이클, 1.5)을 텐서코어 연산으로 은닉하는 것 — 소프트웨어 파이프라이닝이다.

순진한 타일 루프는 "로드 → 동기화 → 계산 → 로드 → …"로 로드하는 동안 텐서코어가 유휴 상태다:

순진(non-pipelined) 루프에서 텐서코어는 load·sync 구간에는 idle, mma 구간에만 busy다.

CUTLASS는 여러 K-스테이지를 미리 prefetch해 둔다(multistage, Stages=N). 다음(들) K 타일을 공유메모리로 비동기 로드(cp.async, Ampere+)하는 동안 현재 타일로 MMA를 실행한다:

파이프라인이 채워지면 텐서코어는 거의 항상 busy 상태를 유지한다.

• cp.async(Ampere+): 글로벌→공유 복사를 레지스터를 거치지 않고 비동기로. 로드가 백그라운드로 진행돼 MMA와 겹친다.
• ldmatrix: 공유메모리의 타일을 텐서코어 MMA가 요구하는 워프 단편 레이아웃으로 레지스터에 적재하는 전용 명령.
• Stages: 버퍼 단계 수. 많을수록 지연을 잘 가리지만 공유메모리를 N배 먹어 점유율↓. 보통 2~4. Hopper(4.2~4.3)는 TMA + warp-specialized 파이프라인으로 이 구조가 또 바뀐다.

이것이 2.7의 "공유메모리 타일 + __syncthreads()"를 production 수준으로 끌어올린 정체다: 동기 이중루프가 아니라 다단계 비동기 prefetch. 직접 작성하기 매우 까다로운 이 파이프라인을 CUTLASS가 템플릿으로 제공한다.

3.4.5 워크스루 A — 타일 계층 수계산

코드 전에 숫자로 계층을 파악해야 한다. M=N=K=4096 FP16 GEMM을 $\mathrm{BM}\times\mathrm{BN}\times\mathrm{BK} = 128\times128\times32$, warp tile $64\times64$, MMA m16n8k16로 분해한다.

1) 그리드(device 계층). 출력 $4096\times4096$을 $128\times128$ 타일로:

• 타일 수 = (4096/128) × (4096/128) = 32 × 32 = 1024개 출력 타일 = 1024 threadblock.
• 각 블록의 K 루프 횟수 = 4096 / 32 = 128번 (K-스테이지 반복).

2) 블록 안 warp 배치(threadblock 계층). $128\times128$을 warp tile $64\times64$로 나누면 2×2 = 4 warp → 블록당 4 warp × 32 = 128 스레드. (실 구성은 warp 배열을 MxN으로 다양하게 설정한다.)

3) warp 안 MMA(warp 계층). warp tile $64\times64$를 MMA $16\times8$로 채우려면 (64/16) × (64/8) = 4 × 8 = 32개 MMA 타일. K는 MMA가 k16이라 BK=32면 K방향 2번. → warp 하나가 K-스테이지마다 32 × 2 = 64회 텐서코어 명령.

4) 공유메모리 한 스테이지. A조각 $128\times32$ + B조각 $32\times128$, FP16(2B):

• A: $128\times32\times2 = 8192 \,\mathrm{B}$, B: $32\times128\times2 = 8192 \,\mathrm{B}$ → 스테이지당 16 KB.
• Stages=3이면 48 KB. SM 공유메모리 한도(세대별 ~100·228 KB 등, 시점 의존·확인 필요)와 비교해 블록당 몇 개 올라가는지 = 점유율(1.6). 48KB면 100KB SM에 블록 2개 → 점유율 제한 가능.

5) 산술 강도 점검(1.10). 이 GEMM의 연산-바이트로 메모리 바운드인지 확인: 총 FLOP $2\cdot M\cdot N\cdot K = 2\cdot4096^3 \approx1.37\mathrm{e}11$. 이만한 크기면 텐서코어 연산 바운드 영역 — 타일을 키워 데이터 재사용을 높일수록 좋다(단 공유메모리·레지스터·점유율과 줄다리기).

체감 포인트: 타일을 키우면(128→256) 재사용↑·HBM 트래픽↓이지만 공유메모리·레지스터를 더 먹어 점유율↓·작은 문제에서 타일 낭비(tile quantization). 작게 잡으면 반대. "정답 타일"은 형상·정밀도·세대마다 달라서, 라이브러리가 여러 타일 구성을 컴파일해 두고 형상별로 고르는(autotuning) 이유가 여기 있다.

3.4.6 워크스루 B — device GEMM 골격 (CUTLASS 2.x 스타일)

CUTLASS 2.x의 device-level API는 타일 계층을 전부 템플릿 파라미터로 받는다. 골격(정확한 타입·디폴트는 설치 버전·examples/로 확인):

#include <cutlass/gemm/device/gemm.h>
using namespace cutlass;

// A(M×K) row-major, B(K×N) col-major, C(M×N) row-major, FP16 입력·FP32 누적
using Gemm = cutlass::gemm::device::Gemm<
    half_t, layout::RowMajor,        // A: element, layout
    half_t, layout::ColumnMajor,     // B
    half_t, layout::RowMajor,        // C
    float,                            // ElementAccumulator (FP32 누적, 1.13)
    arch::OpClassTensorOp,            // 텐서코어 사용(아니면 OpClassSimt)
    arch::Sm80,                       // 타깃 아키텍처(Ampere). 세대마다 바꿈
    gemm::GemmShape<128,128,32>,      // ★ threadblock tile (BM,BN,BK)
    gemm::GemmShape<64,64,32>,        // ★ warp tile
    gemm::GemmShape<16,8,16>,         // ★ MMA instruction tile (세대·정밀도 의존)
    epilogue::thread::LinearCombination<half_t,8,float,float>, // 에필로그: alpha*AB+beta*C
    gemm::threadblock::GemmIdentityThreadblockSwizzle<>,        // 타일→블록 매핑(L2 지역성)
    3 /*Stages*/ >;                   // ★ 파이프라인 단계 수

Gemm gemm_op;
Gemm::Arguments args(
    {M, N, K},                        // 문제 크기
    {dA, lda}, {dB, ldb}, {dC, ldc}, {dD, ldd}, // 텐서·leading dim
    {alpha, beta});                   // 에필로그 스칼라
size_t ws = Gemm::get_workspace_size(args);
uint8_t* workspace; cudaMalloc(&workspace, ws);

cutlass::Status st = gemm_op.initialize(args, workspace); // 디스크립터 준비
st = gemm_op();                                            // 커널 런치
// st == cutlass::Status::kSuccess 확인

읽는 법:

• GemmShape<128,128,32> / <64,64,32> / <16,8,16> — 이 세 줄이 3.4.3의 threadblock·warp·MMA tile. 바꾸면 다른 커널이 컴파일된다(런타임 옵션이 아니라 컴파일 타임 인스턴스화).
• arch::Sm80 — 타깃 SM. 세대를 바꾸면 사용 가능한 MMA 명령·cp.async 유무가 달라진다(2.12 아키텍처 타깃).
• epilogue::...::LinearCombination — alpha·(A·B) + beta·C. 이 자리에 커스텀 에필로그를 삽입해 융합한다(3.5·실전).
• Stages=3 — 3.4.4의 파이프라인 깊이.

함정 미리보기: 이 템플릿 조합이 GPU가 지원하지 못하는 것(예: Sm70에 Ampere 전용 MMA)이면 컴파일은 되어도 런타임 kErrorNotSupported거나 빌드 실패. gemm_op.can_implement(args)로 사전 점검. 또 타일·정밀도 조합마다 별도 코드가 생겨 컴파일 시간·바이너리 크기가 폭발(3.4.8).

3.4.7 워크스루 C — CUTLASS Profiler로 형상 튜닝

코드를 작성하지 않고도 "내 형상에 어떤 타일이 빠른가"를 CUTLASS가 동봉한 프로파일러로 실측할 수 있다(빌드 시 생성, 정확한 인자는 버전 의존):

# CUTLASS 빌드(개략) — 특정 아키텍처만 켜 컴파일 폭발 억제
cmake .. -DCUTLASS_NVCC_ARCHS=80 -DCUTLASS_LIBRARY_KERNELS=cutlass_tensorop_*gemm*
make cutlass_profiler -j

# 4096^3 FP16 GEMM에서 후보 커널들을 실측 비교
./tools/profiler/cutlass_profiler \
    --operation=Gemm --m=4096 --n=4096 --k=4096 \
    --A=f16:row --B=f16:col --C=f16:row --accum=f32

# 출력(개략): 커널 이름마다 Runtime(ms)·GFLOPs. 이름에 타일·stage가 인코딩됨
#   cutlass_tensorop_h16816gemm_128x128_32x3_...   Runtime: ... GFLOPs: ...
#   cutlass_tensorop_h16816gemm_256x128_32x3_...   Runtime: ... GFLOPs: ...

관찰(개략, 환경 의존):

• 같은 GEMM도 타일 구성마다 성능이 갈린다 — 128x128 vs 256x128 vs 128x256이 형상·세대에 따라 순위가 바뀐다.
• 커널 이름의 h16816 = half·MMA m16n8k16, 128x128_32x3 = threadblock $128\times128$, BK=32, Stages=3. 이름만 읽어도 3.4.6의 템플릿 파라미터를 역추적할 수 있다.
• cuBLAS와 비교하면 순수 대형 GEMM은 cuBLAS가 보통 동급 이상 — CUTLASS의 가치는 이 커널에 융합을 끼울 수 있다는 점(3.5·3.6).

3.4.8 중간 점검 (체크포인트)

• [ ] GEMM의 5계층(device·kernel·threadblock·warp·thread, 가장 안쪽이 MMA 명령)을 그리고 각 계층이 어느 메모리/하드웨어(HBM·공유·레지스터·텐서코어)에 사는지 말한다.
• [ ] threadblock tile을 키울 때의 이득(재사용↑·HBM↓)과 비용(공유메모리·레지스터·점유율↓·tile quantization)을 동시에 설명한다.
• [ ] 소프트웨어 파이프라이닝(multistage)이 "무엇의 지연을 무엇으로 은닉하는지"(글로벌 로드를 MMA로) 말하고, Stages↑가 공유메모리·점유율과 줄다리기임을 안다.
• [ ] 타일·정밀도·아키텍처가 전부 컴파일 타임 템플릿 파라미터라 인스턴스화 폭발이 생긴다는 점을 안다.

증상→의심: 텐서코어 활용률 낮은데 메모리 바운드 아님 → 파이프라인 부족(Stages↑·cp.async 경로 확인). 작은 형상에서 GFLOPs 저조 → tile quantization(작은 타일·split-K 고려). 빌드가 수십 분·바이너리 거대 → 켠 타일·아키텍처 조합 과다(CUTLASS_LIBRARY_KERNELS·ARCHS 좁히기). 런타임 kErrorNotSupported → 템플릿이 타깃 SM 미지원(can_implement).

3.4.9 트레이드오프 — CUTLASS 2.x vs 3.x, 그리고 대안

CUTLASS 2.x vs 3.x. 3.x는 CuTe(레이아웃·텐서 대수, 3.5)를 도입하고 API를 collective(threadblock 협력 단위) + kernel layer로 재편했다. Hopper(4.2~4.3)의 TMA·warp-specialization·persistent kernel을 1급으로 표현하려는 재설계다.

견해: 개념·입문은 2.x 계층으로 두고(이 장), 최신 세대(Hopper·Blackwell)에서 실제 성능을 극대화하려면 3.x·CuTe로 간다(3.5). 신규 프로젝트가 Hopper+ 타깃이면 처음부터 3.x 예제를 본다.

무엇으로 GEMM을 구현할 것인가:

비NVIDIA 정직: 같은 자리를 AMD는 Composable Kernel·rocBLAS, 컴파일러 진영은 Triton·XLA·TVM이 채운다. "타일을 하드웨어 계층에 포개고 파이프라인으로 지연을 은닉한다"는 이 장의 사상은 벤더 불문 공통이다 — CUTLASS는 그것을 NVIDIA 텐서코어에 맞춰 템플릿화했을 뿐.

참고 — 2:4 구조적 희소 GEMM(1.14 연결): 1.14에서 본 2:4 구조적 희소성(4원소마다 2개를 0으로)은 별도 커널이 아니라 이 장에서 배운 같은 CUTLASS GEMM 계층의 한 변형으로 지원된다 — 희소 텐서코어 MMA 명령과, 어느 2개가 0인지 나타내는 작은 메타데이터를 함께 쓰는 압축 표현이 축을 하나 더할 뿐, 타일·파이프라인·계층 사상 원리는 동일하다. 이 장은 밀집(dense) GEMM에 집중하고 희소 변형은 개념만 짚는다 — 정확한 API·지원 정밀도·세대는 버전 의존, 확인 필요.

3.4.10 흔한 함정

• cuBLAS보다 빠를 거라는 기대. 순수 대형 GEMM은 cuBLAS가 보통 동급 이상이다(내부에 고도 튜닝 커널·CUTLASS 계열 포함). CUTLASS의 이득은 속도 자체가 아니라 융합·커스텀으로 커널 수를 줄이는 데서 온다. 순수 GEMM을 CUTLASS로 바꿔 더 빠를 거라 단정 금지 — 실측.
• 컴파일 폭발. 타일×정밀도×레이아웃×아키텍처 조합마다 별도 코드가 인스턴스화돼 빌드 수십 분·바이너리 수백 MB. 쓰는 조합만 켠다(CUTLASS_LIBRARY_KERNELS·CUTLASS_NVCC_ARCHS).
• 정렬(alignment) 위반. 텐서코어 벡터화 로드는 포인터·leading dimension이 특정 정렬(예 FP16 8원소=16B)을 요구. 어긋나면 can_implement 실패거나 느린 폴백. 형상·패딩 확인.
• 아키텍처 미스매치. Sm80 템플릿을 Sm70에서 실행하면 cp.async·해당 MMA가 없어 실패·폴백. 타깃 GPU의 SM에 맞춰 인스턴스화(2.12).
• Stages를 맹목적으로 키움. 파이프라인을 깊게 하면 공유메모리를 단계 수만큼 먹어 점유율이 떨어지고 오히려 느려질 수 있다. 형상·세대에 맞춰 실측으로 고른다.
• 2.x 지식으로 Hopper 최적화 시도. Hopper의 TMA·WGMMA·warp-spec은 3.x·CuTe로 표현된다. 2.x 계층으로는 최신 세대 성능을 다 못 낸다 — 3.5·4.2으로.

3.4.11 연습문제

연습 1 (계층 매핑). M=N=K=2048, threadblock $128\times256\times64$, warp $64\times64$, MMA m16n8k16. (a) 출력 타일 수와 threadblock 수, (b) K 루프 반복 수, (c) 블록당 warp 수를 구하라.

연습 2 (공유메모리·점유율). 위 구성에서 FP16 입력, Stages=4. (a) 한 스테이지의 A+B 공유메모리 바이트, (b) 4단계 총량, (c) SM 공유메모리가 100KB(개략)일 때 블록 몇 개 올라가나, 점유율 함의는?

연습 3 (파이프라인). 소프트웨어 파이프라이닝이 없으면 텐서코어가 유휴 상태인 이유를 "글로벌 로드 지연"으로 설명하고, cp.async와 multistage가 각각 무엇을 해결하는지 한 줄씩.

연습 4 (cuBLAS vs CUTLASS 의사결정). 다음 각각에서 cuBLAS와 CUTLASS 중 무엇을 고르고 왜인지: (a) $[8192\times8192]\cdot[8192\times8192]$ 순수 FP16 GEMM, (b) GEMM 출력에 학습 가능한 per-channel scale을 곱하고 즉시 다음 GEMM 입력으로(back-to-back, HBM 안 거치고), (c) 크기가 제각각인 256개 작은 GEMM을 한 런치로(MoE).

연습 5 (컴파일 폭발 진단). CUTLASS 프로파일러 빌드가 40분 걸리고 바이너리가 수백 MB다. 원인과 두 가지 완화책.

3.4.12 요약 · 다음 장

• CUTLASS는 cuBLAS·cuDNN이 표현조차 못 하는 융합·정밀도·형상이 필요할 때 GEMM·conv 커널을 C++ 템플릿으로 조립하는 라이브러리다. 순수 GEMM을 더 빠르게 하려는 도구가 아니라, 커스텀을 삽입하는 도구.
• 핵심 사상: GEMM을 device → threadblock → warp → MMA 타일로 분해해 HBM → 공유메모리 → 레지스터 → 텐서코어 계층에 1:1로 포갠다. 세 타일 크기 + 파이프라인 단계가 곧 튜닝 노브다.
• 소프트웨어 파이프라이닝(multistage·cp.async) 이 글로벌 로드 지연을 MMA로 은닉해 텐서코어를 포화시킨다 — 단 공유메모리·점유율과 줄다리기.
• 타일·정밀도·아키텍처는 컴파일 타임 템플릿 파라미터 → 인스턴스화 폭발에 주의(아키텍처·커널 필터). 최신 세대(Hopper+)는 3.x·CuTe.

다음 장: 3.5 CUTLASS II — CuTe 레이아웃·텐서. 3.x가 타일을 "고정 계층"에서 레이아웃 대수로 해방한 도구가 CuTe다. Layout = Shape ∘ Stride로 임의 텐서 레이아웃을 합성·분할하는 법을 직접 실습한다.