2.1 Warp and the execution model

A deep look at the GPU’s core execution unit: the warp. Starting from the SIMT model, this article covers the cost of warp divergence and how to avoid it, then warp shuffle for zero-latency exchange between threads, so you have the execution-level intuition needed to write fast CUDA kernels.

Contents

• 1. Seeing warps from the hardware side
• 2. The SIMT execution model
• 3. Warp divergence: the cost of branches
• 4. Warp shuffle: exchanging data between threads
• 5. Warp-level primitives and advanced tricks
• 6. Case study: warp-level reduction
• Summary
• Self-check list
• References

---

1. Seeing warps from the hardware side

Picture a classroom with 32 students (threads). The teacher (scheduler) issues one instruction, and every student does the same action at once. That is the core idea of a warp: the smallest unit the GPU schedules and runs, always 32 consecutive threads.

1.1 How warps are formed

When a thread block is assigned to an SM, the hardware automatically groups its threads into warps in threadIdx order:

Key point: Warp partitioning is a hardware behavior; you cannot change it. A 256-thread block becomes 8 warps. If the block size is not a multiple of 32 (say 48 threads), the leftover slots in the last warp are “padded” and do no useful work, which wastes hardware resources.

1.2 Warp scheduler

Each SM has several warp schedulers (modern architectures usually have 4). Every clock cycle a scheduler can pick one ready warp and issue an instruction. That enables a useful latency-hiding pattern:

Warp A issues a memory load → waits for data (~400 cycles)
Warp scheduler switches to Warp B → runs compute
Switches to Warp C → runs compute
...
Warp A's data is ready → scheduled again

Tip: Switching between warps is zero-overhead (zero-overhead context switch), because each warp’s register state stays resident on the SM; there is nothing to save or restore. That is fundamentally different from CPU thread switching.

---

2. The SIMT execution model

2.1 SIMT vs SIMD

NVIDIA calls its model SIMT (Single Instruction, Multiple Threads), not classic SIMD (Single Instruction, Multiple Data). The difference is subtle but important:

SIMT’s idea: you write scalar single-thread code; the hardware broadcasts the same instruction to 32 threads. That lowers the bar for parallel programming, but it also creates a performance trap: branch divergence.

2.2 Execution pipeline

How a warp instruction runs:

graph LR
    A["Fetch"] --> B["Decode"]
    B --> C["Issue"]
    C --> D["Execute"]
    D --> E["Write back"]

Key traits:

• Synchronous execution: the 32 threads in a warp execute the same instruction in the same clock cycle
• Independent state: each thread has its own registers and (logically) its own program counter
• Lockstep: without branches, all threads stay in step

---

3. Warp divergence: the cost of branches

3.1 What warp divergence is

When threads in the same warp hit a conditional (if/else, switch, different loop trip counts), they need different paths. Hardware can issue only one instruction at a time, so it must serialize every path, and threads that do not take a path are temporarily masked off.

It is like telling the class “swimmers do push-ups, non-swimmers do sit-ups,” but gym rules say everyone must do the same move at once. So everyone does push-ups first (non-swimmers stand still), then everyone does sit-ups (swimmers stand still). The time is the sum of both moves, not the parallel time of each group.

3.2 Performance impact of divergence

Consider this code:

__global__ void divergent_kernel(float* data, int N) {
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    if (tid % 2 == 0) {
        // Even threads: path A (10 instructions)
        data[tid] = expf(data[tid]);
    } else {
        // Odd threads: path B (10 instructions)
        data[tid] = logf(data[tid]);
    }
}

Note: In that kernel, half of each warp’s 32 threads take if and half take else. Hardware first runs path A (16 active, 16 idle), then path B (16 active, 16 idle). Total time is roughly 2× a branch-free case.

If the condition is finer (32 threads, 32 different paths), the slowdown can reach 32×.

3.3 Quantifying divergence

You can measure with branch efficiency:

$$\text{Branch Efficiency} = \frac{\text{Non-Divergent Branches}}{\text{Total Branches}} \times 100\%$$

Nsight Compute can show the active-thread fraction of each branch. Ideally every branch is 100% active (32/32).

3.4 Avoidance strategies

Strategy 1: Align branch conditions to warps

// ❌ Not recommended: neighboring threads take different branches
if (threadIdx.x % 2 == 0) { ... }

// ✅ Preferred: threads in the same warp take the same branch
if (threadIdx.x / 32 % 2 == 0) { ... }
// Or branch on warpId
int warpId = threadIdx.x / 32;
if (warpId < 4) { ... }

Strategy 2: Replace branches with arithmetic

