r286741 - [AMDGPU] Add f16 builtin functions (VI+)
Konstantin Zhuravlyov via cfe-commits
cfe-commits at lists.llvm.org
Sat Nov 12 18:37:06 PST 2016
Author: kzhuravl
Date: Sat Nov 12 20:37:05 2016
New Revision: 286741
URL: http://llvm.org/viewvc/llvm-project?rev=286741&view=rev
Log:
[AMDGPU] Add f16 builtin functions (VI+)
Differential Revision: https://reviews.llvm.org/D26476
Added:
cfe/trunk/test/SemaOpenCL/builtins-amdgcn-error-f16.cl
cfe/trunk/test/SemaOpenCL/builtins-amdgcn-error.cl
Removed:
cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-error.cl
Modified:
cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def
cfe/trunk/lib/CodeGen/CGBuiltin.cpp
cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-vi.cl
Modified: cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def?rev=286741&r1=286740&r2=286741&view=diff
==============================================================================
--- cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def (original)
+++ cfe/trunk/include/clang/Basic/BuiltinsAMDGPU.def Sat Nov 12 20:37:05 2016
@@ -84,6 +84,16 @@ BUILTIN(__builtin_amdgcn_ds_swizzle, "ii
// VI+ only builtins.
//===----------------------------------------------------------------------===//
+TARGET_BUILTIN(__builtin_amdgcn_div_fixuph, "hhhh", "nc", "16-bit-insts")
+TARGET_BUILTIN(__builtin_amdgcn_rcph, "hh", "nc", "16-bit-insts")
+TARGET_BUILTIN(__builtin_amdgcn_rsqh, "hh", "nc", "16-bit-insts")
+TARGET_BUILTIN(__builtin_amdgcn_sinh, "hh", "nc", "16-bit-insts")
+TARGET_BUILTIN(__builtin_amdgcn_cosh, "hh", "nc", "16-bit-insts")
+TARGET_BUILTIN(__builtin_amdgcn_ldexph, "hhi", "nc", "16-bit-insts")
+TARGET_BUILTIN(__builtin_amdgcn_frexp_manth, "hh", "nc", "16-bit-insts")
+TARGET_BUILTIN(__builtin_amdgcn_frexp_exph, "ih", "nc", "16-bit-insts")
+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")
//===----------------------------------------------------------------------===//
Modified: cfe/trunk/lib/CodeGen/CGBuiltin.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGBuiltin.cpp?rev=286741&r1=286740&r2=286741&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CGBuiltin.cpp (original)
+++ cfe/trunk/lib/CodeGen/CGBuiltin.cpp Sat Nov 12 20:37:05 2016
@@ -8190,38 +8190,45 @@ Value *CodeGenFunction::EmitAMDGPUBuilti
return emitBinaryBuiltin(*this, E, Intrinsic::amdgcn_ds_swizzle);
case AMDGPU::BI__builtin_amdgcn_div_fixup:
case AMDGPU::BI__builtin_amdgcn_div_fixupf:
+ case AMDGPU::BI__builtin_amdgcn_div_fixuph:
return emitTernaryBuiltin(*this, E, Intrinsic::amdgcn_div_fixup);
case AMDGPU::BI__builtin_amdgcn_trig_preop:
case AMDGPU::BI__builtin_amdgcn_trig_preopf:
return emitFPIntBuiltin(*this, E, Intrinsic::amdgcn_trig_preop);
case AMDGPU::BI__builtin_amdgcn_rcp:
case AMDGPU::BI__builtin_amdgcn_rcpf:
+ case AMDGPU::BI__builtin_amdgcn_rcph:
return emitUnaryBuiltin(*this, E, Intrinsic::amdgcn_rcp);
case AMDGPU::BI__builtin_amdgcn_rsq:
case AMDGPU::BI__builtin_amdgcn_rsqf:
+ case AMDGPU::BI__builtin_amdgcn_rsqh:
return emitUnaryBuiltin(*this, E, Intrinsic::amdgcn_rsq);
case AMDGPU::BI__builtin_amdgcn_rsq_clamp:
case AMDGPU::BI__builtin_amdgcn_rsq_clampf:
return emitUnaryBuiltin(*this, E, Intrinsic::amdgcn_rsq_clamp);
case AMDGPU::BI__builtin_amdgcn_sinf:
+ case AMDGPU::BI__builtin_amdgcn_sinh:
return emitUnaryBuiltin(*this, E, Intrinsic::amdgcn_sin);
case AMDGPU::BI__builtin_amdgcn_cosf:
+ case AMDGPU::BI__builtin_amdgcn_cosh:
return emitUnaryBuiltin(*this, E, Intrinsic::amdgcn_cos);
case AMDGPU::BI__builtin_amdgcn_log_clampf:
return emitUnaryBuiltin(*this, E, Intrinsic::amdgcn_log_clamp);
case AMDGPU::BI__builtin_amdgcn_ldexp:
case AMDGPU::BI__builtin_amdgcn_ldexpf:
+ case AMDGPU::BI__builtin_amdgcn_ldexph:
return emitFPIntBuiltin(*this, E, Intrinsic::amdgcn_ldexp);
case AMDGPU::BI__builtin_amdgcn_frexp_mant:
- case AMDGPU::BI__builtin_amdgcn_frexp_mantf: {
+ case AMDGPU::BI__builtin_amdgcn_frexp_mantf:
+ case AMDGPU::BI__builtin_amdgcn_frexp_manth:
return emitUnaryBuiltin(*this, E, Intrinsic::amdgcn_frexp_mant);
- }
case AMDGPU::BI__builtin_amdgcn_frexp_exp:
- case AMDGPU::BI__builtin_amdgcn_frexp_expf: {
+ case AMDGPU::BI__builtin_amdgcn_frexp_expf:
+ case AMDGPU::BI__builtin_amdgcn_frexp_exph:
return emitUnaryBuiltin(*this, E, Intrinsic::amdgcn_frexp_exp);
- }
case AMDGPU::BI__builtin_amdgcn_fract:
case AMDGPU::BI__builtin_amdgcn_fractf:
+ case AMDGPU::BI__builtin_amdgcn_fracth:
return emitUnaryBuiltin(*this, E, Intrinsic::amdgcn_fract);
case AMDGPU::BI__builtin_amdgcn_lerp:
return emitTernaryBuiltin(*this, E, Intrinsic::amdgcn_lerp);
@@ -8235,6 +8242,7 @@ Value *CodeGenFunction::EmitAMDGPUBuilti
return emitTernaryBuiltin(*this, E, Intrinsic::amdgcn_fcmp);
case AMDGPU::BI__builtin_amdgcn_class:
case AMDGPU::BI__builtin_amdgcn_classf:
+ case AMDGPU::BI__builtin_amdgcn_classh:
return emitFPIntBuiltin(*this, E, Intrinsic::amdgcn_class);
case AMDGPU::BI__builtin_amdgcn_read_exec: {
Removed: cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-error.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-error.cl?rev=286740&view=auto
==============================================================================
--- cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-error.cl (original)
+++ cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-error.cl (removed)
@@ -1,64 +0,0 @@
-// REQUIRES: amdgpu-registered-target
-// RUN: %clang_cc1 -triple amdgcn-unknown-amdhsa -target-cpu tahiti -verify -S -o - %s
-
-// FIXME: We only get one error if the functions are the other order in the
-// file.
-
-#pragma OPENCL EXTENSION cl_khr_fp64 : enable
-typedef unsigned long ulong;
-typedef unsigned int uint;
-
-ulong test_s_memrealtime()
-{
- return __builtin_amdgcn_s_memrealtime(); // expected-error {{'__builtin_amdgcn_s_memrealtime' needs target feature s-memrealtime}}
-}
-
-void test_s_sleep(int x)
-{
- __builtin_amdgcn_s_sleep(x); // expected-error {{argument to '__builtin_amdgcn_s_sleep' must be a constant integer}}
-}
-
-void test_s_incperflevel(int x)
-{
- __builtin_amdgcn_s_incperflevel(x); // expected-error {{argument to '__builtin_amdgcn_s_incperflevel' must be a constant integer}}
-}
-
-void test_s_decperflevel(int x)
-{
- __builtin_amdgcn_s_decperflevel(x); // expected-error {{argument to '__builtin_amdgcn_s_decperflevel' must be a constant integer}}
-}
-
-void test_sicmp_i32(global ulong* out, int a, int b, uint c)
-{
- *out = __builtin_amdgcn_sicmp(a, b, c); // expected-error {{argument to '__builtin_amdgcn_sicmp' must be a constant integer}}
-}
-
-void test_uicmp_i32(global ulong* out, uint a, uint b, uint c)
-{
- *out = __builtin_amdgcn_uicmp(a, b, c); // expected-error {{argument to '__builtin_amdgcn_uicmp' must be a constant integer}}
-}
-
-void test_sicmp_i64(global ulong* out, long a, long b, uint c)
-{
- *out = __builtin_amdgcn_sicmpl(a, b, c); // expected-error {{argument to '__builtin_amdgcn_sicmpl' must be a constant integer}}
-}
-
-void test_uicmp_i64(global ulong* out, ulong a, ulong b, uint c)
-{
- *out = __builtin_amdgcn_uicmpl(a, b, c); // expected-error {{argument to '__builtin_amdgcn_uicmpl' must be a constant integer}}
-}
-
-void test_fcmp_f32(global ulong* out, float a, float b, uint c)
-{
- *out = __builtin_amdgcn_fcmpf(a, b, c); // expected-error {{argument to '__builtin_amdgcn_fcmpf' must be a constant integer}}
-}
-
-void test_fcmp_f64(global ulong* out, double a, double b, uint c)
-{
- *out = __builtin_amdgcn_fcmp(a, b, c); // expected-error {{argument to '__builtin_amdgcn_fcmp' must be a constant integer}}
-}
-
-void test_ds_swizzle(global int* out, int a, int b)
-{
- *out = __builtin_amdgcn_ds_swizzle(a, b); // expected-error {{argument to '__builtin_amdgcn_ds_swizzle' must be a constant integer}}
-}
Modified: cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-vi.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-vi.cl?rev=286741&r1=286740&r2=286741&view=diff
==============================================================================
--- cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-vi.cl (original)
+++ cfe/trunk/test/CodeGenOpenCL/builtins-amdgcn-vi.cl Sat Nov 12 20:37:05 2016
@@ -1,8 +1,79 @@
// REQUIRES: amdgpu-registered-target
// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu tonga -S -emit-llvm -o - %s | FileCheck %s
+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+
typedef unsigned long ulong;
+// CHECK-LABEL: @test_div_fixup_f16
+// CHECK: call half @llvm.amdgcn.div.fixup.f16
+void test_div_fixup_f16(global half* out, half a, half b, half c)
+{
+ *out = __builtin_amdgcn_div_fixuph(a, b, c);
+}
+
+// CHECK-LABEL: @test_rcp_f16
+// CHECK: call half @llvm.amdgcn.rcp.f16
+void test_rcp_f16(global half* out, half a)
+{
+ *out = __builtin_amdgcn_rcph(a);
+}
+
+// CHECK-LABEL: @test_rsq_f16
+// CHECK: call half @llvm.amdgcn.rsq.f16
+void test_rsq_f16(global half* out, half a)
+{
+ *out = __builtin_amdgcn_rsqh(a);
+}
+
+// CHECK-LABEL: @test_sin_f16
+// CHECK: call half @llvm.amdgcn.sin.f16
+void test_sin_f16(global half* out, half a)
+{
+ *out = __builtin_amdgcn_sinh(a);
+}
+
+// CHECK-LABEL: @test_cos_f16
+// CHECK: call half @llvm.amdgcn.cos.f16
+void test_cos_f16(global half* out, half a)
+{
+ *out = __builtin_amdgcn_cosh(a);
+}
+
+// CHECK-LABEL: @test_ldexp_f16
+// CHECK: call half @llvm.amdgcn.ldexp.f16
+void test_ldexp_f16(global half* out, half a, int b)
+{
+ *out = __builtin_amdgcn_ldexph(a, b);
+}
+
+// CHECK-LABEL: @test_frexp_mant_f16
+// CHECK: call half @llvm.amdgcn.frexp.mant.f16
+void test_frexp_mant_f16(global half* out, half a)
+{
+ *out = __builtin_amdgcn_frexp_manth(a);
+}
+
+// CHECK-LABEL: @test_frexp_exp_f16
+// CHECK: call i32 @llvm.amdgcn.frexp.exp.f16
+void test_frexp_exp_f16(global short* out, half a)
+{
+ *out = __builtin_amdgcn_frexp_exph(a);
+}
+
+// CHECK-LABEL: @test_fract_f16
+// CHECK: call half @llvm.amdgcn.fract.f16
+void test_fract_f16(global half* out, half a)
+{
+ *out = __builtin_amdgcn_fracth(a);
+}
+
+// CHECK-LABEL: @test_class_f16
+// CHECK: call i1 @llvm.amdgcn.class.f16
+void test_class_f16(global half* out, half a, int b)
+{
+ *out = __builtin_amdgcn_classh(a, b);
+}
// CHECK-LABEL: @test_s_memrealtime
// CHECK: call i64 @llvm.amdgcn.s.memrealtime()
Added: cfe/trunk/test/SemaOpenCL/builtins-amdgcn-error-f16.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaOpenCL/builtins-amdgcn-error-f16.cl?rev=286741&view=auto
==============================================================================
--- cfe/trunk/test/SemaOpenCL/builtins-amdgcn-error-f16.cl (added)
+++ cfe/trunk/test/SemaOpenCL/builtins-amdgcn-error-f16.cl Sat Nov 12 20:37:05 2016
@@ -0,0 +1,18 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgcn-- -target-cpu tahiti -verify -S -o - %s
+
+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+
+void test_f16(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}}
+ *out = __builtin_amdgcn_rsqh(a); // expected-error {{'__builtin_amdgcn_rsqh' needs target feature 16-bit-insts}}
+ *out = __builtin_amdgcn_sinh(a); // expected-error {{'__builtin_amdgcn_sinh' needs target feature 16-bit-insts}}
+ *out = __builtin_amdgcn_cosh(a); // expected-error {{'__builtin_amdgcn_cosh' needs target feature 16-bit-insts}}
+ *out = __builtin_amdgcn_ldexph(a, b); // expected-error {{'__builtin_amdgcn_ldexph' needs target feature 16-bit-insts}}
+ *out = __builtin_amdgcn_frexp_manth(a); // expected-error {{'__builtin_amdgcn_frexp_manth' needs target feature 16-bit-insts}}
+ *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}}
+}
Added: cfe/trunk/test/SemaOpenCL/builtins-amdgcn-error.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaOpenCL/builtins-amdgcn-error.cl?rev=286741&view=auto
==============================================================================
--- cfe/trunk/test/SemaOpenCL/builtins-amdgcn-error.cl (added)
+++ cfe/trunk/test/SemaOpenCL/builtins-amdgcn-error.cl Sat Nov 12 20:37:05 2016
@@ -0,0 +1,64 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgcn-- -target-cpu tahiti -verify -S -o - %s
+
+// FIXME: We only get one error if the functions are the other order in the
+// file.
+
+#pragma OPENCL EXTENSION cl_khr_fp64 : enable
+typedef unsigned long ulong;
+typedef unsigned int uint;
+
+ulong test_s_memrealtime()
+{
+ return __builtin_amdgcn_s_memrealtime(); // expected-error {{'__builtin_amdgcn_s_memrealtime' needs target feature s-memrealtime}}
+}
+
+void test_s_sleep(int x)
+{
+ __builtin_amdgcn_s_sleep(x); // expected-error {{argument to '__builtin_amdgcn_s_sleep' must be a constant integer}}
+}
+
+void test_s_incperflevel(int x)
+{
+ __builtin_amdgcn_s_incperflevel(x); // expected-error {{argument to '__builtin_amdgcn_s_incperflevel' must be a constant integer}}
+}
+
+void test_s_decperflevel(int x)
+{
+ __builtin_amdgcn_s_decperflevel(x); // expected-error {{argument to '__builtin_amdgcn_s_decperflevel' must be a constant integer}}
+}
+
+void test_sicmp_i32(global ulong* out, int a, int b, uint c)
+{
+ *out = __builtin_amdgcn_sicmp(a, b, c); // expected-error {{argument to '__builtin_amdgcn_sicmp' must be a constant integer}}
+}
+
+void test_uicmp_i32(global ulong* out, uint a, uint b, uint c)
+{
+ *out = __builtin_amdgcn_uicmp(a, b, c); // expected-error {{argument to '__builtin_amdgcn_uicmp' must be a constant integer}}
+}
+
+void test_sicmp_i64(global ulong* out, long a, long b, uint c)
+{
+ *out = __builtin_amdgcn_sicmpl(a, b, c); // expected-error {{argument to '__builtin_amdgcn_sicmpl' must be a constant integer}}
+}
+
+void test_uicmp_i64(global ulong* out, ulong a, ulong b, uint c)
+{
+ *out = __builtin_amdgcn_uicmpl(a, b, c); // expected-error {{argument to '__builtin_amdgcn_uicmpl' must be a constant integer}}
+}
+
+void test_fcmp_f32(global ulong* out, float a, float b, uint c)
+{
+ *out = __builtin_amdgcn_fcmpf(a, b, c); // expected-error {{argument to '__builtin_amdgcn_fcmpf' must be a constant integer}}
+}
+
+void test_fcmp_f64(global ulong* out, double a, double b, uint c)
+{
+ *out = __builtin_amdgcn_fcmp(a, b, c); // expected-error {{argument to '__builtin_amdgcn_fcmp' must be a constant integer}}
+}
+
+void test_ds_swizzle(global int* out, int a, int b)
+{
+ *out = __builtin_amdgcn_ds_swizzle(a, b); // expected-error {{argument to '__builtin_amdgcn_ds_swizzle' must be a constant integer}}
+}
More information about the cfe-commits
mailing list