흔한 GPU parallelism 함정

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

Lock-step ≠ order. Warp 안 thread는 program counter를 공유하지만, 다른 block의 warp들이 interleave해서 print 순서가 어긋나. Ant #128이 Ant #1을 자주 이겨.
개미 너무 많으면 → 줄. 카드가 동시에 host할 수 있는 SM 수보다 block을 더 많이 launch하면 나머지는 차례를 기다려. 대규모 병렬? 응. 무한? 아냐.
Clock 산수로 현실 점검. 2 GHz × 1,000 thread ≈ tick당 명령어 슬롯 200만 — 인상적이지만 한계는 있어. 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