[MPSInductor] Fix nested loop var elimination #156566
Closed
Conversation
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
🔗 Helpful Links🧪 See artifacts and rendered test results at hud.pytorch.org/pr/156566
Note: Links to docs will display an error until the docs builds have been completed. ✅ You can merge normally! (2 Unrelated Failures)As of commit 7e04fb8 with merge base 1d993fa ( BROKEN TRUNK - The following job failed but were present on the merge base:👉 Rebase onto the `viable/strict` branch to avoid these failures
UNSTABLE - The following job is marked as unstable, possibly due to flakiness on trunk:
This comment was automatically generated by Dr. CI and updates every 15 minutes. |
pytorch-bot
bot
added
ciflow/inductor
ciflow/mps
malfet
added
topic: bug fixes
dcci
approved these changes
Jun 22, 2025
|
@pytorchbot merge -f "Lint + MPS are green" |
Merge startedYour change will be merged immediately since you used the force (-f) flag, bypassing any CI checks (ETA: 1-5 minutes). Please use Learn more about merging in the wiki. Questions? Feedback? Please reach out to the PyTorch DevX Team |
pytorchmergebot
pushed a commit
that referenced
this pull request
Jun 23, 2025
From less than max threadgroup size to less or equal to that, which eliminates redundant trivial loops.
I.e. it changes shader code generated for
```python
import torch
def f(x):
var, mean = torch.var_mean(x, dim=2, keepdim = True)
return x / var, var
torch.compile(f)(torch.rand(1, 16, 1024, dtype=torch.float32, device='mps'))
```
from
```metal
[[max_total_threads_per_threadgroup(1024)]]
kernel void generated_kernel(
device float* out_ptr1,
device float* out_ptr2,
constant float* in_ptr0,
uint2 thread_pos [[thread_position_in_grid]],
uint2 group_pos [[thread_position_in_threadgroup]]
) {
auto xindex = thread_pos.x;
auto r0_index = thread_pos.y;
int x0 = xindex;
threadgroup float3 tmp_acc_0[1024];
tmp_acc_0[r0_index * 1] = 0.0;
for(auto r0_1_cnt = 0; r0_1_cnt < 1; ++r0_1_cnt) {
int r0_1 = 1 * r0_index + r0_1_cnt;
auto tmp0 = in_ptr0[r0_1 + 1024*x0];
tmp_acc_0[r0_index * 1] = ::c10::metal::welford_combine(tmp_acc_0[r0_index * 1], float3(tmp0, 0.0, 1.0));
}
auto tmp1 = c10::metal::threadgroup_welford_combine(tmp_acc_0, 1024);
auto tmp2 = 1023.0;
auto tmp3 = tmp1.y / tmp2;
out_ptr1[x0] = static_cast<float>(tmp3);
for(auto r0_1_cnt = 0; r0_1_cnt < 1; ++r0_1_cnt) {
int r0_1 = 1 * r0_index + r0_1_cnt;
auto tmp4 = in_ptr0[r0_1 + 1024*x0];
auto tmp5 = tmp4 / tmp3;
out_ptr2[r0_1 + 1024*x0] = static_cast<float>(tmp5);
}
}
```
to
```metal
[[max_total_threads_per_threadgroup(1024)]]
kernel void generated_kernel(
device float* out_ptr1,
device float* out_ptr2,
constant float* in_ptr0,
uint2 thread_pos [[thread_position_in_grid]],
uint2 group_pos [[thread_position_in_threadgroup]]
) {
auto xindex = thread_pos.x;
auto r0_index = thread_pos.y;
int r0_1 = r0_index;
int x0 = xindex;
threadgroup float tmp_acc_0[1024];
auto tmp0 = in_ptr0[r0_1 + 1024*x0];
tmp_acc_0[r0_index * 1] = tmp0;
auto tmp1 = c10::metal::threadgroup_welford_reduce(tmp_acc_0, 1024);
auto tmp2 = 1023.0;
auto tmp3 = tmp1.y / tmp2;
out_ptr1[x0] = static_cast<float>(tmp3);
auto tmp4 = tmp0 / tmp3;
out_ptr2[r0_1 + 1024*x0] = static_cast<float>(tmp4);
}
``
Pull Request resolved: #156567
Approved by: https://github.com/dcci
ghstack dependencies: #156566
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
Stack from ghstack (oldest at bottom):
As reduction resuts must be kept around
Add regression test that is specific for this issue
Fixes #156426
cc @voznesenskym @penguinwu @EikanWang @jgong5 @Guobing-Chen @XiaobingSuper @zhuhaozhe @blzheng @wenzhe-nrv @jiayisunx @ipiszy @chenyang78 @kadeng @muchulee8 @amjames @chauhang @aakhundov