Sum of absolute differences builtins

Contents

Sum of absolute differences builtins#

The SAD builtins compute the L1 distance between packed byte vectors in a single instruction. They are used in motion estimation, template matching, and image comparison kernels where many small absolute differences must be accumulated quickly.

Common parameters#

All SAD builtins share the same three-operand pattern:

src0 (reference block): Four packed unsigned bytes (or eight bytes for the qsad and mqsad variants). This is the data being compared against a query.
src1 (query block): Four packed unsigned bytes representing the query. For msad_u8 and mqsad_pk_u16_u8, byte positions where src1 is zero are excluded from the accumulation — a masking behavior designed for padded reference windows where positions outside the frame boundary are marked with zero.
src2 (accumulator): Initial accumulator value. Pass 0 for a fresh SAD. Pass the result of a previous call to chain multiple SAD operations into a single running total without additional add instructions.

Architecture availability#

The following table summarizes architecture support for each builtin.

Builtin	CDNA	CDNA2	CDNA3	CDNA4	RDNA2	RDNA3	RDNA3.5	RDNA4
`msad_u8`, `sad_u8`, `sad_hi_u8`, `sad_u16`	Yes	Yes	Yes	Yes	Yes	Yes	Yes	Yes
`qsad_pk_u16_u8`, `mqsad_pk_u16_u8`, `mqsad_u32_u8`	Yes	Yes	Yes	Yes	Yes	Yes	Yes	Yes

Builtin reference#

Each builtin’s full signature, parameters, and return value are documented below.

`__builtin_amdgcn_msad_u8`#

Signature and parameters for this builtin.

unsigned int __builtin_amdgcn_msad_u8(
    unsigned int src0,
    unsigned int src1,
    unsigned int src2);

Computes the sum of absolute differences of four packed unsigned bytes, skipping byte positions where the corresponding src1 byte is zero. This masking behavior is designed for padded reference windows, where positions outside the frame boundary are marked with zero rather than valid pixel data.

Parameter	Type	Description
`src0`	unsigned int	Four packed unsigned bytes: the source data to compare.
`src1`	unsigned int	Four packed unsigned bytes: the reference data. Byte positions where `src1` is zero are excluded from the accumulation.
`src2`	unsigned int	Initial accumulator value. Pass `0u` for a fresh SAD; pass the result of a previous call to accumulate across multiple 32-bit values.

Returns unsigned int – accumulated SAD (\(\text{src2} + \sum_{k=0}^{3} |\text{src0}[k] - \text{src1}[k]|\) for positions where src1[k] != 0).

`__builtin_amdgcn_sad_u8`#

Signature and parameters for this builtin.

unsigned int __builtin_amdgcn_sad_u8(
    unsigned int src0,
    unsigned int src1,
    unsigned int src2);

Computes the sum of absolute differences of four packed unsigned bytes and accumulates the result into bits [15:0] of the destination. Pair with sad_hi_u8 to pack two independent 4-byte SADs into a single 32-bit value.

Parameter	Type	Description
`src0`	unsigned int	Four packed unsigned bytes: the source data.
`src1`	unsigned int	Four packed unsigned bytes: the reference data.
`src2`	unsigned int	Initial accumulator value (bits [15:0] are used).

Returns unsigned int – accumulated SAD in bits [15:0] (\(\text{src2}[15{:}0] + \sum_{k=0}^{3} |\text{src0}[k] - \text{src1}[k]|\)).

`__builtin_amdgcn_sad_hi_u8`#

Signature and parameters for this builtin.

unsigned int __builtin_amdgcn_sad_hi_u8(
    unsigned int src0,
    unsigned int src1,
    unsigned int src2);

Computes the sum of absolute differences of four packed unsigned bytes, shifts each per-byte absolute difference left by 16 bits before accumulating, and places the result in bits [31:16] of the destination. Pair with sad_u8 to pack two independent 4-byte SADs into a single 32-bit value without additional packing instructions.

Parameter	Type	Description
`src0`	unsigned int	Four packed unsigned bytes: the source data.
`src1`	unsigned int	Four packed unsigned bytes: the reference data.
`src2`	unsigned int	Initial accumulator value (bits [31:16] are used).

Returns unsigned int – accumulated SAD in bits [31:16].

`__builtin_amdgcn_sad_u16`#

Signature and parameters for this builtin.

unsigned int __builtin_amdgcn_sad_u16(
    unsigned int src0,
    unsigned int src1,
    unsigned int src2);

Computes the sum of absolute differences of two packed unsigned 16-bit values and accumulates the result into src2.

Parameter	Type	Description
`src0`	unsigned int	Two packed unsigned 16-bit values: the source data.
`src1`	unsigned int	Two packed unsigned 16-bit values: the reference data.
`src2`	unsigned int	Initial accumulator value.

Returns unsigned int – accumulated SAD (\(\text{src2} + \sum_{k=0}^{1} |\text{src0}[k] - \text{src1}[k]|\)).

`__builtin_amdgcn_qsad_pk_u16_u8`#

Signature and parameters for this builtin.

uint64_t __builtin_amdgcn_qsad_pk_u16_u8(
    uint64_t     src0,
    unsigned int src1,
    uint64_t     src2);

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] against the same 4-byte query (src1). The four 16-bit results are packed into the returned 64-bit value.

Parameter	Type	Description
`src0`	uint64_t	Eight packed unsigned bytes: the sliding reference window.
`src1`	unsigned int	Four packed unsigned bytes: the query to compare against each window position.
`src2`	uint64_t	Initial accumulator for all four channels simultaneously.

Returns uint64_t – four 16-bit SAD results packed into a 64-bit value. Extract individual results with (result >> (i * 16)) & 0xffff for window offset i.

`__builtin_amdgcn_mqsad_pk_u16_u8`#

Signature and parameters for this builtin.

uint64_t __builtin_amdgcn_mqsad_pk_u16_u8(
    uint64_t     src0,
    unsigned int src1,
    uint64_t     src2);

Masked variant of qsad_pk_u16_u8. Computes four overlapping 4-byte SADs from an 8-byte sliding window, but skips byte positions where the corresponding src1 byte is zero. Provides the same sliding-window behavior as qsad_pk_u16_u8 with the same masking semantics as msad_u8.

Parameter	Type	Description
`src0`	uint64_t	Eight packed unsigned bytes: the sliding reference window.
`src1`	unsigned int	Four packed unsigned bytes: the query. Byte positions where `src1` is zero are excluded from the accumulation.
`src2`	uint64_t	Initial accumulator for all four channels simultaneously.

Returns uint64_t – four 16-bit masked SAD results packed into a 64-bit value.

`__builtin_amdgcn_mqsad_u32_u8`#

Signature and parameters for this builtin.

uint4 __builtin_amdgcn_mqsad_u32_u8(
    uint64_t     src0,
    unsigned int src1,
    uint4        src2);

Computes four overlapping 4-byte SADs from an 8-byte sliding window, returning four full 32-bit results instead of the 16-bit packed format of the qsad/mqsad variants. Use this when the accumulated SAD values may exceed the 16-bit range (65535).

Parameter	Type	Description
`src0`	uint64_t	Eight packed unsigned bytes: the sliding reference window.
`src1`	unsigned int	Four packed unsigned bytes: the query.
`src2`	uint4	Four-element uint32 accumulator vector.

Returns uint4 – four 32-bit SAD results.