1 /*
2 * Copyright (C) 1999-2007 Brian Paul All Rights Reserved.
3 * Copyright 2015 Philip Taylor <philip@zaynar.co.uk>
4 * Copyright 2018 Advanced Micro Devices, Inc.
5 *
6 * Permission is hereby granted, free of charge, to any person obtaining a
7 * copy of this software and associated documentation files (the "Software"),
8 * to deal in the Software without restriction, including without limitation
9 * the rights to use, copy, modify, merge, publish, distribute, sublicense,
10 * and/or sell copies of the Software, and to permit persons to whom the
11 * Software is furnished to do so, subject to the following conditions:
12 *
13 * The above copyright notice and this permission notice shall be included
14 * in all copies or substantial portions of the Software.
15 *
16 * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS
17 * OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
18 * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL
19 * THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR
20 * OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE,
21 * ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR
22 * OTHER DEALINGS IN THE SOFTWARE.
23 */
24
25 #include <assert.h>
26 #include <math.h>
27
28 #include "igt_halffloat.h"
29 #include "igt_x86.h"
30
31 typedef union { float f; int32_t i; uint32_t u; } fi_type;
32
33 /**
34 * Convert a 4-byte float to a 2-byte half float.
35 *
36 * Not all float32 values can be represented exactly as a float16 value. We
37 * round such intermediate float32 values to the nearest float16. When the
38 * float32 lies exactly between to float16 values, we round to the one with
39 * an even mantissa.
40 *
41 * This rounding behavior has several benefits:
42 * - It has no sign bias.
43 *
44 * - It reproduces the behavior of real hardware: opcode F32TO16 in Intel's
45 * GPU ISA.
46 *
47 * - By reproducing the behavior of the GPU (at least on Intel hardware),
48 * compile-time evaluation of constant packHalf2x16 GLSL expressions will
49 * result in the same value as if the expression were executed on the GPU.
50 */
_float_to_half(float val)51 static inline uint16_t _float_to_half(float val)
52 {
53 const fi_type fi = {val};
54 const int flt_m = fi.i & 0x7fffff;
55 const int flt_e = (fi.i >> 23) & 0xff;
56 const int flt_s = (fi.i >> 31) & 0x1;
57 int s, e, m = 0;
58 uint16_t result;
59
60 /* sign bit */
61 s = flt_s;
62
63 /* handle special cases */
64 if ((flt_e == 0) && (flt_m == 0)) {
65 /* zero */
66 /* m = 0; - already set */
67 e = 0;
68 } else if ((flt_e == 0) && (flt_m != 0)) {
69 /* denorm -- denorm float maps to 0 half */
70 /* m = 0; - already set */
71 e = 0;
72 } else if ((flt_e == 0xff) && (flt_m == 0)) {
73 /* infinity */
74 /* m = 0; - already set */
75 e = 31;
76 } else if ((flt_e == 0xff) && (flt_m != 0)) {
77 /* NaN */
78 m = 1;
79 e = 31;
80 } else {
81 /* regular number */
82 const int new_exp = flt_e - 127;
83 if (new_exp < -14) {
84 /* The float32 lies in the range (0.0, min_normal16) and
85 * is rounded to a nearby float16 value. The result will
86 * be either zero, subnormal, or normal.
87 */
88 e = 0;
89 m = lrintf((1 << 24) * fabsf(fi.f));
90 } else if (new_exp > 15) {
91 /* map this value to infinity */
92 /* m = 0; - already set */
93 e = 31;
94 } else {
95 /* The float32 lies in the range
96 * [min_normal16, max_normal16 + max_step16)
97 * and is rounded to a nearby float16 value. The result
98 * will be either normal or infinite.
99 */
100 e = new_exp + 15;
101 m = lrintf(flt_m / (float)(1 << 13));
102 }
103 }
104
105 assert(0 <= m && m <= 1024);
106 if (m == 1024) {
107 /* The float32 was rounded upwards into the range of the next
108 * exponent, so bump the exponent. This correctly handles the
109 * case where f32 should be rounded up to float16 infinity.
110 */
111 ++e;
112 m = 0;
113 }
114
115 result = (s << 15) | (e << 10) | m;
116 return result;
117 }
118
119 /**
120 * Convert a 2-byte half float to a 4-byte float.
121 * Based on code from:
122 * http://www.opengl.org/discussion_boards/ubb/Forum3/HTML/008786.html
123 */
_half_to_float(uint16_t val)124 static inline float _half_to_float(uint16_t val)
125 {
126 /* XXX could also use a 64K-entry lookup table */
127 const int m = val & 0x3ff;
128 const int e = (val >> 10) & 0x1f;
129 const int s = (val >> 15) & 0x1;
130 int flt_m, flt_e, flt_s;
131 fi_type fi;
132
133 /* sign bit */
134 flt_s = s;
135
136 /* handle special cases */
137 if ((e == 0) && (m == 0)) {
138 /* zero */
139 flt_m = 0;
140 flt_e = 0;
141 } else if ((e == 0) && (m != 0)) {
142 /* denorm -- denorm half will fit in non-denorm single */
143 const float half_denorm = 1.0f / 16384.0f; /* 2^-14 */
144 float mantissa = ((float) (m)) / 1024.0f;
145 float sign = s ? -1.0f : 1.0f;
146 return sign * mantissa * half_denorm;
147 } else if ((e == 31) && (m == 0)) {
148 /* infinity */
149 flt_e = 0xff;
150 flt_m = 0;
151 } else if ((e == 31) && (m != 0)) {
152 /* NaN */
153 flt_e = 0xff;
154 flt_m = 1;
155 } else {
156 /* regular */
157 flt_e = e + 112;
158 flt_m = m << 13;
159 }
160
161 fi.i = (flt_s << 31) | (flt_e << 23) | flt_m;
162 return fi.f;
163 }
164
165 #if defined(__x86_64__) && !defined(__clang__)
166 #pragma GCC push_options
167 #pragma GCC target("f16c")
168
169 #include <immintrin.h>
170
float_to_half_f16c(const float * f,uint16_t * h,unsigned int num)171 static void float_to_half_f16c(const float *f, uint16_t *h, unsigned int num)
172 {
173 for (int i = 0; i < num; i++)
174 h[i] = _cvtss_sh(f[i], 0);
175 }
176
half_to_float_f16c(const uint16_t * h,float * f,unsigned int num)177 static void half_to_float_f16c(const uint16_t *h, float *f, unsigned int num)
178 {
179 for (int i = 0; i < num; i++)
180 f[i] = _cvtsh_ss(h[i]);
181 }
182
183 #pragma GCC pop_options
184
float_to_half(const float * f,uint16_t * h,unsigned int num)185 static void float_to_half(const float *f, uint16_t *h, unsigned int num)
186 {
187 for (int i = 0; i < num; i++)
188 h[i] = _float_to_half(f[i]);
189 }
190
half_to_float(const uint16_t * h,float * f,unsigned int num)191 static void half_to_float(const uint16_t *h, float *f, unsigned int num)
192 {
193 for (int i = 0; i < num; i++)
194 f[i] = _half_to_float(h[i]);
195 }
196
resolve_float_to_half(void)197 static void (*resolve_float_to_half(void))(const float *f, uint16_t *h, unsigned int num)
198 {
199 if (igt_x86_features() & F16C)
200 return float_to_half_f16c;
201
202 return float_to_half;
203 }
204
205 void igt_float_to_half(const float *f, uint16_t *h, unsigned int num)
206 __attribute__((ifunc("resolve_float_to_half")));
207
resolve_half_to_float(void)208 static void (*resolve_half_to_float(void))(const uint16_t *h, float *f, unsigned int num)
209 {
210 if (igt_x86_features() & F16C)
211 return half_to_float_f16c;
212
213 return half_to_float;
214 }
215
216 void igt_half_to_float(const uint16_t *h, float *f, unsigned int num)
217 __attribute__((ifunc("resolve_half_to_float")));
218
219 #else
220
igt_float_to_half(const float * f,uint16_t * h,unsigned int num)221 void igt_float_to_half(const float *f, uint16_t *h, unsigned int num)
222 {
223 for (int i = 0; i < num; i++)
224 h[i] = _float_to_half(f[i]);
225 }
226
igt_half_to_float(const uint16_t * h,float * f,unsigned int num)227 void igt_half_to_float(const uint16_t *h, float *f, unsigned int num)
228 {
229 for (int i = 0; i < num; i++)
230 f[i] = _half_to_float(h[i]);
231 }
232
233 #endif
234
235