제3부 딥러닝 라이브러리

3.5 CUTLASS II — CuTe 레이아웃·텐서

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

3.4의 CUTLASS 2.x는 타일을 고정된 계층 템플릿(GemmShape)으로 표현했다. 3.x는 그 계층을 레이아웃 대수로 해방한다 — 그 도구가 CuTe(CUDA Templates, "cute") 다. CuTe의 단 하나의 핵심 개념은 Layout = (Shape, Stride) 이고, 이것은 "논리 좌표 → 메모리 오프셋"의 함수다. 텐서 분할(threadblock·warp·thread로 나누기), 타일링, 전치, 벡터화가 전부 이 한 추상으로 통일된다. 이 장은 레이아웃을 직접 계산하고 print해 보며 익힌다(실습 장).

배우는 것:

• Layout = Shape ∘ Stride 를 좌표→오프셋 함수로 읽고, 1D·다차원·계층(중첩) 좌표를 직접 오프셋으로 변환한다.
• row-major·col-major·전치·타일을 전부 stride로 표현하고, coalesce·composition의 의미를 안다.
• logical divide(분할) 로 한 텐서를 타일·스레드로 나누는 것이 CUTLASS 분할의 정체임을 본다.
• Tensor = Engine(포인터) + Layout 를 만들고 local_tile·local_partition으로 블록·스레드 몫을 떼는 골격을 읽는다.

선수 지식: 3.4(CUTLASS 계층·타일 — CuTe는 그 계층을 표현하는 대수), 1.7~1.8(공유메모리 뱅크·메모리 병합 — 레이아웃이 곧 접근 패턴), 2.7(타일링). 약간의 함수 합성·인덱싱에 대한 이해. CuTe는 헤더-온리 C++17 템플릿이며 버전마다 API가 바뀐다 — 코드는 개략이며 정확한 이름은 설치 버전 include/cute/·cutlass/media/docs/cute/로 확인. 출력 수치는 정의에 따른 계산이라 결정적(환경 무관).

3.5.2 동기 — 왜 "레이아웃 대수"가 필요한가

GPU 커널의 절반은 "어떤 스레드가 어떤 데이터를 만지는가" 를 인덱스 산수로 푸는 일이다. 2.7의 타일 GEMM에서 row = blockIdx.y*BM + threadIdx.y 같은 줄을 직접 작성했다. 문제는:

• 레이아웃이 바뀌면(전치·타일·벡터화·뱅크 회피 패딩) 그 인덱스 산수를 전부 다시 작성해야 한다.
• 공유메모리 뱅크 충돌 회피(1.7), 메모리 병합(1.8), 텐서코어 단편 레이아웃(1.11)이 전부 다른 인덱싱이고, 섞이면 사람이 추적 불가.

CuTe의 답: 인덱싱을 코드에 흩지 말고, "Layout"이라는 1급 객체로 들고 다니며 대수적으로 합성한다. 좌표→오프셋 매핑을 데이터처럼 다루면, 전치는 stride 교환, 타일링은 layout division, 스레드 분할도 division — 하나의 연산으로 통일된다.

3.5.3 Layout = (Shape, Stride) — 좌표를 오프셋으로

Layout은 함수다: 논리 좌표를 받아 1D 메모리 오프셋을 돌려준다. 두 부분으로 정의된다:

• Shape: 각 차원(mode)의 크기. 예 (4, 6) = 4행 6열.
• Stride: 각 차원에서 좌표가 1 늘 때 오프셋이 얼마 뛰는지. 예 (1, 4).

오프셋 공식(좌표 (i, j)):

$$\mathrm{offset}(i,j) = i \cdot \mathrm{stride}_0 + j \cdot \mathrm{stride}_1$$

예 1 — column-major $4\times6$, stride (1,4):

• $\mathrm{offset}(i,j) = i \cdot 1 + j \cdot 4$
• $(0,0) \to 0 \quad (1,0) \to 1 \quad (2,0) \to 2 \quad (3,0) \to 3$ ← 열이 연속(같은 $j$는 1씩)
• $(0,1) \to 4 \quad (1,1) \to 5 \quad \ldots$ ← 다음 열은 +4
• 메모리: 열0[0,1,2,3] 열1[4,5,6,7] ... (Fortran/cuBLAS 기본, 3.1)

예 2 — row-major $4\times6$, stride (6,1):

