PMPP 기초 재정리

1 Heterogeneous Parallel Computing 정의

• 2003 이후 clock scaling 종료 (열·전력 벽) → multi-core / many-thread 전환
• Multi-core: sequential 성능 유지 · ILP · OoO · large cache
• Many-thread: throughput 극대화 · massive thread · relaxed memory
• von Neumann 모델 · thread = 순차 실행 단위 (독립 PC · register)

2 CPU vs GPU 설계 축

3 Amdahl's Law ★ 1/((1-p)+p/s)

4 Gustafson 보정 weak scaling

• Amdahl = strong scaling (고정 문제, rank ↑)
• Gustafson = weak scaling (rank 당 일감 고정, rank ↑ → 총 문제 ↑)
• ↗ V15 §16 strong/weak scaling · α-β 모델 상세

5 A100 vs CPU 정량 baseline ★

Source: NVIDIA A100 Datasheet v1.3 · 출처 고정값만 표기 · 실측 아님

6 병렬 프로그래밍 4대 난제 알메인동

1. Algorithmic complexity: 병렬화로 work 증가 가능 → work-efficient 확보
2. Memory-bound: 대부분 kernel이 BW 제약 → tiling 필수 ↗ §7
3. Input-sensitive: skewed 분포 → load imbalance
4. Synchronization: barrier · atomic · 결합성

7 Programming Interface 지형

1 Data Parallelism SPMD

• 예: vector add · 픽셀 변환 · GEMM · FFT
• GPU가 빛나는 축: 각 thread 독립 · 동일 코드

2 CUDA 프로그램 5단계 할복실복해

1. 할당 cudaMalloc(&d_A, n*sizeof(float))
2. 복사 H→D cudaMemcpy(d_A, h_A, n, H2D)
3. 실행 kernel<<<G, B>>>(d_A, ...)
4. 복사 D→H cudaMemcpy(h_A, d_A, n, D2H)
5. 해제 cudaFree(d_A)

3 함수 한정자 (Qualifiers)

4 Vector Add ★ canonical kernel

__global__ void vecAdd(const float* A,
                       const float* B,
                       float* C, int n){
  int i = blockIdx.x*blockDim.x + threadIdx.x;
  if (i < n) C[i] = A[i] + B[i];
}
// host launch
int T = 256;
int G = (n + T - 1) / T;   // ceil div
vecAdd<<<G, T>>>(d_A, d_B, d_C, n);

5 Kernel 실행 구성 <<<G, B, shmem, stream>>>

• grid = block 집합 (1D/2D/3D)
• block = thread 집합 (≤1024 thread, 3D)
• block 간 독립 실행 가정 → 동기화 불가 (transparent scalability)
• block 내부 → __syncthreads(), shared memory OK
• optional: dynamic shmem 크기 · CUDA stream 지정

grid[ block(0,0) block(1,0) ... ]
 └─ block[ t(0,0) t(1,0) ... t(31,0)  ← warp 0
            t(0,1) ...  t(31,1)       ← warp 1 ]

6 내장 변수 (Built-in)

7 Error Handling 패턴

#define CUDA_CHECK(x) do{ \
  cudaError_t e = (x); \
  if (e != cudaSuccess){ \
    fprintf(stderr, "CUDA %s:%d: %s\n", \
      __FILE__,__LINE__, cudaGetErrorString(e)); \
    exit(1); \
  } \
} while(0)

CUDA_CHECK(cudaMalloc(&d, n));
kernel<<<G,B>>>(d);
CUDA_CHECK(cudaGetLastError());    // launch
CUDA_CHECK(cudaDeviceSynchronize()); // async

kernel 호출 자체는 async — sync 후에야 kernel 내부 오류 확정.

8 Compilation Flow .cu→PTX→SASS

foo.cu
 └─ nvcc frontend
     ├─ host C++ → g++/clang → host object
     └─ device code
         └─ cicc → NVVM IR → libNVVM → PTX (.ptx)
             └─ ptxas → SASS (cubin, sm_80/sm_90)
                 └─ fatbin → host object 링크

• PTX = forward-compatible virtual ISA ↗ V03 §1
• SASS = architecture-specific 실제 opcode ↗ V04 §12
• JIT: 런타임에 driver가 PTX → 신 SASS 생성 가능
• -arch=sm_80 Ampere · -arch=sm_90 Hopper

9 Memory Management 보조 API

1 왜 Multi-D ?

• 이미지 · 행렬 · volume · tensor = 본질적으로 다차원
• threadIdx·blockIdx 모두 (x, y, z) → 자연스런 mapping
• index 계산 오류 ↓ · 가독성 ↑

dim3 block(16, 16);                        // 256 thread
dim3 grid((W+15)/16, (H+15)/16);     // ceil
kernel<<<grid, block>>>(...);

2 1D Indexing ★

for(int i = blockIdx.x*blockDim.x + threadIdx.x;
         i < N;
         i += gridDim.x*blockDim.x){
  out[i] = f(in[i]);
}

3 2D Indexing ★ row·col

4 3D Indexing volume

dim3 blk(8, 8, 8);          // 512 thread
dim3 grd((W+7)/8,
         (H+7)/8,
         (D+7)/8);

MRI · CT · fluid sim · 3D stencil에 쓰임.

5 Row-major Linearization

• PyTorch 기본 = row-major (C contiguous)
• stride 벡터로 일반화 ↗ V06 §8 CuTe Layout

6 2D Kernel Skeleton ★

__global__ void rgb2gray(uchar* out,
                         const uchar* in,
                         int W, int H){
  int c = blockIdx.x*blockDim.x + threadIdx.x;
  int r = blockIdx.y*blockDim.y + threadIdx.y;
  if (r < H && c < W){
    int i = r*W + c;
    const uchar* p = &in[3*i];
    out[i] = 0.21f*p[0]
           + 0.72f*p[1]
           + 0.07f*p[2];
  }
}

7 Naive Matrix Multiplication

__global__ void matMul(float* C,
                        const float* A,
                        const float* B, int N){
  int r = blockIdx.y*blockDim.y + threadIdx.y;
  int c = blockIdx.x*blockDim.x + threadIdx.x;
  if (r < N && c < N){
    float s = 0.0f;
    for (int k = 0; k < N; ++k)
      s += A[r*N + k] * B[k*N + c];
    C[r*N + c] = s;
  }
}

8 Block Size 경험칙 32의 배수

• thread 수 = 32의 배수 (warp 단위)
• 128 / 256 / 512 에서 탐색 시작
• 16×16 = 256 (warp 8) · 32×32 = 1024 (max)
• 너무 작음 → occupancy ↓ · 너무 큼 → register 부족
• __launch_bounds__(maxT, minB)로 컴파일러 힌트

9 Boundary Pattern

1. grid 계산 = ⌈N/B⌉ 올림 division
2. kernel 진입 직후 if (r<H && c<W) return;
3. padding 로드 시 0 또는 neutral element

1 실행 5-level ★

Grid
 └─ (Hopper+) Cluster        ← cluster.sync · DSMEM
     └─ Block (CTA)          ← __syncthreads
         └─ Warp = 32 thread ← __syncwarp · shfl
             └─ Thread       ← register · PC

• Cluster는 sm_90+ 전용 ↗ V02 §10
• Block 간 독립 가정 → transparent scalability
• Warp = HW scheduling 단위 (scalar thread 아님)

2 SM 내부 구조 A100 예

SM (108개)
 ├─ 4 processing block (partition)
 │   ├─ warp scheduler × 1
 │   ├─ INT32 unit × 16
 │   ├─ FP32 unit × 16
 │   ├─ FP64 unit × 8
 │   ├─ Tensor Core (3rd gen) × 1
 │   ├─ LD/ST unit × 8
 │   └─ SFU × 4     (sin/exp/rsqrt)
 ├─ Register File (64K × 32b = 256KB)
 ├─ L1 / Smem 통합 (192 KB, carveable)
 ├─ L1 instruction cache
 └─ Constant cache (read-only)

Source: NVIDIA A100 GPU Architecture Whitepaper v1.0. Hopper 내부 상세 ↗ V02 §1.

3 SM Count 비교

4 Warp ★ 정의 32 thread · SIMT

• block → warp 자동 분할, x축 우선으로 32씩
• blockDim이 32 배수가 아니면 마지막 warp가 partial (비활성 lane)
• Independent Thread Scheduling (Volta+) → thread별 PC 가능 → __syncwarp() 필수

5 SIMT vs SIMD

6 Warp Scheduler State E·I·S

• Eligible: 의존성 해소, 실행 준비 완료
• Issued: 이번 cycle에 dispatch 됨
• Stalled: barrier / long-latency / scoreboard 대기

7 Thread Hierarchy 매트릭스

8 Transparent Scalability ★

• ∴ block 간 동기화 API가 의도적으로 부재
• grid-level 동기는 kernel 분할로 달성
• 또는 cooperative groups grid_group::sync() (persistent kernel ↗ V05 §9)

9 Warp Formation 규칙

• block (Bx, By, Bz) → 1D linearize t = tx + ty·Bx + tz·Bx·By
• warp id = t / 32 · lane id = t % 32
• 같은 warp ⟺ t / 32 동일
• ∴ 16×16 block → warp 0 = (0..15, 0), warp 1 = (0..15, 1) 등

1 Control Divergence ★ uniform warp

if (tid % 2 == 0)    // ❌ warp ½ mask
  even_work();
else
  odd_work();             // 두 경로 직렬

2 Divergence 비용 공식

• 단, 모든 thread 가 skip 하면 whole-warp skip (cost 0)
• early-return 이 mid-loop break 보다 저렴 (compiler 가 predicate)

3 분기 회피 패턴

1. 조건 산술화: y = c ? a : b → y = a*c + b*(1-c)
2. boundary check는 가장 바깥 block 에서만
3. warp 별 uniform branch 유지 (tid/32 기준)
4. predicate instruction: compiler가 짧은 분기를 @p instr 로 변환

