GPU Architecture Quantitative

1 SM = 4 processing partition Ampere/Hopper 공통 4 quad

• partition당 warp ≤ 16 (GA100: 64 warp/SM · 16×4)
• partition 내 warp scheduler = 1 cycle · 1 warp 발행
• partition 간 공유: L1/Smem, TC, instruction cache, constant cache

출처 NVIDIA GA100 WP v1.0 Fig.5 / GH100 WP v1.04 Fig.7

2 Partition 내부 unit

출처 GA100 WP Fig.6 · GH100 WP §3.1

3 SM 블록도 ASCII

┌──────────────── SM ────────────────┐
│ L1 I-Cache · const cache (shared)  │
├────┬────┬────┬────┐                │
│ P0 │ P1 │ P2 │ P3 │   4 partitions │
│ ┌──┴──┐ ...                         │
│ │warp │ scheduler · dispatch        │
│ │sched│ 1 warp/cyc                  │
│ └──┬──┘                             │
│ ┌──┴──┐                             │
│ │ RF  │ 16K × 32b (partition)       │
│ └──┬──┘                             │
│ ┌──┴──┐                             │
│ │ OC  │ operand collector            │
│ └──┬──┘                             │
│ ┌──┴──────────────────┐              │
│ │FP32·INT·FP64·SFU·LSU│              │
│ │ Tensor Core (4th)   │              │
│ └──┬──────────────────┘              │
├────┴──────── L1/Smem 228 KB ───────┤
│ Tex unit · RT Core (GA10x/AD10x)   │
└────────────────────────────────────┘

4 Dual-issue 조건 partition 1 cyc 2 이슈

• 한 warp는 1 instruction/cyc 만 dispatch
• partition 단위로 서로 다른 pipe에 나눠 이슈 가능 (FP32 + LSU 등)
• compute capability 8.0+: dependency/throughput 조건 충족 시 inst pair

출처 CUDA C PG 12.6 App. H "Compute Capability 8.0/9.0"

5 Instruction / Constant Cache

GA100/GH100 동일 계층. 수치는 CUDA C PG App.H 추정 상한.

6 RT Core 참고 (DC용 GA100/GH100에는 없음)

• BVH traversal + triangle-ray intersect HW
• Ada(RTX 40): 3rd-gen RT, opacity micromap
• DC 카드(A100·H100)에는 없음 ↗ V02 §17

7 핵심 공식: dispatch 상한

1 Register File 용량 ★ 65536 × 4B

출처 CUDA C PG 12.6 Table 21 "Technical Specifications per Compute Capability"

2 Partition 내 sub-bank

• 256 KB RF = 4 partition × 64 KB × (multiple bank)
• register 분할: sub-bank에 RR(round-robin) 배치
• 동일 cycle 에 같은 bank 복수 read → bank conflict

3 Operand Collector 정의

4 FFMA 3-operand 수집 흐름

cycle t:    FFMA r0, r1, r2, r3
             │
  ┌──────────┴──────────┐
  │   Operand Collector │
  │   slot: A B C accum │
  │   ↑ ↑ ↑             │
  └───┬──┬──┬──┬────────┘
      │  │  │  │
     RF RF RF RF  (bank select cycle-staggered)
      │  │  │  │
      └──┴──┴──┴─→ ALU FP32 FMA

개념도 · 출처: Greg Smith GTC 2010 "Fermi SM" · Volta 이후 구조 동일.

5 Register Spill 비용 ★

출처 CUDA C PG App.H "Local Memory" + 표준 micro-benchmark 권장값.

6 Occupancy 제약 공식

출처 CUDA Occupancy Calculator (deprecated) 동일 식 · PG Table 21.

7 Hopper setmaxnreg ★

• warpgroup 내 warp별 register budget 동적 재분배
• setmaxnreg.inc.sync.aligned 232: consumer warp이 큰 reg
• setmaxnreg.dec.sync.aligned 40: producer warp은 최소

8 실무 가이드

1. -maxrregcount or __launch_bounds__로 상한 지정
2. -Xptxas -v로 spill 확인 (stack frame bytes)
3. spill ≠ 항상 나쁨. occupancy 회복이 더 이익이면 ok

