Skip to content

Conversation

@Myrthan
Copy link

@Myrthan Myrthan commented Nov 25, 2025

Summary:
Enable Helion with MTIA for the first time. This is very basic setup to just make showcase it working.

Following diffs will try to enable more modes (e.g. eager) and more tests. I used default mtia machine arch setup for now + set triton_mtia backend + compile it through JIT flow.

Highlights/doubts:

  • Remove RefEager mode from the test, seems to not be used?
  • For some reason i have to call mtia_triton_launcher.init() twice (?), otherwise it result with inaccuracy.. i did not debug it fully
  • There is a check if TYPE_CHECKING: in init.py, not sure how does this work for functions, can i import triton? will CI be able to check this for me?
  • I disabled one test when run on mtia. Test complains about different type interfered from Triton-MTIA compiler (error: IntList but got GenericList)
  • Hardcoded grid size for mtia architectures, I will have to cover this part separately, sadly this is not fully covered in torch by Triton-MTIA, when doing Triton-MTIA update i created hardcoded stub, since it is not being used.

v2
We cannot expose arch names to opensource, hide them then in triton_mtia.
CPU added support for their backend, use their _get_triton_backend instead of calling triton.runtime.driver.active.get_current_target directly

v3
Aligned further to changes made when diff was not landed and i removed custom implementation for handling torch device. Removed calling is_available from torch/mtia/init.py

Differential Revision: D76375133

Summary:
Enable Helion with MTIA for the first time. This is very basic setup to just make showcase it working.

Following diffs will try to enable more modes (e.g. eager) and more tests. I used default mtia machine arch setup for now + set triton_mtia backend + compile it through JIT flow.

Highlights/doubts:
- Remove RefEager mode from the test, seems to not be used?
- For some reason i have to call mtia_triton_launcher.init() twice (?), otherwise it result with inaccuracy.. i did not debug it fully
- There is a check `if TYPE_CHECKING:` in __init__.py, not sure how does this work for functions, can i import triton? will CI be able to check this for me?
- I disabled one test when run on mtia. Test complains about different type interfered from Triton-MTIA compiler (error: IntList but got GenericList)
- Hardcoded grid size for mtia architectures, I will have to cover this part separately, sadly this is not fully covered in torch by Triton-MTIA, when doing Triton-MTIA update i created hardcoded stub, since it is not being used.


v2
We cannot expose arch names to opensource, hide them then in triton_mtia.
CPU added support for their backend, use their _get_triton_backend instead of calling triton.runtime.driver.active.get_current_target directly

v3
Aligned further to changes made when diff was not landed and i removed custom implementation for handling torch device. Removed calling is_available from torch/mtia/__init__.py

Differential Revision: D76375133
@meta-cla meta-cla bot added the CLA Signed This label is managed by the Meta Open Source bot. label Nov 25, 2025
@meta-codesync
Copy link

meta-codesync bot commented Nov 25, 2025

@Myrthan has exported this pull request. If you are a Meta employee, you can view the originating Diff in D76375133.

f"PyTorch version {min_version} or higher required",
)


Copy link
Contributor

Choose a reason for hiding this comment

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

I'd expect this to cause some lints, though for some reason CI isn't running. Can you do ./lint.sh install && ./lint.sh?


init_mtia_device()
# Ignore disk cache. Kernels will still keep an in-memory cache.
os.environ.setdefault("TRITON_ALWAYS_COMPILE", "1")
Copy link
Contributor

Choose a reason for hiding this comment

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

Is this required? Seems a bit odd to be running this code on every kernel launch and not cleaning it up.

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. fb-exported meta-exported

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants