V06 — GEMM 완전판 단권화

1 GEMM 표준형 General Matrix-Matrix Multiply α·A·B+β·C

2 전치 조합 표기 Column-major BLAS 관습

CUTLASS는 LayoutA/LayoutB/LayoutC로 명시. cf. cuBLAS convention.

3 dtype 조합 A · B · accumulator · output ★

dtype 조합·bit layout 근거는 ↗ V09 §3·§4.

4 FLOP · byte 회계

Roofline 계산 원리는 ↗ V18 §3.

5 Compute / Memory bound 경계 M·N 作, K 消

• M, N 크고 K 크면 compute-bound
• M 또는 N 작으면 memory-bound → Split-K/Stream-K 필요 ↗ V05 §7·§8
• K=1 수준 (GEMV) → 완전 memory-bound

6 ScaleType · accumulator 의미

epilogue scale은 출력 직전 FP32 acc에서 수행 (precision 보존).

7 본 권의 좌표 제외 항목

• Attention dispatch → ↗ V07 §4·§10
• Fusion 일반론 → ↗ V13 §8
• MoE grouped GEMM 맥락 → ↗ V08 §3
• PTX mma·ldmatrix → ↗ V03 §7·§8
• WGMMA·TMA PTX → ↗ V04 §5·§7

1 Naive 3-nested loop

// row-major, no tiling
for (m = 0; m < M; ++m)
  for (n = 0; n < N; ++n) {
    acc = 0;
    for (k = 0; k < K; ++k)
      acc += A[m,k] * B[k,n];
    D[m,n] = alpha*acc + beta*C[m,n];
  }

2 3-level 메모리 blocking ★

HBM  ──► L2  ──► SMEM ──► Register
 ↑       ↑       ↑         ↑
 M,N,K   CTA     warp      thread
 전체    tile    tile      fragment

• CTA tile: M_C × N_C × K_C
• Warp tile: M_W × N_W × K_W
• Thread tile: M_T × N_T (reg)

3 Tiled GEMM 의사코드 ★

// CTA(M_C, N_C), stages=S
for (k0 = 0; k0 < K; k0 += K_C) {
  // load to smem (cp.async / TMA)
  load_tile(sA[k0%S], gA[m_blk, k0 : k0+K_C]);
  load_tile(sB[k0%S], gB[k0 : k0+K_C, n_blk]);
  __syncthreads();
  for (kw : warp_K_iter) {
    rA = ldmatrix(sA, kw);
    rB = ldmatrix(sB, kw);
    mma(acc, rA, rB);  // tensor core
  }
}
epilogue(gD, acc, alpha, beta, gC);

4 재사용률 공식

tile이 클수록 HBM traffic ↓ 그러나 SMEM · reg pressure ↑.

5 Double / Multi-stage buffering

6 계층별 BW 비율

A100 whitepaper · SMEM은 SM 당 ~140 GB/s × 108 SM. 계층 상세 ↗ V02 §3·§5.

7 실수 포인트

• K 분할 안 하면 K 전체를 SMEM에 올리려 함 → overflow
• __syncthreads 누락 → stage 간 race
• stride 잘못 → uncoalesced HBM load

1 5-Level 전체 도식 ★ D-K-C-W-T

┌──────────────────────────────────────┐
│ Device   M×N×K (전체 문제)            │  in : A(M,K), B(K,N), C(M,N)
│                                      │  out: D(M,N)
│   ▼ gemm_universal launches grid     │
│                                      │
│ Kernel   grid + TileScheduler        │  in : problem_shape, tile_shape
│                                      │  out: per-CTA (m_blk, n_blk, k_iter)
│   ▼ CTA = tile 1개 담당              │
│                                      │
│ CTA     M_C×N_C × (k0..K) stages      │  in : gA tile, gB tile
│                                      │  out: CTA acc fragment (smem/reg)
│   ▼ 여러 warp(group)가 분담          │
│                                      │
│ Warp    M_W×N_W × K_W                 │  in : sA slice, sB slice
│                                      │  out: warp acc (register)
│   ▼ warp = mma instruction 묶음      │
│                                      │
│ Thread  M_T×N_T (reg fragment)        │  in : rA, rB fragment
│                                      │  out: rC fragment
└──────────────────────────────────────┘

2 각 레벨의 책임

3 입출력 타입 명시

4 Template 계층 이름 규칙

