utils.h 12 KB

123456789101112131415161718192021222324252627282930313233343536373839404142434445464748495051525354555657585960616263646566676869707172737475767778798081828384858687888990919293949596979899100101102103104105106107108109110111112113114115116117118119120121122123124125126127128129130131132133134135136137138139140141142143144145146147148149150151152153154155156157158159160161162163164165166167168169170171172173174175176177178179180181182183184185186187188189190191192193194195196197198199200201202203204205206207208209210211212213214215216217218219220221222223224225226227228229230231232233234235236237238239240241242243244245246247248249250251252253254255256257258259260261262263264265266267268269270271272273274275276277278279280281282283284285286287288289290291292293294295296297298299300301302303304305306307308309310311312313314315316317318319320321322323324325326327328329330331332333334335336337338339340341342343344345346347348349350351352353354355356357358359360361362363364365366367368369370371372373374375376377378379380381382383384385386387388389390391392393394395396397398399400401402403404405406407408409410411412413414415416417418419420421422423424425426427428429430431432433434435436437438439440441442443444445446447448449450451452453454455456457458459460461462463464465466467468469470471472473474475476477478479480481482483484485486
  1. #if !defined(TORCH_STABLE_ONLY) && !defined(TORCH_TARGET_VERSION)
  2. // Metal helper functions
  3. #pragma once
  4. #include <c10/metal/common.h>
  5. #include <metal_stdlib>
  6. namespace c10 {
  7. namespace metal {
  8. namespace detail {
  9. template <typename T>
  10. struct vectypes {};
  11. template <>
  12. struct vectypes<float> {
  13. using type4 = float4;
  14. using type3 = float3;
  15. using type2 = float2;
  16. };
  17. template <>
  18. struct vectypes<half> {
  19. using type4 = half4;
  20. using type3 = half3;
  21. using type2 = half2;
  22. };
  23. template <>
  24. struct vectypes<bfloat> {
  25. using type4 = bfloat4;
  26. using type3 = bfloat3;
  27. using type2 = bfloat2;
  28. };
  29. template <>
  30. struct vectypes<short> {
  31. using type4 = short4;
  32. using type3 = short3;
  33. using type2 = short2;
  34. };
  35. template <>
  36. struct vectypes<int> {
  37. using type4 = int4;
  38. using type3 = int3;
  39. using type2 = int2;
  40. };
  41. template <>
  42. struct vectypes<long> {
  43. using type4 = short4;
  44. using type3 = short3;
  45. using type2 = short2;
  46. };
  47. template <typename T>
  48. struct OpMathType {
  49. using type = T;
  50. };
  51. template <>
  52. struct OpMathType<half> {
  53. using type = float;
  54. };
  55. template <>
  56. struct OpMathType<short> {
  57. using type = int;
  58. };
  59. template <>
  60. struct OpMathType<char> {
  61. using type = int;
  62. };
  63. template <>
  64. struct OpMathType<uchar> {
  65. using type = int;
  66. };
  67. template <>
  68. struct OpMathType<bfloat> {
  69. using type = float;
  70. };
  71. // Type promotion structure for higher precision accumulation
  72. template <typename T>
  73. struct AccumulationType {
  74. using type = T;
  75. };
  76. // Specialization for half - promote to float for accumulation
  77. template <>
  78. struct AccumulationType<half> {
  79. using type = float;
  80. };
  81. // Specialization for bfloat - promote to float for accumulation
  82. template <>
  83. struct AccumulationType<bfloat> {
  84. using type = float;
  85. };
  86. } // namespace detail
  87. template <typename T>
  88. ::metal::enable_if_t<::metal::is_floating_point_v<T>, T> max(T a, T b) {
  89. return ::metal::isunordered(a, b) ? NAN : ::metal::max(a, b);
  90. }
  91. template <typename T, typename U>
  92. ::metal::enable_if_t<::metal::is_integral_v<T>&& ::metal::is_integral_v<U>, T>
  93. max(T a, U b) {
  94. return ::metal::max(a, static_cast<T>(b));
  95. }
  96. template <typename T>
  97. ::metal::enable_if_t<::metal::is_floating_point_v<T>, T> min(T a, T b) {
  98. return ::metal::isunordered(a, b) ? NAN : ::metal::min(a, b);
  99. }
  100. template <typename T, typename U>
  101. ::metal::enable_if_t<::metal::is_integral_v<T>&& ::metal::is_integral_v<U>, T>
  102. min(T a, U b) {
  103. return ::metal::min(a, static_cast<T>(b));
  104. }
  105. template <>
  106. inline bfloat min(bfloat a, bfloat b) {
  107. return bfloat(
  108. ::metal::isunordered(a, b) ? NAN : ::metal::min(float(a), float(b)));
  109. }
  110. template <>
  111. inline bfloat max(bfloat a, bfloat b) {
  112. return bfloat(
  113. ::metal::isunordered(a, b) ? NAN : ::metal::max(float(a), float(b)));
  114. }
  115. template <typename T>
  116. using vec2type_t = typename detail::vectypes<T>::type2;
  117. template <typename T>
  118. using vec4type_t = typename detail::vectypes<T>::type4;
  119. template <typename T>
  120. using opmath_t = typename detail::OpMathType<T>::type;
  121. template <typename T>
  122. using accum_t = typename detail::AccumulationType<T>::type;
  123. // TODO: Move it to type_traits header may be
  124. template <typename F, typename... Args>
  125. using result_of = decltype(::metal::declval<F>()(::metal::declval<Args>()...));
  126. template <typename T>
  127. constexpr constant bool is_complex_v =
  128. ::metal::is_same_v<T, float2> || ::metal::is_same_v<T, half2>;
  129. template <typename T>
  130. constexpr constant bool is_scalar_floating_point_v =
  131. ::metal::is_floating_point_v<T> && ::metal::is_scalar_v<T>;
  132. template <typename T>
  133. constexpr constant bool is_scalar_integral_v =
  134. ::metal::is_integral_v<T> && ::metal::is_scalar_v<T>;
  135. template <typename U, typename V>
  136. using common_dtype = decltype(U(0) + V(0));
  137. // floor_divide
  138. template <
  139. typename T,
  140. typename U,
  141. ::metal::enable_if_t<
  142. is_scalar_integral_v<T> && is_scalar_integral_v<U>,
  143. bool> = true>
  144. inline common_dtype<T, U> floor_divide(T x, U y) {
  145. const auto quot = x / y;
  146. return (x < 0) == (y < 0) ? quot : (x % y != 0) ? quot - 1 : quot;
  147. }
  148. template <
  149. typename T,
  150. typename U,
  151. ::metal::enable_if_t<
  152. is_scalar_floating_point_v<T> && is_scalar_floating_point_v<U>,
  153. bool> = true>
  154. inline common_dtype<T, U> floor_divide(T x, U y) {
  155. return ::metal::floor(x / y);
  156. }
  157. // Workaround for Metal compiler bug: the compiler produces wrong results
  158. // when optimizing fused (x / A) % B expressions for integral types.
  159. template <
  160. typename T,
  161. typename U,
  162. ::metal::enable_if_t<
  163. is_scalar_integral_v<T> && is_scalar_integral_v<U>,
  164. bool> = true>
  165. inline common_dtype<T, U> safe_mod(volatile T x, U y) {
  166. return x % y;
  167. }
  168. // fmod
  169. template <
  170. typename T,
  171. typename U,
  172. ::metal::enable_if_t<
  173. is_scalar_integral_v<T> && is_scalar_integral_v<U>,
  174. bool> = true>
  175. inline common_dtype<T, U> fmod(T x, U y) {
  176. return x % y;
  177. }
  178. template <
  179. typename T,
  180. typename U,
  181. ::metal::enable_if_t<
  182. is_scalar_floating_point_v<T> && is_scalar_floating_point_v<U>,
  183. bool> = true>
  184. inline common_dtype<T, U> fmod(T x, U y) {
  185. return ::metal::fmod(x, y);
  186. }
  187. // cast_to primitives
  188. // - No-op if types as the same
  189. template <
  190. typename T,
  191. typename U,
  192. ::metal::enable_if_t<::metal::is_same_v<U, T>, bool> = true>
  193. inline T cast_to(const U from) {
  194. return from;
  195. }
  196. // - Simple cast between scalar and complex dtypes
  197. template <
  198. typename T,
  199. typename U,
  200. ::metal::enable_if_t<
  201. !::metal::is_same_v<U, T> && (is_complex_v<T> == is_complex_v<U>),
  202. bool> = true>
  203. inline T cast_to(const U from) {
  204. return static_cast<T>(from);
  205. }
  206. // - Scalar to complex
  207. template <
  208. typename T,
  209. typename U,
  210. ::metal::enable_if_t<is_complex_v<T> && !is_complex_v<U>, bool> = true>
  211. inline T cast_to(const U from) {
  212. return T(float(from), 0.0);
  213. }
  214. // - Complex to scalar (should not really be used, but exists for compliteness)
  215. template <
  216. typename T,
  217. typename U,
  218. ::metal::enable_if_t<!is_complex_v<T> && is_complex_v<U>, bool> = true>
  219. inline T cast_to(const U from) {
  220. return static_cast<T>(from.x);
  221. }
  222. // Generalizable math operators (used for both scalar and complex)
  223. template <
  224. typename T,
  225. typename U,
  226. ::metal::enable_if_t<!is_complex_v<T>, bool> = true>
  227. inline common_dtype<T, U> mul(const T x, const U y) {
  228. return x * y;
  229. }
  230. template <
  231. typename T,
  232. typename U,
  233. ::metal::enable_if_t<is_complex_v<T> && is_complex_v<U>, bool> = true>
  234. inline common_dtype<T, U> mul(const T x, const U y) {
  235. return T(x.x * y.x - x.y * y.y, x.x * y.y + x.y * y.x);
  236. }
  237. template <
  238. typename T,
  239. typename U,
  240. ::metal::enable_if_t<!is_complex_v<T>, bool> = true>
  241. inline common_dtype<T, U> div(const T x, const U y) {
  242. return x / y;
  243. }
  244. template <
  245. typename T,
  246. typename U,
  247. ::metal::enable_if_t<is_complex_v<T> && is_complex_v<U>, bool> = true>
  248. inline common_dtype<T, U> div(const T x, const U y) {
  249. return T(::metal::dot(x, y), x.y * y.x - x.x * y.y) / ::metal::dot(y, y);
  250. }
  251. // Remainder operator
  252. template <
  253. typename T,
  254. typename U,
  255. ::metal::enable_if_t<
  256. is_scalar_floating_point_v<T> || is_scalar_floating_point_v<U>,
  257. bool> = true>
  258. inline float remainder(const T x, const U y) {
  259. const auto x_f = static_cast<float>(x);
  260. const auto y_f = static_cast<float>(y);
  261. return x_f - y_f * floor_divide(x_f, y_f);
  262. }
  263. template <
  264. typename T,
  265. typename U,
  266. ::metal::enable_if_t<
  267. is_scalar_integral_v<T> && is_scalar_integral_v<U>,
  268. bool> = true>
  269. inline common_dtype<T, U> remainder(const T x, const U y) {
  270. auto rc = x % y;
  271. return rc == 0 || (x ^ y) > 0 ? rc : rc + y;
  272. }
  273. // Based on aten/src/ATen/native/Pow.h
  274. template <
  275. typename T,
  276. ::metal::enable_if_t<is_scalar_integral_v<T>, bool> = true>
  277. inline T powi_impl(T a, T b) {
  278. T result = 1;
  279. while (b) {
  280. if (b & 1) {
  281. result *= a;
  282. }
  283. b /= 2;
  284. a *= a;
  285. }
  286. return result;
  287. }
  288. template <
  289. typename T,
  290. typename U,
  291. ::metal::enable_if_t<
  292. is_scalar_floating_point_v<T> || is_scalar_floating_point_v<U>,
  293. bool> = true>
  294. inline float pow(T a, U b) {
  295. return ::metal::precise::pow(static_cast<float>(a), static_cast<float>(b));
  296. }
  297. // Complex pow - use polar form: a = r*e^(i*theta)
  298. // a^b = exp(b * log(a)) = exp(b * (log(r) + i*theta))
  299. template <
  300. typename T,
  301. typename U,
  302. ::metal::enable_if_t<is_complex_v<T> && is_complex_v<U>, bool> = true>
  303. inline float2 pow(T a, U b) {
  304. // Convert a to polar form
  305. // Use explicit computation instead of length() due to numerical issues
  306. const auto r = ::metal::precise::sqrt(a.x * a.x + a.y * a.y);
  307. // Special case: if r is 0, return 0
  308. if (r == 0.0) {
  309. return float2(0.0, 0.0);
  310. }
  311. const auto theta = ::metal::precise::atan2(a.y, a.x);
  312. const auto log_r = ::metal::precise::log(r);
  313. // Calculate a^b = r^b * e^(i*theta*b)
  314. // new_r = exp(b.x * log(r) - b.y * theta)
  315. // new_theta = b.x * theta + b.y * log(r)
  316. const auto new_r = ::metal::precise::exp(b.x * log_r - b.y * theta);
  317. const auto new_theta = b.x * theta + b.y * log_r;
  318. return float2(
  319. new_r * ::metal::precise::cos(new_theta),
  320. new_r * ::metal::precise::sin(new_theta));
  321. }
  322. // Integral pow - unsigned types
  323. template <
  324. typename T,
  325. typename U,
  326. ::metal::enable_if_t<
  327. is_scalar_integral_v<T> && !::metal::is_signed_v<T>,
  328. bool> = true>
  329. inline T pow(T a, U b) {
  330. return powi_impl(a, T(b));
  331. }
  332. // Integral pow - signed types
  333. template <
  334. typename T,
  335. typename U,
  336. ::metal::enable_if_t<
  337. is_scalar_integral_v<T>&& ::metal::is_signed_v<T>,
  338. bool> = true>
  339. inline T pow(T a, U b) {
  340. if (b < 0) {
  341. if (a == 1) {
  342. return 1;
  343. } else if (a == -1) {
  344. auto negative = (-b) % static_cast<T>(2);
  345. return negative ? -1 : 1;
  346. } else {
  347. return 0;
  348. }
  349. }
  350. return powi_impl(a, T(b));
  351. }
  352. // Based on algorithm described in
  353. // https://docs.oracle.com/cd/E19957-01/806-3568/ncg_goldberg.html#1202
  354. inline float log1p(float x) {
  355. const auto xp1 = 1.0f + x;
  356. // First two elements of Taylor series for log(1+x) in Horner's form are:
  357. // log(1+x) = x * (1 - x * (.5 ...)), but if 1 + x == x, then it's just x
  358. if (xp1 == 1.0f) {
  359. return x;
  360. }
  361. auto rc = ::metal::precise::log(xp1);
  362. if (x > -.5 && x < .5) {
  363. // Order of operations is important here for higher precision
  364. rc *= x / (xp1 - 1.0f);
  365. }
  366. return rc;
  367. }
  368. // The function is ported from mlx
  369. inline float2 log1p(float2 in) {
  370. float x = in.x;
  371. float y = in.y;
  372. float zabs = ::metal::precise::sqrt(x * x + y * y);
  373. float theta = ::metal::atan2(y, x + 1);
  374. if (zabs < 0.5f) {
  375. float r = x * (2 + x) + y * y;
  376. if (r == 0) { // handle underflow
  377. return {x, theta};
  378. }
  379. return {0.5f * log1p(r), theta};
  380. } else {
  381. auto z0 = ::metal::sqrt((x + 1) * (x + 1) + y * y);
  382. return {::metal::log(z0), theta};
  383. }
  384. }
  385. template <typename T1, typename T2 = T1>
  386. struct pair {
  387. T1 first;
  388. T2 second;
  389. };
  390. template <typename T>
  391. inline T conj(T a) {
  392. return a;
  393. }
  394. template <>
  395. inline half2 conj(half2 a) {
  396. return half2(a.x, -a.y);
  397. }
  398. template <>
  399. inline float2 conj(float2 a) {
  400. return float2(a.x, -a.y);
  401. }
  402. #define INSTANTIATE_FOR_ALL_TYPES(MACRO) \
  403. MACRO(float); \
  404. MACRO(half); \
  405. MACRO(bfloat); \
  406. MACRO(float2); \
  407. MACRO(long); \
  408. MACRO(char); \
  409. MACRO(uchar); \
  410. MACRO(short); \
  411. MACRO(int);
  412. #define INSTANTIATE_FOR_FLOAT_TYPES(MACRO) \
  413. MACRO(float); \
  414. MACRO(half); \
  415. MACRO(bfloat);
  416. } // namespace metal
  417. } // namespace c10
  418. #else
  419. #error "This file should not be included when either TORCH_STABLE_ONLY or TORCH_TARGET_VERSION is defined."
  420. #endif // !defined(TORCH_STABLE_ONLY) && !defined(TORCH_TARGET_VERSION)