|
#pragma once |
|
#ifndef FP16_BITCASTS_H |
|
#define FP16_BITCASTS_H |
|
|
|
#if defined(__cplusplus) && (__cplusplus >= 201103L) |
|
#include <cstdint> |
|
#elif !defined(__OPENCL_VERSION__) |
|
#include <stdint.h> |
|
#endif |
|
|
|
#if defined(__INTEL_COMPILER) |
|
#include <immintrin.h> |
|
#endif |
|
|
|
#if defined(_MSC_VER) && (defined(_M_ARM) || defined(_M_ARM64)) |
|
#include <intrin.h> |
|
#endif |
|
|
|
|
|
static inline float fp32_from_bits(uint32_t w) { |
|
#if defined(__OPENCL_VERSION__) |
|
return as_float(w); |
|
#elif defined(__CUDA_ARCH__) |
|
return __uint_as_float((unsigned int) w); |
|
#elif defined(__INTEL_COMPILER) |
|
return _castu32_f32(w); |
|
#elif defined(_MSC_VER) && (defined(_M_ARM) || defined(_M_ARM64)) |
|
return _CopyFloatFromInt32((__int32) w); |
|
#else |
|
union { |
|
uint32_t as_bits; |
|
float as_value; |
|
} fp32 = { w }; |
|
return fp32.as_value; |
|
#endif |
|
} |
|
|
|
static inline uint32_t fp32_to_bits(float f) { |
|
#if defined(__OPENCL_VERSION__) |
|
return as_uint(f); |
|
#elif defined(__CUDA_ARCH__) |
|
return (uint32_t) __float_as_uint(f); |
|
#elif defined(__INTEL_COMPILER) |
|
return _castf32_u32(f); |
|
#elif defined(_MSC_VER) && (defined(_M_ARM) || defined(_M_ARM64)) |
|
return (uint32_t) _CopyInt32FromFloat(f); |
|
#else |
|
union { |
|
float as_value; |
|
uint32_t as_bits; |
|
} fp32 = { f }; |
|
return fp32.as_bits; |
|
#endif |
|
} |
|
|
|
static inline double fp64_from_bits(uint64_t w) { |
|
#if defined(__OPENCL_VERSION__) |
|
return as_double(w); |
|
#elif defined(__CUDA_ARCH__) |
|
return __longlong_as_double((long long) w); |
|
#elif defined(__INTEL_COMPILER) |
|
return _castu64_f64(w); |
|
#elif defined(_MSC_VER) && (defined(_M_ARM) || defined(_M_ARM64)) |
|
return _CopyDoubleFromInt64((__int64) w); |
|
#else |
|
union { |
|
uint64_t as_bits; |
|
double as_value; |
|
} fp64 = { w }; |
|
return fp64.as_value; |
|
#endif |
|
} |
|
|
|
static inline uint64_t fp64_to_bits(double f) { |
|
#if defined(__OPENCL_VERSION__) |
|
return as_ulong(f); |
|
#elif defined(__CUDA_ARCH__) |
|
return (uint64_t) __double_as_longlong(f); |
|
#elif defined(__INTEL_COMPILER) |
|
return _castf64_u64(f); |
|
#elif defined(_MSC_VER) && (defined(_M_ARM) || defined(_M_ARM64)) |
|
return (uint64_t) _CopyInt64FromDouble(f); |
|
#else |
|
union { |
|
double as_value; |
|
uint64_t as_bits; |
|
} fp64 = { f }; |
|
return fp64.as_bits; |
|
#endif |
|
} |
|
|
|
#endif |
|
|
