[PATCH] D50493: AMDGPU: Add another missing builtin

[Matt Arsenault via Phabricator via cfe-commits]

arsenm created this revision.
arsenm added reviewers: kzhuravl, b-sumner.
Herald added subscribers: t-tye, tpr, dstuttard, yaxunl, nhaehnle, wdng.

https://reviews.llvm.org/D50493

Files:
  include/clang/Basic/BuiltinsAMDGPU.def
  test/CodeGenOpenCL/builtins-amdgcn-ci.cl
  test/SemaOpenCL/builtins-amdgcn-error-ci.cl


Index: test/SemaOpenCL/builtins-amdgcn-error-ci.cl
===================================================================
--- test/SemaOpenCL/builtins-amdgcn-error-ci.cl
+++ test/SemaOpenCL/builtins-amdgcn-error-ci.cl
@@ -1,7 +1,8 @@
 // REQUIRES: amdgpu-registered-target
 // RUN: %clang_cc1 -triple amdgcn-- -target-cpu tahiti -verify -S -o - %s
 
-void test_ci_s_dcache_inv_vol()
+void test_ci_biltins()
 {
   __builtin_amdgcn_s_dcache_inv_vol(); // expected-error 
{{'__builtin_amdgcn_s_dcache_inv_vol' needs target feature ci-insts}}
+  __builtin_amdgcn_buffer_wbinvl1_vol(); // expected-error 
{{'__builtin_amdgcn_buffer_wbinvl1_vol' needs target feature ci-insts}}
 }
Index: test/CodeGenOpenCL/builtins-amdgcn-ci.cl
===================================================================
--- test/CodeGenOpenCL/builtins-amdgcn-ci.cl
+++ test/CodeGenOpenCL/builtins-amdgcn-ci.cl
@@ -10,3 +10,10 @@
   __builtin_amdgcn_s_dcache_inv_vol();
 }
 
+// CHECK-LABEL: @test_buffer_wbinvl1_vol
+// CHECK: call void @llvm.amdgcn.buffer.wbinvl1.vol()
+void test_buffer_wbinvl1_vol()
+{
+  __builtin_amdgcn_buffer_wbinvl1_vol();
+}
+
Index: include/clang/Basic/BuiltinsAMDGPU.def
===================================================================
--- include/clang/Basic/BuiltinsAMDGPU.def
+++ include/clang/Basic/BuiltinsAMDGPU.def
@@ -104,6 +104,7 @@
 // CI+ only builtins.
 
//===----------------------------------------------------------------------===//
 TARGET_BUILTIN(__builtin_amdgcn_s_dcache_inv_vol, "v", "n", "ci-insts")
+TARGET_BUILTIN(__builtin_amdgcn_buffer_wbinvl1_vol, "v", "n", "ci-insts")
 
 
//===----------------------------------------------------------------------===//
 // VI+ only builtins.



Index: test/SemaOpenCL/builtins-amdgcn-error-ci.cl
===================================================================
--- test/SemaOpenCL/builtins-amdgcn-error-ci.cl
+++ test/SemaOpenCL/builtins-amdgcn-error-ci.cl
@@ -1,7 +1,8 @@
 // REQUIRES: amdgpu-registered-target
 // RUN: %clang_cc1 -triple amdgcn-- -target-cpu tahiti -verify -S -o - %s
 
-void test_ci_s_dcache_inv_vol()
+void test_ci_biltins()
 {
   __builtin_amdgcn_s_dcache_inv_vol(); // expected-error {{'__builtin_amdgcn_s_dcache_inv_vol' needs target feature ci-insts}}
+  __builtin_amdgcn_buffer_wbinvl1_vol(); // expected-error {{'__builtin_amdgcn_buffer_wbinvl1_vol' needs target feature ci-insts}}
 }
Index: test/CodeGenOpenCL/builtins-amdgcn-ci.cl
===================================================================
--- test/CodeGenOpenCL/builtins-amdgcn-ci.cl
+++ test/CodeGenOpenCL/builtins-amdgcn-ci.cl
@@ -10,3 +10,10 @@
   __builtin_amdgcn_s_dcache_inv_vol();
 }
 
+// CHECK-LABEL: @test_buffer_wbinvl1_vol
+// CHECK: call void @llvm.amdgcn.buffer.wbinvl1.vol()
+void test_buffer_wbinvl1_vol()
+{
+  __builtin_amdgcn_buffer_wbinvl1_vol();
+}
+
Index: include/clang/Basic/BuiltinsAMDGPU.def
===================================================================
--- include/clang/Basic/BuiltinsAMDGPU.def
+++ include/clang/Basic/BuiltinsAMDGPU.def
@@ -104,6 +104,7 @@
 // CI+ only builtins.
 //===----------------------------------------------------------------------===//
 TARGET_BUILTIN(__builtin_amdgcn_s_dcache_inv_vol, "v", "n", "ci-insts")
+TARGET_BUILTIN(__builtin_amdgcn_buffer_wbinvl1_vol, "v", "n", "ci-insts")
 
 //===----------------------------------------------------------------------===//
 // VI+ only builtins.

_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits