From 063634ba4acf3d1eb9261b6905efe0f9db36b2d2 Mon Sep 17 00:00:00 2001 From: pytorchbot Date: Mon, 31 Aug 2026 09:11:45 +0000 Subject: [PATCH] Bump PyTorch pin to nightly dev20260830 --- .ci/docker/ci_commit_pins/pytorch.txt | 2 +- .../c10/c10/util/BFloat16-math.h | 298 +----------------- .../core/portable_type/c10/c10/util/complex.h | 17 +- .../core/portable_type/c10/c10/util/irange.h | 122 +------ .../c10/c10/util/llvmMathExtras.h | 1 + .../portable_type/c10/c10/util/overflows.h | 17 +- .../c10/c10/util/safe_numerics.h | 8 +- .../c10/torch/headeronly/macros/Macros.h | 9 + .../c10/torch/headeronly/util/Half.h | 2 +- .../torch/headeronly/util/TypeSafeSignMath.h | 90 ++---- .../c10/torch/headeronly/util/complex.h | 62 +++- torch_pin.py | 2 +- 12 files changed, 122 insertions(+), 508 deletions(-) diff --git a/.ci/docker/ci_commit_pins/pytorch.txt b/.ci/docker/ci_commit_pins/pytorch.txt index 401a0594d98..fe3fe3c1590 100644 --- a/.ci/docker/ci_commit_pins/pytorch.txt +++ b/.ci/docker/ci_commit_pins/pytorch.txt @@ -1 +1 @@ -release/2.13 +7b4f739ae84d496739c2a73f8da95e8e829fd5cc diff --git a/runtime/core/portable_type/c10/c10/util/BFloat16-math.h b/runtime/core/portable_type/c10/c10/util/BFloat16-math.h index bce89e8acd9..4410002b9e6 100644 --- a/runtime/core/portable_type/c10/c10/util/BFloat16-math.h +++ b/runtime/core/portable_type/c10/c10/util/BFloat16-math.h @@ -1,299 +1,3 @@ #pragma once -#include -#include - -C10_CLANG_DIAGNOSTIC_PUSH() -#if C10_CLANG_HAS_WARNING("-Wimplicit-float-conversion") -C10_CLANG_DIAGNOSTIC_IGNORE("-Wimplicit-float-conversion") -#endif - -namespace c10 { -template -struct is_reduced_floating_point - : std::integral_constant< - bool, - std::is_same_v || std::is_same_v> {}; - -template -constexpr bool is_reduced_floating_point_v = - is_reduced_floating_point::value; -} // namespace c10 - -namespace std { - -#if !defined(FBCODE_CAFFE2) && !defined(C10_NODEPRECATED) -using c10::is_reduced_floating_point; -using c10::is_reduced_floating_point_v; -#endif // !defined(FBCODE_CAFFE2) && !defined(C10_NODEPRECATED) - -template < - typename T, - typename std::enable_if_t, int> = 0> -inline T acos(T a) { - return std::acos(float(a)); -} -template < - typename T, - typename std::enable_if_t, int> = 0> -inline T asin(T a) { - return std::asin(float(a)); -} -template < - typename T, - typename std::enable_if_t, int> = 0> -inline T atan(T a) { - return std::atan(float(a)); -} -template < - typename T, - typename std::enable_if_t, int> = 0> -inline T atanh(T a) { - return std::atanh(float(a)); -} -template < - typename T, - typename std::enable_if_t, int> = 0> -inline T erf(T a) { - return std::erf(float(a)); -} -template < - typename T, - typename std::enable_if_t, int> = 0> -inline T erfc(T a) { - return std::erfc(float(a)); -} -template < - typename T, - typename std::enable_if_t, int> = 0> -inline T exp(T a) { - return std::exp(float(a)); -} -template < - typename T, - typename std::enable_if_t, int> = 0> -inline T expm1(T a) { - return std::expm1(float(a)); -} -template < - typename T, - typename std::enable_if_t, int> = 0> -inline bool isfinite(T a) { - return std::isfinite(float(a)); -} -template < - typename T, - typename std::enable_if_t, int> = 0> -inline T log(T a) { - return std::log(float(a)); -} -template < - typename T, - typename std::enable_if_t, int> = 0> -inline T log10(T a) { - return std::log10(float(a)); -} -template < - typename T, - typename std::enable_if_t, int> = 0> -inline T log1p(T a) { - return std::log1p(float(a)); -} -template < - typename T, - typename std::enable_if_t, int> = 0> -inline T log2(T a) { - return std::log2(float(a)); -} -template < - typename T, - typename std::enable_if_t, int> = 0> -inline T ceil(T a) { - return std::ceil(float(a)); -} -template < - typename T, - typename std::enable_if_t, int> = 0> -inline T cos(T a) { - return std::cos(float(a)); -} -template < - typename T, - typename std::enable_if_t, int> = 0> -inline T floor(T a) { - return std::floor(float(a)); -} -template < - typename T, - typename std::enable_if_t, int> = 0> -inline T nearbyint(T a) { - return std::nearbyint(float(a)); -} -template < - typename T, - typename std::enable_if_t, int> = 0> -inline T sin(T a) { - return std::sin(float(a)); -} -template < - typename T, - typename std::enable_if_t, int> = 0> -inline T tan(T a) { - return std::tan(float(a)); -} -template < - typename T, - typename std::enable_if_t, int> = 0> -inline T sinh(T a) { - return std::sinh(float(a)); -} -template < - typename T, - typename std::enable_if_t, int> = 0> -inline T cosh(T a) { - return std::cosh(float(a)); -} -template < - typename T, - typename std::enable_if_t, int> = 0> -inline T tanh(T a) { - return std::tanh(float(a)); -} -template < - typename T, - typename std::enable_if_t, int> = 0> -inline T trunc(T a) { - return std::trunc(float(a)); -} -template < - typename T, - typename std::enable_if_t, int> = 0> -inline T lgamma(T a) { - return std::lgamma(float(a)); -} -template < - typename T, - typename std::enable_if_t, int> = 0> -inline T sqrt(T a) { - return std::sqrt(float(a)); -} -template < - typename T, - typename std::enable_if_t, int> = 0> -inline T rsqrt(T a) { - return 1.0f / std::sqrt(float(a)); -} -template < - typename T, - typename std::enable_if_t, int> = 0> -inline T abs(T a) { - return std::abs(float(a)); -} -#if defined(_MSC_VER) && defined(__CUDACC__) -template < - typename T, - typename std::enable_if_t, int> = 0> -inline T pow(T a, double b) { - return std::pow(float(a), float(b)); -} -#else -template < - typename T, - typename std::enable_if_t, int> = 0> -inline T pow(T a, double b) { - return std::pow(float(a), b); -} -#endif -template < - typename T, - typename std::enable_if_t, int> = 0> -inline T pow(T a, T b) { - return std::pow(float(a), float(b)); -} -template < - typename T, - typename std::enable_if_t, int> = 0> -inline T fmod(T a, T b) { - return std::fmod(float(a), float(b)); -} - -/* - The following function is inspired from the implementation in `musl` - Link to License: https://git.musl-libc.org/cgit/musl/tree/COPYRIGHT - ---------------------------------------------------------------------- - Copyright © 2005-2020 Rich Felker, et al. - - Permission is hereby granted, free of charge, to any person obtaining - a copy of this software and associated documentation files (the - "Software"), to deal in the Software without restriction, including - without limitation the rights to use, copy, modify, merge, publish, - distribute, sublicense, and/or sell copies of the Software, and to - permit persons to whom the Software is furnished to do so, subject to - the following conditions: - - The above copyright notice and this permission notice shall be - included in all copies or substantial portions of the Software. - - THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, - EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF - MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. - IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY - CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, - TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE - SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. - ---------------------------------------------------------------------- - */ -template < - typename T, - typename std::enable_if_t, int> = 0> -C10_HOST_DEVICE inline T nextafter(T from, T to) { - // Reference: - // https://git.musl-libc.org/cgit/musl/tree/src/math/nextafter.c - using int_repr_t = uint16_t; - constexpr uint8_t bits = 16; - union { - T f; - int_repr_t i; - } ufrom = {from}, uto = {to}; - - // get a mask to get the sign bit i.e. MSB - int_repr_t sign_mask = int_repr_t{1} << (bits - 1); - - // short-circuit: if either is NaN, return NaN - if (from != from || to != to) { - return from + to; - } - - // short-circuit: if they are exactly the same. - if (ufrom.i == uto.i) { - return from; - } - - // mask the sign-bit to zero i.e. positive - // equivalent to abs(x) - int_repr_t abs_from = ufrom.i & ~sign_mask; - int_repr_t abs_to = uto.i & ~sign_mask; - if (abs_from == 0) { - // if both are zero but with different sign, - // preserve the sign of `to`. - if (abs_to == 0) { - return to; - } - // smallest subnormal with sign of `to`. - ufrom.i = (uto.i & sign_mask) | int_repr_t{1}; - return ufrom.f; - } - - // if abs(from) > abs(to) or sign(from) != sign(to) - if (abs_from > abs_to || ((ufrom.i ^ uto.i) & sign_mask)) { - ufrom.i--; - } else { - ufrom.i++; - } - - return ufrom.f; -} - -} // namespace std - -C10_CLANG_DIAGNOSTIC_POP() +#include diff --git a/runtime/core/portable_type/c10/c10/util/complex.h b/runtime/core/portable_type/c10/c10/util/complex.h index 4e699684bc3..f9849a94ced 100644 --- a/runtime/core/portable_type/c10/c10/util/complex.h +++ b/runtime/core/portable_type/c10/c10/util/complex.h @@ -31,19 +31,11 @@ C10_HOST_DEVICE T abs(const c10::complex& z) { #endif } -#if defined(USE_ROCM) -#define ROCm_Bug(x) -#else -#define ROCm_Bug(x) x -#endif - template C10_HOST_DEVICE T arg(const c10::complex& z) { - return ROCm_Bug(std)::atan2(std::imag(z), std::real(z)); + return std::atan2(std::imag(z), std::real(z)); } -#undef ROCm_Bug - template constexpr T norm(const c10::complex& z) { return z.real() * z.real() + z.imag() * z.imag(); @@ -73,6 +65,9 @@ constexpr c10::complex conj(const c10::complex& z) { #define C10_INTERNAL_INCLUDE_COMPLEX_REMAINING_H // math functions are included in a separate file #include // IWYU pragma: keep -// utilities for complex types -#include // IWYU pragma: keep #undef C10_INTERNAL_INCLUDE_COMPLEX_REMAINING_H + +namespace c10 { +using torch::headeronly::is_complex; +using torch::headeronly::scalar_value_type; +} // namespace c10 diff --git a/runtime/core/portable_type/c10/c10/util/irange.h b/runtime/core/portable_type/c10/c10/util/irange.h index cc52d443ee5..c9dd7e574e1 100644 --- a/runtime/core/portable_type/c10/c10/util/irange.h +++ b/runtime/core/portable_type/c10/c10/util/irange.h @@ -1,123 +1,3 @@ -// Copyright 2004-present Facebook. All Rights Reserved. - #pragma once -#include - -#include -#include -#include -#include - -namespace c10 { - -namespace detail { - -template < - typename I, - bool one_sided = false, - std::enable_if_t, int> = 0> -struct integer_iterator { - using iterator_category = std::input_iterator_tag; - using value_type = I; - using difference_type = std::ptrdiff_t; - using pointer = I*; - using reference = I&; - - explicit constexpr integer_iterator(I val) : value(val) {} - - constexpr I operator*() const { - return value; - } - - constexpr I const* operator->() const { - return &value; - } - - constexpr integer_iterator& operator++() { - ++value; - return *this; - } - - constexpr integer_iterator operator++(int) { - const auto copy = *this; - ++*this; - return copy; - } - - constexpr bool operator==(const integer_iterator& other) const { - if constexpr (one_sided) { - // Range-for loops' end test is `begin != end`, not `begin < - // end`. To handle `c10::irange(n)` where n < 0 (which should be - // empty), we just make `begin != end` fail whenever `end` is - // negative. - return is_negative(other.value) || value == other.value; - } else { - return value == other.value; - } - // Suppress "warning: missing return statement at end of non-void function" - // which Nvidia's Robert Crovella confirms is an NVCC compiler error - // here https://stackoverflow.com/a/64561686/752843 on 2020-10-27 - // `__builtin_unreachable();` would be best here, but it's not - // available with all compilers. So we instead return an arbitrary - // value trusting that this line will, in fact, never be reached. - return false; // Horrible hack - } - - constexpr bool operator!=(const integer_iterator& other) const { - return !(*this == other); - } - - protected: - I value; -}; - -} // namespace detail - -template < - typename I, - bool one_sided = false, - std::enable_if_t, bool> = true> -struct integer_range { - public: - constexpr integer_range(I begin, I end) : begin_(begin), end_(end) {} - using iterator = detail::integer_iterator; - constexpr iterator begin() const { - return begin_; - } - constexpr iterator end() const { - return end_; - } - - private: - iterator begin_; - iterator end_; -}; - -/// Creates an integer range for the half-open interval [begin, end) -/// If end<=begin, then the range is empty. -/// The range has the type of the `end` integer; `begin` integer is -/// cast to this type. -template < - typename Integer1, - typename Integer2, - std::enable_if_t, bool> = true, - std::enable_if_t, bool> = true> -constexpr integer_range irange(Integer1 begin, Integer2 end) { - // If end<=begin then the range is empty; we can achieve this effect by - // choosing the larger of {begin, end} as the loop terminator - return { - static_cast(begin), - std::max(static_cast(begin), end)}; -} - -/// Creates an integer range for the half-open interval [0, end) -/// If end<=begin, then the range is empty -template < - typename Integer, - std::enable_if_t, bool> = true> -constexpr integer_range irange(Integer end) { - return {Integer(), end}; -} - -} // namespace c10 +#include diff --git a/runtime/core/portable_type/c10/c10/util/llvmMathExtras.h b/runtime/core/portable_type/c10/c10/util/llvmMathExtras.h index da297241449..8ae5cde4f02 100644 --- a/runtime/core/portable_type/c10/c10/util/llvmMathExtras.h +++ b/runtime/core/portable_type/c10/c10/util/llvmMathExtras.h @@ -400,6 +400,7 @@ constexpr inline bool isShiftedUInt(uint64_t x) { N + S <= 64, "isShiftedUInt with N + S > 64 is too wide."); // Per the two static_asserts above, S must be strictly less than 64. So // 1 << S is not undefined behavior. + // NOLINTNEXTLINE(bugprone-chained-comparison) return isUInt(x) && (x % (UINT64_C(1) << S) == 0); } diff --git a/runtime/core/portable_type/c10/c10/util/overflows.h b/runtime/core/portable_type/c10/c10/util/overflows.h index 183a2f62a32..348ee5ffa40 100644 --- a/runtime/core/portable_type/c10/c10/util/overflows.h +++ b/runtime/core/portable_type/c10/c10/util/overflows.h @@ -61,13 +61,28 @@ template std::enable_if_t, bool> overflows( From f, bool strict_unsigned [[maybe_unused]] = false) { - using limit = std::numeric_limits::type>; + using ToScalar = typename scalar_value_type::type; + using limit = std::numeric_limits; if (limit::has_infinity && std::isinf(static_cast(f))) { return false; } if (!limit::has_quiet_NaN && (f != f)) { return true; } + if constexpr (std::is_integral_v) { + // limit::max() for wide integer types is NOT exactly representable in + // floating point (e.g. int64 max = 2^63-1 rounds up to 2^63), so `f > + // limit::max()` lets a just-out-of-range value like 2^63 slip through and + // then become INT64_MIN via static_cast. Compare against the + // exactly-representable upper bound max()+1 == 2^digits instead. lowest() + // is 0 or a negated power of two, so it stays exact. (digits-1 keeps the + // shift < 64 for the uint64 case; the *2 recovers 2^digits without a 1<<64 + // overflow.) + constexpr int digits = limit::digits; + constexpr From upper = + static_cast(uint64_t{1} << (digits - 1)) * From{2}; + return f < static_cast(limit::lowest()) || f >= upper; + } return f < limit::lowest() || f > limit::max(); } diff --git a/runtime/core/portable_type/c10/c10/util/safe_numerics.h b/runtime/core/portable_type/c10/c10/util/safe_numerics.h index bfdb968ff96..f695fa28e77 100644 --- a/runtime/core/portable_type/c10/c10/util/safe_numerics.h +++ b/runtime/core/portable_type/c10/c10/util/safe_numerics.h @@ -8,8 +8,8 @@ // GCC has __builtin_mul_overflow from before it supported __has_builtin #ifdef _MSC_VER #define C10_HAS_BUILTIN_OVERFLOW() (0) -#include #include +#include #else #define C10_HAS_BUILTIN_OVERFLOW() (1) #endif @@ -68,9 +68,7 @@ C10_ALWAYS_INLINE bool mul_overflows(T a, T b, T* out) { // This test isn't exact, but avoids doing integer division *out = a * b; constexpr int bits = sizeof(T) * 8; - return ( - (c10::llvm::countLeadingZeros(a) + c10::llvm::countLeadingZeros(b)) < - bits); + return (std::countl_zero(a) + std::countl_zero(b)) < bits; } #endif } @@ -98,7 +96,7 @@ bool safe_multiplies_u64(It first, It last, uint64_t* out) { prod *= x; // log2(0) isn't valid, so need to track it specially is_zero |= (x == 0); - prod_log2 += c10::llvm::Log2_64_Ceil(x); + prod_log2 += std::bit_width(x - 1); } *out = prod; // This test isn't exact, but avoids doing integer division diff --git a/runtime/core/portable_type/c10/torch/headeronly/macros/Macros.h b/runtime/core/portable_type/c10/torch/headeronly/macros/Macros.h index cef99df3f56..08c4e9f1f84 100644 --- a/runtime/core/portable_type/c10/torch/headeronly/macros/Macros.h +++ b/runtime/core/portable_type/c10/torch/headeronly/macros/Macros.h @@ -123,6 +123,15 @@ #define C10_HAS_CPP_ATTRIBUTE(x) (0) #endif +/// Bind a returned reference/pointer's lifetime to a parameter (or *this) so +/// Clang can warn when it would dangle. Expands to nothing on compilers that +/// lack the attribute (e.g. non-clang, older nvcc). +#if C10_HAS_CPP_ATTRIBUTE(clang::lifetimebound) +#define C10_LIFETIMEBOUND [[clang::lifetimebound]] +#else +#define C10_LIFETIMEBOUND +#endif + #ifndef FBCODE_CAFFE2 /// DEPRECATED: Warn if a type or return value is discarded. #define C10_NODISCARD [[nodiscard]] diff --git a/runtime/core/portable_type/c10/torch/headeronly/util/Half.h b/runtime/core/portable_type/c10/torch/headeronly/util/Half.h index e5aa622656c..401472357ec 100644 --- a/runtime/core/portable_type/c10/torch/headeronly/util/Half.h +++ b/runtime/core/portable_type/c10/torch/headeronly/util/Half.h @@ -213,7 +213,7 @@ C10_HOST_DEVICE inline float fp16_ieee_to_fp32_value(uint16_t h) { * 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 + * same mantissa and the half-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 diff --git a/runtime/core/portable_type/c10/torch/headeronly/util/TypeSafeSignMath.h b/runtime/core/portable_type/c10/torch/headeronly/util/TypeSafeSignMath.h index c33a286bc5b..8e897957fee 100644 --- a/runtime/core/portable_type/c10/torch/headeronly/util/TypeSafeSignMath.h +++ b/runtime/core/portable_type/c10/torch/headeronly/util/TypeSafeSignMath.h @@ -14,20 +14,6 @@ C10_CLANG_DIAGNOSTIC_IGNORE("-Wimplicit-int-float-conversion") 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 @@ -35,19 +21,12 @@ inline constexpr bool is_negative(const T& x, std::false_type /*is_unsigned*/) { /// 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)); + if constexpr (std::is_unsigned_v) { + // An unsigned value can never be less than zero. + return false; + } else { + return x < T(0); + } } /// Returns the sign of x as -1, 0, 1 @@ -57,7 +36,11 @@ inline constexpr int signum(const T& x, std::false_type /*is_unsigned*/) { /// However, notably, c10::Half does not :-( template inline constexpr int signum(const T& x) { - return signum(x, std::is_unsigned()); + if constexpr (std::is_unsigned_v) { + return T(0) < x; + } else { + return (T(0) < x) - (x < T(0)); + } } /// Returns true if a and b are not both negative @@ -86,53 +69,22 @@ inline constexpr bool greater_than_max(const T& x) { #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 +/// Returns true if x is less than the lowest value of type Limit /// 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()); + if constexpr (std::is_unsigned_v) { + // x is unsigned, so it can never be below the lowest value of any type. + return false; + } else if constexpr (std::is_unsigned_v) { + // Limit is unsigned, so its lowest value is zero. + return x < T(0); + } else { + return x < std::numeric_limits::lowest(); + } } } // 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 index 733a22d5dbb..265b2506f14 100644 --- a/runtime/core/portable_type/c10/torch/headeronly/util/complex.h +++ b/runtime/core/portable_type/c10/torch/headeronly/util/complex.h @@ -3,6 +3,7 @@ #include #include +#include #include #if defined(__CUDACC__) || defined(__HIPCC__) @@ -160,7 +161,7 @@ struct alignas(sizeof(T) * 2) complex { 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: +// NOTE can not be implemented as follows due to ROCm bug: // explicit C10_HOST_DEVICE complex(const thrust::complex &other): // complex(other.real(), other.imag()) {} #endif @@ -588,6 +589,60 @@ struct alignas(4) complex { } }; +template <> +struct alignas(4) complex { + BFloat16 real_; + BFloat16 imag_; + + // Constructors + complex() = default; + // BFloat16 constructor is not constexpr so the following constructor can't + // be constexpr + C10_HOST_DEVICE explicit inline complex( + const BFloat16& real, + const BFloat16& 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 BFloat16 real() const { + return real_; + } + constexpr C10_HOST_DEVICE BFloat16 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 HIDDEN_NAMESPACE_BEGIN(torch, headeronly) @@ -614,3 +669,8 @@ using c10::complex_literals::operator""_id; HIDDEN_NAMESPACE_END(torch, headeronly) C10_CLANG_DIAGNOSTIC_POP() + +#define C10_INTERNAL_INCLUDE_COMPLEX_REMAINING_H +// utilities for complex types +#include // IWYU pragma: keep +#undef C10_INTERNAL_INCLUDE_COMPLEX_REMAINING_H diff --git a/torch_pin.py b/torch_pin.py index ca593b1ef05..cc1253368f4 100644 --- a/torch_pin.py +++ b/torch_pin.py @@ -1,2 +1,2 @@ TORCH_VERSION = "2.13.0" -# NIGHTLY_VERSION = "dev20260318" Temporarily pinning to stable release candidate. Revert https://github.com/pytorch/executorch/pull/18287 +NIGHTLY_VERSION = "dev20260830"