// ❌ Not recommended: conditional branch
if (x > 0) {
    y = x;
} else {
    y = 0;
}

// ✅ Preferred: branchless ReLU
y = x * (x > 0);  // uses bool → int implicit conversion
// Or a builtin
y = fmaxf(x, 0.0f);

Strategy 3: Data reordering

Gather data that should take the same path into contiguous slots so threads in one warp run the same path:

// Preprocess: gather positives and negatives separately
// Positives in data[0..pos_count-1]
// Negatives in data[pos_count..N-1]
// Warps that process positives then have no divergence

Strategy 4: Unify loop exit conditions In short: turn “when does this thread leave the loop” into “does the body run this iteration” via predication. Every thread walks the same loop structure, which is friendlier to branch prediction, instruction cache, and the GPU’s SIMT model.

// ❌ Not recommended: each thread loops a different number of times
for (int i = 0; i < array_length[tid]; i++) {
    // work
}

// ✅ Preferred: shared trip count, mask inside
int max_len = block_max_length;  // max within the block
for (int i = 0; i < max_len; i++) {
    if (i < array_length[tid]) {
        // work
    }
}

3.5 Independent thread scheduling (Volta+)

Starting with Volta, NVIDIA added independent thread scheduling. Each thread has a truly independent program counter and call stack, which means:

• On divergence, hardware can interleave paths more flexibly
• Thread-level sync primitives (like __syncwarp()) are supported
• Threads in a warp can sit at different code locations

Note: Independent scheduling does not remove the performance cost of divergence; it only makes scheduling more flexible. For best performance you still want to avoid divergence inside a warp.

---

4. Warp shuffle: exchanging data between threads

4.1 Why warp shuffle

The classic way to exchange data between threads goes through shared memory:

Thread A writes shared memory → __syncthreads() → Thread B reads shared memory

Two problems: (1) it burns scarce shared memory; (2) the barrier adds latency.

Warp shuffle lets threads in the same warp read each other’s registers directly, with no shared memory and no explicit sync, like passing notes across the desk instead of putting them on the podium first.

4.2 Four shuffle instructions

Signatures for all shuffle ops:

T __shfl_sync(unsigned mask, T var, int srcLane, int width=32);
T __shfl_up_sync(unsigned mask, T var, unsigned delta, int width=32);
T __shfl_down_sync(unsigned mask, T var, unsigned delta, int width=32);
T __shfl_xor_sync(unsigned mask, T var, int laneMask, int width=32);

Parameters:

• mask: which threads participate; usually 0xFFFFFFFF for all 32
• var: the value the current thread shares
• width: logical sub-warp width (2, 4, 8, 16, or 32)

4.3 Shuffle illustrated

With 8 threads as a simplified demo (real warps have 32):

Initial: Lane[0]=A  Lane[1]=B  Lane[2]=C  Lane[3]=D  Lane[4]=E  Lane[5]=F  Lane[6]=G  Lane[7]=H

__shfl_down_sync(mask, var, 2):
  Lane[0] gets Lane[2]'s value C
  Lane[1] gets Lane[3]'s value D
  Lane[2] gets Lane[4]'s value E
  ...

__shfl_xor_sync(mask, var, 1):  (swap adjacent pairs)
  Lane[0] ↔ Lane[1]
  Lane[2] ↔ Lane[3]
  Lane[4] ↔ Lane[5]
  Lane[6] ↔ Lane[7]

4.4 Example: broadcast inside a warp

__device__ float warp_broadcast(float val, int srcLane) {
    // Every thread gets srcLane's val
    return __shfl_sync(0xFFFFFFFF, val, srcLane);
}

// Usage: let every thread in the warp see lane 0's result
float result = compute_something();
float broadcast_val = warp_broadcast(result, 0);

4.5 Example: sum inside a warp

__device__ float warp_reduce_sum(float val) {
    // Butterfly reduce: log2(32) = 5 steps
    for (int offset = 16; offset > 0; offset >>= 1) {
        val += __shfl_down_sync(0xFFFFFFFF, val, offset);
    }
    return val;  // only lane 0 holds the correct result
}

Reduce visualized (8 threads):

Initial: [1] [2] [3] [4] [5] [6] [7] [8]

Step 1 (offset=4): each thread adds the value 4 lanes away
       [1+5] [2+6] [3+7] [4+8] [5] [6] [7] [8]
       = [6]  [8]  [10]  [12]

Step 2 (offset=2): each thread adds the value 2 lanes away
       [6+10] [8+12] [10] [12] ...
       = [16]  [20]

Step 3 (offset=1): each thread adds the value 1 lane away
       [16+20] [20] ...
       = [36]    ← lane 0 holds the sum of all elements