4 Synchronization 3종

5 Latency Hiding ★

6 Occupancy 정의 ★ active / max

7 Occupancy 제약 4요소 B·T·R·S

8 Occupancy ≠ Performance

• 높은 occupancy = latency hiding 여력
• 너무 높이면 register spill → local memory (DRAM) → 오히려 ↓
• GEMM · attention 은 reg 많이 씀 → 25~50% 가 정상
• ILP / MLP 로 보완 가능 (한 thread 가 여러 독립 op 동시 in-flight)

9 __launch_bounds__

__global__
__launch_bounds__(256, 4)   // T, minBlocks/SM
void k(...){ ... }

• 컴파일러에게 "이 block 크기와 최소 occupancy 맞춰라" 지시
• reg 사용을 제한해 지정 occupancy 달성
• -Xptxas -v 로 실제 reg/smem 확인

1 메모리 계층 정량표 ★★ 전상지공레텍

Source: A100 / H100 Whitepaper. latency는 근사치, 세대별 실측 ↗ V02 §14.

2 Bandwidth 비교

3 A100 vs H100 정량 비교 ★★ 10+ metric

세대별 신기능 표 ↗ V02 §7, §8.

4 Local Memory ⚠ 이름이 속임수

• compiler 가 register 부족 시 자동 allocate
• -Xptxas -v 로 "X bytes stack frame" 확인
• 해법: loop fuse, 배열 크기 compile-time 상수화

5 PTX Storage Class 매핑

PTX state space 상세 ↗ V03 §3.

6 Compute / Memory Ratio ★ CGMA · AI

Roofline 상세 공식과 multi-tier 차트 ↗ V18 §1~3.

1 Tiling 핵심 ★★

1. Load phase: A 타일 · B 타일 → smem (협력 로드)
2. __syncthreads()
3. Compute phase: smem 에서 내적 부분합 누적
4. __syncthreads() · 다음 phase

2 Reuse Factor 분석

3 Tile Size Trade-off

4 Tiled Matmul ★★ shared mem full listing

#define TILE 32

__global__ void matMulTiled(
    float* C, const float* A,
    const float* B, int N){

  __shared__ float As[TILE][TILE];
  __shared__ float Bs[TILE][TILE];

  int bx = blockIdx.x, by = blockIdx.y;
  int tx = threadIdx.x, ty = threadIdx.y;

  int row = by*TILE + ty;
  int col = bx*TILE + tx;
  float acc = 0.0f;

  // phase 순회 (K 방향)
  for(int ph = 0; ph < (N+TILE-1)/TILE; ++ph){

    // ── 1. 협력 로드 ──
    int aCol = ph*TILE + tx;
    int bRow = ph*TILE + ty;
    As[ty][tx] = (row < N && aCol < N)
               ? A[row*N + aCol] : 0.0f;
    Bs[ty][tx] = (bRow < N && col < N)
               ? B[bRow*N + col] : 0.0f;
    __syncthreads();

    // ── 2. 부분 내적 ──
    #pragma unroll
    for(int k = 0; k < TILE; ++k)
      acc += As[ty][k] * Bs[k][tx];

    __syncthreads();
  }

  if(row < N && col < N)
    C[row*N + col] = acc;
}

5 Double Buffering 개념

iter i   : load → B      compute ← A
iter i+1 : load → A      compute ← B
iter i+2 : load → B      compute ← A

• smem 2배 필요 (occupancy ↓ 가능)
• Ampere cp.async 로 자동화 ↗ V03 §6
• Hopper TMA + mbarrier ↗ V04 §4

6 Thread Coarsening

7 Boundary Tile

• N 이 TILE 배수 아닐 때 out-of-range → 0.0f 로드 (neutral)
• 출력 쓰기 직전 if(r<N && c<N)
• zero padding → 수치적 안전 (summation 에 0 기여)

1 Coalesced Access ★★ warp → 1 trans

✓ coalesced:  t0→a[0],  t1→a[1],  ... t31→a[31]   (128 B)
✗ strided  :  t0→a[0],  t1→a[N],  ...             (N trans)
✗ random   :  t0→a[?]                              (최대 32 trans)

2 Coalescing 조건

1. warp 전체가 동일 명령 (not divergent)
2. 주소 = base + stride·lane, stride = element size
3. base 가 128 B 정렬 (preferred)
4. elements 합 = 32 · 4B = 128 B (or 256B for 8B type)

3 Row vs Col Major 접근

4 AoS vs SoA

// AoS (Array of Struct) — ❌ strided per-field
struct Atom { float x, y, z, m; };
Atom A[N];
__global__ void k(Atom* A){
  int i = ...;
  float x = A[i].x;      // stride 16B · 25% eff
}

// SoA (Struct of Array) — ✓ coalesced
struct Atoms {
  float* x; float* y;
  float* z; float* m;
};
float x = A.x[i];         // stride 4B · 100% eff