// Device
cutlass::gemm::device::GemmUniversal<
  Kernel_>;
// Kernel
cutlass::gemm::kernel::GemmUniversal<
  CollectiveMainloop_, CollectiveEpilogue_,
  TileScheduler_>;
// Threadblock (2.x)
cutlass::gemm::threadblock::Mma<
  Shape_, Warp_, Policy_, Stages_>;

5 2.x vs 3.x 차이 ★

6 암기 키

1 CTA tile 구성

2 Mainloop 골격 ★

// CTA scope
acc = zero();
prologue_prefetch(S);  // 첫 S stage
for (k_iter : K/K_C - S) {
  issue_next_stage();   // async load
  wait_oldest_stage();
  warp_mma(acc, sA[oldest], sB[oldest]);
  advance_stage();
}
drain_remaining(S);     // 마지막 tiles
epilogue(acc);

3 SMEM 사용량 ★

4 Predicate & 경계

• M, N, K가 tile 배수 아님 → predicate mask
• out-of-bound load는 0 또는 identity
• 2.x: PredicatedTileIterator · 3.x: TMA descriptor가 처리
• 경계 CTA는 K iteration도 다름 (split-K)

5 Stage count trade-off

6 SMEM 레이아웃 배치

smem:
 ┌─ sA[0] ──┐ ┌─ sA[1] ──┐ ... ┌─ sA[S-1] ──┐
 │ M_C×K_C  │ │ M_C×K_C  │     │ M_C×K_C    │
 ├─ sB[0] ──┤ ├─ sB[1] ──┤ ... ├─ sB[S-1] ──┤
 │ K_C×N_C  │ │ K_C×N_C  │     │ K_C×N_C    │
 └──────────┘ └──────────┘     └────────────┘
 ┌─ acc buffer (for epilogue split) ──┐
 └────────────────────────────────────┘

7 Warp-group 분할 (Hopper) 3.x

• 1 producer WG + 1~2 consumer WG
• producer: TMA load → mbarrier arrive
• consumer: mbarrier wait → wgmma → epilogue
• 상세 스케줄은 ↗ §15 p16

1 Warp tile 구성

2 MMA shape · warp tile 관계 ★

mma shape 전체 목록은 ↗ V03 §7 (Ampere) · ↗ V04 §8 (Hopper).

3 Warp tile → MMA 반복

4 MmaMultistage 구조 (2.x)

// Ampere warp-level multi-stage
class MmaMultistage {
  IteratorA iter_a;   // gmem→smem
  IteratorB iter_b;
  Policy::Warp warp_mma;
  int smem_stage = 0;

  void operator()(acc, k_iter) {
    prologue();
    for (k : k_iter) {
      ldmatrix(rA, sA[stage]);
      ldmatrix(rB, sB[stage]);
      warp_mma(acc, rA, rB, acc);
      advance();
    }
  }
};

5 Tensor Core fragment layout ★

WGMMA는 A를 register 대신 SMEM으로 받을 수 있음 (A-descriptor).

6 ldmatrix ↔ fragment 정합

smem tile (16×16 FP16)
 ┌──────────────────────┐
 │ warp 32 thread 협력   │ ──► ldmatrix.x4
 │ 16B × 32 = 512B       │
 └──────────────────────┘
        │
        ▼
 register fragment (A shape)

ldmatrix.trans variant = 전치 로드. mma 입력 shape에 맞춰 thread-mapping 자동.

7 Accumulator 수명

• K iteration 전체를 register에 누적
• epilogue 직전에만 smem/HBM으로 flush
• register pressure = M_W·N_W / 32 elem/thread · 4B
• e.g. 64×128 FP32 acc → 64·128/32·4 = 1024 B/thread = 256 reg

1 Iterator 개념 2.x 추상

• PredicatedTileIterator : 경계 mask 포함
• RegularTileIterator : 경계 없음 (aligned)
• MaskedTileIterator : 외부 mask 받음

2 IteratorA / IteratorB 템플릿

using IteratorA = PredicatedTileIterator<
  MatrixShape<M_C, K_C>,
  ElementA,
  LayoutA,
  AdvanceRank,   // K 축 전진
  ThreadMap,     // thread → elem 맵
  AccessSize>;   // vector width

3 ThreadMap 역할

4 Thread-level 출력 fragment ★

5 Fragment 크기 계산

6 Vectorized access