1 L1/Smem 통합 ★ unified SRAM

출처 GA100 WP Table 2 · GH100 WP §3.2 · GB100 WP v1.0 §3.1

2 Carve-out 옵션 (GA100)

• 0 / 8 / 16 / 32 / 64 / 100 / 132 / 164 KB smem
• 나머지가 L1 data cache
• cudaFuncSetAttribute(..., cudaFuncAttributePreferredSharedMemoryCarveout, %)

출처 CUDA C PG 12.6 §3.2.5 / GA100 WP Table 3

3 Smem Bank 구조 ★

• 32-bit mode: word = 4B, bank = (addr/4) mod 32
• 64-bit mode: double 용. cudaDeviceSetSharedMemConfig
• broadcast: 동일 주소 read는 1 cycle에 모두 제공

4 Hopper swizzle WGMMA 용

• TMA load 시 128B / 64B / 32B swizzle mode 선택
• WGMMA matrix descriptor에 swizzle 종류 포함
• bank conflict-free access 를 HW 가 보증

상세 PTX는 ↗ V04 §4·§7

5 Distributed Shared Memory ★ DSM = cluster smem

cluster (1..16 CTA, default ≤8)
 ├── CTA0.smem 228 KB ─┐
 ├── CTA1.smem 228 KB ─┤  unified
 ├── CTA2.smem 228 KB ─┼── logical
 └── CTA3.smem 228 KB ─┘  space
addr = (rank, offset)

6 DSM access 경로

• cluster.map_shared_rank (PTX) → remote ptr
• SM-to-SM network (GPC 내부) 로 이동
• read latency ≫ local smem, but ≪ L2

출처 GH100 WP §3.3 "Thread Block Cluster" · PTX ISA 8.x §9.7.12

7 Smem 사용 정석 3종

1. tile buffer (GEMM/conv) — V01 §7, V06
2. reduction scratchpad — V01 §11
3. producer/consumer ring — ↗ V04 §9

1 L2 용량 · 분할 ★

출처 GA100 WP Table 2 · GH100 WP §3.4 · GB200 WP §2.1 (die partitioned)

2 L2 partition 구조 ★ NUMA-like

• GH100: partition A ↔ B link BW ≈ L2 BW 의 부분값
• Blackwell: chiplet 2 die 간 NV-HBI 10 TB/s (공식)

출처 GH100 WP Fig.10 · GB200 WP "NV-HBI" 항목

3 L2 Persistence Window ★

// host
cudaStreamAttrValue attr = {};
attr.accessPolicyWindow.base_ptr = p;
attr.accessPolicyWindow.num_bytes = N;
attr.accessPolicyWindow.hitRatio = 1.0f;
attr.accessPolicyWindow.hitProp =
  cudaAccessPropertyPersisting;
attr.accessPolicyWindow.missProp =
  cudaAccessPropertyStreaming;
cudaStreamSetAttribute(s,
  cudaStreamAttributeAccessPolicyWindow,&attr);

출처 CUDA C PG 12.6 §5.2.6 "L2 Cache Set-Aside"

4 set-aside 카빙

cudaDeviceGetLimit(cudaLimitPersistingL2CacheSize)

5 Compression 무손실 CCS

• GA100+: L2 내 compressible sector HW 압축
• 동일 byte 패턴·제로런 중심 → 효과는 데이터 의존
• 사용자 투명. PTX 에 별도 지시 없음

출처 GA100 WP §2.4 "Compute Data Compression"

6 L2 위치가 중요한 이유

• GEMM: reuse → L2 hit ≈ tile 크기 의존
• attention KV cache: 크기가 L2 > 이면 miss

7 흔한 실수

• persistence window 가 커널 종료 후에도 남는 것 잊음 → clear 필요
• window 범위를 L2 용량보다 크게 설정 → hitRatio 낮춰도 효과 반감

1 HBM 세대 비교 ★ 2e→3→3e

출처 JEDEC JESD238A (HBM3) · NVIDIA GH100/GB200 WP 공식 수치