• $\mathrm{offset}(i,j) = i \cdot 6 + j \cdot 1$
• $(0,0) \to 0 \quad (0,1) \to 1 \quad \ldots \quad (0,5) \to 5$ ← 행이 연속
• $(1,0) \to 6 \quad \ldots$ ← 다음 행은 +6
• 메모리: 행0[0..5] 행1[6..11] ... (C/PyTorch 기본)

핵심 통찰: 전치는 데이터 이동이 아니라 stride 교환이다. $4\times6$ row-major (6,1)를 전치한 $6\times4$는 stride (1,6)로 같은 메모리를 다르게 읽는 레이아웃일 뿐(복사 0). cuBLAS의 "행우선을 전치로 속이는 트릭"(3.1)이 이 사실의 응용이다.

CuTe 표기: Layout = Shape:Stride. 위 둘은 (4,6):(1,4), (4,6):(6,1). 코드:

#include <cute/layout.hpp>
using namespace cute;
auto colmaj = make_layout(make_shape(Int<4>{}, Int<6>{}), make_stride(Int<1>{}, Int<4>{}));
auto rowmaj = make_layout(make_shape(Int<4>{}, Int<6>{}), GenRowMajor{}); // stride 자동 (6,1)
print_layout(colmaj);   // 좌표→오프셋 표를 ASCII로 찍어줌(디버깅 필수 도구)
print(rowmaj(1,2));     // offset(1,2)=1·6+2·1=8 출력

print_layout은 CuTe 학습의 핵심 도구다 — 레이아웃을 머리로 추측하지 말고 출력해서 좌표→오프셋 격자를 눈으로 확인한다.

3.5.4 계층(중첩) 레이아웃 — 타일이 곧 중첩 shape

CuTe가 강력한 이유: Shape와 Stride가 중첩(계층)될 수 있다. mode 하나가 다시 (a,b) 같은 튜플일 수 있고, 그것이 곧 "이 차원은 타일 구조를 가진다"는 뜻이다.

예 — 8 길이를 (2,4)로 본 레이아웃 (2,4):(4,1):

• 좌표 $(i,j)$, $i \in \{0,1\}$, $j \in \{0,1,2,3\}$
• $\mathrm{offset} = i \cdot 4 + j \cdot 1$
• 논리적으로는 8칸을 "2개의 4묶음"으로 본다:
  • $i=0$: [0 1 2 3] $\quad i=1$: [4 5 6 7]

이제 실질적 쓸모 — 타일링을 중첩 shape로. (M,N) 행렬을 (BM,BN) 타일로 나눈 결과를 한 레이아웃으로 표현:

   shape  ((BM, M/BM), (BN, N/BN))      ← 각 축이 (타일내 좌표, 타일 인덱스)
   coord  ((m_in,  m_tile), (n_in, n_tile))
   → 한 좌표로 "몇 번 타일의 안쪽 어디"를 동시에 가리킴

이 중첩이 3.4의 device→threadblock→warp→thread 계층을 한 레이아웃 객체로 표현하는 방법이다. "타일의 타일의 타일"이 중첩 shape의 중첩 shape로 쌓인다. CUTLASS 3.x가 2.x의 고정 GemmShape 계층을 버리고 CuTe로 간 이유가 이것 — 계층을 코드 구조가 아니라 데이터(레이아웃) 로 들고 다닌다.

좌표 선형화(colexicographic). 다차원 좌표를 1D 인덱스로 접을 때 CuTe는 첫 mode가 가장 빨리 순회하는(column-major스러운) 순서를 쓴다. (2,4) shape에서 1D 인덱스 5 → 좌표 (1, 2) (5 = 1 + 2·2... 주의: shape (2,4)의 자연 stride는 (1,2)). 이 규약을 정확히 알아야 print_layout 출력을 읽는다.

3.5.5 워크스루 A — 레이아웃 오프셋 수계산

연습 삼아 종이로 풀고 print_layout으로 검산하는 흐름을 보인다.

문제. Layout (4,2):(2,1). 모든 좌표의 오프셋 표를 채워라.

오프셋 식 offset(i,j) = i·2 + j·1 (i∈0..3, j∈0..1):

메모리 [0..7]을 "행이 2칸씩 연속, 4행"으로 본 것 = row-major 4×2.

검산 코드:

