https://github.com/jayfoad created
https://github.com/llvm/llvm-project/pull/77926
None
>From 3d4b8547514f2315130599230e769a8c73be01c3 Mon Sep 17 00:00:00 2001
From: Jay Foad <jay.f...@amd.com>
Date: Fri, 12 Jan 2024 12:43:16 +0000
Subject: [PATCH] [AMDGPU] Add GFX12 __builtin_amdgcn_s_sleep_var
---
clang/include/clang/Basic/BuiltinsAMDGPU.def | 1 +
clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12.cl | 15 +++++++++++++++
2 files changed, 16 insertions(+)
diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.def
b/clang/include/clang/Basic/BuiltinsAMDGPU.def
index e562ef04a30194..d0c4b664bf0313 100644
--- a/clang/include/clang/Basic/BuiltinsAMDGPU.def
+++ b/clang/include/clang/Basic/BuiltinsAMDGPU.def
@@ -410,6 +410,7 @@ TARGET_BUILTIN(__builtin_amdgcn_cvt_sr_fp8_f32, "ifiiIi",
"nc", "fp8-insts")
// GFX12+ only builtins.
//===----------------------------------------------------------------------===//
+TARGET_BUILTIN(__builtin_amdgcn_s_sleep_var, "vUi", "n", "gfx12-insts")
TARGET_BUILTIN(__builtin_amdgcn_permlane16_var, "UiUiUiUiIbIb", "nc",
"gfx12-insts")
TARGET_BUILTIN(__builtin_amdgcn_permlanex16_var, "UiUiUiUiIbIb", "nc",
"gfx12-insts")
TARGET_BUILTIN(__builtin_amdgcn_s_barrier_signal, "vIi", "n", "gfx12-insts")
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12.cl
b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12.cl
index 2899d9e5c28898..ebd367bba0cdc1 100644
--- a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12.cl
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12.cl
@@ -5,6 +5,21 @@
typedef unsigned int uint;
+// CHECK-LABEL: @test_s_sleep_var(
+// CHECK-NEXT: entry:
+// CHECK-NEXT: [[D_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
+// CHECK-NEXT: store i32 [[D:%.*]], ptr addrspace(5) [[D_ADDR]], align 4
+// CHECK-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(5) [[D_ADDR]], align 4
+// CHECK-NEXT: call void @llvm.amdgcn.s.sleep.var(i32 [[TMP0]])
+// CHECK-NEXT: call void @llvm.amdgcn.s.sleep.var(i32 15)
+// CHECK-NEXT: ret void
+//
+void test_s_sleep_var(int d)
+{
+ __builtin_amdgcn_s_sleep_var(d);
+ __builtin_amdgcn_s_sleep_var(15);
+}
+
// CHECK-LABEL: @test_permlane16_var(
// CHECK-NEXT: entry:
// CHECK-NEXT: [[OUT_ADDR:%.*]] = alloca ptr addrspace(1), align 8,
addrspace(5)
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