LLM tensor layout 은 자연스럽게 SoA (각 tensor 가 별도 배열).

5 Shared Memory Bank 구조

6 Bank Conflict 예시

s[tid]       : bank(tid) → 모두 다른 bank ✓ OK
s[2*tid]     : bank 0,2,4,...,62 → 2-way conflict
s[32*tid]    : 모두 bank 0        → 32-way (완전 직렬)
s[tid ^ pad] : XOR swizzle        → conflict 회피

• A[32][32] col 접근 = 32-way conflict
• A[32][33] padding → conflict 0 (한 칸 offset)
• swizzle: XOR-based 재매핑 ↗ V06 §11

7 Padding Trick

// ❌ 32-way on column read
__shared__ float A[32][32];

// ✓ no conflict · 1 col wasted
__shared__ float A[32][33];

• 비용: smem 3% 낭비 (32→33)
• 이득: 32× throughput
• vectorized load (float4) 시에는 padding 전략 달라짐

8 Vectorized Load/Store

• float4 load = 128 B / 32 thread (1 tx)
• int4, double2 동일 효과
• alignment 필수: __align__(16) 또는 reinterpret_cast<float4*>
• PTX ld.global.v4.f32 → 4-way vectorized

9 Latency Hiding 역산

1 1D · 2D Convolution 정의

• 각 출력점 = 주변 (2R+1) 개 입력의 가중합
• F 는 작고 (5×5, 7×7) 전 thread 가 재사용

2 Naive 1D Convolution

__global__ void conv1d(float* y,
    const float* x, const float* F,
    int N, int R){
  int i = blockIdx.x*blockDim.x + threadIdx.x;
  if(i >= N) return;
  float s = 0.0f;
  for(int k = -R; k <= R; ++k){
    int j = i + k;
    if(j >= 0 && j < N)
      s += x[j] * F[k + R];
  }
  y[i] = s;
}

3 3단계 최적화 ★ 상타할

1. 상수 memory: F 를 __constant__ 로 (broadcast cache)
2. 타일링: input tile 을 smem 에 로드
3. 할로: tile + R 테두리 확장 로드 (neighbor 접근)

4 Constant Memory 패턴

__constant__ float F[25];    // 5×5 filter

// host
cudaMemcpyToSymbol(F, h_F, sizeof(F));

// device
__global__ void conv(...){
  ...
  s += X[...] * F[k + R];   // broadcast
  ...
}

• 64 KB 총량 · warp 전체가 같은 주소 읽으면 1 cycle
• 서로 다른 주소 → 직렬화
• cached → 두 번째 접근부터 L1-cache 수준

5 Tiled + Halo 구조 ★

tile: TILE=16, R=2
┌──────── tile + halo ────────┐
│ halo ▒▒ │  inner tile  │ ▒▒ │
│      R  │   TILE       │  R │
└─────────┴──────────────┴────┘
smem 크기 = (TILE + 2R)
output 범위 = TILE
input 범위 = TILE + 2R

• block thread 수 = TILE (또는 TILE+2R 협력)
• halo thread 가 경계 예외 처리
• 경계 밖 → zero (또는 mirror/clamp/wrap)

6 Halo 로드 패턴

__shared__ float sX[TILE + 2*R];
int tx = threadIdx.x;
int gi = blockIdx.x*TILE + tx;

// 중앙 로드
sX[tx + R] = (gi < N) ? x[gi] : 0;

// 좌측 halo
if(tx < R){
  int li = gi - R;
  sX[tx] = (li >= 0) ? x[li] : 0;
}
// 우측 halo 유사
__syncthreads();

7 Boundary 처리 모드

8 Separable Convolution

• Gaussian · box · Sobel 등 분리 가능
• 임시 버퍼 T = x * f_row, y = T * f_col
• K = 5 → 25 → 10 연산 (2.5× 절약)

9 CGMA 개선 예 2D 5×5

1 Stencil 정의

2 3D Tiled Stencil

3D volume → z-axis slab 순회
slab[z]        : smem (xy-plane, tile + halo)
slab[z-1, z+1] : register (z-neighbor)
                 → smem 절약
┌──────┐
│ halo │   reg: u[z-1], u[z], u[z+1]
│ xy   │   smem: tile(z) + halo
│ halo │
└──────┘

• brick / 3.5D blocking 기법
• read-read-write pattern 유지
• thread coarsening (2~4개 z 담당) 효과 큼

3 안정성 조건 (PDE)

• double buffer: u_old ↔ u_new 교대
• halo 초기화 = Dirichlet / Neumann / periodic 경계 조건

4 Histogram 정의

__global__ void histNaive(
    int* h, const uchar* x, int N){
  int i = blockIdx.x*blockDim.x + threadIdx.x;
  if(i < N) atomicAdd(&h[x[i]], 1);
}

5 Privatization ★ 사프머

1. block 마다 private smem hist 소유
2. thread 가 smem atomic 으로 카운트 (빠름)
3. __syncthreads()
4. block 결과 → global atomic merge