auto L = make_layout(make_shape(Int<4>{},Int<2>{}), make_stride(Int<2>{},Int<1>{}));
print_layout(L);          // 위 표와 일치 확인
assert(L(2,1) == 5);      // 2·2 + 1·1 = 5
assert(size(L) == 8);     // 전체 원소 수
assert(cosize(L) == 8);   // 도달하는 최대 오프셋+1 (메모리 점유 범위)

왜 cosize가 중요한가. size는 논리 원소 수, cosize는 실제 메모리에서 차지하는 범위다. 둘이 다르면(예: 패딩·스트라이드 점프) 메모리에 구멍이 있다는 뜻 — 공유메모리 뱅크 충돌 회피용 패딩(1.7)이 정확히 이 경우다. (4,4):(5,1)처럼 stride를 일부러 4가 아닌 5로 주면 행마다 1칸씩 간격을 두어 뱅크 충돌을 깬다. 그 패딩이 레이아웃의 stride에 자연스럽게 들어간다.

3.5.6 레이아웃 연산 — coalesce·composition·divide

CuTe의 실질적 힘은 레이아웃을 합성하는 데 있다. 세 가지가 핵심이다.

coalesce — 합칠 수 있는 mode를 합친다. (2,4):(1,2)(열우선)을 평탄화하면 stride가 맞아떨어져(인접 mode가 $d1 = s0\cdot d0$) 8:1(연속 8칸)과 동등. 불필요한 계층을 접어 단순화한다. (반대로 위 (2,4):(4,1)은 colex 선형화가 [0,4,1,5,2,6,3,7]로 흩어져 8:1로 합쳐지지 않는다.)

composition — 레이아웃을 함수처럼 합성한다 R = A ∘ B. "B로 좌표를 매핑한 결과를 다시 A로 매핑". 이것이 타일 안에서의 재인덱싱을 표현한다. 예: 전치는 항등 데이터에 stride를 바꾸는 레이아웃을 합성한 것. 벡터화 로드(연속 4원소를 한 번에)는 _4 타일을 합성한 것.

logical divide — 텐서를 타일/스레드로 나눈다(가장 실무적). A를 타일러 B로 나누면, 결과 레이아웃의 한 mode는 "타일 안", 다른 mode는 "타일 인덱스"가 된다(3.5.4의 중첩 구조가 자동 생성).

   logical_divide(A, tiler) → ((타일_안 좌표), (타일_인덱스 좌표))
   예) 8칸 벡터를 4씩 나눔:  8:1  ÷  4  →  (4, 2):(1, 4)
         mode0 = 타일 안 4칸(stride1),  mode1 = 2개 타일(stride4)

이 한 연산이 CUTLASS의 분할 전부를 떠받친다:

• local_tile(전체텐서, 타일shape, 타일좌표): 전체 텐서에서 이 블록이 맡을 타일 하나를 떼낸다(device→threadblock).
• local_partition(타일, 스레드레이아웃, threadIdx): 타일을 스레드들에게 나눠 이 스레드 몫을 떼낸다(threadblock→thread).

즉 2.7에서 직접 작성한 blockIdx·threadIdx 인덱스 산수가 CuTe에선 local_tile·local_partition 두 줄로 끝난다 — 레이아웃이 바뀌어도 코드는 그대로.

3.5.7 워크스루 B — Tensor와 분할 골격

Tensor = 포인터(Engine) + Layout. 레이아웃에 실제 메모리를 붙이면 텐서다.

#include <cute/tensor.hpp>
using namespace cute;

// 1) HBM의 A(M×K) row-major를 텐서로 감싼다 (데이터 이동 없음, 뷰)
auto gA = make_tensor(make_gmem_ptr(dA),
                      make_layout(make_shape(M, K), GenRowMajor{}));  // (M,K):(K,1)

// 2) 이 블록(blockIdx)이 맡을 BM×BK 타일 하나를 떼낸다
auto tile = make_shape(Int<128>{}, Int<32>{});           // BM×BK
auto blk_coord = make_coord(blockIdx.x, _);              // (이 블록의 M타일, K는 전체)
auto gA_blk = local_tile(gA, tile, blk_coord);           // (128,32, K/32) — 이 블록 몫

