diff --git a/.ci/docker/ci_commit_pins/pytorch.txt b/.ci/docker/ci_commit_pins/pytorch.txt index 6305196d2ad..1082cb4d2d1 100644 --- a/.ci/docker/ci_commit_pins/pytorch.txt +++ b/.ci/docker/ci_commit_pins/pytorch.txt @@ -1 +1 @@ -6fc0ad22f0a07b6f38d138861c56a765d5a9bb02 +e7152ff8a6a929a0db7f3f4a72a5b6d471769cd3 diff --git a/install_requirements.py b/install_requirements.py index 40169a17f3b..17753180bf9 100644 --- a/install_requirements.py +++ b/install_requirements.py @@ -71,7 +71,7 @@ def python_is_compatible(): # # NOTE: If you're changing, make the corresponding change in .ci/docker/ci_commit_pins/pytorch.txt # by picking the hash from the same date in https://hud.pytorch.org/hud/pytorch/pytorch/nightly/ -NIGHTLY_VERSION = "dev20250725" +NIGHTLY_VERSION = "dev20250811" def install_requirements(use_pytorch_nightly): diff --git a/runtime/core/portable_type/c10/c10/util/BFloat16-inl.h b/runtime/core/portable_type/c10/c10/util/BFloat16-inl.h index 1ed866f78d9..6d3510cd5be 100644 --- a/runtime/core/portable_type/c10/c10/util/BFloat16-inl.h +++ b/runtime/core/portable_type/c10/c10/util/BFloat16-inl.h @@ -1,340 +1 @@ -#pragma once - -#include -#include - -#include - -C10_CLANG_DIAGNOSTIC_PUSH() -#if C10_CLANG_HAS_WARNING("-Wimplicit-int-float-conversion") -C10_CLANG_DIAGNOSTIC_IGNORE("-Wimplicit-int-float-conversion") -#endif - -#if defined(CL_SYCL_LANGUAGE_VERSION) -#include // for SYCL 1.2.1 -#elif defined(SYCL_LANGUAGE_VERSION) -#include // for SYCL 2020 -#endif - -namespace c10 { - -/// Constructors -inline C10_HOST_DEVICE BFloat16::BFloat16(float value) - : -#if defined(__CUDACC__) && !defined(USE_ROCM) && defined(__CUDA_ARCH__) && \ - __CUDA_ARCH__ >= 800 - x(__bfloat16_as_ushort(__float2bfloat16(value))) -#elif defined(__SYCL_DEVICE_ONLY__) && \ - defined(SYCL_EXT_ONEAPI_BFLOAT16_MATH_FUNCTIONS) - x(c10::bit_cast(sycl::ext::oneapi::bfloat16(value))) -#else - // RNE by default - x(detail::round_to_nearest_even(value)) -#endif -{ -} - -/// Implicit conversions -inline C10_HOST_DEVICE BFloat16::operator float() const { -#if defined(__CUDACC__) && !defined(USE_ROCM) - return __bfloat162float(*reinterpret_cast(&x)); -#elif defined(__SYCL_DEVICE_ONLY__) && \ - defined(SYCL_EXT_ONEAPI_BFLOAT16_MATH_FUNCTIONS) - return float(*reinterpret_cast(&x)); -#else - return detail::f32_from_bits(x); -#endif -} - -#if defined(__CUDACC__) && !defined(USE_ROCM) -inline C10_HOST_DEVICE BFloat16::BFloat16(const __nv_bfloat16& value) { - x = *reinterpret_cast(&value); -} -inline C10_HOST_DEVICE BFloat16::operator __nv_bfloat16() const { - return *reinterpret_cast(&x); -} -#endif - -#if defined(SYCL_EXT_ONEAPI_BFLOAT16_MATH_FUNCTIONS) -inline C10_HOST_DEVICE BFloat16::BFloat16( - const sycl::ext::oneapi::bfloat16& value) { - x = *reinterpret_cast(&value); -} -inline C10_HOST_DEVICE BFloat16::operator sycl::ext::oneapi::bfloat16() const { - return *reinterpret_cast(&x); -} -#endif - -// CUDA intrinsics - -#if defined(__CUDACC__) || defined(__HIPCC__) -inline C10_DEVICE BFloat16 __ldg(const BFloat16* ptr) { -#if !defined(USE_ROCM) && defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 800 - return __ldg(reinterpret_cast(ptr)); -#else - return *ptr; -#endif -} -#endif - -/// Arithmetic - -inline C10_HOST_DEVICE BFloat16 -operator+(const BFloat16& a, const BFloat16& b) { - return static_cast(a) + static_cast(b); -} - -inline C10_HOST_DEVICE BFloat16 -operator-(const BFloat16& a, const BFloat16& b) { - return static_cast(a) - static_cast(b); -} - -inline C10_HOST_DEVICE BFloat16 -operator*(const BFloat16& a, const BFloat16& b) { - return static_cast(a) * static_cast(b); -} - -inline C10_HOST_DEVICE BFloat16 operator/(const BFloat16& a, const BFloat16& b) - __ubsan_ignore_float_divide_by_zero__ { - return static_cast(a) / static_cast(b); -} - -inline C10_HOST_DEVICE BFloat16 operator-(const BFloat16& a) { - return -static_cast(a); -} - -inline C10_HOST_DEVICE BFloat16& operator+=(BFloat16& a, const BFloat16& b) { - a = a + b; - return a; -} - -inline C10_HOST_DEVICE BFloat16& operator-=(BFloat16& a, const BFloat16& b) { - a = a - b; - return a; -} - -inline C10_HOST_DEVICE BFloat16& operator*=(BFloat16& a, const BFloat16& b) { - a = a * b; - return a; -} - -inline C10_HOST_DEVICE BFloat16& operator/=(BFloat16& a, const BFloat16& b) { - a = a / b; - return a; -} - -inline C10_HOST_DEVICE BFloat16& operator|(BFloat16& a, const BFloat16& b) { - a.x = a.x | b.x; - return a; -} - -inline C10_HOST_DEVICE BFloat16& operator^(BFloat16& a, const BFloat16& b) { - a.x = a.x ^ b.x; - return a; -} - -inline C10_HOST_DEVICE BFloat16& operator&(BFloat16& a, const BFloat16& b) { - a.x = a.x & b.x; - return a; -} - -/// Arithmetic with floats - -inline C10_HOST_DEVICE float operator+(BFloat16 a, float b) { - return static_cast(a) + b; -} -inline C10_HOST_DEVICE float operator-(BFloat16 a, float b) { - return static_cast(a) - b; -} -inline C10_HOST_DEVICE float operator*(BFloat16 a, float b) { - return static_cast(a) * b; -} -inline C10_HOST_DEVICE float operator/(BFloat16 a, float b) { - return static_cast(a) / b; -} - -inline C10_HOST_DEVICE float operator+(float a, BFloat16 b) { - return a + static_cast(b); -} -inline C10_HOST_DEVICE float operator-(float a, BFloat16 b) { - return a - static_cast(b); -} -inline C10_HOST_DEVICE float operator*(float a, BFloat16 b) { - return a * static_cast(b); -} -inline C10_HOST_DEVICE float operator/(float a, BFloat16 b) { - return a / static_cast(b); -} - -inline C10_HOST_DEVICE float& operator+=(float& a, const BFloat16& b) { - return a += static_cast(b); -} -inline C10_HOST_DEVICE float& operator-=(float& a, const BFloat16& b) { - return a -= static_cast(b); -} -inline C10_HOST_DEVICE float& operator*=(float& a, const BFloat16& b) { - return a *= static_cast(b); -} -inline C10_HOST_DEVICE float& operator/=(float& a, const BFloat16& b) { - return a /= static_cast(b); -} - -/// Arithmetic with doubles - -inline C10_HOST_DEVICE double operator+(BFloat16 a, double b) { - return static_cast(a) + b; -} -inline C10_HOST_DEVICE double operator-(BFloat16 a, double b) { - return static_cast(a) - b; -} -inline C10_HOST_DEVICE double operator*(BFloat16 a, double b) { - return static_cast(a) * b; -} -inline C10_HOST_DEVICE double operator/(BFloat16 a, double b) { - return static_cast(a) / b; -} - -inline C10_HOST_DEVICE double operator+(double a, BFloat16 b) { - return a + static_cast(b); -} -inline C10_HOST_DEVICE double operator-(double a, BFloat16 b) { - return a - static_cast(b); -} -inline C10_HOST_DEVICE double operator*(double a, BFloat16 b) { - return a * static_cast(b); -} -inline C10_HOST_DEVICE double operator/(double a, BFloat16 b) { - return a / static_cast(b); -} - -/// Arithmetic with ints - -inline C10_HOST_DEVICE BFloat16 operator+(BFloat16 a, int b) { - return a + static_cast(b); -} -inline C10_HOST_DEVICE BFloat16 operator-(BFloat16 a, int b) { - return a - static_cast(b); -} -inline C10_HOST_DEVICE BFloat16 operator*(BFloat16 a, int b) { - return a * static_cast(b); -} -inline C10_HOST_DEVICE BFloat16 operator/(BFloat16 a, int b) { - return a / static_cast(b); -} - -inline C10_HOST_DEVICE BFloat16 operator+(int a, BFloat16 b) { - return static_cast(a) + b; -} -inline C10_HOST_DEVICE BFloat16 operator-(int a, BFloat16 b) { - return static_cast(a) - b; -} -inline C10_HOST_DEVICE BFloat16 operator*(int a, BFloat16 b) { - return static_cast(a) * b; -} -inline C10_HOST_DEVICE BFloat16 operator/(int a, BFloat16 b) { - return static_cast(a) / b; -} - -//// Arithmetic with int64_t - -inline C10_HOST_DEVICE BFloat16 operator+(BFloat16 a, int64_t b) { - return a + static_cast(b); -} -inline C10_HOST_DEVICE BFloat16 operator-(BFloat16 a, int64_t b) { - return a - static_cast(b); -} -inline C10_HOST_DEVICE BFloat16 operator*(BFloat16 a, int64_t b) { - return a * static_cast(b); -} -inline C10_HOST_DEVICE BFloat16 operator/(BFloat16 a, int64_t b) { - return a / static_cast(b); -} - -inline C10_HOST_DEVICE BFloat16 operator+(int64_t a, BFloat16 b) { - return static_cast(a) + b; -} -inline C10_HOST_DEVICE BFloat16 operator-(int64_t a, BFloat16 b) { - return static_cast(a) - b; -} -inline C10_HOST_DEVICE BFloat16 operator*(int64_t a, BFloat16 b) { - return static_cast(a) * b; -} -inline C10_HOST_DEVICE BFloat16 operator/(int64_t a, BFloat16 b) { - return static_cast(a) / b; -} - -// Overloading < and > operators, because std::max and std::min use them. - -inline C10_HOST_DEVICE bool operator>(BFloat16& lhs, BFloat16& rhs) { - return float(lhs) > float(rhs); -} - -inline C10_HOST_DEVICE bool operator<(BFloat16& lhs, BFloat16& rhs) { - return float(lhs) < float(rhs); -} - -} // namespace c10 - -namespace std { - -template <> -class numeric_limits { - public: - static constexpr bool is_signed = true; - static constexpr bool is_specialized = true; - static constexpr bool is_integer = false; - static constexpr bool is_exact = false; - static constexpr bool has_infinity = true; - static constexpr bool has_quiet_NaN = true; - static constexpr bool has_signaling_NaN = true; - static constexpr auto has_denorm = numeric_limits::has_denorm; - static constexpr auto has_denorm_loss = - numeric_limits::has_denorm_loss; - static constexpr auto round_style = numeric_limits::round_style; - static constexpr bool is_iec559 = false; - static constexpr bool is_bounded = true; - static constexpr bool is_modulo = false; - static constexpr int digits = 8; - static constexpr int digits10 = 2; - static constexpr int max_digits10 = 4; - static constexpr int radix = 2; - static constexpr int min_exponent = -125; - static constexpr int min_exponent10 = -37; - static constexpr int max_exponent = 128; - static constexpr int max_exponent10 = 38; - static constexpr auto traps = numeric_limits::traps; - static constexpr auto tinyness_before = - numeric_limits::tinyness_before; - - static constexpr c10::BFloat16 min() { - return c10::BFloat16(0x0080, c10::BFloat16::from_bits()); - } - static constexpr c10::BFloat16 lowest() { - return c10::BFloat16(0xFF7F, c10::BFloat16::from_bits()); - } - static constexpr c10::BFloat16 max() { - return c10::BFloat16(0x7F7F, c10::BFloat16::from_bits()); - } - static constexpr c10::BFloat16 epsilon() { - return c10::BFloat16(0x3C00, c10::BFloat16::from_bits()); - } - static constexpr c10::BFloat16 round_error() { - return c10::BFloat16(0x3F00, c10::BFloat16::from_bits()); - } - static constexpr c10::BFloat16 infinity() { - return c10::BFloat16(0x7F80, c10::BFloat16::from_bits()); - } - static constexpr c10::BFloat16 quiet_NaN() { - return c10::BFloat16(0x7FC0, c10::BFloat16::from_bits()); - } - static constexpr c10::BFloat16 signaling_NaN() { - return c10::BFloat16(0x7F80, c10::BFloat16::from_bits()); - } - static constexpr c10::BFloat16 denorm_min() { - return c10::BFloat16(0x0001, c10::BFloat16::from_bits()); - } -}; - -} // namespace std - -C10_CLANG_DIAGNOSTIC_POP() +#include diff --git a/runtime/core/portable_type/c10/c10/util/BFloat16.h b/runtime/core/portable_type/c10/c10/util/BFloat16.h index 06236df1fc8..6d3510cd5be 100644 --- a/runtime/core/portable_type/c10/c10/util/BFloat16.h +++ b/runtime/core/portable_type/c10/c10/util/BFloat16.h @@ -1,116 +1 @@ -#pragma once - -// Defines the bloat16 type (brain floating-point). This representation uses -// 1 bit for the sign, 8 bits for the exponent and 7 bits for the mantissa. - -#include -#include -#include -#include -#include -#include -#include - -#if defined(__CUDACC__) && !defined(USE_ROCM) -#include -#endif - -#if defined(CL_SYCL_LANGUAGE_VERSION) -#include // for SYCL 1.2.1 -#elif defined(SYCL_LANGUAGE_VERSION) -#include // for SYCL 2020 -#endif - -namespace c10 { - -namespace detail { -inline C10_HOST_DEVICE float f32_from_bits(uint16_t src) { - float res = 0; - uint32_t tmp = src; - tmp <<= 16; - -#if defined(USE_ROCM) && defined(__HIPCC__) - float* tempRes; - - // We should be using memcpy in order to respect the strict aliasing rule - // but it fails in the HIP environment. - tempRes = reinterpret_cast(&tmp); - res = *tempRes; -#else - std::memcpy(&res, &tmp, sizeof(tmp)); -#endif - - return res; -} - -inline C10_HOST_DEVICE uint16_t bits_from_f32(float src) { - uint32_t res = 0; - -#if defined(USE_ROCM) && defined(__HIPCC__) - // We should be using memcpy in order to respect the strict aliasing rule - // but it fails in the HIP environment. - uint32_t* tempRes = reinterpret_cast(&src); - res = *tempRes; -#else - std::memcpy(&res, &src, sizeof(res)); -#endif - - return res >> 16; -} - -inline C10_HOST_DEVICE uint16_t round_to_nearest_even(float src) { -#if defined(USE_ROCM) && defined(__HIPCC__) - if (src != src) { -#elif defined(_MSC_VER) - if (isnan(src)) { -#else - if (std::isnan(src)) { -#endif - return UINT16_C(0x7FC0); - } else { - const uint32_t U32 = c10::bit_cast(src); - uint32_t rounding_bias = ((U32 >> 16) & 1) + UINT32_C(0x7FFF); - return static_cast((U32 + rounding_bias) >> 16); - } -} -} // namespace detail - -struct alignas(2) BFloat16 { - uint16_t x; - - // HIP wants __host__ __device__ tag, CUDA does not -#if defined(USE_ROCM) && defined(__HIPCC__) - C10_HOST_DEVICE BFloat16() = default; -#else - BFloat16() = default; -#endif - - struct from_bits_t {}; - static constexpr C10_HOST_DEVICE from_bits_t from_bits() { - return from_bits_t(); - } - - constexpr C10_HOST_DEVICE BFloat16(unsigned short bits, from_bits_t) - : x(bits) {} - /* implicit */ inline C10_HOST_DEVICE BFloat16(float value); - inline C10_HOST_DEVICE operator float() const; - -#if defined(__CUDACC__) && !defined(USE_ROCM) - inline C10_HOST_DEVICE BFloat16(const __nv_bfloat16& value); - explicit inline C10_HOST_DEVICE operator __nv_bfloat16() const; -#endif - -#if defined(SYCL_EXT_ONEAPI_BFLOAT16_MATH_FUNCTIONS) - inline C10_HOST_DEVICE BFloat16(const sycl::ext::oneapi::bfloat16& value); - explicit inline C10_HOST_DEVICE operator sycl::ext::oneapi::bfloat16() const; -#endif -}; - -inline std::ostream& operator<<(std::ostream& out, const BFloat16& value) { - out << (float)value; - return out; -} - -} // namespace c10 - -#include // IWYU pragma: keep +#include diff --git a/runtime/core/portable_type/c10/c10/util/Half-inl.h b/runtime/core/portable_type/c10/c10/util/Half-inl.h index ae4469e5636..fe66779a0e5 100644 --- a/runtime/core/portable_type/c10/c10/util/Half-inl.h +++ b/runtime/core/portable_type/c10/c10/util/Half-inl.h @@ -1,350 +1 @@ -#pragma once - -#include -#include - -#include -#include - -#ifdef __CUDACC__ -#include -#endif - -#ifdef __HIPCC__ -#include -#endif - -#if defined(CL_SYCL_LANGUAGE_VERSION) -#include // for SYCL 1.2.1 -#elif defined(SYCL_LANGUAGE_VERSION) -#include // for SYCL 2020 -#endif - -#if (defined(CPU_CAPABILITY_AVX2) || defined(CPU_CAPABILITY_AVX512)) && \ - !defined(__APPLE__) -#include -#endif - -C10_CLANG_DIAGNOSTIC_PUSH() -#if C10_CLANG_HAS_WARNING("-Wimplicit-int-float-conversion") -C10_CLANG_DIAGNOSTIC_IGNORE("-Wimplicit-int-float-conversion") -#endif - -namespace c10 { - -#if defined(__aarch64__) && !defined(__CUDACC__) -/// Constructors -inline Half::Half(float16_t value) : x(detail::fp16_to_bits(value)) {} -inline Half::operator float16_t() const { - return detail::fp16_from_bits(x); -} -#else - -inline C10_HOST_DEVICE Half::Half(float value) - : -#if defined(__CUDA_ARCH__) || defined(__HIP_DEVICE_COMPILE__) - x(__half_as_short(__float2half(value))) -#elif defined(__SYCL_DEVICE_ONLY__) - x(c10::bit_cast(sycl::half(value))) -#elif (defined(CPU_CAPABILITY_AVX2) || defined(CPU_CAPABILITY_AVX512)) && \ - !defined(__APPLE__) - x(at::vec::float2half_scalar(value)) -#else - x(detail::fp16_ieee_from_fp32_value(value)) -#endif -{ -} - -/// Implicit conversions - -inline C10_HOST_DEVICE Half::operator float() const { -#if defined(__CUDA_ARCH__) || defined(__HIP_DEVICE_COMPILE__) - return __half2float(*reinterpret_cast(&x)); -#elif defined(__SYCL_DEVICE_ONLY__) - return float(c10::bit_cast(x)); -#elif (defined(CPU_CAPABILITY_AVX2) || defined(CPU_CAPABILITY_AVX512)) && \ - !defined(__APPLE__) - return at::vec::half2float_scalar(x); -#elif defined(__aarch64__) && !defined(__CUDACC__) - return detail::native_fp16_to_fp32_value(x); -#else - return detail::fp16_ieee_to_fp32_value(x); -#endif -} - -#endif /* !defined(__aarch64__) || defined(__CUDACC__) \ - */ - -#if defined(__CUDACC__) || defined(__HIPCC__) -inline C10_HOST_DEVICE Half::Half(const __half& value) { - x = *reinterpret_cast(&value); -} -inline C10_HOST_DEVICE Half::operator __half() const { - return *reinterpret_cast(&x); -} -#endif - -#ifdef SYCL_LANGUAGE_VERSION -inline C10_HOST_DEVICE Half::Half(const sycl::half& value) { - x = *reinterpret_cast(&value); -} -inline C10_HOST_DEVICE Half::operator sycl::half() const { - return *reinterpret_cast(&x); -} -#endif - -// CUDA intrinsics - -#if (defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 350)) || \ - (defined(__clang__) && defined(__CUDA__)) -inline __device__ Half __ldg(const Half* ptr) { - return __ldg(reinterpret_cast(ptr)); -} -#endif - -/// Arithmetic - -inline C10_HOST_DEVICE Half operator+(const Half& a, const Half& b) { - return static_cast(a) + static_cast(b); -} - -inline C10_HOST_DEVICE Half operator-(const Half& a, const Half& b) { - return static_cast(a) - static_cast(b); -} - -inline C10_HOST_DEVICE Half operator*(const Half& a, const Half& b) { - return static_cast(a) * static_cast(b); -} - -inline C10_HOST_DEVICE Half operator/(const Half& a, const Half& b) - __ubsan_ignore_float_divide_by_zero__ { - return static_cast(a) / static_cast(b); -} - -inline C10_HOST_DEVICE Half operator-(const Half& a) { -#if (defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 530) || \ - defined(__HIP_DEVICE_COMPILE__) - return __hneg(a); -#elif defined(__SYCL_DEVICE_ONLY__) - return -c10::bit_cast(a); -#else - return -static_cast(a); -#endif -} - -inline C10_HOST_DEVICE Half& operator+=(Half& a, const Half& b) { - a = a + b; - return a; -} - -inline C10_HOST_DEVICE Half& operator-=(Half& a, const Half& b) { - a = a - b; - return a; -} - -inline C10_HOST_DEVICE Half& operator*=(Half& a, const Half& b) { - a = a * b; - return a; -} - -inline C10_HOST_DEVICE Half& operator/=(Half& a, const Half& b) { - a = a / b; - return a; -} - -/// Arithmetic with floats - -inline C10_HOST_DEVICE float operator+(Half a, float b) { - return static_cast(a) + b; -} -inline C10_HOST_DEVICE float operator-(Half a, float b) { - return static_cast(a) - b; -} -inline C10_HOST_DEVICE float operator*(Half a, float b) { - return static_cast(a) * b; -} -inline C10_HOST_DEVICE float operator/(Half a, float b) - __ubsan_ignore_float_divide_by_zero__ { - return static_cast(a) / b; -} - -inline C10_HOST_DEVICE float operator+(float a, Half b) { - return a + static_cast(b); -} -inline C10_HOST_DEVICE float operator-(float a, Half b) { - return a - static_cast(b); -} -inline C10_HOST_DEVICE float operator*(float a, Half b) { - return a * static_cast(b); -} -inline C10_HOST_DEVICE float operator/(float a, Half b) - __ubsan_ignore_float_divide_by_zero__ { - return a / static_cast(b); -} - -inline C10_HOST_DEVICE float& operator+=(float& a, const Half& b) { - return a += static_cast(b); -} -inline C10_HOST_DEVICE float& operator-=(float& a, const Half& b) { - return a -= static_cast(b); -} -inline C10_HOST_DEVICE float& operator*=(float& a, const Half& b) { - return a *= static_cast(b); -} -inline C10_HOST_DEVICE float& operator/=(float& a, const Half& b) { - return a /= static_cast(b); -} - -/// Arithmetic with doubles - -inline C10_HOST_DEVICE double operator+(Half a, double b) { - return static_cast(a) + b; -} -inline C10_HOST_DEVICE double operator-(Half a, double b) { - return static_cast(a) - b; -} -inline C10_HOST_DEVICE double operator*(Half a, double b) { - return static_cast(a) * b; -} -inline C10_HOST_DEVICE double operator/(Half a, double b) - __ubsan_ignore_float_divide_by_zero__ { - return static_cast(a) / b; -} - -inline C10_HOST_DEVICE double operator+(double a, Half b) { - return a + static_cast(b); -} -inline C10_HOST_DEVICE double operator-(double a, Half b) { - return a - static_cast(b); -} -inline C10_HOST_DEVICE double operator*(double a, Half b) { - return a * static_cast(b); -} -inline C10_HOST_DEVICE double operator/(double a, Half b) - __ubsan_ignore_float_divide_by_zero__ { - return a / static_cast(b); -} - -/// Arithmetic with ints - -inline C10_HOST_DEVICE Half operator+(Half a, int b) { - return a + static_cast(b); -} -inline C10_HOST_DEVICE Half operator-(Half a, int b) { - return a - static_cast(b); -} -inline C10_HOST_DEVICE Half operator*(Half a, int b) { - return a * static_cast(b); -} -inline C10_HOST_DEVICE Half operator/(Half a, int b) { - return a / static_cast(b); -} - -inline C10_HOST_DEVICE Half operator+(int a, Half b) { - return static_cast(a) + b; -} -inline C10_HOST_DEVICE Half operator-(int a, Half b) { - return static_cast(a) - b; -} -inline C10_HOST_DEVICE Half operator*(int a, Half b) { - return static_cast(a) * b; -} -inline C10_HOST_DEVICE Half operator/(int a, Half b) { - return static_cast(a) / b; -} - -//// Arithmetic with int64_t - -inline C10_HOST_DEVICE Half operator+(Half a, int64_t b) { - return a + static_cast(b); -} -inline C10_HOST_DEVICE Half operator-(Half a, int64_t b) { - return a - static_cast(b); -} -inline C10_HOST_DEVICE Half operator*(Half a, int64_t b) { - return a * static_cast(b); -} -inline C10_HOST_DEVICE Half operator/(Half a, int64_t b) { - return a / static_cast(b); -} - -inline C10_HOST_DEVICE Half operator+(int64_t a, Half b) { - return static_cast(a) + b; -} -inline C10_HOST_DEVICE Half operator-(int64_t a, Half b) { - return static_cast(a) - b; -} -inline C10_HOST_DEVICE Half operator*(int64_t a, Half b) { - return static_cast(a) * b; -} -inline C10_HOST_DEVICE Half operator/(int64_t a, Half b) { - return static_cast(a) / b; -} - -/// NOTE: we do not define comparisons directly and instead rely on the implicit -/// conversion from c10::Half to float. - -} // namespace c10 - -namespace std { - -template <> -class numeric_limits { - public: - static constexpr bool is_specialized = true; - static constexpr bool is_signed = true; - static constexpr bool is_integer = false; - static constexpr bool is_exact = false; - static constexpr bool has_infinity = true; - static constexpr bool has_quiet_NaN = true; - static constexpr bool has_signaling_NaN = true; - static constexpr auto has_denorm = numeric_limits::has_denorm; - static constexpr auto has_denorm_loss = - numeric_limits::has_denorm_loss; - static constexpr auto round_style = numeric_limits::round_style; - static constexpr bool is_iec559 = true; - static constexpr bool is_bounded = true; - static constexpr bool is_modulo = false; - static constexpr int digits = 11; - static constexpr int digits10 = 3; - static constexpr int max_digits10 = 5; - static constexpr int radix = 2; - static constexpr int min_exponent = -13; - static constexpr int min_exponent10 = -4; - static constexpr int max_exponent = 16; - static constexpr int max_exponent10 = 4; - static constexpr auto traps = numeric_limits::traps; - static constexpr auto tinyness_before = - numeric_limits::tinyness_before; - static constexpr c10::Half min() { - return c10::Half(0x0400, c10::Half::from_bits()); - } - static constexpr c10::Half lowest() { - return c10::Half(0xFBFF, c10::Half::from_bits()); - } - static constexpr c10::Half max() { - return c10::Half(0x7BFF, c10::Half::from_bits()); - } - static constexpr c10::Half epsilon() { - return c10::Half(0x1400, c10::Half::from_bits()); - } - static constexpr c10::Half round_error() { - return c10::Half(0x3800, c10::Half::from_bits()); - } - static constexpr c10::Half infinity() { - return c10::Half(0x7C00, c10::Half::from_bits()); - } - static constexpr c10::Half quiet_NaN() { - return c10::Half(0x7E00, c10::Half::from_bits()); - } - static constexpr c10::Half signaling_NaN() { - return c10::Half(0x7D00, c10::Half::from_bits()); - } - static constexpr c10::Half denorm_min() { - return c10::Half(0x0001, c10::Half::from_bits()); - } -}; - -} // namespace std - -C10_CLANG_DIAGNOSTIC_POP() +#include diff --git a/runtime/core/portable_type/c10/c10/util/Half.h b/runtime/core/portable_type/c10/c10/util/Half.h index bdcf7458145..98480b22db3 100644 --- a/runtime/core/portable_type/c10/c10/util/Half.h +++ b/runtime/core/portable_type/c10/c10/util/Half.h @@ -1,424 +1,8 @@ -#pragma once +#include -/// Defines the Half type (half-precision floating-point) including conversions -/// to standard C types and basic arithmetic operations. Note that arithmetic -/// operations are implemented by converting to floating point and -/// performing the operation in float32, instead of using CUDA half intrinsics. -/// Most uses of this type within ATen are memory bound, including the -/// element-wise kernels, and the half intrinsics aren't efficient on all GPUs. -/// If you are writing a compute bound kernel, you can use the CUDA half -/// intrinsics directly on the Half type from device code. - -#include -#include -#include -#include -#include - -#if defined(__cplusplus) -#include -#elif !defined(__OPENCL_VERSION__) -#include -#endif - -#ifdef _MSC_VER -#include -#endif - -#include -#include -#include -#include -#include - -#ifdef __CUDACC__ -#include -#endif - -#ifdef __HIPCC__ -#include -#endif - -#if defined(CL_SYCL_LANGUAGE_VERSION) -#include // for SYCL 1.2.1 -#elif defined(SYCL_LANGUAGE_VERSION) -#include // for SYCL 2020 -#endif - -#if defined(__aarch64__) && !defined(__CUDACC__) -#include -#endif - -#if defined(__GNUC__) || defined(__clang__) -#if defined(__x86_64__) || defined(_M_X64) || defined(__i386) || \ - defined(_M_IX86) -#if defined(__F16C__) && \ - !(defined(__CUDA_ARCH__) || defined(__CUDACC__) || \ - defined(__HIP_DEVICE_COMPILE__)) -#define C10_X86_F16 1 -#include // import conversion ops from f16cintrin.h -#endif // defined(__F16C__) && !(defined(__CUDA_ARCH__) || defined(__CUDACC__) - // || defined(__HIP_DEVICE_COMPILE__)) -#endif // __x86_64__ || _M_X64 || __i386 || _M_IX86 -#endif // __GNUC__ || __clang__ - -namespace c10 { - -namespace detail { - -/* - * Convert a 16-bit floating-point number in IEEE half-precision format, in bit - * representation, to a 32-bit floating-point number in IEEE single-precision - * format, in bit representation. - * - * @note The implementation doesn't use any floating-point operations. - */ -inline uint32_t fp16_ieee_to_fp32_bits(uint16_t h) { - /* - * Extend the half-precision floating-point number to 32 bits and shift to the - * upper part of the 32-bit word: - * +---+-----+------------+-------------------+ - * | S |EEEEE|MM MMMM MMMM|0000 0000 0000 0000| - * +---+-----+------------+-------------------+ - * Bits 31 26-30 16-25 0-15 - * - * S - sign bit, E - bits of the biased exponent, M - bits of the mantissa, 0 - * - zero bits. - */ - const uint32_t w = (uint32_t)h << 16; - /* - * Extract the sign of the input number into the high bit of the 32-bit word: - * - * +---+----------------------------------+ - * | S |0000000 00000000 00000000 00000000| - * +---+----------------------------------+ - * Bits 31 0-31 - */ - const uint32_t sign = w & UINT32_C(0x80000000); - /* - * Extract mantissa and biased exponent of the input number into the bits 0-30 - * of the 32-bit word: - * - * +---+-----+------------+-------------------+ - * | 0 |EEEEE|MM MMMM MMMM|0000 0000 0000 0000| - * +---+-----+------------+-------------------+ - * Bits 30 27-31 17-26 0-16 - */ - const uint32_t nonsign = w & UINT32_C(0x7FFFFFFF); - /* - * Renorm shift is the number of bits to shift mantissa left to make the - * half-precision number normalized. If the initial number is normalized, some - * of its high 6 bits (sign == 0 and 5-bit exponent) equals one. In this case - * renorm_shift == 0. If the number is denormalize, renorm_shift > 0. Note - * that if we shift denormalized nonsign by renorm_shift, the unit bit of - * mantissa will shift into exponent, turning the biased exponent into 1, and - * making mantissa normalized (i.e. without leading 1). - */ -#ifdef _MSC_VER - unsigned long nonsign_bsr; - _BitScanReverse(&nonsign_bsr, (unsigned long)nonsign); - uint32_t renorm_shift = (uint32_t)nonsign_bsr ^ 31; -#else - uint32_t renorm_shift = __builtin_clz(nonsign); -#endif - renorm_shift = renorm_shift > 5 ? renorm_shift - 5 : 0; - /* - * Iff half-precision number has exponent of 15, the addition overflows - * it into bit 31, and the subsequent shift turns the high 9 bits - * into 1. Thus inf_nan_mask == 0x7F800000 if the half-precision number - * had exponent of 15 (i.e. was NaN or infinity) 0x00000000 otherwise - */ - const int32_t inf_nan_mask = - ((int32_t)(nonsign + 0x04000000) >> 8) & INT32_C(0x7F800000); - /* - * Iff nonsign is 0, it overflows into 0xFFFFFFFF, turning bit 31 - * into 1. Otherwise, bit 31 remains 0. The signed shift right by 31 - * broadcasts bit 31 into all bits of the zero_mask. Thus zero_mask == - * 0xFFFFFFFF if the half-precision number was zero (+0.0h or -0.0h) - * 0x00000000 otherwise - */ - const int32_t zero_mask = (int32_t)(nonsign - 1) >> 31; - /* - * 1. Shift nonsign left by renorm_shift to normalize it (if the input - * was denormal) - * 2. Shift nonsign right by 3 so the exponent (5 bits originally) - * becomes an 8-bit field and 10-bit mantissa shifts into the 10 high - * bits of the 23-bit mantissa of IEEE single-precision number. - * 3. Add 0x70 to the exponent (starting at bit 23) to compensate the - * different in exponent bias (0x7F for single-precision number less 0xF - * for half-precision number). - * 4. Subtract renorm_shift from the exponent (starting at bit 23) to - * account for renormalization. As renorm_shift is less than 0x70, this - * can be combined with step 3. - * 5. Binary OR with inf_nan_mask to turn the exponent into 0xFF if the - * input was NaN or infinity. - * 6. Binary ANDNOT with zero_mask to turn the mantissa and exponent - * into zero if the input was zero. - * 7. Combine with the sign of the input number. - */ - return sign | - ((((nonsign << renorm_shift >> 3) + ((0x70 - renorm_shift) << 23)) | - inf_nan_mask) & - ~zero_mask); -} - -/* - * Convert a 16-bit floating-point number in IEEE half-precision format, in bit - * representation, to a 32-bit floating-point number in IEEE single-precision - * format. - * - * @note The implementation relies on IEEE-like (no assumption about rounding - * mode and no operations on denormals) floating-point operations and bitcasts - * between integer and floating-point variables. - */ -C10_HOST_DEVICE inline float fp16_ieee_to_fp32_value(uint16_t h) { -#ifdef C10_X86_F16 - return _cvtsh_ss(h); -#else - /* - * Extend the half-precision floating-point number to 32 bits and shift to the - * upper part of the 32-bit word: - * +---+-----+------------+-------------------+ - * | S |EEEEE|MM MMMM MMMM|0000 0000 0000 0000| - * +---+-----+------------+-------------------+ - * Bits 31 26-30 16-25 0-15 - * - * S - sign bit, E - bits of the biased exponent, M - bits of the mantissa, 0 - * - zero bits. - */ - const uint32_t w = (uint32_t)h << 16; - /* - * Extract the sign of the input number into the high bit of the 32-bit word: - * - * +---+----------------------------------+ - * | S |0000000 00000000 00000000 00000000| - * +---+----------------------------------+ - * Bits 31 0-31 - */ - const uint32_t sign = w & UINT32_C(0x80000000); - /* - * Extract mantissa and biased exponent of the input number into the high bits - * of the 32-bit word: - * - * +-----+------------+---------------------+ - * |EEEEE|MM MMMM MMMM|0 0000 0000 0000 0000| - * +-----+------------+---------------------+ - * Bits 27-31 17-26 0-16 - */ - const uint32_t two_w = w + w; - - /* - * Shift mantissa and exponent into bits 23-28 and bits 13-22 so they become - * mantissa and exponent of a single-precision floating-point number: - * - * S|Exponent | Mantissa - * +-+---+-----+------------+----------------+ - * |0|000|EEEEE|MM MMMM MMMM|0 0000 0000 0000| - * +-+---+-----+------------+----------------+ - * Bits | 23-31 | 0-22 - * - * Next, there are some adjustments to the exponent: - * - The exponent needs to be corrected by the difference in exponent bias - * between single-precision and half-precision formats (0x7F - 0xF = 0x70) - * - Inf and NaN values in the inputs should become Inf and NaN values after - * conversion to the single-precision number. Therefore, if the biased - * exponent of the half-precision input was 0x1F (max possible value), the - * biased exponent of the single-precision output must be 0xFF (max possible - * value). We do this correction in two steps: - * - First, we adjust the exponent by (0xFF - 0x1F) = 0xE0 (see exp_offset - * below) rather than by 0x70 suggested by the difference in the exponent bias - * (see above). - * - Then we multiply the single-precision result of exponent adjustment by - * 2**(-112) to reverse the effect of exponent adjustment by 0xE0 less the - * necessary exponent adjustment by 0x70 due to difference in exponent bias. - * The floating-point multiplication hardware would ensure than Inf and - * NaN would retain their value on at least partially IEEE754-compliant - * implementations. - * - * Note that the above operations do not handle denormal inputs (where biased - * exponent == 0). However, they also do not operate on denormal inputs, and - * do not produce denormal results. - */ - constexpr uint32_t exp_offset = UINT32_C(0xE0) << 23; - // const float exp_scale = 0x1.0p-112f; - constexpr uint32_t scale_bits = (uint32_t)15 << 23; - float exp_scale_val = 0; -#if defined(_MSC_VER) && defined(__clang__) - __builtin_memcpy(&exp_scale_val, &scale_bits, sizeof(exp_scale_val)); -#else - std::memcpy(&exp_scale_val, &scale_bits, sizeof(exp_scale_val)); -#endif - - const float exp_scale = exp_scale_val; - const float normalized_value = - fp32_from_bits((two_w >> 4) + exp_offset) * exp_scale; - - /* - * Convert denormalized half-precision inputs into single-precision results - * (always normalized). Zero inputs are also handled here. - * - * In a denormalized number the biased exponent is zero, and mantissa has - * on-zero bits. First, we shift mantissa into bits 0-9 of the 32-bit word. - * - * zeros | mantissa - * +---------------------------+------------+ - * |0000 0000 0000 0000 0000 00|MM MMMM MMMM| - * +---------------------------+------------+ - * Bits 10-31 0-9 - * - * Now, remember that denormalized half-precision numbers are represented as: - * FP16 = mantissa * 2**(-24). - * The trick is to construct a normalized single-precision number with the - * same mantissa and thehalf-precision input and with an exponent which would - * scale the corresponding mantissa bits to 2**(-24). A normalized - * single-precision floating-point number is represented as: FP32 = (1 + - * mantissa * 2**(-23)) * 2**(exponent - 127) Therefore, when the biased - * exponent is 126, a unit change in the mantissa of the input denormalized - * half-precision number causes a change of the constructed single-precision - * number by 2**(-24), i.e. the same amount. - * - * The last step is to adjust the bias of the constructed single-precision - * number. When the input half-precision number is zero, the constructed - * single-precision number has the value of FP32 = 1 * 2**(126 - 127) = - * 2**(-1) = 0.5 Therefore, we need to subtract 0.5 from the constructed - * single-precision number to get the numerical equivalent of the input - * half-precision number. - */ - constexpr uint32_t magic_mask = UINT32_C(126) << 23; - constexpr float magic_bias = 0.5f; - const float denormalized_value = - fp32_from_bits((two_w >> 17) | magic_mask) - magic_bias; - - /* - * - Choose either results of conversion of input as a normalized number, or - * as a denormalized number, depending on the input exponent. The variable - * two_w contains input exponent in bits 27-31, therefore if its smaller than - * 2**27, the input is either a denormal number, or zero. - * - Combine the result of conversion of exponent and mantissa with the sign - * of the input number. - */ - constexpr uint32_t denormalized_cutoff = UINT32_C(1) << 27; - const uint32_t result = sign | - (two_w < denormalized_cutoff ? fp32_to_bits(denormalized_value) - : fp32_to_bits(normalized_value)); - return fp32_from_bits(result); -#endif // C10_X86_F16 -} - -/* - * Convert a 32-bit floating-point number in IEEE single-precision format to a - * 16-bit floating-point number in IEEE half-precision format, in bit - * representation. - * - * @note The implementation relies on IEEE-like (no assumption about rounding - * mode and no operations on denormals) floating-point operations and bitcasts - * between integer and floating-point variables. - */ -inline uint16_t fp16_ieee_from_fp32_value(float f) { -#ifdef C10_X86_F16 - return _cvtss_sh(f, _MM_FROUND_TO_NEAREST_INT); -#else - // const float scale_to_inf = 0x1.0p+112f; - // const float scale_to_zero = 0x1.0p-110f; - constexpr uint32_t scale_to_inf_bits = (uint32_t)239 << 23; - constexpr uint32_t scale_to_zero_bits = (uint32_t)17 << 23; - float scale_to_inf_val = 0, scale_to_zero_val = 0; - std::memcpy(&scale_to_inf_val, &scale_to_inf_bits, sizeof(scale_to_inf_val)); - std::memcpy( - &scale_to_zero_val, &scale_to_zero_bits, sizeof(scale_to_zero_val)); - const float scale_to_inf = scale_to_inf_val; - const float scale_to_zero = scale_to_zero_val; - -#if defined(_MSC_VER) && _MSC_VER == 1916 - float base = ((signbit(f) != 0 ? -f : f) * scale_to_inf) * scale_to_zero; -#else - float base = (fabsf(f) * scale_to_inf) * scale_to_zero; -#endif - - const uint32_t w = fp32_to_bits(f); - const uint32_t shl1_w = w + w; - const uint32_t sign = w & UINT32_C(0x80000000); - uint32_t bias = shl1_w & UINT32_C(0xFF000000); - if (bias < UINT32_C(0x71000000)) { - bias = UINT32_C(0x71000000); - } - - base = fp32_from_bits((bias >> 1) + UINT32_C(0x07800000)) + base; - const uint32_t bits = fp32_to_bits(base); - const uint32_t exp_bits = (bits >> 13) & UINT32_C(0x00007C00); - const uint32_t mantissa_bits = bits & UINT32_C(0x00000FFF); - const uint32_t nonsign = exp_bits + mantissa_bits; - return static_cast( - (sign >> 16) | - (shl1_w > UINT32_C(0xFF000000) ? UINT16_C(0x7E00) : nonsign)); -#endif // C10_X86_F16 -} - -#ifdef C10_X86_F16 -#undef C10_X86_F16 -#endif // C10_X86_F16 - -#if defined(__aarch64__) && !defined(__CUDACC__) -inline float16_t fp16_from_bits(uint16_t h) { - return c10::bit_cast(h); -} - -inline uint16_t fp16_to_bits(float16_t f) { - return c10::bit_cast(f); -} - -// According to https://godbolt.org/z/frExdbsWG it would translate to single -// fcvt s0, h0 -inline float native_fp16_to_fp32_value(uint16_t h) { - return static_cast(fp16_from_bits(h)); -} - -inline uint16_t native_fp16_from_fp32_value(float f) { - return fp16_to_bits(static_cast(f)); -} -#endif - -} // namespace detail - -struct alignas(2) Half { - unsigned short x; - - struct from_bits_t {}; - C10_HOST_DEVICE static constexpr from_bits_t from_bits() { - return from_bits_t(); - } - - // HIP wants __host__ __device__ tag, CUDA does not -#if defined(USE_ROCM) - C10_HOST_DEVICE Half() = default; -#else - Half() = default; -#endif - - constexpr C10_HOST_DEVICE Half(unsigned short bits, from_bits_t) : x(bits) {} -#if defined(__aarch64__) && !defined(__CUDACC__) - inline Half(float16_t value); - inline operator float16_t() const; -#else - inline C10_HOST_DEVICE Half(float value); - inline C10_HOST_DEVICE operator float() const; +// need to keep the following for BC because the APIs in here were exposed +// before migrating Half to torch/headeronly +#if (defined(CPU_CAPABILITY_AVX2) || defined(CPU_CAPABILITY_AVX512)) && \ + !defined(__APPLE__) +#include #endif - -#if defined(__CUDACC__) || defined(__HIPCC__) - inline C10_HOST_DEVICE Half(const __half& value); - inline C10_HOST_DEVICE operator __half() const; -#endif -#ifdef SYCL_LANGUAGE_VERSION - inline C10_HOST_DEVICE Half(const sycl::half& value); - inline C10_HOST_DEVICE operator sycl::half() const; -#endif -}; - -inline std::ostream& operator<<(std::ostream& out, const Half& value) { - out << (float)value; - return out; -} - -} // namespace c10 - -#include // IWYU pragma: keep diff --git a/runtime/core/portable_type/c10/c10/util/TypeSafeSignMath.h b/runtime/core/portable_type/c10/c10/util/TypeSafeSignMath.h index 58c05067830..28520225d4b 100644 --- a/runtime/core/portable_type/c10/c10/util/TypeSafeSignMath.h +++ b/runtime/core/portable_type/c10/c10/util/TypeSafeSignMath.h @@ -1,140 +1 @@ -#pragma once - -#include -#include -#include - -C10_CLANG_DIAGNOSTIC_PUSH() -#if C10_CLANG_HAS_WARNING("-Wstring-conversion") -C10_CLANG_DIAGNOSTIC_IGNORE("-Wstring-conversion") -#endif -#if C10_CLANG_HAS_WARNING("-Wimplicit-int-float-conversion") -C10_CLANG_DIAGNOSTIC_IGNORE("-Wimplicit-int-float-conversion") -#endif - -namespace c10 { - -/// Returns false since we cannot have x < 0 if x is unsigned. -template -inline constexpr bool is_negative( - const T& /*x*/, - std::true_type /*is_unsigned*/) { - return false; -} - -/// Returns true if a signed variable x < 0 -template -inline constexpr bool is_negative(const T& x, std::false_type /*is_unsigned*/) { - return x < T(0); -} - -/// Returns true if x < 0 -/// NOTE: Will fail on an unsigned custom type -/// For the most part it's possible to fix this if -/// the custom type has a constexpr constructor. -/// However, notably, c10::Half does not :-( -template -inline constexpr bool is_negative(const T& x) { - return is_negative(x, std::is_unsigned()); -} - -/// Returns the sign of an unsigned variable x as 0, 1 -template -inline constexpr int signum(const T& x, std::true_type /*is_unsigned*/) { - return T(0) < x; -} - -/// Returns the sign of a signed variable x as -1, 0, 1 -template -inline constexpr int signum(const T& x, std::false_type /*is_unsigned*/) { - return (T(0) < x) - (x < T(0)); -} - -/// Returns the sign of x as -1, 0, 1 -/// NOTE: Will fail on an unsigned custom type -/// For the most part it's possible to fix this if -/// the custom type has a constexpr constructor. -/// However, notably, c10::Half does not :-( -template -inline constexpr int signum(const T& x) { - return signum(x, std::is_unsigned()); -} - -/// Returns true if a and b are not both negative -template -inline constexpr bool signs_differ(const T& a, const U& b) { - return is_negative(a) != is_negative(b); -} - -// Suppress sign compare warning when compiling with GCC -// as later does not account for short-circuit rule before -// raising the warning, see https://godbolt.org/z/Tr3Msnz99 -#ifdef __GNUC__ -#pragma GCC diagnostic push -#pragma GCC diagnostic ignored "-Wsign-compare" -#endif - -/// Returns true if x is greater than the greatest value of the type Limit -template -inline constexpr bool greater_than_max(const T& x) { - constexpr bool can_overflow = - std::numeric_limits::digits > std::numeric_limits::digits; - return can_overflow && x > (std::numeric_limits::max)(); -} - -#ifdef __GNUC__ -#pragma GCC diagnostic pop -#endif - -/// Returns true if x < lowest(Limit). Standard comparison -template -inline constexpr bool less_than_lowest( - const T& x, - std::false_type /*limit_is_unsigned*/, - std::false_type /*x_is_unsigned*/) { - return x < std::numeric_limits::lowest(); -} - -/// Returns false since all the limit is signed and therefore includes -/// negative values but x cannot be negative because it is unsigned -template -inline constexpr bool less_than_lowest( - const T& /*x*/, - std::false_type /*limit_is_unsigned*/, - std::true_type /*x_is_unsigned*/) { - return false; -} - -/// Returns true if x < 0, where 0 is constructed from T. -/// Limit is not signed, so its lower value is zero -template -inline constexpr bool less_than_lowest( - const T& x, - std::true_type /*limit_is_unsigned*/, - std::false_type /*x_is_unsigned*/) { - return x < T(0); -} - -/// Returns false sign both types are unsigned -template -inline constexpr bool less_than_lowest( - const T& /*x*/, - std::true_type /*limit_is_unsigned*/, - std::true_type /*x_is_unsigned*/) { - return false; -} - -/// Returns true if x is less than the lowest value of type T -/// NOTE: Will fail on an unsigned custom type -/// For the most part it's possible to fix this if -/// the custom type has a constexpr constructor. -/// However, notably, c10::Half does not : -template -inline constexpr bool less_than_lowest(const T& x) { - return less_than_lowest( - x, std::is_unsigned(), std::is_unsigned()); -} - -} // namespace c10 - -C10_CLANG_DIAGNOSTIC_POP() +#include diff --git a/runtime/core/portable_type/c10/c10/util/bit_cast.h b/runtime/core/portable_type/c10/c10/util/bit_cast.h index d7d2aa8dd39..49d0822d94f 100644 --- a/runtime/core/portable_type/c10/c10/util/bit_cast.h +++ b/runtime/core/portable_type/c10/c10/util/bit_cast.h @@ -1,46 +1 @@ -#pragma once - -#include -#include - -#include - -#if __has_include() && (defined(__cpp_lib_bit_cast) && __cpp_lib_bit_cast >= 201806L) -#include -#define C10_HAVE_STD_BIT_CAST 1 -#else -#define C10_HAVE_STD_BIT_CAST 0 -#endif // __has_include() && (__cplusplus >= 202002L || - // (defined(__cpp_lib_bit_cast) && __cpp_lib_bit_cast >= 201806L)) - -namespace c10 { - -#if C10_HAVE_STD_BIT_CAST -using std::bit_cast; -#else -// Implementations of std::bit_cast() from C++ 20. -// -// This is a less sketchy version of reinterpret_cast. -// -// See https://en.cppreference.com/w/cpp/numeric/bit_cast for more -// information as well as the source of our implementations. -template -C10_HOST_DEVICE std::enable_if_t< - sizeof(To) == sizeof(From) && std::is_trivially_copyable_v && - std::is_trivially_copyable_v, - To> -// constexpr support needs compiler magic -bit_cast(const From& src) noexcept { - static_assert( - std::is_trivially_constructible_v, - "This implementation additionally requires " - "destination type to be trivially constructible"); - - To dst; - std::memcpy(&dst, &src, sizeof(To)); - return dst; -} -#endif // C10_HAVE_STD_BIT_CAST -#undef C10_HAVE_STD_BIT_CAST - -} // namespace c10 +#include diff --git a/runtime/core/portable_type/c10/c10/util/complex.h b/runtime/core/portable_type/c10/c10/util/complex.h index b63710d9458..4e699684bc3 100644 --- a/runtime/core/portable_type/c10/c10/util/complex.h +++ b/runtime/core/portable_type/c10/c10/util/complex.h @@ -4,531 +4,7 @@ #include #include - -#if defined(__CUDACC__) || defined(__HIPCC__) -#include -#endif - -C10_CLANG_DIAGNOSTIC_PUSH() -#if C10_CLANG_HAS_WARNING("-Wimplicit-float-conversion") -C10_CLANG_DIAGNOSTIC_IGNORE("-Wimplicit-float-conversion") -#endif -#if C10_CLANG_HAS_WARNING("-Wfloat-conversion") -C10_CLANG_DIAGNOSTIC_IGNORE("-Wfloat-conversion") -#endif - -namespace c10 { - -// c10::complex is an implementation of complex numbers that aims -// to work on all devices supported by PyTorch -// -// Most of the APIs duplicates std::complex -// Reference: https://en.cppreference.com/w/cpp/numeric/complex -// -// [NOTE: Complex Operator Unification] -// Operators currently use a mix of std::complex, thrust::complex, and -// c10::complex internally. The end state is that all operators will use -// c10::complex internally. Until then, there may be some hacks to support all -// variants. -// -// -// [Note on Constructors] -// -// The APIs of constructors are mostly copied from C++ standard: -// https://en.cppreference.com/w/cpp/numeric/complex/complex -// -// Since C++14, all constructors are constexpr in std::complex -// -// There are three types of constructors: -// - initializing from real and imag: -// `constexpr complex( const T& re = T(), const T& im = T() );` -// - implicitly-declared copy constructor -// - converting constructors -// -// Converting constructors: -// - std::complex defines converting constructor between float/double/long -// double, -// while we define converting constructor between float/double. -// - For these converting constructors, upcasting is implicit, downcasting is -// explicit. -// - We also define explicit casting from std::complex/thrust::complex -// - Note that the conversion from thrust is not constexpr, because -// thrust does not define them as constexpr ???? -// -// -// [Operator =] -// -// The APIs of operator = are mostly copied from C++ standard: -// https://en.cppreference.com/w/cpp/numeric/complex/operator%3D -// -// Since C++20, all operator= are constexpr. Although we are not building with -// C++20, we also obey this behavior. -// -// There are three types of assign operator: -// - Assign a real value from the same scalar type -// - In std, this is templated as complex& operator=(const T& x) -// with specialization `complex& operator=(T x)` for float/double/long -// double Since we only support float and double, on will use `complex& -// operator=(T x)` -// - Copy assignment operator and converting assignment operator -// - There is no specialization of converting assignment operators, which type -// is -// convertible is solely dependent on whether the scalar type is convertible -// -// In addition to the standard assignment, we also provide assignment operators -// with std and thrust -// -// -// [Casting operators] -// -// std::complex does not have casting operators. We define casting operators -// casting to std::complex and thrust::complex -// -// -// [Operator ""] -// -// std::complex has custom literals `i`, `if` and `il` defined in namespace -// `std::literals::complex_literals`. We define our own custom literals in the -// namespace `c10::complex_literals`. Our custom literals does not follow the -// same behavior as in std::complex, instead, we define _if, _id to construct -// float/double complex literals. -// -// -// [real() and imag()] -// -// In C++20, there are two overload of these functions, one it to return the -// real/imag, another is to set real/imag, they are both constexpr. We follow -// this design. -// -// -// [Operator +=,-=,*=,/=] -// -// Since C++20, these operators become constexpr. In our implementation, they -// are also constexpr. -// -// There are two types of such operators: operating with a real number, or -// operating with another complex number. For the operating with a real number, -// the generic template form has argument type `const T &`, while the overload -// for float/double/long double has `T`. We will follow the same type as -// float/double/long double in std. -// -// [Unary operator +-] -// -// Since C++20, they are constexpr. We also make them expr -// -// [Binary operators +-*/] -// -// Each operator has three versions (taking + as example): -// - complex + complex -// - complex + real -// - real + complex -// -// [Operator ==, !=] -// -// Each operator has three versions (taking == as example): -// - complex == complex -// - complex == real -// - real == complex -// -// Some of them are removed on C++20, but we decide to keep them -// -// [Operator <<, >>] -// -// These are implemented by casting to std::complex -// -// -// -// TODO(@zasdfgbnm): c10::complex is not currently supported, -// because: -// - lots of members and functions of c10::Half are not constexpr -// - thrust::complex only support float and double - -template -struct alignas(sizeof(T) * 2) complex { - using value_type = T; - - T real_ = T(0); - T imag_ = T(0); - - constexpr complex() = default; - C10_HOST_DEVICE constexpr complex(const T& re, const T& im = T()) - : real_(re), imag_(im) {} - template - explicit constexpr complex(const std::complex& other) - : complex(other.real(), other.imag()) {} -#if defined(__CUDACC__) || defined(__HIPCC__) - template - explicit C10_HOST_DEVICE complex(const thrust::complex& other) - : real_(other.real()), imag_(other.imag()) {} -// NOTE can not be implemented as follow due to ROCm bug: -// explicit C10_HOST_DEVICE complex(const thrust::complex &other): -// complex(other.real(), other.imag()) {} -#endif - - // Use SFINAE to specialize casting constructor for c10::complex and - // c10::complex - template - C10_HOST_DEVICE explicit constexpr complex( - const std::enable_if_t, complex>& other) - : real_(other.real_), imag_(other.imag_) {} - template - C10_HOST_DEVICE constexpr complex( - const std::enable_if_t, complex>& other) - : real_(other.real_), imag_(other.imag_) {} - - constexpr complex& operator=(T re) { - real_ = re; - imag_ = 0; - return *this; - } - - constexpr complex& operator+=(T re) { - real_ += re; - return *this; - } - - constexpr complex& operator-=(T re) { - real_ -= re; - return *this; - } - - constexpr complex& operator*=(T re) { - real_ *= re; - imag_ *= re; - return *this; - } - - constexpr complex& operator/=(T re) { - real_ /= re; - imag_ /= re; - return *this; - } - - template - constexpr complex& operator=(const complex& rhs) { - real_ = rhs.real(); - imag_ = rhs.imag(); - return *this; - } - - template - constexpr complex& operator+=(const complex& rhs) { - real_ += rhs.real(); - imag_ += rhs.imag(); - return *this; - } - - template - constexpr complex& operator-=(const complex& rhs) { - real_ -= rhs.real(); - imag_ -= rhs.imag(); - return *this; - } - - template - constexpr complex& operator*=(const complex& rhs) { - // (a + bi) * (c + di) = (a*c - b*d) + (a * d + b * c) i - T a = real_; - T b = imag_; - U c = rhs.real(); - U d = rhs.imag(); - real_ = a * c - b * d; - imag_ = a * d + b * c; - return *this; - } - -#ifdef __APPLE__ -#define FORCE_INLINE_APPLE __attribute__((always_inline)) -#else -#define FORCE_INLINE_APPLE -#endif - template - constexpr FORCE_INLINE_APPLE complex& operator/=(const complex& rhs) - __ubsan_ignore_float_divide_by_zero__ { - // (a + bi) / (c + di) = (ac + bd)/(c^2 + d^2) + (bc - ad)/(c^2 + d^2) i - // the calculation below follows numpy's complex division - T a = real_; - T b = imag_; - U c = rhs.real(); - U d = rhs.imag(); - -#if defined(__GNUC__) && !defined(__clang__) - // std::abs is already constexpr by gcc - auto abs_c = std::abs(c); - auto abs_d = std::abs(d); -#else - auto abs_c = c < 0 ? -c : c; - auto abs_d = d < 0 ? -d : d; -#endif - - if (abs_c >= abs_d) { - if (abs_c == U(0) && abs_d == U(0)) { - /* divide by zeros should yield a complex inf or nan */ - real_ = a / abs_c; - imag_ = b / abs_d; - } else { - auto rat = d / c; - auto scl = U(1.0) / (c + d * rat); - real_ = (a + b * rat) * scl; - imag_ = (b - a * rat) * scl; - } - } else { - auto rat = c / d; - auto scl = U(1.0) / (d + c * rat); - real_ = (a * rat + b) * scl; - imag_ = (b * rat - a) * scl; - } - return *this; - } -#undef FORCE_INLINE_APPLE - - template - constexpr complex& operator=(const std::complex& rhs) { - real_ = rhs.real(); - imag_ = rhs.imag(); - return *this; - } - -#if defined(__CUDACC__) || defined(__HIPCC__) - template - C10_HOST_DEVICE complex& operator=(const thrust::complex& rhs) { - real_ = rhs.real(); - imag_ = rhs.imag(); - return *this; - } -#endif - - template - explicit constexpr operator std::complex() const { - return std::complex(std::complex(real(), imag())); - } - -#if defined(__CUDACC__) || defined(__HIPCC__) - template - C10_HOST_DEVICE explicit operator thrust::complex() const { - return static_cast>(thrust::complex(real(), imag())); - } -#endif - - // consistent with NumPy behavior - explicit constexpr operator bool() const { - return real() || imag(); - } - - C10_HOST_DEVICE constexpr T real() const { - return real_; - } - constexpr void real(T value) { - real_ = value; - } - C10_HOST_DEVICE constexpr T imag() const { - return imag_; - } - constexpr void imag(T value) { - imag_ = value; - } -}; - -namespace complex_literals { - -constexpr complex operator""_if(long double imag) { - return complex(0.0f, static_cast(imag)); -} - -constexpr complex operator""_id(long double imag) { - return complex(0.0, static_cast(imag)); -} - -constexpr complex operator""_if(unsigned long long imag) { - return complex(0.0f, static_cast(imag)); -} - -constexpr complex operator""_id(unsigned long long imag) { - return complex(0.0, static_cast(imag)); -} - -} // namespace complex_literals - -template -constexpr complex operator+(const complex& val) { - return val; -} - -template -constexpr complex operator-(const complex& val) { - return complex(-val.real(), -val.imag()); -} - -template -constexpr complex operator+(const complex& lhs, const complex& rhs) { - complex result = lhs; - return result += rhs; -} - -template -constexpr complex operator+(const complex& lhs, const T& rhs) { - complex result = lhs; - return result += rhs; -} - -template -constexpr complex operator+(const T& lhs, const complex& rhs) { - return complex(lhs + rhs.real(), rhs.imag()); -} - -template -constexpr complex operator-(const complex& lhs, const complex& rhs) { - complex result = lhs; - return result -= rhs; -} - -template -constexpr complex operator-(const complex& lhs, const T& rhs) { - complex result = lhs; - return result -= rhs; -} - -template -constexpr complex operator-(const T& lhs, const complex& rhs) { - complex result = -rhs; - return result += lhs; -} - -template -constexpr complex operator*(const complex& lhs, const complex& rhs) { - complex result = lhs; - return result *= rhs; -} - -template -constexpr complex operator*(const complex& lhs, const T& rhs) { - complex result = lhs; - return result *= rhs; -} - -template -constexpr complex operator*(const T& lhs, const complex& rhs) { - complex result = rhs; - return result *= lhs; -} - -template -constexpr complex operator/(const complex& lhs, const complex& rhs) { - complex result = lhs; - return result /= rhs; -} - -template -constexpr complex operator/(const complex& lhs, const T& rhs) { - complex result = lhs; - return result /= rhs; -} - -template -constexpr complex operator/(const T& lhs, const complex& rhs) { - complex result(lhs, T()); - return result /= rhs; -} - -// Define operators between integral scalars and c10::complex. std::complex does -// not support this when T is a floating-point number. This is useful because it -// saves a lot of "static_cast" when operate a complex and an integer. This -// makes the code both less verbose and potentially more efficient. -#define COMPLEX_INTEGER_OP_TEMPLATE_CONDITION \ - typename std::enable_if_t< \ - std::is_floating_point_v && std::is_integral_v, \ - int> = 0 - -template -constexpr c10::complex operator+(const c10::complex& a, const iT& b) { - return a + static_cast(b); -} - -template -constexpr c10::complex operator+(const iT& a, const c10::complex& b) { - return static_cast(a) + b; -} - -template -constexpr c10::complex operator-(const c10::complex& a, const iT& b) { - return a - static_cast(b); -} - -template -constexpr c10::complex operator-(const iT& a, const c10::complex& b) { - return static_cast(a) - b; -} - -template -constexpr c10::complex operator*(const c10::complex& a, const iT& b) { - return a * static_cast(b); -} - -template -constexpr c10::complex operator*(const iT& a, const c10::complex& b) { - return static_cast(a) * b; -} - -template -constexpr c10::complex operator/(const c10::complex& a, const iT& b) { - return a / static_cast(b); -} - -template -constexpr c10::complex operator/(const iT& a, const c10::complex& b) { - return static_cast(a) / b; -} - -#undef COMPLEX_INTEGER_OP_TEMPLATE_CONDITION - -template -constexpr bool operator==(const complex& lhs, const complex& rhs) { - return (lhs.real() == rhs.real()) && (lhs.imag() == rhs.imag()); -} - -template -constexpr bool operator==(const complex& lhs, const T& rhs) { - return (lhs.real() == rhs) && (lhs.imag() == T()); -} - -template -constexpr bool operator==(const T& lhs, const complex& rhs) { - return (lhs == rhs.real()) && (T() == rhs.imag()); -} - -template -constexpr bool operator!=(const complex& lhs, const complex& rhs) { - return !(lhs == rhs); -} - -template -constexpr bool operator!=(const complex& lhs, const T& rhs) { - return !(lhs == rhs); -} - -template -constexpr bool operator!=(const T& lhs, const complex& rhs) { - return !(lhs == rhs); -} - -template -std::basic_ostream& operator<<( - std::basic_ostream& os, - const complex& x) { - return (os << static_cast>(x)); -} - -template -std::basic_istream& operator>>( - std::basic_istream& is, - complex& x) { - std::complex tmp; - is >> tmp; - x = tmp; - return is; -} - -} // namespace c10 +#include // std functions // @@ -594,72 +70,6 @@ constexpr c10::complex conj(const c10::complex& z) { } // namespace std -namespace c10 { - -template -C10_HOST_DEVICE complex polar(const T& r, const T& theta = T()) { -#if defined(__CUDACC__) || defined(__HIPCC__) - return static_cast>(thrust::polar(r, theta)); -#else - // std::polar() requires r >= 0, so spell out the explicit implementation to - // avoid a branch. - return complex(r * std::cos(theta), r * std::sin(theta)); -#endif -} - -template <> -struct alignas(4) complex { - Half real_; - Half imag_; - - // Constructors - complex() = default; - // Half constructor is not constexpr so the following constructor can't - // be constexpr - C10_HOST_DEVICE explicit inline complex(const Half& real, const Half& imag) - : real_(real), imag_(imag) {} - C10_HOST_DEVICE inline complex(const c10::complex& value) - : real_(value.real()), imag_(value.imag()) {} - - // Conversion operator - inline C10_HOST_DEVICE operator c10::complex() const { - return {real_, imag_}; - } - - constexpr C10_HOST_DEVICE Half real() const { - return real_; - } - constexpr C10_HOST_DEVICE Half imag() const { - return imag_; - } - - C10_HOST_DEVICE complex& operator+=(const complex& other) { - real_ = static_cast(real_) + static_cast(other.real_); - imag_ = static_cast(imag_) + static_cast(other.imag_); - return *this; - } - - C10_HOST_DEVICE complex& operator-=(const complex& other) { - real_ = static_cast(real_) - static_cast(other.real_); - imag_ = static_cast(imag_) - static_cast(other.imag_); - return *this; - } - - C10_HOST_DEVICE complex& operator*=(const complex& other) { - auto a = static_cast(real_); - auto b = static_cast(imag_); - auto c = static_cast(other.real()); - auto d = static_cast(other.imag()); - real_ = a * c - b * d; - imag_ = a * d + b * c; - return *this; - } -}; - -} // namespace c10 - -C10_CLANG_DIAGNOSTIC_POP() - #define C10_INTERNAL_INCLUDE_COMPLEX_REMAINING_H // math functions are included in a separate file #include // IWYU pragma: keep diff --git a/runtime/core/portable_type/c10/c10/util/floating_point_utils.h b/runtime/core/portable_type/c10/c10/util/floating_point_utils.h index b240c4ea232..10aa67c7cb8 100644 --- a/runtime/core/portable_type/c10/c10/util/floating_point_utils.h +++ b/runtime/core/portable_type/c10/c10/util/floating_point_utils.h @@ -1,33 +1 @@ -#pragma once - -#include -#include -#include - -namespace c10::detail { - -C10_HOST_DEVICE inline float fp32_from_bits(uint32_t w) { -#if defined(__OPENCL_VERSION__) - return as_float(w); -#elif defined(__CUDA_ARCH__) || defined(__HIP_DEVICE_COMPILE__) - return __uint_as_float((unsigned int)w); -#elif defined(__INTEL_COMPILER) - return _castu32_f32(w); -#else - return c10::bit_cast(w); -#endif -} - -C10_HOST_DEVICE inline uint32_t fp32_to_bits(float f) { -#if defined(__OPENCL_VERSION__) - return as_uint(f); -#elif defined(__CUDA_ARCH__) || defined(__HIP_DEVICE_COMPILE__) - return (uint32_t)__float_as_uint(f); -#elif defined(__INTEL_COMPILER) - return _castf32_u32(f); -#else - return c10::bit_cast(f); -#endif -} - -} // namespace c10::detail +#include diff --git a/runtime/core/portable_type/c10/torch/headeronly/util/BFloat16.h b/runtime/core/portable_type/c10/torch/headeronly/util/BFloat16.h new file mode 100644 index 00000000000..2c1f805ac7b --- /dev/null +++ b/runtime/core/portable_type/c10/torch/headeronly/util/BFloat16.h @@ -0,0 +1,478 @@ +#pragma once + +// Defines the bloat16 type (brain floating-point). This representation uses +// 1 bit for the sign, 8 bits for the exponent and 7 bits for the mantissa. + +#include +#include + +#include +#include +#include +#include +#include + +#if defined(__CUDACC__) && !defined(USE_ROCM) +#include +#endif + +#if defined(CL_SYCL_LANGUAGE_VERSION) +#include // for SYCL 1.2.1 +#elif defined(SYCL_LANGUAGE_VERSION) +#include // for SYCL 2020 +#endif + +namespace c10 { + +struct alignas(2) BFloat16 { + uint16_t x; + + // HIP wants __host__ __device__ tag, CUDA does not +#if defined(USE_ROCM) && defined(__HIPCC__) + C10_HOST_DEVICE BFloat16() = default; +#else + BFloat16() = default; +#endif + + struct from_bits_t {}; + static constexpr C10_HOST_DEVICE from_bits_t from_bits() { + return from_bits_t(); + } + + constexpr C10_HOST_DEVICE BFloat16(unsigned short bits, from_bits_t) + : x(bits) {} + /* implicit */ inline C10_HOST_DEVICE BFloat16(float value); + inline C10_HOST_DEVICE operator float() const; + +#if defined(__CUDACC__) && !defined(USE_ROCM) + inline C10_HOST_DEVICE BFloat16(const __nv_bfloat16& value); + explicit inline C10_HOST_DEVICE operator __nv_bfloat16() const; +#endif + +#if defined(SYCL_EXT_ONEAPI_BFLOAT16_MATH_FUNCTIONS) + inline C10_HOST_DEVICE BFloat16(const sycl::ext::oneapi::bfloat16& value); + explicit inline C10_HOST_DEVICE operator sycl::ext::oneapi::bfloat16() const; +#endif +}; + +inline std::ostream& operator<<(std::ostream& out, const BFloat16& value) { + out << (float)value; + return out; +} + +namespace detail { +inline C10_HOST_DEVICE float f32_from_bits(uint16_t src) { + float res = 0; + uint32_t tmp = src; + tmp <<= 16; + +#if defined(USE_ROCM) && defined(__HIPCC__) + float* tempRes; + + // We should be using memcpy in order to respect the strict aliasing rule + // but it fails in the HIP environment. + tempRes = reinterpret_cast(&tmp); + res = *tempRes; +#else + std::memcpy(&res, &tmp, sizeof(tmp)); +#endif + + return res; +} + +inline C10_HOST_DEVICE uint16_t bits_from_f32(float src) { + uint32_t res = 0; + +#if defined(USE_ROCM) && defined(__HIPCC__) + // We should be using memcpy in order to respect the strict aliasing rule + // but it fails in the HIP environment. + uint32_t* tempRes = reinterpret_cast(&src); + res = *tempRes; +#else + std::memcpy(&res, &src, sizeof(res)); +#endif + + return res >> 16; +} + +inline C10_HOST_DEVICE uint16_t round_to_nearest_even(float src) { +#if defined(USE_ROCM) && defined(__HIPCC__) + if (src != src) { +#elif defined(_MSC_VER) + if (isnan(src)) { +#else + if (std::isnan(src)) { +#endif + return UINT16_C(0x7FC0); + } else { + const uint32_t U32 = c10::bit_cast(src); + uint32_t rounding_bias = ((U32 >> 16) & 1) + UINT32_C(0x7FFF); + return static_cast((U32 + rounding_bias) >> 16); + } +} + +} // namespace detail + +//-------- the following is copied from c10/util/BFloat16-inl.h ---------// +C10_CLANG_DIAGNOSTIC_PUSH() +#if C10_CLANG_HAS_WARNING("-Wimplicit-int-float-conversion") +C10_CLANG_DIAGNOSTIC_IGNORE("-Wimplicit-int-float-conversion") +#endif + +/// Constructors +inline C10_HOST_DEVICE BFloat16::BFloat16(float value) + : +#if defined(__CUDACC__) && !defined(USE_ROCM) && defined(__CUDA_ARCH__) && \ + __CUDA_ARCH__ >= 800 + x(__bfloat16_as_ushort(__float2bfloat16(value))) +#elif defined(__SYCL_DEVICE_ONLY__) && \ + defined(SYCL_EXT_ONEAPI_BFLOAT16_MATH_FUNCTIONS) + x(c10::bit_cast(sycl::ext::oneapi::bfloat16(value))) +#else + // RNE by default + x(detail::round_to_nearest_even(value)) +#endif +{ +} + +/// Implicit conversions +inline C10_HOST_DEVICE BFloat16::operator float() const { +#if defined(__CUDACC__) && !defined(USE_ROCM) + return __bfloat162float(*reinterpret_cast(&x)); +#elif defined(__SYCL_DEVICE_ONLY__) && \ + defined(SYCL_EXT_ONEAPI_BFLOAT16_MATH_FUNCTIONS) + return float(*reinterpret_cast(&x)); +#else + return detail::f32_from_bits(x); +#endif +} + +#if defined(__CUDACC__) && !defined(USE_ROCM) +inline C10_HOST_DEVICE BFloat16::BFloat16(const __nv_bfloat16& value) { + x = *reinterpret_cast(&value); +} +inline C10_HOST_DEVICE BFloat16::operator __nv_bfloat16() const { + return *reinterpret_cast(&x); +} +#endif + +#if defined(SYCL_EXT_ONEAPI_BFLOAT16_MATH_FUNCTIONS) +inline C10_HOST_DEVICE BFloat16::BFloat16( + const sycl::ext::oneapi::bfloat16& value) { + x = *reinterpret_cast(&value); +} +inline C10_HOST_DEVICE BFloat16::operator sycl::ext::oneapi::bfloat16() const { + return *reinterpret_cast(&x); +} +#endif + +// CUDA intrinsics + +#if defined(__CUDACC__) || defined(__HIPCC__) +inline C10_DEVICE BFloat16 __ldg(const BFloat16* ptr) { +#if !defined(USE_ROCM) && defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 800 + return __ldg(reinterpret_cast(ptr)); +#else + return *ptr; +#endif +} +#endif + +/// Arithmetic + +inline C10_HOST_DEVICE BFloat16 +operator+(const BFloat16& a, const BFloat16& b) { + return static_cast(a) + static_cast(b); +} + +inline C10_HOST_DEVICE BFloat16 +operator-(const BFloat16& a, const BFloat16& b) { + return static_cast(a) - static_cast(b); +} + +inline C10_HOST_DEVICE BFloat16 +operator*(const BFloat16& a, const BFloat16& b) { + return static_cast(a) * static_cast(b); +} + +inline C10_HOST_DEVICE BFloat16 operator/(const BFloat16& a, const BFloat16& b) + __ubsan_ignore_float_divide_by_zero__ { + return static_cast(a) / static_cast(b); +} + +inline C10_HOST_DEVICE BFloat16 operator-(const BFloat16& a) { + return -static_cast(a); +} + +inline C10_HOST_DEVICE BFloat16& operator+=(BFloat16& a, const BFloat16& b) { + a = a + b; + return a; +} + +inline C10_HOST_DEVICE BFloat16& operator-=(BFloat16& a, const BFloat16& b) { + a = a - b; + return a; +} + +inline C10_HOST_DEVICE BFloat16& operator*=(BFloat16& a, const BFloat16& b) { + a = a * b; + return a; +} + +inline C10_HOST_DEVICE BFloat16& operator/=(BFloat16& a, const BFloat16& b) { + a = a / b; + return a; +} + +inline C10_HOST_DEVICE BFloat16& operator|(BFloat16& a, const BFloat16& b) { + a.x = a.x | b.x; + return a; +} + +inline C10_HOST_DEVICE BFloat16& operator^(BFloat16& a, const BFloat16& b) { + a.x = a.x ^ b.x; + return a; +} + +inline C10_HOST_DEVICE BFloat16& operator&(BFloat16& a, const BFloat16& b) { + a.x = a.x & b.x; + return a; +} + +/// Arithmetic with floats + +inline C10_HOST_DEVICE float operator+(BFloat16 a, float b) { + return static_cast(a) + b; +} +inline C10_HOST_DEVICE float operator-(BFloat16 a, float b) { + return static_cast(a) - b; +} +inline C10_HOST_DEVICE float operator*(BFloat16 a, float b) { + return static_cast(a) * b; +} +inline C10_HOST_DEVICE float operator/(BFloat16 a, float b) { + return static_cast(a) / b; +} + +inline C10_HOST_DEVICE float operator+(float a, BFloat16 b) { + return a + static_cast(b); +} +inline C10_HOST_DEVICE float operator-(float a, BFloat16 b) { + return a - static_cast(b); +} +inline C10_HOST_DEVICE float operator*(float a, BFloat16 b) { + return a * static_cast(b); +} +inline C10_HOST_DEVICE float operator/(float a, BFloat16 b) { + return a / static_cast(b); +} + +inline C10_HOST_DEVICE float& operator+=(float& a, const BFloat16& b) { + return a += static_cast(b); +} +inline C10_HOST_DEVICE float& operator-=(float& a, const BFloat16& b) { + return a -= static_cast(b); +} +inline C10_HOST_DEVICE float& operator*=(float& a, const BFloat16& b) { + return a *= static_cast(b); +} +inline C10_HOST_DEVICE float& operator/=(float& a, const BFloat16& b) { + return a /= static_cast(b); +} + +/// Arithmetic with doubles + +inline C10_HOST_DEVICE double operator+(BFloat16 a, double b) { + return static_cast(a) + b; +} +inline C10_HOST_DEVICE double operator-(BFloat16 a, double b) { + return static_cast(a) - b; +} +inline C10_HOST_DEVICE double operator*(BFloat16 a, double b) { + return static_cast(a) * b; +} +inline C10_HOST_DEVICE double operator/(BFloat16 a, double b) { + return static_cast(a) / b; +} + +inline C10_HOST_DEVICE double operator+(double a, BFloat16 b) { + return a + static_cast(b); +} +inline C10_HOST_DEVICE double operator-(double a, BFloat16 b) { + return a - static_cast(b); +} +inline C10_HOST_DEVICE double operator*(double a, BFloat16 b) { + return a * static_cast(b); +} +inline C10_HOST_DEVICE double operator/(double a, BFloat16 b) { + return a / static_cast(b); +} + +/// Arithmetic with ints + +inline C10_HOST_DEVICE BFloat16 operator+(BFloat16 a, int b) { + // NOLINTNEXTLINE(bugprone-narrowing-conversions,cppcoreguidelines-narrowing-conversions) + return a + static_cast(b); +} +inline C10_HOST_DEVICE BFloat16 operator-(BFloat16 a, int b) { + // NOLINTNEXTLINE(bugprone-narrowing-conversions,cppcoreguidelines-narrowing-conversions) + return a - static_cast(b); +} +inline C10_HOST_DEVICE BFloat16 operator*(BFloat16 a, int b) { + // NOLINTNEXTLINE(bugprone-narrowing-conversions,cppcoreguidelines-narrowing-conversions) + return a * static_cast(b); +} +inline C10_HOST_DEVICE BFloat16 operator/(BFloat16 a, int b) { + // NOLINTNEXTLINE(bugprone-narrowing-conversions,cppcoreguidelines-narrowing-conversions) + return a / static_cast(b); +} + +inline C10_HOST_DEVICE BFloat16 operator+(int a, BFloat16 b) { + // NOLINTNEXTLINE(bugprone-narrowing-conversions,cppcoreguidelines-narrowing-conversions) + return static_cast(a) + b; +} +inline C10_HOST_DEVICE BFloat16 operator-(int a, BFloat16 b) { + // NOLINTNEXTLINE(bugprone-narrowing-conversions,cppcoreguidelines-narrowing-conversions) + return static_cast(a) - b; +} +inline C10_HOST_DEVICE BFloat16 operator*(int a, BFloat16 b) { + // NOLINTNEXTLINE(bugprone-narrowing-conversions,cppcoreguidelines-narrowing-conversions) + return static_cast(a) * b; +} +inline C10_HOST_DEVICE BFloat16 operator/(int a, BFloat16 b) { + // NOLINTNEXTLINE(bugprone-narrowing-conversions,cppcoreguidelines-narrowing-conversions) + return static_cast(a) / b; +} + +//// Arithmetic with int64_t + +inline C10_HOST_DEVICE BFloat16 operator+(BFloat16 a, int64_t b) { + // NOLINTNEXTLINE(bugprone-narrowing-conversions,cppcoreguidelines-narrowing-conversions) + return a + static_cast(b); +} +inline C10_HOST_DEVICE BFloat16 operator-(BFloat16 a, int64_t b) { + // NOLINTNEXTLINE(bugprone-narrowing-conversions,cppcoreguidelines-narrowing-conversions) + return a - static_cast(b); +} +inline C10_HOST_DEVICE BFloat16 operator*(BFloat16 a, int64_t b) { + // NOLINTNEXTLINE(bugprone-narrowing-conversions,cppcoreguidelines-narrowing-conversions) + return a * static_cast(b); +} +inline C10_HOST_DEVICE BFloat16 operator/(BFloat16 a, int64_t b) { + // NOLINTNEXTLINE(bugprone-narrowing-conversions,cppcoreguidelines-narrowing-conversions) + return a / static_cast(b); +} + +inline C10_HOST_DEVICE BFloat16 operator+(int64_t a, BFloat16 b) { + // NOLINTNEXTLINE(bugprone-narrowing-conversions,cppcoreguidelines-narrowing-conversions) + return static_cast(a) + b; +} +inline C10_HOST_DEVICE BFloat16 operator-(int64_t a, BFloat16 b) { + // NOLINTNEXTLINE(bugprone-narrowing-conversions,cppcoreguidelines-narrowing-conversions) + return static_cast(a) - b; +} +inline C10_HOST_DEVICE BFloat16 operator*(int64_t a, BFloat16 b) { + // NOLINTNEXTLINE(bugprone-narrowing-conversions,cppcoreguidelines-narrowing-conversions) + return static_cast(a) * b; +} +inline C10_HOST_DEVICE BFloat16 operator/(int64_t a, BFloat16 b) { + // NOLINTNEXTLINE(bugprone-narrowing-conversions,cppcoreguidelines-narrowing-conversions) + return static_cast(a) / b; +} + +// Overloading < and > operators, because std::max and std::min use them. + +inline C10_HOST_DEVICE bool operator>(BFloat16& lhs, BFloat16& rhs) { + return float(lhs) > float(rhs); +} + +inline C10_HOST_DEVICE bool operator<(BFloat16& lhs, BFloat16& rhs) { + return float(lhs) < float(rhs); +} + +C10_CLANG_DIAGNOSTIC_POP() +} // namespace c10 + +namespace torch::headeronly { + +namespace detail { +using c10::detail::bits_from_f32; +using c10::detail::f32_from_bits; +using c10::detail::round_to_nearest_even; +} // namespace detail + +using c10::BFloat16; +using c10::operator+; +using c10::operator-; +using c10::operator*; +using c10::operator/; +using c10::operator+=; +using c10::operator-=; +using c10::operator*=; +using c10::operator/=; +using c10::operator<; +using c10::operator>; +using c10::operator<<; +} // namespace torch::headeronly + +namespace std { + +template <> +class numeric_limits { + public: + static constexpr bool is_signed = true; + static constexpr bool is_specialized = true; + static constexpr bool is_integer = false; + static constexpr bool is_exact = false; + static constexpr bool has_infinity = true; + static constexpr bool has_quiet_NaN = true; + static constexpr bool has_signaling_NaN = true; + static constexpr auto has_denorm = numeric_limits::has_denorm; + static constexpr auto has_denorm_loss = + numeric_limits::has_denorm_loss; + static constexpr auto round_style = numeric_limits::round_style; + static constexpr bool is_iec559 = false; + static constexpr bool is_bounded = true; + static constexpr bool is_modulo = false; + static constexpr int digits = 8; + static constexpr int digits10 = 2; + static constexpr int max_digits10 = 4; + static constexpr int radix = 2; + static constexpr int min_exponent = -125; + static constexpr int min_exponent10 = -37; + static constexpr int max_exponent = 128; + static constexpr int max_exponent10 = 38; + static constexpr auto traps = numeric_limits::traps; + static constexpr auto tinyness_before = + numeric_limits::tinyness_before; + + static constexpr c10::BFloat16 min() { + return c10::BFloat16(0x0080, c10::BFloat16::from_bits()); + } + static constexpr c10::BFloat16 lowest() { + return c10::BFloat16(0xFF7F, c10::BFloat16::from_bits()); + } + static constexpr c10::BFloat16 max() { + return c10::BFloat16(0x7F7F, c10::BFloat16::from_bits()); + } + static constexpr c10::BFloat16 epsilon() { + return c10::BFloat16(0x3C00, c10::BFloat16::from_bits()); + } + static constexpr c10::BFloat16 round_error() { + return c10::BFloat16(0x3F00, c10::BFloat16::from_bits()); + } + static constexpr c10::BFloat16 infinity() { + return c10::BFloat16(0x7F80, c10::BFloat16::from_bits()); + } + static constexpr c10::BFloat16 quiet_NaN() { + return c10::BFloat16(0x7FC0, c10::BFloat16::from_bits()); + } + static constexpr c10::BFloat16 signaling_NaN() { + return c10::BFloat16(0x7F80, c10::BFloat16::from_bits()); + } + static constexpr c10::BFloat16 denorm_min() { + return c10::BFloat16(0x0001, c10::BFloat16::from_bits()); + } +}; + +} // namespace std diff --git a/runtime/core/portable_type/c10/torch/headeronly/util/Half.h b/runtime/core/portable_type/c10/torch/headeronly/util/Half.h new file mode 100644 index 00000000000..59a86f07e33 --- /dev/null +++ b/runtime/core/portable_type/c10/torch/headeronly/util/Half.h @@ -0,0 +1,787 @@ +#pragma once + +/// Defines the Half type (half-precision floating-point) including conversions +/// to standard C types and basic arithmetic operations. Note that arithmetic +/// operations are implemented by converting to floating point and +/// performing the operation in float32, instead of using CUDA half intrinsics. +/// Most uses of this type within ATen are memory bound, including the +/// element-wise kernels, and the half intrinsics aren't efficient on all GPUs. +/// If you are writing a compute bound kernel, you can use the CUDA half +/// intrinsics directly on the Half type from device code. + +#include +#include +#include + +#if defined(__cplusplus) +#include +#elif !defined(__OPENCL_VERSION__) +#include +#endif + +#ifdef _MSC_VER +#include +#endif + +#include +#include +#include + +#ifdef __CUDACC__ +#include +#endif + +#ifdef __HIPCC__ +#include +#endif + +#if defined(CL_SYCL_LANGUAGE_VERSION) +#include // for SYCL 1.2.1 +#elif defined(SYCL_LANGUAGE_VERSION) +#include // for SYCL 2020 +#endif + +#if (defined(CPU_CAPABILITY_AVX2) || defined(CPU_CAPABILITY_AVX512)) && \ + !defined(__APPLE__) +#include +#endif + +#if defined(__aarch64__) && !defined(__CUDACC__) +#include +#endif + +#if defined(__GNUC__) || defined(__clang__) +#if defined(__x86_64__) || defined(_M_X64) || defined(__i386) || \ + defined(_M_IX86) +#if defined(__F16C__) && \ + !(defined(__CUDA_ARCH__) || defined(__CUDACC__) || \ + defined(__HIP_DEVICE_COMPILE__)) +#define C10_X86_F16 1 +#include // import conversion ops from f16cintrin.h +#endif // defined(__F16C__) && !(defined(__CUDA_ARCH__) || defined(__CUDACC__) + // || defined(__HIP_DEVICE_COMPILE__)) +#endif // __x86_64__ || _M_X64 || __i386 || _M_IX86 +#endif // __GNUC__ || __clang__ + +namespace c10 { + +struct alignas(2) Half { + unsigned short x; + + struct from_bits_t {}; + C10_HOST_DEVICE static constexpr from_bits_t from_bits() { + return from_bits_t(); + } + + // HIP wants __host__ __device__ tag, CUDA does not +#if defined(USE_ROCM) + C10_HOST_DEVICE Half() = default; +#else + Half() = default; +#endif + + constexpr C10_HOST_DEVICE Half(unsigned short bits, from_bits_t) : x(bits) {} +#if defined(__aarch64__) && !defined(__CUDACC__) + inline Half(float16_t value); + inline operator float16_t() const; +#else + inline C10_HOST_DEVICE Half(float value); + inline C10_HOST_DEVICE operator float() const; +#endif + +#if defined(__CUDACC__) || defined(__HIPCC__) + inline C10_HOST_DEVICE Half(const __half& value); + inline C10_HOST_DEVICE operator __half() const; +#endif +#ifdef SYCL_LANGUAGE_VERSION + inline C10_HOST_DEVICE Half(const sycl::half& value); + inline C10_HOST_DEVICE operator sycl::half() const; +#endif +}; + +inline std::ostream& operator<<(std::ostream& out, const Half& value) { + out << (float)value; + return out; +} + +namespace detail { +/* + * Convert a 16-bit floating-point number in IEEE half-precision format, in bit + * representation, to a 32-bit floating-point number in IEEE single-precision + * format. + * + * @note The implementation relies on IEEE-like (no assumption about rounding + * mode and no operations on denormals) floating-point operations and bitcasts + * between integer and floating-point variables. + */ +C10_HOST_DEVICE inline float fp16_ieee_to_fp32_value(uint16_t h) { +#ifdef C10_X86_F16 + return _cvtsh_ss(h); +#else + /* + * Extend the half-precision floating-point number to 32 bits and shift to the + * upper part of the 32-bit word: + * +---+-----+------------+-------------------+ + * | S |EEEEE|MM MMMM MMMM|0000 0000 0000 0000| + * +---+-----+------------+-------------------+ + * Bits 31 26-30 16-25 0-15 + * + * S - sign bit, E - bits of the biased exponent, M - bits of the mantissa, 0 + * - zero bits. + */ + const uint32_t w = (uint32_t)h << 16; + /* + * Extract the sign of the input number into the high bit of the 32-bit word: + * + * +---+----------------------------------+ + * | S |0000000 00000000 00000000 00000000| + * +---+----------------------------------+ + * Bits 31 0-31 + */ + const uint32_t sign = w & UINT32_C(0x80000000); + /* + * Extract mantissa and biased exponent of the input number into the high bits + * of the 32-bit word: + * + * +-----+------------+---------------------+ + * |EEEEE|MM MMMM MMMM|0 0000 0000 0000 0000| + * +-----+------------+---------------------+ + * Bits 27-31 17-26 0-16 + */ + const uint32_t two_w = w + w; + + /* + * Shift mantissa and exponent into bits 23-28 and bits 13-22 so they become + * mantissa and exponent of a single-precision floating-point number: + * + * S|Exponent | Mantissa + * +-+---+-----+------------+----------------+ + * |0|000|EEEEE|MM MMMM MMMM|0 0000 0000 0000| + * +-+---+-----+------------+----------------+ + * Bits | 23-31 | 0-22 + * + * Next, there are some adjustments to the exponent: + * - The exponent needs to be corrected by the difference in exponent bias + * between single-precision and half-precision formats (0x7F - 0xF = 0x70) + * - Inf and NaN values in the inputs should become Inf and NaN values after + * conversion to the single-precision number. Therefore, if the biased + * exponent of the half-precision input was 0x1F (max possible value), the + * biased exponent of the single-precision output must be 0xFF (max possible + * value). We do this correction in two steps: + * - First, we adjust the exponent by (0xFF - 0x1F) = 0xE0 (see exp_offset + * below) rather than by 0x70 suggested by the difference in the exponent bias + * (see above). + * - Then we multiply the single-precision result of exponent adjustment by + * 2**(-112) to reverse the effect of exponent adjustment by 0xE0 less the + * necessary exponent adjustment by 0x70 due to difference in exponent bias. + * The floating-point multiplication hardware would ensure than Inf and + * NaN would retain their value on at least partially IEEE754-compliant + * implementations. + * + * Note that the above operations do not handle denormal inputs (where biased + * exponent == 0). However, they also do not operate on denormal inputs, and + * do not produce denormal results. + */ + constexpr uint32_t exp_offset = UINT32_C(0xE0) << 23; + // const float exp_scale = 0x1.0p-112f; + constexpr uint32_t scale_bits = (uint32_t)15 << 23; + float exp_scale_val = 0; +#if defined(_MSC_VER) && defined(__clang__) + __builtin_memcpy(&exp_scale_val, &scale_bits, sizeof(exp_scale_val)); +#else + std::memcpy(&exp_scale_val, &scale_bits, sizeof(exp_scale_val)); +#endif + + const float exp_scale = exp_scale_val; + const float normalized_value = + fp32_from_bits((two_w >> 4) + exp_offset) * exp_scale; + + /* + * Convert denormalized half-precision inputs into single-precision results + * (always normalized). Zero inputs are also handled here. + * + * In a denormalized number the biased exponent is zero, and mantissa has + * on-zero bits. First, we shift mantissa into bits 0-9 of the 32-bit word. + * + * zeros | mantissa + * +---------------------------+------------+ + * |0000 0000 0000 0000 0000 00|MM MMMM MMMM| + * +---------------------------+------------+ + * Bits 10-31 0-9 + * + * Now, remember that denormalized half-precision numbers are represented as: + * FP16 = mantissa * 2**(-24). + * The trick is to construct a normalized single-precision number with the + * same mantissa and thehalf-precision input and with an exponent which would + * scale the corresponding mantissa bits to 2**(-24). A normalized + * single-precision floating-point number is represented as: FP32 = (1 + + * mantissa * 2**(-23)) * 2**(exponent - 127) Therefore, when the biased + * exponent is 126, a unit change in the mantissa of the input denormalized + * half-precision number causes a change of the constructed single-precision + * number by 2**(-24), i.e. the same amount. + * + * The last step is to adjust the bias of the constructed single-precision + * number. When the input half-precision number is zero, the constructed + * single-precision number has the value of FP32 = 1 * 2**(126 - 127) = + * 2**(-1) = 0.5 Therefore, we need to subtract 0.5 from the constructed + * single-precision number to get the numerical equivalent of the input + * half-precision number. + */ + constexpr uint32_t magic_mask = UINT32_C(126) << 23; + constexpr float magic_bias = 0.5f; + const float denormalized_value = + fp32_from_bits((two_w >> 17) | magic_mask) - magic_bias; + + /* + * - Choose either results of conversion of input as a normalized number, or + * as a denormalized number, depending on the input exponent. The variable + * two_w contains input exponent in bits 27-31, therefore if its smaller than + * 2**27, the input is either a denormal number, or zero. + * - Combine the result of conversion of exponent and mantissa with the sign + * of the input number. + */ + constexpr uint32_t denormalized_cutoff = UINT32_C(1) << 27; + const uint32_t result = sign | + (two_w < denormalized_cutoff ? fp32_to_bits(denormalized_value) + : fp32_to_bits(normalized_value)); + return fp32_from_bits(result); +#endif // C10_X86_F16 +} + +/* + * Convert a 32-bit floating-point number in IEEE single-precision format to a + * 16-bit floating-point number in IEEE half-precision format, in bit + * representation. + * + * @note The implementation relies on IEEE-like (no assumption about rounding + * mode and no operations on denormals) floating-point operations and bitcasts + * between integer and floating-point variables. + */ +inline uint16_t fp16_ieee_from_fp32_value(float f) { +#ifdef C10_X86_F16 + return _cvtss_sh(f, _MM_FROUND_TO_NEAREST_INT); +#else + // const float scale_to_inf = 0x1.0p+112f; + // const float scale_to_zero = 0x1.0p-110f; + constexpr uint32_t scale_to_inf_bits = (uint32_t)239 << 23; + constexpr uint32_t scale_to_zero_bits = (uint32_t)17 << 23; + float scale_to_inf_val = 0, scale_to_zero_val = 0; + std::memcpy(&scale_to_inf_val, &scale_to_inf_bits, sizeof(scale_to_inf_val)); + std::memcpy( + &scale_to_zero_val, &scale_to_zero_bits, sizeof(scale_to_zero_val)); + const float scale_to_inf = scale_to_inf_val; + const float scale_to_zero = scale_to_zero_val; + +#if defined(_MSC_VER) && _MSC_VER == 1916 + float base = ((signbit(f) != 0 ? -f : f) * scale_to_inf) * scale_to_zero; +#else + float base = (fabsf(f) * scale_to_inf) * scale_to_zero; +#endif + + const uint32_t w = fp32_to_bits(f); + const uint32_t shl1_w = w + w; + const uint32_t sign = w & UINT32_C(0x80000000); + uint32_t bias = shl1_w & UINT32_C(0xFF000000); + if (bias < UINT32_C(0x71000000)) { + bias = UINT32_C(0x71000000); + } + + base = fp32_from_bits((bias >> 1) + UINT32_C(0x07800000)) + base; + const uint32_t bits = fp32_to_bits(base); + const uint32_t exp_bits = (bits >> 13) & UINT32_C(0x00007C00); + const uint32_t mantissa_bits = bits & UINT32_C(0x00000FFF); + const uint32_t nonsign = exp_bits + mantissa_bits; + return static_cast( + (sign >> 16) | + (shl1_w > UINT32_C(0xFF000000) ? UINT16_C(0x7E00) : nonsign)); +#endif // C10_X86_F16 +} + +/* + * Convert a 16-bit floating-point number in IEEE half-precision format, in bit + * representation, to a 32-bit floating-point number in IEEE single-precision + * format, in bit representation. + * + * @note The implementation doesn't use any floating-point operations. + */ +inline uint32_t fp16_ieee_to_fp32_bits(uint16_t h) { + /* + * Extend the half-precision floating-point number to 32 bits and shift to the + * upper part of the 32-bit word: + * +---+-----+------------+-------------------+ + * | S |EEEEE|MM MMMM MMMM|0000 0000 0000 0000| + * +---+-----+------------+-------------------+ + * Bits 31 26-30 16-25 0-15 + * + * S - sign bit, E - bits of the biased exponent, M - bits of the mantissa, 0 + * - zero bits. + */ + const uint32_t w = (uint32_t)h << 16; + /* + * Extract the sign of the input number into the high bit of the 32-bit word: + * + * +---+----------------------------------+ + * | S |0000000 00000000 00000000 00000000| + * +---+----------------------------------+ + * Bits 31 0-31 + */ + const uint32_t sign = w & UINT32_C(0x80000000); + /* + * Extract mantissa and biased exponent of the input number into the bits 0-30 + * of the 32-bit word: + * + * +---+-----+------------+-------------------+ + * | 0 |EEEEE|MM MMMM MMMM|0000 0000 0000 0000| + * +---+-----+------------+-------------------+ + * Bits 30 27-31 17-26 0-16 + */ + const uint32_t nonsign = w & UINT32_C(0x7FFFFFFF); + /* + * Renorm shift is the number of bits to shift mantissa left to make the + * half-precision number normalized. If the initial number is normalized, some + * of its high 6 bits (sign == 0 and 5-bit exponent) equals one. In this case + * renorm_shift == 0. If the number is denormalize, renorm_shift > 0. Note + * that if we shift denormalized nonsign by renorm_shift, the unit bit of + * mantissa will shift into exponent, turning the biased exponent into 1, and + * making mantissa normalized (i.e. without leading 1). + */ +#ifdef _MSC_VER + unsigned long nonsign_bsr; + _BitScanReverse(&nonsign_bsr, (unsigned long)nonsign); + uint32_t renorm_shift = (uint32_t)nonsign_bsr ^ 31; +#else + uint32_t renorm_shift = __builtin_clz(nonsign); +#endif + renorm_shift = renorm_shift > 5 ? renorm_shift - 5 : 0; + /* + * Iff half-precision number has exponent of 15, the addition overflows + * it into bit 31, and the subsequent shift turns the high 9 bits + * into 1. Thus inf_nan_mask == 0x7F800000 if the half-precision number + * had exponent of 15 (i.e. was NaN or infinity) 0x00000000 otherwise + */ + const int32_t inf_nan_mask = + ((int32_t)(nonsign + 0x04000000) >> 8) & INT32_C(0x7F800000); + /* + * Iff nonsign is 0, it overflows into 0xFFFFFFFF, turning bit 31 + * into 1. Otherwise, bit 31 remains 0. The signed shift right by 31 + * broadcasts bit 31 into all bits of the zero_mask. Thus zero_mask == + * 0xFFFFFFFF if the half-precision number was zero (+0.0h or -0.0h) + * 0x00000000 otherwise + */ + const int32_t zero_mask = (int32_t)(nonsign - 1) >> 31; + /* + * 1. Shift nonsign left by renorm_shift to normalize it (if the input + * was denormal) + * 2. Shift nonsign right by 3 so the exponent (5 bits originally) + * becomes an 8-bit field and 10-bit mantissa shifts into the 10 high + * bits of the 23-bit mantissa of IEEE single-precision number. + * 3. Add 0x70 to the exponent (starting at bit 23) to compensate the + * different in exponent bias (0x7F for single-precision number less 0xF + * for half-precision number). + * 4. Subtract renorm_shift from the exponent (starting at bit 23) to + * account for renormalization. As renorm_shift is less than 0x70, this + * can be combined with step 3. + * 5. Binary OR with inf_nan_mask to turn the exponent into 0xFF if the + * input was NaN or infinity. + * 6. Binary ANDNOT with zero_mask to turn the mantissa and exponent + * into zero if the input was zero. + * 7. Combine with the sign of the input number. + */ + return sign | + ((((nonsign << renorm_shift >> 3) + ((0x70 - renorm_shift) << 23)) | + inf_nan_mask) & + ~zero_mask); +} + +#ifdef C10_X86_F16 +#undef C10_X86_F16 +#endif // C10_X86_F16 + +#if defined(__aarch64__) && !defined(__CUDACC__) +inline float16_t fp16_from_bits(uint16_t h) { + return c10::bit_cast(h); +} + +inline uint16_t fp16_to_bits(float16_t f) { + return c10::bit_cast(f); +} + +// According to https://godbolt.org/z/frExdbsWG it would translate to single +// fcvt s0, h0 +inline float native_fp16_to_fp32_value(uint16_t h) { + return static_cast(fp16_from_bits(h)); +} + +inline uint16_t native_fp16_from_fp32_value(float f) { + return fp16_to_bits(static_cast(f)); +} +#endif + +} // namespace detail + +//---------- below is copied from c10/util/Half-inl.h ----------------// +C10_CLANG_DIAGNOSTIC_PUSH() +#if C10_CLANG_HAS_WARNING("-Wimplicit-int-float-conversion") +C10_CLANG_DIAGNOSTIC_IGNORE("-Wimplicit-int-float-conversion") +#endif + +#if defined(__aarch64__) && !defined(__CUDACC__) +/// Constructors +inline Half::Half(float16_t value) : x(detail::fp16_to_bits(value)) {} +inline Half::operator float16_t() const { + return detail::fp16_from_bits(x); +} +#else + +inline C10_HOST_DEVICE Half::Half(float value) + : +#if defined(__CUDA_ARCH__) || defined(__HIP_DEVICE_COMPILE__) + x(__half_as_short(__float2half(value))) +#elif defined(__SYCL_DEVICE_ONLY__) + x(c10::bit_cast(sycl::half(value))) +#elif (defined(CPU_CAPABILITY_AVX2) || defined(CPU_CAPABILITY_AVX512)) && \ + !defined(__APPLE__) + x(at::vec::float2half_scalar(value)) +#else + x(detail::fp16_ieee_from_fp32_value(value)) +#endif +{ +} + +/// Implicit conversions + +inline C10_HOST_DEVICE Half::operator float() const { +#if defined(__CUDA_ARCH__) || defined(__HIP_DEVICE_COMPILE__) + return __half2float(*reinterpret_cast(&x)); +#elif defined(__SYCL_DEVICE_ONLY__) + return float(c10::bit_cast(x)); +#elif (defined(CPU_CAPABILITY_AVX2) || defined(CPU_CAPABILITY_AVX512)) && \ + !defined(__APPLE__) + return at::vec::half2float_scalar(x); +#elif defined(__aarch64__) && !defined(__CUDACC__) + return detail::native_fp16_to_fp32_value(x); +#else + return detail::fp16_ieee_to_fp32_value(x); +#endif +} + +#endif /* !defined(__aarch64__) || defined(__CUDACC__) \ + */ + +#if defined(__CUDACC__) || defined(__HIPCC__) +inline C10_HOST_DEVICE Half::Half(const __half& value) { + x = *reinterpret_cast(&value); +} +inline C10_HOST_DEVICE Half::operator __half() const { + return *reinterpret_cast(&x); +} +#endif + +#ifdef SYCL_LANGUAGE_VERSION +inline C10_HOST_DEVICE Half::Half(const sycl::half& value) { + x = *reinterpret_cast(&value); +} +inline C10_HOST_DEVICE Half::operator sycl::half() const { + return *reinterpret_cast(&x); +} +#endif + +// CUDA intrinsics + +#if (defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 350)) || \ + (defined(__clang__) && defined(__CUDA__)) +inline __device__ Half __ldg(const Half* ptr) { + return __ldg(reinterpret_cast(ptr)); +} +#endif + +/// Arithmetic + +inline C10_HOST_DEVICE Half operator+(const Half& a, const Half& b) { + return static_cast(a) + static_cast(b); +} + +inline C10_HOST_DEVICE Half operator-(const Half& a, const Half& b) { + return static_cast(a) - static_cast(b); +} + +inline C10_HOST_DEVICE Half operator*(const Half& a, const Half& b) { + return static_cast(a) * static_cast(b); +} + +inline C10_HOST_DEVICE Half operator/(const Half& a, const Half& b) + __ubsan_ignore_float_divide_by_zero__ { + return static_cast(a) / static_cast(b); +} + +inline C10_HOST_DEVICE Half operator-(const Half& a) { +#if (defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 530) || \ + defined(__HIP_DEVICE_COMPILE__) + return __hneg(a); +#elif defined(__SYCL_DEVICE_ONLY__) + return -c10::bit_cast(a); +#else + return -static_cast(a); +#endif +} + +inline C10_HOST_DEVICE Half& operator+=(Half& a, const Half& b) { + a = a + b; + return a; +} + +inline C10_HOST_DEVICE Half& operator-=(Half& a, const Half& b) { + a = a - b; + return a; +} + +inline C10_HOST_DEVICE Half& operator*=(Half& a, const Half& b) { + a = a * b; + return a; +} + +inline C10_HOST_DEVICE Half& operator/=(Half& a, const Half& b) { + a = a / b; + return a; +} + +/// Arithmetic with floats + +inline C10_HOST_DEVICE float operator+(Half a, float b) { + return static_cast(a) + b; +} +inline C10_HOST_DEVICE float operator-(Half a, float b) { + return static_cast(a) - b; +} +inline C10_HOST_DEVICE float operator*(Half a, float b) { + return static_cast(a) * b; +} +inline C10_HOST_DEVICE float operator/(Half a, float b) + __ubsan_ignore_float_divide_by_zero__ { + return static_cast(a) / b; +} + +inline C10_HOST_DEVICE float operator+(float a, Half b) { + return a + static_cast(b); +} +inline C10_HOST_DEVICE float operator-(float a, Half b) { + return a - static_cast(b); +} +inline C10_HOST_DEVICE float operator*(float a, Half b) { + return a * static_cast(b); +} +inline C10_HOST_DEVICE float operator/(float a, Half b) + __ubsan_ignore_float_divide_by_zero__ { + return a / static_cast(b); +} + +inline C10_HOST_DEVICE float& operator+=(float& a, const Half& b) { + return a += static_cast(b); +} +inline C10_HOST_DEVICE float& operator-=(float& a, const Half& b) { + return a -= static_cast(b); +} +inline C10_HOST_DEVICE float& operator*=(float& a, const Half& b) { + return a *= static_cast(b); +} +inline C10_HOST_DEVICE float& operator/=(float& a, const Half& b) { + return a /= static_cast(b); +} + +/// Arithmetic with doubles + +inline C10_HOST_DEVICE double operator+(Half a, double b) { + return static_cast(a) + b; +} +inline C10_HOST_DEVICE double operator-(Half a, double b) { + return static_cast(a) - b; +} +inline C10_HOST_DEVICE double operator*(Half a, double b) { + return static_cast(a) * b; +} +inline C10_HOST_DEVICE double operator/(Half a, double b) + __ubsan_ignore_float_divide_by_zero__ { + return static_cast(a) / b; +} + +inline C10_HOST_DEVICE double operator+(double a, Half b) { + return a + static_cast(b); +} +inline C10_HOST_DEVICE double operator-(double a, Half b) { + return a - static_cast(b); +} +inline C10_HOST_DEVICE double operator*(double a, Half b) { + return a * static_cast(b); +} +inline C10_HOST_DEVICE double operator/(double a, Half b) + __ubsan_ignore_float_divide_by_zero__ { + return a / static_cast(b); +} + +/// Arithmetic with ints + +inline C10_HOST_DEVICE Half operator+(Half a, int b) { + // NOLINTNEXTLINE(cppcoreguidelines-narrowing-conversions,bugprone-narrowing-conversions) + return a + static_cast(b); +} +inline C10_HOST_DEVICE Half operator-(Half a, int b) { + // NOLINTNEXTLINE(cppcoreguidelines-narrowing-conversions,bugprone-narrowing-conversions) + return a - static_cast(b); +} +inline C10_HOST_DEVICE Half operator*(Half a, int b) { + // NOLINTNEXTLINE(cppcoreguidelines-narrowing-conversions,bugprone-narrowing-conversions) + return a * static_cast(b); +} +inline C10_HOST_DEVICE Half operator/(Half a, int b) { + // NOLINTNEXTLINE(cppcoreguidelines-narrowing-conversions,bugprone-narrowing-conversions) + return a / static_cast(b); +} + +inline C10_HOST_DEVICE Half operator+(int a, Half b) { + // NOLINTNEXTLINE(cppcoreguidelines-narrowing-conversions,bugprone-narrowing-conversions) + return static_cast(a) + b; +} +inline C10_HOST_DEVICE Half operator-(int a, Half b) { + // NOLINTNEXTLINE(cppcoreguidelines-narrowing-conversions,bugprone-narrowing-conversions) + return static_cast(a) - b; +} +inline C10_HOST_DEVICE Half operator*(int a, Half b) { + // NOLINTNEXTLINE(cppcoreguidelines-narrowing-conversions,bugprone-narrowing-conversions) + return static_cast(a) * b; +} +inline C10_HOST_DEVICE Half operator/(int a, Half b) { + // NOLINTNEXTLINE(cppcoreguidelines-narrowing-conversions,bugprone-narrowing-conversions) + return static_cast(a) / b; +} + +//// Arithmetic with int64_t + +inline C10_HOST_DEVICE Half operator+(Half a, int64_t b) { + // NOLINTNEXTLINE(cppcoreguidelines-narrowing-conversions,bugprone-narrowing-conversions) + return a + static_cast(b); +} +inline C10_HOST_DEVICE Half operator-(Half a, int64_t b) { + // NOLINTNEXTLINE(cppcoreguidelines-narrowing-conversions,bugprone-narrowing-conversions) + return a - static_cast(b); +} +inline C10_HOST_DEVICE Half operator*(Half a, int64_t b) { + // NOLINTNEXTLINE(cppcoreguidelines-narrowing-conversions,bugprone-narrowing-conversions) + return a * static_cast(b); +} +inline C10_HOST_DEVICE Half operator/(Half a, int64_t b) { + // NOLINTNEXTLINE(cppcoreguidelines-narrowing-conversions,bugprone-narrowing-conversions) + return a / static_cast(b); +} + +inline C10_HOST_DEVICE Half operator+(int64_t a, Half b) { + // NOLINTNEXTLINE(cppcoreguidelines-narrowing-conversions,bugprone-narrowing-conversions) + return static_cast(a) + b; +} +inline C10_HOST_DEVICE Half operator-(int64_t a, Half b) { + // NOLINTNEXTLINE(cppcoreguidelines-narrowing-conversions,bugprone-narrowing-conversions) + return static_cast(a) - b; +} +inline C10_HOST_DEVICE Half operator*(int64_t a, Half b) { + // NOLINTNEXTLINE(cppcoreguidelines-narrowing-conversions,bugprone-narrowing-conversions) + return static_cast(a) * b; +} +inline C10_HOST_DEVICE Half operator/(int64_t a, Half b) { + // NOLINTNEXTLINE(cppcoreguidelines-narrowing-conversions,bugprone-narrowing-conversions) + return static_cast(a) / b; +} + +/// NOTE: we do not define comparisons directly and instead rely on the implicit +/// conversion from c10::Half to float. + +C10_CLANG_DIAGNOSTIC_POP() + +} // namespace c10 + +namespace torch::headeronly { + +using c10::Half; +using c10::operator+; +using c10::operator-; +using c10::operator*; +using c10::operator/; +using c10::operator+=; +using c10::operator-=; +using c10::operator*=; +using c10::operator/=; +using c10::operator<<; + +namespace detail { +#if defined(__aarch64__) && !defined(__CUDACC__) +using c10::detail::fp16_from_bits; +using c10::detail::fp16_to_bits; +using c10::detail::native_fp16_from_fp32_value; +using c10::detail::native_fp16_to_fp32_value; +#endif + +using c10::detail::fp16_ieee_from_fp32_value; +using c10::detail::fp16_ieee_to_fp32_bits; +using c10::detail::fp16_ieee_to_fp32_value; +} // namespace detail + +} // namespace torch::headeronly + +namespace std { + +template <> +class numeric_limits { + public: + static constexpr bool is_specialized = true; + static constexpr bool is_signed = true; + static constexpr bool is_integer = false; + static constexpr bool is_exact = false; + static constexpr bool has_infinity = true; + static constexpr bool has_quiet_NaN = true; + static constexpr bool has_signaling_NaN = true; + static constexpr auto has_denorm = numeric_limits::has_denorm; + static constexpr auto has_denorm_loss = + numeric_limits::has_denorm_loss; + static constexpr auto round_style = numeric_limits::round_style; + static constexpr bool is_iec559 = true; + static constexpr bool is_bounded = true; + static constexpr bool is_modulo = false; + static constexpr int digits = 11; + static constexpr int digits10 = 3; + static constexpr int max_digits10 = 5; + static constexpr int radix = 2; + static constexpr int min_exponent = -13; + static constexpr int min_exponent10 = -4; + static constexpr int max_exponent = 16; + static constexpr int max_exponent10 = 4; + static constexpr auto traps = numeric_limits::traps; + static constexpr auto tinyness_before = + numeric_limits::tinyness_before; + static constexpr c10::Half min() { + return c10::Half(0x0400, c10::Half::from_bits()); + } + static constexpr c10::Half lowest() { + return c10::Half(0xFBFF, c10::Half::from_bits()); + } + static constexpr c10::Half max() { + return c10::Half(0x7BFF, c10::Half::from_bits()); + } + static constexpr c10::Half epsilon() { + return c10::Half(0x1400, c10::Half::from_bits()); + } + static constexpr c10::Half round_error() { + return c10::Half(0x3800, c10::Half::from_bits()); + } + static constexpr c10::Half infinity() { + return c10::Half(0x7C00, c10::Half::from_bits()); + } + static constexpr c10::Half quiet_NaN() { + return c10::Half(0x7E00, c10::Half::from_bits()); + } + static constexpr c10::Half signaling_NaN() { + return c10::Half(0x7D00, c10::Half::from_bits()); + } + static constexpr c10::Half denorm_min() { + return c10::Half(0x0001, c10::Half::from_bits()); + } +}; + +} // namespace std diff --git a/runtime/core/portable_type/c10/torch/headeronly/util/TypeSafeSignMath.h b/runtime/core/portable_type/c10/torch/headeronly/util/TypeSafeSignMath.h new file mode 100644 index 00000000000..561ea0467a0 --- /dev/null +++ b/runtime/core/portable_type/c10/torch/headeronly/util/TypeSafeSignMath.h @@ -0,0 +1,148 @@ +#pragma once + +#include +#include +#include + +C10_CLANG_DIAGNOSTIC_PUSH() +#if C10_CLANG_HAS_WARNING("-Wstring-conversion") +C10_CLANG_DIAGNOSTIC_IGNORE("-Wstring-conversion") +#endif +#if C10_CLANG_HAS_WARNING("-Wimplicit-int-float-conversion") +C10_CLANG_DIAGNOSTIC_IGNORE("-Wimplicit-int-float-conversion") +#endif + +namespace c10 { + +/// Returns false since we cannot have x < 0 if x is unsigned. +template +inline constexpr bool is_negative( + const T& /*x*/, + std::true_type /*is_unsigned*/) { + return false; +} + +/// Returns true if a signed variable x < 0 +template +inline constexpr bool is_negative(const T& x, std::false_type /*is_unsigned*/) { + return x < T(0); +} + +/// Returns true if x < 0 +/// NOTE: Will fail on an unsigned custom type +/// For the most part it's possible to fix this if +/// the custom type has a constexpr constructor. +/// However, notably, c10::Half does not :-( +template +inline constexpr bool is_negative(const T& x) { + return is_negative(x, std::is_unsigned()); +} + +/// Returns the sign of an unsigned variable x as 0, 1 +template +inline constexpr int signum(const T& x, std::true_type /*is_unsigned*/) { + return T(0) < x; +} + +/// Returns the sign of a signed variable x as -1, 0, 1 +template +inline constexpr int signum(const T& x, std::false_type /*is_unsigned*/) { + return (T(0) < x) - (x < T(0)); +} + +/// Returns the sign of x as -1, 0, 1 +/// NOTE: Will fail on an unsigned custom type +/// For the most part it's possible to fix this if +/// the custom type has a constexpr constructor. +/// However, notably, c10::Half does not :-( +template +inline constexpr int signum(const T& x) { + return signum(x, std::is_unsigned()); +} + +/// Returns true if a and b are not both negative +template +inline constexpr bool signs_differ(const T& a, const U& b) { + return is_negative(a) != is_negative(b); +} + +// Suppress sign compare warning when compiling with GCC +// as later does not account for short-circuit rule before +// raising the warning, see https://godbolt.org/z/Tr3Msnz99 +#ifdef __GNUC__ +#pragma GCC diagnostic push +#pragma GCC diagnostic ignored "-Wsign-compare" +#endif + +/// Returns true if x is greater than the greatest value of the type Limit +template +inline constexpr bool greater_than_max(const T& x) { + constexpr bool can_overflow = + std::numeric_limits::digits > std::numeric_limits::digits; + return can_overflow && x > (std::numeric_limits::max)(); +} + +#ifdef __GNUC__ +#pragma GCC diagnostic pop +#endif + +/// Returns true if x < lowest(Limit). Standard comparison +template +inline constexpr bool less_than_lowest( + const T& x, + std::false_type /*limit_is_unsigned*/, + std::false_type /*x_is_unsigned*/) { + return x < std::numeric_limits::lowest(); +} + +/// Returns false since all the limit is signed and therefore includes +/// negative values but x cannot be negative because it is unsigned +template +inline constexpr bool less_than_lowest( + const T& /*x*/, + std::false_type /*limit_is_unsigned*/, + std::true_type /*x_is_unsigned*/) { + return false; +} + +/// Returns true if x < 0, where 0 is constructed from T. +/// Limit is not signed, so its lower value is zero +template +inline constexpr bool less_than_lowest( + const T& x, + std::true_type /*limit_is_unsigned*/, + std::false_type /*x_is_unsigned*/) { + return x < T(0); +} + +/// Returns false sign both types are unsigned +template +inline constexpr bool less_than_lowest( + const T& /*x*/, + std::true_type /*limit_is_unsigned*/, + std::true_type /*x_is_unsigned*/) { + return false; +} + +/// Returns true if x is less than the lowest value of type T +/// NOTE: Will fail on an unsigned custom type +/// For the most part it's possible to fix this if +/// the custom type has a constexpr constructor. +/// However, notably, c10::Half does not : +template +inline constexpr bool less_than_lowest(const T& x) { + return less_than_lowest( + x, std::is_unsigned(), std::is_unsigned()); +} + +} // namespace c10 + +C10_CLANG_DIAGNOSTIC_POP() + +namespace torch::headeronly { +using c10::greater_than_max; +using c10::is_negative; +using c10::less_than_lowest; +using c10::signs_differ; +using c10::signum; +} // namespace torch::headeronly diff --git a/runtime/core/portable_type/c10/torch/headeronly/util/bit_cast.h b/runtime/core/portable_type/c10/torch/headeronly/util/bit_cast.h new file mode 100644 index 00000000000..334ba5b8e5b --- /dev/null +++ b/runtime/core/portable_type/c10/torch/headeronly/util/bit_cast.h @@ -0,0 +1,50 @@ +#pragma once + +#include +#include + +#include + +#if __has_include() && (defined(__cpp_lib_bit_cast) && __cpp_lib_bit_cast >= 201806L) +#include +#define C10_HAVE_STD_BIT_CAST 1 +#else +#define C10_HAVE_STD_BIT_CAST 0 +#endif // __has_include() && (__cplusplus >= 202002L || + // (defined(__cpp_lib_bit_cast) && __cpp_lib_bit_cast >= 201806L)) + +namespace torch::headeronly { + +#if C10_HAVE_STD_BIT_CAST +using std::bit_cast; +#else +// Implementations of std::bit_cast() from C++ 20. +// +// This is a less sketchy version of reinterpret_cast. +// +// See https://en.cppreference.com/w/cpp/numeric/bit_cast for more +// information as well as the source of our implementations. +template +C10_HOST_DEVICE std::enable_if_t< + sizeof(To) == sizeof(From) && std::is_trivially_copyable_v && + std::is_trivially_copyable_v, + To> +// constexpr support needs compiler magic +bit_cast(const From& src) noexcept { + static_assert( + std::is_trivially_constructible_v, + "This implementation additionally requires " + "destination type to be trivially constructible"); + + To dst; + std::memcpy(&dst, &src, sizeof(To)); + return dst; +} +#endif // C10_HAVE_STD_BIT_CAST +#undef C10_HAVE_STD_BIT_CAST + +} // namespace torch::headeronly + +namespace c10 { +using torch::headeronly::bit_cast; +} // namespace c10 diff --git a/runtime/core/portable_type/c10/torch/headeronly/util/complex.h b/runtime/core/portable_type/c10/torch/headeronly/util/complex.h new file mode 100644 index 00000000000..e0a356436ac --- /dev/null +++ b/runtime/core/portable_type/c10/torch/headeronly/util/complex.h @@ -0,0 +1,616 @@ +#pragma once + +#include + +#include +#include + +#if defined(__CUDACC__) || defined(__HIPCC__) +#include +#endif + +C10_CLANG_DIAGNOSTIC_PUSH() +#if C10_CLANG_HAS_WARNING("-Wimplicit-float-conversion") +C10_CLANG_DIAGNOSTIC_IGNORE("-Wimplicit-float-conversion") +#endif +#if C10_CLANG_HAS_WARNING("-Wfloat-conversion") +C10_CLANG_DIAGNOSTIC_IGNORE("-Wfloat-conversion") +#endif + +namespace c10 { + +// c10::complex is an implementation of complex numbers that aims +// to work on all devices supported by PyTorch +// +// Most of the APIs duplicates std::complex +// Reference: https://en.cppreference.com/w/cpp/numeric/complex +// +// [NOTE: Complex Operator Unification] +// Operators currently use a mix of std::complex, thrust::complex, and +// c10::complex internally. The end state is that all operators will use +// c10::complex internally. Until then, there may be some hacks to support all +// variants. +// +// +// [Note on Constructors] +// +// The APIs of constructors are mostly copied from C++ standard: +// https://en.cppreference.com/w/cpp/numeric/complex/complex +// +// Since C++14, all constructors are constexpr in std::complex +// +// There are three types of constructors: +// - initializing from real and imag: +// `constexpr complex( const T& re = T(), const T& im = T() );` +// - implicitly-declared copy constructor +// - converting constructors +// +// Converting constructors: +// - std::complex defines converting constructor between float/double/long +// double, +// while we define converting constructor between float/double. +// - For these converting constructors, upcasting is implicit, downcasting is +// explicit. +// - We also define explicit casting from std::complex/thrust::complex +// - Note that the conversion from thrust is not constexpr, because +// thrust does not define them as constexpr ???? +// +// +// [Operator =] +// +// The APIs of operator = are mostly copied from C++ standard: +// https://en.cppreference.com/w/cpp/numeric/complex/operator%3D +// +// Since C++20, all operator= are constexpr. Although we are not building with +// C++20, we also obey this behavior. +// +// There are three types of assign operator: +// - Assign a real value from the same scalar type +// - In std, this is templated as complex& operator=(const T& x) +// with specialization `complex& operator=(T x)` for float/double/long +// double Since we only support float and double, on will use `complex& +// operator=(T x)` +// - Copy assignment operator and converting assignment operator +// - There is no specialization of converting assignment operators, which type +// is +// convertible is solely dependent on whether the scalar type is convertible +// +// In addition to the standard assignment, we also provide assignment operators +// with std and thrust +// +// +// [Casting operators] +// +// std::complex does not have casting operators. We define casting operators +// casting to std::complex and thrust::complex +// +// +// [Operator ""] +// +// std::complex has custom literals `i`, `if` and `il` defined in namespace +// `std::literals::complex_literals`. We define our own custom literals in the +// namespace `c10::complex_literals`. Our custom literals does not follow the +// same behavior as in std::complex, instead, we define _if, _id to construct +// float/double complex literals. +// +// +// [real() and imag()] +// +// In C++20, there are two overload of these functions, one it to return the +// real/imag, another is to set real/imag, they are both constexpr. We follow +// this design. +// +// +// [Operator +=,-=,*=,/=] +// +// Since C++20, these operators become constexpr. In our implementation, they +// are also constexpr. +// +// There are two types of such operators: operating with a real number, or +// operating with another complex number. For the operating with a real number, +// the generic template form has argument type `const T &`, while the overload +// for float/double/long double has `T`. We will follow the same type as +// float/double/long double in std. +// +// [Unary operator +-] +// +// Since C++20, they are constexpr. We also make them expr +// +// [Binary operators +-*/] +// +// Each operator has three versions (taking + as example): +// - complex + complex +// - complex + real +// - real + complex +// +// [Operator ==, !=] +// +// Each operator has three versions (taking == as example): +// - complex == complex +// - complex == real +// - real == complex +// +// Some of them are removed on C++20, but we decide to keep them +// +// [Operator <<, >>] +// +// These are implemented by casting to std::complex +// +// +// +// TODO(@zasdfgbnm): c10::complex is not currently supported, +// because: +// - lots of members and functions of c10::Half are not constexpr +// - thrust::complex only support float and double + +template +struct alignas(sizeof(T) * 2) complex { + using value_type = T; + + T real_ = T(0); + T imag_ = T(0); + + constexpr complex() = default; + C10_HOST_DEVICE constexpr complex(const T& re, const T& im = T()) + : real_(re), imag_(im) {} + template + explicit constexpr complex(const std::complex& other) + : complex(other.real(), other.imag()) {} +#if defined(__CUDACC__) || defined(__HIPCC__) + template + explicit C10_HOST_DEVICE complex(const thrust::complex& other) + : real_(other.real()), imag_(other.imag()) {} +// NOTE can not be implemented as follow due to ROCm bug: +// explicit C10_HOST_DEVICE complex(const thrust::complex &other): +// complex(other.real(), other.imag()) {} +#endif + + // Use SFINAE to specialize casting constructor for c10::complex and + // c10::complex + template + C10_HOST_DEVICE explicit constexpr complex( + const std::enable_if_t, complex>& other) + : real_(other.real_), imag_(other.imag_) {} + template + C10_HOST_DEVICE constexpr complex( + const std::enable_if_t, complex>& other) + : real_(other.real_), imag_(other.imag_) {} + + constexpr complex& operator=(T re) { + real_ = re; + imag_ = 0; + return *this; + } + + constexpr complex& operator+=(T re) { + real_ += re; + return *this; + } + + constexpr complex& operator-=(T re) { + real_ -= re; + return *this; + } + + constexpr complex& operator*=(T re) { + real_ *= re; + imag_ *= re; + return *this; + } + + constexpr complex& operator/=(T re) { + real_ /= re; + imag_ /= re; + return *this; + } + + template + constexpr complex& operator=(const complex& rhs) { + real_ = rhs.real(); + imag_ = rhs.imag(); + return *this; + } + + template + constexpr complex& operator+=(const complex& rhs) { + real_ += rhs.real(); + imag_ += rhs.imag(); + return *this; + } + + template + constexpr complex& operator-=(const complex& rhs) { + real_ -= rhs.real(); + imag_ -= rhs.imag(); + return *this; + } + + template + constexpr complex& operator*=(const complex& rhs) { + // (a + bi) * (c + di) = (a*c - b*d) + (a * d + b * c) i + T a = real_; + T b = imag_; + U c = rhs.real(); + U d = rhs.imag(); + real_ = a * c - b * d; + imag_ = a * d + b * c; + return *this; + } + +#ifdef __APPLE__ +#define FORCE_INLINE_APPLE __attribute__((always_inline)) +#else +#define FORCE_INLINE_APPLE +#endif + template + constexpr FORCE_INLINE_APPLE complex& operator/=(const complex& rhs) + __ubsan_ignore_float_divide_by_zero__ { + // (a + bi) / (c + di) = (ac + bd)/(c^2 + d^2) + (bc - ad)/(c^2 + d^2) i + // the calculation below follows numpy's complex division + T a = real_; + T b = imag_; + U c = rhs.real(); + U d = rhs.imag(); + +#if defined(__GNUC__) && !defined(__clang__) + // std::abs is already constexpr by gcc + auto abs_c = std::abs(c); + auto abs_d = std::abs(d); +#else + auto abs_c = c < 0 ? -c : c; + auto abs_d = d < 0 ? -d : d; +#endif + + if (abs_c >= abs_d) { + if (abs_c == U(0) && abs_d == U(0)) { + /* divide by zeros should yield a complex inf or nan */ + real_ = a / abs_c; + imag_ = b / abs_d; + } else { + auto rat = d / c; + auto scl = U(1.0) / (c + d * rat); + real_ = (a + b * rat) * scl; + imag_ = (b - a * rat) * scl; + } + } else { + auto rat = c / d; + auto scl = U(1.0) / (d + c * rat); + real_ = (a * rat + b) * scl; + imag_ = (b * rat - a) * scl; + } + return *this; + } +#undef FORCE_INLINE_APPLE + + template + constexpr complex& operator=(const std::complex& rhs) { + real_ = rhs.real(); + imag_ = rhs.imag(); + return *this; + } + +#if defined(__CUDACC__) || defined(__HIPCC__) + template + C10_HOST_DEVICE complex& operator=(const thrust::complex& rhs) { + real_ = rhs.real(); + imag_ = rhs.imag(); + return *this; + } +#endif + + template + explicit constexpr operator std::complex() const { + return std::complex(std::complex(real(), imag())); + } + +#if defined(__CUDACC__) || defined(__HIPCC__) + template + C10_HOST_DEVICE explicit operator thrust::complex() const { + return static_cast>(thrust::complex(real(), imag())); + } +#endif + + // consistent with NumPy behavior + explicit constexpr operator bool() const { + return real() || imag(); + } + + C10_HOST_DEVICE constexpr T real() const { + return real_; + } + constexpr void real(T value) { + real_ = value; + } + C10_HOST_DEVICE constexpr T imag() const { + return imag_; + } + constexpr void imag(T value) { + imag_ = value; + } +}; + +namespace complex_literals { + +constexpr complex operator""_if(long double imag) { + return complex(0.0f, static_cast(imag)); +} + +constexpr complex operator""_id(long double imag) { + return complex(0.0, static_cast(imag)); +} + +constexpr complex operator""_if(unsigned long long imag) { + return complex(0.0f, static_cast(imag)); +} + +constexpr complex operator""_id(unsigned long long imag) { + return complex(0.0, static_cast(imag)); +} + +} // namespace complex_literals + +template +constexpr complex operator+(const complex& val) { + return val; +} + +template +constexpr complex operator-(const complex& val) { + return complex(-val.real(), -val.imag()); +} + +template +constexpr complex operator+(const complex& lhs, const complex& rhs) { + complex result = lhs; + return result += rhs; +} + +template +constexpr complex operator+(const complex& lhs, const T& rhs) { + complex result = lhs; + return result += rhs; +} + +template +constexpr complex operator+(const T& lhs, const complex& rhs) { + return complex(lhs + rhs.real(), rhs.imag()); +} + +template +constexpr complex operator-(const complex& lhs, const complex& rhs) { + complex result = lhs; + return result -= rhs; +} + +template +constexpr complex operator-(const complex& lhs, const T& rhs) { + complex result = lhs; + return result -= rhs; +} + +template +constexpr complex operator-(const T& lhs, const complex& rhs) { + complex result = -rhs; + return result += lhs; +} + +template +constexpr complex operator*(const complex& lhs, const complex& rhs) { + complex result = lhs; + return result *= rhs; +} + +template +constexpr complex operator*(const complex& lhs, const T& rhs) { + complex result = lhs; + return result *= rhs; +} + +template +constexpr complex operator*(const T& lhs, const complex& rhs) { + complex result = rhs; + return result *= lhs; +} + +template +constexpr complex operator/(const complex& lhs, const complex& rhs) { + complex result = lhs; + return result /= rhs; +} + +template +constexpr complex operator/(const complex& lhs, const T& rhs) { + complex result = lhs; + return result /= rhs; +} + +template +constexpr complex operator/(const T& lhs, const complex& rhs) { + complex result(lhs, T()); + return result /= rhs; +} + +// Define operators between integral scalars and c10::complex. std::complex does +// not support this when T is a floating-point number. This is useful because it +// saves a lot of "static_cast" when operate a complex and an integer. This +// makes the code both less verbose and potentially more efficient. +#define COMPLEX_INTEGER_OP_TEMPLATE_CONDITION \ + typename std::enable_if_t< \ + std::is_floating_point_v && std::is_integral_v, \ + int> = 0 + +template +constexpr c10::complex operator+(const c10::complex& a, const iT& b) { + return a + static_cast(b); +} + +template +constexpr c10::complex operator+(const iT& a, const c10::complex& b) { + return static_cast(a) + b; +} + +template +constexpr c10::complex operator-(const c10::complex& a, const iT& b) { + return a - static_cast(b); +} + +template +constexpr c10::complex operator-(const iT& a, const c10::complex& b) { + return static_cast(a) - b; +} + +template +constexpr c10::complex operator*(const c10::complex& a, const iT& b) { + return a * static_cast(b); +} + +template +constexpr c10::complex operator*(const iT& a, const c10::complex& b) { + return static_cast(a) * b; +} + +template +constexpr c10::complex operator/(const c10::complex& a, const iT& b) { + return a / static_cast(b); +} + +template +constexpr c10::complex operator/(const iT& a, const c10::complex& b) { + return static_cast(a) / b; +} + +#undef COMPLEX_INTEGER_OP_TEMPLATE_CONDITION + +template +constexpr bool operator==(const complex& lhs, const complex& rhs) { + return (lhs.real() == rhs.real()) && (lhs.imag() == rhs.imag()); +} + +template +constexpr bool operator==(const complex& lhs, const T& rhs) { + return (lhs.real() == rhs) && (lhs.imag() == T()); +} + +template +constexpr bool operator==(const T& lhs, const complex& rhs) { + return (lhs == rhs.real()) && (T() == rhs.imag()); +} + +template +constexpr bool operator!=(const complex& lhs, const complex& rhs) { + return !(lhs == rhs); +} + +template +constexpr bool operator!=(const complex& lhs, const T& rhs) { + return !(lhs == rhs); +} + +template +constexpr bool operator!=(const T& lhs, const complex& rhs) { + return !(lhs == rhs); +} + +template +std::basic_ostream& operator<<( + std::basic_ostream& os, + const complex& x) { + return (os << static_cast>(x)); +} + +template +std::basic_istream& operator>>( + std::basic_istream& is, + complex& x) { + std::complex tmp; + is >> tmp; + x = tmp; + return is; +} + +template +C10_HOST_DEVICE complex polar(const T& r, const T& theta = T()) { +#if defined(__CUDACC__) || defined(__HIPCC__) + return static_cast>(thrust::polar(r, theta)); +#else + // std::polar() requires r >= 0, so spell out the explicit implementation to + // avoid a branch. + return complex(r * std::cos(theta), r * std::sin(theta)); +#endif +} + +template <> +struct alignas(4) complex { + Half real_; + Half imag_; + + // Constructors + complex() = default; + // Half constructor is not constexpr so the following constructor can't + // be constexpr + C10_HOST_DEVICE explicit inline complex(const Half& real, const Half& imag) + : real_(real), imag_(imag) {} + C10_HOST_DEVICE inline complex(const c10::complex& value) + : real_(value.real()), imag_(value.imag()) {} + + // Conversion operator + inline C10_HOST_DEVICE operator c10::complex() const { + return {real_, imag_}; + } + + constexpr C10_HOST_DEVICE Half real() const { + return real_; + } + constexpr C10_HOST_DEVICE Half imag() const { + return imag_; + } + + C10_HOST_DEVICE complex& operator+=(const complex& other) { + real_ = static_cast(real_) + static_cast(other.real_); + imag_ = static_cast(imag_) + static_cast(other.imag_); + return *this; + } + + C10_HOST_DEVICE complex& operator-=(const complex& other) { + real_ = static_cast(real_) - static_cast(other.real_); + imag_ = static_cast(imag_) - static_cast(other.imag_); + return *this; + } + + C10_HOST_DEVICE complex& operator*=(const complex& other) { + auto a = static_cast(real_); + auto b = static_cast(imag_); + auto c = static_cast(other.real()); + auto d = static_cast(other.imag()); + real_ = a * c - b * d; + imag_ = a * d + b * c; + return *this; + } +}; + +} // namespace c10 + +namespace torch::headeronly { +using c10::complex; +using c10::operator+; +using c10::operator-; +using c10::operator*; +using c10::operator/; +using c10::operator+=; +using c10::operator-=; +using c10::operator*=; +using c10::operator/=; +using c10::operator==; +using c10::operator!=; +using c10::operator<<; +using c10::operator>>; +using c10::polar; + +namespace complex_literals { +using c10::complex_literals::operator""_if; +using c10::complex_literals::operator""_id; +} // namespace complex_literals + +} // namespace torch::headeronly + +C10_CLANG_DIAGNOSTIC_POP() diff --git a/runtime/core/portable_type/c10/torch/headeronly/util/floating_point_utils.h b/runtime/core/portable_type/c10/torch/headeronly/util/floating_point_utils.h new file mode 100644 index 00000000000..c469cc6a4f6 --- /dev/null +++ b/runtime/core/portable_type/c10/torch/headeronly/util/floating_point_utils.h @@ -0,0 +1,38 @@ +#pragma once + +#include +#include +#include + +namespace torch::headeronly::detail { + +C10_HOST_DEVICE inline float fp32_from_bits(uint32_t w) { +#if defined(__OPENCL_VERSION__) + return as_float(w); +#elif defined(__CUDA_ARCH__) || defined(__HIP_DEVICE_COMPILE__) + return __uint_as_float((unsigned int)w); +#elif defined(__INTEL_COMPILER) + return _castu32_f32(w); +#else + return torch::headeronly::bit_cast(w); +#endif +} + +C10_HOST_DEVICE inline uint32_t fp32_to_bits(float f) { +#if defined(__OPENCL_VERSION__) + return as_uint(f); +#elif defined(__CUDA_ARCH__) || defined(__HIP_DEVICE_COMPILE__) + return (uint32_t)__float_as_uint(f); +#elif defined(__INTEL_COMPILER) + return _castf32_u32(f); +#else + return torch::headeronly::bit_cast(f); +#endif +} + +} // namespace torch::headeronly::detail + +namespace c10::detail { +using torch::headeronly::detail::fp32_from_bits; +using torch::headeronly::detail::fp32_to_bits; +} // namespace c10::detail