https://github.com/shiltian updated https://github.com/llvm/llvm-project/pull/181193
>From c503914390e150e8163b11bd0a8da99cb6eadfda Mon Sep 17 00:00:00 2001 From: Shilei Tian <[email protected]> Date: Thu, 12 Feb 2026 12:36:49 -0500 Subject: [PATCH 1/3] [RFC][Docs][Clang][AMDGPU] Add AMDGPU builtins documentation Add comprehensive documentation for AMDGPU target-specific builtins (`AMDGPUBuiltins.rst`) covering argument semantics, restrictions, and lowering notes for all builtin families. This documentation was generated by AI (Claude) by cross-referencing: - `clang/include/clang/Basic/BuiltinsAMDGPU.td` (builtin definitions) - `llvm/include/llvm/IR/IntrinsicsAMDGPU.td` (intrinsic definitions) - `clang/lib/Sema/SemaAMDGPU.cpp` (argument validation/constraints) - `clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp` (lowering logic) I did my best to proofread the parts I'm familiar with, but it would be greatly appreciated if more people could help review it as well. --- clang/docs/AMDGPUBuiltins.rst | 1807 +++++++++++++++++++++++++++++++++ clang/docs/index.rst | 1 + 2 files changed, 1808 insertions(+) create mode 100644 clang/docs/AMDGPUBuiltins.rst diff --git a/clang/docs/AMDGPUBuiltins.rst b/clang/docs/AMDGPUBuiltins.rst new file mode 100644 index 0000000000000..9ca9a2bf3bd32 --- /dev/null +++ b/clang/docs/AMDGPUBuiltins.rst @@ -0,0 +1,1807 @@ +=============== +AMDGPU Builtins +=============== + +.. contents:: + :local: + :depth: 2 + +This document describes the AMDGPU target-specific builtins available in Clang. +Most of these builtins provide direct access to AMDGPU hardware instructions +and intrinsics. They are defined in ``clang/include/clang/Basic/BuiltinsAMDGPU.td`` +and typically lower to LLVM intrinsics defined in +``llvm/include/llvm/IR/IntrinsicsAMDGPU.td``. + +All AMDGPU builtins use the ``__builtin_amdgcn_`` prefix (or ``__builtin_r600_`` +for R600 targets). Arguments marked ``_Constant`` must be compile-time +constant expressions. + +ABI / Special Register Builtins +=============================== + +These builtins provide access to kernel dispatch metadata, work-item and +workgroup identification, and other ABI-level information. They are available +on all SI+ targets. + +Pointer Builtins +---------------- + +.. list-table:: + :header-rows: 1 + :widths: 40 60 + + * - Builtin + - Description + * - ``void __constant * __builtin_amdgcn_dispatch_ptr()`` + - Returns a pointer (in constant address space 4) to the dispatch packet + (``hsa_kernel_dispatch_packet_t``). Used internally to derive workgroup + size, grid size, and other dispatch parameters. + * - ``void __constant * __builtin_amdgcn_kernarg_segment_ptr()`` + - Returns a pointer to the beginning of the kernel argument segment. + * - ``void __constant * __builtin_amdgcn_implicitarg_ptr()`` + - Returns a pointer to the implicit arguments appended after explicit + kernel arguments. Layout depends on the code object version. + * - ``void __constant * __builtin_amdgcn_queue_ptr()`` + - Returns a pointer to the ``hsa_queue_t`` object for the queue executing + the current kernel. + +Work-Item and Workgroup Identification +-------------------------------------- + +All of these are ``Const`` (pure) builtins that take no arguments and return +``unsigned int`` (or ``unsigned short`` for workgroup size). + +.. list-table:: + :header-rows: 1 + :widths: 40 30 30 + + * - Builtin + - Return Type + - Description + * - ``__builtin_amdgcn_workgroup_id_{x,y,z}()`` + - ``unsigned int`` + - Workgroup ID in the specified dimension. + * - ``__builtin_amdgcn_workitem_id_{x,y,z}()`` + - ``unsigned int`` + - Work-item (thread) ID within the workgroup. + * - ``__builtin_amdgcn_workgroup_size_{x,y,z}()`` + - ``unsigned short`` + - Workgroup size in the specified dimension. Lowered via a load from the + dispatch or implicit argument pointer, not a dedicated instruction. + * - ``__builtin_amdgcn_grid_size_{x,y,z}()`` + - ``unsigned int`` + - Total grid size in the specified dimension. Lowered via a load from the + dispatch pointer. + +**GFX1250+ Cluster Identification** (requires ``gfx1250-insts``): + +.. list-table:: + :header-rows: 1 + :widths: 50 50 + + * - Builtin + - Description + * - ``__builtin_amdgcn_cluster_id_{x,y,z}()`` + - Cluster ID in the specified dimension. + * - ``__builtin_amdgcn_cluster_workgroup_id_{x,y,z}()`` + - Workgroup ID within the cluster. + * - ``__builtin_amdgcn_cluster_workgroup_flat_id()`` + - Flat (linearized) workgroup ID within the cluster. + * - ``__builtin_amdgcn_cluster_workgroup_max_id_{x,y,z}()`` + - Maximum workgroup ID within the cluster. + * - ``__builtin_amdgcn_cluster_workgroup_max_flat_id()`` + - Maximum flat workgroup ID within the cluster. + +Other ABI Builtins +------------------ + +.. list-table:: + :header-rows: 1 + :widths: 40 60 + + * - Builtin + - Description + * - ``unsigned int __builtin_amdgcn_mbcnt_lo(unsigned int mask, unsigned int val)`` + - Counts the number of set bits in ``mask`` for lanes lower than the + current lane within the lower 32 bits of the exec mask, adds ``val``. + * - ``unsigned int __builtin_amdgcn_mbcnt_hi(unsigned int mask, unsigned int val)`` + - Same as ``mbcnt_lo`` but for the upper 32 bits of the exec mask. + * - ``uint64_t __builtin_amdgcn_s_memtime()`` + - Returns a 64-bit timestamp counter. Requires ``s-memtime-inst``. + +Instruction Builtins +==================== + +Scalar Instruction Builtins +--------------------------- + +.. list-table:: + :header-rows: 1 + :widths: 40 60 + + * - Builtin + - Description + * - ``unsigned int __builtin_amdgcn_s_getreg(_Constant int hwreg)`` + - Reads a hardware register. ``hwreg`` is an encoded register specifier + (register ID, offset, and width packed into 16 bits). + * - ``void __builtin_amdgcn_s_setreg(_Constant int hwreg, unsigned int val)`` + - Writes ``val`` to a hardware register. ``hwreg`` must be in + range [0, 65535]. + * - ``uint64_t __builtin_amdgcn_s_getpc()`` + - Returns the current program counter. + * - ``void __builtin_amdgcn_s_waitcnt(_Constant int cnt)`` + - Inserts an ``s_waitcnt`` instruction with the encoded wait count. + * - ``void __builtin_amdgcn_s_sendmsg(_Constant int msg, unsigned int gsdata)`` + - Sends message ``msg`` with GS data in ``gsdata``. + * - ``void __builtin_amdgcn_s_sendmsghalt(_Constant int msg, unsigned int gsdata)`` + - Same as ``s_sendmsg`` but also halts the wavefront. + * - ``void __builtin_amdgcn_s_barrier()`` + - Inserts a workgroup barrier. + * - ``void __builtin_amdgcn_s_ttracedata(int data)`` + - Writes ``data`` to the thread trace buffer. + * - ``void __builtin_amdgcn_s_sleep(_Constant int duration)`` + - Sleeps for approximately ``duration`` cycles. + * - ``void __builtin_amdgcn_s_incperflevel(_Constant int level)`` + - Increments the performance counter level. + * - ``void __builtin_amdgcn_s_decperflevel(_Constant int level)`` + - Decrements the performance counter level. + * - ``void __builtin_amdgcn_s_setprio(_Constant short prio)`` + - Sets the wavefront priority. + * - ``void __builtin_amdgcn_s_dcache_inv()`` + - Invalidates the scalar data cache. + * - ``void __builtin_amdgcn_buffer_wbinvl1()`` + - Write-back and invalidate L1 buffer cache. + * - ``unsigned int __builtin_amdgcn_groupstaticsize()`` + - Returns the size of static LDS allocation in the current workgroup. + * - ``unsigned int __builtin_amdgcn_wavefrontsize()`` + - Returns the wavefront size (32 or 64). + * - ``void __builtin_amdgcn_wave_barrier()`` + - Inserts a wave-level barrier hint. + +Division and Math Builtins +-------------------------- + +Division Support +^^^^^^^^^^^^^^^^ + +These builtins implement steps of the iterative double-precision division +algorithm. + +``__builtin_amdgcn_div_scale`` / ``__builtin_amdgcn_div_scalef`` +~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ + +.. code-block:: c + + double __builtin_amdgcn_div_scale(double numer, double denom, bool select_quotient, bool *flag_out); + float __builtin_amdgcn_div_scalef(float numer, float denom, bool select_quotient, bool *flag_out); + +Scales the numerator or denominator for a subsequent iterative division. + +- ``numer``: The numerator. +- ``denom``: The denominator. +- ``select_quotient``: If ``true``, selects the numerator for scaling; if + ``false``, selects the denominator. +- ``flag_out``: Pointer to a ``bool`` where the overflow/underflow flag is + written. + +**Lowering note**: The underlying intrinsic returns ``{result, flag}`` as a +struct. The builtin unpacks this, returning the result and storing the flag +through the pointer. + +``__builtin_amdgcn_div_fmas`` / ``__builtin_amdgcn_div_fmasf`` +~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ + +.. code-block:: c + + double __builtin_amdgcn_div_fmas(double a, double b, double c, bool vcc); + float __builtin_amdgcn_div_fmasf(float a, float b, float c, bool vcc); + +Fused multiply-add for division, with VCC flag input. + +- ``a``, ``b``, ``c``: FMA operands (computes ``a * b + c``). +- ``vcc``: The flag from ``div_scale``. + +**Lowering note**: The integer ``vcc`` argument is converted to ``i1`` via +``IsNotNull`` before passing to the intrinsic. + +``__builtin_amdgcn_div_fixup`` / ``__builtin_amdgcn_div_fixupf`` / ``__builtin_amdgcn_div_fixuph`` +~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ + +.. code-block:: c + + double __builtin_amdgcn_div_fixup(double a, double b, double c); + float __builtin_amdgcn_div_fixupf(float a, float b, float c); + __fp16 __builtin_amdgcn_div_fixuph(__fp16 a, __fp16 b, __fp16 c); // requires 16-bit-insts + +Applies post-division fixup for special values (NaN, Inf, zero). + +Trigonometric Pre-operation +^^^^^^^^^^^^^^^^^^^^^^^^^^^ + +.. code-block:: c + + double __builtin_amdgcn_trig_preop(double src, int segment); + float __builtin_amdgcn_trig_preopf(float src, int segment); + +Looks up ``2.0 / pi`` with segment selector ``segment[4:0]`` for range +reduction before trigonometric operations. + +Single-Argument Math Builtins +^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ + +These builtins compute hardware-precision math operations. The ``f32`` versions +(e.g., ``sinf``, ``logf``) may not handle denormals correctly. The ``h``-suffixed +variants require ``16-bit-insts``. + +.. list-table:: + :header-rows: 1 + :widths: 25 25 25 25 + + * - Operation + - f64 + - f32 + - f16 + * - Reciprocal + - ``__builtin_amdgcn_rcp`` + - ``__builtin_amdgcn_rcpf`` + - ``__builtin_amdgcn_rcph`` + * - Square root + - ``__builtin_amdgcn_sqrt`` + - ``__builtin_amdgcn_sqrtf`` + - ``__builtin_amdgcn_sqrth`` + * - Reciprocal sqrt + - ``__builtin_amdgcn_rsq`` + - ``__builtin_amdgcn_rsqf`` + - ``__builtin_amdgcn_rsqh`` + * - Reciprocal sqrt clamp + - ``__builtin_amdgcn_rsq_clamp`` + - ``__builtin_amdgcn_rsq_clampf`` + - + * - Sine (input: turns) + - + - ``__builtin_amdgcn_sinf`` + - ``__builtin_amdgcn_sinh`` + * - Cosine (input: turns) + - + - ``__builtin_amdgcn_cosf`` + - ``__builtin_amdgcn_cosh`` + * - Log2 + - + - ``__builtin_amdgcn_logf`` + - + * - Log clamp + - + - ``__builtin_amdgcn_log_clampf`` + - + * - Exp2 + - + - ``__builtin_amdgcn_exp2f`` + - + * - Fraction + - ``__builtin_amdgcn_fract`` + - ``__builtin_amdgcn_fractf`` + - ``__builtin_amdgcn_fracth`` + * - Mantissa + - ``__builtin_amdgcn_frexp_mant`` + - ``__builtin_amdgcn_frexp_mantf`` + - ``__builtin_amdgcn_frexp_manth`` + * - Exponent + - ``__builtin_amdgcn_frexp_exp`` + - ``__builtin_amdgcn_frexp_expf`` + - ``__builtin_amdgcn_frexp_exph`` + +Note: ``sinf``/``cosf`` take input in **turns** (1.0 = full circle), not +radians. ``logf`` performs ``log2``. ``exp2f`` performs ``2^x``. The ``frexp_exp`` +variants return ``int`` (or ``short`` for f16). + +Ldexp +^^^^^ + +.. code-block:: c + + double __builtin_amdgcn_ldexp(double x, int exp); + float __builtin_amdgcn_ldexpf(float x, int exp); + __fp16 __builtin_amdgcn_ldexph(__fp16 x, int exp); // requires 16-bit-insts + +Computes ``x * 2^exp``. Lowered to the standard ``llvm.ldexp`` intrinsic. +For the ``h`` variant, the exponent is truncated to ``i16``. + +FP Classify +^^^^^^^^^^^ + +.. code-block:: c + + bool __builtin_amdgcn_class(double x, int mask); + bool __builtin_amdgcn_classf(float x, int mask); + bool __builtin_amdgcn_classh(__fp16 x, int mask); // requires 16-bit-insts + +Tests ``x`` against a bitmask of FP classes. Returns ``true`` if ``x`` matches +any of the selected classes. The ``mask`` bits are: + +- Bit 0: Signaling NaN +- Bit 1: Quiet NaN +- Bit 2: Negative infinity +- Bit 3: Negative normal +- Bit 4: Negative denormal +- Bit 5: Negative zero +- Bit 6: Positive zero +- Bit 7: Positive denormal +- Bit 8: Positive normal +- Bit 9: Positive infinity + +Median +^^^^^^ + +.. code-block:: c + + float __builtin_amdgcn_fmed3f(float a, float b, float c); + __fp16 __builtin_amdgcn_fmed3h(__fp16 a, __fp16 b, __fp16 c); // requires gfx9-insts + +Returns the median (middle value) of three floating-point numbers. + +Cube Map Builtins +^^^^^^^^^^^^^^^^^ + +Require ``cube-insts``. All take three floats (x, y, z direction vector +components) and return a float. + +.. list-table:: + :header-rows: 1 + :widths: 40 60 + + * - Builtin + - Description + * - ``__builtin_amdgcn_cubeid(x, y, z)`` + - Returns the face ID (0-5) of the cube map. + * - ``__builtin_amdgcn_cubesc(x, y, z)`` + - Returns the S coordinate for the cube face. + * - ``__builtin_amdgcn_cubetc(x, y, z)`` + - Returns the T coordinate for the cube face. + * - ``__builtin_amdgcn_cubema(x, y, z)`` + - Returns the major axis value. + +Data Sharing Builtins +--------------------- + +.. list-table:: + :header-rows: 1 + :widths: 40 60 + + * - Builtin + - Description + * - ``int __builtin_amdgcn_ds_swizzle(int data, _Constant int pattern)`` + - Performs a data-parallel swizzle within the wavefront according to the + encoded ``pattern``. + * - ``int __builtin_amdgcn_ds_permute(int addr, int data)`` + - Forward cross-lane permutation. Lane ``i`` gets the ``data`` value from + the lane specified by ``addr / 4``. + * - ``int __builtin_amdgcn_ds_bpermute(int addr, int data)`` + - Backward cross-lane permutation. Lane ``i`` reads from the lane + specified by ``addr / 4``. + * - ``int __builtin_amdgcn_ds_append(int __local *ptr)`` + - Atomically increments the value at ``ptr`` and returns the old value. + The pointer must be in LDS (address space 3). + * - ``int __builtin_amdgcn_ds_consume(int __local *ptr)`` + - Atomically decrements the value at ``ptr`` and returns the new value. + +DS Float Atomics +^^^^^^^^^^^^^^^^ + +.. code-block:: c + + float __builtin_amdgcn_ds_faddf(float __local *ptr, float val, _Constant int ordering, _Constant int scope, _Constant bool isVolatile); + float __builtin_amdgcn_ds_fminf(float __local *ptr, float val, _Constant int ordering, _Constant int scope, _Constant bool isVolatile); + float __builtin_amdgcn_ds_fmaxf(float __local *ptr, float val, _Constant int ordering, _Constant int scope, _Constant bool isVolatile); + +Perform atomic float add/min/max on LDS memory. The ``ordering`` and ``scope`` +arguments are passed through but the operations are lowered to ``AtomicRMW`` +instructions. + +Lane Builtins +------------- + +.. list-table:: + :header-rows: 1 + :widths: 40 60 + + * - Builtin + - Description + * - ``int __builtin_amdgcn_readfirstlane(int val)`` + - Returns the value of ``val`` from the first active lane. + * - ``int __builtin_amdgcn_readlane(int val, int lane)`` + - Returns the value of ``val`` from the specified ``lane``. + +Bit Manipulation +---------------- + +.. list-table:: + :header-rows: 1 + :widths: 40 60 + + * - Builtin + - Description + * - ``unsigned int __builtin_amdgcn_alignbit(unsigned int hi, unsigned int lo, unsigned int shift)`` + - Concatenates ``hi:lo`` as a 64-bit value and extracts 32 bits starting + at bit ``shift``. Lowered to ``llvm.fshr``. + * - ``unsigned int __builtin_amdgcn_alignbyte(unsigned int hi, unsigned int lo, unsigned int shift)`` + - Same as ``alignbit`` but ``shift`` is in bytes. + * - ``unsigned int __builtin_amdgcn_ubfe(unsigned int base, unsigned int offset, unsigned int width)`` + - Unsigned bitfield extract from ``base`` starting at ``offset`` for + ``width`` bits. + * - ``unsigned int __builtin_amdgcn_sbfe(unsigned int base, unsigned int offset, unsigned int width)`` + - Signed bitfield extract. + * - ``unsigned int __builtin_amdgcn_lerp(unsigned int a, unsigned int b, unsigned int c)`` + - Per-byte unsigned linear interpolation. Requires ``lerp-inst``. + * - ``unsigned int __builtin_amdgcn_perm(unsigned int a, unsigned int b, unsigned int sel)`` + - Byte permutation. ``sel`` encodes which byte of the ``a:b`` pair to + select for each byte of the result. Requires ``gfx8-insts``. + +Conversion Builtins +------------------- + +.. list-table:: + :header-rows: 1 + :widths: 40 60 + + * - Builtin + - Description + * - ``half2 __builtin_amdgcn_cvt_pkrtz(float a, float b)`` + - Converts two f32 values to packed f16 with round-to-zero. + * - ``short2 __builtin_amdgcn_cvt_pknorm_i16(float a, float b)`` + - Converts two f32 values to packed normalized i16. Requires + ``cvt-pknorm-vop2-insts``. + * - ``ushort2 __builtin_amdgcn_cvt_pknorm_u16(float a, float b)`` + - Converts two f32 values to packed normalized u16. + * - ``short2 __builtin_amdgcn_cvt_pk_i16(int a, int b)`` + - Packs two i32 values into i16x2. + * - ``ushort2 __builtin_amdgcn_cvt_pk_u16(unsigned int a, unsigned int b)`` + - Packs two u32 values into u16x2. + * - ``unsigned int __builtin_amdgcn_cvt_pk_u8_f32(float val, unsigned int bytesel, unsigned int old)`` + - Converts ``val`` to u8 and inserts at byte ``bytesel`` in ``old``. + * - ``float __builtin_amdgcn_cvt_off_f32_i4(int val)`` + - Converts a 4-bit integer offset to f32. + +SAD (Sum of Absolute Differences) +--------------------------------- + +.. list-table:: + :header-rows: 1 + :widths: 40 60 + + * - Builtin + - Description + * - ``unsigned int __builtin_amdgcn_msad_u8(unsigned int a, unsigned int b, unsigned int c)`` + - Masked sum of absolute differences of unsigned 8-bit values. + * - ``unsigned int __builtin_amdgcn_sad_u8(unsigned int a, unsigned int b, unsigned int c)`` + - Sum of absolute differences of unsigned 8-bit values. Requires + ``sad-insts``. + * - ``unsigned int __builtin_amdgcn_sad_hi_u8(unsigned int a, unsigned int b, unsigned int c)`` + - SAD with result in high 16 bits. Requires ``sad-insts``. + * - ``unsigned int __builtin_amdgcn_sad_u16(unsigned int a, unsigned int b, unsigned int c)`` + - SAD of unsigned 16-bit values. Requires ``sad-insts``. + * - ``uint64_t __builtin_amdgcn_qsad_pk_u16_u8(uint64_t a, unsigned int b, uint64_t c)`` + - Quad SAD packed. Requires ``qsad-insts``. + * - ``uint64_t __builtin_amdgcn_mqsad_pk_u16_u8(uint64_t a, unsigned int b, uint64_t c)`` + - Masked quad SAD packed. + * - ``uint4 __builtin_amdgcn_mqsad_u32_u8(uint64_t a, unsigned int b, uint4 c)`` + - Masked quad SAD returning 4x u32. + +Buffer Resource and Load/Store +============================== + +make_buffer_rsrc +---------------- + +.. code-block:: c + + __amdgpu_buffer_rsrc_t __builtin_amdgcn_make_buffer_rsrc(void *base, short stride, int64_t num_records, int flags); + +Constructs a buffer resource descriptor from the given fields: + +- ``base``: Base pointer. +- ``stride``: Stride of structured buffer (0 for raw). +- ``num_records``: Number of records (bytes for raw buffers). +- ``flags``: SRD flags (DST_SEL, NUM_FORMAT, DATA_FORMAT, etc.). + +Raw Buffer Load/Store +--------------------- + +These builtins load/store data through a buffer resource descriptor. + +.. code-block:: c + + // Stores + void __builtin_amdgcn_raw_buffer_store_b{8,16,32,64,96,128}(data, __amdgpu_buffer_rsrc_t rsrc, int offset, int soffset, _Constant int cachepolicy); + // Loads + T __builtin_amdgcn_raw_buffer_load_b{8,16,32,64,96,128}(__amdgpu_buffer_rsrc_t rsrc, int offset, int soffset, _Constant int cachepolicy); + +Arguments: + +- ``rsrc``: Buffer resource descriptor (128-bit SRD, typically SGPR). +- ``offset``: Byte offset (VGPR or immediate). Included in bounds checking and + swizzling. +- ``soffset``: Scalar byte offset (SGPR or immediate). Excluded from bounds + checking and swizzling. +- ``cachepolicy``: Immediate bitfield controlling cache behavior: + + - Pre-GFX12: bit 0 = GLC, bit 1 = SLC, bit 2 = DLC (gfx10/gfx11), + bit 3 = SWZ, bit 4 = SCC (gfx90a). + - GFX942: bit 0 = SC0, bit 1 = NT, bit 3 = SWZ, bit 4 = SC1. + - GFX12+: bits [0:2] = TH, bits [3:4] = scope, bit 6 = SWZ. + - All: bit 31 = volatile (stripped at lowering). + +The data types for each width are: ``b8`` = ``unsigned char``, +``b16`` = ``unsigned short``, ``b32`` = ``unsigned int``, +``b64`` = ``uint2``, ``b96`` = ``uint3``, ``b128`` = ``uint4``. + +Raw Ptr Buffer Atomics +---------------------- + +.. code-block:: c + + int __builtin_amdgcn_raw_ptr_buffer_atomic_add_i32(int val, __amdgpu_buffer_rsrc_t rsrc, int offset, int soffset, _Constant int cachepolicy); + float __builtin_amdgcn_raw_ptr_buffer_atomic_fadd_f32(...); // requires atomic-fadd-rtn-insts + half2 __builtin_amdgcn_raw_ptr_buffer_atomic_fadd_v2f16(...); // requires atomic-buffer-global-pk-add-f16-insts + float __builtin_amdgcn_raw_ptr_buffer_atomic_fmin_f32(...); // requires atomic-fmin-fmax-global-f32 + float __builtin_amdgcn_raw_ptr_buffer_atomic_fmax_f32(...); // requires atomic-fmin-fmax-global-f32 + double __builtin_amdgcn_raw_ptr_buffer_atomic_fmin_f64(...); // requires atomic-fmin-fmax-global-f64 + double __builtin_amdgcn_raw_ptr_buffer_atomic_fmax_f64(...); // requires atomic-fmin-fmax-global-f64 + +Same argument layout as raw buffer loads: ``(val, rsrc, offset, soffset, cachepolicy)``. + +Buffer-to-LDS Load Builtins +--------------------------- + +.. code-block:: c + + // Raw pointer variants + void __builtin_amdgcn_raw_ptr_buffer_load_lds(__amdgpu_buffer_rsrc_t rsrc, void __local *lds_ptr, _Constant unsigned int size, int offset, int soffset, _Constant int aux, _Constant int cachepolicy); + void __builtin_amdgcn_raw_ptr_buffer_load_async_lds(...); // same args + // Struct pointer variants (extra vindex arg) + void __builtin_amdgcn_struct_ptr_buffer_load_lds(__amdgpu_buffer_rsrc_t rsrc, void __local *lds_ptr, _Constant unsigned int size, int vindex, int offset, int soffset, _Constant int aux, _Constant int cachepolicy); + void __builtin_amdgcn_struct_ptr_buffer_load_async_lds(...); + +All require ``vmem-to-lds-load-insts``. + +**Restriction on ``size``**: Must be a compile-time constant. Valid values are +**1, 2, 4** on most targets, or **1, 2, 4, 12, 16** on targets with +``gfx950-insts``. + +Async / LDS Load Builtins +========================= + +.. code-block:: c + + void __builtin_amdgcn_load_to_lds(void *src, void __local *dst, _Constant unsigned int size, _Constant int offset, _Constant unsigned int aux); + void __builtin_amdgcn_load_async_to_lds(void *src, void __local *dst, _Constant unsigned int size, _Constant int offset, _Constant unsigned int aux); + void __builtin_amdgcn_global_load_lds(void __global *src, void __local *dst, _Constant unsigned int size, _Constant int offset, _Constant unsigned int aux); + void __builtin_amdgcn_global_load_async_lds(void __global *src, void __local *dst, _Constant unsigned int size, _Constant int offset, _Constant unsigned int aux); + +All require ``vmem-to-lds-load-insts``. + +- ``src``: Source pointer (flat or global address space). +- ``dst``: Destination pointer in LDS (address space 3). +- ``size``: Number of bytes to load. **Must be constant**. Valid: 1, 2, 4 + (or 1, 2, 4, 12, 16 with ``gfx950-insts``). +- ``offset``: Byte offset applied to ``dst``. +- ``aux``: Auxiliary cache policy. + +Async Mark Builtins +------------------- + +.. code-block:: c + + void __builtin_amdgcn_asyncmark(); + void __builtin_amdgcn_wait_asyncmark(_Constant unsigned short count); + +Mark and wait for asynchronous operations. Require ``vmem-to-lds-load-insts``. + +Ballot and Wave Builtins +======================== + +Ballot +------ + +.. code-block:: c + + uint32_t __builtin_amdgcn_ballot_w32(bool pred); // requires wavefrontsize32 + uint64_t __builtin_amdgcn_ballot_w64(bool pred); + +Returns a bitmask of active lanes where ``pred`` is ``true``. + +Inverse Ballot +-------------- + +.. code-block:: c + + bool __builtin_amdgcn_inverse_ballot_w32(uint32_t mask); // requires wavefrontsize32 + bool __builtin_amdgcn_inverse_ballot_w64(uint64_t mask); // requires wavefrontsize64 + +Returns ``true`` for the current lane if the corresponding bit in ``mask`` is +set. + +Deprecated Compare Builtins +--------------------------- + +These are deprecated in favor of ``ballot``: + +.. code-block:: c + + uint64_t __builtin_amdgcn_uicmp(unsigned int a, unsigned int b, _Constant int cmp); + uint64_t __builtin_amdgcn_uicmpl(uint64_t a, uint64_t b, _Constant int cmp); + uint64_t __builtin_amdgcn_sicmp(int a, int b, _Constant int cmp); + uint64_t __builtin_amdgcn_sicmpl(int64_t a, int64_t b, _Constant int cmp); + uint64_t __builtin_amdgcn_fcmp(double a, double b, _Constant int cmp); + uint64_t __builtin_amdgcn_fcmpf(float a, float b, _Constant int cmp); + +The ``cmp`` constant selects the comparison predicate (uses LLVM ICmp/FCmp +predicate encoding). + +Wave Reduction Builtins +----------------------- + +Perform a reduction across all active lanes in the wavefront. All take the +form: + +.. code-block:: c + + T __builtin_amdgcn_wave_reduce_<op>_<type>(T val, _Constant int32_t strategy); + +Available operations and types: + +.. list-table:: + :header-rows: 1 + :widths: 20 80 + + * - Operation + - Types + * - ``add``, ``sub`` + - ``u32``, ``u64`` + * - ``min``, ``max`` + - ``i32``, ``u32``, ``i64``, ``u64`` + * - ``and``, ``or``, ``xor`` + - ``b32``, ``b64`` + * - ``fadd``, ``fsub``, ``fmin``, ``fmax`` + - ``f32``, ``f64`` + +The ``strategy`` argument is an implementation strategy hint. + +Flat Addressing Builtins +======================== + +.. code-block:: c + + bool __builtin_amdgcn_is_shared(const void *ptr); + bool __builtin_amdgcn_is_private(const void *ptr); + +Tests whether the given flat pointer actually points to shared (LDS) or private +(scratch) memory. + +GWS (Global Wave Sync) Builtins +=============================== + +All require ``gws``. + +.. code-block:: c + + void __builtin_amdgcn_ds_gws_init(unsigned int val, unsigned int id); + void __builtin_amdgcn_ds_gws_barrier(unsigned int val, unsigned int id); + void __builtin_amdgcn_ds_gws_sema_v(unsigned int id); + void __builtin_amdgcn_ds_gws_sema_br(unsigned int val, unsigned int id); + void __builtin_amdgcn_ds_gws_sema_p(unsigned int id); + void __builtin_amdgcn_ds_gws_sema_release_all(unsigned int id); // requires ci-insts + +- ``val``: Value or count for the GWS operation. +- ``id``: GWS resource ID. + +CI+ Cache Control +----------------- + +.. code-block:: c + + void __builtin_amdgcn_s_dcache_inv_vol(); // requires ci-insts + void __builtin_amdgcn_buffer_wbinvl1_vol(); // requires ci-insts + +Invalidate volatile data cache and buffer L1 cache. + +Interpolation Builtins +====================== + +.. code-block:: c + + float __builtin_amdgcn_interp_p1(float i, unsigned int attr_chan, unsigned int attr, unsigned int m0); + float __builtin_amdgcn_interp_p2(float p1, float j, unsigned int attr_chan, unsigned int attr, unsigned int m0); + float __builtin_amdgcn_interp_mov(unsigned int param, unsigned int attr_chan, unsigned int attr, unsigned int m0); + float __builtin_amdgcn_interp_p1_f16(float i, unsigned int attr_chan, unsigned int attr, bool high, unsigned int m0); + __fp16 __builtin_amdgcn_interp_p2_f16(float p1, float j, unsigned int attr_chan, unsigned int attr, bool high, unsigned int m0); + +- ``i``, ``j``: Interpolation coordinates (barycentric). +- ``attr_chan``: Attribute channel (0-3 for x/y/z/w). +- ``attr``: Attribute index. +- ``param``: For ``interp_mov``, selects the parameter to copy (P0, P10, P20). +- ``high``: For f16 variants, selects the high or low half of the attribute. +- ``m0``: The M0 register value (parameter base offset). + +DPP (Data-Parallel Primitives) Builtins +======================================= + +Require ``dpp`` feature. + +``__builtin_amdgcn_mov_dpp`` +---------------------------- + +.. code-block:: c + + int __builtin_amdgcn_mov_dpp(int src, _Constant int dpp_ctrl, _Constant int row_mask, _Constant int bank_mask, _Constant bool bound_ctrl); + +Performs a DPP data movement with the given control pattern. The data arguments +must be arithmetic (non-complex) types. + +- ``src``: Source data. +- ``dpp_ctrl``: DPP control word encoding the permutation pattern. +- ``row_mask``: 4-bit mask selecting which rows participate. +- ``bank_mask``: 4-bit mask selecting which banks participate. +- ``bound_ctrl``: If true, out-of-bounds lanes get zero instead of their own + value. + +**Lowering note**: A poison value is prepended as the "old" operand. For types +smaller than 32 bits, the value is zero-extended to i32 for the intrinsic, then +truncated back. + +``__builtin_amdgcn_update_dpp`` +------------------------------- + +.. code-block:: c + + int __builtin_amdgcn_update_dpp(int old, int src, _Constant int dpp_ctrl, _Constant int row_mask, _Constant int bank_mask, _Constant bool bound_ctrl); + +Like ``mov_dpp`` but with an explicit ``old`` value for out-of-bounds or +masked-off lanes. Sema validates that ``old`` and ``src`` have the same type +(or compatible signed/unsigned pair). + +``__builtin_amdgcn_mov_dpp8`` +----------------------------- + +.. code-block:: c + + unsigned int __builtin_amdgcn_mov_dpp8(unsigned int src, _Constant unsigned int sel); + +Requires ``gfx10-insts``. Performs a DPP8 permutation with ``sel`` encoding +8 3-bit lane selectors. + +Permute / Lane Builtins +======================= + +GFX10+ Permute Lane +------------------- + +.. code-block:: c + + unsigned int __builtin_amdgcn_permlane16(unsigned int old, unsigned int src, unsigned int src1, unsigned int src2, _Constant bool fi, _Constant bool bc); + unsigned int __builtin_amdgcn_permlanex16(unsigned int old, unsigned int src, unsigned int src1, unsigned int src2, _Constant bool fi, _Constant bool bc); + +Requires ``gfx10-insts``. + +- ``old``: Value for inactive lanes. +- ``src``: Source value. +- ``src1``, ``src2``: Lane index selects (SGPR). +- ``fi``: Fetch inactive -- if true, fetch from inactive lanes. +- ``bc``: Bound control -- if true, out-of-bound lanes get their own value. + +GFX12+ Variable Permute Lane +---------------------------- + +.. code-block:: c + + unsigned int __builtin_amdgcn_permlane16_var(unsigned int old, unsigned int src, unsigned int lane_sel, _Constant bool fi, _Constant bool bc); + unsigned int __builtin_amdgcn_permlanex16_var(unsigned int old, unsigned int src, unsigned int lane_sel, _Constant bool fi, _Constant bool bc); + +Requires ``gfx12-insts``. Like the above but ``lane_sel`` is a VGPR. + +GFX11+ permlane64 +----------------- + +.. code-block:: c + + unsigned int __builtin_amdgcn_permlane64(unsigned int src); // requires gfx11-insts + +Swaps data between lane ``i`` and lane ``i + 32`` (i.e., between the two +halves of a wave64). This is a no-op in wave32. + +GFX950 Permute Lane Swap +------------------------ + +.. code-block:: c + + uint2 __builtin_amdgcn_permlane16_swap(unsigned int old, unsigned int src, _Constant bool fi, _Constant bool bc); + uint2 __builtin_amdgcn_permlane32_swap(unsigned int old, unsigned int src, _Constant bool fi, _Constant bool bc); + +Requires ``permlane16-swap`` / ``permlane32-swap``. Returns a 2-element vector: +element 0 is the current lane result, element 1 is the swapped lane result. + +**Lowering note**: The intrinsic returns ``{i32, i32}``; the builtin packs +this into ``<2 x i32>``. + +GFX1250 Permute Lane Builtins +----------------------------- + +Require ``gfx1250-insts,wavefrontsize32``. + +.. code-block:: c + + int __builtin_amdgcn_permlane_bcast(int old, int src, int lane_sel); + int __builtin_amdgcn_permlane_up(int old, int src, int lane_sel); + int __builtin_amdgcn_permlane_down(int old, int src, int lane_sel); + int __builtin_amdgcn_permlane_xor(int old, int src, int lane_sel); + int __builtin_amdgcn_permlane_idx_gen(int src, int lane_sel); + +Deep Learning / Dot Product Builtins +==================================== + +These builtins perform dot product accumulate operations used in deep learning +workloads. + +Float Dot Product +----------------- + +.. code-block:: c + + float __builtin_amdgcn_fdot2(half2 a, half2 b, float c, _Constant bool clamp); // dot10-insts + __fp16 __builtin_amdgcn_fdot2_f16_f16(half2 a, half2 b, __fp16 c); // dot9-insts + short __builtin_amdgcn_fdot2_bf16_bf16(short2 a, short2 b, short c); // dot9-insts + float __builtin_amdgcn_fdot2_f32_bf16(short2 a, short2 b, float c, _Constant bool clamp); // dot12-insts + float __builtin_amdgcn_fdot2c_f32_bf16(bf16x2 a, bf16x2 b, float c, _Constant bool clamp); // dot13-insts + +Computes ``dot(a, b) + c``. The ``clamp`` flag, when present and true, clamps +the result to [0.0, 1.0]. + +Integer Dot Product +------------------- + +.. code-block:: c + + int __builtin_amdgcn_sdot2(short2 a, short2 b, int c, _Constant bool clamp); // dot2-insts + unsigned int __builtin_amdgcn_udot2(ushort2 a, ushort2 b, unsigned int c, _Constant bool clamp); // dot2-insts + int __builtin_amdgcn_sdot4(int a, int b, int c, _Constant bool clamp); // dot1-insts + unsigned int __builtin_amdgcn_udot4(unsigned int a, unsigned int b, unsigned int c, _Constant bool clamp); // dot7-insts + int __builtin_amdgcn_sdot8(int a, int b, int c, _Constant bool clamp); // dot1-insts + unsigned int __builtin_amdgcn_udot8(unsigned int a, unsigned int b, unsigned int c, _Constant bool clamp); // dot7-insts + +For ``sdot4``/``udot4``, each 32-bit argument is treated as four packed 8-bit +integers. For ``sdot8``/``udot8``, they are eight packed 4-bit integers. The +``clamp`` flag saturates on overflow. + +Mixed Sign Dot Product +^^^^^^^^^^^^^^^^^^^^^^ + +.. code-block:: c + + int __builtin_amdgcn_sudot4(_Constant bool a_sign, int a, _Constant bool b_sign, int b, int c, _Constant bool clamp); // dot8-insts + int __builtin_amdgcn_sudot8(_Constant bool a_sign, int a, _Constant bool b_sign, int b, int c, _Constant bool clamp); // dot8-insts + +- ``a_sign``, ``b_sign``: If true, treat the corresponding operand as signed. + If false, treat as unsigned. + +FP8/BF8 Dot Product +^^^^^^^^^^^^^^^^^^^ + +.. code-block:: c + + float __builtin_amdgcn_dot4_f32_fp8_fp8(unsigned int a, unsigned int b, float c); // dot11-insts + float __builtin_amdgcn_dot4_f32_fp8_bf8(unsigned int a, unsigned int b, float c); + float __builtin_amdgcn_dot4_f32_bf8_fp8(unsigned int a, unsigned int b, float c); + float __builtin_amdgcn_dot4_f32_bf8_bf8(unsigned int a, unsigned int b, float c); + +Each 32-bit input encodes 4 packed 8-bit floats (FP8 or BF8). + +Raytracing / BVH Builtins +========================= + +GFX10+ BVH Intersect Ray +------------------------ + +.. code-block:: c + + uint4 __builtin_amdgcn_image_bvh_intersect_ray(uint node, float ray_extent, float4 ray_origin, float4 ray_dir, float4 ray_inv_dir, uint4 texture_desc); + uint4 __builtin_amdgcn_image_bvh_intersect_ray_h(uint node, float ray_extent, float4 ray_origin, half4 ray_dir, half4 ray_inv_dir, uint4 texture_desc); + uint4 __builtin_amdgcn_image_bvh_intersect_ray_l(uint64_t node, float ray_extent, float4 ray_origin, float4 ray_dir, float4 ray_inv_dir, uint4 texture_desc); + uint4 __builtin_amdgcn_image_bvh_intersect_ray_lh(uint64_t node, float ray_extent, float4 ray_origin, half4 ray_dir, half4 ray_inv_dir, uint4 texture_desc); + +Requires ``gfx10-insts``. Naming convention: ``l`` = 64-bit node pointer; +``h`` = half-precision direction vectors. + +- ``node``: BVH node pointer. +- ``ray_extent``: Maximum ray distance. +- ``ray_origin``: Ray origin (only xyz used). +- ``ray_dir``: Ray direction (only xyz used). +- ``ray_inv_dir``: Inverse ray direction (only xyz used). +- ``texture_desc``: Image descriptor for the BVH. + +**Lowering note**: The builtin takes ``float4``/``half4`` vectors but only +the first 3 components are used. The lowering uses a shuffle to extract +xyz (vec3) before passing to the intrinsic. + +GFX12+ BVH Builtins +------------------- + +.. code-block:: c + + uint10 __builtin_amdgcn_image_bvh8_intersect_ray(uint64_t node, float ray_extent, unsigned char ray_mask, float3 ray_origin, float3 ray_dir, unsigned int offset, uint4 texture_desc, float3 *ray_origin_out, float3 *ray_dir_out); + uint10 __builtin_amdgcn_image_bvh_dual_intersect_ray(uint64_t node, float ray_extent, unsigned char ray_mask, float3 ray_origin, float3 ray_dir, uint2 offsets, uint4 texture_desc, float3 *ray_origin_out, float3 *ray_dir_out); + +Requires ``gfx12-insts``. The intrinsic returns ``{vdata, ray_origin, ray_dir}`` +as a struct. The builtin returns ``vdata`` and writes the updated ray origin +and direction through the pointer arguments. + +DS BVH Stack Builtins +--------------------- + +.. code-block:: c + + uint2 __builtin_amdgcn_ds_bvh_stack_rtn(unsigned int addr, unsigned int data, uint4 input, _Constant int controls); // gfx11-insts + uint2 __builtin_amdgcn_ds_bvh_stack_push4_pop1_rtn(unsigned int addr, unsigned int data, uint4 input, _Constant int controls); // gfx11-insts + uint2 __builtin_amdgcn_ds_bvh_stack_push8_pop1_rtn(unsigned int addr, unsigned int data, uint8 input, _Constant int controls); // gfx12-insts + uint64_t2 __builtin_amdgcn_ds_bvh_stack_push8_pop2_rtn(unsigned int addr, unsigned int data, uint8 input, _Constant int controls); // gfx12-insts + +The ``push8_pop2`` variant returns ``<2 x i64>``; the intrinsic's second +``i32`` return value is zero-extended to ``i64``. + +MFMA (Matrix Fused Multiply-Add) Builtins +========================================= + +MFMA builtins perform matrix multiply-accumulate operations. They require +``mai-insts``, ``fp8-insts``, or ``gfx950-insts`` depending on the variant. + +General Signature +----------------- + +.. code-block:: c + + vecN __builtin_amdgcn_mfma_<outtype>_MxNxK<intype>(A a, B b, vecN c, _Constant int cbsz, _Constant int abid, _Constant int blgp); + +- ``a``, ``b``: Input matrix fragments. +- ``c``: Accumulator matrix fragment. +- ``cbsz``: Control broadcast size (broadcast mode for A operand). +- ``abid``: Accumulator broadcast ID (selects which accumulator portion + to use). +- ``blgp``: B-matrix lane group pattern. + +All three control arguments (``cbsz``, ``abid``, ``blgp``) must be compile-time +constants. + +The matrix dimensions (M, N, K) and input/output types are encoded in the +builtin name. For example, ``mfma_f32_32x32x1f32`` computes a 32x32x1 matrix +multiply with f32 inputs and f32 accumulator. + +SMFMAC (Sparse Matrix FMA) +-------------------------- + +.. code-block:: c + + vecN __builtin_amdgcn_smfmac_<outtype>_MxNxK<intype>(A a, B b, vecN c, int idx, _Constant int cbsz, _Constant int abid); + +- ``a``, ``b``: Sparse input matrix fragments. +- ``c``: Accumulator. +- ``idx``: Sparse index (non-constant, VGPR). +- ``cbsz``, ``abid``: Constant control parameters. + +GFX950 MFMA Scale +----------------- + +.. code-block:: c + + vec4 __builtin_amdgcn_mfma_scale_f32_16x16x128_f8f6f4(int8 a, int8 b, vec4 c, _Constant int cbsz, _Constant int abid, _Constant int blgp, int a_scale, _Constant int a_scale_fmt, int b_scale); + vec16 __builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4(int8 a, int8 b, vec16 c, _Constant int cbsz, _Constant int abid, _Constant int blgp, int a_scale, _Constant int a_scale_fmt, int b_scale); + +Requires ``gfx950-insts``. Adds per-operand scaling factors: + +- ``a_scale``, ``b_scale``: Scale factors for A and B matrices. +- ``a_scale_fmt``: Scale format selector (constant). + +FP8 Conversion Builtins +======================= + +These builtins convert between FP8/BF8 formats and standard float types. +Require ``fp8-conversion-insts`` or related features. + +.. code-block:: c + + float __builtin_amdgcn_cvt_f32_bf8(int src, _Constant int byte_sel); // fp8-conversion-insts + float __builtin_amdgcn_cvt_f32_fp8(int src, _Constant int byte_sel); + float __builtin_amdgcn_cvt_f32_fp8_e5m3(int src, _Constant int byte_sel); // fp8e5m3-insts + +- ``src``: Packed 4x FP8/BF8 values in an ``int``. +- ``byte_sel``: Selects which byte (0-3) to convert. + +.. code-block:: c + + float2 __builtin_amdgcn_cvt_pk_f32_bf8(int src, _Constant bool word_sel); + float2 __builtin_amdgcn_cvt_pk_f32_fp8(int src, _Constant bool word_sel); + +- ``word_sel``: If false, converts bytes [0:1]; if true, converts bytes [2:3]. + +.. code-block:: c + + int __builtin_amdgcn_cvt_pk_bf8_f32(float a, float b, int old, _Constant bool word_sel); + int __builtin_amdgcn_cvt_pk_fp8_f32(float a, float b, int old, _Constant bool word_sel); + +- ``a``, ``b``: Two f32 values to pack into FP8. +- ``old``: Previous packed value (the unselected half is preserved). +- ``word_sel``: Selects which half of the result to write. + +.. code-block:: c + + int __builtin_amdgcn_cvt_sr_bf8_f32(float val, int rng, int old, _Constant int byte_sel); + int __builtin_amdgcn_cvt_sr_fp8_f32(float val, int rng, int old, _Constant int byte_sel); + +- ``val``: Input float. +- ``rng``: Random bits for stochastic rounding. +- ``old``: Previous packed result (unselected bytes preserved). +- ``byte_sel``: Which byte position (0-3) to write. + +WMMA (Wave Matrix Multiply-Accumulate) Builtins +=============================================== + +WMMA builtins perform cooperative matrix multiply operations across a wavefront. +The ``_w32`` suffix indicates wave32 requirement, ``_w64`` for wave64. + +GFX11 WMMA +---------- + +Requires ``gfx11-insts`` and the appropriate wavefront size. + +**Float/Half WMMA** (16x16x16): + +.. code-block:: c + + // Wave32 variants + float8 __builtin_amdgcn_wmma_f32_16x16x16_f16_w32(half16 a, half16 b, float8 c); + float8 __builtin_amdgcn_wmma_f32_16x16x16_bf16_w32(short16 a, short16 b, float8 c); + half16 __builtin_amdgcn_wmma_f16_16x16x16_f16_w32(half16 a, half16 b, half16 c, _Constant bool opsel); + short16 __builtin_amdgcn_wmma_bf16_16x16x16_bf16_w32(short16 a, short16 b, short16 c, _Constant bool opsel); + +- ``a``, ``b``: Input matrix fragments. +- ``c``: Accumulator. +- ``opsel``: Operation select -- controls which half of the accumulator to + write (for f16/bf16 output variants). + +The ``_tied`` variants (e.g., ``wmma_f16_16x16x16_f16_tied_w32``) have the +same signature but indicate the output is tied to the input accumulator +register. + +**Integer WMMA**: + +.. code-block:: c + + int8 __builtin_amdgcn_wmma_i32_16x16x16_iu8_w32(_Constant bool a_sign, int4 a, _Constant bool b_sign, int4 b, int8 c, _Constant bool clamp); + int8 __builtin_amdgcn_wmma_i32_16x16x16_iu4_w32(_Constant bool a_sign, int2 a, _Constant bool b_sign, int2 b, int8 c, _Constant bool clamp); + +- ``a_sign``, ``b_sign``: If true, treat operand as signed. +- ``clamp``: If true, saturate the result. + +Wave64 variants have the same argument pattern but smaller vector sizes +(e.g., ``float4`` instead of ``float8``). + +GFX12 WMMA +---------- + +Requires ``gfx12-insts``. Distinguished by ``_gfx12`` suffix. Smaller input +vectors (no A/B replication required). Adds FP8/BF8 variants: + +.. code-block:: c + + float8 __builtin_amdgcn_wmma_f32_16x16x16_fp8_fp8_w32_gfx12(int2 a, int2 b, float8 c); + // ... similar for fp8_bf8, bf8_fp8, bf8_bf8 + +GFX1250 WMMA +------------ + +Requires ``gfx1250-insts,wavefrontsize32``. Extended variants with additional +control: + +.. code-block:: c + + float8 __builtin_amdgcn_wmma_f32_16x16x4_f32(_Constant bool a_neg, float2 a, _Constant bool b_neg, float2 b, _Constant short matrix_fmt, float8 c, _Constant bool neg_lo, _Constant bool neg_hi); + +- ``a_neg``, ``b_neg``: Negate input operands. +- ``matrix_fmt``: Matrix format selector. +- ``neg_lo``, ``neg_hi``: Negate low/high halves of accumulator. + +The ``_iu8`` variant accepts an optional clamp argument (7 or 8 args): + +.. code-block:: c + + int8 __builtin_amdgcn_wmma_i32_16x16x64_iu8(_Constant bool a_sign, int8 a, _Constant bool b_sign, int8 b, int8 c, _Constant bool neg_lo, _Constant bool neg_hi[, _Constant bool clamp]); + +**Restriction**: When the optional ``clamp`` argument is present, it must be +a compile-time constant boolean. + +SWMMAC (Sparse Wave Matrix Multiply-Accumulate) +=============================================== + +GFX12 SWMMAC +------------ + +Requires ``gfx12-insts``. The ``index`` argument provides sparsity +metadata. + +.. code-block:: c + + float8 __builtin_amdgcn_swmmac_f32_16x16x32_f16_w32(half8 a, half16 b, float8 c, int index); + float8 __builtin_amdgcn_swmmac_f32_16x16x32_bf16_w32(short8 a, short16 b, float8 c, int index); + half8 __builtin_amdgcn_swmmac_f16_16x16x32_f16_w32(half8 a, half16 b, half8 c, int index); + short8 __builtin_amdgcn_swmmac_bf16_16x16x32_bf16_w32(short8 a, short16 b, short8 c, int index); + +**Integer SWMMAC**: + +.. code-block:: c + + int8 __builtin_amdgcn_swmmac_i32_16x16x32_iu8_w32(_Constant bool a_sign, int2 a, _Constant bool b_sign, int4 b, int8 c, int index, _Constant bool clamp); + int8 __builtin_amdgcn_swmmac_i32_16x16x32_iu4_w32(_Constant bool a_sign, int a, _Constant bool b_sign, int2 b, int8 c, int index, _Constant bool clamp); + +**FP8/BF8 SWMMAC**: + +.. code-block:: c + + float8 __builtin_amdgcn_swmmac_f32_16x16x32_fp8_fp8_w32(int2 a, int4 b, float8 c, int index); + // ... similar for fp8_bf8, bf8_fp8, bf8_bf8 + +GFX1250 SWMMAC +-------------- + +Requires ``gfx1250-insts,wavefrontsize32``. Extended variants with negation +control and additional ``neg_lo``, ``neg_hi`` flags. + +The ``_iu8`` variant also accepts an optional clamp argument: + +.. code-block:: c + + int8 __builtin_amdgcn_swmmac_i32_16x16x128_iu8(_Constant bool a_sign, int8 a, _Constant bool b_sign, int16 b, int8 c, int index, _Constant bool neg_lo, _Constant bool neg_hi[, _Constant bool clamp]); + +Atomic Builtins +=============== + +Global Atomics +-------------- + +.. code-block:: c + + double __builtin_amdgcn_global_atomic_fadd_f64(double __global *ptr, double val); // gfx90a-insts + float __builtin_amdgcn_global_atomic_fadd_f32(float __global *ptr, float val); // atomic-fadd-rtn-insts + half2 __builtin_amdgcn_global_atomic_fadd_v2f16(half2 __global *ptr, half2 val); // atomic-buffer-global-pk-add-f16-insts + double __builtin_amdgcn_global_atomic_fmin_f64(double __global *ptr, double val); // gfx90a-insts + double __builtin_amdgcn_global_atomic_fmax_f64(double __global *ptr, double val); // gfx90a-insts + short2 __builtin_amdgcn_global_atomic_fadd_v2bf16(short2 __global *ptr, short2 val); // atomic-global-pk-add-bf16-inst + +Flat Atomics +------------ + +.. code-block:: c + + double __builtin_amdgcn_flat_atomic_fadd_f64(double *ptr, double val); // gfx90a-insts + double __builtin_amdgcn_flat_atomic_fmin_f64(double *ptr, double val); // gfx90a-insts + double __builtin_amdgcn_flat_atomic_fmax_f64(double *ptr, double val); // gfx90a-insts + float __builtin_amdgcn_flat_atomic_fadd_f32(float *ptr, float val); // gfx940-insts + half2 __builtin_amdgcn_flat_atomic_fadd_v2f16(half2 *ptr, half2 val); // atomic-flat-pk-add-16-insts + short2 __builtin_amdgcn_flat_atomic_fadd_v2bf16(short2 *ptr, short2 val); // atomic-flat-pk-add-16-insts + +DS (LDS) Atomics +---------------- + +.. code-block:: c + + double __builtin_amdgcn_ds_atomic_fadd_f64(double __local *ptr, double val); // gfx90a-insts + float __builtin_amdgcn_ds_atomic_fadd_f32(float __local *ptr, float val); // gfx8-insts + short2 __builtin_amdgcn_ds_atomic_fadd_v2bf16(short2 __local *ptr, short2 val); // atomic-ds-pk-add-16-insts + half2 __builtin_amdgcn_ds_atomic_fadd_v2f16(half2 __local *ptr, half2 val); // atomic-ds-pk-add-16-insts + +**Lowering note**: All global, flat, and DS atomic builtins are lowered to +``AtomicRMW`` IR instructions (not intrinsic calls). + +Atomic Inc/Dec +-------------- + +.. code-block:: c + + uint32_t __builtin_amdgcn_atomic_inc32(volatile uint32_t *ptr, uint32_t val, unsigned int ordering, const char *scope); + uint64_t __builtin_amdgcn_atomic_inc64(volatile uint64_t *ptr, uint64_t val, unsigned int ordering, const char *scope); + uint32_t __builtin_amdgcn_atomic_dec32(volatile uint32_t *ptr, uint32_t val, unsigned int ordering, const char *scope); + uint64_t __builtin_amdgcn_atomic_dec64(volatile uint64_t *ptr, uint64_t val, unsigned int ordering, const char *scope); + +- ``ptr``: Pointer to the atomic variable. +- ``val``: Comparand value. +- ``ordering``: C11 atomic memory ordering constant. +- ``scope``: Synchronization scope string (e.g., ``"workgroup"``). + +**Restrictions**: ``ordering`` must be a valid C11/C++11 atomic ordering. +``scope`` must be a string literal that is a valid synchronization scope. + +Fence +===== + +.. code-block:: c + + void __builtin_amdgcn_fence(unsigned int ordering, const char *scope, ...); + +Inserts a memory fence with the given ordering and synchronization scope. + +- ``ordering``: C11 atomic ordering. ``relaxed`` and ``consume`` are not + allowed for fences. +- ``scope``: Synchronization scope string literal (e.g., ``"workgroup"``, + ``"agent"``, ``"system"``). +- Optional variadic arguments specify address spaces affected. + +**Lowering note**: Lowered to an LLVM ``fence`` instruction, not an intrinsic +call. + +Scheduling / Wait Builtins +========================== + +Scheduling Barriers +------------------- + +``__builtin_amdgcn_sched_barrier`` +~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ + +.. code-block:: c + + void __builtin_amdgcn_sched_barrier(_Constant int mask); + +Controls which instruction types may be reordered across this point: + +- ``0x0000``: No instructions may cross. +- ``0x0001``: ALU (non-memory, non-side-effect). +- ``0x0002``: VALU. +- ``0x0004``: SALU. +- ``0x0008``: MFMA/WMMA. +- ``0x0010``: All VMEM. +- ``0x0020``: VMEM reads. +- ``0x0040``: VMEM writes. +- ``0x0080``: All DS. +- ``0x0100``: DS reads. +- ``0x0200``: DS writes. + +``__builtin_amdgcn_sched_group_barrier`` +~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ + +.. code-block:: c + + void __builtin_amdgcn_sched_group_barrier(_Constant int mask, _Constant int count, _Constant int group_id); + +- ``mask``: Same encoding as ``sched_barrier``. +- ``count``: Number of instructions in this scheduling group. +- ``group_id``: Identifier for synchronizing with other sched_group_barriers. + +``__builtin_amdgcn_iglp_opt`` +~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ + +.. code-block:: c + + void __builtin_amdgcn_iglp_opt(_Constant int strategy); + +Scheduler optimization hint. ``strategy = 0`` enables small GEMM optimization. + +Wait Events +----------- + +.. code-block:: c + + void __builtin_amdgcn_s_wait_event_export_ready(); // gfx11-insts + void __builtin_amdgcn_s_wait_event(_Constant short mask); // gfx11-insts + +**``s_wait_event`` restrictions**: The ``mask`` argument has target-specific +valid bits: + +- GFX11: Bit 0 should be 0; setting bit 0 generates a warning. +- GFX12+: Bit 1 should be 0; clearing bit 1 generates a warning. +- The suggested value for ``export_ready`` is 2 on both GFX11 and GFX12+. +- Values > 2 generate a warning. + +Sleep +----- + +.. code-block:: c + + void __builtin_amdgcn_s_sleep(_Constant int duration); + void __builtin_amdgcn_s_sleep_var(unsigned int duration); // gfx12-insts + +The ``_var`` variant allows a non-constant duration. + +GFX11+ Message Return +--------------------- + +.. code-block:: c + + unsigned int __builtin_amdgcn_s_sendmsg_rtn(_Constant unsigned int msg); // gfx11-insts + uint64_t __builtin_amdgcn_s_sendmsg_rtnl(_Constant unsigned int msg); // gfx11-insts + +Send a message and return a value. + +Special Builtins +================ + +.. code-block:: c + + uint64_t __builtin_amdgcn_read_exec(); + unsigned int __builtin_amdgcn_read_exec_lo(); + unsigned int __builtin_amdgcn_read_exec_hi(); + +Read the current EXEC mask. These are **not** lowered to dedicated intrinsics +but synthesized via ``ballot(true)``. ``read_exec_lo`` returns the lower 32 +bits; ``read_exec_hi`` shifts right by 32 and truncates. + +.. code-block:: c + + void __builtin_amdgcn_endpgm(); + +Terminates the shader/kernel. Marked ``NoReturn``. + +.. code-block:: c + + uint64_t __builtin_amdgcn_get_fpenv(); + void __builtin_amdgcn_set_fpenv(uint64_t env); + +Get/set the floating-point environment (MODE register). Lowered to +``llvm.get.fpenv`` / ``llvm.set.fpenv``. + +VI+ Miscellaneous +----------------- + +.. code-block:: c + + uint64_t __builtin_amdgcn_s_memrealtime(); // requires s-memrealtime + void __builtin_amdgcn_s_dcache_wb(); // requires gfx8-insts + +GFX12+ Barrier Builtins +======================= + +Require ``gfx12-insts``. + +.. code-block:: c + + void __builtin_amdgcn_s_barrier_signal(_Constant int bar_id); + void __builtin_amdgcn_s_barrier_signal_var(void *bar, int member_count); + bool __builtin_amdgcn_s_barrier_signal_isfirst(_Constant int bar_id); + void __builtin_amdgcn_s_barrier_wait(_Constant short signal_count); + void __builtin_amdgcn_s_barrier_init(void *bar, int member_count); + void __builtin_amdgcn_s_barrier_join(void *bar); + void __builtin_amdgcn_s_barrier_leave(_Constant short flags); + unsigned int __builtin_amdgcn_s_get_barrier_state(int bar_id); + unsigned int __builtin_amdgcn_s_get_named_barrier_state(void *bar); + +- ``bar_id``: Barrier ID (constant for signal/wait, variable for get_state). +- ``bar``: Pointer to named barrier object. +- ``member_count``: Number of wavefronts participating. +- ``signal_count``: Number of signals to wait for. + +Prefetch Builtins +----------------- + +.. code-block:: c + + void __builtin_amdgcn_s_prefetch_data(const void *ptr, unsigned int length); // gfx12-insts + void __builtin_amdgcn_s_buffer_prefetch_data(__amdgpu_buffer_rsrc_t rsrc, _Constant int offset, unsigned int length); // gfx12-insts + +GFX12+ Transport Loads +---------------------- + +.. code-block:: c + + int2 __builtin_amdgcn_global_load_tr_b64_v2i32(int2 __global *ptr); // gfx12-insts,wavefrontsize32 + short8 __builtin_amdgcn_global_load_tr_b128_v8i16(short8 __global *ptr); // gfx12-insts,wavefrontsize32 + // ... also half8, bf16x8 variants + int __builtin_amdgcn_global_load_tr_b64_i32(int __global *ptr); // gfx12-insts,wavefrontsize64 + short4 __builtin_amdgcn_global_load_tr_b128_v4i16(short4 __global *ptr); // gfx12-insts,wavefrontsize64 + +Transport loads for matrix operands. Pointer must be in global address space. + +.. code-block:: c + + int __builtin_amdgcn_ds_bpermute_fi_b32(int addr, int data); // gfx12-insts + +DS backward permutation with fetch-inactive. + +GFX12+ Bitwise Ternary Operation +-------------------------------- + +.. code-block:: c + + int __builtin_amdgcn_bitop3_b32(int a, int b, int c, _Constant unsigned int lut); // bitop3-insts + short __builtin_amdgcn_bitop3_b16(short a, short b, short c, _Constant unsigned int lut); // bitop3-insts + +Performs a bitwise ternary operation using ``lut`` as a truth table. For each +bit position, the corresponding bits of ``a``, ``b``, ``c`` form a 3-bit index +into the 8-bit ``lut``. + +GFX12+ Stochastic Rounding Conversion +------------------------------------- + +.. code-block:: c + + bf16x2 __builtin_amdgcn_cvt_sr_bf16_f32(bf16x2 old, float val, unsigned int rng, _Constant bool sel); // f32-to-f16bf16-cvt-sr-insts + half2 __builtin_amdgcn_cvt_sr_f16_f32(half2 old, float val, unsigned int rng, _Constant bool sel); + +- ``old``: Previous packed result. +- ``val``: Input f32 value. +- ``rng``: Random bits for stochastic rounding. +- ``sel``: Selects which half to write. + +GFX12+ Scaled Conversion Builtins +--------------------------------- + +Large family of scaled conversion builtins for FP8/BF8/FP6/BF6/FP4 formats +with scale factors. Representative patterns: + +**Unscaled pack to FP6/BF6** (``f16bf16-to-fp6bf6-cvt-scale-insts``): + +.. code-block:: c + + uint6 __builtin_amdgcn_cvt_scalef32_pk32_fp6_f16(half32 src, float scale); + uint6 __builtin_amdgcn_cvt_scalef32_pk32_bf6_bf16(bf16x32 src, float scale); + +**Scaled element conversion** (``fp8-cvt-scale-insts`` / ``bf8-cvt-scale-insts``): + +.. code-block:: c + + half2 __builtin_amdgcn_cvt_scalef32_f16_fp8(half2 old, int src, float scale, _Constant int byte_sel, _Constant bool word_sel); + float __builtin_amdgcn_cvt_scalef32_f32_fp8(int src, float scale, _Constant int byte_sel); + short2 __builtin_amdgcn_cvt_scalef32_pk_fp8_f32(short2 old, float a, float b, float scale, _Constant bool word_sel); + float2 __builtin_amdgcn_cvt_scalef32_pk_f32_fp8(unsigned int src, float scale, _Constant bool word_sel); + +**Scaled pack FP4** (``fp4-cvt-scale-insts``): + +.. code-block:: c + + float2 __builtin_amdgcn_cvt_scalef32_pk_f32_fp4(unsigned int src, float scale, _Constant int nibble_sel); + unsigned int __builtin_amdgcn_cvt_scalef32_pk_fp4_f32(unsigned int old, float a, float b, float scale, _Constant int nibble_sel); + +**Stochastic rounding scaled** variants add a ``unsigned int rng`` argument +for random bits. + +**Scaled pack to FP6/BF6 from f32** (``fp6bf6-cvt-scale-insts``): + +.. code-block:: c + + float32 __builtin_amdgcn_cvt_scalef32_pk32_f32_fp6(uint6 src, float scale); + half32 __builtin_amdgcn_cvt_scalef32_pk32_f16_fp6(uint6 src, float scale); + +GFX1250+ Only Builtins +====================== + +Require ``gfx1250-insts`` unless otherwise noted. + +Cluster Barrier +--------------- + +.. code-block:: c + + void __builtin_amdgcn_s_cluster_barrier(); + +Barrier synchronization across all workgroups in the cluster. + +Prefetch +-------- + +.. code-block:: c + + void __builtin_amdgcn_flat_prefetch(const void *ptr, _Constant int length); // vmem-pref-insts + void __builtin_amdgcn_global_prefetch(const void __global *ptr, _Constant int length); // vmem-pref-insts + +Monitor Loads +------------- + +.. code-block:: c + + int __builtin_amdgcn_global_load_monitor_b32(int __global *ptr, _Constant int ordering, _Constant int scope); + int2 __builtin_amdgcn_global_load_monitor_b64(int2 __global *ptr, _Constant int ordering, _Constant int scope); + int4 __builtin_amdgcn_global_load_monitor_b128(int4 __global *ptr, _Constant int ordering, _Constant int scope); + int __builtin_amdgcn_flat_load_monitor_b32(int *ptr, _Constant int ordering, _Constant int scope); + int2 __builtin_amdgcn_flat_load_monitor_b64(int2 *ptr, _Constant int ordering, _Constant int scope); + int4 __builtin_amdgcn_flat_load_monitor_b128(int4 *ptr, _Constant int ordering, _Constant int scope); + +- ``ordering``: Atomic ordering. Must be valid for a load (no ``release`` + or ``acq_rel``). +- ``scope``: Synchronization scope constant. + +Cluster Loads +------------- + +Require ``mcast-load-insts,wavefrontsize32``. + +.. code-block:: c + + int __builtin_amdgcn_cluster_load_b32(int __global *ptr, _Constant int flags, int mask); + int2 __builtin_amdgcn_cluster_load_b64(int2 __global *ptr, _Constant int flags, int mask); + int4 __builtin_amdgcn_cluster_load_b128(int4 __global *ptr, _Constant int flags, int mask); + +Async LDS Load/Store (GFX1250) +------------------------------ + +.. code-block:: c + + void __builtin_amdgcn_cluster_load_async_to_lds_b{8,32,64,128}(T __global *src, T __local *dst, _Constant int flags, _Constant int offset, int mask); + void __builtin_amdgcn_global_load_async_to_lds_b{8,32,64,128}(T __global *src, T __local *dst, _Constant int flags, _Constant int offset); + void __builtin_amdgcn_global_store_async_from_lds_b{8,32,64,128}(T __global *dst, T __local *src, _Constant int flags, _Constant int offset); + +DS Atomic Barriers +------------------ + +.. code-block:: c + + void __builtin_amdgcn_ds_atomic_async_barrier_arrive_b64(long __local *ptr); + long __builtin_amdgcn_ds_atomic_barrier_arrive_rtn_b64(long __local *ptr, long val); + +Tensor Load/Store +----------------- + +.. code-block:: c + + void __builtin_amdgcn_tensor_load_to_lds(int4 coords, int8 desc, int4 slice0, int4 slice1, _Constant int flags); + void __builtin_amdgcn_tensor_load_to_lds_d2(int4 coords, int8 desc, _Constant int flags); + void __builtin_amdgcn_tensor_store_from_lds(int4 coords, int8 desc, int4 slice0, int4 slice1, _Constant int flags); + void __builtin_amdgcn_tensor_store_from_lds_d2(int4 coords, int8 desc, _Constant int flags); + +Transpose Loads (GFX1250) +------------------------- + +Various transpose load builtins for different data widths, from both global +and LDS address spaces. These are used to load matrix fragments in transposed +layout for WMMA operations. + +.. code-block:: c + + int2 __builtin_amdgcn_global_load_tr4_b64_v2i32(int2 __global *ptr); // transpose-load-f4f6-insts,wavefrontsize32 + int2 __builtin_amdgcn_global_load_tr8_b64_v2i32(int2 __global *ptr); // gfx1250-insts,wavefrontsize32 + int3 __builtin_amdgcn_global_load_tr6_b96_v3i32(int3 __global *ptr); // transpose-load-f4f6-insts,wavefrontsize32 + short8 __builtin_amdgcn_global_load_tr16_b128_v8i16(short8 __global *ptr); // gfx1250-insts,wavefrontsize32 + // ... also half8, bf16x8 variants + // DS (LDS) variants: + int2 __builtin_amdgcn_ds_load_tr4_b64_v2i32(int2 __local *ptr); + // ... etc. + +GFX950 also has DS-only transpose reads: + +.. code-block:: c + + int2 __builtin_amdgcn_ds_read_tr4_b64_v2i32(int2 __local *ptr); // gfx950-insts + int3 __builtin_amdgcn_ds_read_tr6_b96_v3i32(int3 __local *ptr); + int2 __builtin_amdgcn_ds_read_tr8_b64_v2i32(int2 __local *ptr); + short4 __builtin_amdgcn_ds_read_tr16_b64_v4i16(short4 __local *ptr); + // ... also half4, bf16x4 variants + +Miscellaneous GFX1250 Builtins +------------------------------ + +.. code-block:: c + + void __builtin_amdgcn_s_setprio_inc_wg(_Constant short prio); // setprio-inc-wg-inst + void __builtin_amdgcn_s_monitor_sleep(_Constant short duration); + void __builtin_amdgcn_s_wakeup_barrier(void *bar); // s-wakeup-barrier-inst + void __builtin_amdgcn_s_wait_asynccnt(_Constant unsigned short count); + void __builtin_amdgcn_s_wait_tensorcnt(_Constant unsigned short count); + +GFX1250 Math Builtins +--------------------- + +.. code-block:: c + + float __builtin_amdgcn_tanhf(float x); // tanh-insts + __fp16 __builtin_amdgcn_tanhh(__fp16 x); // tanh-insts + __bf16 __builtin_amdgcn_tanh_bf16(__bf16 x); // bf16-trans-insts + __bf16 __builtin_amdgcn_rcp_bf16(__bf16 x); + __bf16 __builtin_amdgcn_sqrt_bf16(__bf16 x); + __bf16 __builtin_amdgcn_rsq_bf16(__bf16 x); + __bf16 __builtin_amdgcn_log_bf16(__bf16 x); + __bf16 __builtin_amdgcn_exp2_bf16(__bf16 x); + __bf16 __builtin_amdgcn_sin_bf16(__bf16 x); + __bf16 __builtin_amdgcn_cos_bf16(__bf16 x); + +Hardware bf16 transcendental functions. All require ``bf16-trans-insts`` +(except tanh which requires ``tanh-insts``). + +GFX1250 Packed Arithmetic +------------------------- + +.. code-block:: c + + unsigned short __builtin_amdgcn_ashr_pk_i8_i32(unsigned int a, unsigned int b, unsigned int c); // ashr-pk-insts + unsigned short __builtin_amdgcn_ashr_pk_u8_i32(unsigned int a, unsigned int b, unsigned int c); + + unsigned short __builtin_amdgcn_sat_pk4_i4_i8(unsigned int src); // gfx1250-insts + unsigned short __builtin_amdgcn_sat_pk4_u4_u8(unsigned int src); + +GFX1250 Add-Min/Max +------------------- + +.. code-block:: c + + int __builtin_amdgcn_add_max_i32(int a, int b, int c, _Constant bool clamp); // add-min-max-insts + unsigned int __builtin_amdgcn_add_max_u32(unsigned int a, unsigned int b, unsigned int c, _Constant bool clamp); + int __builtin_amdgcn_add_min_i32(int a, int b, int c, _Constant bool clamp); + unsigned int __builtin_amdgcn_add_min_u32(unsigned int a, unsigned int b, unsigned int c, _Constant bool clamp); + +Computes ``max(a + b, c)`` (add_max) or ``min(a + b, c)`` (add_min). +The ``clamp`` flag saturates the result. + +Packed 16-bit variants (require ``pk-add-min-max-insts``): + +.. code-block:: c + + short2 __builtin_amdgcn_pk_add_max_i16(short2 a, short2 b, short2 c, _Constant bool clamp); + ushort2 __builtin_amdgcn_pk_add_max_u16(ushort2 a, ushort2 b, ushort2 c, _Constant bool clamp); + short2 __builtin_amdgcn_pk_add_min_i16(short2 a, short2 b, short2 c, _Constant bool clamp); + ushort2 __builtin_amdgcn_pk_add_min_u16(ushort2 a, ushort2 b, ushort2 c, _Constant bool clamp); + +GFX1250 Permutation +------------------- + +Require ``tensor-cvt-lut-insts``. + +.. code-block:: c + + uint2 __builtin_amdgcn_perm_pk16_b4_u4(unsigned int src, unsigned int idx, uint2 lut); + uint3 __builtin_amdgcn_perm_pk16_b6_u4(unsigned int src, unsigned long idx, uint2 lut); + uint4 __builtin_amdgcn_perm_pk16_b8_u4(unsigned long src, unsigned long idx, uint2 lut); + +GFX1250 Scaled Conversions +-------------------------- + +Additional ``cvt_scale_pk*`` builtins for FP8/BF8/FP6/BF6/FP4 conversions +with scaling, specific to GFX1250: + +.. code-block:: c + + half8 __builtin_amdgcn_cvt_scale_pk8_f16_fp8(uint2 src, unsigned int scale, _Constant unsigned int byte_sel); + bf16x8 __builtin_amdgcn_cvt_scale_pk8_bf16_fp8(uint2 src, unsigned int scale, _Constant unsigned int byte_sel); + // ... similar for bf8, fp4 source types, and f32 output + +**Restriction on ``byte_sel``**: Must be a compile-time constant in the +range [0, 15]. + +GFX1250 Cooperative Atomics +--------------------------- + +128-byte cooperative atomic load/store operations. Require +``gfx1250-insts,wavefrontsize32``. + +.. code-block:: c + + // 32x4B = 128 bytes, 1 int per lane, 32 lanes + int __builtin_amdgcn_cooperative_atomic_load_32x4B(int *ptr, _Constant int ordering, const char *scope); + void __builtin_amdgcn_cooperative_atomic_store_32x4B(int *ptr, int val, _Constant int ordering, const char *scope); + + // 16x8B = 128 bytes, 2 ints per lane, 16 lanes + int2 __builtin_amdgcn_cooperative_atomic_load_16x8B(int2 *ptr, _Constant int ordering, const char *scope); + void __builtin_amdgcn_cooperative_atomic_store_16x8B(int2 *ptr, int2 val, _Constant int ordering, const char *scope); + + // 8x16B = 128 bytes, 4 ints per lane, 8 lanes + int4 __builtin_amdgcn_cooperative_atomic_load_8x16B(int4 *ptr, _Constant int ordering, const char *scope); + void __builtin_amdgcn_cooperative_atomic_store_8x16B(int4 *ptr, int4 val, _Constant int ordering, const char *scope); + +**Restrictions**: + +- ``ptr`` must be in **global** or **generic** address space. +- ``ordering``: Must be a valid C11 atomic ordering. Loads cannot use + ``release`` or ``acq_rel``. Stores cannot use ``acquire`` or ``acq_rel``. +- ``scope``: Must be a **string literal** specifying a valid synchronization + scope. + +Image Builtins +============== + +Image builtins provide access to texture image load, store, and sample +operations. All require ``image-insts`` or ``extended-image-insts``. + +The builtin names encode the operation, dimensionality, return type, and +coordinate type: +``__builtin_amdgcn_image_<op>_<dim>_<rettype>_<coordtype>``. + +Image Load +---------- + +.. code-block:: c + + float4 __builtin_amdgcn_image_load_<dim>_v4f32_i32(int dmask, int coord..., __amdgpu_texture_t rsrc, int texfailctrl, int cachepolicy); + half4 __builtin_amdgcn_image_load_<dim>_v4f16_i32(int dmask, int coord..., __amdgpu_texture_t rsrc, int texfailctrl, int cachepolicy); + +Supported dimensions: ``1d``, ``1darray``, ``2d``, ``2darray``, ``3d``, +``cube``. Mip variants: ``mip_1d``, ``mip_1darray``, ``mip_2d``, +``mip_2darray``, ``mip_3d``, ``mip_cube``. Some 2D variants also have +scalar ``float``/``f32`` return. + +Arguments: + +- ``dmask``: Data mask (constant) selecting which channels to load. +- ``coord...``: Integer coordinates. Count depends on dimensionality + (1D=1, 2D=2, 3D=3, array=+1, mip=+1). +- ``rsrc``: Texture resource descriptor (``__amdgpu_texture_t``, constant). +- ``texfailctrl``: Texture fail control (constant). +- ``cachepolicy``: Cache policy (constant). + +**Restriction**: ``dmask``, ``rsrc``, ``texfailctrl``, and ``cachepolicy`` +must all be compile-time constants. + +Image Store +----------- + +.. code-block:: c + + void __builtin_amdgcn_image_store_<dim>_v4f32_i32(float4 data, int dmask, int coord..., __amdgpu_texture_t rsrc, int texfailctrl, int cachepolicy); + void __builtin_amdgcn_image_store_<dim>_v4f16_i32(half4 data, int dmask, int coord..., ...); + +Same dimensionality and mip variants as loads. The ``data`` argument precedes +the other arguments. + +**Restriction**: ``dmask``, ``rsrc``, ``texfailctrl``, and ``cachepolicy`` +must all be compile-time constants. + +Image Sample +------------ + +.. code-block:: c + + float4 __builtin_amdgcn_image_sample_<dim>_v4f32_f32(int dmask, float coord..., __amdgpu_texture_t rsrc, int4 sampler, bool unorm, int texfailctrl, int cachepolicy); + +Requires ``image-insts``. Supported dimensions: ``1d``, ``1darray``, ``2d``, +``2darray``, ``3d``, ``cube``. Also scalar float return for 2D variants. + +Additional arguments compared to image_load: + +- ``sampler``: Sampler state (``int4``). +- ``unorm``: If true, coordinates are unnormalized. + +Sample Variants +^^^^^^^^^^^^^^^ + +Require ``extended-image-insts``: + +- **``sample_lz``**: Sample at LOD zero. Same args as ``sample``. +- **``sample_l``**: Sample at explicit LOD. Adds one ``float lod`` coordinate. +- **``sample_d``**: Sample with explicit derivatives. Adds derivative + coordinates (count depends on dimensionality: 2 for 1D, 4 for 2D, 6 for + 3D). +- **``gather4_lz_2d``**: 4-sample gather at LOD zero (2D only). + +R600 Builtins +============= + +These builtins are for the older R600/Northern Islands GPU targets. + +.. code-block:: c + + unsigned char __constant * __builtin_r600_implicitarg_ptr(); + + unsigned int __builtin_r600_read_tgid_x(); + unsigned int __builtin_r600_read_tgid_y(); + unsigned int __builtin_r600_read_tgid_z(); + unsigned int __builtin_r600_read_tidig_x(); + unsigned int __builtin_r600_read_tidig_y(); + unsigned int __builtin_r600_read_tidig_z(); + + double __builtin_r600_recipsqrt_ieee(double x); + float __builtin_r600_recipsqrt_ieeef(float x); + +Miscellaneous +============= + +.. code-block:: c + + unsigned int __builtin_amdgcn_prng_b32(unsigned int src); // prng-inst + +Pseudo-random number generation. + +.. code-block:: c + + void __builtin_amdgcn_s_ttracedata_imm(_Constant short data); // gfx10-insts + +Write immediate data to thread trace. + +GFX950 Scaled Conversion Builtins +--------------------------------- + +.. code-block:: c + + uint6 __builtin_amdgcn_cvt_scalef32_2xpk16_fp6_f32(float16 a, float16 b, float scale); // gfx950-insts + uint6 __builtin_amdgcn_cvt_scalef32_2xpk16_bf6_f32(float16 a, float16 b, float scale); + +Converts two sets of 16 f32 values to packed FP6/BF6 with scale. diff --git a/clang/docs/index.rst b/clang/docs/index.rst index 5cdbd52a2398d..6f809f45f113e 100644 --- a/clang/docs/index.rst +++ b/clang/docs/index.rst @@ -64,6 +64,7 @@ Using Clang as a Compiler APINotes DebuggingCoroutines AMDGPUSupport + AMDGPUBuiltins CXXTypeAwareAllocators CommandGuide/index FAQ >From 075a657eaf504444a5bd59a81b47cc76d168d9b3 Mon Sep 17 00:00:00 2001 From: Shilei Tian <[email protected]> Date: Thu, 12 Feb 2026 13:09:39 -0500 Subject: [PATCH 2/3] add a warning --- clang/docs/AMDGPUBuiltins.rst | 5 +++++ 1 file changed, 5 insertions(+) diff --git a/clang/docs/AMDGPUBuiltins.rst b/clang/docs/AMDGPUBuiltins.rst index 9ca9a2bf3bd32..57a5c4f234b3a 100644 --- a/clang/docs/AMDGPUBuiltins.rst +++ b/clang/docs/AMDGPUBuiltins.rst @@ -12,6 +12,11 @@ and intrinsics. They are defined in ``clang/include/clang/Basic/BuiltinsAMDGPU.t and typically lower to LLVM intrinsics defined in ``llvm/include/llvm/IR/IntrinsicsAMDGPU.td``. +.. warning:: + + These builtins, including their names, arguments, and target requirements, + are all subject to change without warning across LLVM releases. + All AMDGPU builtins use the ``__builtin_amdgcn_`` prefix (or ``__builtin_r600_`` for R600 targets). Arguments marked ``_Constant`` must be compile-time constant expressions. >From 96393b81af18e771b0782166d54dde1fa012922f Mon Sep 17 00:00:00 2001 From: Shilei Tian <[email protected]> Date: Thu, 12 Feb 2026 13:50:08 -0500 Subject: [PATCH 3/3] fix doc build error --- clang/docs/AMDGPUBuiltins.rst | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/clang/docs/AMDGPUBuiltins.rst b/clang/docs/AMDGPUBuiltins.rst index 57a5c4f234b3a..8c4a8a0cfd9fe 100644 --- a/clang/docs/AMDGPUBuiltins.rst +++ b/clang/docs/AMDGPUBuiltins.rst @@ -1246,7 +1246,7 @@ Scheduling Barriers ------------------- ``__builtin_amdgcn_sched_barrier`` -~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ +^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ .. code-block:: c @@ -1267,7 +1267,7 @@ Controls which instruction types may be reordered across this point: - ``0x0200``: DS writes. ``__builtin_amdgcn_sched_group_barrier`` -~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ +^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ .. code-block:: c @@ -1278,7 +1278,7 @@ Controls which instruction types may be reordered across this point: - ``group_id``: Identifier for synchronizing with other sched_group_barriers. ``__builtin_amdgcn_iglp_opt`` -~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ +^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ .. code-block:: c _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
