Arithmetic and packing builtins for AMD GPUs

Arithmetic and packing builtins give you direct access to dedicated units for the sum of absolute differences, integer dot products, FP16 (16-bit half-precision) dot products, and type conversion and packing. For a complete listing of all builtins with their signatures and supported architectures, see Arithmetic and packing builtin reference.

The complete source file is available for download:

• builtins_cross_arch_arithmetic.hip

Architecture availability

Most arithmetic and packing builtins are available on all AMD Instinct (CDNA) and AMD Radeon (RDNA) architectures. Some builtins — particularly mixed-sign dot products, FP8/BF8 conversions, and stochastic rounding — are limited to specific generations. For per-builtin availability, see the architecture tables in the reference pages below.

Sum of absolute differences

The sum of absolute differences (SAD) family computes the L1 distance between two byte vectors, producing a scalar accumulator that measures the similarity between the two blocks of data. The fundamental use case is block matching in video encoding and image processing. Each candidate reference block is compared against the current block by accumulating the per-byte absolute differences, and the candidate with the lowest SAD is the best match. The three builtins discussed below cover single 4-byte SAD, dual-SAD packing, and four overlapping 8-byte SADs, giving you progressively higher throughput as your data layout permits.

Masked SAD with msad_u8

The __builtin_amdgcn_msad_u8 intrinsic computes the SAD of four packed bytes, excluding byte positions where the reference byte (src1) is zero from the accumulation. This masking behavior is designed for padded reference windows, where the encoder marks positions outside the frame boundary with zero rather than valid pixel data. Accumulating those positions would artificially inflate the SAD and bias the match; skipping them keeps the comparison honest regardless of the padding layout.

// msad_u8: SAD of four packed bytes, skipping positions where the src1 byte is zero.
// Useful for block matching over padded reference windows where padding is coded as zero.
__global__ __launch_bounds__(block_size)
void msad_u8(const unsigned int *src0,
             const unsigned int *src1,
             unsigned int       *output,
             int                 n)
{
    int gid = blockIdx.x * blockDim.x + threadIdx.x;
    if (gid >= n)
    {
        return;
    }
    output[gid] = __builtin_amdgcn_msad_u8(src0[gid], src1[gid], 0u);
}

The third argument is the initial accumulator value. Passing 0u gives a fresh SAD for each element; passing the result of a previous call accumulates across multiple 32-bit values.

For the full signature and parameter details, see msad_u8 in the SAD builtins reference.

Dual-SAD packing with sad_u8 and sad_hi_u8

__builtin_amdgcn_sad_u8 computes the SAD of four packed bytes and accumulates the result into bits [15:0] of the destination 32-bit value. __builtin_amdgcn_sad_hi_u8 computes the same byte SAD but shifts each per-byte absolute difference left by 16 bits before accumulating, placing its result in bits [31:16]. Because the two results occupy different halves of the same 32-bit value, one call to each packs two independent 4-byte SADs into a single 32-bit value with no additional packing instructions, storing two SAD results in one register.

// sad_u8 accumulates into bits [15:0]; sad_hi_u8 shifts each byte difference
// left 16 before accumulating, placing the result in bits [31:16]. One call
// to each packs two independent 4-byte SADs into one register.
__global__ __launch_bounds__(block_size)
void sad_variants(const unsigned int *src0_lo,
                  const unsigned int *src0_hi,
                  const unsigned int *src1_lo,
                  const unsigned int *src1_hi,
                  unsigned int       *output,
                  int                 n)
{
    int gid = blockIdx.x * blockDim.x + threadIdx.x;
    if (gid >= n)
    {
        return;
    }
    unsigned int packed = __builtin_amdgcn_sad_u8(src0_lo[gid], src1_lo[gid], 0u);
    packed              = __builtin_amdgcn_sad_hi_u8(src0_hi[gid], src1_hi[gid], packed);
    output[gid]         = packed;
}

The low SAD accumulates from src0_lo against src1_lo; the high SAD accumulates from src0_hi against src1_hi into the same 32-bit output value. The maximum value of a 4-byte SAD is 4 × 255 = 1020, which fits in 10 bits, so neither half overflows into the other regardless of the input values.

For the full signatures and parameter details, see sad_u8 and sad_hi_u8 in the SAD builtins reference.

Overlapping SAD with qsad_pk_u16_u8

__builtin_amdgcn_qsad_pk_u16_u8 computes four overlapping 4-byte SADs from a single 8-byte reference window in one instruction. The 8-byte source (src0) is treated as a sliding window: the four SADs compare bytes [0:3], [1:4], [2:5], and [3:6] of the window against the same 4-byte query (src1). The four 16-bit results are packed into the returned 64-bit value, making this the natural tool for a full search over a scan line where the reference window advances one byte at a time.