---

5. Warp-level primitives and advanced tricks

5.1 Warp vote functions

Warp vote lets the whole warp “vote” on a boolean:

// true if every thread's predicate is true
int __all_sync(unsigned mask, int predicate);

// true if any thread's predicate is true
int __any_sync(unsigned mask, int predicate);

// 32-bit bitmask; each bit is one lane's predicate
unsigned __ballot_sync(unsigned mask, int predicate);

Useful pattern: early exit

__global__ void search_kernel(int* data, int target, int* found) {
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    int pred = (data[tid] == target);

    // If no thread in the warp found it, the whole warp can skip follow-up work
    if (__any_sync(0xFFFFFFFF, pred)) {
        // At least one thread found the target
        if (pred) {
            atomicExch(found, tid);
        }
    }
}

5.2 __activemask() and __syncwarp()

// Mask of currently active threads (mainly meaningful on Volta+)
unsigned mask = __activemask();

// Warp barrier (ensure named threads finished prior work)
__syncwarp(0xFFFFFFFF);

Tip: On pre-Volta architectures __syncwarp() is usually a no-op (warps were already lockstep). From Volta on, with independent scheduling, explicit warp sync becomes necessary.

5.3 Warp match functions (Volta+)

// Mask of threads that hold the same val
unsigned __match_any_sync(unsigned mask, T val);

// If every participating thread has the same val, return the full mask; else 0
unsigned __match_all_sync(unsigned mask, T val, int* pred);

Use case: auto-group threads with the same key for warp-level group-by.

5.4 Using the width parameter

The shuffle width argument can split a warp into logical sub-warps:

// Treat a 32-thread warp as 4 independent 8-thread subgroups
// Reduce independently inside each subgroup
float sub_warp_sum(float val) {
    for (int offset = 4; offset > 0; offset >>= 1) {
        val += __shfl_down_sync(0xFFFFFFFF, val, offset, 8);
    }
    return val;
}
// Lanes 0, 8, 16, 24 each hold their subgroup sum

---

6. Case study: warp-level reduction

6.1 Full block-level reduce

Combine warp shuffle with shared memory for an efficient block sum:

__device__ float warp_reduce_sum(float val) {
    for (int offset = 16; offset > 0; offset >>= 1) {
        val += __shfl_down_sync(0xFFFFFFFF, val, offset);
    }
    return val;
}

__global__ void block_reduce_kernel(float* input, float* output, int N) {
    int tid = threadIdx.x + blockIdx.x * blockDim.x;

    // Each thread loads one element
    float val = (tid < N) ? input[tid] : 0.0f;

    // Step 1: reduce inside the warp
    val = warp_reduce_sum(val);

    // Step 2: each warp's lane 0 writes its result to shared memory
    __shared__ float warp_sums[32];  // at most 32 warps (1024 threads / 32)
    int laneId = threadIdx.x % 32;
    int warpId = threadIdx.x / 32;

    if (laneId == 0) {
        warp_sums[warpId] = val;
    }
    __syncthreads();

    // Step 3: the first warp reduces warp_sums
    int numWarps = blockDim.x / 32;
    val = (threadIdx.x < numWarps) ? warp_sums[threadIdx.x] : 0.0f;

    if (warpId == 0) {
        val = warp_reduce_sum(val);
    }

    // Lane 0 writes this block's result
    if (threadIdx.x == 0) {
        output[blockIdx.x] = val;
    }
}

6.2 Performance comparison

Advantages of the warp-shuffle version:

• Far less shared memory (free it for other uses)
• Fewer barriers
• Register-to-register traffic with very low latency

---

Summary

---

Self-check list

• Explain why a warp is size 32 and how warps are formed
• State the core difference between SIMT and SIMD
• Spot warp divergence in code and estimate its cost
• Rewrite an if/else kernel to be branchless or warp-aligned
• Use all four warp shuffle ops (shfl / shfl_up / shfl_down / shfl_xor) correctly
• Implement warp-level reduce (sum/max/min) with warp shuffle
• Explain width and implement a sub-warp reduce
• Combine shuffle and shared memory for a block-level reduce
• Use __ballot_sync for warp-level filtering
• Explain how Volta independent thread scheduling affects programming

---

References

• NVIDIA CUDA C++ Programming Guide - Warp Shuffle Functions
• NVIDIA CUDA C++ Programming Guide - SIMT Architecture
• NVIDIA Volta Architecture Whitepaper
• Using CUDA Warp-Level Primitives - NVIDIA Developer Blog
• Faster Parallel Reductions on Kepler - NVIDIA Developer Blog