Skip to content

[metal] Reuse Inductor's MetalOverrides for MSL expression emission#1853

Merged
aditvenk merged 1 commit intomainfrom
aditvenk/stack/19
Apr 8, 2026
Merged

[metal] Reuse Inductor's MetalOverrides for MSL expression emission#1853
aditvenk merged 1 commit intomainfrom
aditvenk/stack/19

Conversation

@aditvenk
Copy link
Copy Markdown
Contributor

@aditvenk aditvenk commented Mar 27, 2026

Stacked PRs:


[metal] Reuse Inductor's MetalOverrides for MSL expression emission

Subclass Inductor's torch._inductor.codegen.mps.MetalOverrides instead
of reimplementing ~170 lines of math/cast/comparison overrides. This
brings NaN-safe math (c10::metal::max/min), IEEE-compliant precision
(metal::precise::*), and type-safe casts (decltype) for free.

The C++ namespace syntax (::) in Inductor's expression output is
replaced with . before Python AST parsing, then restored to :: by the
MSL walker. Override where() to use Python ternary (parseable by AST),
and _special_unary/_special_binary to skip V.kernel.headers dependency.

@aditvenk aditvenk force-pushed the aditvenk/stack/19 branch from b3e1da8 to e8fe296 Compare March 27, 2026 22:35
@meta-cla meta-cla bot added the CLA Signed This label is managed by the Meta Open Source bot. label Mar 27, 2026
@aditvenk aditvenk marked this pull request as draft March 27, 2026 22:47
@aditvenk aditvenk changed the base branch from aditvenk/stack/18 to main March 27, 2026 22:47
@aditvenk aditvenk force-pushed the aditvenk/stack/19 branch from e8fe296 to 8ef44cd Compare March 27, 2026 22:48
@aditvenk aditvenk changed the base branch from main to aditvenk/stack/18 March 27, 2026 22:48
@aditvenk aditvenk marked this pull request as ready for review March 27, 2026 22:48
@aditvenk aditvenk requested review from jansel, malfet and oulgen and removed request for oulgen March 27, 2026 22:54
@aditvenk aditvenk marked this pull request as draft March 28, 2026 00:11
@aditvenk aditvenk changed the base branch from aditvenk/stack/18 to main March 28, 2026 00:11
@aditvenk aditvenk force-pushed the aditvenk/stack/19 branch from 8ef44cd to 6bd2521 Compare March 28, 2026 00:12
@aditvenk aditvenk changed the base branch from main to aditvenk/stack/18 March 28, 2026 00:13
@aditvenk aditvenk marked this pull request as ready for review March 28, 2026 00:14
@aditvenk aditvenk marked this pull request as draft March 28, 2026 01:23
@aditvenk aditvenk changed the base branch from aditvenk/stack/18 to main March 28, 2026 01:23
@aditvenk aditvenk force-pushed the aditvenk/stack/19 branch from 6bd2521 to 73bed39 Compare March 28, 2026 01:24
@aditvenk aditvenk changed the base branch from main to aditvenk/stack/18 March 28, 2026 01:26
@aditvenk aditvenk marked this pull request as ready for review March 28, 2026 01:26
@aditvenk aditvenk marked this pull request as draft March 28, 2026 01:37
@aditvenk aditvenk changed the base branch from main to aditvenk/stack/18 March 28, 2026 02:45
@aditvenk aditvenk marked this pull request as ready for review March 28, 2026 02:45
@aditvenk aditvenk marked this pull request as draft March 28, 2026 04:44
@aditvenk aditvenk changed the base branch from aditvenk/stack/18 to main March 28, 2026 04:44
@aditvenk aditvenk force-pushed the aditvenk/stack/19 branch from 2c95aa5 to 20ba09a Compare March 28, 2026 04:44
@aditvenk aditvenk changed the base branch from main to aditvenk/stack/18 March 28, 2026 04:45
@aditvenk aditvenk marked this pull request as ready for review March 28, 2026 04:45
@aditvenk aditvenk marked this pull request as draft March 28, 2026 04:47
@aditvenk aditvenk changed the base branch from aditvenk/stack/18 to main March 28, 2026 04:47
@aditvenk aditvenk force-pushed the aditvenk/stack/19 branch from 20ba09a to 849c5c3 Compare March 28, 2026 04:47
@aditvenk aditvenk changed the base branch from main to aditvenk/stack/18 March 28, 2026 04:47
@aditvenk aditvenk marked this pull request as ready for review March 28, 2026 04:48
@aditvenk aditvenk marked this pull request as draft March 28, 2026 06:46
@aditvenk aditvenk changed the base branch from aditvenk/stack/18 to main March 28, 2026 06:46
@aditvenk aditvenk force-pushed the aditvenk/stack/19 branch from 849c5c3 to 53876fa Compare March 28, 2026 06:46
@aditvenk aditvenk changed the base branch from main to aditvenk/stack/18 March 28, 2026 06:46
@aditvenk aditvenk marked this pull request as ready for review March 28, 2026 06:47
@aditvenk aditvenk marked this pull request as draft March 28, 2026 17:18
@aditvenk aditvenk changed the base branch from aditvenk/stack/18 to main March 28, 2026 17:18
@aditvenk aditvenk changed the base branch from main to aditvenk/stack/18 March 28, 2026 17:19
@aditvenk aditvenk marked this pull request as ready for review March 28, 2026 17:19
Copy link
Copy Markdown
Contributor

@jansel jansel left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

@aditvenk
Copy link
Copy Markdown
Contributor Author

Can you reuse the version of this in inductor: https://github.com/pytorch/pytorch/blob/47fa87a257f72a30c2028e1200c1673ed36125b8/torch/_inductor/codegen/mps.py#L185

Will take a look at it. From skimming it seems like it pulls in Inductor dependencies (V.kernel, CSEVariable) and also emits MSL differently.

@aditvenk aditvenk marked this pull request as draft March 31, 2026 04:49
@aditvenk aditvenk changed the base branch from aditvenk/stack/18 to main March 31, 2026 04:49
@aditvenk
Copy link
Copy Markdown
Contributor Author

Can you reuse the version of this in inductor: https://github.com/pytorch/pytorch/blob/47fa87a257f72a30c2028e1200c1673ed36125b8/torch/_inductor/codegen/mps.py#L185

Refactored to use Inductor's override for 99%. A couple of methods need to be handled slightly differently because Helion parses override strings as Python AST. While this works completely for Triton/Cute/Pallas because those are valid python syntax. For Metal, since it is C++, some things like ::, C++ ternary operator don't parse.

Subclass Inductor's torch._inductor.codegen.mps.MetalOverrides instead
of reimplementing ~170 lines of math/cast/comparison overrides. This
brings NaN-safe math (c10::metal::max/min), IEEE-compliant precision
(metal::precise::*), and type-safe casts (decltype) for free.

The C++ namespace syntax (::) in Inductor's expression output is
replaced with . before Python AST parsing, then restored to :: by the
MSL walker. Override where() to use Python ternary (parseable by AST),
and _special_unary/_special_binary to skip V.kernel.headers dependency.

stack-info: PR: #1853, branch: aditvenk/stack/19
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

CLA Signed This label is managed by the Meta Open Source bot.

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants