diff options
| -rw-r--r-- | clang/include/clang/Basic/BuiltinsAMDGPU.def | 6 | ||||
| -rw-r--r-- | clang/lib/Basic/Targets.cpp | 3 | ||||
| -rw-r--r-- | clang/lib/CodeGen/CGBuiltin.cpp | 1 | ||||
| -rw-r--r-- | clang/test/CodeGenOpenCL/builtins-amdgcn-gfx9.cl | 11 | ||||
| -rw-r--r-- | clang/test/SemaOpenCL/builtins-amdgcn-error-f16.cl | 6 | ||||
| -rw-r--r-- | clang/test/SemaOpenCL/builtins-amdgcn-error-gfx9.cl | 9 |
6 files changed, 34 insertions, 2 deletions
diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.def b/clang/include/clang/Basic/BuiltinsAMDGPU.def index 2c4efe92dc8..48080d71aaa 100644 --- a/clang/include/clang/Basic/BuiltinsAMDGPU.def +++ b/clang/include/clang/Basic/BuiltinsAMDGPU.def @@ -100,6 +100,12 @@ 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. //===----------------------------------------------------------------------===// BUILTIN(__builtin_amdgcn_read_exec, "LUi", "nc") diff --git a/clang/lib/Basic/Targets.cpp b/clang/lib/Basic/Targets.cpp index dd1cf451073..7e373187793 100644 --- a/clang/lib/Basic/Targets.cpp +++ b/clang/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/clang/lib/CodeGen/CGBuiltin.cpp b/clang/lib/CodeGen/CGBuiltin.cpp index 9f0cfa7db3b..446f2a039f8 100644 --- a/clang/lib/CodeGen/CGBuiltin.cpp +++ b/clang/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/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx9.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-gfx9.cl new file mode 100644 index 00000000000..333b610f81b --- /dev/null +++ b/clang/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/clang/test/SemaOpenCL/builtins-amdgcn-error-f16.cl b/clang/test/SemaOpenCL/builtins-amdgcn-error-f16.cl index 7fa47179cb7..3487b1a5a80 100644 --- a/clang/test/SemaOpenCL/builtins-amdgcn-error-f16.cl +++ b/clang/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/clang/test/SemaOpenCL/builtins-amdgcn-error-gfx9.cl b/clang/test/SemaOpenCL/builtins-amdgcn-error-gfx9.cl new file mode 100644 index 00000000000..c9fd8ab2cae --- /dev/null +++ b/clang/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}} +} |
