Through the Pascal generation, a warp had a single program counter and a single call stack, shared by all 32 of its threads. The hardware tracked an “active mask” of which lanes were enabled, executed each side of a divergent branch in turn while the other side sat idle, and reconverged the warp implicitly when the divergent paths rejoined. Lockstep execution was an architectural fact, and idioms such as the warp-synchronous reduction leaned on it.
The Volta architecture (2017) replaced this model with independent thread scheduling: every thread has its own program counter and call stack, and the scheduler may interleave the divergent paths of a warp at instruction granularity, guaranteeing forward progress for every thread in the warp. Algorithms that were previously impossible become possible – one lane may hold a lock while another lane of the same warp spins waiting for it, a pattern that deadlocked on pre-Volta hardware because the spinning side of the branch never yielded to the side that would release the lock. The warp is still the unit of execution – converged threads issue instructions together, and the hardware works to keep threads converged for efficiency – but convergence became an optimization rather than a guarantee.
The corollary is that code may no longer assume lanes
execute in lockstep. Reading a value that a neighboring lane just wrote,
safe by construction on pre-Volta hardware, now requires explicit
synchronization. CUDA 9 accordingly added __syncwarp() and
_sync-suffixed forms of every warp intrinsic – __ballot_sync(),
__shfl_sync(), __any_sync(), and kin – each taking a mask of the lanes
participating in the operation. The unsuffixed forms were deprecated in
CUDA 9 and have since been removed. Warp-synchronous idioms remain
valid, but only when written in terms of these explicit primitives.
A common but incorrect shortcut is to reach for volatile
instead. Applied to a shared-memory pointer, the qualifier instructs the
compiler not to elide the memory traffic that an access implies: because
a read through a volatile pointer is deemed to have
hardware side effects, the compiler may not substitute a cached,
register-held copy of the value, and must issue the load. (The same
property makes volatile useful in benchmarking, to keep the
compiler from optimizing away work whose result would otherwise go
unused.) But the optimizations volatile suppresses are
largely orthogonal to the ones that must be constrained for a
lane-to-lane handoff through shared memory to remain correct. In
particular, volatile says nothing about where the
compiler may not reorder instructions – it is not a memory barrier
– so the compiler remains free to move operations across the very points
at which the code assumes memory has settled. Only a barrier conveys
those boundaries: __syncthreads() for a thread block, __syncwarp() for a
warp. Warp-synchronous code that omits __syncwarp() is incorrect under
Independent Thread Scheduling whether or not its shared-memory pointers
are declared volatile.