PMPP 단권화

1 Heterogeneous Parallel Computing

• 2003 이후 clock ↑ 벽 (열·전력) → multi-core 전환
• 두 설계 trajectory:
  • Multi-core: sequential 성능 유지 (Intel 24-core 등)
  • Many-thread: throughput 극대화 (NVIDIA GPU)
• von Neumann 모델 · thread = 순차 실행 단위

2 CPU vs GPU 설계 철학

3 A100 성능 수치 기준 기종

cf. H100 FP64 ~34 TF (SM 수↑, TMA, async copy)

4 Amdahl's Law ★ 병렬화 상한 s·p·1

• sequential 부분이 peach pit · parallel 부분이 flesh
• straightforward parallel → DRAM bandwidth saturate → ~10× 수준
• on-chip memory 활용 최적화 → 100×+ 가능

5 Speed 수요 동기

1. Scientific simulation (bio molecule, climate)
2. Video/Image processing (HDR, view synth)
3. UI · computer vision · 음성
4. Deep Learning ★ : 대규모 labeled data + GPU throughput
5. Digital twin · realistic physics

6 병렬 프로그래밍 4대 난제

1. Algorithmic complexity: parallel = more work 가능 → work efficiency 확보
2. Memory-bound 다수 → Ch5, Ch6 최적화
3. Input-sensitive: 분포 불균형 → uneven work
4. Synchronization 필요 ↔ embarrassingly parallel

7 Programming Interfaces 비교

HPC 실무: MPI + CUDA 혼합 (multi-GPU · NCCL) — Ch20

8 CUDA 등장 의의 2007

• G80: GPGPU API(OpenGL/D3D) → 범용 parallel 인터페이스
• 그래픽 pixel 제약 제거 · C/C++로 직접
• 설치 기반 >10억대 · 경제성 확보

9 책의 3대 목표

1. Performance: GPU HW 이해 기반 최적화
2. Correctness: barrier·atomic·memory consistency
3. Scalability: 차세대 HW 확장성

10 책 구성 4부

1. Part I: 기초 (Ch2~6)
2. Part II: primitive patterns (Ch7~12)
3. Part III: advanced patterns (Ch13~19)
4. Part IV: MPI+CUDA, dynamic parallelism (Ch20~21)

1 Data Parallelism 정의

• 예: vector add · 이미지 픽셀 변환 · 행렬곱
• task parallelism과 구분: 서로 다른 task 동시 실행

2 CUDA 프로그램 구조 Host + Device 할복실복해

1. 할당: cudaMalloc(&d_A, size)
2. 복사 H→D: cudaMemcpy(d_A, h_A, n, H2D)
3. 실행: kernel<<<grid, block>>>(...)
4. 복사 D→H: cudaMemcpy(h_A, d_A, n, D2H)
5. 해제: cudaFree(d_A)

3 함수 한정자

__global__ 반드시 void return

4 Vector Add 커널 ★ i = bx·bD + tx

__global__ void vecAdd(float* A, 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, B=(n+T-1)/T;
vecAdd<<<B,T>>>(d_A,d_B,d_C,n);

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

• grid: block의 집합 (최대 3D)
• block: thread의 집합 (최대 3D, ≤1024 thr)
• block 간 독립 실행 — 동기화 불가
• block 내부: __syncthreads(), shared mem OK

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

6 내장 변수

7 cudaError 처리

cudaError_t e = cudaMalloc(&d,n);
if(e != cudaSuccess){
  printf("%s\n",cudaGetErrorString(e));
  exit(-1);
}

실무 매크로: CUDA_CHECK(x) 래퍼 권장

8 Unified / Pinned / Streams 간략

• cudaMallocManaged: UM, page fault로 자동 migrate
• cudaMallocHost: pinned host mem → H↔D 대역폭 ↑
• cudaStream_t: 비동기 복사/커널 겹침 (H2D || kernel || D2H)

9 워크플로 체크리스트

1. host 입력 준비
2. device 메모리 할당 크기 = N·sizeof(T)
3. H→D 복사
4. 그리드 차원 = ⌈n/T⌉ · boundary if
5. 커널 오류 check + cudaDeviceSynchronize()
6. D→H 복사 / cudaFree

1 2D/3D 그리드 왜?

• 이미지·행렬·볼륨 데이터 = 다차원 → 자연스러운 mapping
• threadIdx·blockIdx 모두 (x,y,z) 사용 가능
• 커널 내부 인덱스 계산을 명확·오류↓하게

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

2 2D 인덱스 공식 ★ row·col

3 Row-major 선형화

4 ColorToGrayscale 예 ★

__global__ void rgb2gray(uchar* out, 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;
    uchar* p = &in[3*i];
    out[i] = 0.21f*p[0]+0.72f*p[1]+0.07f*p[2];
  }
}

