Hopper PTX & SASS 단권화

1 왜 Hopper가 ISA를 늘렸나 비동기의 정식화

• sm_90: Hopper 공통 feature
• sm_90a: architecture-specific (TMA/WGMMA). forward-compat 포기 대신 full 기능

2 신기능 5종 TWC·mD

1. TMA — Tensor Memory Accelerator. 1-thread가 tile copy 발행
2. WGMMA — warp-group async MMA (128T = 4 warp)
3. Cluster — block들의 집합, 같은 GPC에 co-schedule
4. DSM — Distributed Shared Memory (cluster 내 remote smem)
5. mbarrier — smem 64-bit HW barrier (phase bit)

3 비동기 ISA 지도 ★

        Ampere              Hopper
 copy    cp.async (warp)     cp.async.bulk.tensor (1T)
 MMA     mma.sync (sync)     wgmma.mma_async (async)
 sync    bar.sync            mbarrier.{arrive|wait}
 scope   block               cluster (block ×N)
 shmem   local only          + remote (DSM)

공통: cp.async·mma.sync는 Hopper에서도 유효 ↗ V03 §6·§7

4 Producer–Consumer 계층

producer WG0  →  [sA,sB] smem  →  consumer WG1
   TMA issue     mbarrier full       WGMMA async
      ↑______ mbarrier empty ________|

5 도입 이유 3가지

출처: GH100 Whitepaper §3, PTX ISA 8.x §9.7

6 Async Proxy의 등장

• generic proxy: 일반 ld/st
• async proxy: TMA/WGMMA가 쓰는 경로
• 두 proxy 사이에 reorder 가능 → fence.proxy.async 필요

7 Out-of-Scope

• SASS 전수 opcode — 핵심 7개만
• 실전 decompile 감각 — 코드 훈련 영역
• CUTLASS 내부 레이아웃 ↗ V06
• FA3 스케줄 디테일 ↗ V07 §9·§10

1 Cluster 정의 ★ block의 block

• 최대 크기: portable 8 / non-portable up to 16 blocks
• CTA 간 cluster.sync로 집합적 동기화
• DSM으로 remote smem 접근 가능 ↗ §3 p.4

2 선언 문법

// C++ attribute
__global__ __cluster_dims__(2,1,1)
void kernel(...){ ... }

// PTX directive
.reqnctapercluster 2, 1, 1
// launch API (dynamic)
cudaLaunchKernelEx(&cfg, kernel, args);

cluster shape는 kernel-time 또는 launch-time 결정. cudaFuncSetAttribute(..., NonPortableClusterSize, 1)로 16 unlock.

3 Built-in 좌표

Cooperative Groups: cluster_group cg = this_cluster()

4 PTX instruction

// cluster-wide barrier
barrier.cluster.arrive;
barrier.cluster.wait;
// or fused
barrier.cluster.sync;

// rank query
mov.u32 %r0, %cluster_ctarank;
mov.u32 %r1, %cluster_nctaid.x;

5 Cluster 활용 이유

1. TMA multicast: 1 load → N block shmem 동시 broadcast
2. DSM reuse: 인접 block이 서로의 smem 읽음 → L2 BW 절약
3. Cluster launch control: co-location 보장 → latency 예측

6 전형 cluster shape

출처: CUTLASS 3.x SM90 collective.

7 Launch 조건 배수

1 DSM 정의 ★ remote smem

• address space: .shared::cluster (확장된 smem window)
• block의 smem은 cluster rank 별로 mapped region에 나타남
• HBM 거치지 않음 — L2보다 훨씬 낮은 latency

2 Latency 대략값

정확 수치는 V02 §14 ↗ V02. 여기는 비율만 기억.

3 Address mapping PTX

// local smem ptr → remote smem ptr
// target rank = r
mapa.shared::cluster.u32
  %r_remote, %r_local, %r_rank;
// 이후 일반 ld/st.shared::cluster 사용
ld.shared::cluster.f32 %f0,[%r_remote];
st.shared::cluster.f32 [%r_remote],%f1;

