[MPSInductor] Speedup sum/prod reductions

[malfet]

Stack from ghstack (oldest at bottom):

• -> [MPSInductor] Speedup sum/prod reductions #150566

By using cooperative simd_sum/simd_product instead of a C-style for loop for threadgroup reductions. This also allows significantly reduce amount of shared memory needed to perform those reductions

Using such reduction increases the torch.compile performance for gpt-fast using stories110M from 29 tokens/sec to 630 tokens/sec on M4 and changes perf of torch.rand as follows:

size	before	after
512x512	202.1	131.8
1024x1024	780.6	176.9
2048x2048	1423.4	339.9
4096x4097	2982.2	1047.2

Unfortunately, none of the SIMDgroup operations are available for 64-bit integers, but one can simulate the behavior using using simd_shuffle_down of 64-bit values represented as int2 types, that yields reduction in $log_2(threadgroup\_size)$ steps. `mlx/kernels/reduction/ops.h contains an implementation of such algorithm, but alas it yields wrong results on M1/M2(and may be M3 machines) if not all threads in the simdgroup are active which could be observed by running

import torch
lib=torch.mps.compile_shader("""
kernel void do_sum(device int* out, constant int* in, uint idx [[thread_position_in_grid]]) {
  out[idx] = metal::simd_shuffle_down(in[idx], 8);
}
""")
x=torch.arange(22, device='mps', dtype=torch.int32)
y=torch.empty_like(x)
lib.do_sum(y, x)
print(y)

that returns following on M4

tensor([ 8,  9, 10, 11, 12, 13, 14, 15, 16, 17, 18, 19, 20, 21,  0,  0,  0,  0, 0,  0,  0,  0], device='mps:0', dtype=torch.int32)

but same kernel running on M1 returns

tensor([ 8,  9, 10, 11, 12, 13, 14, 15, 16, 17, 18, 19, 20, 21, 14, 15, 16, 17, 18, 19, 20, 21], device='mps:0', dtype=torch.int32)

This discrepancy in behavior can be addressed by using simd_shuffle_and_fill_down, but any kernels using simd_shuffle_and_fill_down cause an internal compiler error on MacOS-13.2. Considering that OS is to be EOL soon, skip the offending tests.

cc @voznesenskym @penguinwu @EikanWang @jgong5 @Guobing-Chen @XiaobingSuper @zhuhaozhe @blzheng @wenzhe-nrv @jiayisunx @ipiszy @chenyang78 @kadeng @muchulee8 @amjames @chauhang @aakhundov

[pytorch-bot]

🔗 Helpful Links

🧪 See artifacts and rendered test results at hud.pytorch.org/pr/150566

• 📄 Preview Python docs built from this PR
• 📄 Preview C++ docs built from this PR
• ❓ Need help or want to give feedback on the CI? Visit the bot commands wiki or our office hours

Note: Links to docs will display an error until the docs builds have been completed.

❌ 1 New Failure, 22 Pending, 1 Unrelated Failure

As of commit 7525368 with merge base 15dbad2 ():

NEW FAILURE - The following job has failed:

• pull / cuda12.4-py3.10-gcc9-sm75 / test (pr_time_benchmarks, 1, 1, linux.g4dn.metal.nvidia.gpu) (gh)
  REGRESSION: benchmark ('aotdispatcher_training_subclass_cpu', 'compile_time_instruction_count') failed, actual result 9987959924 is 1.61% higher than expected 9830000000 ±+1.50% if this is an expected regression, please update the expected results.

BROKEN TRUNK - The following job failed but were present on the merge base:

👉 Rebase onto the `viable/strict` branch to avoid these failures

• pull / linux-jammy-xpu-2025.0-py3.9 / build (gh) (trunk failure)
  /usr/lib/gcc/x86_64-linux-gnu/11/../../../../include/c++/11/sstream:152:52: error: expected value in expression

This comment was automatically generated by Dr. CI and updates every 15 minutes.

[malfet]

@pytorchbot merge -f "Lint + MPS are finally green"

[pytorchmergebot]

Merge started

Your change will be merged immediately since you used the force (-f) flag, bypassing any CI checks (ETA: 1-5 minutes). Please use -f as last resort and instead consider -i/--ignore-current to continue the merge ignoring current failures. This will allow currently pending tests to finish and report signal before the merge.

Learn more about merging in the wiki.

Questions? Feedback? Please reach out to the PyTorch DevX Team

Advanced Debugging

Check the merge workflow status
here