expm1f.h 3.8 KB

123456789101112131415161718192021222324252627282930313233343536373839404142434445464748495051525354555657585960616263646566676869707172737475767778798081828384858687888990919293949596979899100101102
  1. #if !defined(TORCH_STABLE_ONLY) && !defined(TORCH_TARGET_VERSION)
  2. // Copy-and-pasted from:
  3. // https://github.com/ml-explore/mlx/blob/99c33d011d63174f50cea37c3eede002958be6d3/mlx/backend/metal/kernels/expm1f.h
  4. #pragma once
  5. #include <metal_math>
  6. // Original license copied below:
  7. // Copyright (c) 2015-2023 Norbert Juffa
  8. // All rights reserved.
  9. //
  10. // Redistribution and use in source and binary forms, with or without
  11. // modification, are permitted provided that the following conditions
  12. // are met:
  13. //
  14. // 1. Redistributions of source code must retain the above copyright
  15. // notice, this list of conditions and the following disclaimer.
  16. //
  17. // 2. Redistributions in binary form must reproduce the above copyright
  18. // notice, this list of conditions and the following disclaimer in the
  19. // documentation and/or other materials provided with the distribution.
  20. //
  21. // THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS
  22. // "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT
  23. // LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR
  24. // A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT
  25. // HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL,
  26. // SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT
  27. // LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE,
  28. // DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY
  29. // THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
  30. // (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
  31. // OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
  32. namespace c10 {
  33. namespace metal {
  34. /* Compute exponential base e minus 1. Maximum ulp error = 0.997458
  35. i = rint(a/log(2)), f = a-i*log(2). Then expm1(a) = 2**i * (expm1(f)+1) - 1.
  36. Compute r = expm1(f). Then expm1(a)= 2 * (0.5 * 2**i * r + 0.5 * 2**i - 0.5).
  37. With t = 0.5*2**i, expm1(a) = 2*(r * t + t-0.5). However, for best accuracy,
  38. when i == 1, expm1(a)= 2*(r + 0.5), and when i == 0, expm1(a) = r.
  39. NOTE: Scale factor b is only applied if i < 0 or i > 1 (should be power of 2)
  40. */
  41. inline float expm1f_scaled_unchecked(float a, float b) {
  42. float f, j, r, s, t, u, v, x, y;
  43. int i;
  44. // exp(a) = 2**i * exp(f); i = rintf (a / log(2))
  45. j = ::metal::fma(1.442695f, a, 12582912.f); // 0x1.715476p0, 0x1.8p23
  46. j = j - 12582912.0f; // 0x1.8p23
  47. i = (int)j;
  48. f = ::metal::fma(j, -6.93145752e-1f, a);
  49. // approximate r = exp(f)-1 on interval [-log(2)/2, +log(2)/2]
  50. s = f * f;
  51. if (a == 0.0f)
  52. s = a; // ensure -0 is passed through
  53. // err = 0.997458 ulp1 = 11081805
  54. r = 1.97350979e-4f; // 0x1.9de000p-13
  55. r = ::metal::fma(r, f, 1.39309070e-3f); // 0x1.6d30bcp-10
  56. r = ::metal::fma(r, f, 8.33343994e-3f); // 0x1.1111f6p-7
  57. r = ::metal::fma(r, f, 4.16668020e-2f); // 0x1.55559ep-5
  58. r = ::metal::fma(r, f, 1.66666716e-1f); // 0x1.55555cp-3
  59. r = ::metal::fma(r, f, 4.99999970e-1f); // 0x1.fffffep-2
  60. u = (j == 1) ? (f + 0.5f) : f;
  61. v = ::metal::fma(r, s, u);
  62. s = 0.5f * b;
  63. t = ::metal::ldexp(s, i);
  64. y = t - s;
  65. x = (t - y) - s; // double-float canonicalization of difference
  66. r = ::metal::fma(v, t, x) + y;
  67. r = r + r;
  68. if (j == 0)
  69. r = v;
  70. if (j == 1)
  71. r = v + v;
  72. return r;
  73. }
  74. /* Compute exponential base e minus 1. max ulp err = 0.99746 */
  75. inline float expm1f(float a) {
  76. float r;
  77. r = expm1f_scaled_unchecked(a, 1.0f);
  78. /* handle severe overflow and underflow */
  79. if (::metal::abs(a - 1.0f) > 88.0f) {
  80. r = ::metal::pow(2, a);
  81. r = ::metal::fma(r, r, -1.0f);
  82. }
  83. return r;
  84. }
  85. } // namespace metal
  86. } // namespace c10
  87. #else
  88. #error "This file should not be included when either TORCH_STABLE_ONLY or TORCH_TARGET_VERSION is defined."
  89. #endif // !defined(TORCH_STABLE_ONLY) && !defined(TORCH_TARGET_VERSION)