• grid = ⌈W/16⌉ × ⌈H/16⌉
• block 16×16 = 256 thread (warp 8개)

5 행렬곱 naive baseline

__global__ void matMul(float* C,float* A,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;
    for(int k=0;k<N;++k)
      s += A[r*N+k]*B[k*N+c];
    C[r*N+c]=s;
  }
}

6 block 크기 선택 경험칙

• 16×16 = 256 or 32×32 = 1024 (max)
• thread 수 = 32의 배수 (warp 단위)
• 너무 작음 → occupancy↓ / 너무 큼 → 레지스터 부족

7 경계 처리 패턴 if r<H && c<W

1. 내부 thread: 실제 계산
2. 초과 thread: early return
3. grid = ⌈N/B⌉ · 올림 division

8 3D 예 volume

dim3 blk(8,8,8);
dim3 grd((W+7)/8,(H+7)/8,(D+7)/8);
i = z*(H*W) + y*W + x;

MRI·CT·유체 시뮬에 사용

1 GPU 계층 ★ 전체 구조

GPU
 └─ SM (Streaming Multiprocessor) × N
     ├─ warp scheduler × 2~4
     ├─ SP/CUDA core × 64~128
     ├─ Tensor Core × 4
     ├─ Register File (64K × 32b)
     ├─ Shared Mem / L1 (128~192 KB)
     └─ LD/ST, SFU

2 Warp 정의 ★ 32 thread · SIMT

• block → warp 단위로 자동 분할 (x축 먼저)
• warp = HW 스케줄링 단위 (scalar thread ×)
• SIMT ≈ SIMD + 개별 thread 제어 흐름 허용

3 SIMT vs SIMD

4 Control Divergence ★ 발산

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

5 분기 회피 패턴

1. 조건 산술화: y = c ? a : b → y = a*c + b*(1-c)
2. 경계 조건은 마지막 block에서만 검사
3. warp당 uniform branch 유지

6 Latency Hiding ★

7 Transparent Scalability

• block 간 실행 순서·동시성 가정 ×
• 같은 바이너리가 더 많은 SM 가진 차세대 GPU에서 자동으로 빨라짐
• ∴ block 간 sync가 금지된 이유

1 Warp Scheduling

• 각 SM = multi warp scheduler (A100: 4개)
• 매 cycle: ready warp 중 선택
• issue stall 원인: memory wait · dependency · barrier

cycle: w0┐ w1┐ w2┐ w0┐
            │   │   │   │
          issue issue issue issue
        (scoreboard-free warps rotate)

2 __syncthreads() ★

3 Atomic Ch9·10 대비

• atomicAdd / Max / CAS 등
• read-modify-write 원자성 보장
• 동일 주소 경합 → 직렬화 (throughput ↓)
• shared mem atomic이 global보다 빠름

4 Occupancy ★ active / max

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

6 Occupancy ≠ 성능

• 높은 occupancy = latency hiding 여력
• 너무 높이면 레지스터 spill → 오히려 ↓
• ILP(instruction-level) / MLP로 보완 가능

7 튜닝 체크리스트

1. block size = 128~512 에서 탐색
2. -Xptxas -v로 레지스터/smem 확인
3. __launch_bounds__로 컴파일러 힌트
4. Nsight Compute로 stall reason 분석

8 Volatile / Fence

• __threadfence(): 메모리 쓰기 가시성 보장
• volatile: 레지스터 캐시 금지 (polling시)
• Independent Thread Scheduling (Volta+) → __syncwarp() 필수

1 메모리 계층표 ★★ 필수 암기 전상지공레로

2 Local ≠ 가깝다 ⚠

3 Shared Memory 패턴

__global__ void k(...){
  __shared__ float s[256];
  int t = threadIdx.x;
  s[t] = d_in[...];        // load
  __syncthreads();
  // 여러 thread가 s[] 공유 사용
}