mapa = map address. 단일 instruction, smem offset 보존.

4 Cooperative Groups API

auto cg = cg::this_cluster();
int  myr = cg.block_rank();
int  N   = cg.num_blocks();
float* remote =
  cg.map_shared_rank(local_ptr, peer_rank);
// remote는 유효한 .shared::cluster ptr
cg.sync();   // cluster barrier

5 Atomic도 가능

• atom.shared::cluster.add 등
• cluster scope atomic의 가시성: 같은 cluster 내 thread에 대해 순서 보장

6 DSM 활용 패턴

7 동기화 규칙 ★

8 제한 사항

• cluster 내부만 — cross-cluster는 L2 경유
• cluster resident 동안만 유효 (block eviction 없음)
• 총 smem 사용량 = |cluster| × per-block smem 고려

1 TMA란 ★★ 1T·N-D

• 발행 단위: 1 thread (vs cp.async = 32T)
• 방향: G→S (load) · S→G (store) · S→S (cluster multicast)
• 차원: 1D–5D tensor, box(tile) 단위 copy

2 Tensor Descriptor 구조

3 Descriptor 생성 (host)

CUtensorMap desc;
cuTensorMapEncodeTiled(
  &desc,
  CU_TENSOR_MAP_DATA_TYPE_FLOAT16,
  /*rank=*/2,
  g_ptr, /*dims=*/{M,N},
  /*strides=*/{N*2, 2},
  /*box=*/{BM, BN},
  /*elemStride=*/{1,1},
  CU_TENSOR_MAP_INTERLEAVE_NONE,
  CU_TENSOR_MAP_SWIZZLE_128B,
  CU_TENSOR_MAP_L2_PROMOTION_L2_128B,
  CU_TENSOR_MAP_FLOAT_OOB_FILL_NONE);

128-byte 정렬 · const로 kernel에 넘김 (__grid_constant__ CUtensorMap).

4 Tile mode vs Im2col mode

im2col: descriptor의 interleave·lower/upper offset·kernel strides로 conv 변환.

5 Swizzle 선택 bank conflict 회피

• 32-bank smem과 WGMMA fragment layout 일치
• 잘못 설정 시 silent 성능 저하 (정답은 맞으나 BW 반감)

6 발행 → 완료 흐름도 ★

[thread 0]          [TMA engine]         [mbarrier]
  init desc               ─                init(1)
  expect_tx(bytes)────────────────────────▶ set TX count
  cp.async.bulk.tensor ──▶ fetch tile
                          │
                          ▼
                        write smem
                        arrive.complete_tx──▶ TX count↓
  try_wait.parity ◀──────────────────────── phase flip
  consume smem

7 OOB 처리

• NONE: 호출자가 bounds 보장
• ZERO: OOB 영역 0으로 채움 (attention padding에 유용)
• NaN: float만 해당

1 명령 형식 ★

2 G→S 2D load ★★

// thread 0만 발행
if (threadIdx.x == 0) {
  asm volatile(
    "cp.async.bulk.tensor.2d"
    ".shared::cluster.global"
    ".mbarrier::complete_tx::bytes"
    " [%0], [%1, {%2,%3}], [%4];"
    ::
    "r"(smem_ptr),   // dst
    "l"(tma_desc),   // CUtensorMap*
    "r"(coord_x),
    "r"(coord_y),
    "r"(bar_ptr));   // mbarrier
}

• completion 방식: complete_tx::bytes (mbarrier TX counter에 bytes 추가)
• 다른 옵션: bulk_group (commit_group/wait_group 기반)

3 전체 lifecycle pseudo-code ★★

// === setup (thread 0) ===
if (tid == 0) {
  mbarrier_init(bar, 1);
  fence_mbarrier_init();
}
__syncthreads();

// === phase 0: issue ===
if (tid == 0) {
  mbarrier_expect_tx(bar, BM*BN*2);
  cp_async_bulk_tensor_2d(
      sA, &tma_desc, {cx, cy}, bar);
}

// === wait (all threads) ===
uint32_t phase = 0;
while (!mbarrier_try_wait_parity(
              bar, phase)) { }

// === consume ===
wgmma_fence();
wgmma(acc, sA, sB);
wgmma_commit_group();
wgmma_wait_group(0);

// === next iter: toggle phase ===
phase ^= 1;

4 변종 table

5 Multicast mask

// N-bit mask: cluster rank별 수신 여부
asm("cp.async.bulk.tensor.2d"
    ".shared::cluster.global"
    ".multicast::cluster"
    ".mbarrier::complete_tx::bytes"
    " [%0],[%1,{%2,%3}],[%4],%5;"
    :: ... , "h"(mask));

mask = 0b0011 → rank 0·1 수신.

6 주의점

1 mbarrier 구조 ★

[63]    phase bit
[62:44] expected arrive count
[43:25] current arrive count
[24: 0] pending TX bytes  (TMA용)

정확한 bit layout은 구현 dependent. 의미만 사용.

2 Lifecycle

1. init — expected count 설정 + phase=0
2. expect_tx — TMA가 들어올 byte 수 예약 (optional)
3. arrive — count++ (thread 단위)
4. arrive.expect_tx — count++ & TX 예약
5. try_wait.parity — phase 반전 때까지 spin
6. phase flip: count==expected && TX==0 충족 시 자동

3 PTX 명령 표

4 기본 사용 패턴

// init (1 thread)
asm("mbarrier.init.shared.b64"
    " [%0], %1;"
  :: "r"(bar), "r"(N_ARRIVE));
fence_mbarrier_init_release();

// arrive (각 thread)
asm("mbarrier.arrive.shared.b64"
    " _, [%0];"
  :: "r"(bar));

// wait
uint32_t p = 0;
while(!try_wait_parity(bar, p)){}
p ^= 1;

5 Phase bit 개념 ★

iter 0:  init→arrives→flip(0→1)  wait p=0
iter 1:      arrives→flip(1→0)  wait p=1
iter 2:      arrives→flip(0→1)  wait p=0
...  phase ^= 1 매 iter

6 짝짓기 규칙

7 함정

1 warpgroup 정의 ★

• Block당 warpgroup 수 ≤ 4 (typical 2: producer + consumer)
• 모든 128T가 같은 wgmma.mma_async 호출 — lane divergence 금지

2 Async 의미론

wgmma.fence         ; 이전 reg writes flush
wgmma.mma_async     ; 발행 후 즉시 다음 명령 가능
wgmma.commit_group  ; 지금까지 발행된 mma 그룹 닫기
wgmma.wait_group N  ; N개 group 이하만 in-flight 유지

commit/wait 쌍으로 pipeline 제어. N=0이면 모든 in-flight 완료 대기.

3 Operand 배치

4 명령 형식 ★★

5 사용 template ★★★

// === setup ===
uint64_t descA = make_desc(sA, stride, swizzle);
uint64_t descB = make_desc(sB, stride, swizzle);
float d[64];   // accumulator fragment

// === mainloop ===
asm("wgmma.fence.sync.aligned;");
asm volatile(
 "wgmma.mma_async.sync.aligned"
 ".m64n256k16"
 ".f32.f16.f16"
 " {%0,%1,...,%127},"   // D: 128 regs
 " %128, %129,"            // descA, descB
 " %130, 1, 1, 1, 1;"      // scaleD,A,B,tnsA,tnsB
 : "+f"(d[0]), "+f"(d[1]), ...
 : "l"(descA), "l"(descB),
   "r"(scaleD));
asm("wgmma.commit_group.sync.aligned;");
asm("wgmma.wait_group.sync.aligned 0;");
// d[] 이제 consumption 가능

6 Matrix descriptor bit layout ★

[63:62] base offset       (smem alignment)
[61:49] leading dim BO    (byte offset between 8·k tiles)
[48:46] swizzle           (0=none,1=32B,2=64B,3=128B)
[45:32] stride dim BO     (byte offset between groups of 8 rows)
[31:14] reserved
[13: 0] start address     (smem offset / 16)

