Wavefront-level builtins for AMD GPUs

Wavefront-level operations give you direct access to the hardware mechanisms that move and aggregate data within a wavefront. The builtins in this topic work across both CDNA (AMD Instinct) and RDNA (AMD Radeon) architectures, and cover three areas: lane operations, wavefront reductions, and wavefront voting.

Architecture availability

Most wavefront-level builtins are available on all AMD Instinct (CDNA) and AMD Radeon (RDNA) architectures. Some builtins — particularly Data Parallel Primitives (DPP) DPP8, cross-row permutations, and ballot width variants — are limited to specific generations. For per-builtin availability, see the architecture tables in the reference pages below.

The complete source file is available for download:

• wavefront_builtins.hip

Lane operations

Within a wavefront, lanes execute the same instruction simultaneously, but each holds its own register values. Many algorithms require lanes to exchange or replicate those values - to share a computed result, apply a cyclic shift, or reorganize data before the next computation step. The builtins in this topic cover the distinct ways of expressing that communication: reading from a named lane, moving data in a cyclic pattern, applying a compile-time-fixed permutation, and broadcasting within a fixed-size group. Each fills a different niche, and choosing the right one affects both correctness and performance.

Broadcasting from a specific lane with readlane

When a single lane holds a value that every other lane needs - a wavefront-wide maximum, a shared configuration parameter, a count computed by one thread - the most direct way to share it is __builtin_amdgcn_readlane. Unlike a shuffle, which requires every lane to participate with a relative offset, readlane reads from one named lane regardless of who is calling, making it the natural choice when the source is fixed rather than relative. The example here uses it to broadcast a wavefront maximum after a __shfl_down reduction, which is a representative case: one lane holds the answer, and all others need it.

__builtin_amdgcn_readlane(val, lane) returns the value of val held by the lane specified by lane. Unlike readfirstlane, which always reads the first active lane, readlane accepts a runtime lane index. The lane index must be uniform — the same value across all lanes in the wavefront at the point of the call. If the index is derived from per-lane data, use readfirstlane to promote it to a uniform value first. readfirstlane reads from the first active lane of the wavefront, which is a deterministic choice, so every lane in the wavefront receives the same index value and the uniformity requirement is satisfied.

__global__ void broadcast_readlane(const int *input, int *output, int n)
{
    int gid  = blockIdx.x * blockDim.x + threadIdx.x;
    int lane = threadIdx.x % warpSize;
    int val  = (gid < n) ? input[gid] : 0;

    // Find the warp maximum via __shfl_down reduction.
    int max_val = val;
    for (int offset = warpSize / 2; offset > 0; offset >>= 1)
    {
        max_val = max(max_val, __shfl_down(max_val, offset));
    }

    // Broadcast the maximum from lane 0 to all lanes.
    // The lane index must be uniform; use readfirstlane to promote a
    // runtime-derived index to a uniform value.
    int broadcast = __builtin_amdgcn_readlane(max_val, 0);

    if (gid < n)
    {
        output[gid] = broadcast;
    }

    (void)lane;
}

After a __shfl_down reduction, lane 0 holds the wavefront maximum. Passing the literal 0 as the lane index satisfies the uniformity requirement and broadcasts that value to every lane in the wavefront.

For the full signatures and parameter details, see readlane and readfirstlane in the shuffle and lane access reference.

Wavefront rotation using mov_dpp and ds_bpermute

Rotation shifts every lane’s value one position around the wavefront, so lane i receives what lane i - 1 held, and lane 0 wraps around to receive what the last lane held. This cyclic movement is a building block for algorithms that scan or pipeline values across lanes, such as prefix operations or producer-consumer patterns within a wavefront. Data Parallel Primitives (DPP) is the hardware-native way to express rotation on AMD GPUs, executing in a dedicated unit without consuming the general instruction pipeline. On CDNA-based GPUs, a single wave_ror:1 instruction covers the entire wave64 atomically. On RDNA-based GPUs, where wave_ror is not available on gfx11xx and later, the same result requires two steps: a row_ror:1 rotates within each 16-lane row, and ds_bpermute supplies the wrap-around value at the boundary of each row.

// wave_ror:1 (CDNA): rotates the entire warp right by one lane.
// row_ror:1 (RDNA): rotates within each row of 16; ds_bpermute fixes wrap-around.
#ifdef HAS_DPP_BROADCAST
constexpr int dpp_wave_ror1 = 0x134;
#else
constexpr int dpp_row_ror1  = 0x121;
#endif
__device__ int warp_rotate_right(int val)
{
#ifdef HAS_DPP_BROADCAST
    // CDNA: single wave_ror:1 rotates all lanes atomically.
    return __builtin_amdgcn_mov_dpp(val, dpp_wave_ror1, 0xf, 0xf, false);
#else
    // RDNA: row_ror:1 rotates within each row of 16.
    int rotated = __builtin_amdgcn_mov_dpp(val, dpp_row_ror1, 0xf, 0xf, false);

    // ds_bpermute (byte addressing: lane k -> addr k*4) fixes up the
    // wrap-around value at lane 0 of each row.
    int lane    = threadIdx.x % warpSize;
    int addr    = (lane == 0) ? 124 : (lane == 16) ? 60 : lane * 4;
    int wrapped = __builtin_amdgcn_ds_bpermute(addr, val);

    return (lane == 0 || lane == 16) ? wrapped : rotated;
#endif
}

__global__ void rotate_warp(const int *input, int *output, int n)
{
    int gid = blockIdx.x * blockDim.x + threadIdx.x;
    int val = (gid < n) ? input[gid] : 0;

    // Rotate right warpSize times, XORing each step with the step index.
    for (int i = 0; i < warpSize; ++i)
    {
        val = warp_rotate_right(val);
        val ^= i;
    }

    if (gid < n)
    {
        output[gid] = val;
    }
}

ds_bpermute uses byte addressing: to read from lane k, pass k * 4 as the index. Lane 0 of the first row reads from lane 31 (address 124), and lane 0 of the second row reads from lane 15 (address 60). All other lanes keep the row_ror result unchanged.

For the full signatures and parameter details, see mov_dpp and ds_bpermute in the DPP and data-share permutation reference.

Wavefront rotation using __shfl

The DPP rotation above is the preferred approach on AMD hardware, but the same pattern can be expressed using __shfl. This is useful when the performance difference is not significant for your workload. Each lane computes its source index as (lane - 1 + warpSize) % warpSize and because __shfl accepts any uniform lane index, the modular arithmetic produces the wrap-around without the architecture-specific workaround that the DPP path requires.

__device__ int warp_rotate_right_shfl(int val)
{
    int lane = threadIdx.x % warpSize;
    // Rotate right: each lane reads from (lane - 1 + warpSize) % warpSize.
    return __shfl(val, (lane - 1 + warpSize) % warpSize);
}

__global__ void rotate_warp_shfl(const int *input, int *output, int n)
{
    int gid = blockIdx.x * blockDim.x + threadIdx.x;
    int val = (gid < n) ? input[gid] : 0;

    // Rotate right warpSize times, XORing each step with the step index.
    // Produces the same result as rotate_warp.
    for (int i = 0; i < warpSize; ++i)
    {
        val = warp_rotate_right_shfl(val);
        val ^= i;
    }

    if (gid < n)
    {
        output[gid] = val;
    }
}

The DPP and __shfl kernels produce identical output and can be verified against the same CPU reference.

The two paths compile to different Instruction Set Architecture (ISA) on CDNA-based GPUs: the DPP path emits v_mov_b32_dpp wave_rol:1, a pure register move that executes in the dedicated DPP unit without touching memory, while the __shfl path emits ds_bpermute_b32, which routes through the data-share unit. On RDNA-based GPUs, where wave_ror is unavailable, the DPP path uses v_mov_b32_dpp row_ror:1 for the within-row lanes and ds_bpermute_b32 for the cross-row fixup; the __shfl path uses only ds_bpermute_b32. The ISA output for both can be inspected by compiling with -save-temps and examining the generated .s file.

Lane swap using ds_swizzle

