https://github.com/shiltian updated https://github.com/llvm/llvm-project/pull/181574
>From d9e6c2ec14e97dabd32f982cc9043f6a81dbe4fc Mon Sep 17 00:00:00 2001 From: Shilei Tian <[email protected]> Date: Sun, 15 Feb 2026 16:00:43 -0500 Subject: [PATCH] [Clang][AMDGPU][Docs] Add builtin documentation for AMDGPU builtins 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. --- clang/docs/CMakeLists.txt | 1 + clang/docs/index.rst | 1 + clang/include/clang/Basic/BuiltinsAMDGPU.td | 127 ++++++++++--- .../include/clang/Basic/BuiltinsAMDGPUDocs.td | 172 ++++++++++++++++++ 4 files changed, 274 insertions(+), 27 deletions(-) create mode 100644 clang/include/clang/Basic/BuiltinsAMDGPUDocs.td diff --git a/clang/docs/CMakeLists.txt b/clang/docs/CMakeLists.txt index e3233a0b2d96c9..e69d4750aeb4cb 100644 --- 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 diff --git a/clang/docs/index.rst b/clang/docs/index.rst index 9647d1cd2fae9e..99b56e65dd3ead 100644 --- 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 diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.td b/clang/include/clang/Basic/BuiltinsAMDGPU.td index 1e8e8f359a2182..2a7d7e0c419f10 100644 --- 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 = [DocABIWorkgroupIdXYZ]; +} +def __builtin_amdgcn_workgroup_id_y : AMDGPUBuiltin<"unsigned int()", [Const]> { + let Documentation = [DocABIWorkgroupIdXYZ]; +} +def __builtin_amdgcn_workgroup_id_z : AMDGPUBuiltin<"unsigned int()", [Const]> { + let Documentation = [DocABIWorkgroupIdXYZ]; +} -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 = [DocABIClusterIdXYZ]; +} +def __builtin_amdgcn_cluster_id_y + : AMDGPUBuiltin<"unsigned int()", [Const], "gfx1250-insts"> { + let Documentation = [DocABIClusterIdXYZ]; +} +def __builtin_amdgcn_cluster_id_z + : AMDGPUBuiltin<"unsigned int()", [Const], "gfx1250-insts"> { + let Documentation = [DocABIClusterIdXYZ]; +} -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 = [DocABIClusterWorkgroupIdXYZ]; +} +def __builtin_amdgcn_cluster_workgroup_id_y + : AMDGPUBuiltin<"unsigned int()", [Const], "gfx1250-insts"> { + let Documentation = [DocABIClusterWorkgroupIdXYZ]; +} +def __builtin_amdgcn_cluster_workgroup_id_z + : AMDGPUBuiltin<"unsigned int()", [Const], "gfx1250-insts"> { + let Documentation = [DocABIClusterWorkgroupIdXYZ]; +} +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 = [DocABIClusterWorkgroupMaxIdXYZ]; +} +def __builtin_amdgcn_cluster_workgroup_max_id_y + : AMDGPUBuiltin<"unsigned int()", [Const], "gfx1250-insts"> { + let Documentation = [DocABIClusterWorkgroupMaxIdXYZ]; +} +def __builtin_amdgcn_cluster_workgroup_max_id_z + : AMDGPUBuiltin<"unsigned int()", [Const], "gfx1250-insts"> { + let Documentation = [DocABIClusterWorkgroupMaxIdXYZ]; +} +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 = [DocABIWorkitemIdXYZ]; +} +def __builtin_amdgcn_workitem_id_y : AMDGPUBuiltin<"unsigned int()", [Const]> { + let Documentation = [DocABIWorkitemIdXYZ]; +} +def __builtin_amdgcn_workitem_id_z : AMDGPUBuiltin<"unsigned int()", [Const]> { + let Documentation = [DocABIWorkitemIdXYZ]; +} -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 = [DocABIWorkgroupSizeXYZ]; +} +def __builtin_amdgcn_workgroup_size_y + : AMDGPUBuiltin<"unsigned short()", [Const]> { + let Documentation = [DocABIWorkgroupSizeXYZ]; +} +def __builtin_amdgcn_workgroup_size_z + : AMDGPUBuiltin<"unsigned short()", [Const]> { + let Documentation = [DocABIWorkgroupSizeXYZ]; +} -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 = [DocABIGridSizeXYZ]; +} +def __builtin_amdgcn_grid_size_y : AMDGPUBuiltin<"unsigned int()", [Const]> { + let Documentation = [DocABIGridSizeXYZ]; +} +def __builtin_amdgcn_grid_size_z : AMDGPUBuiltin<"unsigned int()", [Const]> { + let Documentation = [DocABIGridSizeXYZ]; +} 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]>; diff --git a/clang/include/clang/Basic/BuiltinsAMDGPUDocs.td b/clang/include/clang/Basic/BuiltinsAMDGPUDocs.td new file mode 100644 index 00000000000000..6d001a5771c0dc --- /dev/null +++ b/clang/include/clang/Basic/BuiltinsAMDGPUDocs.td @@ -0,0 +1,172 @@ +//===--- 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 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) +}]; +} + +//===----------------------------------------------------------------------===// +// 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. +}]; +} + +//===----------------------------------------------------------------------===// +// 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 object for the queue executing the current kernel. +}]; +} + +def DocABIWorkgroupIdXYZ : Documentation { + let Category = DocCatAMDGPUABI; + let Content = [{ +Returns the workgroup ID in the X/Y/Z dimension. +}]; +} + +def DocABIClusterIdXYZ : Documentation { + let Category = DocCatAMDGPUABI; + let Content = [{ +Returns the cluster ID in the X/Y/Z dimension. +Requires ``gfx1250-insts``. +}]; +} + +def DocABIClusterWorkgroupIdXYZ : Documentation { + let Category = DocCatAMDGPUABI; + let Content = [{ +Returns the workgroup ID within the cluster in the X/Y/Z dimension. +Requires ``gfx1250-insts``. +}]; +} + +def DocABIClusterWorkgroupFlatId : Documentation { + let Category = DocCatAMDGPUABI; + let Content = [{ +Returns the flat (linearized) workgroup ID within the cluster. +Requires ``gfx1250-insts``. +}]; +} + +def DocABIClusterWorkgroupMaxIdXYZ : Documentation { + let Category = DocCatAMDGPUABI; + let Content = [{ +Returns the maximum workgroup ID within the cluster in the X/Y/Z dimension. +Requires ``gfx1250-insts``. +}]; +} + +def DocABIClusterWorkgroupMaxFlatId : Documentation { + let Category = DocCatAMDGPUABI; + let Content = [{ +Returns the maximum flat (linearized) workgroup ID within the cluster. +Requires ``gfx1250-insts``. +}]; +} + +def DocABIWorkitemIdXYZ : Documentation { + let Category = DocCatAMDGPUABI; + let Content = [{ +Returns the work-item (thread) ID within the workgroup in the X/Y/Z dimension. +}]; +} + +def DocABIWorkgroupSizeXYZ : Documentation { + let Category = DocCatAMDGPUABI; + let Content = [{ +Returns the workgroup size in the X/Y/Z dimension. +}]; +} + +def DocABIGridSizeXYZ : Documentation { + let Category = DocCatAMDGPUABI; + let Content = [{ +Returns the total grid size in the X/Y/Z dimension. +}]; +} _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