2 GPU별 HBM 구성

출처 NVIDIA A100 WP Table 2 · H100 WP Table 2 · H200 datasheet · B200 GTC'24 spec

3 Memory Controller 계층

SM → L2 slice → MC → HBM stack
                   ↑
              per-stack controller
                   ↑
              channel pseudo-channel
                   ↑
              bank group · bank · row

• GH100: 12 × 512-bit HBM3 MC (총 6144-bit)
• MC 단위로 access coalesce 및 row-open 관리

출처 GH100 WP §3.4 "Memory Subsystem"

4 Coalescing 조건 복습

• sector = 32 byte (L2 ↔ HBM 최소 단위)
• 128 B = 4 sector = warp FP32 load 크기

PG §5.3.2 "Device Memory Access"

5 ECC Overhead ★

출처 NVIDIA HBM ECC 기술문서 · nvidia-smi man.

6 GPU Memory Partition =GPC · uGPU

• GA100: 8 GPC × (2 TPC × SM + L2 slice)
• GH100: 8 GPC (enabled 7 on H100 SXM5 : 132 SM)
• MIG 는 GPC + L2 slice + HBM ch 묶음으로 자른다 ↗ V02 §18

7 PCIe / NVLink 외부 링크

출처 A100/H100 WP · GB200 GTC'24 session

1 세대 개요

출처 NVIDIA 공식 세대 표기 (각 arch WP).

2 dtype 지원 매트릭스 ★★ 16→BF→TF→8→4

출처 GV100/GA100/GH100/GB200 WP + PTX ISA 8.x "mma" 허용 dtype 목록

3 Dense TC peak (SXM / W=TDP) 16·19·31

단위 TF/s · dense (2:4 sparsity 아님) · 출처 A100/H100 WP Table 2 · B200 GTC'24 keynote 공식 슬라이드

4 2:4 Sparsity

• A100·H100·B200 지원
• FP8/FP16 모두 적용 가능 (Hopper+)
• 실측 ×1.3~1.5 (×2 이론 상한)

출처 GA100 WP §2.3 "Sparsity"

5 세대별 의미론 변화

6 accumulator 규칙

• FP16/BF16/FP8 → accum FP32
• FP16 × FP16 → FP16 accum 도 있음 (정확도 ↓)
• TF32 → FP32 accum (항상)
• FP64 TC → FP64 accum

PTX ISA 8.x §9.7.13 "Matrix multiply-accumulate"

7 쉐이프·피연산자 배치

• Ampere mma: A·B·C·D 모두 register fragment
• Hopper wgmma: A=reg or smem, B 항상 smem, D=reg
• 자세한 layout ↗ V03 §7 · V04 §7

1 Die 구성 ★

출처 NVIDIA GA100 WP v1.0 §2.1 / Table 2

2 총 연산 자원 (A100 SXM4)

3 메모리 계층 요약

출처 GA100 WP Table 2 · latency는 public micro-benchmark 권장값 ↗ V02 §14

4 Ampere 신기능 ★

• cp.async — G→Smem 직결, reg 경유 제거
• 3rd-gen TC: TF32/BF16/FP64/INT8
• 2:4 Sparsity — ×2 가속
• MIG — 최대 7 instance
• Async copy barrier (mbarrier 초기 버전)
• 3rd-gen NVLink 600 GB/s

출처 GA100 WP §2

5 연산 peak 요약 (SXM4)

출처 A100 WP Table 2 · boost clock 기준

6 A100 variant

NVIDIA A100 Datasheet (2021)

7 GA100 vs GA10x

• GA100: FP64 1:2 of FP32, Smem 164 KB, NVLink
• GA102 (RTX 3090): FP64 1:64, 128 KB L1+Smem, RT Core
• 자세한 분화 ↗ V02 §17

1 Die 구성 ★

출처 NVIDIA GH100 WP v1.04 §2.1 / Table 2

2 총 연산 자원 (H100 SXM5)

3 메모리 계층 요약

