[mps/inductor] Add support for fmod(). (pytorch#144449)

397 -> 395 tests failing. `static_cast<>` is because there are several overloads of `fmod()` that's otherwise ambiguous. I wonder if we should take in account NaN propagation (maybe it's not tested).

Pull Request resolved: pytorch#144449
Approved by: https://github.com/malfet

Author: dcci

test/inductor/test_mps_basic.py

@@ -43,6 +43,8 @@ class MPSBasicTests(TestCase):
     test_addmm = CommonTemplate.test_addmm
     test_cat_empty = CommonTemplate.test_cat_empty
     test_floordiv = CommonTemplate.test_floordiv
+    test_fmod = CommonTemplate.test_fmod
+    test_fmod_zero_dim = CommonTemplate.test_fmod_zero_dim
     test_inf = CommonTemplate.test_inf
     test_isinf = CommonTemplate.test_isinf
     test_isinf2 = CommonTemplate.test_isinf2

torch/_inductor/codegen/mps.py

@@ -206,6 +206,12 @@ def floor(x: CSEVariable) -> str:
     def sign(x: CSEVariable) -> str:
         return f"metal::sign({x})"
 
+    @staticmethod
+    def fmod(a: CSEVariable, b: CSEVariable) -> str:
+        typecast_a = f"static_cast<decltype({a}+{b})>({a})"
+        typecast_b = f"static_cast<decltype({a}+{b})>({b})"
+        return f"metal::fmod({typecast_a}, {typecast_b})"
+
 
 class MetalKernel(SIMDKernel):
     overrides = MetalOverrides  # type: ignore[assignment]