[clang] [RFC][Docs][Clang][AMDGPU] Add AMDGPU builtins documentation (PR #181193)

[Shilei Tian via cfe-commits]

================
@@ -0,0 +1,1990 @@
+===============
+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.
+
+.. note::
+
+   This document was generated with AI assistance, cross-referencing the
+   following sources:
+
+   - ``clang/include/clang/Basic/BuiltinsAMDGPU.td`` (builtin definitions)
+   - ``llvm/include/llvm/IR/IntrinsicsAMDGPU.td`` (intrinsic definitions)
+   - ``clang/lib/Sema/SemaAMDGPU.cpp`` (argument validation and constraints)
+   - ``clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp`` (lowering logic)
+   - `GPUOpen Machine-Readable ISA <https://gpuopen.com/machine-readable-isa/>`_
+     (ISA documents)
+
+.. 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). Some arguments must be compile-time constant expressions;
+this is noted in the descriptions where applicable.
+
+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 read-only pointer to the dispatch packet, which contains
+       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 queue_t object for the queue executing the
+       current kernel.
+
+Work-Item and Workgroup Identification
+--------------------------------------
+
+These builtins take no arguments and have no side effects. They 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.
+   * - ``__builtin_amdgcn_grid_size_{x,y,z}()``
+     - ``unsigned int``
+     - Total grid size in the specified dimension.
+
+**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
----------------
shiltian wrote:

Yeah, indeed, but not sure what "kind" of instructions they are.

https://github.com/llvm/llvm-project/pull/181193