6 Privatized Histogram Kernel

__global__ void histP(int* h,
    const uchar* x, int N, int B){
  __shared__ int s[256];
  for(int t = threadIdx.x; t < B; t += blockDim.x)
    s[t] = 0;
  __syncthreads();

  // 각 thread grid-stride
  for(int i = blockIdx.x*blockDim.x + threadIdx.x;
       i < N;
       i += gridDim.x*blockDim.x)
    atomicAdd(&s[x[i]], 1);
  __syncthreads();

  for(int t = threadIdx.x; t < B; t += blockDim.x)
    atomicAdd(&h[t], s[t]);
}

7 Aggregation 기법

• thread local counter 유지, bin 이 같으면 누적
• bin 이 바뀔 때만 atomic
• image spatial locality 클수록 이득 ↑
• MoE / sort-then-reduce 계열로 확장 ↗ V05 §5

int curBin = -1, count = 0;
for(i = ...; i < N; i += stride){
  int b = x[i];
  if(b == curBin) ++count;
  else{
    if(curBin >= 0)
      atomicAdd(&s[curBin], count);
    curBin = b; count = 1;
  }
}

8 Large Bin 전략

• bin 수 > smem 용량 (예 65536) → 단일 block 에 못 넣음
• 해법 1: block 별 부분 bin 담당 (bin 범위 분할)
• 해법 2: 2-pass (coarse hist → fine hist)
• 해법 3: register privatization (적은 bin)

9 Atomic API 요약

PTX scope qualifier (.cta / .gpu / .sys) ↗ V03 §11.

1 Reduction 정의

부동소수는 결합법칙 근사적 성립 — 순서 다르면 미세 차이. ↗ V09 §7

2 Stage 1 — Naive Divergent ★

// ❌ warp 내 절반씩 mask → divergence ↑
__shared__ float s[BS];
s[tid] = input[i]; __syncthreads();
for(int st = 1; st < blockDim.x; st *= 2){
  if(tid % (2*st) == 0)
    s[tid] += s[tid+st];
  __syncthreads();
}

3 Stage 2 — Interleaved Addressing

// 연속 thread가 동작 — divergence ↓
for(int st = 1; st < blockDim.x; st *= 2){
  int idx = 2*st*tid;
  if(idx < blockDim.x)
    s[idx] += s[idx+st];
  __syncthreads();
}

• tid 연속 → warp 앞쪽만 활성
• 단 bank conflict 발생 (stride 2, 4, ... 증가)

4 Stage 3 — Sequential Addressing ★

// ✓ divergence 0 + bank conflict 0
for(int st = blockDim.x/2; st > 0; st >>= 1){
  if(tid < st)
    s[tid] += s[tid+st];
  __syncthreads();
}

• active thread 가 warp 의 앞쪽에 몰림 → divergence 없음
• stride = 1 단위 → bank conflict 없음
• thread 절반은 idle (해결 Stage 4)

5 Stage 4 — First Add on Load

// load 단계에서 2개 합산
int i = blockIdx.x*(blockDim.x*2)
      + threadIdx.x;
s[tid] = input[i] + input[i + blockDim.x];
__syncthreads();
// 이후 sequential reduction

• grid 크기 ½ → block 수 반감
• thread 0 번째 load 에서 ½ work 이미 수행
• DRAM throughput ↑

6 Stage 5 — Unroll Last Warp

// warp (32 thread) 는 implicit sync
if(tid < 32){
  volatile float* vs = s;
  vs[tid] += vs[tid+32];
  vs[tid] += vs[tid+16];
  vs[tid] += vs[tid+8];
  vs[tid] += vs[tid+4];
  vs[tid] += vs[tid+2];
  vs[tid] += vs[tid+1];
}

7 Stage 6 — Warp Shuffle ★★

// smem 불필요 · 레지스터만
float v = s[tid];
if(tid < 32){
  v += __shfl_down_sync(0xffffffff, v, 16);
  v += __shfl_down_sync(0xffffffff, v, 8);
  v += __shfl_down_sync(0xffffffff, v, 4);
  v += __shfl_down_sync(0xffffffff, v, 2);
  v += __shfl_down_sync(0xffffffff, v, 1);
  if(tid == 0) out[blockIdx.x] = v;
}

• lane 간 direct register exchange
• smem 사용 0 · __syncthreads 불필요
• __shfl_down_sync(mask, v, offset)

8 Stage 7 — Multi-block / Library

1. block level: warp shuffle 로 partial sum → smem
2. block 0 (warp 0) 이 block sum 들을 다시 reduce
3. 또는 atomicAdd(&out[0], blockSum)
4. 실무: cub::DeviceReduce / thrust::reduce

9 7-Stage 진화 요약표

1 Scan 정의 ★

• ⊕ = 임의의 결합 연산 (+, max, *, string concat …)
• stream compaction · radix sort · sparse 의 기반

2 Kogge-Stone work = N log N

