[clang] [AMDGPU] Programmatically port old `.def` clang builtins to `.td` (PR #175873)
Joseph Huber via cfe-commits
cfe-commits at lists.llvm.org
Wed Jan 14 07:51:05 PST 2026
================
@@ -0,0 +1,1108 @@
+//===--- BuiltinsAMDGPU.td - AMDGPU Builtin function defs -------*- 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 the AMDGPU-specific builtin function database.
+//
+//===----------------------------------------------------------------------===//
+
+include "clang/Basic/BuiltinsBase.td"
+
+//===----------------------------------------------------------------------===//
+// AMDGPU builtin base classes
+//===----------------------------------------------------------------------===//
+
+class OtherBuiltin<string prototype, list<Attribute> Attr = []> : TargetBuiltin {
+ let Spellings = ["__builtin_"#NAME];
+ let Prototype = prototype;
+ let Attributes = !listconcat([NoThrow], Attr);
+}
+
+class AMDGPUBuiltin<string prototype, list<Attribute> Attr = []> : TargetBuiltin {
+ let Spellings = ["__builtin_amdgcn_"#NAME];
+ let Prototype = prototype;
+ let Attributes = !listconcat([NoThrow], Attr);
+}
+
+//===----------------------------------------------------------------------===//
+// SI+ only builtins.
+//===----------------------------------------------------------------------===//
+
+def dispatch_ptr : AMDGPUBuiltin<"void address_space<4>*()", [Const]>;
+def kernarg_segment_ptr : AMDGPUBuiltin<"void address_space<4>*()", [Const]>;
+def implicitarg_ptr : AMDGPUBuiltin<"void address_space<4>*()", [Const]>;
+def queue_ptr : AMDGPUBuiltin<"void address_space<4>*()", [Const]>;
+
+let Features = "gfx1250-insts" in {
+ def cluster_id_x : AMDGPUBuiltin<"unsigned int()", [Const]>;
+ def cluster_id_y : AMDGPUBuiltin<"unsigned int()", [Const]>;
+ def cluster_id_z : AMDGPUBuiltin<"unsigned int()", [Const]>;
+
+ def cluster_workgroup_id_x : AMDGPUBuiltin<"unsigned int()", [Const]>;
+ def cluster_workgroup_id_y : AMDGPUBuiltin<"unsigned int()", [Const]>;
+ def cluster_workgroup_id_z : AMDGPUBuiltin<"unsigned int()", [Const]>;
+ def cluster_workgroup_flat_id : AMDGPUBuiltin<"unsigned int()", [Const]>;
+
+ def cluster_workgroup_max_id_x : AMDGPUBuiltin<"unsigned int()", [Const]>;
+ def cluster_workgroup_max_id_y : AMDGPUBuiltin<"unsigned int()", [Const]>;
+ def cluster_workgroup_max_id_z : AMDGPUBuiltin<"unsigned int()", [Const]>;
+ def cluster_workgroup_max_flat_id : AMDGPUBuiltin<"unsigned int()", [Const]>;
+}
+
+def workgroup_id_x : AMDGPUBuiltin<"unsigned int()", [Const]>;
+def workgroup_id_y : AMDGPUBuiltin<"unsigned int()", [Const]>;
+def workgroup_id_z : AMDGPUBuiltin<"unsigned int()", [Const]>;
+
+def workitem_id_x : AMDGPUBuiltin<"unsigned int()", [Const]>;
+def workitem_id_y : AMDGPUBuiltin<"unsigned int()", [Const]>;
+def workitem_id_z : AMDGPUBuiltin<"unsigned int()", [Const]>;
+
+def workgroup_size_x : AMDGPUBuiltin<"unsigned short()", [Const]>;
+def workgroup_size_y : AMDGPUBuiltin<"unsigned short()", [Const]>;
+def workgroup_size_z : AMDGPUBuiltin<"unsigned short()", [Const]>;
+
+def grid_size_x : AMDGPUBuiltin<"unsigned int()", [Const]>;
+def grid_size_y : AMDGPUBuiltin<"unsigned int()", [Const]>;
+def grid_size_z : AMDGPUBuiltin<"unsigned int()", [Const]>;
+
+def mbcnt_hi : AMDGPUBuiltin<"unsigned int(unsigned int, unsigned int)", [Const]>;
+def mbcnt_lo : AMDGPUBuiltin<"unsigned int(unsigned int, unsigned int)", [Const]>;
+
+let Features = "s-memtime-inst" in { def s_memtime : AMDGPUBuiltin<"uint64_t()">; }
+
+//===----------------------------------------------------------------------===//
+// Instruction builtins.
+//===----------------------------------------------------------------------===//
+def s_getreg : AMDGPUBuiltin<"unsigned int(_Constant int)">;
+def s_setreg : AMDGPUBuiltin<"void(_Constant int, unsigned int)">;
+def s_getpc : AMDGPUBuiltin<"uint64_t()">;
+def s_waitcnt : AMDGPUBuiltin<"void(_Constant int)">;
+def s_sendmsg : AMDGPUBuiltin<"void(_Constant int, unsigned int)">;
+def s_sendmsghalt : AMDGPUBuiltin<"void(_Constant int, unsigned int)">;
+def s_barrier : AMDGPUBuiltin<"void()">;
+def s_ttracedata : AMDGPUBuiltin<"void(int)">;
+def wave_barrier : AMDGPUBuiltin<"void()">;
+def sched_barrier : AMDGPUBuiltin<"void(_Constant int)">;
+def sched_group_barrier : AMDGPUBuiltin<"void(_Constant int, _Constant int, _Constant int)">;
+def iglp_opt : AMDGPUBuiltin<"void(_Constant int)">;
+def s_dcache_inv : AMDGPUBuiltin<"void()">;
+def buffer_wbinvl1 : AMDGPUBuiltin<"void()">;
+def fence : AMDGPUBuiltin<"void(unsigned int, char const *, ...)">;
+def groupstaticsize : AMDGPUBuiltin<"unsigned int()">;
+def wavefrontsize : AMDGPUBuiltin<"unsigned int()", [Const]>;
+
+def atomic_inc32 : AMDGPUBuiltin<"uint32_t(uint32_t volatile *, uint32_t, unsigned int, char const *)">;
+def atomic_inc64 : AMDGPUBuiltin<"uint64_t(uint64_t volatile *, uint64_t, unsigned int, char const *)">;
+
+def atomic_dec32 : AMDGPUBuiltin<"uint32_t(uint32_t volatile *, uint32_t, unsigned int, char const *)">;
+def atomic_dec64 : AMDGPUBuiltin<"uint64_t(uint64_t volatile *, uint64_t, unsigned int, char const *)">;
+
+// FIXME: Need to disallow constant address space.
+def div_scale : AMDGPUBuiltin<"double(double, double, bool, bool *)">;
+def div_scalef : AMDGPUBuiltin<"float(float, float, bool, bool *)">;
+def div_fmas : AMDGPUBuiltin<"double(double, double, double, bool)", [Const]>;
+def div_fmasf : AMDGPUBuiltin<"float(float, float, float, bool)", [Const]>;
+def div_fixup : AMDGPUBuiltin<"double(double, double, double)", [Const]>;
+def div_fixupf : AMDGPUBuiltin<"float(float, float, float)", [Const]>;
+def trig_preop : AMDGPUBuiltin<"double(double, int)", [Const]>;
+def trig_preopf : AMDGPUBuiltin<"float(float, int)", [Const]>;
+def rcp : AMDGPUBuiltin<"double(double)", [Const]>;
+def rcpf : AMDGPUBuiltin<"float(float)", [Const]>;
+def sqrt : AMDGPUBuiltin<"double(double)", [Const]>;
+def sqrtf : AMDGPUBuiltin<"float(float)", [Const]>;
+def rsq : AMDGPUBuiltin<"double(double)", [Const]>;
+def rsqf : AMDGPUBuiltin<"float(float)", [Const]>;
+def rsq_clamp : AMDGPUBuiltin<"double(double)", [Const]>;
+def rsq_clampf : AMDGPUBuiltin<"float(float)", [Const]>;
+def sinf : AMDGPUBuiltin<"float(float)", [Const]>;
+def cosf : AMDGPUBuiltin<"float(float)", [Const]>;
+def logf : AMDGPUBuiltin<"float(float)", [Const]>;
+def exp2f : AMDGPUBuiltin<"float(float)", [Const]>;
+def log_clampf : AMDGPUBuiltin<"float(float)", [Const]>;
+def ldexp : AMDGPUBuiltin<"double(double, int)", [Const]>;
+def ldexpf : AMDGPUBuiltin<"float(float, int)", [Const]>;
+def frexp_mant : AMDGPUBuiltin<"double(double)", [Const]>;
+def frexp_mantf : AMDGPUBuiltin<"float(float)", [Const]>;
+def frexp_exp : AMDGPUBuiltin<"int(double)", [Const]>;
+def frexp_expf : AMDGPUBuiltin<"int(float)", [Const]>;
+def fract : AMDGPUBuiltin<"double(double)", [Const]>;
+def fractf : AMDGPUBuiltin<"float(float)", [Const]>;
+def amdgcn_class : OtherBuiltin<"bool(double, int)", [Const]>;
+def classf : AMDGPUBuiltin<"bool(float, int)", [Const]>;
+def s_sleep : AMDGPUBuiltin<"void(_Constant int)">;
+def s_incperflevel : AMDGPUBuiltin<"void(_Constant int)">;
+def s_decperflevel : AMDGPUBuiltin<"void(_Constant int)">;
+def s_setprio : AMDGPUBuiltin<"void(_Constant short)">;
+def ds_swizzle : AMDGPUBuiltin<"int(int, _Constant int)", [Const]>;
+def ds_permute : AMDGPUBuiltin<"int(int, int)", [Const]>;
+def ds_bpermute : AMDGPUBuiltin<"int(int, int)", [Const]>;
+def readfirstlane : AMDGPUBuiltin<"int(int)", [Const]>;
+def readlane : AMDGPUBuiltin<"int(int, int)", [Const]>;
+def fmed3f : AMDGPUBuiltin<"float(float, float, float)", [Const]>;
+def ds_faddf : AMDGPUBuiltin<"float(float address_space<3>*, float, _Constant int, _Constant int, _Constant bool)">;
+def ds_fminf : AMDGPUBuiltin<"float(float address_space<3>*, float, _Constant int, _Constant int, _Constant bool)">;
----------------
jhuber6 wrote:
Look at `BUILTIN(__builtin_amdgcn_mqsad_u32_u8, "V4UiWUiUiV4Ui", "nc")` for example, the documentation states:
```
// V -> Vector, followed by the number of elements and the base type.
// E -> ext_vector, followed by the number of elements and the base type.
```
https://github.com/llvm/llvm-project/pull/175873
More information about the cfe-commits
mailing list