While readlane and rotation address individual lanes or shift the whole wavefront by one, some algorithms need a fixed, symmetric rearrangement of groups - for example, swapping pairs of lanes, or exchanging two halves of a tile (a fixed-size contiguous sub-group of lanes within a wavefront). __builtin_amdgcn_ds_swizzle is designed exactly for this: the permutation is encoded entirely in a compile-time mask, so that the hardware can apply it in a single instruction with no per-lane index computation. The mask specifies a bitwise transformation of each lane index to its source, making it well-suited to power-of-two group swaps that appear in butterfly networks, data rearrangement before matrix operations, or paired-lane exchanges. For the full signature, see ds_swizzle in the DPP and data-share permutation reference.

The 16-bit mask encodes the permutation: bits [14:10] hold and_mask, bits [9:5] hold or_mask, and bits [4:0] hold xor_mask. The source lane is computed as (lane & and_mask) | or_mask ^ xor_mask. Setting and_mask = 0x1F and or_mask = 0, and placing a single set bit in xor_mask, swaps neighboring groups whose size is determined by the position of that bit. The example demonstrates the swap pattern; the mask value and how it determines group size are explained in the prose below.

// ds_swizzle mask layout: bits [14:10]=and_mask, [9:5]=or_mask, [4:0]=xor_mask.
// Source lane = (lane & and_mask) | or_mask ^ xor_mask.
// 0x101F: and_mask=0x1F, xor_mask=0x04 - swaps neighboring groups of 4.
constexpr int swizzle_swap_groups4 = 0x101F;
__global__ void swizzle_swap_groups(const int *input, int *output, int n)
{
    int gid = blockIdx.x * blockDim.x + threadIdx.x;
    int val = (gid < n) ? input[gid] : 0;

    int swapped = __builtin_amdgcn_ds_swizzle(val, swizzle_swap_groups4);

    if (gid < n)
    {
        output[gid] = swapped;
    }
}

The mask 0x101F sets xor_mask = 0x04, which applies a bitwise XOR between each lane index and 4. Lanes 0–3 read from lanes 4–7 and vice versa, lanes 8–11 read from lanes 12–15, and so on across the wavefront. Shifting the set bit changes the group size: 0x041F swaps neighboring pairs, 0x081F swaps groups of two, 0x201F swaps groups of eight, and 0x401F swaps groups of sixteen.

For the full signature and parameter details, see ds_swizzle in the DPP and data-share permutation reference.

Wavefront reductions

A wavefront reduction combines a value held by each lane into a single scalar result. This pattern appears constantly in GPU kernels - summing partial products, finding a maximum across a tile, counting active threads - and the efficiency of the reduction matters because it often sits in the critical path of a kernel. The three implementations below all express the same butterfly pattern, where lanes exchange values with increasingly distant partners until one lane holds the aggregate. Still, they differ significantly in how directly they map to hardware. The first two reduce floating-point values and show the progression from shuffle-based code to hardware-native DPP. The third uses an unsigned integer reduction to demonstrate wave_reduce_add_u32, the single-instruction form currently available in the compiler. Understanding all three lets you choose the right level of abstraction for your target and performance requirements.

Reduction using __shfl_down

__shfl_down is the most readable starting point for wavefront reductions. It reads from a lane a fixed offset ahead of the current lane, which maps directly onto the butterfly pattern when called in a loop that halves the offset on each step. The code structure closely mirrors the algorithm’s logical structure, making it straightforward to reason about and maintain. The kernel is templated on WarpSize, a compile-time constant passed at the call site, which allows the compiler to fully unroll the reduction loop and makes the code adapt to wave32 and wave64 without modification. This makes __shfl_down the right choice when clarity and simplicity matter, or as a reference implementation against which a more hardware-specific version can be verified.

template<int WarpSize>
__device__ float shfl_down_reduce(float val)
{
    // Halve the active offset each step until all lanes have contributed.
    // WarpSize is a compile-time constant, allowing the compiler to fully
    // unroll this loop.
    for (int offset = WarpSize / 2; offset > 0; offset >>= 1)
    {
        val += __shfl_down(val, offset);
    }
    return val;
}

