[Clang][AMDGPU][Docs] Add builtin documentation for AMDGPU builtins (#181574)

Use the documentation generation infrastructure to document the AMDGPU builtins. This PR starts with the ABI / Special Register builtins. Documentation for the remaining builtin categories will be added incrementally in follow-up patches.
2026-02-23 13:35:35 -05:00 · 2026-02-23 13:35:35 -05:00 · e6f30334de
commit e6f30334de
parent 77b31b90d7
4 changed files with 370 additions and 27 deletions
--- a/clang/docs/CMakeLists.txt
+++ b/clang/docs/CMakeLists.txt
@ -132,6 +132,7 @@ if (LLVM_ENABLE_SPHINX)
    # Generated files
    gen_rst_file_from_td(AttributeReference.rst -gen-attr-docs ../include/clang/Basic/Attr.td "${docs_targets}")
    gen_rst_file_from_td(DiagnosticsReference.rst -gen-diag-docs ../include/clang/Basic/Diagnostic.td "${docs_targets}")
+    gen_rst_file_from_td(AMDGPUBuiltinReference.rst -gen-builtin-docs ../include/clang/Basic/BuiltinsAMDGPU.td "${docs_targets}")
    gen_rst_file_from_td(ClangCommandLineReference.rst -gen-opt-docs ../include/clang/Options/ClangOptionDocs.td "${docs_targets}")

    # Another generated file from a different source
--- a/clang/docs/index.rst
+++ b/clang/docs/index.rst
@ -22,6 +22,7 @@ Using Clang as a Compiler
   ClangCommandLineReference
   AttributeReference
   DiagnosticsReference
+   AMDGPUBuiltinReference
   WarningSuppressionMappings
   CrossCompilation
   ClangStaticAnalyzer
--- a/clang/include/clang/Basic/BuiltinsAMDGPU.td
+++ b/clang/include/clang/Basic/BuiltinsAMDGPU.td
@ -11,6 +11,7 @@
 //===----------------------------------------------------------------------===//

 include "clang/Basic/BuiltinsBase.td"
+include "clang/Basic/BuiltinsAMDGPUDocs.td"

 //===----------------------------------------------------------------------===//
 // AMDGPU builtin base classes
@ -27,40 +28,112 @@ class AMDGPUBuiltin<string prototype, list<Attribute> Attr = [], string Feat = "
 // SI+ only builtins.
 //===----------------------------------------------------------------------===//

-def __builtin_amdgcn_dispatch_ptr : AMDGPUBuiltin<"void address_space<4> *()", [Const]>;
-def __builtin_amdgcn_kernarg_segment_ptr : AMDGPUBuiltin<"void address_space<4> *()", [Const]>;
-def __builtin_amdgcn_implicitarg_ptr : AMDGPUBuiltin<"void address_space<4> *()", [Const]>;
-def __builtin_amdgcn_queue_ptr : AMDGPUBuiltin<"void address_space<4> *()", [Const]>;
+def __builtin_amdgcn_dispatch_ptr
+    : AMDGPUBuiltin<"void address_space<4> *()", [Const]> {
+  let Documentation = [DocABIDispatchPtr];
+}
+def __builtin_amdgcn_kernarg_segment_ptr
+    : AMDGPUBuiltin<"void address_space<4> *()", [Const]> {
+  let Documentation = [DocABIKernargSegmentPtr];
+}
+def __builtin_amdgcn_implicitarg_ptr
+    : AMDGPUBuiltin<"void address_space<4> *()", [Const]> {
+  let Documentation = [DocABIImplicitargPtr];
+}
+def __builtin_amdgcn_queue_ptr
+    : AMDGPUBuiltin<"void address_space<4> *()", [Const]> {
+  let Documentation = [DocABIQueuePtr];
+}

-def __builtin_amdgcn_workgroup_id_x : AMDGPUBuiltin<"unsigned int()", [Const]>;
-def __builtin_amdgcn_workgroup_id_y : AMDGPUBuiltin<"unsigned int()", [Const]>;
-def __builtin_amdgcn_workgroup_id_z : AMDGPUBuiltin<"unsigned int()", [Const]>;
+def __builtin_amdgcn_workgroup_id_x : AMDGPUBuiltin<"unsigned int()", [Const]> {
+  let Documentation = [DocABIWorkgroupIdX];
+}
+def __builtin_amdgcn_workgroup_id_y : AMDGPUBuiltin<"unsigned int()", [Const]> {
+  let Documentation = [DocABIWorkgroupIdY];
+}
+def __builtin_amdgcn_workgroup_id_z : AMDGPUBuiltin<"unsigned int()", [Const]> {
+  let Documentation = [DocABIWorkgroupIdZ];
+}

-def __builtin_amdgcn_cluster_id_x : AMDGPUBuiltin<"unsigned int()", [Const], "gfx1250-insts">;
-def __builtin_amdgcn_cluster_id_y : AMDGPUBuiltin<"unsigned int()", [Const], "gfx1250-insts">;
-def __builtin_amdgcn_cluster_id_z : AMDGPUBuiltin<"unsigned int()", [Const], "gfx1250-insts">;
+def __builtin_amdgcn_cluster_id_x
+    : AMDGPUBuiltin<"unsigned int()", [Const], "gfx1250-insts"> {
+  let Documentation = [DocABIClusterIdX];
+}
+def __builtin_amdgcn_cluster_id_y
+    : AMDGPUBuiltin<"unsigned int()", [Const], "gfx1250-insts"> {
+  let Documentation = [DocABIClusterIdY];
+}
+def __builtin_amdgcn_cluster_id_z
+    : AMDGPUBuiltin<"unsigned int()", [Const], "gfx1250-insts"> {
+  let Documentation = [DocABIClusterIdZ];
+}

-def __builtin_amdgcn_cluster_workgroup_id_x : AMDGPUBuiltin<"unsigned int()", [Const], "gfx1250-insts">;
-def __builtin_amdgcn_cluster_workgroup_id_y : AMDGPUBuiltin<"unsigned int()", [Const], "gfx1250-insts">;
-def __builtin_amdgcn_cluster_workgroup_id_z : AMDGPUBuiltin<"unsigned int()", [Const], "gfx1250-insts">;
-def __builtin_amdgcn_cluster_workgroup_flat_id : AMDGPUBuiltin<"unsigned int()", [Const], "gfx1250-insts">;
+def __builtin_amdgcn_cluster_workgroup_id_x
+    : AMDGPUBuiltin<"unsigned int()", [Const], "gfx1250-insts"> {
+  let Documentation = [DocABIClusterWorkgroupIdX];
+}
+def __builtin_amdgcn_cluster_workgroup_id_y
+    : AMDGPUBuiltin<"unsigned int()", [Const], "gfx1250-insts"> {
+  let Documentation = [DocABIClusterWorkgroupIdY];
+}
+def __builtin_amdgcn_cluster_workgroup_id_z
+    : AMDGPUBuiltin<"unsigned int()", [Const], "gfx1250-insts"> {
+  let Documentation = [DocABIClusterWorkgroupIdZ];
+}
+def __builtin_amdgcn_cluster_workgroup_flat_id
+    : AMDGPUBuiltin<"unsigned int()", [Const], "gfx1250-insts"> {
+  let Documentation = [DocABIClusterWorkgroupFlatId];
+}

-def __builtin_amdgcn_cluster_workgroup_max_id_x : AMDGPUBuiltin<"unsigned int()", [Const], "gfx1250-insts">;
-def __builtin_amdgcn_cluster_workgroup_max_id_y : AMDGPUBuiltin<"unsigned int()", [Const], "gfx1250-insts">;
-def __builtin_amdgcn_cluster_workgroup_max_id_z : AMDGPUBuiltin<"unsigned int()", [Const], "gfx1250-insts">;
-def __builtin_amdgcn_cluster_workgroup_max_flat_id : AMDGPUBuiltin<"unsigned int()", [Const], "gfx1250-insts">;
+def __builtin_amdgcn_cluster_workgroup_max_id_x
+    : AMDGPUBuiltin<"unsigned int()", [Const], "gfx1250-insts"> {
+  let Documentation = [DocABIClusterWorkgroupMaxIdX];
+}
+def __builtin_amdgcn_cluster_workgroup_max_id_y
+    : AMDGPUBuiltin<"unsigned int()", [Const], "gfx1250-insts"> {
+  let Documentation = [DocABIClusterWorkgroupMaxIdY];
+}
+def __builtin_amdgcn_cluster_workgroup_max_id_z
+    : AMDGPUBuiltin<"unsigned int()", [Const], "gfx1250-insts"> {
+  let Documentation = [DocABIClusterWorkgroupMaxIdZ];
+}
+def __builtin_amdgcn_cluster_workgroup_max_flat_id
+    : AMDGPUBuiltin<"unsigned int()", [Const], "gfx1250-insts"> {
+  let Documentation = [DocABIClusterWorkgroupMaxFlatId];
+}

-def __builtin_amdgcn_workitem_id_x : AMDGPUBuiltin<"unsigned int()", [Const]>;
-def __builtin_amdgcn_workitem_id_y : AMDGPUBuiltin<"unsigned int()", [Const]>;
-def __builtin_amdgcn_workitem_id_z : AMDGPUBuiltin<"unsigned int()", [Const]>;
+def __builtin_amdgcn_workitem_id_x : AMDGPUBuiltin<"unsigned int()", [Const]> {
+  let Documentation = [DocABIWorkitemIdX];
+}
+def __builtin_amdgcn_workitem_id_y : AMDGPUBuiltin<"unsigned int()", [Const]> {
+  let Documentation = [DocABIWorkitemIdY];
+}
+def __builtin_amdgcn_workitem_id_z : AMDGPUBuiltin<"unsigned int()", [Const]> {
+  let Documentation = [DocABIWorkitemIdZ];
+}

-def __builtin_amdgcn_workgroup_size_x : AMDGPUBuiltin<"unsigned short()", [Const]>;
-def __builtin_amdgcn_workgroup_size_y : AMDGPUBuiltin<"unsigned short()", [Const]>;
-def __builtin_amdgcn_workgroup_size_z : AMDGPUBuiltin<"unsigned short()", [Const]>;
+def __builtin_amdgcn_workgroup_size_x
+    : AMDGPUBuiltin<"unsigned short()", [Const]> {
+  let Documentation = [DocABIWorkgroupSizeX];
+}
+def __builtin_amdgcn_workgroup_size_y
+    : AMDGPUBuiltin<"unsigned short()", [Const]> {
+  let Documentation = [DocABIWorkgroupSizeY];
+}
+def __builtin_amdgcn_workgroup_size_z
+    : AMDGPUBuiltin<"unsigned short()", [Const]> {
+  let Documentation = [DocABIWorkgroupSizeZ];
+}

-def __builtin_amdgcn_grid_size_x : AMDGPUBuiltin<"unsigned int()", [Const]>;
-def __builtin_amdgcn_grid_size_y : AMDGPUBuiltin<"unsigned int()", [Const]>;
-def __builtin_amdgcn_grid_size_z : AMDGPUBuiltin<"unsigned int()", [Const]>;
+def __builtin_amdgcn_grid_size_x : AMDGPUBuiltin<"unsigned int()", [Const]> {
+  let Documentation = [DocABIGridSizeX];
+}
+def __builtin_amdgcn_grid_size_y : AMDGPUBuiltin<"unsigned int()", [Const]> {
+  let Documentation = [DocABIGridSizeY];
+}
+def __builtin_amdgcn_grid_size_z : AMDGPUBuiltin<"unsigned int()", [Const]> {
+  let Documentation = [DocABIGridSizeZ];
+}

 def __builtin_amdgcn_mbcnt_hi : AMDGPUBuiltin<"unsigned int(unsigned int, unsigned int)", [Const]>;
 def __builtin_amdgcn_mbcnt_lo : AMDGPUBuiltin<"unsigned int(unsigned int, unsigned int)", [Const]>;
--- a/clang/include/clang/Basic/BuiltinsAMDGPUDocs.td
+++ b/clang/include/clang/Basic/BuiltinsAMDGPUDocs.td
@ -0,0 +1,268 @@
+//===--- BuiltinsAMDGPUDocs.td - AMDGPU Builtin Documentation ---*- C++ -*-===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+//
+// This file defines documentation records for AMDGPU builtins. It is included
+// by BuiltinsAMDGPU.td and used by the -gen-builtin-docs TableGen backend to
+// generate AMDGPUBuiltinReference.rst.
+//
+//===----------------------------------------------------------------------===//
+
+//===----------------------------------------------------------------------===//
+// Global introduction
+//===----------------------------------------------------------------------===//
+
+def GlobalDocumentation {
+  code Intro = [{..
+  -------------------------------------------------------------------
+  NOTE: This file is automatically generated by running clang-tblgen
+  -gen-builtin-docs. Do not edit this file by hand!!
+  -------------------------------------------------------------------
+
+===============
+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.
+
+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.
+
+.. warning::
+
+   These builtins, including their names, arguments, and target requirements,
+   are all subject to change without warning across LLVM releases.
+
+.. note::
+
+   This document is a work in progress. Not all builtins are fully documented
+   yet. The initial descriptions were 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)
+}];
+}
+
+//===----------------------------------------------------------------------===//
+// Documentation categories
+//===----------------------------------------------------------------------===//
+
+def DocCatAMDGPUABI : DocumentationCategory<"ABI / Special Register Builtins"> {
+  let Content = [{
+These builtins provide access to kernel dispatch metadata, work-item and
+workgroup identification, and other ABI-level information. They are available
+on all AMDGCN targets unless otherwise noted.
+}];
+}
+
+//===----------------------------------------------------------------------===//
+// ABI / Special Register Builtins — Documentation records
+//===----------------------------------------------------------------------===//
+
+def DocABIDispatchPtr : Documentation {
+  let Category = DocCatAMDGPUABI;
+  let Content = [{
+Returns a read-only pointer to the dispatch packet, which contains
+workgroup size, grid size, and other dispatch parameters.
+}];
+}
+
+def DocABIKernargSegmentPtr : Documentation {
+  let Category = DocCatAMDGPUABI;
+  let Content = [{
+Returns a pointer to the beginning of the kernel argument segment.
+}];
+}
+
+def DocABIImplicitargPtr : Documentation {
+  let Category = DocCatAMDGPUABI;
+  let Content = [{
+Returns a pointer to the implicit arguments appended after explicit
+kernel arguments. Layout depends on the code object version.
+}];
+}
+
+def DocABIQueuePtr : Documentation {
+  let Category = DocCatAMDGPUABI;
+  let Content = [{
+Returns a pointer to the ``queue_t`` object for the queue executing the
+current kernel.
+}];
+}
+
+def DocABIWorkgroupIdX : Documentation {
+  let Category = DocCatAMDGPUABI;
+  let Content = [{
+Returns the workgroup ID in the X dimension.
+}];
+}
+
+def DocABIWorkgroupIdY : Documentation {
+  let Category = DocCatAMDGPUABI;
+  let Content = [{
+Returns the workgroup ID in the Y dimension.
+}];
+}
+
+def DocABIWorkgroupIdZ : Documentation {
+  let Category = DocCatAMDGPUABI;
+  let Content = [{
+Returns the workgroup ID in the Z dimension.
+}];
+}
+
+def DocABIClusterIdX : Documentation {
+  let Category = DocCatAMDGPUABI;
+  let Content = [{
+Returns the cluster ID in the X dimension.
+}];
+}
+
+def DocABIClusterIdY : Documentation {
+  let Category = DocCatAMDGPUABI;
+  let Content = [{
+Returns the cluster ID in the Y dimension.
+}];
+}
+
+def DocABIClusterIdZ : Documentation {
+  let Category = DocCatAMDGPUABI;
+  let Content = [{
+Returns the cluster ID in the Z dimension.
+}];
+}
+
+def DocABIClusterWorkgroupIdX : Documentation {
+  let Category = DocCatAMDGPUABI;
+  let Content = [{
+Returns the workgroup ID within the cluster in the X dimension.
+}];
+}
+
+def DocABIClusterWorkgroupIdY : Documentation {
+  let Category = DocCatAMDGPUABI;
+  let Content = [{
+Returns the workgroup ID within the cluster in the Y dimension.
+}];
+}
+
+def DocABIClusterWorkgroupIdZ : Documentation {
+  let Category = DocCatAMDGPUABI;
+  let Content = [{
+Returns the workgroup ID within the cluster in the Z dimension.
+}];
+}
+
+def DocABIClusterWorkgroupFlatId : Documentation {
+  let Category = DocCatAMDGPUABI;
+  let Content = [{
+Returns the flat (linearized) workgroup ID within the cluster.
+}];
+}
+
+def DocABIClusterWorkgroupMaxIdX : Documentation {
+  let Category = DocCatAMDGPUABI;
+  let Content = [{
+Returns the maximum workgroup ID within the cluster in the X dimension.
+}];
+}
+
+def DocABIClusterWorkgroupMaxIdY : Documentation {
+  let Category = DocCatAMDGPUABI;
+  let Content = [{
+Returns the maximum workgroup ID within the cluster in the Y dimension.
+}];
+}
+
+def DocABIClusterWorkgroupMaxIdZ : Documentation {
+  let Category = DocCatAMDGPUABI;
+  let Content = [{
+Returns the maximum workgroup ID within the cluster in the Z dimension.
+}];
+}
+
+def DocABIClusterWorkgroupMaxFlatId : Documentation {
+  let Category = DocCatAMDGPUABI;
+  let Content = [{
+Returns the maximum flat (linearized) workgroup ID within the cluster.
+}];
+}
+
+def DocABIWorkitemIdX : Documentation {
+  let Category = DocCatAMDGPUABI;
+  let Content = [{
+Returns the work-item (thread) ID within the workgroup in the X dimension.
+}];
+}
+
+def DocABIWorkitemIdY : Documentation {
+  let Category = DocCatAMDGPUABI;
+  let Content = [{
+Returns the work-item (thread) ID within the workgroup in the Y dimension.
+}];
+}
+
+def DocABIWorkitemIdZ : Documentation {
+  let Category = DocCatAMDGPUABI;
+  let Content = [{
+Returns the work-item (thread) ID within the workgroup in the Z dimension.
+}];
+}
+
+def DocABIWorkgroupSizeX : Documentation {
+  let Category = DocCatAMDGPUABI;
+  let Content = [{
+Returns the workgroup size in the X dimension.
+}];
+}
+
+def DocABIWorkgroupSizeY : Documentation {
+  let Category = DocCatAMDGPUABI;
+  let Content = [{
+Returns the workgroup size in the Y dimension.
+}];
+}
+
+def DocABIWorkgroupSizeZ : Documentation {
+  let Category = DocCatAMDGPUABI;
+  let Content = [{
+Returns the workgroup size in the Z dimension.
+}];
+}
+
+def DocABIGridSizeX : Documentation {
+  let Category = DocCatAMDGPUABI;
+  let Content = [{
+Returns the total grid size in the X dimension.
+}];
+}
+
+def DocABIGridSizeY : Documentation {
+  let Category = DocCatAMDGPUABI;
+  let Content = [{
+Returns the total grid size in the Y dimension.
+}];
+}
+
+def DocABIGridSizeZ : Documentation {
+  let Category = DocCatAMDGPUABI;
+  let Content = [{
+Returns the total grid size in the Z dimension.
+}];
+}