__shared__ float s[BS];
s[tid] = x[i]; __syncthreads();
for(int o = 1; o < BS; o *= 2){
  float v = (tid >= o) ? s[tid - o] : 0;
  __syncthreads();
  if(tid >= o) s[tid] += v;
  __syncthreads();
}

3 Brent-Kung ★ up·down

1. Up-sweep: pairwise 합을 binary tree 위로 (reduction 과 동일)
2. tree 맨 위 원소를 identity (= 0) 로 초기화
3. Down-sweep: 왼쪽·오른쪽 합 교환·누적하며 내려옴

4 Brent-Kung Up-Sweep

for(int st = 1; st < BS; st *= 2){
  int idx = (tid + 1) * 2 * st - 1;
  if(idx < BS)
    s[idx] += s[idx - st];
  __syncthreads();
}
// 맨 위 원소 초기화 (exclusive 의 경우)
if(tid == 0) s[BS-1] = 0;

5 Brent-Kung Down-Sweep

for(int st = BS/2; st > 0; st /= 2){
  __syncthreads();
  int idx = (tid + 1) * 2 * st - 1;
  if(idx < BS){
    float t = s[idx - st];
    s[idx - st] = s[idx];
    s[idx] += t;
  }
}

6 3-Phase Hierarchical Scan

1. 각 block 이 local scan (shared mem)
2. block sum 수집 → aux[]
3. aux 를 scan
4. block i 의 모든 원소에 aux[i−1] 더함

┌ blk0 ┐┌ blk1 ┐┌ blk2 ┐
 local   local   local
   Σ0      Σ1      Σ2
   └────── scan aux ─────┘
   +0      +Σ0     +Σ0+Σ1

7 Scan 알고리즘 비교표 ★

고급 변종 (segmented · warp-level · CUB) ↗ V05 §4.

8 Scan 응용 ★

• Stream compaction: flag[i] ∈ {0,1} scan → 밀집 output index
• Radix sort: bit 별 scan + scatter
• CSR build: row 길이 → row_ptr (exclusive scan) ↗ §14
• quicksort partition · histogram CDF
• LLM MoE token permutation ↗ V08 §2

9 Warp-level Scan

// inclusive scan within warp
float v = x;
for(int o = 1; o < 32; o *= 2){
  float t = __shfl_up_sync(
             0xffffffff, v, o);
  if(lane >= o) v += t;
}

1 Parallel Merge 정의

2 Co-rank Binary Search

int coRank(int k, const T* A, int m,
                   const T* B, int n){
  int iLow = max(0, k - n);
  int iHigh = min(k, m);
  while(1){
    int i = (iLow + iHigh)/2;
    int j = k - i;
    if(i > 0 && j < n && A[i-1] > B[j])
      iHigh = i - 1;
    else if(j > 0 && i < m && B[j-1] > A[i])
      iLow = i + 1;
    else return i;
  }
}

O(log(m+n)) per thread · total O((m+n) log(m+n)) work.

3 Merge Path

A →
│A[0] A[1] A[2] A[3]
│  ╱╱╱
B↓ ╱╱╱
B[0]  → 대각선 = merge path
B[1]
B[2]

• 각 대각선 위의 점 → 하나의 출력 구간
• thread block 이 하나의 tile 담당
• block 간 완전 독립 → grid scale

4 Sort 방식 비교

실무: cub::DeviceRadixSort · thrust::sort. 직접 구현은 특수 요건일 때만.

5 Radix Sort 원리 ★

1. LSB (또는 MSB) 부터 b-bit 씩
2. 각 pass: 각 key 의 그 b-bit 추출
3. histogram (2^b bin) + scan → scatter index
4. key 를 output 위치에 scatter
5. 상위 bit 로 반복

6 Radix Sort 핵심 Kernel

// 1 pass of b-bit radix
int key = in[i];
int digit = (key >> shift) & MASK;

// block-level hist + scan
atomicAdd(&hist[digit], 1);
__syncthreads();
blockExclusiveScan(hist);

// scatter
int dst = hist[digit]
          + blockOffset(digit);
out[dst] = key;

7 Bitonic Sort

• SIMT 친화적 (branch-free)
• 작은 N · power-of-2 · warp 단위 sort 에 유리
• top-k selection 기반 구현

bitonic network log²N step
stage 1: compare (i, i±1)
stage 2: compare (i, i±2)
...
stage log N: compare (i, i±N/2)
각 stage = log N substep

8 Sort 선택 결정 트리

• integer key · N 큼 → radix sort
• float32 · stable 필요 없음 → radix (IEEE 754 비트 tweaking)
• 일반 compare-sort · 작은 N → bitonic / block sort
• top-k · 큰 K → radix top-k ↗ V05 §6
• top-k · 작은 K → bitonic top-k

9 Implementation Notes

1 Sparse Matrix Format 비교 ★

2 CSR 구조

A = [1 0 2 0]     rp=[0,2,3,5,6]
    [0 0 3 0]     ci=[0,2,2,1,3,0]
    [0 4 5 0]     val=[1,2,3,4,5,6]
    [6 0 0 0]