• block 내 thread 간 공유·재사용
• 동적 크기: extern __shared__ float s[];
• 32 bank로 분할 (4B stride)

4 Constant Mem 작고 자주 읽는 것

• 64 KB, 모든 thread 같은 값 읽기 → 1 cycle broadcast
• 필터 계수 · 상수 테이블
• 동일 warp가 다른 주소 읽으면 직렬화

__constant__ float F[25];
cudaMemcpyToSymbol(F, h_F, sizeof(F));

5 Register vs Spill

• thread당 255 reg 상한 (HW)
• 초과 시 → local(DRAM) spill ⚠
• __launch_bounds__(T, minBlocks)로 상한 지시

6 Compute/Memory Ratio ★ CGMA / AI

7 Roofline 요약

GFLOPS
  ▲          _______ peak compute
  │        /
  │      /  BW·AI line
  │    /
  └─────────────────▶ AI (FLOP/B)
     ↑knee = peak/BW

AI < knee → memory-bound · AI ≥ knee → compute-bound

1 Tiling 핵심 ★★

1. Phase: A·B 타일 shared mem에 로드 (협력 로딩)
2. __syncthreads()
3. shared에서 내적 부분합
4. __syncthreads() · 다음 타일

2 효과 분석

3 Tiled MatMul 커널 ★

#define T 16
__global__ void mm(float* C,float* A,float* B,int N){
  __shared__ float As[T][T], Bs[T][T];
  int bx=blockIdx.x, by=blockIdx.y;
  int tx=threadIdx.x, ty=threadIdx.y;
  int r=by*T+ty, c=bx*T+tx;
  float s=0;
  for(int ph=0; ph<N/T; ++ph){
    As[ty][tx]=A[r*N + ph*T+tx];   // coalesced
    Bs[ty][tx]=B[(ph*T+ty)*N + c];
    __syncthreads();
    for(int k=0;k<T;++k) s += As[ty][k]*Bs[k][tx];
    __syncthreads();
  }
  C[r*N+c]=s;
}

4 타일 크기 trade-off

5 Boundary Tile 처리

• N이 TILE 배수가 아닌 경우
• out-of-range 위치 → 0.0f 로드 (padding)
• 출력 쓰기 전 if(r<N && c<N) 검사

if(r<N && ph*T+tx<N)
  As[ty][tx]=A[r*N+ph*T+tx];
else As[ty][tx]=0.0f;

6 Thread Coarsening

• 레지스터 많이 씀 → occupancy ↓ 주의
• Volta+ tensor core 매트릭스에서 자연스러움

1 Coalesced Access ★★ warp → 1 transaction

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

2 행·열 우선 차이

3 Bank Conflict ★ Shared Memory 32 banks · 4B

4 Memory Latency Hiding 복기

5 성능 최적화 체크리스트 ★ 병코재분레

1. 병렬성 충분? (warp ↑)
2. 코alesced load/store
3. 재사용 via shared/register (tiling)
4. 분기 발산 최소화 (uniform warp)
5. 레지스터·shared 균형 (occupancy)

6 툴

• nvprof (구) / Nsight Compute: kernel 프로파일
• Nsight Systems: 타임라인 (stream 겹침)
• cuda-memcheck: OOB, race

7 흔한 실수 top 5

1. H→D 복사 누락 / 방향 플래그
2. boundary if 빠짐 → OOB
3. __syncthreads() 분기 안쪽
4. shared 초기화 race (sync 없음)
5. strided global access → bandwidth 낭비

1 1D/2D Convolution 정의

• F = filter/mask (작고 재사용)
• 출력 각 점은 주변 R개와 가중합

2 Naive 커널

