3.5Research: Scalarization and Affine Execution

The SIMT model launches thousands of “independent” scalar threads — but in the common case they march down the same paths doing nearly the same thing. Beyond grouping them into warps, research exploits their value structure (Kim et al.): a uniformuniform variableSame value for every thread in a kernel; storable in one scalar register.Glossary → variable holds the same value in every thread (the parameter a, literals like THRESHOLD); an affineaffine variableValue that is base + stride·threadID; storable as a (base, stride) pair.Glossary → variable is a linear function of thread ID — most famously the address &y[idx] = &y[0] + 4·threadIdx.x, storable as one (base, stride) pair instead of a 32-wide vector. ScalarizationscalarizationExecuting uniform/affine work once on a scalar path (or one gated lane) instead of per-lane.Glossary → executes such work once instead of per-lane.

3.5.1 Detecting uniform and affine variables

Compiler-driven. CUDA/OpenCL already tell the compiler which values are constants and which register holds the thread ID; control-dependency analysis propagates from there. AMD GCN relies entirely on this — its scalar unit and scalar register file (§2.2.2) execute compiler-identified scalar work. Asanovic et al.’s convergent/variant analysis (stepped through in §3.4.4) marks whole regions eligible for scalarization; Decoupled Affine Computation (DAC, Wang & Lin) adds a divergent affine analysis to rip entire affine strands out of the kernel into a separate warp, fed to the main kernel through a hardware queue.

Hardware detection works with unmodified ISAs and catches what static analysis misses. The classic design is tag-based (Collange et al.): every register carries a U/A/V tag, seeded at launch and propagated by simple rules — deliberately conservative ones:

Watch the tags flow through the book’s running example — including the two subtleties (addresses vs. data, and what divergence does to affine registers):

__global__ void vsadd(int y[], int a) {
  int idx = threadIdx.x;      // ← A (stride 1)
  y[idx] = y[idx] + a;        // a ← U
  if (y[idx] > THRESHOLD)     // THRESHOLD ← U
    y[idx] = Y_MAX_VALUE;     // Y_MAX_VALUE ← U
}

More aggressive detectors move to the write-back stage: Gilani et al. compare all lanes’ values on every vector write-back and reroute uniform ones to a scalar register file; Lee et al. compress incoming vectors as ⟨base, delta, immediate⟩ (BDI); Warp Approximation (Wong et al.) measures d-similarity (shared most-significant bits) and, above a threshold, runs approximate scalar execution on a representative lane; G-Scalar (Liu et al.) uses a simpler common-bytes compressor — and uniquely keeps scalarizing under divergence, by checking only the active lanes’ operand values.

3.5.2 Exploiting them

• Compressed register storage. Scalar/affine register files (GCN; FG-SIMT’s ASRF; Gilani et al.) hold (base, stride) pairs; the reclaimed vector-RF space sustains more in-flight warps. BDI-compressed registers (Lee et al.) still occupy a full slot but touch only a subset of banks — less read energy; Warp Approximation reads just the representative lane; DAC’s affine warp packs multiple warps’ bases and deltas into the lanes of shared vector registers.
• Scalarized operations. Execute once: on GCN’s dedicated scalar pipeline or FG-SIMT’s control processor (which also absorb control flow, predication, and host communication — freeing the SIMD path from implementing the full ISA), or on one clock-gated lane of the existing datapath (Warp Approximation, G-Scalar — no duplicated function units, and G-Scalar even scalarizes SFU instructions cheaply). DAC aggregates affine ops from many warps into its single affine warp.
• Memory access acceleration. An affine address means a known stride: coalescing becomes simple hardware (FG-SIMT’s address-generation unit expands base+stride directly), bulk transfers replace load loops — and DAC runs its affine warp ahead of the others as a prefetcher, parking data in the L1 for a later dequeue.