dnrops
/
clang
mirror of https://github.com/microsoft/clang.git
1
0
Fork 0
Code Issues Packages Projects Releases Wiki Activity

AMDGPU: Add fmed3 half builtin 

git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@295874 91177308-0d34-0410-b5e6-96231b3b80d8
This commit is contained in:
Matt Arsenault 2017-02-22 20:55:59 +00:00
parent 9d5a4118c5
commit 521e2b23fa
6 changed files with 34 additions and 2 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

6
include/clang/Basic/BuiltinsAMDGPU.def
View File

@ -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.
//===----------------------------------------------------------------------===//

3
lib/Basic/Targets.cpp
View File

@ -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;

1
lib/CodeGen/CGBuiltin.cpp
View File

@ -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>(

11
test/CodeGenOpenCL/builtins-amdgcn-gfx9.cl Normal file
View File

@ -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);
}

6
test/SemaOpenCL/builtins-amdgcn-error-f16.cl
View File

@ -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}}
}

9
test/SemaOpenCL/builtins-amdgcn-error-gfx9.cl Normal file
View File

@ -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}}
}