• 128-bit (16 B) per thread가 이상적 (coalesced)
• FP16 → 8 elem/thread vector load
• FP32 → 4 elem/thread
• access size < 128 bit이면 transaction 낭비

7 AccessSize trade-off

8 Register 압박 진단

9 2.x → 3.x 전환

다음 페이지 §7부터 CuTe 본론.

1 문제의식 ★

• IteratorA, IteratorB, SmemIterator, AccumulatorIterator ...
• each × (layout × dtype × shape × stage) 조합
• indexing 수식을 template으로 인코딩 → 난독

2 CuTe의 한 줄 요약 shape+stride=Layout

3 Layout의 위력

4 Hierarchical shape ★

shape = (4, (2, 3))    → 4×(2·3) = 24 elem
stride = (6, (3, 1))   → col-major outer, row-major inner
                         4 그룹 × (2×3 tile)

5 같은 데이터, 여러 View

6 Compile-time vs runtime

static 부분은 template 메타로 fold 되어 런타임 cost 0.

7 요구 지식

• C++17 template + CTAD
• integer_sequence / tuple 메타
• GCD · 나눗셈 · coprime 개념
• modular arithmetic (swizzle)

1 Layout 정의 ★

Layout<(M,N), (1,M)>     → col-major  M×N
Layout<(M,N), (N,1)>     → row-major
Layout<(M,N), (1,0)>     → N축 broadcast
Layout<(M,N), (0,1)>     → M축 broadcast

2 make_* helper

auto s = make_shape(_8{}, _16{});
auto d = make_stride(_16{}, _1{});
auto L = make_layout(s, d);   // row-major 8×16

// 간이 ctor
auto L2 = make_layout(
   make_shape(_8{}, _16{}),
   LayoutRight{});              // row-major
auto L3 = make_layout(s, LayoutLeft{});  // col-major

3 size · cosize · rank

4 Injective / Surjective

• size == cosize → bijective (1:1)
• size < cosize → "sparse" (내부 갭)
• size > cosize → broadcast (동일 elem 공유, stride=0)

5 Hierarchical shape 예 ★

shape=((4,2),(8,4))
stride=((1,16),(4,32))

좌표 ((i,j),(k,l)):
L = i·1 + j·16 + k·4 + l·32
  → col-major 8×8 tile in row-major 4×8 super-grid

6 Column/Row major 레시피

7 print_layout 출력 이해

// Layout<(4,3),(1,4)> col-major 4×3
     0   1   2
0    0   4   8
1    1   5   9
2    2   6  10
3    3   7  11
// cell 값 = L(i,j) offset

print_layout(L)·print_tensor(t) 디버그 필수.

1 연산 전체 표 ★★

2 composition ∘ 함수 합성

3 logical_divide ★

4 logical_product

= replicate L tile. GEMM의 CTA → 전체 grid 복원에 사용.

5 coalesce · complement 보조

• coalesce: stride가 정확히 연속되는 축 병합 → 연산 단순화
• complement: bijective 유지하며 빈 stride 채움

6 zipped_divide 사용 예 ★

// CTA tile → thread partition
auto cta_tile = Layout<Shape<_128,_64>,
                         Stride<_64,_1>>{};
auto thr_tile = Layout<Shape<_32,_8>,
                         Stride<_8,_1>>{};
auto ztile = zipped_divide(
      cta_tile, thr_tile);
// ((thr_M, thr_N), (rest_M, rest_N))
// 각 thread가 자신의 rest tile 받음

7 조합 법칙 요약

8 디버깅 습관

1 Tensor = Engine + Layout

2 make_tensor 사용

auto gA = make_tensor(
   make_gmem_ptr(A_ptr),
   make_shape(M, K),
   make_stride(K, _1{}));

auto sA = make_tensor(
   make_smem_ptr(smemA),
   Layout<Shape<_128,_64>,
          Stride<_64,_1>>{});

auto rA = make_tensor<half>(
   Layout<Shape<_8>>{});   // register

3 local_tile ★

// 128×64 CTA tile 한 개
auto tA = local_tile(gA,
   make_shape(_128{}, _64{}),
   make_coord(m_blk, k_blk));

4 local_partition ★

// 32×8 thread layout (256 thr/CTA)
auto thr_layout = Layout<
   Shape<_32,_8>,
   Stride<_8,_1>>{};