// 3) 블록 안 256 스레드에게 타일을 나눠 이 스레드 몫을 떼낸다
auto thr_layout = make_layout(make_shape(Int<16>{}, Int<16>{}));  // 16×16=256 스레드 배치
auto tAgA = local_partition(gA_blk, thr_layout, threadIdx.x);     // 이 스레드가 만질 원소들

// 4) 공유메모리 목적지도 레이아웃으로 — 뱅크 회피 패딩을 stride로
__shared__ half smem[128*33];   // 패딩 stride(33) 반영 → cosize=128*33(=4224), [128*32]면 오버런
auto sA = make_tensor(make_smem_ptr(smem),
                      make_layout(make_shape(Int<128>{},Int<32>{}),
                                  make_stride(Int<33>{},Int<1>{}))); // ★ 33 패딩→뱅크충돌 회피
auto tAsA = local_partition(sA, thr_layout, threadIdx.x);

// 5) 복사: 스레드 몫을 글로벌→공유로 (cute::copy가 벡터화·cp.async 선택)
copy(tAgA, tAsA);
__syncthreads();
// 이후 MMA: cute::gemm(tiled_mma, ...) 로 텐서코어 호출(3.x collective가 감쌈)

읽는 법:

• make_tensor는 메모리를 복사하지 않는다 — 포인터 + 레이아웃 뷰일 뿐. 전치·타일·재배열이 전부 레이아웃 조작이라 0복사.
• local_tile/local_partition이 2.7의 손 인덱싱을 대체. 레이아웃만 바꾸면 같은 코드가 다른 분할로 동작.
• (33,1) 패딩 stride 한 줄로 공유메모리 뱅크 충돌(1.7)을 깬다 — 인덱스 산수에 패딩을 흩지 않고 레이아웃에 고정한다.
• copy·gemm은 텐서의 레이아웃을 보고 cp.async·ldmatrix·MMA 명령을 알아서 고른다(3.4의 파이프라인 빌딩블록).

이것이 CUTLASS 3.x가 Hopper(4.2~4.3)의 TMA·WGMMA·warp-specialization을 표현하는 방식이다 — 전부 "어떤 레이아웃의 텐서를 어디로 복사하고 어떤 MMA로 곱하나"의 레이아웃 문제로 환원된다.

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

• [ ] Layout = (Shape, Stride)를 좌표→오프셋 함수로 읽고, row-major·col-major·전치를 stride로 쓴다.
• [ ] 중첩 shape ((BM,M/BM),(BN,N/BN))가 왜 "타일링을 한 레이아웃으로" 표현하는지 설명한다.
• [ ] logical_divide/local_tile/local_partition이 2.7의 blockIdx·threadIdx 인덱스 산수를 대체함을 안다.
• [ ] 공유메모리 패딩(뱅크 회피)을 stride로 표현하고 size vs cosize로 메모리 구멍을 진단한다.

증상→의심: print_layout 출력이 예상과 다름 → colexicographic 선형화·stride 오해(수계산 재검). 공유메모리 접근이 느림 → 패딩 stride 누락(뱅크 충돌, cosize 확인). 분할 결과 크기 안 맞음 → tiler shape가 텐서 shape를 나누지 못함(나눠떨어지는지·divisibility). 빌드 시 템플릿 에러 폭포 → static Int<> vs 동적 int 혼용·shape 불일치.

3.5.9 트레이드오프 — CuTe를 직접 쓸 것인가

견해: CuTe를 직접 쓸 일은 커널 연구자·CUTLASS 기여자·"라이브러리에 없는 융합을 Hopper에서 최대 성능으로"가 필요한 소수다. 대부분은 ① 표준이면 cuBLAS/cuDNN, ② 커스텀이 필요하면 Triton으로 먼저 시도하고, Triton으로도 부족할 때 CuTe로 내려간다. 단 CUTLASS 3.x 예제·FlashAttention 3(3.6)·최신 커널을 읽으려면 CuTe 독해는 필수다 — 그래서 이 장이 "쓰기"보다 "읽기"에 무게를 둔다.

비NVIDIA 정직: 레이아웃을 1급 객체로 다루는 발상은 CuTe만의 것이 아니다 — Triton의 block pointer, XLA/MLIR의 layout/affine map, Halide의 schedule이 같은 문제(데이터 배치와 연산 분리)를 다른 문법으로 푼다. "인덱싱을 코드에서 분리해 레이아웃으로 추상화한다"는 사상이 본질이고 CuTe는 그 NVIDIA 텐서코어용 구현이다.

