| /external/eigen/test/ |
| D | half_float.cpp | 16 struct half;
19 using Eigen::half;
26 VERIFY_IS_EQUAL(half(1.0f).x, 0x3c00); in test_conversion()
27 VERIFY_IS_EQUAL(half(0.5f).x, 0x3800); in test_conversion()
28 VERIFY_IS_EQUAL(half(0.33333f).x, 0x3555); in test_conversion()
29 VERIFY_IS_EQUAL(half(0.0f).x, 0x0000); in test_conversion()
30 VERIFY_IS_EQUAL(half(-0.0f).x, 0x8000); in test_conversion()
31 VERIFY_IS_EQUAL(half(65504.0f).x, 0x7bff); in test_conversion()
32 VERIFY_IS_EQUAL(half(65536.0f).x, 0x7c00); // Becomes infinity. in test_conversion()
35 VERIFY_IS_EQUAL(half(-5.96046e-08f).x, 0x8001); in test_conversion()
[all …]
|
| /external/swiftshader/third_party/llvm-7.0/llvm/test/Transforms/SLPVectorizer/AMDGPU/ |
| D | packed-math.ll | 8 ; GFX89: load <2 x half>, <2 x half> addrspace(3)*
9 ; GFX89: load <2 x half>, <2 x half> addrspace(3)*
10 ; GFX89: fmul <2 x half>
11 ; GFX89: store <2 x half> %{{.*}}, <2 x half> addrspace(3)* %
13 define amdgpu_kernel void @test1_as_3_3_3_v2f16(half addrspace(3)* %a, half addrspace(3)* %b, half …
14 %i0 = load half, half addrspace(3)* %a, align 2
15 %i1 = load half, half addrspace(3)* %b, align 2
16 %mul = fmul half %i0, %i1
17 %arrayidx3 = getelementptr inbounds half, half addrspace(3)* %a, i64 1
18 %i3 = load half, half addrspace(3)* %arrayidx3, align 2
[all …]
|
| /external/swiftshader/third_party/llvm-7.0/llvm/test/CodeGen/AArch64/ |
| D | fp16-vector-load-store.ll | 4 define <4 x half> @load_64(<4 x half>* nocapture readonly %a) #0 {
8 %0 = load <4 x half>, <4 x half>* %a, align 8
9 ret <4 x half> %0
13 define <8 x half> @load_128(<8 x half>* nocapture readonly %a) #0 {
17 %0 = load <8 x half>, <8 x half>* %a, align 16
18 ret <8 x half> %0
22 define <4 x half> @load_dup_64(half* nocapture readonly %a) #0 {
26 %0 = load half, half* %a, align 2
27 %1 = insertelement <4 x half> undef, half %0, i32 0
28 %2 = shufflevector <4 x half> %1, <4 x half> undef, <4 x i32> zeroinitializer
[all …]
|
| D | fp16_intrinsic_lane.ll | 3 declare half @llvm.aarch64.neon.fmulx.f16(half, half)
4 declare <4 x half> @llvm.aarch64.neon.fmulx.v4f16(<4 x half>, <4 x half>)
5 declare <8 x half> @llvm.aarch64.neon.fmulx.v8f16(<8 x half>, <8 x half>)
6 declare <4 x half> @llvm.fma.v4f16(<4 x half>, <4 x half>, <4 x half>)
7 declare <8 x half> @llvm.fma.v8f16(<8 x half>, <8 x half>, <8 x half>)
8 declare half @llvm.fma.f16(half, half, half) #1
10 define dso_local <4 x half> @t_vfma_lane_f16(<4 x half> %a, <4 x half> %b, <4 x half> %c, i32 %lane…
16 %lane1 = shufflevector <4 x half> %c, <4 x half> undef, <4 x i32> zeroinitializer
17 %fmla3 = tail call <4 x half> @llvm.fma.v4f16(<4 x half> %b, <4 x half> %lane1, <4 x half> %a)
18 ret <4 x half> %fmla3
[all …]
|
| D | fp16_intrinsic_vector_2op.ll | 3 declare <4 x half> @llvm.aarch64.neon.fmulx.v4f16(<4 x half>, <4 x half>)
4 declare <8 x half> @llvm.aarch64.neon.fmulx.v8f16(<8 x half>, <8 x half>)
5 declare <4 x half> @llvm.aarch64.neon.fminnmp.v4f16(<4 x half>, <4 x half>)
6 declare <8 x half> @llvm.aarch64.neon.fminnmp.v8f16(<8 x half>, <8 x half>)
7 declare <4 x half> @llvm.aarch64.neon.fmaxnmp.v4f16(<4 x half>, <4 x half>)
8 declare <8 x half> @llvm.aarch64.neon.fmaxnmp.v8f16(<8 x half>, <8 x half>)
9 declare <4 x half> @llvm.aarch64.neon.fabd.v4f16(<4 x half>, <4 x half>)
10 declare <8 x half> @llvm.aarch64.neon.fabd.v8f16(<8 x half>, <8 x half>)
11 declare <4 x half> @llvm.fabs.v4f16(<4 x half>)
12 declare <8 x half> @llvm.fabs.v8f16(<8 x half>)
[all …]
|
| D | fp16-vector-shuffle.ll | 4 define <4 x half> @select_64(<4 x half> %a, <4 x half> %b, <4 x i16> %c) #0 {
8 %0 = bitcast <4 x half> %a to <4 x i16>
9 %1 = bitcast <4 x half> %b to <4 x i16>
14 %3 = bitcast <4 x i16> %vbsl5.i to <4 x half>
15 ret <4 x half> %3
19 define <8 x half> @select_128(<8 x half> %a, <8 x half> %b, <8 x i16> %c) #0 {
23 %0 = bitcast <8 x half> %a to <8 x i16>
24 %1 = bitcast <8 x half> %b to <8 x i16>
29 %3 = bitcast <8 x i16> %vbsl5.i to <8 x half>
30 ret <8 x half> %3
[all …]
|
| /external/llvm/test/CodeGen/AArch64/ |
| D | fp16-vector-load-store.ll | 4 define <4 x half> @load_64(<4 x half>* nocapture readonly %a) #0 {
8 %0 = load <4 x half>, <4 x half>* %a, align 8
9 ret <4 x half> %0
13 define <8 x half> @load_128(<8 x half>* nocapture readonly %a) #0 {
17 %0 = load <8 x half>, <8 x half>* %a, align 16
18 ret <8 x half> %0
22 define <4 x half> @load_dup_64(half* nocapture readonly %a) #0 {
26 %0 = load half, half* %a, align 2
27 %1 = insertelement <4 x half> undef, half %0, i32 0
28 %2 = shufflevector <4 x half> %1, <4 x half> undef, <4 x i32> zeroinitializer
[all …]
|
| D | fp16-vector-shuffle.ll | 4 define <4 x half> @select_64(<4 x half> %a, <4 x half> %b, <4 x i16> %c) #0 {
8 %0 = bitcast <4 x half> %a to <4 x i16>
9 %1 = bitcast <4 x half> %b to <4 x i16>
14 %3 = bitcast <4 x i16> %vbsl5.i to <4 x half>
15 ret <4 x half> %3
19 define <8 x half> @select_128(<8 x half> %a, <8 x half> %b, <8 x i16> %c) #0 {
23 %0 = bitcast <8 x half> %a to <8 x i16>
24 %1 = bitcast <8 x half> %b to <8 x i16>
29 %3 = bitcast <8 x i16> %vbsl5.i to <8 x half>
30 ret <8 x half> %3
[all …]
|
| D | f16-instructions.ll | 11 define half @test_fadd(half %a, half %b) #0 {
12 %r = fadd half %a, %b
13 ret half %r
22 define half @test_fsub(half %a, half %b) #0 {
23 %r = fsub half %a, %b
24 ret half %r
33 define half @test_fmul(half %a, half %b) #0 {
34 %r = fmul half %a, %b
35 ret half %r
44 define half @test_fdiv(half %a, half %b) #0 {
[all …]
|
| /external/eigen/Eigen/src/Core/arch/CUDA/ |
| D | Half.h | 48 struct half;
76 struct half : public half_impl::half_base { struct
81 EIGEN_DEVICE_FUNC half() {} in half() function
83 EIGEN_DEVICE_FUNC half(const __half& h) : half_impl::half_base(h) {} in half() function
84 EIGEN_DEVICE_FUNC half(const half& h) : half_impl::half_base(h) {} in half() function
86 explicit EIGEN_DEVICE_FUNC half(bool b) in half() argument
89 explicit EIGEN_DEVICE_FUNC half(const T& val) in half() function
91 explicit EIGEN_DEVICE_FUNC half(float f) in half() argument
135 EIGEN_DEVICE_FUNC half& operator=(const half& other) { argument
150 __device__ half operator + (const half& a, const half& b) {
[all …]
|
| /external/tensorflow/tensorflow/compiler/xla/tests/ |
| D | half_test.cc | 43 std::function<half(half)> compute_func;
51 std::vector<half> x({half(1.4), half(-2.3), half(3.2), half(-4.1), half(9.0), in XLA_TEST_P()
52 half(42.0), half(-9.0), half(-100.0)}); in XLA_TEST_P()
55 auto x_data = CreateR1Parameter<half>(x, /*parameter_number=*/0, "x", in XLA_TEST_P()
58 std::function<half(half)> compute_func = GetParam().compute_func; in XLA_TEST_P()
59 std::vector<half> expected; in XLA_TEST_P()
67 ComputeAndCompareR1<half>(&builder, expected, {x_data.get()}, error_spec_); in XLA_TEST_P()
70 half sign_imp(half value) { in sign_imp()
72 return half((x < .0) ? -1 : (x > .0)); in sign_imp()
75 half round_imp(half value) { in round_imp()
[all …]
|
| /external/tensorflow/tensorflow/lite/delegates/gpu/cl/kernels/ |
| D | concat_test.cc | 37 src0.data = {half(0.0f), half(-1.0f), half(-0.05f), half(0.045f)}; in TEST_F()
39 src1.data = {half(1.0f), half(-1.2f), half(-0.45f), half(1.045f), in TEST_F()
40 half(1.1f), half(-1.3f), half(-0.55f), half(2.045f)}; in TEST_F()
60 {half(0.0f), half(-1.0f), half(1.0f), half(-1.2f), in TEST_F()
61 half(-0.45f), half(1.045f), half(-0.05f), half(0.045f), in TEST_F()
62 half(1.1f), half(-1.3f), half(-0.55f), half(2.045f)})); in TEST_F()
70 src0.data = {half(0.0f), half(-1.0f), half(-0.05f), half(0.045f)}; in TEST_F()
72 src1.data = {half(1.0f), half(-1.2f)}; in TEST_F()
91 Pointwise(FloatNear(0.0f), {half(0.0f), half(-1.0f), half(-0.05f), in TEST_F()
92 half(0.045f), half(1.0f), half(-1.2f)})); in TEST_F()
[all …]
|
| /external/swiftshader/third_party/llvm-7.0/llvm/test/CodeGen/AMDGPU/ |
| D | llvm.amdgcn.fmad.ftz.f16.ll | 5 declare half @llvm.amdgcn.fmad.ftz.f16(half %a, half %b, half %c)
11 half addrspace(1)* %r,
12 half addrspace(1)* %a,
13 half addrspace(1)* %b,
14 half addrspace(1)* %c) {
15 %a.val = load half, half addrspace(1)* %a
16 %b.val = load half, half addrspace(1)* %b
17 %c.val = load half, half addrspace(1)* %c
18 %r.val = call half @llvm.amdgcn.fmad.ftz.f16(half %a.val, half %b.val, half %c.val)
19 store half %r.val, half addrspace(1)* %r
[all …]
|
| D | packed-op-sel.ll | 14 …kernel void @fma_vector_vector_scalar_lo(<2 x half> addrspace(1)* %out, <2 x half> addrspace(3)* %…
16 %lds.gep1 = getelementptr inbounds <2 x half>, <2 x half> addrspace(3)* %lds, i32 1
18 %vec0 = load volatile <2 x half>, <2 x half> addrspace(3)* %lds, align 4
19 %vec1 = load volatile <2 x half>, <2 x half> addrspace(3)* %lds.gep1, align 4
20 %scalar0 = load volatile half, half addrspace(3)* %arg2, align 2
22 %scalar0.vec = insertelement <2 x half> undef, half %scalar0, i32 0
23 …%scalar0.broadcast = shufflevector <2 x half> %scalar0.vec, <2 x half> undef, <2 x i32> zeroinitia…
25 …%result = tail call <2 x half> @llvm.fma.v2f16(<2 x half> %vec0, <2 x half> %vec1, <2 x half> %sca…
26 store <2 x half> %result, <2 x half> addrspace(1)* %out, align 4
42 …ma_vector_vector_neg_broadcast_scalar_lo(<2 x half> addrspace(1)* %out, <2 x half> addrspace(3)* %…
[all …]
|
| D | fmuladd.f16.ll | 8 declare half @llvm.fmuladd.f16(half, half, half) #1
9 declare half @llvm.fabs.f16(half) #1
15 define amdgpu_kernel void @fmuladd_f16(half addrspace(1)* %out, half addrspace(1)* %in1,
16 half addrspace(1)* %in2, half addrspace(1)* %in3) #0 {
17 %r0 = load half, half addrspace(1)* %in1
18 %r1 = load half, half addrspace(1)* %in2
19 %r2 = load half, half addrspace(1)* %in3
20 %r3 = tail call half @llvm.fmuladd.f16(half %r0, half %r1, half %r2)
21 store half %r3, half addrspace(1)* %out
33 define amdgpu_kernel void @fmuladd_2.0_a_b_f16(half addrspace(1)* %out, half addrspace(1)* %in) #0 {
[all …]
|
| D | fcanonicalize.f16.ll | 5 declare half @llvm.fabs.f16(half) #0
6 declare half @llvm.canonicalize.f16(half) #0
7 declare <2 x half> @llvm.fabs.v2f16(<2 x half>) #0
8 declare <2 x half> @llvm.canonicalize.v2f16(<2 x half>) #0
18 define amdgpu_kernel void @v_test_canonicalize_var_f16(half addrspace(1)* %out) #1 {
19 %val = load half, half addrspace(1)* %out
20 %canonicalized = call half @llvm.canonicalize.f16(half %val)
21 store half %canonicalized, half addrspace(1)* undef
28 define amdgpu_kernel void @s_test_canonicalize_var_f16(half addrspace(1)* %out, i16 zeroext %val.ar…
29 %val = bitcast i16 %val.arg to half
[all …]
|
| D | v_mac_f16.ll | 18 half addrspace(1)* %r,
19 half addrspace(1)* %a,
20 half addrspace(1)* %b,
21 half addrspace(1)* %c) #0 {
23 %a.val = load half, half addrspace(1)* %a
24 %b.val = load half, half addrspace(1)* %b
25 %c.val = load half, half addrspace(1)* %c
27 %t.val = fmul half %a.val, %b.val
28 %r.val = fadd half %t.val, %c.val
30 store half %r.val, half addrspace(1)* %r
[all …]
|
| D | fneg-fabs.f16.ll | 12 define amdgpu_kernel void @fneg_fabs_fadd_f16(half addrspace(1)* %out, half %x, half %y) {
13 %fabs = call half @llvm.fabs.f16(half %x)
14 %fsub = fsub half -0.0, %fabs
15 %fadd = fadd half %y, %fsub
16 store half %fadd, half addrspace(1)* %out, align 2
30 define amdgpu_kernel void @fneg_fabs_fmul_f16(half addrspace(1)* %out, half %x, half %y) {
31 %fabs = call half @llvm.fabs.f16(half %x)
32 %fsub = fsub half -0.0, %fabs
33 %fmul = fmul half %y, %fsub
34 store half %fmul, half addrspace(1)* %out, align 2
[all …]
|
| D | pk_max_f16_literal.ll | 5 define amdgpu_kernel void @test_pk_max_f16_literal_0_1(<2 x half> addrspace(1)* nocapture %arg) {
9 %tmp2 = getelementptr inbounds <2 x half>, <2 x half> addrspace(1)* %arg, i64 %tmp1
10 %tmp3 = load <2 x half>, <2 x half> addrspace(1)* %tmp2, align 4
11 …%tmp4 = tail call <2 x half> @llvm.maxnum.v2f16(<2 x half> %tmp3, <2 x half> <half 0xH0000, half 0…
12 store <2 x half> %tmp4, <2 x half> addrspace(1)* %tmp2, align 4
18 define amdgpu_kernel void @test_pk_max_f16_literal_1_0(<2 x half> addrspace(1)* nocapture %arg) {
22 %tmp2 = getelementptr inbounds <2 x half>, <2 x half> addrspace(1)* %arg, i64 %tmp1
23 %tmp3 = load <2 x half>, <2 x half> addrspace(1)* %tmp2, align 4
24 …%tmp4 = tail call <2 x half> @llvm.maxnum.v2f16(<2 x half> %tmp3, <2 x half> <half 0xH3C00, half 0…
25 store <2 x half> %tmp4, <2 x half> addrspace(1)* %tmp2, align 4
[all …]
|
| D | select.f16.ll | 21 half addrspace(1)* %r,
22 half addrspace(1)* %a,
23 half addrspace(1)* %b,
24 half addrspace(1)* %c,
25 half addrspace(1)* %d) {
27 %a.val = load volatile half, half addrspace(1)* %a
28 %b.val = load volatile half, half addrspace(1)* %b
29 %c.val = load volatile half, half addrspace(1)* %c
30 %d.val = load volatile half, half addrspace(1)* %d
31 %fcmp = fcmp olt half %a.val, %b.val
[all …]
|
| D | llvm.amdgcn.div.fixup.f16.ll | 3 declare half @llvm.amdgcn.div.fixup.f16(half %a, half %b, half %c)
13 half addrspace(1)* %r,
14 half addrspace(1)* %a,
15 half addrspace(1)* %b,
16 half addrspace(1)* %c) {
18 %a.val = load volatile half, half addrspace(1)* %a
19 %b.val = load volatile half, half addrspace(1)* %b
20 %c.val = load volatile half, half addrspace(1)* %c
21 %r.val = call half @llvm.amdgcn.div.fixup.f16(half %a.val, half %b.val, half %c.val)
22 store half %r.val, half addrspace(1)* %r
[all …]
|
| /external/swiftshader/third_party/llvm-7.0/llvm/test/CodeGen/ARM/ |
| D | fp16-vminmaxnm-safe.ll | 4 ; TODO: we can't pass half-precision arguments as "half" types yet. We do
7 ; is the shortest way to get a half type. But when we can pass half types, we
10 define half @fp16_vminnm_o(i16 signext %a, i16 signext %b) {
14 %0 = bitcast i16 %a to half
15 %1 = bitcast i16 %b to half
16 %cmp = fcmp olt half %0, %1
17 %cond = select i1 %cmp, half %0, half %1
18 ret half %cond
21 define half @fp16_vminnm_o_rev(i16 signext %a, i16 signext %b) {
25 %0 = bitcast i16 %a to half
[all …]
|
| D | fp16-vminmaxnm.ll | 4 ; TODO: we can't pass half-precision arguments as "half" types yet. We do
7 ; is the shortest way to get a half type. But when we can pass half types, we
10 define half @fp16_vminnm_o(i16 signext %a, i16 signext %b) {
16 %0 = bitcast i16 %a to half
17 %1 = bitcast i16 %b to half
18 %cmp = fcmp fast olt half %0, %1
19 %cond = select i1 %cmp, half %0, half %1
20 ret half %cond
23 define half @fp16_vminnm_o_rev(i16 signext %a, i16 signext %b) {
29 %0 = bitcast i16 %a to half
[all …]
|
| D | fp16-promote.ll | 17 define void @test_fadd(half* %p, half* %q) #0 {
18 %a = load half, half* %p, align 2
19 %b = load half, half* %q, align 2
20 %r = fadd half %a, %b
21 store half %r, half* %p
34 define void @test_fsub(half* %p, half* %q) #0 {
35 %a = load half, half* %p, align 2
36 %b = load half, half* %q, align 2
37 %r = fsub half %a, %b
38 store half %r, half* %p
[all …]
|
| /external/llvm/test/CodeGen/ARM/ |
| D | fp16-promote.ll | 17 define void @test_fadd(half* %p, half* %q) #0 {
18 %a = load half, half* %p, align 2
19 %b = load half, half* %q, align 2
20 %r = fadd half %a, %b
21 store half %r, half* %p
34 define void @test_fsub(half* %p, half* %q) #0 {
35 %a = load half, half* %p, align 2
36 %b = load half, half* %q, align 2
37 %r = fsub half %a, %b
38 store half %r, half* %p
[all …]
|