[llvm-branch-commits] [clang] [Clang][AMDGPU][Docs] Add builtin documentation for AMDGPU builtins (PR #181574)
Shilei Tian via llvm-branch-commits
llvm-branch-commits at lists.llvm.org
Sun Feb 15 13:01:16 PST 2026
https://github.com/shiltian created https://github.com/llvm/llvm-project/pull/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.
>From 5d5ee4646490c28dbf79b046731841ff2d7393c1 Mon Sep 17 00:00:00 2001
From: Shilei Tian <i at tianshilei.me>
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/include/clang/Basic/BuiltinsAMDGPU.td | 144 +++++++--
.../include/clang/Basic/BuiltinsAMDGPUDocs.td | 291 ++++++++++++++++++
3 files changed, 406 insertions(+), 30 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/include/clang/Basic/BuiltinsAMDGPU.td b/clang/include/clang/Basic/BuiltinsAMDGPU.td
index 78443ac291f31..afb1083e3d31d 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,45 +28,128 @@ 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]>;
+def __builtin_amdgcn_mbcnt_hi
+ : AMDGPUBuiltin<"unsigned int(unsigned int, unsigned int)", [Const]> {
+ let Documentation = [DocABIMbcntHi];
+ let ArgNames = ["mask", "val"];
+}
+def __builtin_amdgcn_mbcnt_lo
+ : AMDGPUBuiltin<"unsigned int(unsigned int, unsigned int)", [Const]> {
+ let Documentation = [DocABIMbcntLo];
+ let ArgNames = ["mask", "val"];
+}
-def __builtin_amdgcn_s_memtime : AMDGPUBuiltin<"uint64_t()", [], "s-memtime-inst">;
+def __builtin_amdgcn_s_memtime
+ : AMDGPUBuiltin<"uint64_t()", [], "s-memtime-inst"> {
+ let Documentation = [DocABISMemtime];
+}
//===----------------------------------------------------------------------===//
// Instruction builtins.
diff --git a/clang/include/clang/Basic/BuiltinsAMDGPUDocs.td b/clang/include/clang/Basic/BuiltinsAMDGPUDocs.td
new file mode 100644
index 0000000000000..91dd71752c30c
--- /dev/null
+++ b/clang/include/clang/Basic/BuiltinsAMDGPUDocs.td
@@ -0,0 +1,291 @@
+//===--- 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.
+
+.. 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.
+}];
+}
+
+//===----------------------------------------------------------------------===//
+// 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. Requires ``gfx1250-insts``.
+}];
+}
+
+def DocABIClusterIdY : Documentation {
+ let Category = DocCatAMDGPUABI;
+ let Content = [{
+Returns the cluster ID in the Y dimension. Requires ``gfx1250-insts``.
+}];
+}
+
+def DocABIClusterIdZ : Documentation {
+ let Category = DocCatAMDGPUABI;
+ let Content = [{
+Returns the cluster ID in the Z dimension. Requires ``gfx1250-insts``.
+}];
+}
+
+def DocABIClusterWorkgroupIdX : Documentation {
+ let Category = DocCatAMDGPUABI;
+ let Content = [{
+Returns the workgroup ID within the cluster in the X dimension.
+Requires ``gfx1250-insts``.
+}];
+}
+
+def DocABIClusterWorkgroupIdY : Documentation {
+ let Category = DocCatAMDGPUABI;
+ let Content = [{
+Returns the workgroup ID within the cluster in the Y dimension.
+Requires ``gfx1250-insts``.
+}];
+}
+
+def DocABIClusterWorkgroupIdZ : Documentation {
+ let Category = DocCatAMDGPUABI;
+ let Content = [{
+Returns the workgroup ID within the cluster in the 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 DocABIClusterWorkgroupMaxIdX : Documentation {
+ let Category = DocCatAMDGPUABI;
+ let Content = [{
+Returns the maximum workgroup ID within the cluster in the X dimension.
+Requires ``gfx1250-insts``.
+}];
+}
+
+def DocABIClusterWorkgroupMaxIdY : Documentation {
+ let Category = DocCatAMDGPUABI;
+ let Content = [{
+Returns the maximum workgroup ID within the cluster in the Y dimension.
+Requires ``gfx1250-insts``.
+}];
+}
+
+def DocABIClusterWorkgroupMaxIdZ : Documentation {
+ let Category = DocCatAMDGPUABI;
+ let Content = [{
+Returns the maximum workgroup ID within the cluster in the 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 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.
+}];
+}
+
+def DocABIMbcntLo : Documentation {
+ let Category = DocCatAMDGPUABI;
+ let Content = [{
+Counts the number of bits set in ``mask`` among lanes 0--31 that are
+below the current lane, and adds ``val``. Lanes 32--63 count all 32
+bits.
+}];
+}
+
+def DocABIMbcntHi : Documentation {
+ let Category = DocCatAMDGPUABI;
+ let Content = [{
+Counts the number of bits set in ``mask`` among lanes 32--63 that are
+below the current lane, and adds ``val``. Lanes 0--31 add zero.
+
+**Typical usage**: To count the number of matching lanes below the
+current lane across a full 64-lane wavefront, call ``mbcnt_lo`` first
+and pass its result as ``val`` to ``mbcnt_hi``.
+}];
+}
+
+def DocABISMemtime : Documentation {
+ let Category = DocCatAMDGPUABI;
+ let Content = [{
+Returns the current 64-bit timestamp.
+}];
+}
More information about the llvm-branch-commits
mailing list