[clang] 4c65787 - [AMDGPU] Add GFX12 __builtin_amdgcn_s_sleep_var (#77926)

[via cfe-commits]

Author: Jay Foad
Date: 2024-01-18T10:14:01Z
New Revision: 4c65787f1e45199713f71f63817651ff2decd96c

URL: https://github.com/llvm/llvm-project/commit/4c65787f1e45199713f71f63817651ff2decd96c
DIFF: https://github.com/llvm/llvm-project/commit/4c65787f1e45199713f71f63817651ff2decd96c.diff

LOG: [AMDGPU] Add GFX12 __builtin_amdgcn_s_sleep_var (#77926)

Added: 
    clang/test/CodeGenOpenCL/builtins-amdgcn-gfx11-param-err.cl

Modified: 
    clang/include/clang/Basic/BuiltinsAMDGPU.def
    clang/test/CodeGenOpenCL/builtins-amdgcn-gfx11-err.cl
    clang/test/CodeGenOpenCL/builtins-amdgcn-gfx12.cl

Removed: 
    


################################################################################
diff  --git a/clang/include/clang/Basic/BuiltinsAMDGPU.def b/clang/include/clang/Basic/BuiltinsAMDGPU.def
index f02b4d321328fe..f80d182ec08908 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-conversion-
 // 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-gfx11-err.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx11-err.cl
index 00ecf32d949238..622e9dd2eed42f 100644
--- a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx11-err.cl
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx11-err.cl
@@ -2,10 +2,7 @@
 
 // RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1100 -verify -S -emit-llvm -o - %s
 
-typedef unsigned int uint;
-typedef uint uint2 __attribute__((ext_vector_type(2)));
-typedef uint uint4 __attribute__((ext_vector_type(4)));
-
-kernel void builtins_amdgcn_bvh_err(global uint2* out, uint addr, uint data, uint4 data1, uint offset) {
-  *out = __builtin_amdgcn_ds_bvh_stack_rtn(addr, data, data1, offset); // expected-error {{'__builtin_amdgcn_ds_bvh_stack_rtn' must be a constant integer}}
+void test_s_sleep_var(int d)
+{
+  __builtin_amdgcn_s_sleep_var(d); // expected-error {{'__builtin_amdgcn_s_sleep_var' needs target feature gfx12-insts}}
 }

diff  --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx11-param-err.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx11-param-err.cl
new file mode 100644
index 00000000000000..00ecf32d949238
--- /dev/null
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx11-param-err.cl
@@ -0,0 +1,11 @@
+// REQUIRES: amdgpu-registered-target
+
+// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1100 -verify -S -emit-llvm -o - %s
+
+typedef unsigned int uint;
+typedef uint uint2 __attribute__((ext_vector_type(2)));
+typedef uint uint4 __attribute__((ext_vector_type(4)));
+
+kernel void builtins_amdgcn_bvh_err(global uint2* out, uint addr, uint data, uint4 data1, uint offset) {
+  *out = __builtin_amdgcn_ds_bvh_stack_rtn(addr, data, data1, offset); // expected-error {{'__builtin_amdgcn_ds_bvh_stack_rtn' must be a constant integer}}
+}

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)