The CUDA Handbook ▸ Chapter 8. Streaming Multiprocessors

8.6 Miscellaneous Instructions

8.6.1 Warp-Level Primitives

It did not take long for the importance of warps as a primitive unit of execution (naturally residing between threads and threadblocks) to become evident to CUDA programmers. Starting with SM 1.x, NVIDIA began adding instructions that specifically operate on warps.

For most of CUDA’s history, these primitives operated implicitly on whichever threads of the warp happened to be executing together: __any(), __ballot(), __shfl() and the rest took no argument describing which lanes were participating. Volta’s independent thread scheduling (Section 7.5) invalidated that assumption – the hardware no longer guarantees that the lanes of a warp advance in lockstep, so “whichever threads happen to be converged” ceases to be well defined. CUDA 9 therefore introduced a _sync variant of every warp primitive that takes an explicit 32-bit mask naming the lanes expected to participate, and deprecated the older implicit forms. All the primitives below are shown in their _sync form; the pre-Volta names survive for source compatibility but should not be used in new code.

The Participation Mask

The mask is a 32-bit bitmask, one bit per lane: bit i stands for lane i of the warp, where a thread’s lane is its index within the warp (threadIdx.x % 32 in a one-dimensional block). A set bit declares that the corresponding lane participates in the collective. Every participating thread passes the mask as the first argument, and each must set its own lane’s bit – the set of threads that actually call must be exactly the set the mask names.

The mask is more than a hint. When the named lanes are not already converged, the primitive first synchronizes them: it waits until every lane named in the mask has reached this call executing the same primitive with the same mask, then performs the operation and releases them. That is what makes the exchange well defined under independent thread scheduling – the hardware no longer assumes lockstep, so each collective re-establishes it explicitly over the lanes it was told to expect. Two consequences follow, and both are easy to get wrong:

• Name exactly the lanes that will arrive. If the mask names a lane that has branched away or already exited, that lane can never reach the call, and the lanes waiting on it hang (formally, the result is undefined). Conversely, if two threads that should cooperate pass different masks, or a participating thread omits its own bit, the result is undefined.

• You can read only from a lane you named. For the data-exchange primitives – __shfl_*_sync() and the match primitives – a thread may read only from a lane that is itself in the mask. Reading from a lane that is not participating returns an undefined value, not zero.

Where does the mask come from? If the whole warp is active – no divergence, and the block is a multiple of the warp size – pass 0xffffffff (often spelled FULL_MASK), naming all 32 lanes. Inside divergent control flow, where only some lanes are live, the correct mask is dictated by the program’s logic. When it cannot be written down at compile time, __activemask() returns the set of lanes currently converged with the caller – but __activemask() and the mask argument answer different questions, and conflating them is a classic bug. __activemask() observes which lanes happen to be together at this instant; the mask you pass asserts which lanes must be. If your algorithm requires a specific set of lanes to cooperate, compute that set from the algorithm; do not ask the hardware which lanes showed up and then trust the answer.

Vote

The VOTE instruction (first available in SM 1.2) evaluates a condition across a warp and broadcasts a one-bit summary to every participating lane.

• __all_sync(mask, predicate) returns nonzero if predicate is nonzero for all lanes named in mask.

• __any_sync(mask, predicate) returns nonzero if predicate is nonzero for any named lane.

• __ballot_sync(mask, predicate) returns a 32-bit value whose ith bit is the (Boolean) predicate of lane i for each lane named in mask; bits for lanes outside the mask are zero. The Fermi-era VOTE.BALLOT it wraps was the first variant to pass back per-lane results rather than a single aggregate.

Shuffle

Kepler added shuffle instructions (SHFL) that exchange data directly between the lanes of a warp without staging it through shared memory. They execute with latency comparable to shared memory but perform the exchange in a single operation – no separate store and load – and free the shared memory they would otherwise consume. The four exchange patterns, defined in sm_30_intrinsics.h and overloaded for 32- and 64-bit integer and floating-point types, are:

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

The width parameter, which defaults to the warp size of 32, must be a power of two in the range 2..32; a smaller value subdivides the warp into independent segments, each behaving as a warp of that width with its own logical lane 0. A thread may exchange data only within its own segment.

• __shfl_sync() returns the value of var held by lane srcLane. If srcLane is outside the range 0..width-1, it is reduced modulo width. Passing the same srcLane to every lane broadcasts one lane’s value across the (sub)warp.

• __shfl_up_sync() reads from the lane delta below the caller (laneID - delta); lanes near the bottom, whose source would be negative, return their own var.

• __shfl_down_sync() reads from the lane delta above the caller (laneID + delta), with the analogous behavior at the top. Together, __shfl_up_sync() and __shfl_down_sync() are the building blocks of warp-level scans.

• __shfl_xor_sync() reads from lane laneID ^ laneMask. Because XOR with a fixed mask is its own inverse, every lane simultaneously reads and is read from, which makes it ideal for butterfly reductions: sweeping laneMask through 16, 8, 4, 2, 1 leaves every lane holding the warp-wide reduction of an associative operator in log2(width) steps.

Match

Volta (SM 7.0) added the match primitives (MATCH), which find the lanes of a warp that hold a given value:

unsigned __match_any_sync(unsigned mask, T value);
unsigned __match_all_sync(unsigned mask, T value, int *pred);

__match_any_sync() returns the set of lanes among mask whose value equals the caller’s – a warp-wide “group by value” in one instruction, useful for resolving which lanes address the same histogram bin or for building conflict-free scatters. __match_all_sync() returns mask and sets *pred to true if every named lane holds the same value, and returns 0 otherwise. Both are overloaded for 32- and 64-bit types.

Reduce

