r339109 - AMDGPU: Add builtin for s_dcache_inv_vol

[Matt Arsenault via cfe-commits]

Author: arsenm
Date: Tue Aug  7 00:49:04 2018
New Revision: 339109

URL: http://llvm.org/viewvc/llvm-project?rev=339109&view=rev
Log:
AMDGPU: Add builtin for s_dcache_inv_vol

Added:
    cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-ci.cl
    cfe/trunk/test/SemaOpenCL/builtins-amdgcn-error-ci.cl
Modified:
    cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def
    cfe/trunk/lib/Basic/Targets/AMDGPU.cpp
    cfe/trunk/test/CodeGenOpenCL/amdgpu-features.cl

Modified: cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def?rev=339109&r1=339108&r2=339109&view=diff
==============================================================================
--- cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def (original)
+++ cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def Tue Aug  7 00:49:04 2018
@@ -101,6 +101,11 @@ BUILTIN(__builtin_amdgcn_ds_fminf, "ff*3
 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.
 //===----------------------------------------------------------------------===//
 

Modified: cfe/trunk/lib/Basic/Targets/AMDGPU.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Basic/Targets/AMDGPU.cpp?rev=339109&r1=339108&r2=339109&view=diff
==============================================================================
--- cfe/trunk/lib/Basic/Targets/AMDGPU.cpp (original)
+++ cfe/trunk/lib/Basic/Targets/AMDGPU.cpp Tue Aug  7 00:49:04 2018
@@ -148,12 +148,14 @@ bool AMDGPUTargetInfo::initFeatureMap(
       Features["16-bit-insts"] = true;
       Features["dpp"] = true;
       Features["s-memrealtime"] = true;
-      break;
+      LLVM_FALLTHROUGH;
     case GK_GFX704:
     case GK_GFX703:
     case GK_GFX702:
     case GK_GFX701:
     case GK_GFX700:
+      Features["ci-insts"] = true;
+      LLVM_FALLTHROUGH;
     case GK_GFX601:
     case GK_GFX600:
       break;

Modified: cfe/trunk/test/CodeGenOpenCL/amdgpu-features.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCL/amdgpu-features.cl?rev=339109&r1=339108&r2=339109&view=diff
==============================================================================
--- cfe/trunk/test/CodeGenOpenCL/amdgpu-features.cl (original)
+++ cfe/trunk/test/CodeGenOpenCL/amdgpu-features.cl Tue Aug  7 00:49:04 2018
@@ -5,8 +5,16 @@
 
 // RUN: %clang_cc1 -triple amdgcn -target-cpu gfx904 -S -emit-llvm -o - %s | FileCheck --check-prefix=GFX904 %s
 // RUN: %clang_cc1 -triple amdgcn -target-cpu gfx906 -S -emit-llvm -o - %s | FileCheck --check-prefix=GFX906 %s
+// RUN: %clang_cc1 -triple amdgcn -target-cpu gfx801 -S -emit-llvm -o - %s | FileCheck --check-prefix=GFX801 %s
+// RUN: %clang_cc1 -triple amdgcn -target-cpu gfx700 -S -emit-llvm -o - %s | FileCheck --check-prefix=GFX700 %s
+// 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,+dpp,+fp32-denormals,+fp64-fp16-denormals,+gfx9-insts,+s-memrealtime"
-// GFX906: "target-features"="+16-bit-insts,+dl-insts,+dpp,+fp32-denormals,+fp64-fp16-denormals,+gfx9-insts,+s-memrealtime"
+// 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"
+// 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"
 
 kernel void test() {}

Added: cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-ci.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-ci.cl?rev=339109&view=auto
==============================================================================
--- cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-ci.cl (added)
+++ cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-ci.cl Tue Aug  7 00:49:04 2018
@@ -0,0 +1,12 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu hawaii -S -emit-llvm -o - %s | FileCheck %s
+// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu fiji -S -emit-llvm -o - %s | FileCheck %s
+// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx906 -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();
+}
+

Added: cfe/trunk/test/SemaOpenCL/builtins-amdgcn-error-ci.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaOpenCL/builtins-amdgcn-error-ci.cl?rev=339109&view=auto
==============================================================================
--- cfe/trunk/test/SemaOpenCL/builtins-amdgcn-error-ci.cl (added)
+++ cfe/trunk/test/SemaOpenCL/builtins-amdgcn-error-ci.cl Tue Aug  7 00:49:04 2018
@@ -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}}
+}