auto tAgA = local_partition(
   tA, thr_layout, threadIdx.x);

5 copy / gemm 호출

// Copy atom으로 gmem → smem
copy(copy_atom,
     thr_copy.partition_S(gA_tile),
     thr_copy.partition_D(sA));
cp_async_fence();
cp_async_wait<0>();

// TiledMMA로 연산
auto tCrA = thr_mma.partition_A(sA);
auto tCrB = thr_mma.partition_B(sB);
auto tCrC = thr_mma.partition_C(rC);
gemm(tiled_mma, tCrA, tCrB, tCrC);

6 Iterator 패턴

7 Helper 라이브러리

• make_tensor / make_layout
• partition_A/B/C (TiledMMA slice)
• partition_S/D (TiledCopy)
• recast<T> (dtype 재해석, e.g. uint32 → 2×half)
• flatten, group_modes
• print_tensor, print_layout

1 왜 필요한가 bank conflict 회피

• row-major FP16 row = 128 B = 32 word → 한 row가 32 bank 전부 hit
• 같은 col 접근은 stride 32 word → 모두 bank 0에 몰림
• ldmatrix / wgmma는 bank conflict 없는 layout 요구

2 Swizzle 함수 정의 ★

3 파라미터 의미 표 ★

Swizzle<3,3,3> = 128 B line 내 8×8 block 재배치. WGMMA A/B 기본.

4 시각적 예시 ★ Swizzle<2,0,2>

원본 offset(6-bit):
 bit:  5 4 3 2 1 0
 ─────────────────
 S=2, B=2, M=0
 → XOR region [3:2] with [1:0]

offset  bin     swizzled
 0      000000  000000  (row 0, col 0)
 1      000001  000001  (row 0, col 1)
 4      000100  000100  (row 1, col 0)
 5      000101  000101  (row 1, col 1)
 8      001000  001010  (row 2, col 0 → 2)
 9      001001  001011  (row 2, col 1 → 3)
12      001100  001110
16      010000  010100
20      010100  010000  (swapped!)

5 표준 preset (CUTLASS)

preset 이름의 숫자는 보통 swizzle chunk 크기 바이트.

6 Layout composition

// 기본 SMEM layout 위에 swizzle 덮어쓰기
using SmemLayoutA = decltype(
  composition(
    Swizzle<3, 3, 3>{},
    Layout<Shape<_64,_64>,
           Stride<_64,_1>>{}));

// 좌표 (i,j) → swizzled offset
offset = SmemLayoutA(i, j);
// ldmatrix / wgmma descriptor에도 swizzle 인코딩

7 WGMMA descriptor 제약

1 TMA 요약 Tensor Memory Accelerator

• PTX: cp.async.bulk.tensor.{1,2,3,4,5}d
• mode: tile, im2col
• 완료: mbarrier arrive (async proxy)

TMA PTX 전체는 ↗ V04 §5·§6.

2 Tensor descriptor 구조

3 CuTe → TMA 생성 ★

// CuTe Layout이 TMA descriptor로 직변환
auto tma_a = make_tma_copy(
   SM90_TMA_LOAD{},
   gA_tensor,                    // global Tensor
   SmemLayoutA{},                // swizzled smem
   make_shape(_128{}, _64{}));   // box

// kernel 내에서 partition
auto thr_tma = tma_a.get_slice(block_id);
auto tAgA = thr_tma.partition_S(gA);
auto tAsA = thr_tma.partition_D(sA);

4 Launch 흐름

Host :
 ─ make_tma_copy  ──► TMA desc struct
                      (global tensor map)
 ─ kernel<<<...>>>(desc, ...)

Device :
 thread 0 만 발행:
 ─ cp.async.bulk.tensor (desc, coord, smem, bar)
 ─ HW가 shape/stride/swizzle 해석
 ─ 완료시 bar.arrive(expected_bytes)
 모든 thread:
 ─ bar.wait(parity)

5 expected_bytes 수식

여러 TMA가 같은 bar 사용시 exp_bytes는 합산.

6 multicast TMA ★

auto tma_a_mc = make_tma_copy(
   SM90_TMA_LOAD_MULTICAST{},
   gA, SmemLayoutA{},
   CtaShape{},
   cluster_shape);

multicast는 cluster ≥ 2 + same A tile CTAs.

7 제약 · 함정

