1 #pragma once 2 #ifndef FP16_BITCASTS_H 3 #define FP16_BITCASTS_H 4 5 #if defined(__cplusplus) && (__cplusplus >= 201103L) 6 #include <cstdint> 7 #elif !defined(__OPENCL_VERSION__) 8 #include <stdint.h> 9 #endif 10 11 #if defined(__INTEL_COMPILER) 12 #include <immintrin.h> 13 #endif 14 15 #if defined(_MSC_VER) && (defined(_M_ARM) || defined(_M_ARM64)) 16 #include <intrin.h> 17 #endif 18 19 fp32_from_bits(uint32_t w)20static inline float fp32_from_bits(uint32_t w) { 21 #if defined(__OPENCL_VERSION__) 22 return as_float(w); 23 #elif defined(__CUDA_ARCH__) 24 return __uint_as_float((unsigned int) w); 25 #elif defined(__INTEL_COMPILER) 26 return _castu32_f32(w); 27 #elif defined(_MSC_VER) && (defined(_M_ARM) || defined(_M_ARM64)) 28 return _CopyFloatFromInt32((__int32) w); 29 #else 30 union { 31 uint32_t as_bits; 32 float as_value; 33 } fp32 = { w }; 34 return fp32.as_value; 35 #endif 36 } 37 fp32_to_bits(float f)38static inline uint32_t fp32_to_bits(float f) { 39 #if defined(__OPENCL_VERSION__) 40 return as_uint(f); 41 #elif defined(__CUDA_ARCH__) 42 return (uint32_t) __float_as_uint(f); 43 #elif defined(__INTEL_COMPILER) 44 return _castf32_u32(f); 45 #elif defined(_MSC_VER) && (defined(_M_ARM) || defined(_M_ARM64)) 46 return (uint32_t) _CopyInt32FromFloat(f); 47 #else 48 union { 49 float as_value; 50 uint32_t as_bits; 51 } fp32 = { f }; 52 return fp32.as_bits; 53 #endif 54 } 55 fp64_from_bits(uint64_t w)56static inline double fp64_from_bits(uint64_t w) { 57 #if defined(__OPENCL_VERSION__) 58 return as_double(w); 59 #elif defined(__CUDA_ARCH__) 60 return __longlong_as_double((long long) w); 61 #elif defined(__INTEL_COMPILER) 62 return _castu64_f64(w); 63 #elif defined(_MSC_VER) && (defined(_M_ARM) || defined(_M_ARM64)) 64 return _CopyDoubleFromInt64((__int64) w); 65 #else 66 union { 67 uint64_t as_bits; 68 double as_value; 69 } fp64 = { w }; 70 return fp64.as_value; 71 #endif 72 } 73 fp64_to_bits(double f)74static inline uint64_t fp64_to_bits(double f) { 75 #if defined(__OPENCL_VERSION__) 76 return as_ulong(f); 77 #elif defined(__CUDA_ARCH__) 78 return (uint64_t) __double_as_longlong(f); 79 #elif defined(__INTEL_COMPILER) 80 return _castf64_u64(f); 81 #elif defined(_MSC_VER) && (defined(_M_ARM) || defined(_M_ARM64)) 82 return (uint64_t) _CopyInt64FromDouble(f); 83 #else 84 union { 85 double as_value; 86 uint64_t as_bits; 87 } fp64 = { f }; 88 return fp64.as_bits; 89 #endif 90 } 91 92 #endif /* FP16_BITCASTS_H */ 93