// qsad_pk_u16_u8: four overlapping 8-byte SADs in one instruction.
// src0 is a 64-bit sliding reference window; src1 is the 4-byte query.
// The four results are packed as uint16 values in the returned uint64.
__global__ __launch_bounds__(block_size)
void qsad_pk_u16_u8(const unsigned long long *src0,
                    const unsigned int        *src1,
                    unsigned long long        *output,
                    int                        n)
{
    int gid = blockIdx.x * blockDim.x + threadIdx.x;
    if (gid >= n)
    {
        return;
    }
    output[gid] = __builtin_amdgcn_qsad_pk_u16_u8(src0[gid], src1[gid], 0ull);
}

Each of the four packed uint16 results can be extracted with a shift and mask: (result >> (i * 16)) & 0xffff gives the SAD for window offset i. The third argument is the initial accumulator for all four channels simultaneously.

For the full signature and parameter details, see qsad_pk_u16_u8 in the SAD builtins reference.

Integer dot products

The dot product builtins compute the inner product of two short integer vectors in a single instruction, accumulating the result into a wider integer. Using the builtin guarantees the corresponding hardware instruction is emitted regardless of compiler heuristics, and makes the intent explicit in the source code.

Unsigned 4-element dot product with udot4

__builtin_amdgcn_udot4 computes the dot product of two vectors of four unsigned bytes, accumulating into a 32-bit unsigned integer. The builtin takes its inputs as packed 32-bit values. The kernel is responsible for packing the individual bytes before the call. The fourth argument is a Boolean that negates the result of the dot product before accumulation. Passing false gives a plain accumulate.

// udot4: dot product of four unsigned byte pairs accumulated into uint32 (dot7-insts).
// Bytes are packed inline per thread before calling the intrinsic.
__global__ __launch_bounds__(block_size)
void udot4(const unsigned char *a, const unsigned char *b, unsigned int *output, int n)
{
    int gid  = blockIdx.x * blockDim.x + threadIdx.x;
    int base = gid * 4;
    if (base + 3 >= n)
    {
        return;
    }
    unsigned int a_word = (unsigned int)a[base + 0]
                        | ((unsigned int)a[base + 1] << 8)
                        | ((unsigned int)a[base + 2] << 16)
                        | ((unsigned int)a[base + 3] << 24);
    unsigned int b_word = (unsigned int)b[base + 0]
                        | ((unsigned int)b[base + 1] << 8)
                        | ((unsigned int)b[base + 2] << 16)
                        | ((unsigned int)b[base + 3] << 24);
    output[gid] = __builtin_amdgcn_udot4(a_word, b_word, 0u, false);
}

Each thread packs four consecutive bytes from each input array into a 32-bit value by using shifts and bitwise OR, then calls the builtin with an accumulator of 0. The packing is the only per-thread overhead; the dot product itself is one instruction. For signed integer dot products, see sdot4 in the integer dot product reference.

For the full signature and parameter details, see udot4 in the integer dot product reference.

FP16 dot product

The FP16 dot product builtin computes the inner product of two pairs of half-precision values and accumulates into a single-precision float. The primary motivation is throughput: FP16 multiply-add has higher theoretical throughput than FP32 (32-bit single-precision) on hardware with dedicated mixed-precision units, and accumulating into FP32 preserves the dynamic range needed for deep learning and signal processing workloads.

FP16-to-FP32 accumulation with fdot2

__builtin_amdgcn_fdot2 computes a[0] * b[0] + a[1] * b[1] in FP16 precision and accumulates the result into a FP32 accumulator. Inputs are passed as half2 two-element vectors. The fourth argument negates the dot product before accumulation; passing false gives a plain accumulate.

// fdot2: dot product of two FP16 pairs accumulated into FP32 (dot10-insts).
// Each thread loads one pair of half2 vectors from consecutive elements.
__global__ __launch_bounds__(block_size)
void fdot2(const __fp16 *a, const __fp16 *b, float *output, int n)
{
    int gid  = blockIdx.x * blockDim.x + threadIdx.x;
    int base = gid * 2;
    if (base + 1 >= n)
    {
        return;
    }
    using half2 = __attribute__((vector_size(4))) __fp16;
    half2 va    = *reinterpret_cast<const half2 *>(&a[base]);
    half2 vb    = *reinterpret_cast<const half2 *>(&b[base]);
    output[gid] = __builtin_amdgcn_fdot2(va, vb, 0.0f, false);
}