CuTe의 make_gmma_desc가 자동 생성. 수동 생성은 silent 오답 유발.

7 scale·tns 파라미터

8 기존 mma와 차이

1 Shape 축 의미

D[m64, nN] = A[m64, kK] · B[kK, nN]
     ↑             ↑              ↑
  accumulator   SM-input-A   SM-input-B (smem)

• m = 64 고정 (warpgroup 단일 축)
• N = 8, 16, 24, ..., 256 (8 단위)
• K는 dtype별로 결정

2 N 값 허용 셋

출처: PTX ISA 8.x §9.7.14.

3 권장 N (H100 SXM5)

N 클수록 발행 횟수 ↓ → issue overhead ↓.

4 dtype × K 완전표 ★★★

FP8: E4M3=더 정밀·낮은 range, E5M2=높은 range·낮은 정밀 ↗ V09 §5.

5 Accum register 수

128T × 128 reg = 16K reg/WG → consumer에 몰아주기 필요.

6 Register budget 제약

7 Sparsity variant

• wgmma.mma_async.sp.sync.aligned
• A가 2:4 sparse → 실효 K 2배
• metadata: 2-bit per 4 elem, smem에 동반 load

1 WS 정의 ★ 역할 분담

• Cooperative Groups 모델에서 벗어나 imperative async 설계
• 목적: WGMMA를 TMA·softmax와 시간적으로 overlap

2 WG 분할 전형

E=epilogue. FA3 커널은 1P+2C. CUTLASS GEMM은 config 선택.

3 Producer / Consumer pseudo-code ★★★

// === shared mbarriers ===
__shared__ uint64_t bar_full[S];
__shared__ uint64_t bar_empty[S];

// === init (1 thread) ===
if (tid == 0) {
  for (i=0; i<S; i++){
    mbarrier_init(&bar_full[i],  1);
    mbarrier_init(&bar_empty[i], 2);  // 2 cons WG
  }
  fence_mbarrier_init_cluster();
}
__syncthreads();

// === PRODUCER warp-group (WG0) ===
if (wg_idx == 0) {
  setmaxnreg_dec(40);   // 240 → 40
  uint32_t p = 0;
  for (k=0; k<K_TILES; k++){
    int slot = k % S;
    mbarrier_wait_parity(
        &bar_empty[slot], p);
    if (lane==0) {
      mbarrier_expect_tx(
          &bar_full[slot],
          BM*BK*2 + BK*BN*2);
      cp_async_bulk_tensor(
          &sA[slot], &desc_A,
          {k*BK, m0}, &bar_full[slot]);
      cp_async_bulk_tensor(
          &sB[slot], &desc_B,
          {n0, k*BK}, &bar_full[slot]);
    }
    if ((k%S) == S-1) p ^= 1;
  }
}

4 Consumer pseudo-code ★★★

// === CONSUMER warp-groups (WG1, WG2) ===
else {
  setmaxnreg_inc(232);  // 40 → 232
  float acc[128] = {0};
  uint32_t p = 0;
  for (k=0; k<K_TILES; k++){
    int slot = k % S;
    mbarrier_wait_parity(
        &bar_full[slot], p);
    wgmma_fence();
    wgmma_mma_async(
        acc,
        make_desc(&sA[slot]),
        make_desc(&sB[slot]));
    wgmma_commit_group();
    wgmma_wait_group(1);  // 1 in-flight
    mbarrier_arrive(
        &bar_empty[slot]);
    if ((k%S) == S-1) p ^= 1;
  }
  wgmma_wait_group(0);
  store_epilogue(acc);
}

5 패턴 비교

상세 trade-off ↗ V06 §15, ↗ V07 §10.

1 Proxy 개념 ★

thread  ──generic──▶ smem/L2/HBM
TMA     ──async ────▶ smem/L2/HBM
WGMMA   ──async ────▶ RF/smem(read)

