https://github.com/shiltian updated https://github.com/llvm/llvm-project/pull/181574
>From 0ef4b210aeb777b0c0ed6d07026a536f82e98f53 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 | 268 ++++++++++++++++++ 4 files changed, 370 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 e3233a0b2d96c..e69d4750aeb4c 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 9647d1cd2fae9..99b56e65dd3ea 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 1e8e8f359a218..6e8bf388a9e5a 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 = [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]>; diff --git a/clang/include/clang/Basic/BuiltinsAMDGPUDocs.td b/clang/include/clang/Basic/BuiltinsAMDGPUDocs.td new file mode 100644 index 0000000000000..74f3e4637f277 --- /dev/null +++ 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 SI+ 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. +}]; +} _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