The kernel casts two consecutive __fp16 array elements to a half2 vector via reinterpret_cast before the call. The builtin is most valuable when the surrounding code is complex enough that the compiler does not automatically generate the equivalent hardware instruction.

For the full signature and parameter details, see fdot2 in the floating-point dot product reference.

Conversion and packing

The SAD and dot product builtins expect their inputs packed as byte fields within 32-bit values. When your source data is stored as floats, converting and packing manually would require several steps. __builtin_amdgcn_cvt_pk_u8_f32 handles both in one instruction: it clamps the float to [0, 255], converts it to an unsigned byte, and inserts it into one of the four byte positions of an existing destination 32-bit value, leaving the other three bytes unchanged.

Byte packing with cvt_pk_u8_f32

__builtin_amdgcn_cvt_pk_u8_f32 takes the source float, a compile-time byte selector (0–3), and the current destination 32-bit value. Four calls, advancing the selector from 0 to 3 and threading the output of each call into the next, fill all four byte slots of the output 32-bit value in four instructions with no intermediate registers.

// cvt_pk_u8_f32: clamp-and-convert one FP32 to uint8 and insert it at byte position i of dst.
// Four calls pack a full word suitable for use with sad_u8, udot4, or similar.
__global__ __launch_bounds__(block_size)
void cvt_pk_u8_f32(const float *input, unsigned int *output, int n)
{
    int gid  = blockIdx.x * blockDim.x + threadIdx.x;
    int base = gid * 4;
    if (base + 3 >= n)
    {
        return;
    }
    unsigned int word = 0u;
    word = __builtin_amdgcn_cvt_pk_u8_f32(input[base + 0], 0, word);
    word = __builtin_amdgcn_cvt_pk_u8_f32(input[base + 1], 1, word);
    word = __builtin_amdgcn_cvt_pk_u8_f32(input[base + 2], 2, word);
    word = __builtin_amdgcn_cvt_pk_u8_f32(input[base + 3], 3, word);
    output[gid] = word;
}

The packed 32-bit value is directly usable as the src0 argument to sad_u8 or as a component of the packed inputs to udot4, making this builtin the natural bridge between float pipelines and the byte-oriented arithmetic builtins in this chapter.

For the full signature and parameter details, see cvt_pk_u8_f32 in the conversion and packing reference.

Compile and run:

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

Note

The example above targets CDNA3 (gfx942). Replace --offload-arch with the appropriate target for your GPU — for example, gfx90a for CDNA2 or gfx1200 for RDNA4. All builtins used in the examples are available on every supported architecture.

Naming convention

All hardware arithmetic and packing builtins use the prefix __builtin_amdgcn_ followed by an operation name and type suffix:

__builtin_amdgcn_<operation>[_<qualifier>]_<type>

operation

The core operation (sad, msad, qsad, udot, sdot, sudot, fdot, cvt).

qualifier (optional)

A modifier that refines the operation. hi places the result in the upper half (sad_hi_u8). pk indicates a pack/unpack conversion (cvt_pk_u8_f32). sr indicates stochastic rounding (cvt_sr_fp8_f32).

type

The element type. u8 = unsigned 8-bit, u16 = unsigned 16-bit, f32 = FP32, f16 = FP16, bf16 = BF16, fp8 = FP8 (E4M3), bf8 = BF8 (E5M2), i4 = 4-bit fixed-point. When two types appear, the first is the input and the second is the output (cvt_pk_u8_f32: FP32 input, uint8 output).

The trailing digit on dot-product names indicates the number of element pairs per call: udot4 computes a 4-element dot product, udot8 an 8-element dot product, fdot2 a 2-element dot product.

Arithmetic and packing builtin reference

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

Family	Description
Sum of absolute differences	L1 distance between packed byte vectors (msad_u8, sad_u8, sad_hi_u8, sad_u16, qsad_pk_u16_u8, mqsad_pk_u16_u8, mqsad_u32_u8)
Integer dot products	Inner products of packed integer vectors (udot4, udot8, sdot4, sdot8, sdot2, udot2, sudot4, sudot8)
Floating-point dot products	Inner products of packed FP16, BF16, FP8, and BF8 vectors (fdot2, fdot2_f16_f16, fdot2_bf16_bf16, fdot2_f32_bf16, dot4_f32_fp8_*, fdot2c_f32_bf16)
Conversion and packing	Float-to-byte, byte-to-float, and pack/unpack conversions (cvt_pk_*, cvt_pkrtz, cvt_pknorm_*, cvt_sr_*, cvt_f32_*)