[PATCH] D50320: AMDGPU: Add builtin for s_dcache_inv_vol

[Matt Arsenault via Phabricator via cfe-commits]

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

https://reviews.llvm.org/D50320

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
===================================================================
--- /dev/null
+++ test/SemaOpenCL/builtins-amdgcn-error-ci.cl
@@ -0,0 +1,7 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgcn-- -target-cpu tahiti -verify -S -o - %s
+
+void test_ci_s_dcache_inv_vol()
+{
+  __builtin_amdgcn_s_dcache_inv_vol(); // expected-error 
{{'__builtin_amdgcn_s_dcache_inv_vol' needs target feature ci-insts}}
+}
Index: test/CodeGenOpenCL/builtins-amdgcn-ci.cl
===================================================================
--- /dev/null
+++ test/CodeGenOpenCL/builtins-amdgcn-ci.cl
@@ -0,0 +1,10 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu hawaii -S 
-emit-llvm -o - %s | FileCheck %s
+
+// CHECK-LABEL: @test_s_dcache_inv_vol
+// CHECK: call void @llvm.amdgcn.s.dcache.inv.vol(
+void test_s_dcache_inv_vol()
+{
+  __builtin_amdgcn_s_dcache_inv_vol();
+}
+
Index: include/clang/Basic/BuiltinsAMDGPU.def
===================================================================
--- include/clang/Basic/BuiltinsAMDGPU.def
+++ include/clang/Basic/BuiltinsAMDGPU.def
@@ -100,6 +100,11 @@
 BUILTIN(__builtin_amdgcn_ds_fminf, "ff*3fIiIiIb", "n")
 BUILTIN(__builtin_amdgcn_ds_fmaxf, "ff*3fIiIiIb", "n")
 
+//===----------------------------------------------------------------------===//
+// CI+ only builtins.
+//===----------------------------------------------------------------------===//
+TARGET_BUILTIN(__builtin_amdgcn_s_dcache_inv_vol, "v", "n", "ci-insts")
+
 
//===----------------------------------------------------------------------===//
 // VI+ only builtins.
 
//===----------------------------------------------------------------------===//



Index: test/SemaOpenCL/builtins-amdgcn-error-ci.cl
===================================================================
--- /dev/null
+++ test/SemaOpenCL/builtins-amdgcn-error-ci.cl
@@ -0,0 +1,7 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgcn-- -target-cpu tahiti -verify -S -o - %s
+
+void test_ci_s_dcache_inv_vol()
+{
+  __builtin_amdgcn_s_dcache_inv_vol(); // expected-error {{'__builtin_amdgcn_s_dcache_inv_vol' needs target feature ci-insts}}
+}
Index: test/CodeGenOpenCL/builtins-amdgcn-ci.cl
===================================================================
--- /dev/null
+++ test/CodeGenOpenCL/builtins-amdgcn-ci.cl
@@ -0,0 +1,10 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu hawaii -S -emit-llvm -o - %s | FileCheck %s
+
+// CHECK-LABEL: @test_s_dcache_inv_vol
+// CHECK: call void @llvm.amdgcn.s.dcache.inv.vol(
+void test_s_dcache_inv_vol()
+{
+  __builtin_amdgcn_s_dcache_inv_vol();
+}
+
Index: include/clang/Basic/BuiltinsAMDGPU.def
===================================================================
--- include/clang/Basic/BuiltinsAMDGPU.def
+++ include/clang/Basic/BuiltinsAMDGPU.def
@@ -100,6 +100,11 @@
 BUILTIN(__builtin_amdgcn_ds_fminf, "ff*3fIiIiIb", "n")
 BUILTIN(__builtin_amdgcn_ds_fmaxf, "ff*3fIiIiIb", "n")
 
+//===----------------------------------------------------------------------===//
+// CI+ only builtins.
+//===----------------------------------------------------------------------===//
+TARGET_BUILTIN(__builtin_amdgcn_s_dcache_inv_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