GPU 내부 구조 Deep Dive: SIMT부터 HBM까지

처음에는 GPU의 CUDA Core 수를 보고 작은 CPU core가 수천 개 들어 있다고 이해했다. 그런데 그 설명으로는 왜 thread가 warp로 묶이는지, branch에서 lane이 쉬는지, occupancy가 높아도 느릴 수 있는지를 설명할 수 없었다.

그래서 이번에는 제품 사양표의 코어 수에서 시작하지 않았다. CPU가 command를 제출한 뒤 block이 SM/CU에 배치되고, warp/wave 명령이 register와 execution pipeline을 거쳐 L2와 GDDR/HBM까지 이동하는 경로를 순서대로 따라갔다. 그래픽의 rasterizer·ROP, AI의 Tensor/Matrix Core, RT·DMA·video engine도 같은 전체 구조 안에 놓았다.

확인: NVIDIA CUDA Programming Guide v13.3와 Best Practices Guide, Blackwell·Hopper·Turing 공식 설명, AMD CDNA 4·CDNA 3 whitepaper와 HIP 7.15.0 문서를 대조했다. 세대마다 달라지는 lane 수, issue width, cache 크기는 보편 법칙처럼 고정하지 않았다. 출처는 글 맨 아래에 모았다.

결론

GPU는 독립적인 소형 CPU 수천 개를 모아 놓은 장치가 아니다. GPU의 본질은 다음 네 가지다.

1. 같은 명령을 많은 데이터 lane에 적용하는 SIMT/SIMD 실행 구조
2. 수많은 warp/wave의 실행 상태를 on-chip에 유지하는 대규모 하드웨어 멀티스레딩
3. 한 작업의 지연을 줄이기보다 다른 작업을 실행해 지연을 숨기는 처리량 중심 스케줄링
4. 높은 연산량을 계속 공급하기 위한 register, shared memory, cache, NoC, GDDR/HBM 메모리 계층

여기에 그래픽 GPU는 rasterizer, texture unit, ROP 같은 고정 기능 장치를 추가하고, AI·HPC GPU는 Tensor/Matrix Core, HBM, 고속 GPU interconnect를 강화한다. 따라서 GPU 성능은 단순한 코어 수가 아니라 실행 lane의 실제 사용률, warp divergence, occupancy, 데이터 재사용률, 메모리 접근 패턴과 대역폭이 함께 결정한다.

1. 먼저 바로잡아야 할 표현

1.1 CUDA Core는 CPU Core가 아니다

CPU core 하나는 보통 다음 기능을 포함하는 독립적인 범용 프로세서다.

• 독립 instruction fetch/decode
• 정교한 branch predictor
• out-of-order scheduler와 reorder buffer
• speculative execution
• load/store queue
• 여러 정수·부동소수점 실행 포트
• private cache와 prefetcher
• privilege와 interrupt 처리

반면 NVIDIA의 CUDA Core나 AMD의 Stream Processor라는 수치는 대체로 SM/CU 내부의 FP32 또는 vector arithmetic lane 수에 가깝다. 각 lane이 독립적으로 명령어를 fetch하고 완전히 다른 프로그램을 실행하는 것은 아니다.

따라서 비교 단위는 다음과 같이 잡는 편이 정확하다.

CPU 전체 core                ↔ GPU의 SM 또는 CU
CPU vector ALU의 SIMD lane   ↔ CUDA Core 또는 Stream Processor에 가까운 실행 lane
CPU hardware thread          ↔ GPU thread와 정확한 1:1 대응 관계 없음

GPU thread가 특정 CUDA Core 하나를 계속 소유하는 것도 아니다. Warp 명령이 준비된 execution pipeline에 issue되고, 여러 lane이 그 순간의 thread operand를 처리한 뒤 다음 warp가 같은 pipeline을 사용한다.

1.2 GPU가 메모리 지연 자체를 없애는 것은 아니다

VRAM 접근은 여전히 길다. GPU는 수많은 warp/wave를 동시에 resident 상태로 유지하고, 하나가 memory 또는 dependency를 기다릴 때 다른 warp를 실행한다. 이것이 latency hiding이다. 지연을 제거하는 것이 아니라 유용한 다른 작업으로 가리는 것이다. [S1][S2]

1.3 GPU는 전부 programmable한 것도, 전부 graphics 전용인 것도 아니다

• Shader와 compute kernel은 SM/CU에서 프로그램으로 실행된다.
• Rasterization, texture filtering, depth/stencil, blending, video codec 등은 고정 기능 하드웨어가 맡을 수 있다.
• 데이터센터 GPU는 display와 raster 기능을 제거하거나 크게 축소할 수 있다.
• Tensor/Matrix Core와 RT Core도 별도의 완전한 프로세서라기보다 특정 연산을 빠르게 처리하는 전용 pipeline이다.

2. GPU 패키지와 다이의 전체 구조

GPU를 가장 바깥쪽부터 보면 다음과 같다.

CPU / OS / Runtime / Driver
              │
              ▼
PCIe / NVLink / Infinity Fabric / SoC Interconnect
              │
              ▼
Command Front-end ── Graphics·Compute·Copy Queues
              │
              ▼
Global Scheduler / Work Distributor
              │
              ▼
┌──────────────────────────── GPU Package ────────────────────────────┐
│  Processing Array                                                  │
│  ┌────────┐ ┌────────┐ ┌────────┐       ┌──────────────────────┐   │
│  │ SM/CU  │ │ SM/CU  │ │ SM/CU  │  ...  │ 수십~수백 SM/CU/WGP │   │
│  └───┬────┘ └───┬────┘ └───┬────┘       └──────────┬───────────┘   │
│      └───────────┴──────────┴──────── NoC / Crossbar ──────────────┤
│                                  │                                  │
│                 Distributed L2 / Last-Level Cache                   │
│                                  │                                  │
│                    Memory Controllers + PHY                         │
│                                  │                                  │
│                             GDDR / HBM                              │
│                                                                     │
│  Rasterizer·Texture·ROP    Tensor·Matrix·RT    Copy·DMA             │
│  Video·Display             GPU MMU·TLB         Firmware·Power       │
└─────────────────────────────────────────────────────────────────────┘

논리적으로는 다음 영역으로 나뉜다.

1. CPU와 연결되는 host interface
2. 명령을 받는 command front-end
3. block/workgroup을 분배하는 scheduler
4. 실제 shader/kernel을 실행하는 SM/CU array
5. cache, NoC와 외부 메모리 subsystem
6. rasterizer, ROP 같은 graphics fixed-function unit
7. Tensor/Matrix, RT 같은 accelerator
8. DMA, video, display, 보안, 전력 관리용 보조 엔진

NVIDIA의 대형 GPU는 GPC, TPC, SM 같은 반복 계층으로 processing array를 구성한다. AMD는 architecture에 따라 Shader Engine, Shader Array, CU 또는 두 CU를 묶는 WGP를 사용한다. 데이터센터용 AMD CDNA 3는 compute용 XCD chiplet과 memory/I/O용 IOD를 Infinity Fabric으로 연결한다. 명칭과 정확한 계층은 세대별로 달라도, 작은 실행 클러스터를 반복 배치하고 분산 cache와 on-chip network를 통해 memory controller에 연결한다는 원리는 같다. [S3][S4][S5][S6]

3. CPU가 GPU에 작업을 제출하는 경로

3.1 API 호출은 GPU 명령어 한 개가 아니다

예를 들어 다음 kernel launch가 있다고 하자.

vectorAdd<<<4096, 256>>>(a, b, c);

일반적인 제출 흐름은 다음과 같다.