2 Proxy fence PTX

3 언제 쓰는가

• TMA write → WGMMA read: TMA complete_tx가 함께 proxy fence 역할
• TMA write → thread read(ld.shared): 명시 fence 필요
• WGMMA write RF → epilogue: wgmma.wait_group이 release 역할

4 setmaxnreg 정의 ★

setmaxnreg.dec.sync.aligned.u32 %R;  // reduce to R
setmaxnreg.inc.sync.aligned.u32 %R;  // raise to R

5 허용 값 / 규칙

PTX ISA 8.x. __nv_cvta_set_maxnreg_dec_sync_aligned_u32 builtin.

6 재분배 전형 값 ★

FA3 (Hopper) 기본값: prod 40, cons 232 (H100 SM).

7 호출 위치 규칙

• kernel 초반, role-분기 직후 즉시 호출
• mainloop 진입 전 완료
• 중간 변경 금지 (deadlock risk)

if (wg_idx == 0) {
  setmaxnreg_dec(40);
  producer_loop();
} else {
  setmaxnreg_inc(232);
  consumer_loop();
}

8 launch_bounds 와의 관계

9 제한

• setmaxnreg는 WG 전원이 aligned — divergent call 불가
• cluster launch 시 각 block별 독립 적용
• debug: cuobjdump --dump-resource-usage로 per-kernel reg 확인

1 DPX 동기 DP 가속

• Hopper 신설. 대상: bioinformatics / graph / RL path
• 스루풋: Ampere 대비 약 7× (vendor 공식 자료)

출처: GH100 Whitepaper §3.3, CUDA Math API.

2 3-operand min/max

3 복합 relu+add+max

// y = max(x+bias, 0) + c, 단일 inst
asm("viaddmax.s32.relu"
    " %0, %1, %2, %3;"
  : "=r"(y)
  : "r"(x), "r"(bias), "r"(c));

// 3-arg max with saturate
asm("vimax3.s32"
    " %0, %1, %2, %3;"
  : "=r"(m)
  : "r"(a), "r"(b), "r"(c));

4 SW 업데이트 예

5 CUDA C intrinsic

unsigned, signed, packed 버전 합쳐 약 20개.

6 활용 영역

• Bioinformatics: Smith-Waterman, DP 정렬
• Robotics: path planning, dynamic programming
• RL: value iteration
• Graph: Floyd-Warshall all-pairs shortest

7 한계

1 SASS 정의 ★ 실제 opcode

• 각 세대마다 encoding 상이 (sm_80 ≠ sm_90)
• JIT: PTX → ptxas → SASS
• Reverse: nvdisasm, cuobjdump

2 변환 파이프라인

.cu    ─[nvcc front]─▶ .ptx
.ptx   ─[ptxas]──────▶ .cubin (SASS + metadata)
.cubin ─[driver JIT]─▶ (native) → launch

embedded PTX는 driver가 최종 SM에 대해 runtime compile (fat binary).

3 세대별 encoding

출처: Jia et al. 2018 / Dissecting Volta · Ampere reverse-engineering 논문.

4 왜 SASS를 읽나

1. register spill 확인 (LDL/STL 존재)
2. stall 이유 추적 (control code)
3. instruction mix 분석 (FFMA vs HMMA 비율)
4. compiler 스케줄링 의도 역추론

5 nvdisasm 기본

# cubin에서 SASS 덤프
nvdisasm --print-code kernel.cubin

# control code 함께
nvdisasm --print-instruction-encoding \
         --print-line-info kernel.cubin

# cuobjdump from fatbin
cuobjdump --dump-sass a.out

6 PTX ↔ SASS 관계 ★

7 SASS 기본 형식

  [label]   OPCODE.MOD  [P?] dst, src1, src2 [, src3] ;
  예:
  /*0010*/  FFMA R3, R2, R5, R3 ;
  /*0020*/  @P0 BRA  0x50 ;
  /*0030*/  STG.E [R6], R4 ;

• @P0: predicate (0~6)
• .E: effective addr (64-bit pointer)
• 레지스터: R0..R255 + RZ(=0), URx(uniform)

8 Hopper 특유 SASS

opcode 이름은 nvdisasm 버전에 따라 약간 다를 수 있음.

9 범위 밖 의도적 제외

• 전수 opcode 리스트 — 공식 미공개
• 실제 cubin 디컴파일 훈련 — 코드 영역
• 세대별 encoding bit 분석

1 산술 ★

2 Memory

3 Tensor Core ★

HMMA는 warp-level, GMMA는 warpgroup-level async.

4 Control flow

5 Sync / Atomic

6 대표 modifier

7 실행 예 (Hopper GEMM 발췌)

/*0100*/ IMAD.WIDE R4, R3, 0x4, R0 ;   // addr calc
/*0108*/ ULDC.64 UR4, c[0x0][0x160] ; // desc
/*0110*/ LDS.128 R8, [R4+0x0] ;        // smem load
/*0118*/ GMMA.64x256x16.F32.F16 R16, desc_A, desc_B, R16 ;
/*0120*/ GMMA.FENCE ;
/*0128*/ GMMA.COMMIT_GROUP ;
/*0130*/ GMMA.WAIT_GROUP 0x0 ;
/*0138*/ STG.E.128 [R20], R16 ;

주소·opcode는 예시용. 실제 nvdisasm 출력과 다를 수 있음.

8 Spill 진단

1 Control code 개요 ★

128-bit SASS inst = [64b opcode] [64b operand/ctrl]
       control field ≈ 16-bit subset

2 구성 요소

SB = ScoreBoard. 6개 slot. fixed-latency는 stall만, variable(LDG/HMMA)는 SB 사용.

3 의미 해석 ★

• stall = k: 이 inst issue 후 k cycle 뒤 다음 inst (fixed latency 의존성)
• write SB#n: 이 inst 완료 시 SB#n 해제. 해제 전까지 SB#n 보유자는 대기.
• wait 000100: SB#2 기다림. 그 이전 LDG가 완료되어야 진행
• yield: eligible warp 재평가 trigger (occupancy 도움)
• reuse: 직전 inst의 operand register를 operand collector에서 재사용

4 nvdisasm 출력 예

/*0200*/    LDG.E.SB0 R4, [R2] ;       // write SB0
/*0208*/    IADD3 R5, R5, R0, RZ ;     // stall=1
/*0210*/  @SB0 FFMA R6, R4, R7, R6 ;   // wait SB0
/*0218*/    FFMA.REUSE R8, R4, R9, R8; // reuse R4

@SB0 표기는 툴에 따라 WAIT(0) 등으로 나타남.

5 Hopper 추가

• GMMA completion: 별도 GMMA scoreboard
• TMA completion: mbarrier로 signal (SB 미사용)
• GMMA.WAIT_GROUP이 GMMA SB drain

6 Latency 분류 표

7 Stall 값 의미

8 읽기 팁

1. WAIT 표기가 많다 → 메모리 의존성
2. REUSE 표기가 적다 → operand collector 압박
3. YIELD 없음 → occupancy low 의심
4. stall≥8 연속 → fixed-lat 직렬화 (dep chain 길이)

9 프로파일러와 관계

Nsight Compute의 stall reason 분류는 이 control code를 집계한 것. ↗ V18 §4.

1 Dual Issue 정의 ★

• Ampere: 1 inst/cycle/scheduler × 4 scheduler
• Hopper: 2 inst/cycle/scheduler 가능 (dual issue) × 4 scheduler GH100 wp §3.2

2 쌍 조건 표

실제 규칙은 compiler 전담. 사용자는 관찰만.

2 Scoreboard 개념 ★

warp scoreboard:  [SB0 SB1 SB2 SB3 SB4 SB5]
                    ▲                   ▲
                    |                   |
                 LDG.E → set SB0     HMMA → set SB4
                 @SB0 FFMA waits for SB0 == 0

