1 // This file is part of Eigen, a lightweight C++ template library 2 // for linear algebra. 3 // 4 // This Source Code Form is subject to the terms of the Mozilla 5 // Public License v. 2.0. If a copy of the MPL was not distributed 6 // with this file, You can obtain one at http://mozilla.org/MPL/2.0/. 7 // 8 // The conversion routines are Copyright (c) Fabian Giesen, 2016. 9 // The original license follows: 10 // 11 // Copyright (c) Fabian Giesen, 2016 12 // All rights reserved. 13 // Redistribution and use in source and binary forms, with or without 14 // modification, are permitted. 15 // THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS 16 // "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT 17 // LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR 18 // A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT 19 // HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, 20 // SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT 21 // LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, 22 // DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY 23 // THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT 24 // (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE 25 // OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. 26 27 28 // Standard 16-bit float type, mostly useful for GPUs. Defines a new 29 // type Eigen::half (inheriting from CUDA's __half struct) with 30 // operator overloads such that it behaves basically as an arithmetic 31 // type. It will be quite slow on CPUs (so it is recommended to stay 32 // in float32_bits for CPUs, except for simple parameter conversions, I/O 33 // to disk and the likes), but fast on GPUs. 34 35 36 #ifndef EIGEN_HALF_CUDA_H 37 #define EIGEN_HALF_CUDA_H 38 39 #if __cplusplus > 199711L 40 #define EIGEN_EXPLICIT_CAST(tgt_type) explicit operator tgt_type() 41 #else 42 #define EIGEN_EXPLICIT_CAST(tgt_type) operator tgt_type() 43 #endif 44 45 #include <sstream> 46 47 namespace Eigen { 48 49 struct half; 50 51 namespace half_impl { 52 53 #if !defined(EIGEN_HAS_CUDA_FP16) 54 // Make our own __half_raw definition that is similar to CUDA's. 55 struct __half_raw { __half_raw__half_raw56 EIGEN_DEVICE_FUNC __half_raw() : x(0) {} __half_raw__half_raw57 explicit EIGEN_DEVICE_FUNC __half_raw(unsigned short raw) : x(raw) {} 58 unsigned short x; 59 }; 60 #elif defined(EIGEN_CUDACC_VER) && EIGEN_CUDACC_VER < 90000 61 // In CUDA < 9.0, __half is the equivalent of CUDA 9's __half_raw 62 typedef __half __half_raw; 63 #endif 64 65 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC __half_raw raw_uint16_to_half(unsigned short x); 66 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC __half_raw float_to_half_rtne(float ff); 67 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC float half_to_float(__half_raw h); 68 69 struct half_base : public __half_raw { half_basehalf_base70 EIGEN_DEVICE_FUNC half_base() {} half_basehalf_base71 EIGEN_DEVICE_FUNC half_base(const half_base& h) : __half_raw(h) {} half_basehalf_base72 EIGEN_DEVICE_FUNC half_base(const __half_raw& h) : __half_raw(h) {} 73 #if defined(EIGEN_HAS_CUDA_FP16) && defined(EIGEN_CUDACC_VER) && EIGEN_CUDACC_VER >= 90000 half_basehalf_base74 EIGEN_DEVICE_FUNC half_base(const __half& h) : __half_raw(*(__half_raw*)&h) {} 75 #endif 76 }; 77 78 } // namespace half_impl 79 80 // Class definition. 81 struct half : public half_impl::half_base { 82 #if !defined(EIGEN_HAS_CUDA_FP16) || (defined(EIGEN_CUDACC_VER) && EIGEN_CUDACC_VER < 90000) 83 typedef half_impl::__half_raw __half_raw; 84 #endif 85 halfhalf86 EIGEN_DEVICE_FUNC half() {} 87 halfhalf88 EIGEN_DEVICE_FUNC half(const __half_raw& h) : half_impl::half_base(h) {} halfhalf89 EIGEN_DEVICE_FUNC half(const half& h) : half_impl::half_base(h) {} 90 #if defined(EIGEN_HAS_CUDA_FP16) && defined(EIGEN_CUDACC_VER) && EIGEN_CUDACC_VER >= 90000 halfhalf91 EIGEN_DEVICE_FUNC half(const __half& h) : half_impl::half_base(h) {} 92 #endif 93 halfhalf94 explicit EIGEN_DEVICE_FUNC half(bool b) 95 : half_impl::half_base(half_impl::raw_uint16_to_half(b ? 0x3c00 : 0)) {} 96 template<class T> halfhalf97 explicit EIGEN_DEVICE_FUNC half(const T& val) 98 : half_impl::half_base(half_impl::float_to_half_rtne(static_cast<float>(val))) {} halfhalf99 explicit EIGEN_DEVICE_FUNC half(float f) 100 : half_impl::half_base(half_impl::float_to_half_rtne(f)) {} 101 EIGEN_EXPLICIT_CASThalf102 EIGEN_DEVICE_FUNC EIGEN_EXPLICIT_CAST(bool) const { 103 // +0.0 and -0.0 become false, everything else becomes true. 104 return (x & 0x7fff) != 0; 105 } EIGEN_EXPLICIT_CASThalf106 EIGEN_DEVICE_FUNC EIGEN_EXPLICIT_CAST(signed char) const { 107 return static_cast<signed char>(half_impl::half_to_float(*this)); 108 } EIGEN_EXPLICIT_CASThalf109 EIGEN_DEVICE_FUNC EIGEN_EXPLICIT_CAST(unsigned char) const { 110 return static_cast<unsigned char>(half_impl::half_to_float(*this)); 111 } EIGEN_EXPLICIT_CASThalf112 EIGEN_DEVICE_FUNC EIGEN_EXPLICIT_CAST(short) const { 113 return static_cast<short>(half_impl::half_to_float(*this)); 114 } EIGEN_EXPLICIT_CASThalf115 EIGEN_DEVICE_FUNC EIGEN_EXPLICIT_CAST(unsigned short) const { 116 return static_cast<unsigned short>(half_impl::half_to_float(*this)); 117 } EIGEN_EXPLICIT_CASThalf118 EIGEN_DEVICE_FUNC EIGEN_EXPLICIT_CAST(int) const { 119 return static_cast<int>(half_impl::half_to_float(*this)); 120 } EIGEN_EXPLICIT_CASThalf121 EIGEN_DEVICE_FUNC EIGEN_EXPLICIT_CAST(unsigned int) const { 122 return static_cast<unsigned int>(half_impl::half_to_float(*this)); 123 } EIGEN_EXPLICIT_CASThalf124 EIGEN_DEVICE_FUNC EIGEN_EXPLICIT_CAST(long) const { 125 return static_cast<long>(half_impl::half_to_float(*this)); 126 } EIGEN_EXPLICIT_CASThalf127 EIGEN_DEVICE_FUNC EIGEN_EXPLICIT_CAST(unsigned long) const { 128 return static_cast<unsigned long>(half_impl::half_to_float(*this)); 129 } 130 #if EIGEN_HAS_CXX11 EIGEN_EXPLICIT_CASThalf131 EIGEN_DEVICE_FUNC EIGEN_EXPLICIT_CAST(long long) const { 132 return static_cast<long long>(half_impl::half_to_float(*this)); 133 } EIGEN_EXPLICIT_CASThalf134 EIGEN_DEVICE_FUNC EIGEN_EXPLICIT_CAST(unsigned long long) const { 135 return static_cast<unsigned long long>(half_to_float(*this)); 136 } 137 #endif EIGEN_EXPLICIT_CASThalf138 EIGEN_DEVICE_FUNC EIGEN_EXPLICIT_CAST(float) const { 139 return half_impl::half_to_float(*this); 140 } EIGEN_EXPLICIT_CASThalf141 EIGEN_DEVICE_FUNC EIGEN_EXPLICIT_CAST(double) const { 142 return static_cast<double>(half_impl::half_to_float(*this)); 143 } 144 145 EIGEN_DEVICE_FUNC half& operator=(const half& other) { 146 x = other.x; 147 return *this; 148 } 149 }; 150 151 } // end namespace Eigen 152 153 namespace std { 154 template<> 155 struct numeric_limits<Eigen::half> { 156 static const bool is_specialized = true; 157 static const bool is_signed = true; 158 static const bool is_integer = false; 159 static const bool is_exact = false; 160 static const bool has_infinity = true; 161 static const bool has_quiet_NaN = true; 162 static const bool has_signaling_NaN = true; 163 static const float_denorm_style has_denorm = denorm_present; 164 static const bool has_denorm_loss = false; 165 static const std::float_round_style round_style = std::round_to_nearest; 166 static const bool is_iec559 = false; 167 static const bool is_bounded = false; 168 static const bool is_modulo = false; 169 static const int digits = 11; 170 static const int digits10 = 3; // according to http://half.sourceforge.net/structstd_1_1numeric__limits_3_01half__float_1_1half_01_4.html 171 static const int max_digits10 = 5; // according to http://half.sourceforge.net/structstd_1_1numeric__limits_3_01half__float_1_1half_01_4.html 172 static const int radix = 2; 173 static const int min_exponent = -13; 174 static const int min_exponent10 = -4; 175 static const int max_exponent = 16; 176 static const int max_exponent10 = 4; 177 static const bool traps = true; 178 static const bool tinyness_before = false; 179 180 static Eigen::half (min)() { return Eigen::half_impl::raw_uint16_to_half(0x400); } 181 static Eigen::half lowest() { return Eigen::half_impl::raw_uint16_to_half(0xfbff); } 182 static Eigen::half (max)() { return Eigen::half_impl::raw_uint16_to_half(0x7bff); } 183 static Eigen::half epsilon() { return Eigen::half_impl::raw_uint16_to_half(0x0800); } 184 static Eigen::half round_error() { return Eigen::half(0.5); } 185 static Eigen::half infinity() { return Eigen::half_impl::raw_uint16_to_half(0x7c00); } 186 static Eigen::half quiet_NaN() { return Eigen::half_impl::raw_uint16_to_half(0x7e00); } 187 static Eigen::half signaling_NaN() { return Eigen::half_impl::raw_uint16_to_half(0x7e00); } 188 static Eigen::half denorm_min() { return Eigen::half_impl::raw_uint16_to_half(0x1); } 189 }; 190 191 // If std::numeric_limits<T> is specialized, should also specialize 192 // std::numeric_limits<const T>, std::numeric_limits<volatile T>, and 193 // std::numeric_limits<const volatile T> 194 // https://stackoverflow.com/a/16519653/ 195 template<> 196 struct numeric_limits<const Eigen::half> : numeric_limits<Eigen::half> {}; 197 template<> 198 struct numeric_limits<volatile Eigen::half> : numeric_limits<Eigen::half> {}; 199 template<> 200 struct numeric_limits<const volatile Eigen::half> : numeric_limits<Eigen::half> {}; 201 } // end namespace std 202 203 namespace Eigen { 204 205 namespace half_impl { 206 207 #if defined(EIGEN_HAS_CUDA_FP16) && defined(EIGEN_CUDA_ARCH) && EIGEN_CUDA_ARCH >= 530 208 209 // Intrinsics for native fp16 support. Note that on current hardware, 210 // these are no faster than float32_bits arithmetic (you need to use the half2 211 // versions to get the ALU speed increased), but you do save the 212 // conversion steps back and forth. 213 214 EIGEN_STRONG_INLINE __device__ half operator + (const half& a, const half& b) { 215 return __hadd(a, b); 216 } 217 EIGEN_STRONG_INLINE __device__ half operator * (const half& a, const half& b) { 218 return __hmul(a, b); 219 } 220 EIGEN_STRONG_INLINE __device__ half operator - (const half& a, const half& b) { 221 return __hsub(a, b); 222 } 223 EIGEN_STRONG_INLINE __device__ half operator / (const half& a, const half& b) { 224 float num = __half2float(a); 225 float denom = __half2float(b); 226 return __float2half(num / denom); 227 } 228 EIGEN_STRONG_INLINE __device__ half operator - (const half& a) { 229 return __hneg(a); 230 } 231 EIGEN_STRONG_INLINE __device__ half& operator += (half& a, const half& b) { 232 a = a + b; 233 return a; 234 } 235 EIGEN_STRONG_INLINE __device__ half& operator *= (half& a, const half& b) { 236 a = a * b; 237 return a; 238 } 239 EIGEN_STRONG_INLINE __device__ half& operator -= (half& a, const half& b) { 240 a = a - b; 241 return a; 242 } 243 EIGEN_STRONG_INLINE __device__ half& operator /= (half& a, const half& b) { 244 a = a / b; 245 return a; 246 } 247 EIGEN_STRONG_INLINE __device__ bool operator == (const half& a, const half& b) { 248 return __heq(a, b); 249 } 250 EIGEN_STRONG_INLINE __device__ bool operator != (const half& a, const half& b) { 251 return __hne(a, b); 252 } 253 EIGEN_STRONG_INLINE __device__ bool operator < (const half& a, const half& b) { 254 return __hlt(a, b); 255 } 256 EIGEN_STRONG_INLINE __device__ bool operator <= (const half& a, const half& b) { 257 return __hle(a, b); 258 } 259 EIGEN_STRONG_INLINE __device__ bool operator > (const half& a, const half& b) { 260 return __hgt(a, b); 261 } 262 EIGEN_STRONG_INLINE __device__ bool operator >= (const half& a, const half& b) { 263 return __hge(a, b); 264 } 265 266 #else // Emulate support for half floats 267 268 // Definitions for CPUs and older CUDA, mostly working through conversion 269 // to/from float32_bits. 270 271 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half operator + (const half& a, const half& b) { 272 return half(float(a) + float(b)); 273 } 274 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half operator * (const half& a, const half& b) { 275 return half(float(a) * float(b)); 276 } 277 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half operator - (const half& a, const half& b) { 278 return half(float(a) - float(b)); 279 } 280 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half operator / (const half& a, const half& b) { 281 return half(float(a) / float(b)); 282 } 283 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half operator - (const half& a) { 284 half result; 285 result.x = a.x ^ 0x8000; 286 return result; 287 } 288 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half& operator += (half& a, const half& b) { 289 a = half(float(a) + float(b)); 290 return a; 291 } 292 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half& operator *= (half& a, const half& b) { 293 a = half(float(a) * float(b)); 294 return a; 295 } 296 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half& operator -= (half& a, const half& b) { 297 a = half(float(a) - float(b)); 298 return a; 299 } 300 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half& operator /= (half& a, const half& b) { 301 a = half(float(a) / float(b)); 302 return a; 303 } 304 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bool operator == (const half& a, const half& b) { 305 return numext::equal_strict(float(a),float(b)); 306 } 307 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bool operator != (const half& a, const half& b) { 308 return numext::not_equal_strict(float(a), float(b)); 309 } 310 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bool operator < (const half& a, const half& b) { 311 return float(a) < float(b); 312 } 313 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bool operator <= (const half& a, const half& b) { 314 return float(a) <= float(b); 315 } 316 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bool operator > (const half& a, const half& b) { 317 return float(a) > float(b); 318 } 319 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bool operator >= (const half& a, const half& b) { 320 return float(a) >= float(b); 321 } 322 323 #endif // Emulate support for half floats 324 325 // Division by an index. Do it in full float precision to avoid accuracy 326 // issues in converting the denominator to half. 327 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half operator / (const half& a, Index b) { 328 return half(static_cast<float>(a) / static_cast<float>(b)); 329 } 330 331 // Conversion routines, including fallbacks for the host or older CUDA. 332 // Note that newer Intel CPUs (Haswell or newer) have vectorized versions of 333 // these in hardware. If we need more performance on older/other CPUs, they are 334 // also possible to vectorize directly. 335 336 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC __half_raw raw_uint16_to_half(unsigned short x) { 337 __half_raw h; 338 h.x = x; 339 return h; 340 } 341 342 union float32_bits { 343 unsigned int u; 344 float f; 345 }; 346 347 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC __half_raw float_to_half_rtne(float ff) { 348 #if defined(EIGEN_HAS_CUDA_FP16) && defined(EIGEN_CUDA_ARCH) && EIGEN_CUDA_ARCH >= 300 349 __half tmp_ff = __float2half(ff); 350 return *(__half_raw*)&tmp_ff; 351 352 #elif defined(EIGEN_HAS_FP16_C) 353 __half_raw h; 354 h.x = _cvtss_sh(ff, 0); 355 return h; 356 357 #else 358 float32_bits f; f.f = ff; 359 360 const float32_bits f32infty = { 255 << 23 }; 361 const float32_bits f16max = { (127 + 16) << 23 }; 362 const float32_bits denorm_magic = { ((127 - 15) + (23 - 10) + 1) << 23 }; 363 unsigned int sign_mask = 0x80000000u; 364 __half_raw o; 365 o.x = static_cast<unsigned short>(0x0u); 366 367 unsigned int sign = f.u & sign_mask; 368 f.u ^= sign; 369 370 // NOTE all the integer compares in this function can be safely 371 // compiled into signed compares since all operands are below 372 // 0x80000000. Important if you want fast straight SSE2 code 373 // (since there's no unsigned PCMPGTD). 374 375 if (f.u >= f16max.u) { // result is Inf or NaN (all exponent bits set) 376 o.x = (f.u > f32infty.u) ? 0x7e00 : 0x7c00; // NaN->qNaN and Inf->Inf 377 } else { // (De)normalized number or zero 378 if (f.u < (113 << 23)) { // resulting FP16 is subnormal or zero 379 // use a magic value to align our 10 mantissa bits at the bottom of 380 // the float. as long as FP addition is round-to-nearest-even this 381 // just works. 382 f.f += denorm_magic.f; 383 384 // and one integer subtract of the bias later, we have our final float! 385 o.x = static_cast<unsigned short>(f.u - denorm_magic.u); 386 } else { 387 unsigned int mant_odd = (f.u >> 13) & 1; // resulting mantissa is odd 388 389 // update exponent, rounding bias part 1 390 f.u += ((unsigned int)(15 - 127) << 23) + 0xfff; 391 // rounding bias part 2 392 f.u += mant_odd; 393 // take the bits! 394 o.x = static_cast<unsigned short>(f.u >> 13); 395 } 396 } 397 398 o.x |= static_cast<unsigned short>(sign >> 16); 399 return o; 400 #endif 401 } 402 403 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC float half_to_float(__half_raw h) { 404 #if defined(EIGEN_HAS_CUDA_FP16) && defined(EIGEN_CUDA_ARCH) && EIGEN_CUDA_ARCH >= 300 405 return __half2float(h); 406 407 #elif defined(EIGEN_HAS_FP16_C) 408 return _cvtsh_ss(h.x); 409 410 #else 411 const float32_bits magic = { 113 << 23 }; 412 const unsigned int shifted_exp = 0x7c00 << 13; // exponent mask after shift 413 float32_bits o; 414 415 o.u = (h.x & 0x7fff) << 13; // exponent/mantissa bits 416 unsigned int exp = shifted_exp & o.u; // just the exponent 417 o.u += (127 - 15) << 23; // exponent adjust 418 419 // handle exponent special cases 420 if (exp == shifted_exp) { // Inf/NaN? 421 o.u += (128 - 16) << 23; // extra exp adjust 422 } else if (exp == 0) { // Zero/Denormal? 423 o.u += 1 << 23; // extra exp adjust 424 o.f -= magic.f; // renormalize 425 } 426 427 o.u |= (h.x & 0x8000) << 16; // sign bit 428 return o.f; 429 #endif 430 } 431 432 // --- standard functions --- 433 434 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bool (isinf)(const half& a) { 435 return (a.x & 0x7fff) == 0x7c00; 436 } 437 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bool (isnan)(const half& a) { 438 #if defined(EIGEN_HAS_CUDA_FP16) && defined(EIGEN_CUDA_ARCH) && EIGEN_CUDA_ARCH >= 530 439 return __hisnan(a); 440 #else 441 return (a.x & 0x7fff) > 0x7c00; 442 #endif 443 } 444 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bool (isfinite)(const half& a) { 445 return !(isinf EIGEN_NOT_A_MACRO (a)) && !(isnan EIGEN_NOT_A_MACRO (a)); 446 } 447 448 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half abs(const half& a) { 449 half result; 450 result.x = a.x & 0x7FFF; 451 return result; 452 } 453 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half exp(const half& a) { 454 #if EIGEN_CUDACC_VER >= 80000 && defined EIGEN_CUDA_ARCH && EIGEN_CUDA_ARCH >= 530 455 return half(hexp(a)); 456 #else 457 return half(::expf(float(a))); 458 #endif 459 } 460 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half log(const half& a) { 461 #if defined(EIGEN_HAS_CUDA_FP16) && EIGEN_CUDACC_VER >= 80000 && defined(EIGEN_CUDA_ARCH) && EIGEN_CUDA_ARCH >= 530 462 return half(::hlog(a)); 463 #else 464 return half(::logf(float(a))); 465 #endif 466 } 467 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half log1p(const half& a) { 468 return half(numext::log1p(float(a))); 469 } 470 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half log10(const half& a) { 471 return half(::log10f(float(a))); 472 } 473 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half sqrt(const half& a) { 474 #if EIGEN_CUDACC_VER >= 80000 && defined EIGEN_CUDA_ARCH && EIGEN_CUDA_ARCH >= 530 475 return half(hsqrt(a)); 476 #else 477 return half(::sqrtf(float(a))); 478 #endif 479 } 480 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half pow(const half& a, const half& b) { 481 return half(::powf(float(a), float(b))); 482 } 483 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half sin(const half& a) { 484 return half(::sinf(float(a))); 485 } 486 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half cos(const half& a) { 487 return half(::cosf(float(a))); 488 } 489 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half tan(const half& a) { 490 return half(::tanf(float(a))); 491 } 492 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half tanh(const half& a) { 493 return half(::tanhf(float(a))); 494 } 495 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half floor(const half& a) { 496 #if EIGEN_CUDACC_VER >= 80000 && defined EIGEN_CUDA_ARCH && EIGEN_CUDA_ARCH >= 300 497 return half(hfloor(a)); 498 #else 499 return half(::floorf(float(a))); 500 #endif 501 } 502 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half ceil(const half& a) { 503 #if EIGEN_CUDACC_VER >= 80000 && defined EIGEN_CUDA_ARCH && EIGEN_CUDA_ARCH >= 300 504 return half(hceil(a)); 505 #else 506 return half(::ceilf(float(a))); 507 #endif 508 } 509 510 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half (min)(const half& a, const half& b) { 511 #if defined(EIGEN_HAS_CUDA_FP16) && defined(EIGEN_CUDA_ARCH) && EIGEN_CUDA_ARCH >= 530 512 return __hlt(b, a) ? b : a; 513 #else 514 const float f1 = static_cast<float>(a); 515 const float f2 = static_cast<float>(b); 516 return f2 < f1 ? b : a; 517 #endif 518 } 519 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half (max)(const half& a, const half& b) { 520 #if defined(EIGEN_HAS_CUDA_FP16) && defined(EIGEN_CUDA_ARCH) && EIGEN_CUDA_ARCH >= 530 521 return __hlt(a, b) ? b : a; 522 #else 523 const float f1 = static_cast<float>(a); 524 const float f2 = static_cast<float>(b); 525 return f1 < f2 ? b : a; 526 #endif 527 } 528 529 EIGEN_ALWAYS_INLINE std::ostream& operator << (std::ostream& os, const half& v) { 530 os << static_cast<float>(v); 531 return os; 532 } 533 534 } // end namespace half_impl 535 536 // import Eigen::half_impl::half into Eigen namespace 537 // using half_impl::half; 538 539 namespace internal { 540 541 template<> 542 struct random_default_impl<half, false, false> 543 { 544 static inline half run(const half& x, const half& y) 545 { 546 return x + (y-x) * half(float(std::rand()) / float(RAND_MAX)); 547 } 548 static inline half run() 549 { 550 return run(half(-1.f), half(1.f)); 551 } 552 }; 553 554 template<> struct is_arithmetic<half> { enum { value = true }; }; 555 556 } // end namespace internal 557 558 template<> struct NumTraits<Eigen::half> 559 : GenericNumTraits<Eigen::half> 560 { 561 enum { 562 IsSigned = true, 563 IsInteger = false, 564 IsComplex = false, 565 RequireInitialization = false 566 }; 567 568 EIGEN_DEVICE_FUNC static EIGEN_STRONG_INLINE Eigen::half epsilon() { 569 return half_impl::raw_uint16_to_half(0x0800); 570 } 571 EIGEN_DEVICE_FUNC static EIGEN_STRONG_INLINE Eigen::half dummy_precision() { return Eigen::half(1e-2f); } 572 EIGEN_DEVICE_FUNC static EIGEN_STRONG_INLINE Eigen::half highest() { 573 return half_impl::raw_uint16_to_half(0x7bff); 574 } 575 EIGEN_DEVICE_FUNC static EIGEN_STRONG_INLINE Eigen::half lowest() { 576 return half_impl::raw_uint16_to_half(0xfbff); 577 } 578 EIGEN_DEVICE_FUNC static EIGEN_STRONG_INLINE Eigen::half infinity() { 579 return half_impl::raw_uint16_to_half(0x7c00); 580 } 581 EIGEN_DEVICE_FUNC static EIGEN_STRONG_INLINE Eigen::half quiet_NaN() { 582 return half_impl::raw_uint16_to_half(0x7c01); 583 } 584 }; 585 586 } // end namespace Eigen 587 588 // C-like standard mathematical functions and trancendentals. 589 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC Eigen::half fabsh(const Eigen::half& a) { 590 Eigen::half result; 591 result.x = a.x & 0x7FFF; 592 return result; 593 } 594 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC Eigen::half exph(const Eigen::half& a) { 595 return Eigen::half(::expf(float(a))); 596 } 597 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC Eigen::half logh(const Eigen::half& a) { 598 #if EIGEN_CUDACC_VER >= 80000 && defined(EIGEN_CUDA_ARCH) && EIGEN_CUDA_ARCH >= 530 599 return Eigen::half(::hlog(a)); 600 #else 601 return Eigen::half(::logf(float(a))); 602 #endif 603 } 604 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC Eigen::half sqrth(const Eigen::half& a) { 605 return Eigen::half(::sqrtf(float(a))); 606 } 607 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC Eigen::half powh(const Eigen::half& a, const Eigen::half& b) { 608 return Eigen::half(::powf(float(a), float(b))); 609 } 610 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC Eigen::half floorh(const Eigen::half& a) { 611 return Eigen::half(::floorf(float(a))); 612 } 613 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC Eigen::half ceilh(const Eigen::half& a) { 614 return Eigen::half(::ceilf(float(a))); 615 } 616 617 namespace std { 618 619 #if __cplusplus > 199711L 620 template <> 621 struct hash<Eigen::half> { 622 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE std::size_t operator()(const Eigen::half& a) const { 623 return static_cast<std::size_t>(a.x); 624 } 625 }; 626 #endif 627 628 } // end namespace std 629 630 631 // Add the missing shfl_xor intrinsic 632 #if defined(EIGEN_CUDA_ARCH) && EIGEN_CUDA_ARCH >= 300 633 __device__ EIGEN_STRONG_INLINE Eigen::half __shfl_xor(Eigen::half var, int laneMask, int width=warpSize) { 634 #if EIGEN_CUDACC_VER < 90000 635 return static_cast<Eigen::half>(__shfl_xor(static_cast<float>(var), laneMask, width)); 636 #else 637 return static_cast<Eigen::half>(__shfl_xor_sync(0xFFFFFFFF, static_cast<float>(var), laneMask, width)); 638 #endif 639 } 640 #endif 641 642 // ldg() has an overload for __half_raw, but we also need one for Eigen::half. 643 #if defined(EIGEN_CUDA_ARCH) && EIGEN_CUDA_ARCH >= 350 644 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC Eigen::half __ldg(const Eigen::half* ptr) { 645 return Eigen::half_impl::raw_uint16_to_half( 646 __ldg(reinterpret_cast<const unsigned short*>(ptr))); 647 } 648 #endif 649 650 651 #if defined(EIGEN_CUDA_ARCH) 652 namespace Eigen { 653 namespace numext { 654 655 template<> 656 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE 657 bool (isnan)(const Eigen::half& h) { 658 return (half_impl::isnan)(h); 659 } 660 661 template<> 662 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE 663 bool (isinf)(const Eigen::half& h) { 664 return (half_impl::isinf)(h); 665 } 666 667 template<> 668 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE 669 bool (isfinite)(const Eigen::half& h) { 670 return (half_impl::isfinite)(h); 671 } 672 673 } // namespace Eigen 674 } // namespace numext 675 #endif 676 677 #endif // EIGEN_HALF_CUDA_H 678