Dispatch unary_strided* as a 2D grid (grid.x = innermost dim after TensorIterator reorder, grid.y = product of outer dims) instead of a flat 1D grid. Kernel reads pos[0] straight from thread_position_in_grid.x and skip one div/mod per element in the coord decomposition. Adds a uint2 overload of pos_from_thread_index so the kernels stay one-liners.
TG sizing matters: when inner < maxThreadsPerThreadgroup we pack several outer rows per TG along y (tg_y = maxTG/tg_x). Without that, small inner dims (e.g. transposed 256x256) regress because TG count goes up and SIMD occupancy drops; with it, occupancy matches the 1D path and the saved div/mod is pure profit.
The kernel reads thread_position_in_grid as uint-per-axis so each dim must fit in uint32, but the product can exceed UINT32_MAX -- 2D dispatch addresses >4G total threads natively as long as neither axis alone overflows. TensorIterator's with_32bit_indexing() still decomposes when iter.numel() > INT32_MAX, so today the wider envelope isn't exercised, but a >4 GiB strided unary op already runs correctly end-to-end through the split path. Per-axis range is the caller's responsibility (mtl_dispatch1DJob trusts the same way).
Authored with Claude.
Test plan: torch.abs(x, out=y) on a permuted 256x256x256 float tensor:
import time, torch
x = torch.randn(256, 256, 256, device='mps').permute(1, 0, 2)
out = torch.empty(x.shape, device='mps')
for _ in range(20):
torch.abs(x, out=out)
torch.mps.synchronize()
t = time.perf_counter()
for _ in range(300):
torch.abs(x, out=out)
torch.mps.synchronize()
print(f"{(time.perf_counter()-t)/300*1e6:.1f} us")Before: ~651 us. After: ~380 us (~1.7x). Also re-ran the MPS test_copy / test_slice / test_view sweep -- all pass.
Full perf sweep across 3D permute patterns (abs into preallocated out)
```
pattern old_us new_us speedup
abs_p(1, 0, 2)|1024x1024x4 167.55 84.06 1.99x
abs_p(1, 0, 2)|128x128x128 87.32 45.30 1.93x
abs_p(1, 0, 2)|256x256x256 651.46 380.48 1.71x
abs_p(1, 0, 2)|256x256x32 88.03 45.03 1.96x
abs_p(1, 0, 2)|256x32x256 87.27 45.19 1.93x
abs_p(1, 0, 2)|32x256x256 87.52 47.67 1.84x
abs_p(1, 0, 2)|4x1024x1024 167.82 84.31 1.99x
abs_p(1, 0, 2)|512x512x8 87.93 46.25 1.90x
abs_p(1, 0, 2)|64x64x64 15.01 8.66 1.73x
abs_p(1, 0, 2)|8x512x512 87.46 45.51 1.92x
abs_p(2, 0, 1)|1024x1024x4 116.71 57.15 2.04x
abs_p(2, 0, 1)|128x128x128 61.78 55.86 1.11x
abs_p(2, 0, 1)|256x256x256 633.08 456.48 1.39x
abs_p(2, 0, 1)|256x256x32 68.86 55.73 1.24x
abs_p(2, 0, 1)|256x32x256 60.52 55.93 1.08x
abs_p(2, 0, 1)|32x256x256 61.29 59.31 1.03x
abs_p(2, 0, 1)|4x1024x1024 117.15 109.07 1.07x
abs_p(2, 0, 1)|512x512x8 67.96 33.10 2.05x
abs_p(2, 0, 1)|64x64x64 10.70 10.16 1.05x
abs_p(2, 0, 1)|8x512x512 60.73 55.68 1.09x
abs_p(2, 1, 0)|1024x1024x4 167.52 116.02 1.44x
abs_p(2, 1, 0)|128x128x128 87.07 58.02 1.50x
abs_p(2, 1, 0)|256x256x256 813.28 610.72 1.33x
abs_p(2, 1, 0)|256x256x32 91.84 65.53 1.40x
abs_p(2, 1, 0)|256x32x256 87.27 55.90 1.56x
abs_p(2, 1, 0)|32x256x256 87.07 56.14 1.55x
abs_p(2, 1, 0)|4x1024x1024 167.15 112.05 1.49x
abs_p(2, 1, 0)|512x512x8 91.22 55.07 1.66x
abs_p(2, 1, 0)|64x64x64 14.23 11.70 1.22x
abs_p(2, 1, 0)|8x512x512 86.99 56.32 1.54x
```
Binary strided kernels follow the same 1D-dispatch pattern and are unchanged here; the same treatment would likely yield a comparable win and is a natural follow-up.
Pull Request resolved: https://github.com//pull/185291
Approved by: https://github.com/Skylion007, https://github.com/kurtamohler
Community-Analysen & Experten-Meinungen 0
Verwandte Story-Cluster & Quellen (Vektor-KI)
Ähnliche Beiträge
Auch interessante Nachrichten trunk/63cc967a5743d037987d867b7ae38a73e0a3b3e6: [MPS] 2D dispatch for strided unary kernels (#185291)
Thematisch verwandte Begriffe: trunk63cc967a5743d037987d867b7ae38a73e0a3b3e6, dispatch, strided, unary · 6 Treffer
Telerik UI Flaws Let Attackers Chain AES-CBC Padding Oracle to Unauthenticated RCE
🚨 Typosquatting packages on PyPI are targeting 𝗿𝗲𝗾𝘂𝗲𝘀𝘁𝘀, one of the most widely used Python packages. Catch them before compromise.
Reverse Engineering / Authorized Penetration Testing / Security Research Skill Router Pack AI-powered routing + On-demand toolchain bootstrapping + Self-evolving knowledge base Supports Claude Code, Kiro, Cursor, Cline, and other AI coding clients
kaspersky kernel vuln windows 11
ICEDID Configuration Extractor
Videos werden geladen ...
Beiträge werden geladen ...
Videos werden geladen ...
Beiträge werden geladen ...
Videos werden geladen ...
Beiträge werden geladen ...
Videos werden geladen ...
Beiträge werden geladen ...
Videos werden geladen ...
SOCIAL SHARE CARD GENERATOR