https://github.com/CRobeck updated https://github.com/llvm/llvm-project/pull/88076
>From 1e2cab61cbf46e5cc73d7ee6523dcce1a75c7549 Mon Sep 17 00:00:00 2001 From: Corbin Robeck <corbin.rob...@amd.com> Date: Mon, 8 Apr 2024 19:58:57 -0500 Subject: [PATCH 1/4] 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>>]>; >From 8ad802d418ee9f2aea4a1ce56e3b67f750f59a4d Mon Sep 17 00:00:00 2001 From: Corbin Robeck <corbin.rob...@amd.com> Date: Mon, 8 Apr 2024 21:42:23 -0500 Subject: [PATCH 2/4] fix type issue --- clang/include/clang/Basic/BuiltinsAMDGPU.def | 4 ++-- .../CodeGenOpenCL/builtins-amdgcn-gfx12.cl | 21 +++++++++++++++++++ 2 files changed, 23 insertions(+), 2 deletions(-) diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.def b/clang/include/clang/Basic/BuiltinsAMDGPU.def index d2912d271d4005..302888adb47a17 100644 --- a/clang/include/clang/Basic/BuiltinsAMDGPU.def +++ b/clang/include/clang/Basic/BuiltinsAMDGPU.def @@ -61,8 +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_s_ttracedata, "vi", "n") +BUILTIN(__builtin_amdgcn_s_ttracedata_imm, "vIs", "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/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12.cl index ebd367bba0cdc1..61fdd9ae135033 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12.cl @@ -235,3 +235,24 @@ unsigned test_s_get_barrier_state(int a) unsigned State = __builtin_amdgcn_s_get_barrier_state(a); return State; } + +// CHECK-LABEL: @test_s_ttracedata( +// CHECK-NEXT: entry: +// CHECK-NEXT: call void @llvm.amdgcn.s.ttracedata() +// CHECK-NEXT: ret void +// +void test_s_ttracedata() +{ + __builtin_amdgcn_s_ttracedata(); +} + + +// CHECK-LABEL: @test_s_ttracedata_imm( +// CHECK-NEXT: entry: +// CHECK-NEXT: call void @llvm.amdgcn.s.ttracedata_imm() +// CHECK-NEXT: ret void +// +void test_s_ttracedata_imm() +{ + __builtin_amdgcn_s_ttracedata_imm(); +} >From af4f8e63bc58e1e4ae1d532d55909659ed080aa1 Mon Sep 17 00:00:00 2001 From: Corbin Robeck <corbin.rob...@amd.com> Date: Mon, 8 Apr 2024 21:43:06 -0500 Subject: [PATCH 3/4] add ttrace builtin test --- clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12.cl | 11 ++++++----- 1 file changed, 6 insertions(+), 5 deletions(-) diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12.cl index 61fdd9ae135033..26c0ee48306237 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12.cl @@ -238,21 +238,22 @@ unsigned test_s_get_barrier_state(int a) // CHECK-LABEL: @test_s_ttracedata( // CHECK-NEXT: entry: -// CHECK-NEXT: call void @llvm.amdgcn.s.ttracedata() +// CHECK-NEXT: call void @llvm.amdgcn.s.ttracedata(i32 1) // CHECK-NEXT: ret void // void test_s_ttracedata() { - __builtin_amdgcn_s_ttracedata(); + __builtin_amdgcn_s_ttracedata(1); } - // CHECK-LABEL: @test_s_ttracedata_imm( // CHECK-NEXT: entry: -// CHECK-NEXT: call void @llvm.amdgcn.s.ttracedata_imm() +// CHECK-NEXT: call void @llvm.amdgcn.s.ttracedata.imm(i16 1) // CHECK-NEXT: ret void // void test_s_ttracedata_imm() { - __builtin_amdgcn_s_ttracedata_imm(); + __builtin_amdgcn_s_ttracedata_imm(1); } + + >From ac48dc99a67e118162461e839df2b344c71e1d98 Mon Sep 17 00:00:00 2001 From: Corbin Robeck <corbin.rob...@amd.com> Date: Thu, 11 Apr 2024 14:21:04 -0500 Subject: [PATCH 4/4] move s_ttracedata_imm to be target bultin for gfx10+ --- clang/include/clang/Basic/BuiltinsAMDGPU.def | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.def b/clang/include/clang/Basic/BuiltinsAMDGPU.def index 302888adb47a17..3e21a2fe2ac6b3 100644 --- a/clang/include/clang/Basic/BuiltinsAMDGPU.def +++ b/clang/include/clang/Basic/BuiltinsAMDGPU.def @@ -62,7 +62,6 @@ 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, "vi", "n") -BUILTIN(__builtin_amdgcn_s_ttracedata_imm, "vIs", "n") BUILTIN(__builtin_amdgcn_wave_barrier, "v", "n") BUILTIN(__builtin_amdgcn_sched_barrier, "vIi", "n") BUILTIN(__builtin_amdgcn_sched_group_barrier, "vIiIiIi", "n") @@ -269,6 +268,7 @@ TARGET_BUILTIN(__builtin_amdgcn_dot4_f32_bf8_bf8, "fUiUif", "nc", "dot11-insts") TARGET_BUILTIN(__builtin_amdgcn_permlane16, "UiUiUiUiUiIbIb", "nc", "gfx10-insts") TARGET_BUILTIN(__builtin_amdgcn_permlanex16, "UiUiUiUiUiIbIb", "nc", "gfx10-insts") TARGET_BUILTIN(__builtin_amdgcn_mov_dpp8, "UiUiIUi", "nc", "gfx10-insts") +TARGET_BUILTIN(__builtin_amdgcn_s_ttracedata_imm, "vIs", "n", "gfx10-insts") //===----------------------------------------------------------------------===// // Raytracing builtins. _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits