AMDGPU: Add builtin for s_dcache_wb

git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@339110 91177308-0d34-0410-b5e6-96231b3b80d8

diff --git a/include/clang/Basic/BuiltinsAMDGPU.def b/include/clang/Basic/BuiltinsAMDGPU.def
index 6d694ebf67cf32561d99b02025fb59f5bb856286..73204dc747c32daaacf585b1a643c7cbd0c51946 100644 (file)
--- a/include/clang/Basic/BuiltinsAMDGPU.def
+++ b/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/lib/Basic/Targets/AMDGPU.cpp b/lib/Basic/Targets/AMDGPU.cpp
index 310c08acb3a80abde4d8d13af4cf04eb6664d8e0..3f147d3d01eb034f7e370dea1a504c528a21f6e5 100644 (file)
--- a/lib/Basic/Targets/AMDGPU.cpp
+++ b/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/test/CodeGenOpenCL/amdgpu-features.cl b/test/CodeGenOpenCL/amdgpu-features.cl
index feeea9285c3f50bb4edf270184f2b8de7de9acec..5e022483286b1d608043eeab228217058bb4ad16 100644 (file)
--- a/test/CodeGenOpenCL/amdgpu-features.cl
+++ b/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/test/CodeGenOpenCL/builtins-amdgcn-vi.cl b/test/CodeGenOpenCL/builtins-amdgcn-vi.cl
index afa312cfcb1273a8a43b904e972777b1ae674105..cf2969f45746beeb1cd3988b2060e17925655d10 100644 (file)
--- a/test/CodeGenOpenCL/builtins-amdgcn-vi.cl
+++ b/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/test/SemaOpenCL/builtins-amdgcn-error-vi.cl b/test/SemaOpenCL/builtins-amdgcn-error-vi.cl
new file mode 100644 (file)
index 0000000..877a040
--- /dev/null
+++ b/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}}
+}