V=4, nnz=6
row r 의 원소 = ci, val [rp[r]..rp[r+1])

• row_ptr 크기 = V+1 (exclusive scan of row 길이)
• col_idx 크기 = nnz
• val 크기 = nnz

3 SpMV (CSR)

__global__ void spmv_csr(
    float* y, const int* rp,
    const int* ci, const float* val,
    const float* x, int V){
  int r = blockIdx.x*blockDim.x + threadIdx.x;
  if(r >= V) return;
  float s = 0;
  for(int k = rp[r]; k < rp[r+1]; ++k)
    s += val[k] * x[ci[k]];
  y[r] = s;
}

vector · CSR-stream · merge-based kernel 전략 ↗ V05 §13.

4 CSR Build = Scan !

• step 1: row 별 nnz 카운트
• step 2: exclusive scan → row_ptr
• step 3: (r, c) 쌍을 row_ptr 위치로 scatter
• §12 scan 직접 활용 ★

5 Graph = Sparse Matrix

V=5, E=7 (directed)
row_ptr: [0, 2, 3, 5, 6, 7]
col_idx: [1,2, 3, 0,4, 4, 2]
(v=0 neighbors: 1,2 · v=1: 3 · v=2: 0,4 · ...)

6 BFS — Level-Sync ★

• source 로부터 level 단위 동심원 확장
• 각 level = 1 kernel launch (grid-level sync = kernel 경계)
• depth = O(diameter) ≪ V

7 Vertex-Centric BFS (naive)

__global__ void bfsV(
    const int* rp, const int* ci,
    int* lvl, int L, int* done){
  int v = blockIdx.x*blockDim.x + threadIdx.x;
  if(lvl[v] == L){
    for(int k = rp[v]; k < rp[v+1]; ++k){
      int u = ci[k];
      if(lvl[u] == INF){
        lvl[u] = L + 1;
        *done = 0;
      }
    }
  }
}

8 Frontier-Based BFS ★ 프론

1. 현재 frontier 큐 유지 (크기 ≪ V)
2. thread 당 frontier 원소 → 이웃 탐색
3. 새 vertex 는 atomic으로 next frontier push
4. frontier swap · level++

9 Direction-Optimizing (Beamer)

• social graph (small-world) 에서 큰 이득
• heuristic: |frontier| / V 임계값 (≈ 1/20)
• Privatization → block smem next-frontier → global atomic 1회

1 학습 파이프라인

1. Forward: x → (Conv / FC / Norm / Act)… → ŷ
2. Loss: L(ŷ, y)
3. Backward: ∂L/∂w · ∂L/∂x (chain rule)
4. Update: w ← w − η · ∂L/∂w (SGD / Adam)

mini-batch = data-parallel 병렬 축 · distributed 학습은 ↗ V17

2 FC Layer = GEMM

• §7 tiled matmul 그대로 사용
• cuBLAS / CUTLASS ↗ V06
• Tensor Core FP16/BF16 input · FP32 accumulate

3 Convolution Layer

5-nested loop → naive kernel CGMA 낮음 → 재구성 필요.

4 im2col + GEMM ★ 펴서 곱하기

Conv = (unfold X) × (reshape W)
X_col : [C·Kh·Kw, N·Oh·Ow]  ← patch 중복 저장
W_row : [M, C·Kh·Kw]
Y     : W_row · X_col       ← GEMM!

• X_col 중복 저장 → memory ↑, GEMM 재사용 큰 이득
• cuDNN 이 내부 알고리즘 자동 선택

5 Conv 알고리즘 선택

cuDNN findAlgorithm 자동 벤치.

6 Backward = 3 GEMM

• conv backward: transposed conv 로 구현
• recompute (gradient checkpoint) → memory ↓ ↗ V17 §7

7 Mixed Precision ★

• BF16: FP32 와 동일 exponent → loss scaling 불필요
• FP16: loss scaling 필요 (underflow 방지)
• FP8: per-tensor amax history · delayed scaling
• dtype 상세 ↗ V09 §2~5

8 Normalization Kernel

• 2-pass (mean → var) → 1-pass Welford
• warp reduction 으로 row-wise stats
• residual · dropout fused 가능 ↗ V08 §8

9 Softmax · Activation

• 2-pass → online softmax 1-pass ↗ V07 §2
• ReLU / GELU / SwiGLU: element-wise · fuse into prev GEMM epilogue ↗ V06 §16
• attention 은 softmax + matmul 의 합성 kernel ↗ V07

1 왜 Multi-GPU?

• 단일 노드 GPU 메모리 한계 (A100/H100 80 GB)
• 대형 과학 시뮬 · LLM 학습 · 추론
• 노드 간: MPI (message passing)
• 노드 내: CUDA + NCCL (shared · NVLink)

Node0 [GPU0 GPU1 ...] ─┐
Node1 [GPU0 GPU1 ...] ─┼─ IB / NVLink Switch
...                    │
  └─ MPI_COMM_WORLD ───┘

2 Device Management

int n;
cudaGetDeviceCount(&n);        // 노드 내 GPU 수
cudaSetDevice(rank % n);       // rank → GPU 매핑

cudaDeviceProp p;
cudaGetDeviceProperties(&p, id);
// p.name, p.totalGlobalMem, p.multiProcessorCount ...

• 프로세스당 1 GPU 권장 (context 충돌 회피)
• SLURM / PBS 가 rank → GPU 배정

3 Peer Access

int canAccess;
cudaDeviceCanAccessPeer(&canAccess,
                          src, dst);
if(canAccess){
  cudaSetDevice(src);
  cudaDeviceEnablePeerAccess(dst, 0);
  // 이후 dst 포인터 직접 read/write OK
}
cudaMemcpyPeer(d_dst, dst, d_src, src, n);

NVLink 통한 P2P DMA · PCIe 도 가능 (BW 낮음).

4 MPI 핵심 API

5 Domain Decomposition ★

┌──────┬──────┬──────┐
│rank0 │rank1 │rank2 │
│  ▒▒▒ │▒▒▒▒▒▒│ ▒▒▒  │  ← halo 교환
└──────┴──────┴──────┘
1D · 2D · 3D 분할 (surface/volume ratio 고려)

6 전형 Loop

for(step){
  MPI_Isend(halo_right, right_rank, ..., &req[0]);
  MPI_Irecv(halo_left,  left_rank,  ..., &req[1]);

  compute_interior<<<>>>(...);   // overlap

  MPI_Waitall(2, req, status);
  compute_boundary<<<>>>(...);
}

7 CUDA-aware MPI ★

• GPU 포인터를 MPI_Send 에 직접 전달 가능
• 내부적으로 GPUDirect RDMA → CPU 경유 ×
• OpenMPI · MVAPICH2 · HPC-X 지원
• 컴파일 확인: ompi_info --all | grep mca_btl_openib_cuda

8 NCCL vs MPI

collective · topology · ring/tree 상세 ↗ V15.

9 Scaling 개념

• α-β 통신 모델 · ring allreduce 알고리즘 ↗ V15 §3
• TP / PP / DP / EP / SP / CP 병렬화 축 ↗ V15 §7~13

1 7 Parallel Pattern 1-liner

2 Indexing 공식 5종 ★

3 Occupancy 3제약

4 메모리 6종 1-line

5 흔한 실수 10선 ★

1. host 포인터를 kernel에 전달 → segfault
2. cudaMemcpy 방향 플래그 혼동
3. launch 직후 오류 체크 누락 (cudaGetLastError)
4. boundary if 빠짐 → OOB
5. __syncthreads() 를 분기 안쪽
6. shared mem 초기화 race (sync 누락)
7. strided global access → BW 낭비
8. shared mem bank conflict (32-way)
9. auto 배열 → local spill → DRAM
10. FP16 accumulate → overflow · 결과 망가짐

6 수식 Cheat

7 니모닉 총정리

8 패턴 뱃지

공통 변환: tile + smem + coalesce + atomic 최소화 + scan 기반.

9 18권 전체 의존 지도

V01 (PMPP)
 │
 ├─► V02 (HW)    ─► V03 (PTX) ─► V04 (Hopper)
 │                                   │
 ├─► V05 (Primitive)                 │
 │     │                             │
 │     └─► V06 (GEMM) ─► V07 (Attn)
 │                            │
 ├─► V09 (Numeric) ─► V10 (Quant)
 │                                   │
 └─► V15 (Distrib)                   │
           ├─► V16 (Serving)         │
           └─► V17 (Training) ◄──────┘

V11–V14: Triton · MLIR · Inductor · XLA/TVM
V18: Profiling (모든 권 뒤)

10 다음 권으로 이동

• 실리콘 정량 · 세대별 기능 → ↗ V02
• PTX ISA · mma · cp.async → ↗ V03
• Hopper TMA · WGMMA · SASS → ↗ V04
• 병렬 primitive 고급 · Stream-K → ↗ V05
• CUTLASS · CuTe · layout algebra → ↗ V06
• FlashAttention v1/v2/v3 · Paged → ↗ V07
• MoE · RoPE · RMSNorm · speculative → ↗ V08
• FP32/TF32/BF16/FP16/FP8/FP4 → ↗ V09
• GPTQ · AWQ · SmoothQuant · KV quant → ↗ V10
• Triton compiler · Layout inference → ↗ V11
• MLIR · NVPTX backend → ↗ V12
• TorchInductor · FX graph → ↗ V13
• XLA · TVM · Polyhedral → ↗ V14
• NCCL · TP/PP/DP/EP/SP/CP → ↗ V15
• vLLM · SGLang · TRT-LLM → ↗ V16
• FSDP · ZeRO · checkpointing → ↗ V17
• Roofline · stall · metric → ↗ V18

11 완성 체크리스트 ★