From 502b3bfc6a713e5b6640faf48e72de08d7cb0aba Mon Sep 17 00:00:00 2001 From: Stanislav Mekhanoshin Date: Wed, 24 Feb 2021 15:14:30 -0800 Subject: [PATCH] [AMDGPU] require s-memtime-inst for __builtin_amdgcn_s_memtime Differential Revision: https://reviews.llvm.org/D97420 --- clang/include/clang/Basic/BuiltinsAMDGPU.def | 3 ++- clang/test/CodeGenOpenCL/builtins-amdgcn-ci.cl | 8 ++++++++ clang/test/CodeGenOpenCL/builtins-amdgcn-gfx10.cl | 8 ++++++++ clang/test/CodeGenOpenCL/builtins-amdgcn-gfx9.cl | 8 ++++++++ clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl | 7 +++++++ clang/test/CodeGenOpenCL/builtins-amdgcn.cl | 7 ------- clang/test/SemaOpenCL/builtins-amdgcn-error-gfx1030.cl | 7 +++++++ 7 files changed, 40 insertions(+), 8 deletions(-) create mode 100644 clang/test/SemaOpenCL/builtins-amdgcn-error-gfx1030.cl diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.def b/clang/include/clang/Basic/BuiltinsAMDGPU.def index 3544abe..415d8cb 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 b3bf368..56da7ec 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 a90d55f..a9bbaa9 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 344af2b..420506e 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 4408b0432..49faf06 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 20edaf2..3769149 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 0000000..3414914 --- /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}} +} -- 2.7.4