From 31c895ecdfadab302ec1465e4511285637a4e6e0 Mon Sep 17 00:00:00 2001 From: Matt Arsenault Date: Tue, 7 Aug 2018 07:49:13 +0000 Subject: [PATCH] AMDGPU: Add builtin for s_dcache_wb llvm-svn: 339110 --- clang/include/clang/Basic/BuiltinsAMDGPU.def | 1 + clang/lib/Basic/Targets/AMDGPU.cpp | 1 + clang/test/CodeGenOpenCL/amdgpu-features.cl | 6 +++--- clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl | 7 +++++++ clang/test/SemaOpenCL/builtins-amdgcn-error-vi.cl | 8 ++++++++ 5 files changed, 20 insertions(+), 3 deletions(-) create mode 100644 clang/test/SemaOpenCL/builtins-amdgcn-error-vi.cl diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.def b/clang/include/clang/Basic/BuiltinsAMDGPU.def index 6d694ebf67cf..73204dc747c3 100644 --- a/clang/include/clang/Basic/BuiltinsAMDGPU.def +++ b/clang/include/clang/Basic/BuiltinsAMDGPU.def @@ -121,6 +121,7 @@ TARGET_BUILTIN(__builtin_amdgcn_fracth, "hh", "nc", "16-bit-insts") TARGET_BUILTIN(__builtin_amdgcn_classh, "bhi", "nc", "16-bit-insts") TARGET_BUILTIN(__builtin_amdgcn_s_memrealtime, "LUi", "n", "s-memrealtime") TARGET_BUILTIN(__builtin_amdgcn_mov_dpp, "iiIiIiIiIb", "nc", "dpp") +TARGET_BUILTIN(__builtin_amdgcn_s_dcache_wb, "v", "n", "vi-insts") //===----------------------------------------------------------------------===// // GFX9+ only builtins. diff --git a/clang/lib/Basic/Targets/AMDGPU.cpp b/clang/lib/Basic/Targets/AMDGPU.cpp index 310c08acb3a8..3f147d3d01eb 100644 --- a/clang/lib/Basic/Targets/AMDGPU.cpp +++ b/clang/lib/Basic/Targets/AMDGPU.cpp @@ -145,6 +145,7 @@ bool AMDGPUTargetInfo::initFeatureMap( case GK_GFX803: case GK_GFX802: case GK_GFX801: + Features["vi-insts"] = true; Features["16-bit-insts"] = true; Features["dpp"] = true; Features["s-memrealtime"] = true; diff --git a/clang/test/CodeGenOpenCL/amdgpu-features.cl b/clang/test/CodeGenOpenCL/amdgpu-features.cl index feeea9285c3f..5e022483286b 100644 --- a/clang/test/CodeGenOpenCL/amdgpu-features.cl +++ b/clang/test/CodeGenOpenCL/amdgpu-features.cl @@ -10,9 +10,9 @@ // RUN: %clang_cc1 -triple amdgcn -target-cpu gfx600 -S -emit-llvm -o - %s | FileCheck --check-prefix=GFX600 %s // RUN: %clang_cc1 -triple amdgcn -target-cpu gfx601 -S -emit-llvm -o - %s | FileCheck --check-prefix=GFX601 %s -// GFX904: "target-features"="+16-bit-insts,+ci-insts,+dpp,+fp32-denormals,+fp64-fp16-denormals,+gfx9-insts,+s-memrealtime" -// GFX906: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dpp,+fp32-denormals,+fp64-fp16-denormals,+gfx9-insts,+s-memrealtime" -// GFX801: "target-features"="+16-bit-insts,+ci-insts,+dpp,+fp32-denormals,+fp64-fp16-denormals,+s-memrealtime" +// GFX904: "target-features"="+16-bit-insts,+ci-insts,+dpp,+fp32-denormals,+fp64-fp16-denormals,+gfx9-insts,+s-memrealtime,+vi-insts" +// GFX906: "target-features"="+16-bit-insts,+ci-insts,+dl-insts,+dpp,+fp32-denormals,+fp64-fp16-denormals,+gfx9-insts,+s-memrealtime,+vi-insts" +// GFX801: "target-features"="+16-bit-insts,+ci-insts,+dpp,+fp32-denormals,+fp64-fp16-denormals,+s-memrealtime,+vi-insts" // GFX700: "target-features"="+ci-insts,+fp64-fp16-denormals,-fp32-denormals" // GFX600: "target-features"="+fp32-denormals,+fp64-fp16-denormals" // GFX601: "target-features"="+fp64-fp16-denormals,-fp32-denormals" diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl index afa312cfcb12..cf2969f45746 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl @@ -82,6 +82,13 @@ void test_s_memrealtime(global ulong* out) *out = __builtin_amdgcn_s_memrealtime(); } +// CHECK-LABEL: @test_s_dcache_wb() +// CHECK: call void @llvm.amdgcn.s.dcache.wb() +void test_s_dcache_wb() +{ + __builtin_amdgcn_s_dcache_wb(); +} + // CHECK-LABEL: @test_mov_dpp // CHECK: call i32 @llvm.amdgcn.mov.dpp.i32(i32 %src, i32 0, i32 0, i32 0, i1 false) void test_mov_dpp(global int* out, int src) diff --git a/clang/test/SemaOpenCL/builtins-amdgcn-error-vi.cl b/clang/test/SemaOpenCL/builtins-amdgcn-error-vi.cl new file mode 100644 index 000000000000..877a04015e56 --- /dev/null +++ b/clang/test/SemaOpenCL/builtins-amdgcn-error-vi.cl @@ -0,0 +1,8 @@ +// REQUIRES: amdgpu-registered-target +// RUN: %clang_cc1 -triple amdgcn-- -target-cpu tahiti -verify -S -o - %s +// RUN: %clang_cc1 -triple amdgcn-- -target-cpu hawaii -verify -S -o - %s + +void test_vi_s_dcache_wb() +{ + __builtin_amdgcn_s_dcache_wb(); // expected-error {{'__builtin_amdgcn_s_dcache_wb' needs target feature vi-insts}} +} -- 2.34.1