AMDGPU: Add fmed3 half builtin

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

diff --git a/include/clang/Basic/BuiltinsAMDGPU.def b/include/clang/Basic/BuiltinsAMDGPU.def
index 2c4efe92dc87015815da833db83346fe8046ead1..48080d71aaa02d4aa47fe09053e8002732b210dc 100644 (file)
--- a/include/clang/Basic/BuiltinsAMDGPU.def
+++ b/include/clang/Basic/BuiltinsAMDGPU.def
@@ -99,6 +99,12 @@ 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")
 
+//===----------------------------------------------------------------------===//
+// GFX9+ only builtins.
+//===----------------------------------------------------------------------===//
+
+TARGET_BUILTIN(__builtin_amdgcn_fmed3h, "hhhh", "nc", "gfx9-insts")
+
 //===----------------------------------------------------------------------===//
 // Special builtins.
 //===----------------------------------------------------------------------===//

diff --git a/lib/Basic/Targets.cpp b/lib/Basic/Targets.cpp
index dd1cf45107382e67990f2ffcb5aeaa5a84098f2e..7e373187793453cb9d8c8383b3489fef267b03dd 100644 (file)
--- a/lib/Basic/Targets.cpp
+++ b/lib/Basic/Targets.cpp
@@ -2355,6 +2355,9 @@ bool AMDGPUTargetInfo::initFeatureMap(
     case GK_GFX7:
       break;
 
+    case GK_GFX9:
+      Features["gfx9-insts"] = true;
+      LLVM_FALLTHROUGH;
     case GK_GFX8:
       Features["s-memrealtime"] = true;
       Features["16-bit-insts"] = true;

diff --git a/lib/CodeGen/CGBuiltin.cpp b/lib/CodeGen/CGBuiltin.cpp
index 9f0cfa7db3b6adadabefa53c04f8b4b2664500a8..446f2a039f84d94eb6943fb8b515c20a5a29c912 100644 (file)
--- a/lib/CodeGen/CGBuiltin.cpp
+++ b/lib/CodeGen/CGBuiltin.cpp
@@ -8445,6 +8445,7 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
   case AMDGPU::BI__builtin_amdgcn_classh:
     return emitFPIntBuiltin(*this, E, Intrinsic::amdgcn_class);
   case AMDGPU::BI__builtin_amdgcn_fmed3f:
+  case AMDGPU::BI__builtin_amdgcn_fmed3h:
     return emitTernaryBuiltin(*this, E, Intrinsic::amdgcn_fmed3);
   case AMDGPU::BI__builtin_amdgcn_read_exec: {
     CallInst *CI = cast<CallInst>(

diff --git a/test/CodeGenOpenCL/builtins-amdgcn-gfx9.cl b/test/CodeGenOpenCL/builtins-amdgcn-gfx9.cl
new file mode 100644 (file)
index 0000000..333b610
--- /dev/null
+++ b/test/CodeGenOpenCL/builtins-amdgcn-gfx9.cl
@@ -0,0 +1,11 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx900 -S -emit-llvm -o - %s | FileCheck %s
+
+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+
+// CHECK-LABEL: @test_fmed3_f16
+// CHECK: call half @llvm.amdgcn.fmed3.f16(half %a, half %b, half %c)
+void test_fmed3_f16(global half* out, half a, half b, half c)
+{
+  *out = __builtin_amdgcn_fmed3h(a, b, c);
+}

diff --git a/test/SemaOpenCL/builtins-amdgcn-error-f16.cl b/test/SemaOpenCL/builtins-amdgcn-error-f16.cl
index 7fa47179cb71cca21dd165da9c07f0e702ed20f8..3487b1a5a80383ffc81de408178179679d304067 100644 (file)
--- a/test/SemaOpenCL/builtins-amdgcn-error-f16.cl
+++ b/test/SemaOpenCL/builtins-amdgcn-error-f16.cl
@@ -1,9 +1,10 @@
 // REQUIRES: amdgpu-registered-target
-// RUN: %clang_cc1 -triple amdgcn-- -target-cpu tahiti -verify -S -o - %s
+// RUN: %clang_cc1 -triple amdgcn-- -verify -S -o - %s
 
 #pragma OPENCL EXTENSION cl_khr_fp16 : enable
 
-void test_f16(global half *out, half a, half b, half c)
+__attribute__((target("arch=tahiti")))
+void test_f16_tahiti(global half *out, half a, half b, half c)
 {
   *out = __builtin_amdgcn_div_fixuph(a, b, c); // expected-error {{'__builtin_amdgcn_div_fixuph' needs target feature 16-bit-insts}}
   *out = __builtin_amdgcn_rcph(a); // expected-error {{'__builtin_amdgcn_rcph' needs target feature 16-bit-insts}}
@@ -15,4 +16,5 @@ void test_f16(global half *out, half a, half b, half c)
   *out = __builtin_amdgcn_frexp_exph(a); // expected-error {{'__builtin_amdgcn_frexp_exph' needs target feature 16-bit-insts}}
   *out = __builtin_amdgcn_fracth(a); // expected-error {{'__builtin_amdgcn_fracth' needs target feature 16-bit-insts}}
   *out = __builtin_amdgcn_classh(a, b); // expected-error {{'__builtin_amdgcn_classh' needs target feature 16-bit-insts}}
+  *out = __builtin_amdgcn_fmed3h(a, b, c); // expected-error {{'__builtin_amdgcn_fmed3h' needs target feature gfx9-insts}}
 }

diff --git a/test/SemaOpenCL/builtins-amdgcn-error-gfx9.cl b/test/SemaOpenCL/builtins-amdgcn-error-gfx9.cl
new file mode 100644 (file)
index 0000000..c9fd8ab
--- /dev/null
+++ b/test/SemaOpenCL/builtins-amdgcn-error-gfx9.cl
@@ -0,0 +1,9 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgcn-- -target-cpu fiji -verify -S -o - %s
+
+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+
+void test_gfx9_fmed3h(global half *out, half a, half b, half c)
+{
+  *out = __builtin_amdgcn_fmed3h(a, b, c); // expected-error {{'__builtin_amdgcn_fmed3h' needs target feature gfx9-insts}}
+}