Application
  → CUDA / HIP / Vulkan / DirectX runtime
  → User-mode driver
  → GPU용 command buffer 작성
  → Kernel/Shader code와 resource descriptor 준비
  → Queue에 command buffer 주소 등록
  → MMIO doorbell 기록
  → GPU command processor가 명령 fetch

CPU가 GPU에 전달하는 정보에는 다음이 포함될 수 있다.

• 실행할 kernel/shader 주소
• grid, block/workgroup 크기
• vertex/index buffer 주소
• descriptor, texture, sampler 정보
• render target과 depth buffer
• pipeline state
• memory copy 명령
• semaphore, fence와 barrier
• preemption 및 context 관련 상태

3.2 Command processor와 hardware queue

Command processor는 제출된 command stream을 해석하고 적절한 engine에 전달한다.

• Graphics queue: draw/dispatch와 graphics state
• Compute queue: 독립 compute dispatch
• Copy queue: DMA 전송
• Video queue: encode/decode 작업

조건이 맞으면 graphics, compute, copy가 겹쳐 실행될 수 있다. 이를 asynchronous compute나 copy/compute overlap으로 활용한다. 다만 queue가 분리되어도 L2, NoC, VRAM bandwidth, SM 같은 물리 자원을 공유하면 서로 성능에 영향을 준다.

3.3 Context, firmware와 preemption

GPU는 여러 프로세스와 API context를 구분해야 한다.

• process별 virtual address space
• queue와 execution context
• register 및 graphics state
• 접근 권한과 fault isolation
• scheduling priority
• watchdog와 hang recovery

GPU 내부 firmware/microcontroller가 queue scheduling, power management, security validation과 일부 context 관리에 관여할 수 있다. Preemption은 architecture에 따라 command, draw, thread group 또는 instruction에 가까운 서로 다른 granularity를 가질 수 있으며, fine-grained preemption일수록 저장해야 할 실행 상태가 많다.

4. Programming model이 hardware로 매핑되는 방식

4.1 Grid, block, thread

CUDA 기준 계층은 다음과 같다.

Kernel launch
  └─ Grid
      ├─ Thread Block 0
      │   ├─ Thread 0
      │   ├─ Thread 1
      │   └─ ...
      ├─ Thread Block 1
      └─ ...

AMD/OpenCL 계열 용어는 대략 다음과 대응한다.

일반적인 block/workgroup은 하나의 SM/CU에 배치되어 실행을 마칠 때까지 그곳의 register와 shared memory를 사용한다. 최신 architecture에는 SM cluster나 distributed shared memory 같은 명시적 확장이 있지만, 기본 모델은 block-local cooperation이다.

4.2 Warp와 wavefront

NVIDIA는 thread 32개를 warp 하나로 구성한다. AMD Radeon RDNA는 기본적으로 Wave32, AMD Instinct CDNA 계열은 전통적으로 Wave64를 사용하며 일부 RDNA ISA는 Wave64 실행도 지원한다. [S1][S6][S7][S8]

NVIDIA:
Warp 0 = thread 0  ~ 31
Warp 1 = thread 32 ~ 63

AMD RDNA 기본:
Wave 0 = work-item 0  ~ 31

AMD CDNA 전통적 구성:
Wave 0 = work-item 0  ~ 63

Thread block 크기가 warp 크기의 배수가 아니면 마지막 warp 일부 lane은 처음부터 inactive하다.

4.3 논리 thread와 물리 lane

각 thread는 프로그래밍 모델에서 다음 독립 상태를 가진다.

• thread ID
• register 값
• local address space
• 논리적 program counter/control state
• predicate와 call state

하지만 hardware는 같은 warp/wave의 active thread에 공통 명령을 issue한다. 즉, 개별 thread 모델을 제공하면서 실행 backend는 vector/SIMD 형태로 묶는 것이 SIMT다.

5. SM/CU 하나를 확대해서 보기

Instruction Cache
       │
       ▼
Warp/Wave Fetch·Decode
       │
       ▼
Warp/Wave Schedulers ── Scoreboard(의존성 추적)
       │
       ▼
Operand Collectors / Crossbar ◀──▶ Banked Register Files
       │
       ├── FP32·FP64 / FMA ────────────────┐
       ├── Integer / Branch / Scalar ──────┤
       ├── Special Function Unit ──────────┤──▶ Register writeback
       ├── Tensor·Matrix MMA ──────────────┤
       ├── Texture Unit ───────────────────┤
       └── Load·Store / Address Generation ┘
                    │
                    ├── Shared Memory / LDS
                    └── Coalescer → L1 → NoC → L2 → VRAM

5.1 Instruction cache와 fetch/decode

여러 warp가 같은 kernel의 서로 다른 데이터 구간을 실행하므로 instruction locality가 높다. SM/CU의 front-end는 instruction cache에서 코드를 가져와 decode하고, warp/wave별 다음 명령을 scheduler에 제공한다.

Architecture에 따라 instruction cache나 scalar front-end 일부를 인접 CU끼리 공유할 수 있다. 예를 들어 CDNA 3 문서는 두 CU가 instruction cache를 공유하는 구성을 설명한다. [S5]

5.2 Warp/wave slot

SM/CU는 여러 resident warp의 상태를 보관한다.

• 현재 control-flow 위치
• active lane mask
• register allocation 정보
• pending memory operation
• barrier 도착 상태
• exception/fault 관련 상태
• instruction dependency 상태

이 상태가 on-chip에 이미 있으므로 scheduler가 Warp A에서 Warp B로 issue 대상을 바꿀 때 CPU OS context switch처럼 전체 register를 memory에 저장했다가 복원하지 않는다.

5.3 Scheduler와 dispatch unit

Scheduler는 매 issue 기회에 ready warp를 고른다. Ready가 되려면 대략 다음 조건을 만족해야 한다.

ready(warp) =
  실행이 끝나지 않았고
  barrier에서 대기하지 않고
  source operand가 준비됐고
  필요한 execution pipeline이 사용 가능하고
  구조적·메모리 의존성이 해결됨

Architecture는 SM/CU를 여러 subpartition으로 나누고 각 partition에 scheduler, register slice와 execution pipeline을 배치할 수 있다. 정확한 scheduler 수, issue width, dual-issue 조건은 세대별로 달라 일반화하면 안 된다.

5.4 Scoreboard

Scoreboard는 명령어와 register 사이의 RAW/WAR/WAW 및 memory dependency를 추적한다.

1. LOAD R1, [address]
2. ADD  R2, R1, R3

LOAD 결과가 아직 돌아오지 않았다면 R1은 not-ready다. 해당 warp의 ADD를 실행하지 않고 다른 ready warp를 선택한다.

NVIDIA SM은 공식 CUDA 모델에서 CPU처럼 큰 reorder window와 branch speculation으로 한 thread의 명령을 공격적으로 재배치하기보다, warp 단위 in-order issue와 hardware multithreading을 사용한다. [S1]

5.5 Register file

각 thread는 논리적으로 자기 register를 갖지만, 물리적으로는 SM/CU 내부의 거대한 banked register file을 공유한다.

Thread 0  → R0, R1, R2, ...
Thread 1  → R0, R1, R2, ...
...
Thread 31 → R0, R1, R2, ...

Register는 다음 역할을 맡는다.

• arithmetic operand와 결과 저장
• loop/index 변수
• pointer와 predicate
• matrix fragment와 accumulator
• resident warp context 유지

Register는 매우 빠르지만 용량과 port가 제한된다. Thread당 register 사용량이 늘면 동시에 resident 가능한 block 수가 감소한다. Register를 억지로 줄이면 spill이 발생해 오히려 local/global memory traffic이 증가할 수 있다.

AMD architecture는 일반적으로 uniform 값을 위한 scalar register/ALU 경로와 lane별 값을 위한 vector register/VALU 경로를 구분한다. Uniform address나 loop 상태를 scalar 경로에서 한 번 계산하면 vector lane을 절약할 수 있다. 정확한 구성은 architecture family에 따라 다르다. [S5][S6]