• box_shape가 descriptor의 global_shape 안에 fit 해야
• SMEM alignment: 16 B (WGMMA) · 128 B (swizzle chunk)
• per-tensor descriptor 수 제한 (RT API cuTensorMapEncodeTiled)
• cluster launch 실패시 swizzle 조합 재검토

1 3.x의 재구성

Kernel<...>
 ├ CollectiveMainloop  (K iter · stages · wgmma)
 ├ CollectiveEpilogue  (acc → epilogue → gD)
 └ TileScheduler       (persistent · streamk · default)

2 Dispatch Policy ★

sm_90a에서만 TmaWarpSpecialized 선택 가능.

3 Builder 패턴 ★

// High-level auto-select
using Mainloop =
  cutlass::gemm::collective::CollectiveBuilder<
    ArchTag,             // SM90
    OperatorClass,       // TensorOp
    ElementA, LayoutA, AlignA,
    ElementB, LayoutB, AlignB,
    ElementAccumulator,
    TileShape_MNK,
    ClusterShape_MNK,
    StageCountAuto,
    KernelScheduleAuto
  >::CollectiveOp;

4 대칭 Epilogue

using Epilogue =
  cutlass::epilogue::collective::CollectiveBuilder<
    ArchTag, OperatorClass,
    TileShape_MNK, ClusterShape_MNK,
    EpilogueTileType,
    ElementAccumulator, ElementCompute,
    ElementC, LayoutC, AlignC,
    ElementD, LayoutD, AlignD,
    EpilogueScheduleAuto,
    FusionCallbacks        // EVT 연결
  >::CollectiveOp;

5 TileScheduler 선택

6 Kernel 조립

using GemmKernel =
  cutlass::gemm::kernel::GemmUniversal<
    ProblemShape,
    Mainloop,
    Epilogue,
    TileScheduler
  >;

using Gemm =
  cutlass::gemm::device::GemmUniversalAdapter<
    GemmKernel>;

7 장점 요약

• 2.x 템플릿 대비 1/3 수준 boilerplate
• CuTe layout이 warp·thread까지 관통
• 새 기능(WGMMA·TMA·cluster) 1곳에 캡슐화
• autoselection: Auto tag

1 PipelineTmaAsync 개념 ★

stage:  0   1   2   ...  S-1   (ring)
 full :  F   E   E         E   (producer arrive)
 empty:  E   F   F         F   (consumer arrive)

 producer: wait(empty) → TMA → arrive(full)
 consumer: wait(full)  → wgmma → arrive(empty)

2 mbarrier 짝 full/empty

mbarrier 상세는 ↗ V04 §7.

3 Stage count 선택

4 SMEM 한계와 균형

• S ↑ → latency hide ↑ · smem ↑
• S ↑ → occupancy ↓ 가능 (CTA/SM ↓)
• heuristic: ↗ §17 p18

5 Producer / Consumer 분리

setmaxnreg PTX · producer 레지스터 해방 → consumer 더 받음.

6 Async proxy fence

• WGMMA는 async proxy: 일반 SMEM write와 visibility 별개
• producer의 TMA 완료 → mbarrier arrive 는 fence 역할
• consumer는 wait 후 wgmma 가능
• fence.proxy.async.shared::cta 필요 케이스 주의

7 공용 Pipeline 변형

1 왜 2 consumer WG인가

• WGMMA는 async → 발행 후 return 즉시
• 한 WG만 있으면 accumulator 비어있는 중 epilogue 대기
• 2 WG이 교대 또는 분할하면 compute hole 제거

2 Pingpong 구조 ★

Time →
WG1: [load+mma Mi   ][epilogue ]
WG2:           [load+mma Mi+1][epilogue]
WG1:                      [load+mma Mi+2][epi]
WG2:                                 [mma Mi+3][epi]
         ↑ 교대: 서로 다른 M tile

epilogue와 mainloop가 다른 WG에서 겹침. 큰 epilogue에 유리.

3 Cooperative 구조 ★

Time →
WG1: [mma Mi, N0..N/2-1    ][epi L]
WG2: [mma Mi, N/2..N-1      ][epi R]
     ↑ 같은 M tile을 N 축으로 분할
     acc 합쳐서 epilogue 동시

같은 M tile 처리. 큰 N에 유리 (N/2씩 담당).

4 비교 표 ★

5 OrderedSequenceBarrier

// WG1 먼저 epi, WG2 나중
osb.wait(tok_epi, wg_id);
epilogue(acc);
osb.arrive(tok_epi);

6 선택 결정 트리 ★

시작
 │
 ├─ N < 128 ?  ──Y──► Pingpong
 │     (큰 epi가 relative 비용 ↑)
 │
 ├─ epilogue gemma fuse 복잡?
 │     ──Y──► Pingpong (WG별 독립)
 │
 ├─ N ≥ 256 & 작은 batch?
 │     ──Y──► Cooperative
 │
 └─ default ──► Cooperative
               (CUTLASS 권장)

7 주의 ★

1 Epilogue의 자리 α·AB+β·C→fuse

• 입력: FP32 acc fragment
• 출력: dtype_D (BF16/FP8/...)
• fuse가 별도 kernel launch 제거

2 표준 chain

acc (FP32)
  │ · α
  ▼
  + β · C   (load gmem)
  │
  ▼
  + bias    (load broadcast)
  │
  ▼
  activation (ReLU/GELU/Sigmoid)
  │
  ▼
  quantize/cast → dtype_D
  │
  ▼
  store D (gmem)

3 일반적인 fused ops 표 ★

Fusion 일반론은 ↗ V13 §8. 여기선 CUTLASS epilogue 한정.

4 CTA 내 flush 단계

1. acc fragment → smem 전송 (stride 재정렬)
2. smem → reg 재로드 (output layout)
3. op chain 적용
4. gmem store (TMA_STORE or vector st)

smem-staged epilogue가 output coalescing과 layout 재배열에 유리.

5 Register vs SMEM epilogue

6 순서 규칙

• residual add도 FP32 또는 BF16 (dtype 맞출 것)
• quantize scale은 ReLU/GELU 후

7 FP8 GEMM의 epilogue

acc(FP32)
  ├ · alpha (FP32)
  ├ + bias (FP32)
  ├ activation
  ├ · output_scale (FP32)  ← per-tensor
  └ cast → E4M3 (clamp ±448)

amax history update도 epilogue 한 노드.

1 원칙 ★

• 큰 tile = HBM 재사용 ↑, SMEM 압박 ↑, 작은 GEMM 맞추기 어려움
• 작은 tile = 작은 GEMM 유연, 재사용 ↓
• "M·N가 tile의 배수가 되는 최소 tile" 선호
• stage ≥ 3 기본 (H100 5~7)

2 A100 (sm_80) 권장 표 ★

3 H100 (sm_90a) 권장 표 ★

4 Cluster shape 규칙

5 M, N 크기별 휴리스틱 ★

6 L2 capacity 고려

같은 B가 여러 CTA 재사용 되면 L2 hit 기대. tile 순서 (swizzle scheduler)가 이 hit에 영향.

7 실전 결정 흐름

1. M, N, K 구하기
2. tile 후보 2~3개 고르기
   · H100: (128,256,64), (128,128,64), (64,256,64)
3. cluster (1,1) or (2,1)
4. SMEM budget 확인 → stage 조정
5. Coop vs Pingpong
6. autotune 시 이 4~6 조합 시도

1 Batched GEMM 정의

iter b = 0..B-1 :
   D[b] = α·A[b]·B[b] + β·C[b]

2 Grouped GEMM 정의 ★

MoE 맥락 상세는 ↗ V08 §3·§4.

3 자료구조 비교 표 ★

4 GroupedGemm kernel 구조

// host
problem_sizes = [(M0,N0,K0), ...];
ptr_A = [A_0, A_1, ...];
ptr_B, ptr_C, ptr_D = ...;
ld_A, ld_B, ld_C, ld_D = ...;

// kernel 내
int g = tile_scheduler.group_id();
auto [M,N,K] = problem_sizes[g];
run_gemm(ptr_A[g], ptr_B[g], ..., M,N,K);

5 TileScheduler 분배

• 각 group의 tile 수 = ⌈M/bM⌉·⌈N/bN⌉
• 전체 tile을 sequential ID로 매핑
• persistent kernel이 tile → group 역매핑
• 작은 group이 많으면 tail loss ↓

6 이점 · 한계

7 권당 한계

1 Mixed-dtype의 의미

• FP8(E4M3) × BF16 — decode path
• INT4 × FP16 — W4A16 (AWQ/GPTQ)
• FP4 × BF16 — Blackwell 추론

양자화 알고리즘 상세는 ↗ V10 §5·§6·§7.

2 수식 흐름

3 양자화 granularity 표 ★

4 Dequant 삽입 위치

SMEM (quant)
  │
  │ ldmatrix → reg (quant)
  ▼
  dequant(scale, zp) ★  ← 여기 (inline, before mma)
  │
  ▼
 reg (bf16 / fp16 / fp8)
  │
  ▼
 mma(acc, Â, B)

scale fetch는 별도 TMA/cp.async로 prefetch.

5 Scale layout 이슈

• per-group scale: (K/G, N) tensor · N축 broadcast
• SMEM에서 K 단위로 묶어 저장 → iter마다 1회 load
• swizzle 불필요 (작은 tile)
• scale dtype: E8M0 (MXFP8 exp) / FP8 / FP16 다양

6 INT4 packing

packed (uint8):  0xBA
                   │└── low nibble: A (int4)
                   └── high nibble: B (int4)

7 FP8 GEMM 특이점 ★

FP8 bit 상세 ↗ V09 §4.

1 EVT의 동기

• LinearCombination 만으로 안 되는 복합 fusion
• multi-output (D, activation stats 등)
• auxiliary load (bias, residual, scale)

2 3 종 노드

3 Tree 구조 예 ★

          Store(D)
            │
         Compute(cast→D)
            │
         Compute(scale_D ·)
            │
         Compute(GELU)
            │
         Compute(+)
          ┌─┴─┐
       Compute    Load(bias)
        (·α)
          │
        acc         

4 DSL 표기

// pseudo-syntax
using EVT = Sm90EVT<
  Sm90Compute<cast, D_t, F32>,
  Sm90EVT<
    Sm90Compute<mul, F32, F32>,
    Sm90EVT<
      Sm90Compute<gelu, F32, F32>,
      Sm90EVT<
        Sm90Compute<add, F32, F32>,
        Sm90ScalarBroadcast<alpha>,
        Sm90RowBroadcast<bias>
      >
    >,
    Sm90ScalarBroadcast<scale_D>
  >
>;

5 자주 쓰는 사전 조합

6 FP8 fused scale 예

// 입력: FP8 acc, 출력: FP8 + amax_next
using EVT_FP8 = Sm90EVT<
  Sm90Compute<cast, E4M3, F32>,
  Sm90EVT<
    Sm90Compute<mul, F32, F32>,  // · scale_D
    Sm90Compute<amax_update,     // side-effect
                F32, F32>
  >>;

7 제약

• tree는 compile-time 고정
• 동적 op 선택 시 다중 kernel 생성
• aux load는 SMEM 증가 → stage 재계산
• 노드 간 dtype 일관성은 수동 cast 명시

1 Layout 조작 10선 ★

2 수식 상자

3 CUTLASS 3.x 빌더 template ★

using Mainloop = CollectiveBuilder<
  SM90, TensorOp,
  half, RowMajor, 8,
  half, ColMajor, 8,
  float,
  Shape<_128,_256,_64>,
  Shape<_2,_1,_1>,
  StageCountAuto,
  KernelTmaWarpSpecializedCooperative
>::CollectiveOp;

using Epilogue = EpilogueBuilder<
  SM90, TensorOp,
  Shape<_128,_256,_64>,
  Shape<_2,_1,_1>,
  EpilogueTileAuto,
  float, float,
  half, RowMajor, 8,
  half, RowMajor, 8,
  EpilogueScheduleAuto,
  LinearCombination
>::CollectiveOp;

using Kernel = GemmUniversal<
  Shape<int,int,int>,
  Mainloop, Epilogue,
  PersistentScheduler>;

4 5-Level D·K·C·W·T Device·Kernel·CTA·Warp·Thread

Device  M·N·K      gemm_universal
 Kernel  grid      TileScheduler
  CTA    bM·bN·bK  mainloop · stage
   Warp  wM·wN·wK  (wg)mma iter
    Thread fragA/B/C reg

5 스케줄 결정 1-liner

• 큰 N·작은 GEMM수 → Cooperative
• 큰 epilogue / 작은 N → Pingpong
• skinny M·N → StreamK
• default → Persistent + Coop

6 디버그 경로

1. print_layout / print_tensor
2. compile-time assert: shape / alignment / dtype
3. silent wrong result → swizzle 조합 확인
4. hang → mbarrier expected_bytes / arrive count

7 함정 체크