__global__ void conv1d(float* y,float* x,float* F,int N,int R){
  int i=blockIdx.x*blockDim.x+threadIdx.x;
  if(i>=N) return;
  float s=0;
  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. Constant mem에 filter 저장 (brodcast cache)
2. Tiled + halo: block이 입력 타일 + R 테두리 로드 (shared)
3. L2 cache 의존: 최신 GPU는 halo도 L2 hit — 단순 tile + constant filter

4 Halo 패턴 그림

┌──────────── tile+halo ────────────┐
│  halo ▒▒▒▒▒▒▒  내부 tile  ▒▒▒▒▒▒▒ │
│       ↑R cells              R cells↑│
└──────────────────────────────────┘
smem 크기 = (TILE + 2R)

5 경계 처리

• zero-padding (기본)
• mirror (영상) / clamp (에지 복제)
• wrap (circular)

6 CGMA 개선 2D 5×5 예

7 Separable Convolution

1 Stencil Ch8

2 Tiled Stencil 패턴

• Conv와 동일하게 halo 포함 타일 로드
• thread당 2~4개 출력(coarsening)으로 smem 재사용↑
• 3D는 z축 slab 순회로 smem 절약

z0 slab → z1 slab (register reuse)
 ┌──────┐
 │ +R   │   smem: xy-plane
 │ core │   reg : z-1, z, z+1
 │ +R   │
 └──────┘

3 안정성·경계

• CFL 조건: Δt ≤ Δx²/(2D) (diffusion)
• halo를 Dirichlet/Neumann로 초기화
• double buffer (u_old ↔ u_new) 필수

4 Histogram Ch9

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

5 Privatization ★ 사프머

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

6 Privatized Hist 커널

__global__ void histP(int* h,uchar* x,int N,int B){
  __shared__ int s[256];
  for(int t=threadIdx.x;t<B;t+=blockDim.x) s[t]=0;
  __syncthreads();
  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 로컬 카운터로 연속 동일 bin 적재
• bin이 바뀔 때만 atomicAdd
• 이미지 공간 지역성이 클수록 효과 ↑

1 Reduction 정의

2 Naive tree (분기 ↓)

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

3 개선 ① Contiguous thread ★

for(int st=blockDim.x/2; st>0; st>>=1){
  if(tid < st)
    s[tid] += s[tid+st];       // ✓ warp 경계 uniform
  __syncthreads();
}

• 활성 thread가 warp의 앞쪽에 몰림 → divergence ↓
• bank conflict 없음 (stride=1)

4 개선 ② thread당 여러 원소

• load 단계에서 s[tid] = in[i] + in[i+BS] (2개 합산 후 저장)
• → grid 크기 ½, global load bandwidth ↑

5 개선 ③ warp-level primitives

// 마지막 warp (tid<32) unroll
if(tid<32){
  v = s[tid] + s[tid+32];
  for(int o=16; o>0; o/=2)
    v += __shfl_down_sync(0xffffffff, v, o);
  if(tid==0) out[blk] = v;
}

__shfl_down_sync: shared mem 없이 warp 내 교환 → 레지스터만 사용.

6 전체 Reduction 흐름

1. block-level partial sum → global[blk]
2. 두 번째 kernel로 다시 reduce (작아짐)
3. 또는 atomicAdd(&out[0], blockSum)
4. CUB/thrust의 DeviceReduce 사용 권장

7 수치 주의

• 부동소수는 결합법칙 성립X → 동일 입력·다른 순서 = 다른 값
• Kahan / pairwise 합산으로 오차 ↓
• FP16 합산은 overflow 위험 → FP32 accumulator

1 Scan 정의 ★

• stream compaction · radix sort · sparse 연산의 기반

2 Kogge-Stone work = N log N

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

depth log N · 병렬 ↑ · 중복 작업 많음

3 Brent-Kung ★ work-efficient up·down

1. Up-sweep: pairwise 합을 binary tree 위로 (reduction)
2. 최상단 원소를 identity로 초기화
3. Down-sweep: 좌우 합을 교환·누적하며 내려옴

4 Segmented / Large Scan

1. block 내부 scan (shared)
2. 각 block 총합을 수집 → 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

5 응용 ★

• Stream compaction: flag[i]∈{0,1} scan → 밀집 배열의 출력 index
• Radix sort: bit별 scan으로 재배치
• CSR build: row 길이 → row_ptr (exclusive scan)
• quicksort partition · histogram CDF

6 Scan 알고리즘 비교

1 Parallel Merge Ch12

1. thread k의 출력 구간 [k·CEIL .. (k+1)·CEIL)
2. co-rank(k) → (i, j), i+j=k
3. 구간만 시퀀셜 merge

2 Sort Ch13

실무: cub::DeviceRadixSort · thrust::sort

3 Sparse Matrix Ch14

4 CSR 생성 = scan!

• row별 nnz 카운트 → exclusive scan → row_ptr
• Ch11 Scan과 직접 연결 ★

5 최종 체크리스트 ★★ 단권화 점검

6 패턴 요약

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

1 Graph 표현 CSR

V=5, E=7
row_ptr: [0,2,3,5,6,7]
col_idx: [1,2, 3, 0,4, 4, 2]

Ch14의 SpMV와 동일 구조 → Scan으로 빌드 (Ch11 재활용)

2 BFS = level-sync ★

• source에서 level 단위 동심원으로 확장
• 각 level = 1 kernel launch
• depth = O(diameter) ≪ V

3 Vertex-centric (naive)

__global__ void bfsV(int* rp,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; }
    }
  }
}