template<int WarpSize>
__global__ __launch_bounds__(block_size)
void reduce_shfl_down(const float *input, float *output, int n)
{
    int gid    = blockIdx.x * blockDim.x + threadIdx.x;
    int lane   = threadIdx.x % WarpSize;
    int warpid = threadIdx.x / WarpSize;

    float val = (gid < n) ? input[gid] : 0.0f;
    val = shfl_down_reduce<WarpSize>(val);

    // Collect one partial result per warp into shared memory, then reduce
    // those with the first warp.
    extern __shared__ float sdata[];
    int num_warps = blockDim.x / WarpSize;
    if (lane == 0)
    {
        sdata[warpid] = val;
    }
    __syncthreads();

    if (warpid == 0)
    {
        val = (lane < num_warps) ? sdata[lane] : 0.0f;
        val = shfl_down_reduce<WarpSize>(val);
    }

    if (threadIdx.x == 0)
    {
        output[blockIdx.x] = val;
    }
}

After the loop, lane 0 holds the sum of all lanes in the wavefront. The kernel uses a two-phase structure: each wavefront reduces its slice independently, lane 0 writes its partial result to shared memory, and the first wavefront reduces those partials to the block result.

Reduction using mov_dpp

When targeting supported AMD GPUs specifically, Data Parallel Primitives offer a faster path. __builtin_amdgcn_mov_dpp moves data between lanes according to a DPP control word, executing in a dedicated hardware unit rather than the general instruction pipeline used by shuffle operations. The same butterfly pattern as above can be expressed with DPP control words, and the result is a reduction that makes better use of the hardware’s data movement capabilities. The trade-off is that DPP instructions expose architectural differences directly: broadcast instructions are available on gfx9xx targets (CDNA), but not on gfx10xx and later (RDNA, wave32), so the HAS_DPP_BROADCAST macro selects the appropriate path at compile time.

// DPP control words for butterfly reduction steps.
// quad_perm and row_ror steps work on both CDNA and RDNA.
// row_bcast steps are CDNA only; ds_swizzle mask 0x1e0 is the RDNA fallback.
constexpr int dpp_quad_perm_1032 = 0xb1;
constexpr int dpp_quad_perm_2301 = 0x4e;
constexpr int dpp_row_ror4       = 0x124;
constexpr int dpp_row_ror8       = 0x128;
#ifdef HAS_DPP_BROADCAST
constexpr int dpp_row_bcast15    = 0x142;
constexpr int dpp_row_bcast31    = 0x143;
#else
constexpr int swizzle_bcast15    = 0x1e0;
#endif
template<int WarpSize>
__device__ float dpp_warp_reduce(float val)
{
    // Reinterpret float bits as int for DPP and swizzle intrinsics.
    auto as_int   = [](float f) { return __builtin_bit_cast(int, f); };
    auto as_float = [](int i)   { return __builtin_bit_cast(float, i); };

    // Steps 1-2: combine adjacent pairs within groups of 4.
    val += as_float(__builtin_amdgcn_mov_dpp(as_int(val), dpp_quad_perm_1032, 0xf, 0xf, false));
    val += as_float(__builtin_amdgcn_mov_dpp(as_int(val), dpp_quad_perm_2301, 0xf, 0xf, false));
    // Steps 3-4: combine groups of 4 into a row of 16.
    val += as_float(__builtin_amdgcn_mov_dpp(as_int(val), dpp_row_ror4, 0xf, 0xf, false));
    val += as_float(__builtin_amdgcn_mov_dpp(as_int(val), dpp_row_ror8, 0xf, 0xf, false));

#ifdef HAS_DPP_BROADCAST
    // Steps 5-6 (CDNA): combine two rows of 16 into the full wave64 result.
    val += as_float(__builtin_amdgcn_mov_dpp(as_int(val), dpp_row_bcast15, 0xa, 0xf, false));
    val += as_float(__builtin_amdgcn_mov_dpp(as_int(val), dpp_row_bcast31, 0xc, 0xf, false));
#else
    // Step 5 (RDNA): ds_swizzle combines the two rows of the wave32 warp.
    val += as_float(__builtin_amdgcn_ds_swizzle(as_int(val), swizzle_bcast15));
#endif

    return val;
}

