Update on "[MPS][BE][EZ] Aggregate macros"

malfet · malfet · commit 5ac55f9ebfe7 · 2025-02-28T12:33:40.000-08:00
Refactor `INSTANTIATE_UPSAMPLE_BILINEAR2D(DTYPE)`, `INSTANTIATE_UPSAMPLE_BICUBIC2D(DTYPE)` and `INSTANTIATE_UPSAMPLE_BILINEAR2DAA(DTYPE)` use common `INSTANTIATE_UPSAMPLE2D`
Then combine multiple invocations into `INSTANTIATE_UPSAMPLE_ALL`

I.e. functionally it's a no-op, but achieves the same with fewer lines of code

[ghstack-poisoned]
diff --git a/aten/src/ATen/native/mps/kernels/UpSample.metal b/aten/src/ATen/native/mps/kernels/UpSample.metal
@@ -456,17 +456,17 @@ kernel void upsample_bicubic2d_backward(
           constant bool& align_corners [[buffer(7)]],              \
           uint thread_index [[thread_position_in_grid]])
 
-#define INSTANTIATE_UPSAMPLE_2D_BACKWARD(NAME, DTYPE)                      \
-  template [[host_name("upsample" #NAME "_backward_" #DTYPE)]] kernel void \
-      upsample_##NAME##_backward<DTYPE>(                                   \
-          device AtomicType_t<DTYPE> * gradInputData [[buffer(0)]],        \
-          constant DTYPE * gradOutputData [[buffer(1)]],                   \
-          constant ulong4 & input_strides [[buffer(2)]],                   \
-          constant ulong4 & output_strides [[buffer(3)]],                  \
-          constant long4 & input_sizes [[buffer(4)]],                      \
-          constant long4 & output_sizes [[buffer(5)]],                     \
-          constant float2 & scales [[buffer(6)]],                          \
-          constant bool& align_corners [[buffer(7)]],                      \
+#define INSTANTIATE_UPSAMPLE_2D_BACKWARD(NAME, DTYPE)                       \
+  template [[host_name("upsample_" #NAME "_backward_" #DTYPE)]] kernel void \
+      upsample_##NAME##_backward<DTYPE>(                                    \
+          device AtomicType_t<DTYPE> * gradInputData [[buffer(0)]],         \
+          constant DTYPE * gradOutputData [[buffer(1)]],                    \
+          constant ulong4 & input_strides [[buffer(2)]],                    \
+          constant ulong4 & output_strides [[buffer(3)]],                   \
+          constant long4 & input_sizes [[buffer(4)]],                       \
+          constant long4 & output_sizes [[buffer(5)]],                      \
+          constant float2 & scales [[buffer(6)]],                           \
+          constant bool& align_corners [[buffer(7)]],                       \
           uint thread_index [[thread_position_in_grid]])
 
 #define INSTANTIATE_UPSAMPLE_LINEAR(DTYPE)                        \
