흔한 GPU parallelism 함정

모든 CUDA / Metal 초보 무는 다섯 가지

Lock-step ≠ order. Warp 안 thread는 program counter 공유하지만 다른 block의 warp들이 interleave해서 print 순서 어긋남. Ant #128이 Ant #1 자주 이김.
개미 너무 많으면 → 줄. 카드가 동시에 host할 수 있는 SM 수보다 block 더 많이 launch하면 나머지는 차례 기다림. 대규모 병렬? 응. 무한? 아냐.
Clock 산수 reality check. 2 GHz × 1,000 thread ≈ tick당 명령어 슬롯 200만 — 인상적인데 bound 있음. Pipeline stall, memory wait, bank conflict가 표제 숫자 깎음.
스케줄러가 부사관. 어느 warp 언제 돌지 하드웨어 결정. 코드는 briefing 쓰고 스케줄러가 순서 픽. Divergent branch, register 압박, shared-memory 한계 다 입력으로 들어감.
Warp divergence가 throughput 죽임. 32-thread warp가 if (thread_id % 2 == 0) ... else ... 실행하면 두 path 직렬화 — warp 절반이 idle인 동안 나머지가 각 branch 돌림. 사실상 throughput 절반, 심한 형태는 더 깊이 깎음.

이 중 초보가 자기도 모르게 가장 자주 디자인해 들어가는 게 warp divergence. fix는 if 통째 회피가 아니라 — 분기 결정을 warp 안에서 uniform하게 유지해서 warp 안 32 thread가 같은 path 가게.

Code

Bad: per-thread divergent branch (warp-divergent)·cuda

__global__ void bad_branch(float* out, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i >= n) return;

    // warp 안 thread마다 다른 path 픽 → divergence
    if (i % 2 == 0) {
        out[i] = sqrtf(out[i]);          // 짝수 thread는 path A
    } else {
        out[i] = out[i] * out[i];        // 홀수 thread는 path B
    }
}

Better: mask로 branch-free, 또는 warp-uniform 값으로 branch·cuda

// Branch-free: thread 무관하게 같은 수학
__global__ void good_no_branch(float* out, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i >= n) return;
    float v = out[i];
    out[i] = sqrtf(v) * (i & 1) + v * v * (1 - (i & 1));
    // 모든 thread가 두 op 다 하고 mask로 픽. 작업 더 많고 divergence 0.
}

// 또는: branch 결정이 warp별 uniform 보장
__global__ void good_warp_uniform(float* out, int n, int mode) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i >= n) return;
    // 'mode'는 이번 launch 모든 thread에 같음 → divergence 없음
    if (mode == 0) out[i] = sqrtf(out[i]);
    else            out[i] = out[i] * out[i];
}

Exercise

최근 쓰거나 읽은 CUDA 커널 아무거나 골라서 if 다 점검: 단일 warp 안에서 condition이 threadIdx에 따라 변해? 그렇다면 그게 divergence point. 각 branch 얼마나 자주 taken되는지 추정 — 50/50 정도면 비용 두 배 내고 있음. 그 다음 데이터 재구성해서 branch warp-uniform 만들거나 mask 기반 branch-free 정식으로 다시 짜. hot loop에 그런 fix 하나만 해도 1.5–2× 가속 자주.

모든 CUDA / Metal 초보 무는 다섯 가지

Code

External links

Exercise

Progress

댓글 0