![]() |
Eigen
3.3.3
|
00001 // This file is part of Eigen, a lightweight C++ template library 00002 // for linear algebra. 00003 // 00004 // This Source Code Form is subject to the terms of the Mozilla 00005 // Public License v. 2.0. If a copy of the MPL was not distributed 00006 // with this file, You can obtain one at http://mozilla.org/MPL/2.0/. 00007 // 00008 // The conversion routines are Copyright (c) Fabian Giesen, 2016. 00009 // The original license follows: 00010 // 00011 // Copyright (c) Fabian Giesen, 2016 00012 // All rights reserved. 00013 // Redistribution and use in source and binary forms, with or without 00014 // modification, are permitted. 00015 // THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS 00016 // “AS IS” AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT 00017 // LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR 00018 // A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT 00019 // HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, 00020 // SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT 00021 // LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, 00022 // DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY 00023 // THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT 00024 // (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE 00025 // OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. 00026 00027 00028 // Standard 16-bit float type, mostly useful for GPUs. Defines a new 00029 // type Eigen::half (inheriting from CUDA's __half struct) with 00030 // operator overloads such that it behaves basically as an arithmetic 00031 // type. It will be quite slow on CPUs (so it is recommended to stay 00032 // in fp32 for CPUs, except for simple parameter conversions, I/O 00033 // to disk and the likes), but fast on GPUs. 00034 00035 00036 #ifndef EIGEN_HALF_CUDA_H 00037 #define EIGEN_HALF_CUDA_H 00038 00039 #if __cplusplus > 199711L 00040 #define EIGEN_EXPLICIT_CAST(tgt_type) explicit operator tgt_type() 00041 #else 00042 #define EIGEN_EXPLICIT_CAST(tgt_type) operator tgt_type() 00043 #endif 00044 00045 00046 namespace Eigen { 00047 00048 struct half; 00049 00050 namespace half_impl { 00051 00052 #if !defined(EIGEN_HAS_CUDA_FP16) 00053 00054 // Make our own __half definition that is similar to CUDA's. 00055 struct __half { 00056 EIGEN_DEVICE_FUNC __half() {} 00057 explicit EIGEN_DEVICE_FUNC __half(unsigned short raw) : x(raw) {} 00058 unsigned short x; 00059 }; 00060 00061 #endif 00062 00063 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC __half raw_uint16_to_half(unsigned short x); 00064 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC __half float_to_half_rtne(float ff); 00065 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC float half_to_float(__half h); 00066 00067 struct half_base : public __half { 00068 EIGEN_DEVICE_FUNC half_base() {} 00069 EIGEN_DEVICE_FUNC half_base(const half_base& h) : __half(h) {} 00070 EIGEN_DEVICE_FUNC half_base(const __half& h) : __half(h) {} 00071 }; 00072 00073 } // namespace half_impl 00074 00075 // Class definition. 00076 struct half : public half_impl::half_base { 00077 #if !defined(EIGEN_HAS_CUDA_FP16) 00078 typedef half_impl::__half __half; 00079 #endif 00080 00081 EIGEN_DEVICE_FUNC half() {} 00082 00083 EIGEN_DEVICE_FUNC half(const __half& h) : half_impl::half_base(h) {} 00084 EIGEN_DEVICE_FUNC half(const half& h) : half_impl::half_base(h) {} 00085 00086 explicit EIGEN_DEVICE_FUNC half(bool b) 00087 : half_impl::half_base(half_impl::raw_uint16_to_half(b ? 0x3c00 : 0)) {} 00088 template<class T> 00089 explicit EIGEN_DEVICE_FUNC half(const T& val) 00090 : half_impl::half_base(half_impl::float_to_half_rtne(static_cast<float>(val))) {} 00091 explicit EIGEN_DEVICE_FUNC half(float f) 00092 : half_impl::half_base(half_impl::float_to_half_rtne(f)) {} 00093 00094 EIGEN_DEVICE_FUNC EIGEN_EXPLICIT_CAST(bool) const { 00095 // +0.0 and -0.0 become false, everything else becomes true. 00096 return (x & 0x7fff) != 0; 00097 } 00098 EIGEN_DEVICE_FUNC EIGEN_EXPLICIT_CAST(signed char) const { 00099 return static_cast<signed char>(half_impl::half_to_float(*this)); 00100 } 00101 EIGEN_DEVICE_FUNC EIGEN_EXPLICIT_CAST(unsigned char) const { 00102 return static_cast<unsigned char>(half_impl::half_to_float(*this)); 00103 } 00104 EIGEN_DEVICE_FUNC EIGEN_EXPLICIT_CAST(short) const { 00105 return static_cast<short>(half_impl::half_to_float(*this)); 00106 } 00107 EIGEN_DEVICE_FUNC EIGEN_EXPLICIT_CAST(unsigned short) const { 00108 return static_cast<unsigned short>(half_impl::half_to_float(*this)); 00109 } 00110 EIGEN_DEVICE_FUNC EIGEN_EXPLICIT_CAST(int) const { 00111 return static_cast<int>(half_impl::half_to_float(*this)); 00112 } 00113 EIGEN_DEVICE_FUNC EIGEN_EXPLICIT_CAST(unsigned int) const { 00114 return static_cast<unsigned int>(half_impl::half_to_float(*this)); 00115 } 00116 EIGEN_DEVICE_FUNC EIGEN_EXPLICIT_CAST(long) const { 00117 return static_cast<long>(half_impl::half_to_float(*this)); 00118 } 00119 EIGEN_DEVICE_FUNC EIGEN_EXPLICIT_CAST(unsigned long) const { 00120 return static_cast<unsigned long>(half_impl::half_to_float(*this)); 00121 } 00122 EIGEN_DEVICE_FUNC EIGEN_EXPLICIT_CAST(long long) const { 00123 return static_cast<long long>(half_impl::half_to_float(*this)); 00124 } 00125 EIGEN_DEVICE_FUNC EIGEN_EXPLICIT_CAST(unsigned long long) const { 00126 return static_cast<unsigned long long>(half_to_float(*this)); 00127 } 00128 EIGEN_DEVICE_FUNC EIGEN_EXPLICIT_CAST(float) const { 00129 return half_impl::half_to_float(*this); 00130 } 00131 EIGEN_DEVICE_FUNC EIGEN_EXPLICIT_CAST(double) const { 00132 return static_cast<double>(half_impl::half_to_float(*this)); 00133 } 00134 00135 EIGEN_DEVICE_FUNC half& operator=(const half& other) { 00136 x = other.x; 00137 return *this; 00138 } 00139 }; 00140 00141 namespace half_impl { 00142 00143 #if defined(EIGEN_HAS_CUDA_FP16) && defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 530 00144 00145 // Intrinsics for native fp16 support. Note that on current hardware, 00146 // these are no faster than fp32 arithmetic (you need to use the half2 00147 // versions to get the ALU speed increased), but you do save the 00148 // conversion steps back and forth. 00149 00150 __device__ half operator + (const half& a, const half& b) { 00151 return __hadd(a, b); 00152 } 00153 __device__ half operator * (const half& a, const half& b) { 00154 return __hmul(a, b); 00155 } 00156 __device__ half operator - (const half& a, const half& b) { 00157 return __hsub(a, b); 00158 } 00159 __device__ half operator / (const half& a, const half& b) { 00160 float num = __half2float(a); 00161 float denom = __half2float(b); 00162 return __float2half(num / denom); 00163 } 00164 __device__ half operator - (const half& a) { 00165 return __hneg(a); 00166 } 00167 __device__ half& operator += (half& a, const half& b) { 00168 a = a + b; 00169 return a; 00170 } 00171 __device__ half& operator *= (half& a, const half& b) { 00172 a = a * b; 00173 return a; 00174 } 00175 __device__ half& operator -= (half& a, const half& b) { 00176 a = a - b; 00177 return a; 00178 } 00179 __device__ half& operator /= (half& a, const half& b) { 00180 a = a / b; 00181 return a; 00182 } 00183 __device__ bool operator == (const half& a, const half& b) { 00184 return __heq(a, b); 00185 } 00186 __device__ bool operator != (const half& a, const half& b) { 00187 return __hne(a, b); 00188 } 00189 __device__ bool operator < (const half& a, const half& b) { 00190 return __hlt(a, b); 00191 } 00192 __device__ bool operator <= (const half& a, const half& b) { 00193 return __hle(a, b); 00194 } 00195 __device__ bool operator > (const half& a, const half& b) { 00196 return __hgt(a, b); 00197 } 00198 __device__ bool operator >= (const half& a, const half& b) { 00199 return __hge(a, b); 00200 } 00201 00202 #else // Emulate support for half floats 00203 00204 // Definitions for CPUs and older CUDA, mostly working through conversion 00205 // to/from fp32. 00206 00207 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half operator + (const half& a, const half& b) { 00208 return half(float(a) + float(b)); 00209 } 00210 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half operator * (const half& a, const half& b) { 00211 return half(float(a) * float(b)); 00212 } 00213 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half operator - (const half& a, const half& b) { 00214 return half(float(a) - float(b)); 00215 } 00216 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half operator / (const half& a, const half& b) { 00217 return half(float(a) / float(b)); 00218 } 00219 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half operator - (const half& a) { 00220 half result; 00221 result.x = a.x ^ 0x8000; 00222 return result; 00223 } 00224 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half& operator += (half& a, const half& b) { 00225 a = half(float(a) + float(b)); 00226 return a; 00227 } 00228 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half& operator *= (half& a, const half& b) { 00229 a = half(float(a) * float(b)); 00230 return a; 00231 } 00232 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half& operator -= (half& a, const half& b) { 00233 a = half(float(a) - float(b)); 00234 return a; 00235 } 00236 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half& operator /= (half& a, const half& b) { 00237 a = half(float(a) / float(b)); 00238 return a; 00239 } 00240 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bool operator == (const half& a, const half& b) { 00241 return float(a) == float(b); 00242 } 00243 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bool operator != (const half& a, const half& b) { 00244 return float(a) != float(b); 00245 } 00246 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bool operator < (const half& a, const half& b) { 00247 return float(a) < float(b); 00248 } 00249 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bool operator <= (const half& a, const half& b) { 00250 return float(a) <= float(b); 00251 } 00252 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bool operator > (const half& a, const half& b) { 00253 return float(a) > float(b); 00254 } 00255 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bool operator >= (const half& a, const half& b) { 00256 return float(a) >= float(b); 00257 } 00258 00259 #endif // Emulate support for half floats 00260 00261 // Division by an index. Do it in full float precision to avoid accuracy 00262 // issues in converting the denominator to half. 00263 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half operator / (const half& a, Index b) { 00264 return half(static_cast<float>(a) / static_cast<float>(b)); 00265 } 00266 00267 // Conversion routines, including fallbacks for the host or older CUDA. 00268 // Note that newer Intel CPUs (Haswell or newer) have vectorized versions of 00269 // these in hardware. If we need more performance on older/other CPUs, they are 00270 // also possible to vectorize directly. 00271 00272 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC __half raw_uint16_to_half(unsigned short x) { 00273 __half h; 00274 h.x = x; 00275 return h; 00276 } 00277 00278 union FP32 { 00279 unsigned int u; 00280 float f; 00281 }; 00282 00283 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC __half float_to_half_rtne(float ff) { 00284 #if defined(EIGEN_HAS_CUDA_FP16) && defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 300 00285 return __float2half(ff); 00286 00287 #elif defined(EIGEN_HAS_FP16_C) 00288 __half h; 00289 h.x = _cvtss_sh(ff, 0); 00290 return h; 00291 00292 #else 00293 FP32 f; f.f = ff; 00294 00295 const FP32 f32infty = { 255 << 23 }; 00296 const FP32 f16max = { (127 + 16) << 23 }; 00297 const FP32 denorm_magic = { ((127 - 15) + (23 - 10) + 1) << 23 }; 00298 unsigned int sign_mask = 0x80000000u; 00299 __half o; 00300 o.x = static_cast<unsigned short>(0x0u); 00301 00302 unsigned int sign = f.u & sign_mask; 00303 f.u ^= sign; 00304 00305 // NOTE all the integer compares in this function can be safely 00306 // compiled into signed compares since all operands are below 00307 // 0x80000000. Important if you want fast straight SSE2 code 00308 // (since there's no unsigned PCMPGTD). 00309 00310 if (f.u >= f16max.u) { // result is Inf or NaN (all exponent bits set) 00311 o.x = (f.u > f32infty.u) ? 0x7e00 : 0x7c00; // NaN->qNaN and Inf->Inf 00312 } else { // (De)normalized number or zero 00313 if (f.u < (113 << 23)) { // resulting FP16 is subnormal or zero 00314 // use a magic value to align our 10 mantissa bits at the bottom of 00315 // the float. as long as FP addition is round-to-nearest-even this 00316 // just works. 00317 f.f += denorm_magic.f; 00318 00319 // and one integer subtract of the bias later, we have our final float! 00320 o.x = static_cast<unsigned short>(f.u - denorm_magic.u); 00321 } else { 00322 unsigned int mant_odd = (f.u >> 13) & 1; // resulting mantissa is odd 00323 00324 // update exponent, rounding bias part 1 00325 f.u += ((unsigned int)(15 - 127) << 23) + 0xfff; 00326 // rounding bias part 2 00327 f.u += mant_odd; 00328 // take the bits! 00329 o.x = static_cast<unsigned short>(f.u >> 13); 00330 } 00331 } 00332 00333 o.x |= static_cast<unsigned short>(sign >> 16); 00334 return o; 00335 #endif 00336 } 00337 00338 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC float half_to_float(__half h) { 00339 #if defined(EIGEN_HAS_CUDA_FP16) && defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 300 00340 return __half2float(h); 00341 00342 #elif defined(EIGEN_HAS_FP16_C) 00343 return _cvtsh_ss(h.x); 00344 00345 #else 00346 const FP32 magic = { 113 << 23 }; 00347 const unsigned int shifted_exp = 0x7c00 << 13; // exponent mask after shift 00348 FP32 o; 00349 00350 o.u = (h.x & 0x7fff) << 13; // exponent/mantissa bits 00351 unsigned int exp = shifted_exp & o.u; // just the exponent 00352 o.u += (127 - 15) << 23; // exponent adjust 00353 00354 // handle exponent special cases 00355 if (exp == shifted_exp) { // Inf/NaN? 00356 o.u += (128 - 16) << 23; // extra exp adjust 00357 } else if (exp == 0) { // Zero/Denormal? 00358 o.u += 1 << 23; // extra exp adjust 00359 o.f -= magic.f; // renormalize 00360 } 00361 00362 o.u |= (h.x & 0x8000) << 16; // sign bit 00363 return o.f; 00364 #endif 00365 } 00366 00367 // --- standard functions --- 00368 00369 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bool (isinf)(const half& a) { 00370 return (a.x & 0x7fff) == 0x7c00; 00371 } 00372 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bool (isnan)(const half& a) { 00373 #if defined(EIGEN_HAS_CUDA_FP16) && defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 530 00374 return __hisnan(a); 00375 #else 00376 return (a.x & 0x7fff) > 0x7c00; 00377 #endif 00378 } 00379 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC bool (isfinite)(const half& a) { 00380 return !(isinf EIGEN_NOT_A_MACRO (a)) && !(isnan EIGEN_NOT_A_MACRO (a)); 00381 } 00382 00383 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half abs(const half& a) { 00384 half result; 00385 result.x = a.x & 0x7FFF; 00386 return result; 00387 } 00388 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half exp(const half& a) { 00389 return half(::expf(float(a))); 00390 } 00391 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half log(const half& a) { 00392 #if defined(EIGEN_HAS_CUDA_FP16) && defined __CUDACC_VER__ && __CUDACC_VER__ >= 80000 && defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 530 00393 return Eigen::half(::hlog(a)); 00394 #else 00395 return half(::logf(float(a))); 00396 #endif 00397 } 00398 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half log1p(const half& a) { 00399 return half(numext::log1p(float(a))); 00400 } 00401 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half log10(const half& a) { 00402 return half(::log10f(float(a))); 00403 } 00404 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half sqrt(const half& a) { 00405 return half(::sqrtf(float(a))); 00406 } 00407 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half pow(const half& a, const half& b) { 00408 return half(::powf(float(a), float(b))); 00409 } 00410 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half sin(const half& a) { 00411 return half(::sinf(float(a))); 00412 } 00413 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half cos(const half& a) { 00414 return half(::cosf(float(a))); 00415 } 00416 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half tan(const half& a) { 00417 return half(::tanf(float(a))); 00418 } 00419 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half tanh(const half& a) { 00420 return half(::tanhf(float(a))); 00421 } 00422 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half floor(const half& a) { 00423 return half(::floorf(float(a))); 00424 } 00425 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half ceil(const half& a) { 00426 return half(::ceilf(float(a))); 00427 } 00428 00429 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half (min)(const half& a, const half& b) { 00430 #if defined(EIGEN_HAS_CUDA_FP16) && defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 530 00431 return __hlt(b, a) ? b : a; 00432 #else 00433 const float f1 = static_cast<float>(a); 00434 const float f2 = static_cast<float>(b); 00435 return f2 < f1 ? b : a; 00436 #endif 00437 } 00438 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC half (max)(const half& a, const half& b) { 00439 #if defined(EIGEN_HAS_CUDA_FP16) && defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 530 00440 return __hlt(a, b) ? b : a; 00441 #else 00442 const float f1 = static_cast<float>(a); 00443 const float f2 = static_cast<float>(b); 00444 return f1 < f2 ? b : a; 00445 #endif 00446 } 00447 00448 EIGEN_ALWAYS_INLINE std::ostream& operator << (std::ostream& os, const half& v) { 00449 os << static_cast<float>(v); 00450 return os; 00451 } 00452 00453 } // end namespace half_impl 00454 00455 // import Eigen::half_impl::half into Eigen namespace 00456 // using half_impl::half; 00457 00458 namespace internal { 00459 00460 template<> 00461 struct random_default_impl<half, false, false> 00462 { 00463 static inline half run(const half& x, const half& y) 00464 { 00465 return x + (y-x) * half(float(std::rand()) / float(RAND_MAX)); 00466 } 00467 static inline half run() 00468 { 00469 return run(half(-1.f), half(1.f)); 00470 } 00471 }; 00472 00473 template<> struct is_arithmetic<half> { enum { value = true }; }; 00474 00475 } // end namespace internal 00476 00477 template<> struct NumTraits<Eigen::half> 00478 : GenericNumTraits<Eigen::half> 00479 { 00480 EIGEN_DEVICE_FUNC static EIGEN_STRONG_INLINE Eigen::half epsilon() { 00481 return half_impl::raw_uint16_to_half(0x0800); 00482 } 00483 EIGEN_DEVICE_FUNC static EIGEN_STRONG_INLINE Eigen::half dummy_precision() { return Eigen::half(1e-2f); } 00484 EIGEN_DEVICE_FUNC static EIGEN_STRONG_INLINE Eigen::half highest() { 00485 return half_impl::raw_uint16_to_half(0x7bff); 00486 } 00487 EIGEN_DEVICE_FUNC static EIGEN_STRONG_INLINE Eigen::half lowest() { 00488 return half_impl::raw_uint16_to_half(0xfbff); 00489 } 00490 EIGEN_DEVICE_FUNC static EIGEN_STRONG_INLINE Eigen::half infinity() { 00491 return half_impl::raw_uint16_to_half(0x7c00); 00492 } 00493 EIGEN_DEVICE_FUNC static EIGEN_STRONG_INLINE Eigen::half quiet_NaN() { 00494 return half_impl::raw_uint16_to_half(0x7c01); 00495 } 00496 }; 00497 00498 } // end namespace Eigen 00499 00500 // C-like standard mathematical functions and trancendentals. 00501 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC Eigen::half fabsh(const Eigen::half& a) { 00502 Eigen::half result; 00503 result.x = a.x & 0x7FFF; 00504 return result; 00505 } 00506 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC Eigen::half exph(const Eigen::half& a) { 00507 return Eigen::half(::expf(float(a))); 00508 } 00509 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC Eigen::half logh(const Eigen::half& a) { 00510 #if defined __CUDACC_VER__ && __CUDACC_VER__ >= 80000 && defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 530 00511 return Eigen::half(::hlog(a)); 00512 #else 00513 return Eigen::half(::logf(float(a))); 00514 #endif 00515 } 00516 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC Eigen::half sqrth(const Eigen::half& a) { 00517 return Eigen::half(::sqrtf(float(a))); 00518 } 00519 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC Eigen::half powh(const Eigen::half& a, const Eigen::half& b) { 00520 return Eigen::half(::powf(float(a), float(b))); 00521 } 00522 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC Eigen::half floorh(const Eigen::half& a) { 00523 return Eigen::half(::floorf(float(a))); 00524 } 00525 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC Eigen::half ceilh(const Eigen::half& a) { 00526 return Eigen::half(::ceilf(float(a))); 00527 } 00528 00529 namespace std { 00530 00531 #if __cplusplus > 199711L 00532 template <> 00533 struct hash<Eigen::half> { 00534 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE std::size_t operator()(const Eigen::half& a) const { 00535 return static_cast<std::size_t>(a.x); 00536 } 00537 }; 00538 #endif 00539 00540 } // end namespace std 00541 00542 00543 // Add the missing shfl_xor intrinsic 00544 #if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 300 00545 __device__ EIGEN_STRONG_INLINE Eigen::half __shfl_xor(Eigen::half var, int laneMask, int width=warpSize) { 00546 return static_cast<Eigen::half>(__shfl_xor(static_cast<float>(var), laneMask, width)); 00547 } 00548 #endif 00549 00550 // ldg() has an overload for __half, but we also need one for Eigen::half. 00551 #if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 350 00552 EIGEN_STRONG_INLINE EIGEN_DEVICE_FUNC Eigen::half __ldg(const Eigen::half* ptr) { 00553 return Eigen::half_impl::raw_uint16_to_half( 00554 __ldg(reinterpret_cast<const unsigned short*>(ptr))); 00555 } 00556 #endif 00557 00558 00559 #if defined(__CUDA_ARCH__) 00560 namespace Eigen { 00561 namespace numext { 00562 00563 template<> 00564 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE 00565 bool (isnan)(const Eigen::half& h) { 00566 return (half_impl::isnan)(h); 00567 } 00568 00569 template<> 00570 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE 00571 bool (isinf)(const Eigen::half& h) { 00572 return (half_impl::isinf)(h); 00573 } 00574 00575 template<> 00576 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE 00577 bool (isfinite)(const Eigen::half& h) { 00578 return (half_impl::isfinite)(h); 00579 } 00580 00581 } // namespace Eigen 00582 } // namespace numext 00583 #endif 00584 00585 #endif // EIGEN_HALF_CUDA_H