3.4.3–3.4.4Research: MIMD Capability & Complexity-Effective Divergence

3.4.3 Adding MIMD Capability

The boldest response to branch divergencebranch divergenceThreads of one warp taking different targets at a data-dependent branch, forcing serialized execution of each path.Glossary →: build hardware that can stop being SIMD when divergence strikes. Every proposal in this family runs two modes — SIMD when the warp is coherent (one instruction, all threads, shared fetch: peak efficiency) and MIMD when it diverges (per-thread or per-slice instruction supply: lower efficiency, but far below a pure-SIMD divergence penalty).

• Vector-Thread (VT) architecture (Krashinsky et al.): lanes normally march in lockstep, fed from a shared L1 I-cache; on divergence, each lane detaches and runs at its own pace from a private L0 cache. Lee et al.’s comparison: parity with SIMT on regular code, much better on irregular parallelism.
• Temporal SIMT (Keckler, Krashinsky): rotate the warp 90° — instead of 32 threads across 32 lanes in one cycle, stream the warp through one lane over 32 cycles, one instruction fetch per warp. Each lane hosts its own warps and may drop into MIMD independently — control overhead amortized over time rather than space.
• Variable Warp Sizing (VWS, Rogers et al.): a middle path — eight 4-wide slices, each with its own fetch/decode and 256-byte L0, whose narrow warps are ganged into wide execution entities. Coherent code runs ganged (one fetch); divergence splits the gang; matching PCs re-form it:

Even for divergent workloads, those tiny L0s filter most instruction fetches, so the shared L1 needs only ~2× baseline bandwidth.

• Simultaneous Branch Interweaving (SBI, Brunie et al.) belongs here too: issue two different instructions per cycle to fill divergence holes (also in §3.4.1).

3.4.4 Complexity-Effective Divergence Management

The opposite instinct: divergence is rare in typical GPU code, so spend less hardware on it, not more. The baseline stack costs 32×64 bits per warp (down to 6×64 with tricks) — small, but it scales with every in-flight warp and sits idle in coherent kernels.

• The SIMT stack in scalar registers (AMD GCN): no divergence hardware at all — the compiler emulates the stack using the scalar register file and exec-mask arithmetic (you watched exactly this in §2.2.2), with special instructions to accelerate it. The elegant bound: always run the less-populated target first, and worst-case divergence needs only log₂(threads/wavefront) scalar registers — 6 for a 64-thread wavefront. Kernels with no divergent branches reclaim those registers for real scalar work.
• Thread frontiersthread frontiersPer-thread PCs + topologically sorted code; always run the lowest-PC threads so laggards catch up — no stack needed.Glossary → (Diamos et al., from §3.4.2) double as a cheap design: per-thread PCs live in the register file, allocated only in kernels that can actually diverge — everywhere else the storage boosts occupancy instead.
• Stackless SIMT with syncwarp (Asanovic et al.): on a temporal-SIMT substrate, let diverged threads run on their private PCs and have the compiler place a syncwarp at each divergent branch’s reconvergence point. No nested-reconvergence capture — but comparable performance when divergence is rare, at near-zero hardware cost. The enabling compiler machinery is a combined convergent/variant analysis, worth stepping through because §3.5 reuses it for scalarization:

idx  = threadIdx.x;      // invariant? (for now)
v    = y[idx] + a;       // invariant? (for now)
if (v > THRESHOLD)       // uniform branch? (for now)
    y[idx] = Y_MAX;      // invariant? (for now)

• Thread-aware predication (Lee et al.): skip divergence hardware entirely — extend the control-dependence graph with predicate nodes and compile arbitrary control flow to predicationpredicationExecuting an instruction conditionally under a predicate register (@%p1 in PTX, @P0 in SASS) instead of branching.Glossary →, with two rescue optimizations: static branch-uniformity (provably uniform branch → plain uniform branch instruction) and runtime branch-uniformity via consensual branches (cbranch.ifnone skips code whose predicate is null — restoring the stack’s whole-path skipping). Honest trade-offs: predicated code lets the compiler interleave both targets (free TLP!), but raises register pressure (occupancy ↓) and can inflate dynamic instruction count; find_unique (serialize indirect calls) and cbranch.ifany round out the proposed ISA support.