[BE][MPS] Infer results of functor (pytorch#147182)

Do not assume that functor will return the same results as its arguments, but rather dynamically infer it using `decltype` and `::metal::declval`
This is a no-op that prepares for migration of `copysign` of integral arguments, that would return a float
Pull Request resolved: pytorch#147182
Approved by: https://github.com/dcci

Author: malfet

aten/src/ATen/native/mps/kernels/BinaryKernel.metal

@@ -66,16 +66,22 @@ struct nextafter_functor {
   }
 };
 
+// Future BinaryTensorIterator
+template <typename T, typename F>
+using result_of = decltype(::metal::declval<F>()(
+    ::metal::declval<T>(),
+    ::metal::declval<T>()));
+
 template <typename T, typename F>
 kernel void binary_indexing(
     constant void* input_ [[buffer(0)]],
     constant void* other_ [[buffer(1)]],
     device void* out_ [[buffer(2)]],
     constant uint3* offsets [[buffer(3)]],
     uint tid [[thread_position_in_grid]]) {
-  device T* out = (device T*)((device uint8_t*)out_ + offsets[tid].x);
-  constant T* input = (constant T*)((constant uint8_t*)input_ + offsets[tid].y);
-  constant T* other = (constant T*)((constant uint8_t*)other_ + offsets[tid].z);
+  auto out = (device result_of<T, F>*)((device uint8_t*)out_ + offsets[tid].x);
+  auto input = (constant T*)((constant uint8_t*)input_ + offsets[tid].y);
+  auto other = (constant T*)((constant uint8_t*)other_ + offsets[tid].z);
   F f;
   *out = f(*input, *other);
 }
@@ -84,7 +90,7 @@ template <typename T, typename F>
 kernel void binary_dense(
     constant T* input [[buffer(0)]],
     constant T* other [[buffer(1)]],
-    device T* out [[buffer(2)]],
+    device result_of<T, F>* out [[buffer(2)]],
     uint tid [[thread_position_in_grid]]) {
   F f;
   out[tid] = f(input[tid], other[tid]);
@@ -102,7 +108,7 @@ kernel void binary_dense(
   binary_dense<DTYPE, NAME##_functor>(                       \
       constant DTYPE * input_,                               \
       constant DTYPE * other_,                               \
-      device DTYPE * out_,                                   \
+      device result_of<DTYPE, NAME##_functor> * out_,        \
       uint tid)
 
 template <typename T>