| 123456789101112131415161718192021222324252627282930313233343536373839404142434445464748495051525354555657585960616263646566676869707172737475767778798081828384858687888990919293949596979899100101102103104105106107108109110111112113114115116117118119120121122123124125126127128129130131132133134135136137138139140141142143144145146147148149150151152153154155156157158159160161162163164165166167168169170171172173174175176177178179180181182183184185186187188189190191192193194195196197198199200201202203204205206207208209210211212213214215216217218219220221222223224225226227228229230231232233234235236237238239240241242243244245246247248249250251252253254255256257258259260261262263264265266267268269270271272273274275276277278279280281282283284285286287288289290291292293294295296297298299300301302303304305306307308309310311312313314315316317318319320321322323324325326327328329330331332333334335336337338339340341342343344345346347348349350351352353354355356357358359360361362363364365366367368369370371372373374375376377378379380381382383384385386387388389390391392393394395396397398399400401402403404405406407408409410411412413414415416417418419420421422423424425426427428429430431432433434435436437438439440441442443444445446447448449450451452453454455456457458459460461462463464465466467468469470471472473474475476477478479480481482483484485486 |
- #if !defined(TORCH_STABLE_ONLY) && !defined(TORCH_TARGET_VERSION)
- // Metal helper functions
- #pragma once
- #include <c10/metal/common.h>
- #include <metal_stdlib>
- namespace c10 {
- namespace metal {
- namespace detail {
- template <typename T>
- struct vectypes {};
- template <>
- struct vectypes<float> {
- using type4 = float4;
- using type3 = float3;
- using type2 = float2;
- };
- template <>
- struct vectypes<half> {
- using type4 = half4;
- using type3 = half3;
- using type2 = half2;
- };
- template <>
- struct vectypes<bfloat> {
- using type4 = bfloat4;
- using type3 = bfloat3;
- using type2 = bfloat2;
- };
- template <>
- struct vectypes<short> {
- using type4 = short4;
- using type3 = short3;
- using type2 = short2;
- };
- template <>
- struct vectypes<int> {
- using type4 = int4;
- using type3 = int3;
- using type2 = int2;
- };
- template <>
- struct vectypes<long> {
- using type4 = short4;
- using type3 = short3;
- using type2 = short2;
- };
- template <typename T>
- struct OpMathType {
- using type = T;
- };
- template <>
- struct OpMathType<half> {
- using type = float;
- };
- template <>
- struct OpMathType<short> {
- using type = int;
- };
- template <>
- struct OpMathType<char> {
- using type = int;
- };
- template <>
- struct OpMathType<uchar> {
- using type = int;
- };
- template <>
- struct OpMathType<bfloat> {
- using type = float;
- };
- // Type promotion structure for higher precision accumulation
- template <typename T>
- struct AccumulationType {
- using type = T;
- };
- // Specialization for half - promote to float for accumulation
- template <>
- struct AccumulationType<half> {
- using type = float;
- };
- // Specialization for bfloat - promote to float for accumulation
- template <>
- struct AccumulationType<bfloat> {
- using type = float;
- };
- } // namespace detail
- template <typename T>
- ::metal::enable_if_t<::metal::is_floating_point_v<T>, T> max(T a, T b) {
- return ::metal::isunordered(a, b) ? NAN : ::metal::max(a, b);
- }
- template <typename T, typename U>
- ::metal::enable_if_t<::metal::is_integral_v<T>&& ::metal::is_integral_v<U>, T>
- max(T a, U b) {
- return ::metal::max(a, static_cast<T>(b));
- }
- template <typename T>
- ::metal::enable_if_t<::metal::is_floating_point_v<T>, T> min(T a, T b) {
- return ::metal::isunordered(a, b) ? NAN : ::metal::min(a, b);
- }
- template <typename T, typename U>
- ::metal::enable_if_t<::metal::is_integral_v<T>&& ::metal::is_integral_v<U>, T>
- min(T a, U b) {
- return ::metal::min(a, static_cast<T>(b));
- }
- template <>
- inline bfloat min(bfloat a, bfloat b) {
- return bfloat(
- ::metal::isunordered(a, b) ? NAN : ::metal::min(float(a), float(b)));
- }
- template <>
- inline bfloat max(bfloat a, bfloat b) {
- return bfloat(
- ::metal::isunordered(a, b) ? NAN : ::metal::max(float(a), float(b)));
- }
- template <typename T>
- using vec2type_t = typename detail::vectypes<T>::type2;
- template <typename T>
- using vec4type_t = typename detail::vectypes<T>::type4;
- template <typename T>
- using opmath_t = typename detail::OpMathType<T>::type;
- template <typename T>
- using accum_t = typename detail::AccumulationType<T>::type;
- // TODO: Move it to type_traits header may be
- template <typename F, typename... Args>
- using result_of = decltype(::metal::declval<F>()(::metal::declval<Args>()...));
- template <typename T>
- constexpr constant bool is_complex_v =
- ::metal::is_same_v<T, float2> || ::metal::is_same_v<T, half2>;
- template <typename T>
- constexpr constant bool is_scalar_floating_point_v =
- ::metal::is_floating_point_v<T> && ::metal::is_scalar_v<T>;
- template <typename T>
- constexpr constant bool is_scalar_integral_v =
- ::metal::is_integral_v<T> && ::metal::is_scalar_v<T>;
- template <typename U, typename V>
- using common_dtype = decltype(U(0) + V(0));
- // floor_divide
- template <
- typename T,
- typename U,
- ::metal::enable_if_t<
- is_scalar_integral_v<T> && is_scalar_integral_v<U>,
- bool> = true>
- inline common_dtype<T, U> floor_divide(T x, U y) {
- const auto quot = x / y;
- return (x < 0) == (y < 0) ? quot : (x % y != 0) ? quot - 1 : quot;
- }
- template <
- typename T,
- typename U,
- ::metal::enable_if_t<
- is_scalar_floating_point_v<T> && is_scalar_floating_point_v<U>,
- bool> = true>
- inline common_dtype<T, U> floor_divide(T x, U y) {
- return ::metal::floor(x / y);
- }
- // Workaround for Metal compiler bug: the compiler produces wrong results
- // when optimizing fused (x / A) % B expressions for integral types.
- template <
- typename T,
- typename U,
- ::metal::enable_if_t<
- is_scalar_integral_v<T> && is_scalar_integral_v<U>,
- bool> = true>
- inline common_dtype<T, U> safe_mod(volatile T x, U y) {
- return x % y;
- }
- // fmod
- template <
- typename T,
- typename U,
- ::metal::enable_if_t<
- is_scalar_integral_v<T> && is_scalar_integral_v<U>,
- bool> = true>
- inline common_dtype<T, U> fmod(T x, U y) {
- return x % y;
- }
- template <
- typename T,
- typename U,
- ::metal::enable_if_t<
- is_scalar_floating_point_v<T> && is_scalar_floating_point_v<U>,
- bool> = true>
- inline common_dtype<T, U> fmod(T x, U y) {
- return ::metal::fmod(x, y);
- }
- // cast_to primitives
- // - No-op if types as the same
- template <
- typename T,
- typename U,
- ::metal::enable_if_t<::metal::is_same_v<U, T>, bool> = true>
- inline T cast_to(const U from) {
- return from;
- }
- // - Simple cast between scalar and complex dtypes
- template <
- typename T,
- typename U,
- ::metal::enable_if_t<
- !::metal::is_same_v<U, T> && (is_complex_v<T> == is_complex_v<U>),
- bool> = true>
- inline T cast_to(const U from) {
- return static_cast<T>(from);
- }
- // - Scalar to complex
- template <
- typename T,
- typename U,
- ::metal::enable_if_t<is_complex_v<T> && !is_complex_v<U>, bool> = true>
- inline T cast_to(const U from) {
- return T(float(from), 0.0);
- }
- // - Complex to scalar (should not really be used, but exists for compliteness)
- template <
- typename T,
- typename U,
- ::metal::enable_if_t<!is_complex_v<T> && is_complex_v<U>, bool> = true>
- inline T cast_to(const U from) {
- return static_cast<T>(from.x);
- }
- // Generalizable math operators (used for both scalar and complex)
- template <
- typename T,
- typename U,
- ::metal::enable_if_t<!is_complex_v<T>, bool> = true>
- inline common_dtype<T, U> mul(const T x, const U y) {
- return x * y;
- }
- template <
- typename T,
- typename U,
- ::metal::enable_if_t<is_complex_v<T> && is_complex_v<U>, bool> = true>
- inline common_dtype<T, U> mul(const T x, const U y) {
- return T(x.x * y.x - x.y * y.y, x.x * y.y + x.y * y.x);
- }
- template <
- typename T,
- typename U,
- ::metal::enable_if_t<!is_complex_v<T>, bool> = true>
- inline common_dtype<T, U> div(const T x, const U y) {
- return x / y;
- }
- template <
- typename T,
- typename U,
- ::metal::enable_if_t<is_complex_v<T> && is_complex_v<U>, bool> = true>
- inline common_dtype<T, U> div(const T x, const U y) {
- return T(::metal::dot(x, y), x.y * y.x - x.x * y.y) / ::metal::dot(y, y);
- }
- // Remainder operator
- template <
- typename T,
- typename U,
- ::metal::enable_if_t<
- is_scalar_floating_point_v<T> || is_scalar_floating_point_v<U>,
- bool> = true>
- inline float remainder(const T x, const U y) {
- const auto x_f = static_cast<float>(x);
- const auto y_f = static_cast<float>(y);
- return x_f - y_f * floor_divide(x_f, y_f);
- }
- template <
- typename T,
- typename U,
- ::metal::enable_if_t<
- is_scalar_integral_v<T> && is_scalar_integral_v<U>,
- bool> = true>
- inline common_dtype<T, U> remainder(const T x, const U y) {
- auto rc = x % y;
- return rc == 0 || (x ^ y) > 0 ? rc : rc + y;
- }
- // Based on aten/src/ATen/native/Pow.h
- template <
- typename T,
- ::metal::enable_if_t<is_scalar_integral_v<T>, bool> = true>
- inline T powi_impl(T a, T b) {
- T result = 1;
- while (b) {
- if (b & 1) {
- result *= a;
- }
- b /= 2;
- a *= a;
- }
- return result;
- }
- template <
- typename T,
- typename U,
- ::metal::enable_if_t<
- is_scalar_floating_point_v<T> || is_scalar_floating_point_v<U>,
- bool> = true>
- inline float pow(T a, U b) {
- return ::metal::precise::pow(static_cast<float>(a), static_cast<float>(b));
- }
- // Complex pow - use polar form: a = r*e^(i*theta)
- // a^b = exp(b * log(a)) = exp(b * (log(r) + i*theta))
- template <
- typename T,
- typename U,
- ::metal::enable_if_t<is_complex_v<T> && is_complex_v<U>, bool> = true>
- inline float2 pow(T a, U b) {
- // Convert a to polar form
- // Use explicit computation instead of length() due to numerical issues
- const auto r = ::metal::precise::sqrt(a.x * a.x + a.y * a.y);
- // Special case: if r is 0, return 0
- if (r == 0.0) {
- return float2(0.0, 0.0);
- }
- const auto theta = ::metal::precise::atan2(a.y, a.x);
- const auto log_r = ::metal::precise::log(r);
- // Calculate a^b = r^b * e^(i*theta*b)
- // new_r = exp(b.x * log(r) - b.y * theta)
- // new_theta = b.x * theta + b.y * log(r)
- const auto new_r = ::metal::precise::exp(b.x * log_r - b.y * theta);
- const auto new_theta = b.x * theta + b.y * log_r;
- return float2(
- new_r * ::metal::precise::cos(new_theta),
- new_r * ::metal::precise::sin(new_theta));
- }
- // Integral pow - unsigned types
- template <
- typename T,
- typename U,
- ::metal::enable_if_t<
- is_scalar_integral_v<T> && !::metal::is_signed_v<T>,
- bool> = true>
- inline T pow(T a, U b) {
- return powi_impl(a, T(b));
- }
- // Integral pow - signed types
- template <
- typename T,
- typename U,
- ::metal::enable_if_t<
- is_scalar_integral_v<T>&& ::metal::is_signed_v<T>,
- bool> = true>
- inline T pow(T a, U b) {
- if (b < 0) {
- if (a == 1) {
- return 1;
- } else if (a == -1) {
- auto negative = (-b) % static_cast<T>(2);
- return negative ? -1 : 1;
- } else {
- return 0;
- }
- }
- return powi_impl(a, T(b));
- }
- // Based on algorithm described in
- // https://docs.oracle.com/cd/E19957-01/806-3568/ncg_goldberg.html#1202
- inline float log1p(float x) {
- const auto xp1 = 1.0f + x;
- // First two elements of Taylor series for log(1+x) in Horner's form are:
- // log(1+x) = x * (1 - x * (.5 ...)), but if 1 + x == x, then it's just x
- if (xp1 == 1.0f) {
- return x;
- }
- auto rc = ::metal::precise::log(xp1);
- if (x > -.5 && x < .5) {
- // Order of operations is important here for higher precision
- rc *= x / (xp1 - 1.0f);
- }
- return rc;
- }
- // The function is ported from mlx
- inline float2 log1p(float2 in) {
- float x = in.x;
- float y = in.y;
- float zabs = ::metal::precise::sqrt(x * x + y * y);
- float theta = ::metal::atan2(y, x + 1);
- if (zabs < 0.5f) {
- float r = x * (2 + x) + y * y;
- if (r == 0) { // handle underflow
- return {x, theta};
- }
- return {0.5f * log1p(r), theta};
- } else {
- auto z0 = ::metal::sqrt((x + 1) * (x + 1) + y * y);
- return {::metal::log(z0), theta};
- }
- }
- template <typename T1, typename T2 = T1>
- struct pair {
- T1 first;
- T2 second;
- };
- template <typename T>
- inline T conj(T a) {
- return a;
- }
- template <>
- inline half2 conj(half2 a) {
- return half2(a.x, -a.y);
- }
- template <>
- inline float2 conj(float2 a) {
- return float2(a.x, -a.y);
- }
- #define INSTANTIATE_FOR_ALL_TYPES(MACRO) \
- MACRO(float); \
- MACRO(half); \
- MACRO(bfloat); \
- MACRO(float2); \
- MACRO(long); \
- MACRO(char); \
- MACRO(uchar); \
- MACRO(short); \
- MACRO(int);
- #define INSTANTIATE_FOR_FLOAT_TYPES(MACRO) \
- MACRO(float); \
- MACRO(half); \
- MACRO(bfloat);
- } // namespace metal
- } // namespace c10
- #else
- #error "This file should not be included when either TORCH_STABLE_ONLY or TORCH_TARGET_VERSION is defined."
- #endif // !defined(TORCH_STABLE_ONLY) && !defined(TORCH_TARGET_VERSION)
|