[clang] 502b3bf - [AMDGPU] require s-memtime-inst for __builtin_amdgcn_s_memtime
Stanislav Mekhanoshin via cfe-commits
cfe-commits at lists.llvm.org
Thu Feb 25 08:32:15 PST 2021
Author: Stanislav Mekhanoshin
Date: 2021-02-25T08:31:59-08:00
New Revision: 502b3bfc6a713e5b6640faf48e72de08d7cb0aba
URL: https://github.com/llvm/llvm-project/commit/502b3bfc6a713e5b6640faf48e72de08d7cb0aba
DIFF: https://github.com/llvm/llvm-project/commit/502b3bfc6a713e5b6640faf48e72de08d7cb0aba.diff
LOG: [AMDGPU] require s-memtime-inst for __builtin_amdgcn_s_memtime
Differential Revision: https://reviews.llvm.org/D97420
Added:
clang/test/SemaOpenCL/builtins-amdgcn-error-gfx1030.cl
Modified:
clang/include/clang/Basic/BuiltinsAMDGPU.def
clang/test/CodeGenOpenCL/builtins-amdgcn-ci.cl
clang/test/CodeGenOpenCL/builtins-amdgcn-gfx10.cl
clang/test/CodeGenOpenCL/builtins-amdgcn-gfx9.cl
clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl
clang/test/CodeGenOpenCL/builtins-amdgcn.cl
Removed:
################################################################################
diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.def b/clang/include/clang/Basic/BuiltinsAMDGPU.def
index 3544abe35896..415d8cb3e73a 100644
--- a/clang/include/clang/Basic/BuiltinsAMDGPU.def
+++ b/clang/include/clang/Basic/BuiltinsAMDGPU.def
@@ -44,6 +44,8 @@ BUILTIN(__builtin_amdgcn_grid_size_z, "Ui", "nc")
BUILTIN(__builtin_amdgcn_mbcnt_hi, "UiUiUi", "nc")
BUILTIN(__builtin_amdgcn_mbcnt_lo, "UiUiUi", "nc")
+TARGET_BUILTIN(__builtin_amdgcn_s_memtime, "LUi", "n", "s-memtime-inst")
+
//===----------------------------------------------------------------------===//
// Instruction builtins.
//===----------------------------------------------------------------------===//
@@ -105,7 +107,6 @@ BUILTIN(__builtin_amdgcn_cubeid, "ffff", "nc")
BUILTIN(__builtin_amdgcn_cubesc, "ffff", "nc")
BUILTIN(__builtin_amdgcn_cubetc, "ffff", "nc")
BUILTIN(__builtin_amdgcn_cubema, "ffff", "nc")
-BUILTIN(__builtin_amdgcn_s_memtime, "LUi", "n")
BUILTIN(__builtin_amdgcn_s_sleep, "vIi", "n")
BUILTIN(__builtin_amdgcn_s_incperflevel, "vIi", "n")
BUILTIN(__builtin_amdgcn_s_decperflevel, "vIi", "n")
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-ci.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-ci.cl
index b3bf3681b935..56da7eceb6de 100644
--- a/clang/test/CodeGenOpenCL/builtins-amdgcn-ci.cl
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-ci.cl
@@ -5,6 +5,7 @@
// RUN: %clang_cc1 -cl-std=CL2.0 -O0 -triple amdgcn-unknown-unknown -target-cpu gfx1010 -S -emit-llvm -o - %s | FileCheck %s
typedef unsigned int uint;
+typedef unsigned long ulong;
// CHECK-LABEL: @test_s_dcache_inv_vol
// CHECK: call void @llvm.amdgcn.s.dcache.inv.vol(
@@ -27,6 +28,13 @@ void test_gws_sema_release_all(uint id)
__builtin_amdgcn_ds_gws_sema_release_all(id);
}
+// CHECK-LABEL: @test_s_memtime
+// CHECK: call i64 @llvm.amdgcn.s.memtime()
+void test_s_memtime(global ulong* out)
+{
+ *out = __builtin_amdgcn_s_memtime();
+}
+
// CHECK-LABEL: @test_is_shared(
// CHECK: [[CAST:%[0-9]+]] = bitcast i32* %{{[0-9]+}} to i8*
// CHECK: call i1 @llvm.amdgcn.is.shared(i8* [[CAST]]
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx10.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx10.cl
index a90d55fa8a78..a9bbaa9c3f54 100644
--- a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx10.cl
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx10.cl
@@ -4,6 +4,7 @@
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1012 -S -emit-llvm -o - %s | FileCheck %s
typedef unsigned int uint;
+typedef unsigned long ulong;
// CHECK-LABEL: @test_permlane16(
// CHECK: call i32 @llvm.amdgcn.permlane16(i32 %a, i32 %b, i32 %c, i32 %d, i1 false, i1 false)
@@ -22,3 +23,10 @@ void test_permlanex16(global uint* out, uint a, uint b, uint c, uint d) {
void test_mov_dpp8(global uint* out, uint a) {
*out = __builtin_amdgcn_mov_dpp8(a, 1);
}
+
+// CHECK-LABEL: @test_s_memtime
+// CHECK: call i64 @llvm.amdgcn.s.memtime()
+void test_s_memtime(global ulong* out)
+{
+ *out = __builtin_amdgcn_s_memtime();
+}
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx9.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx9.cl
index 344af2b27670..420506ec083c 100644
--- a/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx9.cl
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx9.cl
@@ -3,6 +3,7 @@
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1010 -S -emit-llvm -o - %s | FileCheck %s
#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+typedef unsigned long ulong;
// CHECK-LABEL: @test_fmed3_f16
// CHECK: call half @llvm.amdgcn.fmed3.f16(half %a, half %b, half %c)
@@ -10,3 +11,10 @@ void test_fmed3_f16(global half* out, half a, half b, half c)
{
*out = __builtin_amdgcn_fmed3h(a, b, c);
}
+
+// CHECK-LABEL: @test_s_memtime
+// CHECK: call i64 @llvm.amdgcn.s.memtime()
+void test_s_memtime(global ulong* out)
+{
+ *out = __builtin_amdgcn_s_memtime();
+}
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl
index 4408b043296a..49faf069c8c6 100644
--- a/clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl
@@ -130,3 +130,10 @@ void test_ds_fminf(local float *out, float src) {
void test_ds_fmaxf(local float *out, float src) {
*out = __builtin_amdgcn_ds_fmaxf(out, src, 0, 0, false);
}
+
+// CHECK-LABEL: @test_s_memtime
+// CHECK: call i64 @llvm.amdgcn.s.memtime()
+void test_s_memtime(global ulong* out)
+{
+ *out = __builtin_amdgcn_s_memtime();
+}
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl
index 20edaf2aae3f..3769149c8c6d 100644
--- a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl
@@ -396,13 +396,6 @@ void test_wave_barrier()
__builtin_amdgcn_wave_barrier();
}
-// CHECK-LABEL: @test_s_memtime
-// CHECK: call i64 @llvm.amdgcn.s.memtime()
-void test_s_memtime(global ulong* out)
-{
- *out = __builtin_amdgcn_s_memtime();
-}
-
// CHECK-LABEL: @test_s_sleep
// CHECK: call void @llvm.amdgcn.s.sleep(i32 1)
// CHECK: call void @llvm.amdgcn.s.sleep(i32 15)
diff --git a/clang/test/SemaOpenCL/builtins-amdgcn-error-gfx1030.cl b/clang/test/SemaOpenCL/builtins-amdgcn-error-gfx1030.cl
new file mode 100644
index 000000000000..34149148a5f3
--- /dev/null
+++ b/clang/test/SemaOpenCL/builtins-amdgcn-error-gfx1030.cl
@@ -0,0 +1,7 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgcn-- -target-cpu gfx1030 -verify -S -o - %s
+
+void test_gfx1030_s_memtime()
+{
+ __builtin_amdgcn_s_memtime(); // expected-error {{'__builtin_amdgcn_s_memtime' needs target feature s-memtime-inst}}
+}
More information about the cfe-commits
mailing list