3 할당 규칙

• compiler가 inst 발행 시 SB# 선택 (write barrier)
• 같은 SB# 재사용 전에 반드시 해제
• 과도 할당 → stall 누적, 과소 할당 → false-dep

4 Hopper SB 분리

5 Read-after-Write 분석

6 Compiler scheduling 영향

• ptxas의 instruction scheduling pass가 write/wait 배치 결정
• -Xptxas -O3: aggressive reorder
• #pragma unroll / loop hoist가 체인 길이 결정
• inline asm 직후 : 인접 SASS 순서 제약 가능

7 Pipeline hazard 유형

8 Issue slot 최적화 요점

1 툴 표 ★

2 기본 워크플로

# 1. arch 지정 컴파일 + lineinfo
nvcc -arch=sm_90a -lineinfo \
     -Xptxas -v -o k.cubin k.cu

# 2. PTX
nvcc -arch=sm_90a -ptx k.cu -o k.ptx

# 3. SASS + line
nvdisasm --print-line-info \
         --print-instruction-encoding \
         k.cubin > k.sass

# 4. PTX ↔ SASS side-by-side
diff -y k.ptx k.sass | less

3 PTX ↔ SASS 대응 패턴

4 주석 달기 규칙

1. 범위 표시: mainloop 시작/끝 라벨
2. 역할 표기: producer/consumer 섹션 분리
3. SB 지도: 각 SB#가 어느 load를 커버하는지
4. spill flag: LDL/STL에 ! 표시

5 진단 체크리스트

6 Loop identify

.L_LOOP_START:
  @SB0 LDS ... ;
       GMMA ... ;
       GMMA.COMMIT_GROUP ;
       GMMA.WAIT_GROUP 1 ;
       MBARRIER.ARRIVE ... ;
       BRA .L_LOOP_START ;

mainloop는 보통 BSSY/BRA 쌍 + stage pipeline. MBARRIER가 경계.

7 Source↔SASS mapping

• -lineinfo 필수 (-G는 debug, 성능 왜곡)
• Nsight Compute의 Source View가 같은 정보
• 한 source line ↔ 다수 SASS (N:M)

8 종료 체크

1 PTX 추가분 5종 1-liner ★

2 TMA 템플릿

init(bar, 1); fence_init();
expect_tx(bar, BYTES);
cp.async.bulk.tensor(dst, desc, {x,y}, bar);
while(!try_wait_parity(bar, p)){}
p ^= 1;

3 WGMMA 템플릿

wgmma.fence;
wgmma.mma_async.m64n256k16.f32.f16.f16
      {D[0..127]}, descA, descB,
      scaleD, 1, 1, 1, 1;
wgmma.commit_group;
wgmma.wait_group 0;

4 WS skeleton

if (wg == 0) {      // producer
  setmaxnreg_dec(40);
  for(k : K_TILES) {
    wait_empty(k);
    tma_load_A(k);
    tma_load_B(k);   // → bar_full arrive
  }
} else {               // consumer
  setmaxnreg_inc(232);
  for(k : K_TILES) {
    wait_full(k);
    wgmma(acc, sA[k], sB[k]);
    wgmma_wait(1);
    arrive_empty(k);
  }
  epilogue(acc);
}

5 WGMMA shape 요약

6 전형 reg 분배

7 SASS 핵심 10선 ★

8 함정 7선 ★

1. -arch=sm_90 ≠ sm_90a (TMA/WGMMA 미생성)
2. mbarrier init 후 fence 미호출 → UB
3. TMA descriptor 128B 미정렬 → crash
4. cluster shape 배수 조건 위반 → launch 실패
5. WGMMA B operand RF에 두면 컴파일 에러
6. setmaxnreg 총합 초과 → launch 실패
7. async proxy fence 누락 → silent 잘못된 결과

9 한 장 요약

관련 권: V02 HW 정량, V03 Ampere PTX, V06 CUTLASS, V07 FA3. ↗ V02·V03·V06·V07