A butterfly __shfl_xor_sync() reduction over a full warp takes log2(32) = 5 steps. Ampere (SM 8.0) folds an integer warp reduction into a single instruction, REDUX, exposed as:

unsigned __reduce_add_sync(unsigned mask, unsigned value);
// __reduce_min_sync, __reduce_max_sync (signed & unsigned)
// __reduce_and_sync, __reduce_or_sync, __reduce_xor_sync

Each returns the reduction of value over the lanes named in mask. The arithmetic reductions (add, min, max) come in signed and unsigned overloads; the bitwise ones (and, or, xor) operate on the raw bits. These dispatch the common case of an integer warp reduction without a shuffle-and-accumulate loop; for other types or operators, the __shfl_xor_sync() butterfly remains the general tool.

Warp Synchronization

Just as __syncthreads() is a barrier for a threadblock, __syncwarp(mask = 0xffffffff) is a barrier for a warp: it forces the named lanes to reconverge and acts as a memory fence among them, ordering their accesses to shared and global memory. Under independent thread scheduling, it is the supported way to re-establish, at a chosen point, the lockstep execution that pre-Volta code assumed everywhere – for instance between the write and read phases of a shared-memory exchange within a warp. Unlike the _sync data primitives, __syncwarp() moves no data; it only synchronizes.

8.6.2 Block-Level Primitives

The __syncthreads() intrinsic serves as a barrier: it causes all threads to wait until every thread in the threadblock has arrived at the __syncthreads().

The Fermi instruction set (SM 2.x) added several new block-level barriers that aggregate information about the threads in the threadblock:

• __syncthreads_count(): evaluates a predicate and returns the sum of threads for which the predicate was true;

• __syncthreads_or(): returns the OR of all the inputs across the threadblock;

• __syncthreads_and(): returns the AND of all the inputs across the threadblock.

8.6.3 Performance Counter

Developers can define their own set of performance counters, and increment them in live code with the __prof_trigger() intrinsic:

void __prof_trigger(int counter);

Calling this function increments the corresponding counter by 1 per warp. counter must be in the range 0..7; counters 8..15 are reserved. The value of the counters may be obtained by listing prof_trigger_00..prof_trigger_07 in the profiler configuration file.

8.6.4 Video Instructions

The video instructions described in this section are accessible only via the inline PTX assembler. Their basic functionality is described here, to enable developers to decide whether they might be beneficial for their application. Anyone intending to use these instructions, however, should consult the PTX ISA specification.

Scalar Video Instructions

The scalar video instructions, added with SM 2.0 hardware¹⁷, enable efficient operations on the short (8- and 16-bit) integer types needed for video processing. As described in the PTX 3.1 ISA Specification, the format of these instructions is as follows:

vop.dtype.atype.btype{.sat} d, a{.asel}, b{.bsel};

vop.dtype.atype.btype{.sat}.secop d, a{.asel}, b{.bsel}, c;

The source and destination operands are all 32-bit registers. dtype, atype, and btype may be .u32 or .s32 for unsigned and signed 32-bit integers, respectively. The asel/bsel specifiers select which 8- or 16-bit value to extract from the source operands: b0, b1, b2 and b3 select bytes (numbering from the least significant) and h0/h1 select the least significant and most significant 16 bits, respectively.

Once the input values are extracted, they are sign- or zero- extended internally to signed 33-bit integers and the primary operation is performed, producing a 34-bit intermediate result whose sign depends on dtype. Finally, the result is clamped to the output range and one of the following operations performed:

1. apply a second operation (add, min or max) to the intermediate result and a third operand, or

2. truncate the intermediate result to an 8- or 16-bit value and merge into a specified position in the third operand to produce the final result.

The lower 32-bits are then written to the destination operand.

The vset instruction performs a comparison between the 8-, 16- or 32-bit input operands and generates the corresponding predicate (1 or 0) as output.

The PTX scalar video instructions, and the corresponding operations, are given in Table 8-16.

Table 8-16. Scalar Video Instructions.

Vector (SIMD-Within-a-Register) Video Instructions

These instructions pack two 16-bit values or four 8-bit values into a 32-bit register and operate on all of them at once: like the scalar video instructions, they promote the inputs to a canonical integer format, perform the core operation, then clamp and optionally merge the output. Dedicated hardware for them arrived with SM 3.0 (Kepler). On later architectures, the packed operations are emulated as short sequences of ordinary instructions, so the intrinsics remain available on every GPU – just without the single-instruction implementation Kepler gave them.

Table 8-17 summarizes the PTX instructions and corresponding operations implemented by these instructions. They are most useful for video processing and certain image processing operations (such as the median filter).

Table 8-17. Vector Video Instructions

Each mnemonic has a C intrinsic form built from a __v prefix and a 2 or 4 suffix – __vadd2/__vadd4, __vsub2/__vsub4, __vavgu2/__vavgu4, __vabsdiff2/__vabsdiff4, and the __vcmp* comparisons – declared in the SIMD-within-a-register intrinsics header. The most useful member is __vsadu4(), which sums the absolute differences of four unsigned byte pairs in a single call; it is the inner operation of block-matching motion estimation and of image-similarity metrics, exactly the work the packed byte formats were meant for.

8.6.5 Special Registers

Many special registers are accessed by referencing the built-in variables threadIdx, blockIdx, blockDim, and gridDim. These are 3-dimensional pseudo-structures that specify the thread ID, block ID, thread count, and block count, respectively.

Besides those, another special register is the SM’s clock register, which increments with each clock cycle. This counter can be read with the __clock() or __clock64() intrinsic. The counters are separately tracked for each SM and, like the time stamp counters on CPUs, are most useful for measuring relative performance of different code sequences and best avoided when trying to calculate wall clock times.