template<int WarpSize>
__global__ __launch_bounds__(block_size)
void reduce_dpp(const float *input, float *output, int n)
{
    int gid    = blockIdx.x * blockDim.x + threadIdx.x;
    int lane   = threadIdx.x % WarpSize;
    int warpid = threadIdx.x / WarpSize;

    float val = (gid < n) ? input[gid] : 0.0f;
    val = dpp_warp_reduce<WarpSize>(val);

    extern __shared__ float sdata[];
    int num_warps = blockDim.x / WarpSize;
    if (lane == WarpSize - 1)
    {
        sdata[warpid] = val;
    }
    __syncthreads();

    if (warpid == 0)
    {
        val = (lane < num_warps) ? sdata[lane] : 0.0f;
        val = dpp_warp_reduce<WarpSize>(val);
    }

    if (threadIdx.x == WarpSize - 1)
    {
        output[blockIdx.x] = val;
    }
}

The first two steps use quad_perm control words, which exchange lanes within groups of four. Steps three and four use row_ror rather than row_shr: rotation wraps lanes within the row boundary, whereas a shift leaves destination registers undefined when the source is out of range. On RDNA, where DPP broadcasts are unavailable, ds_swizzle with mask 0x1e0 replicates the value held by lane 15 across lanes 16–31, combining the two rows of the wave32 wavefront in a single instruction. The control word constants used in the example (0xb1, 0x4e, 0x124, and so on) are user-defined values derived from the DPP control word encoding; they are not predefined by HIP or LLVM. For the full DPP control word encoding and available patterns, refer to the Data Parallel Primitives chapter of the CDNA3 ISA or RDNA3 ISA.

Reduction using wave_reduce_add_u32

Both implementations above express the butterfly pattern explicitly in code, requiring you to select control words, handle the architecture split, and manage the reduction steps yourself. __builtin_amdgcn_wave_reduce_add_u32 encapsulates all of that: you supply the value and a strategy hint, and the builtin handles the reduction without you needing to implement or maintain the underlying pattern. This represents the most hardware-specific end of the spectrum - the implementation details are handled by the compiler and runtime rather than by your code. Note that the wave reduce builtins currently cover integer and bitwise operations; for float workloads, the DPP path above remains the recommended approach.

template<int WarpSize>
__global__ __launch_bounds__(block_size)
void reduce_wave(const unsigned int *input, unsigned int *output, int n)
{
    int gid    = blockIdx.x * blockDim.x + threadIdx.x;
    int lane   = threadIdx.x % WarpSize;
    int warpid = threadIdx.x / WarpSize;

    unsigned int val = (gid < n) ? input[gid] : 0u;

    // Reduces val across all active lanes in a single instruction.
    // The result is broadcast to every lane.
    unsigned int sum = __builtin_amdgcn_wave_reduce_add_u32(val, 0);

    extern __shared__ unsigned int sdataui[];
    int num_warps = blockDim.x / WarpSize;
    if (lane == 0)
    {
        sdataui[warpid] = sum;
    }
    __syncthreads();

    if (warpid == 0)
    {
        sum = (lane < num_warps) ? sdataui[lane] : 0u;
        sum = __builtin_amdgcn_wave_reduce_add_u32(sum, 0);
    }

    if (threadIdx.x == 0)
    {
        output[blockIdx.x] = sum;
    }
}

The second argument is a strategy hint: 0 lets the compiler choose, 1 requests the iterative strategy, and 2 requests the DPP-based strategy.

For the full signature and parameter details, see wave_reduce_add_u32 in the wavefront reduction reference.

Wavefront voting

Lane operations move data between lanes, and reductions aggregate it. Wavefront voting answers a different question: which lanes satisfy a condition, and what can the wavefront do with that information collectively? The ballot builtin captures the answer as a bitmask - one bit per lane - which can then be inspected, counted, or used to coordinate writes. The mbcnt builtin builds on this by counting the number of lanes before the current one that have their bit set, giving each qualifying lane a unique sequential index. Together, ballot and mbcnt is the standard mechanism for stream compaction within a wavefront: filtering a set of values to only those that pass a predicate, and writing them contiguously to an output buffer without gaps or collisions, all without shared memory or barriers.

Compaction index using ballot and mbcnt