3.5.10 흔한 함정

• 선형화 규약 혼동. CuTe는 첫 mode가 가장 빨리 순회하는 colexicographic. row-major 관점((i,j)에서 j가 빨리)으로 읽으면 오프셋이 어긋난다. print_layout으로 항상 검산.
• static vs dynamic 정수 혼용. Int<128>{}(컴파일 타임)와 int(런타임)를 섞으면 컴파일러가 최적화를 못 하거나 템플릿 에러. 가능하면 컴파일 타임 상수로(레지스터 인덱싱·언롤 이득).
• divisibility 위반. tiler가 텐서 shape를 나눠떨어뜨리지 못하면 분할이 깨지거나 경계 타일이 잘못된다. 경계 처리(predication)는 별도 — 직접 작성한 경계검사(2.2)가 CuTe에선 predication 텐서로.
• size vs cosize 무시. 패딩·전치로 메모리에 구멍이 생긴 레이아웃을 연속이라 가정하고 memcpy하면 깨진다. 공유메모리 할당은 cosize 기준.
• 0복사 뷰를 데이터로 착각. make_tensor·전치·타일은 뷰다 — 원본을 수정하면 같이 바뀐다. 독립 복사가 필요하면 명시적 copy.
• CuTe를 표준 GEMM에 끌어옴. 표준 GEMM을 CuTe로 직접 작성해 cuBLAS를 이기려는 시도는 보통 무익한 작업이다. CuTe는 라이브러리에 없는 것을 위한 도구.

3.5.11 연습문제

연습 1 (오프셋 수계산). Layout (3,4):(1,3)의 좌표→오프셋 표를 채우고, 이것이 row-major인지 col-major인지, 메모리 연속인지 말하라.

연습 2 (전치). 위 (3,4):(1,3)를 전치한 $4\times3$ 레이아웃의 shape·stride는? 메모리는 한 바이트라도 움직이는가?

연습 3 (뱅크 회피 패딩). 공유메모리 $32\times32$ FP32 타일을 col-major로 두면 한 행을 32 스레드가 읽을 때 모두 같은 뱅크에 떨어져 32-way 충돌(1.7)이 난다. 레이아웃 stride로 어떻게 깨나? size와 cosize는?

연습 4 (분할 = divide). 8칸 1D 텐서 8:1을 4스레드가 나눠 갖는다. (a) 각 스레드가 연속 2칸을 갖는 분할, (b) 각 스레드가 stride-4로 흩어진 2칸을 갖는(병합 친화) 분할의 레이아웃을 각각 쓰고, 어느 쪽이 글로벌 메모리 병합(1.8)에 유리한지.

연습 5 (왜 CuTe인가). 같은 GEMM 커널이 (입력 전치 여부 × FP16/BF16/FP8 × 뱅크 패딩 유무)로 수십 변형이 필요하다. 전통적 손 인덱싱 대비 CuTe가 코드량·버그를 줄이는 메커니즘을 한 문단으로.

3.5.12 요약 · 다음 장

• CuTe의 핵심은 Layout = (Shape, Stride) 한 추상 — 좌표를 메모리 오프셋으로 보내는 함수. row-major·col-major·전치·패딩·벡터화가 전부 stride로 통일된다.
• 중첩 shape가 3.4의 device→block→warp→thread 계층을 한 레이아웃 객체로 표현하고, logical divide / local_tile / local_partition 이 손 인덱싱을 대체한다.
• Tensor = 포인터 + Layout(0복사 뷰). 전치·타일은 데이터 이동 없이 레이아웃 조작. 뱅크 회피·메모리 병합이 레이아웃의 stride 선택으로 결정된다.
• CuTe를 직접 쓸 일은 소수(커널 연구·최신 세대 최적화)지만, CUTLASS 3.x·FlashAttention 최신 커널을 읽으려면 필수.

다음 장: 3.6 어텐션 I — 메모리 문제·FlashAttention 원리. 지금까지의 GEMM·융합·레이아웃 도구가 모이는 첫 큰 응용이다. 어텐션이 왜 메모리 바운드가 되는지($S=QK^T$의 $O(N^2)$ 중간값), 그것을 타일링 + online softmax로 HBM에 안 쓰고 푸는 FlashAttention의 원리를 본다.