출처 GH100 WP Table 2 · latency는 NVIDIA H100 micro-benchmark (nvidia GTC'22)

4 Hopper 신기능 5 ★★ TWCDM

• TMA: Tensor Memory Accelerator
• WGMMA: warp-group async MMA
• Cluster: Thread Block Cluster
• DSM: Distributed Shared Memory
• TE: Transformer Engine (FP8)

상세 ↗ V02 §9 · V04 §1–§7

5 연산 peak (H100 SXM5)

출처 H100 WP Table 2 · SXM5 / boost clock

6 H100 variant

H100 datasheet / H200 datasheet (2024)

7 DPX instruction Dynamic Programming eXt

• vimax3·vimin3: 3-input SIMD max/min + relu
• Smith-Waterman/Needleman-Wunsch 가속 (유전체)
• compare-and-add 결합 감축 1 cycle

출처 GH100 WP §2.6 "DPX"

1 TMA — Tensor Memory Accelerator ★

• descriptor: cuTensorMapEncodeTiled host 생성
• PTX: cp.async.bulk.tensor
• mode: tile / im2col
• OOB 정책: zero / clamp / NaN

출처 GH100 WP §2.4 · PTX 8.x §9.7.8 · 상세 ↗ V04 §4·§5

2 왜 TMA 가 필요한가

3 WGMMA — Warpgroup MMA ★

• shape: m64nNk16 (FP16/BF16), m64nNk32 (FP8), N ∈ {8,16,…,256}
• A: register or smem / B: smem 필수 / D: register
• operand: 64-bit matrix descriptor (smem base+stride+swizzle)
• throughput: H100 1 WGMMA/cyc/SM (FP16 m64n256k16)

출처 PTX ISA 8.x §9.7.13.2 "wgmma" · CUTLASS 3.x Hopper kernel

4 Cluster — Thread Block Cluster ★

• portable 최대 8 block/cluster (opt-in 16)
• __cluster_dims__(x,y,z) 또는 launch attribute
• cluster.sync / cluster.arrive

출처 GH100 WP §3.3 · CUDA C PG §7.27

5 DSM — Distributed Shared Memory ★

• PTX: st.shared::cluster, ld.shared::cluster
• atomic: atom.shared::cluster
• 8 block cluster: 228×8 = ~1.78 MB 공유 가능

출처 PTX ISA 8.x §9.7.12 · GH100 WP §3.3

6 TE — Transformer Engine ★

• E4M3: weight/activation (dynamic range ↓, 정밀 ↑)
• E5M2: gradient (range ↑, 정밀 ↓)
• amax history 로 scale 업데이트
• PyTorch/JAX 통합 라이브러리 제공

출처 GH100 WP §2.3 "Transformer Engine" · NVIDIA TE github docs

7 5 기능 상호관계

TMA  ─────┐         ┌─── smem ─── WGMMA
          ├─(mbar)──┤
producer  │         │   consumer
warps     ↓         ↓   (warp-specialized)
Cluster + DSM: 여러 CTA 간 smem 공유
TE: FP8 layer 주변 wrapper

1 Die 구성 ★

출처 NVIDIA GB200 NVL72 WP · GTC'24 Keynote spec slide

2 메모리

출처 GB200 WP §2 · B200 공식 데이터시트 (2024)

3 신기능 ★ FP4·NV-HBI·TE2

• 5th-gen TC: FP4 E2M1 · FP6 E3M2/E2M3
• Micro-scale block FP: per-block scale factor HW 지원
• 2nd-gen Transformer Engine: FP4/FP6 자동 선택
• NV-HBI 10 TB/s dual-die fusion
• NVLink 5 1800 GB/s / GPU
• 5th-gen RAS: Dedicated RAS engine

출처 GB200 NVL72 WP §2.2 · GTC'24

4 연산 peak (B200 공식)

단위 PF/s (PetaFLOP/s) · 출처 NVIDIA B200 공식 spec sheet · dense = non-sparsity

5 3세대 비교 종합표 ★★ SM·HBM·NVL·TC

출처 A100/H100/GB200 공식 WP · SXM form · boost clock · dense

1 스케줄러 기본 동작 select → issue

• GV100 이후 warp pool: partition 당 최대 16
• 선택 정책: NVIDIA 비공개 (근사 GTO / LRR heuristic)
• queue empty ≠ 항상 bad: bubble 이 다른 warp 으로 메꿔짐

2 eligible 조건

1. PC 의 instruction fetched
2. source operand 준비 (scoreboard clear)
3. destination reg slot free
4. functional unit 사용 가능
5. barrier/mbar 없음

3 Stall Reason 분류 ★ L·S·W·D·MIO·T·B

출처 Nsight Compute "Warp State Statistics" 섹션 · 상세 실측 해석 ↗ V18 §6

4 대표 원인 → 대처 매핑

• Long Scoreboard ↑ → coalesce · tiling · cp.async/TMA
• Short Scoreboard ↑ → smem bank · TC fragment 재사용
• Barrier ↑ → block size ↓ / pipeline 깊게
• Not Selected ↑ → 이미 occupancy 充分 (딴 bottleneck 탐색)

5 Latency Hiding 공식

출처 Little's law 적용. NVIDIA GTC'16 "GPU Performance Analysis" 참고.

6 Issue Slot 수식

• Nsight Compute: sm__inst_issued section
• FMA FP32 peak: partitions × (FP32_core / 16 lane / cyc) = 동일 식

7 Hopper 개선

• async instruction (WGMMA/TMA): issue 후 바로 다음
• scoreboard 수 증가로 ILP 확장
• warp specialization 으로 stall domain 분리

1 5 단계 abstract pipeline F·D·OC·EX·WB

[F]  fetch         I-cache → IB
[D]  decode        opcode + operand spec
[OC] operand coll. RF bank read staging
[EX] execute       ALU / LSU / TC / SFU
[WB] write-back    RF write

• 각 단계는 여러 cycle 일 수 있음 (HW 세부 비공개)
• pipeline 깊이는 public parameter 아님. 여기는 논리 모델.

출처 Greg Smith "Fermi SM" GTC'10 · Volta/Turing 백서 종합.

2 실행 파이프 분류

latency 는 NVIDIA micro-benchmark (Jia et al. 2019 / 2021) 통계 권장값 · 세대별 변동 ↗ V02 §14

3 Hazard 종류

• RAW (read-after-write): 이전 write 결과를 read. scoreboard 로 stall.
• WAW: same reg 에 연속 write. HW in-order 로 해소.
• WAR: write 이전에 old read. in-order 로 해소.
• Structural: 같은 unit 동시 요청. dispatch 막힘.

4 Scoreboard HW 의존성 추적

• short (fixed-latency) vs long (variable-latency) 스코어보드 분리
• SASS control code 에 scoreboard alloc/wait 인코딩

SASS 관련 ↗ V04 §14·§15

5 Bypass / Forwarding

• EX → EX direct forwarding: register write 전에 다음 instruction 에 값 전달
• FMA chain: back-to-back FFMA 사이 bypass 로 dependent-step latency ↓
• HW 구현 세부는 비공개. 공식 모델은 4-cycle FMA 을 기본 가정.

6 In-order Issue · Dual Issue

• dual-issue: partition 내 다른 pipe 에 동시 dispatch 조건 만족 시
• compiler (ptxas) 가 pair 가 될 수 있게 스케줄

7 Async pipeline (Hopper)

• WGMMA / TMA / cp.async 는 issue 후 즉시 retire queue 로
• 완료는 mbarrier 또는 wgmma.wait_group
• pipeline 상에서 별도 async proxy 경유 → generic proxy 와 fence 필요

Async proxy 상세 ↗ V04 §10

1 기본 모델 relaxed + scope

• sm_70+ 부터 scoped atomic / fence 명시
• CUDA Memory Model은 PTX ISA 에 공식화됨 (§8)

출처 PTX ISA 8.x §8 "Memory Consistency Model" (2021~)

2 Scope 계층 ★ cta·cluster·gpu·sys

scope 은 작은 → 큰. .sys 는 가장 비싸며 UM/peer 동기화에서 필요.

3 Ordering qualifier

PTX ISA 8.x §8.4 "Memory Operation Ordering"

4 Acquire·Release 패턴

// producer
st.global         data, %v;
st.release.gpu    flag, 1;

// consumer
ld.acquire.gpu    %f, flag;
@p bra spin_if_zero;
ld.global         %r, data; // see v

5 Fence 명령

6 Proxy 구분 (Hopper)

• TMA load 후 일반 ld 로 읽기 전에 fence
• WGMMA 결과를 store 하기 전에 fence

PTX ISA 8.x §9.7.13.3 "async proxy fence"

7 Atomic 종류

• atom.global.add.s32 · ·.min·.max·.and·.or·.xor·.exch·.cas
• scope: .cta/.gpu/.sys
• float atomic: atom.add.f32 · sm_60+ f64

상세 PTX atomic ↗ V03 §11

1 연산 latency (cyc) 4·20·200·400

출처 Jia et al. "Dissecting GPU Architectures" (2019, 2021) · 권장값. HW 세부는 NVIDIA 미공개.

2 메모리 access latency (cyc)

출처 Jia 2019/2021 · NVIDIA GTC'22 micro-benchmark 슬라이드. clock cycle 단위.

3 HBM latency 의 의미

• HBM 물리 access ~100 ns + NoC + L2 + MC overhead
• cycle 수 ≠ 절대 시간 (clock 세대마다 다름)

4 Warp 단위 hiding 예

Little's law · 공식 논문 Volkov 2010

5 TC async 명령 (Hopper)

출처 CUTLASS 3.x Hopper kernel timing · NVIDIA GTC'22

6 Rule of Thumb

• smem ≈ 20 cyc
• L2 ≈ 200~260 cyc
• HBM ≈ 400~500 cyc
• FFMA ≈ 4 cyc
• HMMA ≈ 16 cyc

7 주의

1 계층별 이론 BW ★

출처 A100/H100/GB200 WP + GTC'22, GTC'24 공식 슬라이드. SM 내부 BW 는 per-SM × SM 수 합산.

2 Peak vs Sustained

• 완전 coalesced, 대용량 stream 에서 HBM peak 의 80~90% 달성
• random access, stride ≠ 1 면 절반 이하로 추락

3 Intra-SM 대역 계산

4 Roofline 입력 FLOP/Byte

ridge = 모든 workload 가 compute-bound 이기 위한 최소 AI · ↗ V18 §2

5 PCIe Gen4 vs Gen5

64 GB/s per direction = 128 GB/s bidirectional aggregate (표기 관행 차이)

6 NVLink 세대 요약

NVSwitch 와 NVL 위상 ↗ V15 §3

7 실효 ↓ 주 원인

• stride access → sector 낭비
• partial warp load → merge 없음
• row thrashing → HBM row-open/close
• ECC off 가 가능하더라도 권장 X

1 TDP 세대

출처 각 세대 datasheet · GB200 NVL72 WP §3.1

2 DVFS 동적 주파수·전압

• boost clock 은 ideal case 상한
• 지속 100% load 에서 base clock 근처로 수렴
• nvidia-smi -lgc 로 clock 고정 가능 (DC)

3 Thermal Throttling

• T_slowdown (A100/H100 ≈ 88~90°C)
• liquid-cooled GB200 은 50~55°C 수준 유지 목표

출처 NVIDIA DCGM user guide · datasheet thermal spec

4 Power Instrument

• nvidia-smi --query-gpu=power.draw,power.limit
• NVML API: nvmlDeviceGetPowerUsage
• sampling rate 수 Hz (ms 단위 이벤트 불가)

5 Power efficiency 지표

6 Clock Gating

• idle pipe 의 clock 자동 중단 → 누수전력 ↓
• 특정 warp 가 SFU 만 쓰면 FP64 pipe 비활성
• SW 수준에서 조작 불가, HW 자동

7 프로파일링 상 의미

• power 포화 = DVFS throttling 가능성 → 커널별 clock 차이 발생
• 짧은 burst kernel 은 boost 유지, long-running 은 base 근처
• 커널 단독 profile 과 실제 cluster run 에서 clock 이 다를 수 있음

실측 해석/Nsight ↗ V18 §7

1 GA100 vs GA102 ★

출처 GA100 WP · GA102 WP (2020)

2 H100 vs L40 (Ada DC)

출처 H100/L40S 공식 datasheet

3 왜 DC 는 FP64·NVLink 특화 ★

• HPC (CFD/기후/원자) 는 FP64 요구
• 대형 모델/분산 학습 은 NVLink·NVSwitch BW 가 병목
• RTX 카드에서 FP64 는 시장 구분선

4 주요 lineage

arch 명은 same·chip(die)만 다름. PMPP 4e App.

5 DC / RTX 공통

• PTX ISA / CUDA API 거의 동일
• compute capability 로 기능 분기
• RTX 도 TC 탑재 (같은 세대)

6 실무 주의

• 개발 PC 에서 RTX 로 prototype → DC 로 배포
• sm_86 (Ampere RTX) vs sm_80 (A100) PTX 호환되나 최적 수치 다름

1 MIG 정의 slice = GPC + L2 + HBM

• A100: 최대 7 slice · H100: 최대 7 slice
• H100 PCIe / H200 / A100 에서 지원. RTX 불가.

출처 NVIDIA MIG User Guide (2023) · A100/H100 WP "MIG" 섹션

2 GPU Instance (GI) 프로파일 (A100 80GB)

출처 NVIDIA A100 MIG User Guide Table 2

3 H100 MIG 프로파일

출처 NVIDIA H100 MIG User Guide · 2024

4 Compute Instance (CI)

• GI = memory 격리 · CI = SM 그룹 분할
• 같은 GI 내 CI 는 L2/HBM 공유

5 격리 수준 ★

• SM 자체 분리: 다른 MIG 가 stall 시켜도 영향 X
• L2 slice 분리: cache pollution 차단
• HBM 주소 분리: memory BW 분할
• QoS: bandwidth 보장
• 단, NVLink/NVDEC 등 일부 자원은 공유 가능

출처 A100 MIG User Guide §2.3

6 사용 시나리오

7 한계

• 재구성은 kernel 실행 중 불가 (nvidia-smi MIG idle 상태에서만)
• slice 내 peak FLOPS = 전체 × (SM fraction) 근사

1 GPU arch 계보 ★ K·M·P·V·T·A·Ada·H·B

출처 NVIDIA CUDA C PG App. H 세대 표 · 각 WP 출시 연도

2 세대별 "한 줄" 특징

• Volta: 1st Tensor Core, NVLink 2
• Turing: INT TC + RT Core
• Ampere: 3rd TC, TF32, MIG, cp.async
• Ada: RTX 전용 (DC=L40), shader execution reorder
• Hopper: TMA·WGMMA·Cluster·FP8·TE
• Blackwell: FP4/FP6·dual-die·NVLink 5

3 Memory 한눈에

A100 / H100 / B200 순 · 단위 cycle 또는 TB/s

4 Hopper 니모닉

5 compute capability 퀵 룩업

TMA/WGMMA는 sm_90a 필수 (아키별 특수). 상세 ↗ V04 §1

6 3세대 종합 ★★

7 실수 방지 10선

1. sm_90 만으로는 TMA/WGMMA 불가 → sm_90a
2. reg 256 KB/SM 이지만 1 thread 상한 255
3. smem 동기화는 __syncthreads, TMA 는 mbarrier
4. cluster ≤ 8 block portable / 16 opt-in
5. DSM 은 같은 cluster 내 에서만
6. L2 persistence window 는 stream 수명 따라감
7. MIG 는 NVLink 일부 공유 · cluster 제약
8. ECC on 기준 peak BW 가 공식값
9. FP16 accum FP16 ≠ FP32 (정확도 차이)
10. Cycle 추정은 micro-benchmark, 절대 시간 아님

8 학습 경로

• 다음: PTX ISA ↗ V03
• Hopper 심화: ↗ V04
• SM 수치 → 커널 설계: ↗ V01 §4·V06
• 측정·검증: ↗ V18