5.6 Operand collector와 register bank

실행 명령은 source operand를 register file에서 가져와야 한다.

FMA R0, R1, R2, R3
→ R1, R2, R3 read
→ R1 × R2 + R3
→ R0 writeback

Operand collector와 crossbar가 register bank에서 데이터를 모아 execution unit으로 전달한다. 같은 bank 또는 port로 요청이 몰리면 추가 cycle이 필요할 수 있다. 따라서 arithmetic pipeline 수만 늘려도 register file이 데이터를 충분히 공급하지 못하면 성능이 오르지 않는다.

5.7 Execution pipeline

SM/CU에는 서로 다른 instruction class를 처리하는 pipeline이 있다.

• FP32/FP64/FP16 arithmetic와 FMA
• integer arithmetic, shift, bit operation
• branch와 predicate
• special function: reciprocal, square root, transcendental 근사
• load/store와 address generation
• texture sampling
• Tensor/Matrix multiply-accumulate
• shuffle, ballot, vote 같은 cross-lane operation

한 pipeline의 latency와 throughput은 다르다. 예를 들어 어떤 명령이 결과를 내기까지 여러 cycle이 걸려도 pipeline이 매 cycle 새 명령을 받을 수 있다면 latency는 길고 throughput은 높을 수 있다.

6. 명령어 하나가 실행되는 세부 경로

일반적인 arithmetic 명령은 다음 단계를 거친다.

Instruction cache fetch
  → decode
  → warp scheduler가 eligible warp 선택
  → scoreboard dependency 확인
  → dispatch port 선택
  → operand collector가 register read
  → execution pipeline
  → result writeback
  → scoreboard가 destination ready 처리

Load 명령은 더 긴 경로를 갖는다.

LD instruction issue
  → lane별 effective address 계산
  → warp 주소 coalescing
  → virtual address translation / TLB
  → L1 tag lookup
  → miss request queue / MSHR
  → on-chip network
  → L2 slice lookup
  → memory partition 선택
  → memory controller가 DRAM command scheduling
  → GDDR/HBM channel·bank·row access
  → L2/L1 fill
  → register writeback
  → scoreboard ready

이 긴 경로를 기다리는 동안 다른 warp가 execution pipeline을 사용한다.

7. SIMT, divergence와 reconvergence

7.1 SIMT와 SIMD의 차이

SIMD는 vector width가 instruction과 programming model에 직접 드러나는 경우가 많다.

[A0 A1 A2 A3] + [B0 B1 B2 B3]

SIMT는 프로그래머에게 scalar thread를 제공한다.

int i = blockIdx.x * blockDim.x + threadIdx.x;
c[i] = a[i] + b[i];

그러나 hardware는 warp 전체에 공통 명령을 issue한다.

LOAD a[i]  × active lanes
LOAD b[i]  × active lanes
ADD        × active lanes
STORE      × active lanes

즉, 프로그래밍 인터페이스는 thread, 실행 backend는 vector에 가까운 구조다. [S1]

7.2 Branch divergence

다음 코드를 생각해 보자.

if ((threadIdx.x & 1) == 0) {
    pathA();
} else {
    pathB();
}

하나의 warp 안에서 짝수 lane과 홀수 lane의 경로가 갈린다.

Path A 실행: active mask = 10101010...
Path B 실행: active mask = 01010101...
Reconverge:  active mask = 11111111...

일반적으로 서로 다른 경로는 mask를 바꿔 순차적으로 실행되므로 비활성 lane의 연산 capacity가 낭비된다. 서로 다른 warp가 다른 분기를 타는 것은 문제가 아니다. 문제는 같은 warp 내부의 divergence다.

짧은 분기는 compiler가 branch 대신 predication으로 변환할 수 있다. 이 경우 양쪽 명령을 실행하되 predicate가 false인 lane은 결과를 기록하지 않는다.

7.3 Independent Thread Scheduling의 의미

NVIDIA Volta 이후에는 thread별 execution state와 sub-warp 단위 reconvergence가 강화되었다. 그러나 이것이 warp의 lane들이 서로 다른 instruction을 같은 순간에 모두 실행하는 32개 독립 CPU가 되었다는 뜻은 아니다. SIMT execution과 active grouping은 여전히 유지된다. 이전 세대의 암묵적인 warp 동기화를 가정한 코드는 명시적인 warp synchronization이 필요하다. [S1]

7.4 Divergence가 생기는 다른 경우

• loop 반복 횟수가 lane마다 다름
• ray마다 다른 object/material을 만남
• sparse graph의 vertex degree가 다름
• 일부 thread만 조기 return
• block 크기가 warp size의 배수가 아님
• exception 또는 page fault 처리 경로가 달라짐

8. Hardware multithreading과 latency hiding

다음과 같은 상태를 생각할 수 있다.

Warp 0: global memory load 대기
Warp 1: FP32 FMA 실행 가능
Warp 2: shared memory 접근 가능
Warp 3: barrier 대기
Warp 4: integer address 계산 가능

Scheduler는 Warp 1, 2, 4 중 issue 가능한 명령을 고른다.

시간 →
Warp 0: LOAD ───────────────── 결과 도착 ─ ADD
Warp 1:      FMA FMA FMA
Warp 2:                  LDS LD
Warp 4:                         INT INT

GPU의 대규모 register file과 warp state storage는 단순 저장 공간이 아니라 latency를 숨기기 위한 실행 대기열이다.

하지만 resident warp가 많다는 사실만으로 충분하지 않다.

• 많은 warp가 모두 같은 memory request를 기다릴 수 있다.
• 특정 FP/INT/Tensor pipeline이 이미 포화될 수 있다.
• register dependency chain 때문에 eligible warp가 없을 수 있다.
• barrier에 대부분의 warp가 묶일 수 있다.
• grid가 작아 일부 SM이 빈 상태일 수 있다.

따라서 active warp 수와 실제로 issue 가능한 eligible warp 수를 구분해야 한다.

9. Resource allocation과 occupancy

9.1 Block이 배치될 때 예약되는 자원

SM/CU가 block/workgroup 하나를 받으면 보통 다음 자원을 할당한다.

• block 전체 thread의 register
• block의 shared memory/LDS
• warp/wave slot
• block slot
• barrier 및 synchronization state

하나라도 부족하면 해당 SM에 block을 더 배치할 수 없다.

9.2 Occupancy 계산 원리

Resident block 수는 개념적으로 다음 최솟값으로 제한된다.

residentBlocks = min(
  architectureBlockLimit,
  floor(maxThreadsPerSM / threadsPerBlock),
  floor(registersPerSM / allocatedRegistersPerBlock),
  floor(sharedMemoryPerSM / sharedMemoryPerBlock),
  warpSlotLimit,
  barrierAndOtherResourceLimits
)

warpsPerBlock = ceil(threadsPerBlock / warpSize)
residentWarps = residentBlocks × warpsPerBlock
occupancy = residentWarps / maxWarpsPerSM

실제 register/shared-memory allocation은 architecture별 granularity로 반올림되므로 단순 나눗셈 결과와 다를 수 있다. [S1][S2]

9.3 Occupancy 100%가 목표는 아니다

Occupancy가 너무 낮으면 latency hiding이 어렵지만, 높다고 항상 빠르지는 않다.

• 연산 pipeline이 이미 포화되면 warp를 더 올려도 이득이 없다.
• register를 줄여 occupancy를 높였지만 spill이 늘 수 있다.
• shared memory tile을 줄여 occupancy를 높였지만 데이터 재사용률이 떨어질 수 있다.
• 큰 block 하나보다 작은 block 여러 개가 tail effect를 줄일 수 있다.

