Eigen  3.3.3
Half.h
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
 All Classes Functions Variables Typedefs Enumerations Enumerator Friends