[clang] [llvm] [AMDGPU] Add cCang builtins for amdgcn s_ttrace intrinsics (PR #88076)

[Corbin Robeck via llvm-commits]

https://github.com/CRobeck created https://github.com/llvm/llvm-project/pull/88076

None

>From 1e2cab61cbf46e5cc73d7ee6523dcce1a75c7549 Mon Sep 17 00:00:00 2001
From: Corbin Robeck <corbin.robeck at amd.com>
Date: Mon, 8 Apr 2024 19:58:57 -0500
Subject: [PATCH] add clang builtins for amdgcn s_ttrace intrinsics

---
 clang/include/clang/Basic/BuiltinsAMDGPU.def | 2 ++
 llvm/include/llvm/IR/IntrinsicsAMDGPU.td     | 3 +++
 2 files changed, 5 insertions(+)

diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.def b/clang/include/clang/Basic/BuiltinsAMDGPU.def
index c660582cc98e66..d2912d271d4005 100644
--- a/clang/include/clang/Basic/BuiltinsAMDGPU.def
+++ b/clang/include/clang/Basic/BuiltinsAMDGPU.def
@@ -61,6 +61,8 @@ BUILTIN(__builtin_amdgcn_s_waitcnt, "vIi", "n")
 BUILTIN(__builtin_amdgcn_s_sendmsg, "vIiUi", "n")
 BUILTIN(__builtin_amdgcn_s_sendmsghalt, "vIiUi", "n")
 BUILTIN(__builtin_amdgcn_s_barrier, "v", "n")
+BUILTIN(__builtin_amdgcn_s_ttracedata, "v", "n")
+BUILTIN(__builtin_amdgcn_s_ttracedata_imm, "v", "n")
 BUILTIN(__builtin_amdgcn_wave_barrier, "v", "n")
 BUILTIN(__builtin_amdgcn_sched_barrier, "vIi", "n")
 BUILTIN(__builtin_amdgcn_sched_group_barrier, "vIiIiIi", "n")
diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
index 6bbc13f1de86e2..ee9a5d7a343980 100644
--- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
+++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
@@ -1887,9 +1887,12 @@ def int_amdgcn_s_setprio :
                                 IntrHasSideEffects]>;
 
 def int_amdgcn_s_ttracedata :
+  ClangBuiltin<"__builtin_amdgcn_s_ttracedata">,
   DefaultAttrsIntrinsic<[], [llvm_i32_ty],
                         [IntrNoMem, IntrHasSideEffects]>;
+
 def int_amdgcn_s_ttracedata_imm :
+  ClangBuiltin<"__builtin_amdgcn_s_ttracedata_imm">,
   DefaultAttrsIntrinsic<[], [llvm_i16_ty],
                         [IntrNoMem, IntrHasSideEffects, ImmArg<ArgIndex<0>>]>;