따라서 목표는 occupancy 최대화가 아니라 충분한 eligible warp와 높은 pipeline utilization을 확보하는 것이다.

10. GPU 메모리 계층

가까움 / 작음 / 높은 bandwidth
────────────────────────────────
Register
Shared Memory / LDS
L1 Data / Texture / Constant·Scalar Cache
L2 / Infinity Cache
GDDR / HBM VRAM
PCIe 또는 coherent link 뒤의 System Memory
────────────────────────────────
멀어짐 / 큼 / 높은 접근 비용

10.1 Register

• thread-private 논리 상태
• on-chip banked storage
• arithmetic operand와 accumulator 저장
• 용량과 port 제한
• 사용량이 occupancy에 직접 영향

Register access도 무조건 비용이 없는 것은 아니다. Register dependency latency, bank/port 충돌과 execution pipeline latency가 존재한다.

10.2 Local memory

CUDA의 local memory는 이름과 달리 일반적으로 SM에 붙은 private SRAM이 아니다. Thread별 private address space일 뿐, 물리 접근은 global memory/cache 계층을 사용할 수 있다.

Local memory 사용 원인:

• register spill
• 큰 thread-local array
• compiler가 register로 배치하기 어려운 동적 index 배열
• call stack 또는 ABI 상태 일부

Local access
  → L1/L2 hit 시 cache 처리
  → miss 시 VRAM 접근

10.3 Shared memory / LDS

Shared memory는 programmer-managed on-chip scratchpad다.

__shared__ float tile[32][32];

주요 용도:

• block 내부 thread 간 데이터 공유
• global memory tile caching
• matrix multiplication
• reduction과 scan
• histogram
• stencil과 convolution
• producer-consumer pipeline

Shared memory는 cache와 달리 어떤 데이터를 언제 적재하고 교체할지 프로그램이 결정한다. 일부 architecture는 L1과 shared memory가 물리 SRAM capacity나 port를 공유하며 설정에 따라 비율을 조정한다. [S1][S2]

10.4 Shared-memory bank conflict

Shared memory/LDS는 여러 bank로 나뉜다. 개념적으로 32개 bank, bank width 4 byte인 경우:

address 0   → bank 0
address 4   → bank 1
...
address 124 → bank 31
address 128 → bank 0

Warp의 lane들이 서로 다른 bank를 접근하면 병렬 처리할 수 있다. 서로 다른 주소가 같은 bank에 몰리면 여러 cycle로 직렬화될 수 있다. 모든 lane이 같은 주소를 읽는 broadcast는 별도로 최적화될 수 있다. 실제 bank 수, width와 conflict 규칙은 architecture별로 확인해야 한다. [S6]

10.5 L1 data cache

• SM/CU에 가깝다.
• global/local load를 caching한다.
• architecture에 따라 shared memory와 capacity/port를 공유한다.
• write policy와 coherence 범위가 세대별로 다르다.
• CPU처럼 모든 private L1 사이의 강한 coherence를 자동 가정하면 안 된다.

10.6 Constant/scalar cache

Warp 전체가 같은 주소를 읽는 uniform access는 한 값을 여러 lane에 broadcast할 수 있다. Constant 또는 scalar cache는 이런 접근에 유리하다. 반대로 같은 cache에서 lane마다 서로 다른 주소를 읽으면 직렬화되거나 여러 transaction이 필요할 수 있다.

10.7 Texture cache와 texture unit

Texture unit은 단순 read cache가 아니다.

• texture coordinate addressing
• wrap, clamp와 border mode
• format conversion
• mip level 선택
• bilinear/trilinear filtering
• anisotropic sampling 일부
• 공간 지역성에 맞춘 cache

Shader가 texture sample 명령을 issue하면 texture unit이 여러 texel을 fetch하고 filtering한 결과를 반환한다.

10.8 L2 / Last-level cache

L2는 여러 SM/CU와 engine이 공유하는 중심 계층이다.

• SM/CU 요청 흡수
• VRAM traffic 감소
• GPU-wide 데이터 공유의 중심점
• DMA, graphics, compute engine 간 데이터 전달
• global atomic 지원
• memory controller 앞에서 request coalescing

큰 GPU에서는 하나의 중앙 SRAM보다 여러 L2 slice를 memory partition 근처에 분산 배치하고, 주소를 hashing/interleaving하여 slice를 선택하는 방식이 일반적이다. 주소 분포가 치우치면 특정 partition에 traffic이 집중되는 partition camping이 발생할 수 있다.

10.9 GDDR와 HBM

GDDR

• GPU board 주변에 개별 DRAM package 배치
• 비교적 적은 pin에서 높은 signaling rate 사용
• 소비자·graphics GPU에 일반적
• memory controller 수와 channel width가 전체 bus width를 구성

HBM

• 여러 DRAM die를 stack
• interposer 또는 advanced package로 GPU와 연결
• 매우 넓은 interface
• 높은 총 bandwidth와 좋은 energy/bit 특성
• AI/HPC accelerator에서 일반적

Memory controller는 다음을 담당한다.

• read/write request queue
• channel, pseudo-channel, bank와 row 선택
• row-buffer locality를 고려한 scheduling
• read/write turnaround 관리
• refresh
• ECC와 오류 보고
• QoS와 traffic arbitration

10.10 Integrated GPU와 unified memory

Integrated GPU/APU는 CPU와 같은 DRAM을 물리적으로 사용할 수 있다. 이 경우 discrete GPU처럼 매번 PCIe로 buffer를 복사하지 않아도 되지만 다음 비용은 남는다.

• CPU/GPU cache coherence
• shared DRAM bandwidth 경쟁
• address translation
• synchronization
• memory layout와 locality

Unified virtual address, unified memory migration, 물리적으로 동일한 DRAM은 서로 다른 개념이다.

11. Memory coalescing

GPU memory access는 thread 하나가 아니라 warp/wave 전체 주소 패턴으로 봐야 한다.

11.1 연속 접근

float x = data[globalThreadId];

lane 0  → A + 0
lane 1  → A + 4
...
lane 31 → A + 124

32개 lane이 연속된 float 32개, 총 128 byte를 요청한다. NVIDIA compute capability 6.0 이상에 대한 공식 예시에서는 정렬된 이 접근을 네 개의 32-byte transaction으로 처리할 수 있다. [S2]

11.2 Strided 접근

float x = data[globalThreadId * 32];

lane 0 → A
lane 1 → A + 128
lane 2 → A + 256
...

주소가 넓게 흩어져 더 많은 transaction과 불필요한 cache-line fetch를 발생시킨다.

11.3 Misalignment와 partial warp

• 시작 주소가 transaction 경계에 맞지 않으면 추가 segment가 필요할 수 있다.
• 일부 lane만 접근해도 포함된 segment 전체가 전송될 수 있다.
• Cache가 인접 warp의 중복 transaction 비용을 일부 흡수할 수 있다.
• 정확한 segment 크기와 규칙은 architecture별로 다르다.

11.4 AoS와 SoA

struct Particle { float x, y, z, mass; };

특정 field만 처리할 때 Array of Structures는 불필요한 field까지 cache line에 포함할 수 있다. Structure of Arrays는 같은 field를 연속 배치하여 coalescing과 vector load에 유리할 수 있다.

AoS: [x y z m][x y z m][x y z m]...
SoA: [x x x x...][y y y y...][z z z z...]

항상 SoA가 정답은 아니다. 실제 access pattern, cache reuse와 후속 연산을 함께 봐야 한다.

12. TLB, GPU MMU와 가상 메모리

현대 GPU는 process별 virtual address와 protection을 지원한다.

GPU virtual address
  → L1 TLB
  → L2/shared TLB
  → page-table walker
  → physical VRAM / system memory / peer memory

주요 구성 요소:

• process address-space identifier
• TLB hierarchy
• page-table walker
• access permission 검사
• page fault와 replay
• CPU/GPU shared virtual addressing
• peer GPU memory mapping
• PCIe BAR mapping
• IOMMU와 system memory 접근

대용량 데이터에서 access pattern이 여러 page에 무작위로 흩어지면 cache miss뿐 아니라 TLB miss와 page-table walk가 병목이 될 수 있다.

Managed/Unified Memory가 oversubscription을 허용하는 환경에서는 page가 CPU memory와 GPU memory 사이를 migration할 수 있다. 편리하지만 잘못된 access pattern은 page thrashing과 interconnect traffic을 발생시킨다. [S1]

13. Synchronization, memory ordering과 atomic

13.1 Warp-level synchronization

Shuffle, ballot, vote는 warp lane 사이에서 register 값을 교환하거나 predicate를 집계한다. Shared memory를 거치지 않아 reduction과 scan에 유리하다. 단, warp size와 active mask를 정확히 처리해야 한다.

13.2 Block barrier

__syncthreads() 같은 block barrier는 block의 participating thread가 해당 지점에 도착할 때까지 기다리고, programming model이 정한 shared/global memory visibility를 제공한다.

조건 분기 안에서 일부 thread만 barrier에 진입하면 deadlock 또는 undefined behavior가 발생할 수 있다.

13.3 Grid-wide synchronization

서로 다른 block은 일반적으로 독립적으로 scheduling되며 실행 순서를 가정할 수 없다. 전체 grid 동기화는 보통 다음 방법을 사용한다.

• kernel 종료 후 다음 kernel launch
• cooperative launch/grid synchronization
• atomic counter와 memory fence를 사용한 제한적 protocol
• 여러 kernel을 CUDA Graph/API dependency로 연결

13.4 Memory ordering

Barrier, visibility와 ordering은 같은 개념이 아니다.

• Execution barrier: 참여 thread의 진행을 맞춤
• Memory fence: 특정 scope에서 read/write 순서를 강제
• Atomic: read-modify-write의 원자성과 선택한 memory order/scope 제공

GPU는 CPU보다 relaxed한 memory model과 cache coherence 범위를 가질 수 있다. Correctness는 API/ISA가 정의한 scope, fence와 atomic 규칙에 따라 작성해야 한다. CDNA 3 문서도 vector L1 cache가 relaxed coherency model을 사용하며 강한 ordering에는 명시적 synchronization이 필요하다고 설명한다. [S5]

13.5 Atomic contention

여러 lane이 같은 주소를 atomic update하면 operation이 직렬화될 수 있다.

atomicAdd(globalCounter, 1);

개선 패턴:

1. thread별 partial result
2. warp-level reduction
3. block-level shared-memory aggregation
4. 마지막에 global atomic 횟수 최소화

Shared-memory atomic, L2/global atomic의 처리 위치와 throughput은 architecture별로 다르다.

14. 주요 실행 유닛

14.1 FP32/FP64와 FMA

일반 arithmetic pipeline은 add, multiply, fused multiply-add, compare와 conversion을 처리한다.

D = A × B + C

FMA는 통계상 multiply와 add를 합쳐 2 FLOP으로 계산하는 경우가 많다.

Peak FLOPS ≈ active arithmetic lanes × operations/cycle × clock

이론 peak는 다음이 모두 맞아야 접근할 수 있다.

• 올바른 datatype과 instruction mix
• 충분한 독립 work
• dependency stall 최소화
• register bandwidth 충분
• memory bottleneck 없음
• thermal/power throttling 없음

Consumer GPU는 FP64 lane을 적게 배치할 수 있고, HPC GPU는 FP64 처리량을 크게 강화할 수 있다. 따라서 FP32 코어 수만으로 과학 계산 성능을 추정하면 안 된다.

14.2 Integer와 address-generation pipeline

GPU kernel에도 integer 연산이 많이 필요하다.

• global index 계산
• pointer arithmetic
• loop counter
• bit packing
• branch predicate
• hash와 compression

FP 연산과 INT 연산이 서로 다른 pipeline을 사용할 수 있지만 동시 issue 조건과 throughput은 generation마다 다르다.

14.3 SFU

Special Function Unit은 reciprocal, reciprocal-square-root, trigonometric/exponential 근사 등 복잡한 함수를 처리한다. 정확도 모드와 compiler flag에 따라 hardware approximation, library sequence 또는 여러 instruction 조합이 사용될 수 있다.

14.4 Load/store와 address-generation unit

Load/store unit은 lane별 effective address를 계산하고, coalescer와 cache/memory subsystem에 request를 제출한다. 연산 unit 수가 많아도 load/store port나 address-generation throughput이 부족하면 memory instruction issue가 병목이 된다.

14.5 Cross-lane unit

• shuffle
• permute
• ballot
• all/any vote
• reduction 일부

Warp lane 사이의 데이터 교환을 register 경로에서 지원한다. 지원 범위와 lane grouping은 architecture별로 다르다.

15. Tensor Core와 Matrix Core

Tensor/Matrix unit은 작은 matrix tile에 대한 multiply-accumulate를 높은 처리량으로 수행한다.

D = A × B + C

일반적으로 warp/wave가 matrix instruction을 issue하고 여러 lane이 operand fragment와 accumulator를 분담한다.

Global Memory
  → Shared Memory / LDS tile
  → Register fragments
  → Tensor/Matrix MMA pipeline
  → Register accumulator
  → Global Memory

높은 성능의 이유:

• matrix operand의 반복 사용
• 내부 data reuse로 register read와 data movement 감소
• 많은 multiply-add를 하나의 instruction으로 표현
• FP16/BF16/FP8/INT8 같은 좁은 datatype으로 더 많은 연산기 배치
• accumulator를 더 높은 precision으로 유지 가능

하지만 다음 조건에서 peak와 멀어진다.

• matrix shape가 tile에 맞지 않음
• alignment/layout가 맞지 않음
• 작은 matrix로 launch overhead가 큼
• shared memory/register 공급 부족
• datatype conversion 비용이 큼
• sparsity 조건을 만족하지 못함
• memory bandwidth가 먼저 포화됨

Tensor Core는 별도의 독립 AI 컴퓨터가 아니다. Warp scheduler, register file, shared memory, cache와 memory bandwidth를 일반 pipeline과 공유한다. [S3][S5]

16. Graphics pipeline

현대 GPU는 programmable shader와 fixed-function unit을 결합한다.

Draw Command
    │
    ▼
Vertex / Index Fetch
    │
    ▼
Vertex Shader
    │
    ▼
Tessellation / Geometry / Mesh Shader
    │
    ▼
Primitive Assembly / Clip / Cull
    │
    ▼
Rasterizer
    │
    ▼
Pixel / Fragment Shader
    │
    ▼
ROP: Depth·Stencil·Blend
    │
    ▼
Framebuffer

16.1 Vertex/index fetch

Vertex와 index buffer를 memory에서 읽고 format을 변환한다. Vertex cache는 같은 vertex가 여러 triangle에서 재사용될 때 shader 실행을 줄인다.

16.2 Vertex shader

SM/CU에서 실행된다.

• model/view/projection transform
• skinning
• normal transform
• vertex attribute 계산

16.3 Tessellation, geometry와 mesh processing

Architecture와 API 기능에 따라 다음 단계가 존재할 수 있다.

• tessellation control/hull shader
• fixed-function tessellator
• tessellation evaluation/domain shader
• geometry shader
• task/amplification shader
• mesh shader

Mesh shader는 전통적인 vertex/geometry front-end 일부를 programmable workgroup 모델로 재구성한다.

16.4 Primitive assembly, clipping과 culling