4 Edge-centric 개선

• edge 단위로 thread 배정 → degree 편차 무관
• frontier에 속한 source의 edge만 필터링 (scan)
• work-efficient ↑

5 Frontier-based ★ 프론

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

6 Direction-Optimizing (Beamer) ★

• social graph 등 small-world에서 대폭 속도 ↑
• heuristic: |frontier| / V 임계값

1 학습 파이프라인

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

mini-batch = data-parallel 병렬 축 ★

2 FC Layer = GEMM

• cuBLAS gemm 호출이 실무 표준
• Tensor Core FP16/BF16 input · FP32 accumulate

3 Convolution Layer

5중 loop → naive 커널은 CGMA 낮음. 재구성 필요.

4 im2col + GEMM ★ 펴서 곱하기

Conv = (unfold X) × (reshape W)
X_col: [K·Kh·Kw, N·Oh·Ow]
W_row: [M, K·Kh·Kw]
Y    : W_row · X_col   → GEMM!

• X_col은 중복 저장 → memory ↑ but GEMM 재사용 큰 이득
• cuDNN이 내부에서 선택: direct / im2col / FFT / Winograd

5 알고리즘 선택

cuDNN findAlgorithm이 자동 벤치

6 Backward

• dW = X^T·dY (GEMM)
• dX = dY·W^T (GEMM) — conv: transposed conv
• 세 개의 GEMM 모두 같은 tile 커널로 처리

7 Mixed Precision ★

• Loss scaling: underflow 방지
• BF16: 지수 범위 FP32와 동일 → scaling 불필요

8 Batch 병렬 + 분산

• Data parallel: batch 분할 → N GPU, gradient AllReduce (NCCL)
• Model parallel: layer 분할 (모델이 한 GPU에 못 들어갈 때)
• Pipeline parallel: stage 파이프라인 (GPipe, 1F1B)

9 라이브러리 실무

직접 커널: fused activation/norm 정도. 핵심 GEMM/Conv는 벤더 라이브러리.

1 왜 MPI + CUDA?

• 단일 노드 GPU 메모리 한계 (80 GB) 초과
• 대형 과학 시뮬·LLM 학습
• 노드 간: MPI (message passing) · 노드 내: CUDA (shared)

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

2 MPI 핵심 API

3 Domain Decomposition ★

┌──────┬──────┬──────┐
│rank0 │rank1 │rank2 │
│   ▒▒▒│▒▒▒▒▒▒│▒▒▒   │  ← halo 교환
└──────┴──────┴──────┘

• 1D / 2D / 3D 분할 (surface/volume 비 고려)
• Stencil(Ch8)의 노드간 확장판

4 전형 루프

for(step){
  MPI_Isend(halo_right, right_rank);  // 비동기
  MPI_Irecv(halo_left,  left_rank);
  compute_interior<<<>>>(...);        // 통신과 겹침
  MPI_Waitall(...);
  compute_boundary<<<>>>(...);
}

5 CUDA-aware MPI ★

• GPU 포인터를 MPI_Send에 직접 전달 OK
• GPUDirect RDMA: PCIe/IB로 GPU↔NIC 직접 DMA (CPU 경유 ×)
• OpenMPI, MVAPICH2 등 지원

6 NCCL DL용 Collective

• GPU 특화 all-reduce / broadcast / all-gather
• ring / tree 알고리즘 자동 선택
• DL에서 MPI_Allreduce 대체재로 선호

PyTorch DDP, Horovod 내부 엔진

7 Scaling 개념

Amdahl (Ch1) ↔ Gustafson 대응