[PATCH] D30241: AMDGPU: Add fmed3 half builtin
Matt Arsenault via Phabricator via cfe-commits
cfe-commits at lists.llvm.org
Tue Feb 21 21:18:16 PST 2017
arsenm created this revision.
Herald added subscribers: tpr, dstuttard, tony-tye, yaxunl, nhaehnle, wdng, kzhuravl.
https://reviews.llvm.org/D30241
Files:
include/clang/Basic/BuiltinsAMDGPU.def
lib/Basic/Targets.cpp
lib/CodeGen/CGBuiltin.cpp
test/CodeGenOpenCL/builtins-amdgcn-gfx9.cl
test/SemaOpenCL/builtins-amdgcn-error-f16.cl
test/SemaOpenCL/builtins-amdgcn-error-gfx9.cl
Index: test/SemaOpenCL/builtins-amdgcn-error-gfx9.cl
===================================================================
--- /dev/null
+++ 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}}
+}
Index: test/SemaOpenCL/builtins-amdgcn-error-f16.cl
===================================================================
--- test/SemaOpenCL/builtins-amdgcn-error-f16.cl
+++ 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 @@
*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}}
}
Index: test/CodeGenOpenCL/builtins-amdgcn-gfx9.cl
===================================================================
--- /dev/null
+++ 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);
+}
Index: lib/CodeGen/CGBuiltin.cpp
===================================================================
--- lib/CodeGen/CGBuiltin.cpp
+++ lib/CodeGen/CGBuiltin.cpp
@@ -8445,6 +8445,7 @@
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>(
Index: lib/Basic/Targets.cpp
===================================================================
--- lib/Basic/Targets.cpp
+++ lib/Basic/Targets.cpp
@@ -2355,6 +2355,9 @@
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;
Index: include/clang/Basic/BuiltinsAMDGPU.def
===================================================================
--- include/clang/Basic/BuiltinsAMDGPU.def
+++ include/clang/Basic/BuiltinsAMDGPU.def
@@ -100,6 +100,12 @@
TARGET_BUILTIN(__builtin_amdgcn_s_memrealtime, "LUi", "n", "s-memrealtime")
//===----------------------------------------------------------------------===//
+// GFX9+ only builtins.
+//===----------------------------------------------------------------------===//
+
+TARGET_BUILTIN(__builtin_amdgcn_fmed3h, "hhhh", "nc", "gfx9-insts")
+
+//===----------------------------------------------------------------------===//
// Special builtins.
//===----------------------------------------------------------------------===//
BUILTIN(__builtin_amdgcn_read_exec, "LUi", "nc")
-------------- next part --------------
A non-text attachment was scrubbed...
Name: D30241.89320.patch
Type: text/x-patch
Size: 4269 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20170222/e58ed987/attachment.bin>
More information about the cfe-commits
mailing list