Vertex를 triangle/line/point로 조립하고 view volume 바깥 primitive를 제거하거나 잘라낸다. Back-face culling, viewport transform, primitive setup이 이어진다.

16.5 Rasterizer

Triangle이 덮는 pixel/sample을 결정한다.

• edge equation
• coverage mask
• sample 위치
• perspective-correct interpolation 준비
• tile/region 분배
• hierarchical depth rejection 연계

16.6 Pixel/fragment shader

생성된 fragment를 warp/wave로 묶어 SM/CU에서 실행한다.

• material과 lighting
• texture sampling
• normal mapping
• alpha/discard
• color와 depth 계산

인접 pixel을 묶어 실행하면 texture gradient 계산과 cache locality에 유리하지만, triangle 경계나 discard 때문에 helper/inactive lane이 생길 수 있다.

16.7 Early-Z와 hierarchical Z

Depth test를 pixel shader 전에 수행할 수 있으면 가려진 fragment의 shader 실행을 피한다. 다만 shader가 depth를 변경하거나 discard/side effect를 사용하면 early test가 제한될 수 있다.

16.8 ROP

ROP 또는 render backend는 최종 framebuffer 단계에 관여한다.

• depth test와 update
• stencil test와 update
• color blending
• multisample resolve 일부
• render-target compression
• framebuffer read/write

Pixel shader가 빨라도 ROP 처리량이나 framebuffer bandwidth가 부족하면 fill-rate bound가 된다. NVIDIA Turing의 구체적 예에서는 GPC별 raster engine, SM별 texture unit, memory partition과 연결된 ROP/L2 구성을 확인할 수 있다. [S4][S9]

16.9 Tile-based rendering

모바일 GPU는 화면을 tile로 나눈 뒤 tile 내부 color/depth를 on-chip memory에 보관하는 tile-based/deferred 방식을 많이 사용한다.

Geometry binning
  → 각 primitive를 tile 목록에 분류
  → tile별 rendering
  → on-chip tile buffer에서 depth/blend
  → 완성된 tile만 DRAM에 기록

목적은 제한된 전력과 DRAM bandwidth에서 framebuffer 왕복을 줄이는 것이다. 데스크톱 GPU도 tile/region 단위 최적화와 compression을 사용할 수 있지만 전체 pipeline 조직은 모바일 TBDR과 동일하다고 단정할 수 없다.

17. Ray-tracing unit

RT Core 같은 전용 unit은 ray tracing 전체를 대신하지 않는다. 주로 다음 작업을 가속한다.

• BVH node traversal
• ray–axis-aligned bounding box test
• ray–triangle intersection
• 후보 hit 보고

실행 흐름은 대략 다음과 같다.

Shader가 ray 생성
  → RT unit이 BVH traversal/intersection
  → 후보 또는 최종 hit 반환
  → Shader가 any-hit / closest-hit / miss 처리
  → material·texture·lighting 계산
  → 필요하면 secondary ray 생성

Ray마다 traversal 경로, hit object와 material이 달라질 수 있어 divergence와 불규칙 memory access가 크다. RT unit은 traversal과 intersection 비용을 줄이지만 shader divergence, texture fetch와 scene-memory bandwidth까지 없애지는 않는다. [S4]

18. DMA, video와 display engine

18.1 Copy/DMA engine

Copy engine은 SM에서 byte-copy kernel을 실행하지 않고 다음 전송을 수행할 수 있다.

• host memory ↔ VRAM
• VRAM ↔ VRAM
• GPU ↔ GPU peer memory
• buffer clear/fill 일부

Pinned host memory, multiple streams와 double buffering을 사용하면 copy와 compute를 겹칠 수 있다. 실제 overlap 여부는 engine 수, bus와 memory bandwidth 경쟁에 달려 있다.

18.2 Video engine

영상 codec은 정형화된 계산이 많아 전용 fixed-function unit이 전력 효율 면에서 유리하다.

• H.264/H.265/AV1 decode
• video encode
• entropy coding
• motion estimation 일부
• color conversion/post-processing 일부

18.3 Display engine

• framebuffer scanout
• display timing
• plane composition 일부
• HDR/output format
• HDMI/DisplayPort 연결

데이터센터 compute GPU는 display engine이 없을 수 있다.

19. On-chip network, memory partition과 chiplet

19.1 NoC와 crossbar

SM/CU, cache slice, memory controller, DMA와 graphics unit 사이에는 대규모 interconnect가 필요하다.

NoC는 다음 traffic을 운반한다.

• cache request/response
• texture fetch
• atomic
• cache-coherence message
• DMA
• page-table walk
• graphics render-target traffic

대역폭, hop latency, arbitration과 congestion이 성능에 영향을 준다. ALU가 남아 있어도 NoC가 포화되면 데이터를 공급할 수 없다.

19.2 Distributed cache와 address mapping

L2와 memory controller를 여러 partition으로 나누고 주소 bit/hash로 mapping하면 parallel bandwidth를 얻는다. 반면 특정 stride가 같은 partition에 집중되면 불균형이 생길 수 있다.

19.3 Chiplet GPU

대형 monolithic die는 수율과 reticle 크기에 제약을 받는다. Chiplet은 기능을 여러 die로 나눠 package 내부 interconnect로 연결한다.

가능한 분할:

• compute die
• cache/memory-controller die
• I/O die
• CPU chiplet과 GPU chiplet
• HBM stack

장점:

• 제조 수율 개선
• 공정 노드별 최적화
• 제품 규모 확장
• compute와 memory capacity 조합 유연성

비용:

• die 간 latency
• interconnect bandwidth와 전력
• cache coherence 복잡도
• 물리 배치와 열 밀도
• software가 관찰할 수 있는 NUMA 성격

CDNA 3는 여러 XCD와 IOD를 Infinity Fabric으로 연결하고 cache/coherence 계층을 구성하는 대표적인 예다. [S5]

2025년에 공개된 더 최신 사례도 같은 방향을 보여 준다. NVIDIA Blackwell Ultra는 두 reticle-sized die를 NV-HBI로 연결하면서 하나의 CUDA accelerator로 노출한다. AMD CDNA 4는 8개 XCD와 2개 IOD, 8개 HBM3E stack을 package 내부 Infinity Fabric으로 묶는다. Chiplet은 GPU programming model을 여러 장치로 쪼갠다는 뜻이 아니라, 하나의 논리 GPU 내부 구현을 여러 die로 분할할 수 있다는 뜻이다. [S10][S11]

20. 전력, clock, 열과 수율

GPU의 모든 unit이 항상 peak clock으로 동시에 작동할 수 있는 것은 아니다.

20.1 Clock/power domain

• graphics/compute core clock
• memory clock
• media/display clock
• interconnect clock
• independent power gating domain

유휴 block은 clock gating 또는 power gating으로 전력을 줄일 수 있다.

20.2 DVFS와 boost

Clock은 다음 조건에 따라 변한다.

• 전력 제한
• 온도
• 전압 안정성
• workload instruction mix
• VRM와 board limit
• 데이터센터의 reliability policy

따라서 명목 boost clock과 장시간 sustained clock은 다를 수 있다. Tensor/FP 연산 밀도가 높은 workload는 power limit에 먼저 도달할 수도 있다.

20.3 Dark silicon과 동시 사용 제약

다이에 모든 기능을 넣어도 전력·열 한계 때문에 모든 transistor를 최대 속도로 동시에 켜지 못할 수 있다. Specialized unit은 특정 작업에서 general-purpose logic보다 높은 performance/watt를 제공하기 위해 존재한다.

20.4 수율과 disabled unit

GPU는 반복 unit 구조이므로 결함이 있는 일부 SM/CU, cache slice 또는 memory partition을 fuse로 비활성화해 다른 SKU로 판매할 수 있다. 따라서 architecture의 full-chip 구성과 실제 제품의 active unit 수는 다를 수 있다. [S3][S4]

20.5 Reliability

데이터센터 GPU는 다음 기능을 강화할 수 있다.

• HBM/cache/register ECC
• parity
• error containment
• row remapping
• thermal/voltage monitoring
• secure boot와 firmware validation
• partitioning/MIG 같은 isolation

ECC도 저장 공간, bandwidth와 power overhead가 있으므로 모든 consumer 제품에서 같은 수준으로 제공되는 것은 아니다.

21. Vector addition 실행 추적

다음 kernel을 실행한다고 하자.

__global__ void add(const float* a, const float* b, float* c) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    c[i] = a[i] + b[i];
}

add<<<4096, 256>>>(a, b, c);

총 thread 수는 다음과 같다.

4096 blocks × 256 threads = 1,048,576 threads

NVIDIA warp 32개 thread를 기준으로 block당 warp는 8개다.

단계 1: 제출

CPU runtime/driver가 kernel 주소, argument와 grid 설정을 command queue에 기록한다.

단계 2: Block dispatch

Global scheduler가 실행 자원이 있는 SM에 block을 배치한다. Block 실행 순서는 소스의 block index 순서와 같다고 가정할 수 없다.

단계 3: Resource allocation

SM은 block에 필요한 register, shared memory, warp slot과 barrier state를 예약한다. 자원이 부족하면 추가 block은 대기한다.

단계 4: 주소 계산과 load issue

Warp scheduler가 ready warp를 골라 index 계산과 a[i], b[i] load를 issue한다.

단계 5: Coalescing과 cache lookup

32 lane이 연속 float를 읽으면 배열 하나당 128-byte 논리 요청이 된다. NVIDIA의 앞서 설명한 조건에서는 32-byte transaction 네 개로 구성될 수 있다.

단계 6: Latency hiding

L1/L2 miss가 발생한 warp는 scoreboard에서 대기한다. Scheduler는 다른 resident warp의 arithmetic/load 명령을 issue한다.

단계 7: Add

두 load 결과가 register에 도착하면 scoreboard가 operand를 ready로 표시한다. Scheduler가 FP add를 issue하고 active lane들이 각자 a[i] + b[i]를 계산한다.

단계 8: Store

결과 store도 warp 주소를 모아 coalescing한 후 memory subsystem에 전달한다. Store 완료의 관찰 시점은 cache policy와 synchronization 규칙에 따라 다르다.

단계 9: Block retire

8개 warp가 모두 종료되면 block 자원이 해제되고 waiting block이 들어온다.

이 kernel은 element당 load 8 byte, store 4 byte에 비해 add 한 번만 수행하므로 보통 arithmetic보다 memory bandwidth 영향을 크게 받는다.

22. Tiled matrix multiplication 실행 추적

행렬 곱의 기본식은 다음과 같다.

C[i][j] += A[i][k] × B[k][j]

Naive 구현은 A/B 원소를 VRAM에서 반복해서 읽는다. Optimized kernel은 tile을 사용한다.

Global Memory
  → block이 A/B tile을 cooperative load
Shared Memory / LDS
  → thread가 register fragment로 load
Register / Tensor·Matrix Core
  → accumulator 재사용
  → 다음 tile을 asynchronous prefetch
  → 최종 C만 Global Memory에 store

핵심 최적화:

1. Global memory 접근 coalescing
2. Shared memory에서 A/B tile 재사용
3. Register accumulator 유지
4. Shared-memory padding으로 bank conflict 감소
5. Double buffering으로 copy와 compute overlap
6. Tensor/Matrix instruction에 맞는 tile/layout 사용
7. Edge tile의 inactive lane과 분기 최소화

행렬 곱이 GPU에 잘 맞는 이유는 같은 입력 값을 여러 multiply-add에 재사용해 arithmetic intensity를 크게 만들 수 있기 때문이다.

23. 성능을 해석하는 모델

23.1 Roofline 관점

Arithmetic intensity = 수행 연산량 / 외부 memory에서 이동한 byte

대략적인 성능 상한:

attainablePerformance = min(
  peakCompute,
  memoryBandwidth × arithmeticIntensity
)

• Arithmetic intensity가 낮으면 memory-bound
• 충분히 높으면 compute-bound로 이동

실제 성능에는 cache, instruction issue, occupancy, divergence와 synchronization도 추가로 영향을 준다.

23.2 주요 병목 분류

Compute-bound

• FP/Tensor pipeline 포화
• 높은 data reuse
• memory bandwidth 여유

Memory-bandwidth-bound

• load/store가 많음
• arithmetic intensity가 낮음
• 이미 coalescing되어 있지만 전체 byte가 너무 큼

Memory-latency-bound

• pointer chasing
• random access
• resident/eligible warp 부족
• TLB miss와 dependency chain

Instruction/issue-bound

• address와 control instruction 비중이 큼
• 특정 dispatch port에 명령이 몰림
• dependency가 짧은 간격으로 반복됨

Divergence-bound

• active lane 비율이 낮음
• warp별 실행 시간이 크게 다름
• ray/graph/sparse workload

Synchronization-bound

• barrier에서 오래 대기
• global atomic contention
• block 간 load imbalance

Graphics fixed-function-bound

• raster/geometry rate
• texture filtering rate
• pixel fill/ROP rate
• framebuffer bandwidth

Transfer-bound

• CPU↔GPU PCIe 복사가 실제 compute보다 큼
• 작은 kernel을 반복 launch
• unified-memory page migration/thrashing

23.3 Tail effect

Grid 마지막에 남은 block 수가 SM 수보다 적으면 일부 SM이 먼저 유휴 상태가 된다. Block 실행 시간 편차가 크면 한 SM의 긴 block 때문에 kernel 전체 완료가 지연된다. Block 크기와 work distribution을 조정하거나 persistent-work queue를 사용할 수 있다.

23.4 Peak 수치 해석 주의사항

제품 표의 TFLOPS는 다음을 확인해야 한다.

• FP64, FP32, TF32, FP16, BF16, FP8, INT8 중 어떤 datatype인지
• vector core인지 matrix core인지
• FMA를 2 operation으로 계산했는지
• dense인지 structured sparsity 조건인지
• boost clock 기준인지 sustained clock 기준인지
• memory bandwidth가 실제 workload를 공급할 수 있는지

서로 다른 precision과 sparsity 조건의 peak 숫자를 직접 비교하면 잘못된 결론을 낼 수 있다.

24. CPU와 GPU 비교

GPU에 불리한 작업:

• 병렬 작업 수가 적음
• 짧은 kernel을 매우 자주 launch
• data transfer가 계산보다 큼
• 같은 warp 안의 제어 흐름이 심하게 갈림
• global synchronization이 빈번함
• linked-list/tree pointer chasing
• block/workgroup별 작업량 편차가 큼
• 작은 데이터에 복잡한 제어만 많음

25. Architecture별 차이를 읽는 방법

다음 값은 고정된 GPU 보편 상수가 아니다.

• SM/CU당 lane 수
• warp/wave size
• scheduler와 dispatch unit 수
• FP/INT 동시 issue 조건
• register file 크기와 allocation granularity
• shared memory/L1 capacity와 partition 방식
• shared-memory bank 수와 width
• L1/L2 cache line과 transaction size
• Tensor/Matrix tile과 지원 datatype
• RT unit 배치
• cache coherence와 memory scope
• memory-controller/ROP 결합 방식

공식 architecture 문서를 읽을 때는 다음 순서가 유용하다.

1. Full-chip block diagram
2. SM/CU block diagram
3. Warp/wave size와 issue model
4. Register/shared-memory resource limit
5. Cache hierarchy와 transaction granularity
6. Execution pipeline별 throughput/latency
7. Product별 active unit 수와 memory 구성
8. Compiler/ISA가 실제로 생성한 instruction

Architecture whitepaper의 full-chip 수치는 실제 판매 제품의 fuse 구성과 다를 수 있다.

26. 자주 하는 오해

오해 1: GPU core가 10,000개면 CPU core 10,000개와 같다

아니다. 대부분 독립 fetch/decode와 branch predictor를 가진 범용 core가 아니라 vector arithmetic lane이다.

오해 2: Thread 1개가 CUDA Core 1개에 고정된다

아니다. Thread는 warp의 lane으로 실행되고, 여러 warp가 같은 physical pipeline을 시간 분할해 사용한다.

오해 3: GPU는 context switching이 없으므로 thread가 모두 동시에 실행된다

Resident warp는 빠르게 교대하지만, launch된 모든 thread가 동시에 물리적으로 resident한 것은 아니다. Block이 wave 단위로 SM/CU에 들어오고 나간다.

오해 4: Branch가 있으면 GPU에서 실행할 수 없다

실행할 수 있다. 같은 warp 안에서 경로가 갈릴 때 lane 활용률과 실행 시간이 나빠질 수 있을 뿐이다.

오해 5: Shared memory는 자동 cache다

아니다. Programmer가 적재, 동기화, 재사용을 직접 관리하는 scratchpad다.

오해 6: Local memory는 register 다음으로 가까운 memory다

아니다. Thread-private 주소 공간이지만 register spill 등이 global/cache memory path를 사용할 수 있다.

오해 7: Occupancy가 높을수록 항상 빠르다

아니다. 충분히 latency를 숨긴 뒤에는 pipeline 포화, data reuse와 spill이 더 중요할 수 있다.

오해 8: Tensor Core가 있으면 모든 AI 코드가 자동으로 빨라진다

아니다. 지원 datatype, shape, layout, tile, memory supply와 library/kernel 구현이 맞아야 한다.

오해 9: Unified Memory면 CPU/GPU 데이터 이동 비용이 사라진다

아니다. 주소 공간만 통합되거나 page migration이 자동화된 것일 수 있다. 물리 이동과 coherence 비용은 남는다.

오해 10: GPU bandwidth가 높으므로 모든 memory 접근이 빠르다

아니다. 높은 bandwidth와 낮은 단일 접근 latency는 다르다. Random dependent load는 GPU에서도 어렵다.

27. 면접용 답변 구조

짧게 설명해야 한다면 다음 순서로 답하면 된다.

GPU는 처리량 중심의 대규모 병렬 프로세서입니다. CPU가 제출한 command를 front-end가 해석하고, scheduler가 thread block을 여러 SM/CU에 분배합니다. SM/CU는 thread를 warp/wave로 묶어 같은 명령을 여러 execution lane에 issue하는 SIMT 방식으로 실행합니다. 각 SM에는 warp scheduler, scoreboard, 큰 register file, shared memory/L1, FP·INT·load/store·Tensor 같은 pipeline이 있습니다. 한 warp가 memory를 기다리면 상태를 이미 on-chip에 둔 다른 warp를 실행해 latency를 숨깁니다. Memory는 register, shared/L1, L2, GDDR/HBM 계층이며 warp 전체 주소가 coalescing되는지가 중요합니다. 그래픽 GPU는 rasterizer, texture unit, ROP를 추가하고 최신 GPU는 RT와 matrix accelerator를 포함합니다. 따라서 실제 성능은 코어 수보다 divergence, occupancy, 데이터 재사용, cache hit, coalescing과 memory bandwidth에 의해 결정됩니다.

추가 질문에는 다음 키워드로 확장한다.

• 실행: SIMT, warp/wave, scoreboard, latency hiding
• 자원: register pressure, shared memory, occupancy
• memory: coalescing, bank conflict, L1/L2/HBM, TLB
• graphics: rasterizer, fragment shader, ROP
• AI: Tensor/Matrix Core, tiling, mixed precision
• 성능: arithmetic intensity, roofline, divergence, tail effect

출처

2026-07-21에 NVIDIA와 AMD의 공식 문서를 직접 대조했다. CUDA Programming Guide는 v13.3, AMD HIP 문서는 7.15.0 기준이다. 구체적인 unit 수, cache 크기와 issue 규칙은 제품과 세대마다 달라 대상 architecture 문서를 다시 확인해야 한다.

• [S1] NVIDIA, CUDA Programming Guide SIMT execution model, 32-thread warp, independent thread scheduling, hardware multithreading, thread/block/grid와 memory hierarchy. https://docs.nvidia.com/cuda/cuda-programming-guide/

• [S2] NVIDIA, CUDA C++ Best Practices Guide Global-memory coalescing, memory space 특성, occupancy, register/shared-memory 자원과 latency hiding. https://docs.nvidia.com/cuda/cuda-c-best-practices-guide/

• [S3] NVIDIA, Hopper Architecture In-Depth GH100 full-chip의 GPC/TPC/SM 계층, L2, HBM controller, Tensor Core와 제품별 active-unit 구성. https://developer.nvidia.com/blog/nvidia-hopper-architecture-in-depth/

• [S4] NVIDIA, Turing Architecture In-Depth Graphics용 GPC/raster engine/TPC/SM, texture unit, ROP/L2/memory partition, Tensor Core와 RT Core의 구체적 배치 예시. https://developer.nvidia.com/blog/nvidia-turing-architecture-in-depth/

• [S5] AMD, Introducing AMD CDNA 3 Architecture Whitepaper XCD/IOD chiplet, Infinity Fabric, ACE/CU/L2 계층, CU의 scheduler·shader·matrix·L1·LDS 구조와 cache coherence 특성. https://www.amd.com/content/dam/amd/en/documents/instinct-tech-docs/white-papers/amd-cdna-3-white-paper.pdf

• [S6] AMD ROCm, HIP Hardware Implementation AMD CU/WGP, SIMD pipeline, RDNA/CDNA 차이, LDS와 bank 구조, hardware scheduling. https://rocm.docs.amd.com/projects/HIP/en/develop/understand/hardware_implementation.html

• [S7] AMD ROCm, Introduction to the HIP Programming Model Thread/workgroup/grid, warp/wavefront와 RDNA/CDNA wave size, resource에 따른 resident wave 제한. https://rocm.docs.amd.com/projects/HIP/en/develop/understand/programming_model.html

• [S8] AMD GPUOpen, RDNA Architecture Presentation RDNA의 WGP, Wave32/Wave64, cache hierarchy와 graphics/compute shader 실행 특성. https://gpuopen.com/wp-content/uploads/2019/08/RDNA_Architecture_public.pdf

• [S9] NVIDIA GPU Gems, Graphics Pipeline Performance Vertex processing, raster/fragment 단계와 ROP의 depth·stencil·blend·framebuffer 역할. https://developer.nvidia.com/gpugems/gpugems/part-v-performance-and-practicalities/chapter-28-graphics-pipeline-performance

• [S10] NVIDIA, Inside NVIDIA Blackwell Ultra 두 reticle-sized die를 NV-HBI로 연결하면서 하나의 CUDA accelerator와 coherent shared L2로 구성하는 최신 multi-die 사례. https://developer.nvidia.com/blog/inside-nvidia-blackwell-ultra-the-chip-powering-the-ai-factory-era/

• [S11] AMD, Introducing AMD CDNA 4 Architecture Whitepaper 8개 XCD, 2개 IOD, HBM3E와 Infinity Fabric으로 구성된 heterogeneous chiplet GPU의 최신 사례. https://www.amd.com/content/dam/amd/en/documents/instinct-tech-docs/white-papers/amd-cdna-4-architecture-whitepaper.pdf