__global__ void compaction_index(const int *input, int *output, int n)
{
    int gid    = blockIdx.x * blockDim.x + threadIdx.x;
    int lane   = threadIdx.x % warpSize;
    int warpid = threadIdx.x / warpSize;
    int val    = (gid < n) ? input[gid] : 0;

    bool keep  = (gid < n) && (val > 0);
    int  index = 0;

#ifdef HAS_DPP_BROADCAST
    // wave64 (CDNA): chain mbcnt_lo and mbcnt_hi over the 64-bit mask.
    unsigned long long mask = __builtin_amdgcn_ballot_w64(keep);
    unsigned int lo         = static_cast<unsigned int>(mask);
    unsigned int hi         = static_cast<unsigned int>(mask >> 32);
    index                   = __builtin_amdgcn_mbcnt_lo(lo, 0);
    index                   = __builtin_amdgcn_mbcnt_hi(hi, index);
#else
    // wave32 (RDNA): mbcnt_lo over the 32-bit mask suffices.
    unsigned int mask = __builtin_amdgcn_ballot_w32(keep);
    index             = __builtin_amdgcn_mbcnt_lo(mask, 0);
#endif

    // Compute per-warp base offsets for contiguous compacted output.
    extern __shared__ int sdatai[];
    int num_warps  = blockDim.x / warpSize;
    int warp_count = __builtin_amdgcn_readlane(index + (keep ? 1 : 0), warpSize - 1);
    if (lane == 0)
    {
        sdatai[warpid] = warp_count;
    }
    __syncthreads();

    int warp_base = 0;
    for (int w = 0; w < warpid; ++w)
    {
        warp_base += sdatai[w];
    }

    if (keep)
    {
        output[blockIdx.x * blockDim.x + warp_base + index] = val;
    }

    (void)num_warps;
}

mbcnt_hi takes the upper 32 bits of the mask and the result of mbcnt_lo as its accumulator, so the final index is the count of passing lanes strictly below the current lane across the full 64-bit mask. Lanes where the predicate is false receive an index, too, but should not write to the output.

For the full signatures and parameter details, see ballot_w64 and mbcnt_lo / mbcnt_hi in the wavefront voting reference.

Compile and run:

amdclang++ -O3 -std=c++17 --offload-arch=gfx942 \
    wavefront_builtins.hip -o wavefront_builtins
./wavefront_builtins

Note

The example above targets CDNA3 (gfx942, wave64). Replace --offload-arch with the appropriate target for your GPU — for example, gfx90a for CDNA2 or gfx1200 for RDNA4 (wave32). The __shfl* builtins work on all architectures; the DPP-specific code paths are guarded by architecture macros in the example file.

Naming convention

Wavefront-level builtins come in two families:

HIP wrappers use the __shfl prefix:

__shfl[_up|_down|_xor](val, offset_or_lane, width)

The suffix indicates the direction or mode: _up reads from a lower lane, _down from a higher lane, _xor reads from a lane whose index is the bitwise XOR of the current lane index and a mask, and no suffix reads from an absolute lane index.

Compiler builtins use the __builtin_amdgcn_ prefix and map directly to ISA instructions:

__builtin_amdgcn_<operation>(args...)

Key operation names:

• readlane / readfirstlane / writelane – named lane access

• mov_dpp / update_dpp / mov_dpp8 – Data Parallel Primitives

• ds_swizzle / ds_permute / ds_bpermute – data-share permutations

• permlane16 / permlanex16 / permlane64 – cross-row permutations

• ballot_w64 / ballot_w32 – wavefront voting

• mbcnt_lo / mbcnt_hi – masked bit count (compaction index)

• wave_reduce_<op>_<type> – single-instruction wavefront reductions, where <op> is add, sub, min, max, and, or, or xor, and <type> is u32, u64, i32, i64, b32, or b64

Use the __shfl* family for standard shuffle operations. Use the __builtin_amdgcn_* intrinsics when you need a specific hardware feature (DPP, ds_swizzle, wavefront voting) or when the ISA instruction gives measurable performance benefit.

Wavefront builtin reference

Each reference page documents the full signature, parameter details, and architecture support for every builtin in that family.

Family	Description
Shuffle and lane access	__shfl* wrappers and hardware readlane, readfirstlane, writelane
DPP and data-share permutations	mov_dpp, update_dpp, mov_dpp8, ds_swizzle, ds_permute, ds_bpermute
Cross-row permutations	permlane16, permlanex16, permlane64, and runtime/swap variants
Wavefront reductions	wave_reduce_<op>_<type> for add, sub, min, max, and, or, xor
Wavefront voting and synchronization	ballot, inverse_ballot, mbcnt, wave_barrier, wave_id