Common GPU Parallelism Gotchas

Five things that bite every CUDA / Metal beginner

Lock-step ≠ order. Threads in a warp share a program counter, but printouts can finish out of sequence because warps from different blocks interleave. Ant #128 often beats Ant #1.
Too many ants → queue. Launch more blocks than your card has SMs to host concurrently and the extras simply wait their turn. Massive parallel? Yes. Unlimited? No.
Clock math reality check. 2 GHz × 1,000 threads ≈ 2 million instruction slots per tick — impressive but bounded. Pipeline stalls, memory waits, and bank conflicts erode the headline number.
Scheduler is the drill sergeant. Hardware decides which warps run when. Your code writes the briefing; the scheduler picks the order. Divergent branches, register pressure, and shared-memory limits all feed into it.
Warp divergence kills throughput. A 32-thread warp executing if (thread_id % 2 == 0) ... else ... serializes the two paths — half the warp is idle while the other half runs each branch. Effectively half throughput; severe forms can cut deeper.

Of these, warp divergence is the one beginners design themselves into accidentally most often. The fix isn't to avoid if entirely — it's to keep the branch decision uniform within a warp, so all 32 threads in a warp take the same 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;

    // Each thread inside a warp picks a different path → divergence
    if (i % 2 == 0) {
        out[i] = sqrtf(out[i]);          // even threads do path A
    } else {
        out[i] = out[i] * out[i];        // odd threads do path B
    }
}

Better: branch-free using mask, or branch on warp-uniform value·cuda

// Branch-free: same math regardless of 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));
    // Each thread does both ops, picks via mask. Higher work, zero divergence.
}

// Or: ensure the branch decision is uniform per warp
__global__ void good_warp_uniform(float* out, int n, int mode) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i >= n) return;
    // 'mode' is the same for all threads in this launch → no divergence
    if (mode == 0) out[i] = sqrtf(out[i]);
    else            out[i] = out[i] * out[i];
}

Exercise

Take any CUDA kernel you've written or read recently and check every if: does its condition vary by threadIdx within a single warp? If yes, that's a divergence point. Estimate how often each branch is taken; if it's roughly 50/50 you're paying double cost. Then rewrite either by restructuring data so the branch becomes warp-uniform, or by using a mask-based branch-free formulation. Even one such fix in a hot loop can be a 1.5–2× speedup.

Five things that bite every CUDA / Metal beginner

Code

External links

Exercise

Progress

Comments 0