From: Jan Vesely <jan.vesely@rutgers.edu>
Date: Sat, 25 Feb 2017 04:20:22 +0000 (+0000)
Subject: AMDGPU: export l1 cache invalidation intrinsics
X-Git-Url: https://granicus.if.org/sourcecode?a=commitdiff_plain;h=77ab7730c9bbf3e562b5866b472e5ed5a132ef03;p=clang

AMDGPU: export l1 cache invalidation intrinsics

Differential Revision: https://reviews.llvm.org/D30360

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

diff --git a/include/clang/Basic/BuiltinsAMDGPU.def b/include/clang/Basic/BuiltinsAMDGPU.def
index f39a24157a..d57be9a6d1 100644
--- a/include/clang/Basic/BuiltinsAMDGPU.def
+++ b/include/clang/Basic/BuiltinsAMDGPU.def
@@ -39,6 +39,8 @@ BUILTIN(__builtin_amdgcn_s_getreg, "UiIi", "n")
 BUILTIN(__builtin_amdgcn_s_waitcnt, "vIi", "n")
 BUILTIN(__builtin_amdgcn_s_barrier, "v", "n")
 BUILTIN(__builtin_amdgcn_wave_barrier, "v", "n")
+BUILTIN(__builtin_amdgcn_s_dcache_inv, "v", "n")
+BUILTIN(__builtin_amdgcn_buffer_wbinvl1, "v", "n")
 BUILTIN(__builtin_amdgcn_div_scale, "dddbb*", "n")
 BUILTIN(__builtin_amdgcn_div_scalef, "fffbb*", "n")
 BUILTIN(__builtin_amdgcn_div_fmas, "ddddb", "nc")
diff --git a/test/CodeGenOpenCL/builtins-amdgcn.cl b/test/CodeGenOpenCL/builtins-amdgcn.cl
index e9e765b521..7b4ef08ce3 100644
--- a/test/CodeGenOpenCL/builtins-amdgcn.cl
+++ b/test/CodeGenOpenCL/builtins-amdgcn.cl
@@ -263,6 +263,20 @@ void test_class_f64(global double* out, double a, int b)
   *out = __builtin_amdgcn_class(a, b);
 }
 
+// CHECK-LABEL: @test_buffer_wbinvl1
+// CHECK: call void @llvm.amdgcn.buffer.wbinvl1(
+void test_buffer_wbinvl1()
+{
+  __builtin_amdgcn_buffer_wbinvl1();
+}
+
+// CHECK-LABEL: @test_s_dcache_inv
+// CHECK: call void @llvm.amdgcn.s.dcache.inv(
+void test_s_dcache_inv()
+{
+  __builtin_amdgcn_s_dcache_inv();
+}
+
 // CHECK-LABEL: @test_s_waitcnt
 // CHECK: call void @llvm.amdgcn.s.waitcnt(
 void test_s_waitcnt()