Add Apple GPU backend for Triton

[imperatormk]

Out-of-tree Apple GPU (Metal) backend with:

C++ MLIR passes (loaded via TRITON_PASS_PLUGIN_PATH):

• AccelerateAppleMatmul: tt.dot → AppleMmaEncoding (8x8 simdgroup MMA)
• SimplifyGatherLayout: strip efficient_layout for Metal JIT
• ConvertTritonAppleGPUToLLVM: lower to LLVM IR with simdgroup intrinsics
• LowerGPUToAIR: gpu.thread_id/block_dim → AIR intrinsics

Python backend (pip installable, entry_points discovery):

• compiler.py: TTIR → TTGIR → LLVM IR → metallib pipeline
• driver.py: MPS GPU dispatch, buffer binding, scalar packing
• metal_utils.m: ObjC++ Metal bridge (compiled at install time)

metal-ir-pipeline (git submodule):

• LLVM IR → Metal AIR transforms → v1 typed bitcode → metallib
• 26 transformation passes for Metal GPU JIT compatibility

[imperatormk]

In combination with this work I was able to compile Qwen 3.5 with fla through torch inductor

[plotfi]

@imperatormk Are you able to include the metal-ir-pipeline code without a submodule? I think we want to be able to review everything involved in the stack.

[plotfi]

I also wonder if we couldn't start with a much simpler PR for an initial Metal Hello World backend?

I have an example of a really tiny backend here (does enough codeine for the vector-add kernel on amdgpu):

https://github.com/plotfi/triton/tree/plotfi-nano-backend/third_party/nano

[imperatormk]

1. Included it directly without submodule so you can have a clearer overview although still not sure what would be the best way to vendor this in the end, it's just a lot of code/files
2. You're correct that it's better to work on this incrementally, opened Add minimal Apple GPU backend #64 which also should be somewhat easier to review although the infra overhead is still there so it's still not 100% nano

ps. I had small issues when trying to build the upstream Triton on macOS, will quickly list them here because I am not sure they could have been avoided in another way without making cmake changes:

• third_party/nvidia/CMakeLists.txt tries to compile GSanLibrary.cu with the LLVM-built clang, which can't find libc++ on macOS. Needs if(NOT APPLE) guard
• examples/plugins/TritonPlugin.cpp fails to link on macOS — undefined LLVM symbols from libtriton.so. Also if(NOT APPLE) guard needed

Let me know what you think!

[imperatormk]

@plotfi I got an idea; any chance this code can be brought to a separate branch so someone can still use it and report potential issues before deciding to merge into main? I can track the main branch regularly given this repo is under active development. Let me know!

[plotfi]

@imperatormk @CRobeck @abrown What do yall think about making a dir or a branch for Darco here?

I'd be for it, but I don't feel very easy with any external dependencies / git sub modules myself. Assuming everything is incorporated in one branch or directory thats fine by me.

[imperatormk]

No submodule/ext dependencies as per the latest code. However the LLVM metal-ir-pipeline part bloats up the project size significantly, you can decide to rehost it separately on your end in the future

[enjustli]

I think you could refer to the code structure under triton-lang/triton third_party/nvidia and adjust your file organization accordingly.

current

backend/
├── AppleGPU/
│   ├── include/
│   │   ├── Dialect/
│   │   ├── TritonAppleGPUToLLVM/
│   │   ├── TritonAppleGPUTransforms/
│   │   └── CMakeLists.txt
│   ├── lib/
│   ├── metal-ir-pipeline/
│   ├── python/
│   ├── CMakeLists.txt
│   ├── ExportAppleGPU.cpp
│   ├── README.md
│   └── triton-ext.conf
├── CMakeLists.txt
└── .gitignore

pseudocode

backend/
├── AppleGPU/
│   ├── backend/
│   │   ├── lib/
│   │   ├── __init__.py
│   │   ├── driver.c
│   │   ├── driver.py
│   │   └── compiler.py
│   ├── include/
│   │   ├── Dialect/
│   │   ├── TritonAppleGPUToLLVM/
│   │   ├── TritonAppleGPUTransforms/
│   │   └── CMakeLists.txt
│   ├── lib/
│   ├── metal-ir-pipeline
│   ├── language/apple 
│   ├── CMakeLists.txt
│   ├──  triton_apple.cpp 
│   ├── README.md
│   └── triton-ext.conf
└── CMakeLists.txt

i'm also an macbook user 😊

[imperatormk]

@enjustli Sure I'll attempt together with next changes related to tracking the latest code from main.

ps: Not sure about what M series you have at hand but help needed to test all this on M3+ because I fear there might be subtle differences of what their GPU ISA expects compared to the M1/M2 models that I have tested on!

[enjustli]

i use M2 apple silicon

[enjustli]

Maybe it could also be worth considering contributing this Apple backend implementation to the Triton ecosystem under the triton-lang organization, perhaps as a separate repository like triton-apple.

Maintaining it as an independent repository might make collaboration easier, especially since there seems to be growing interest in running LLM workloads on Apple hardware. I think there could be developers in the community who would be interested in contributing as well.

The process would probably not be easy, but I still feel the attempt itself could be meaningful. With the current progress of coding agents and LLM-assisted development, the barrier to contributing to compiler infrastructure is also becoming lower than before.

@imperatormk

[imperatormk]

Thank you @enjustli yes that's definitely a cleaner deal however the maintainers here had an idea to exercise this repo by appending a whole backend here and work on shaping it better. I think it IS possible to maintain it like that as some kind of subsystem, similar to how torch has inductor

[enjustli]

That said, I’m still a bit concerned about the current pace of progress in this repository. Since the LLM ecosystem is evolving very quickly right now, I feel that getting feedback and iteration cycles earlier and more frequently could also be quite important.

[imperatormk]

@plotfi @abrown 😄

[enjustli]

I found an issue discussing this problem already. From the issue history, it looks like three different people each made their own implementation.
triton-lang/triton#4824

[enjustli]

I think we should consider proposing an independent Triton-Metal backend under the triton-lang organization. Regardless of whether Apple eventually open-sources their own implementation, simply waiting probably isn't very meaningful.

Having a community-driven repository could help consolidate development efforts, avoid duplicated work, and make collaboration much easier. Right now, one of the biggest difficulties is that information and experiments are scattered across many separate projects, which makes the ecosystem harder to move forward efficiently.

@imperatormk

[imperatormk]

@enjustli See triton-lang/triton#9701 - they were open to making a separate repo so I suppose well can return to that option if they still like (this would be my preference at this point). There is triton-cpu already so this is a good adjacent to it. BTW I'd like to see this integrated to torch inductor eventually (see above that I managed to pull off a good POC for that) so I don't know which approach would facilitate it best

[enjustli]

OK

[enjustli]

If it becomes a separate repository, one possible approach could be to package Triton-Metal as a standalone wheel and distribute it through pip.

From the user side, the experience could remain very simple — after installation, users would just use import triton as usual, without requiring any additional initialization or backend-specific APIs.

This is similar to how the current triton-shared project works, where the entire backend is packaged and distributed as a complete wheel. https://github.com/facebookincubator/triton-shared/blob/b32f76945570b198279226cf9e89f9ee7be58da7/README.md?plain=1#L27-L41

[enjustli]

My English is bad. so I mostly rely on AI-assisted translation. Please excuse me if some of my wording or phrasing sounds unnatural.

[imperatormk]

Sounds like a good plan, let me try to get ahold of the devs here. And I'd trade my English for your Chinese any day of the week 💯

[abrown]

@imperatormk, ok, I'm back looking at this PR! Apologies for not reviewing this sooner but I've been pulled in a few other directions lately and I had put this PR in a mental bucket of "nothing needed from me." I did reiterate your request for access to the Triton Slack workspace; if we get that going, that will be a good way to have discussions on this kind of thing. Thanks for this contribution and being patient!

So, a high-level review:

• source location: I want to agree with @plotfi's idea up above to add a directory for this work to live in triton-ext; this seems experimental in nature but that shouldn't prevent it from being upstreamed, like TLX; however, if you agree with @enjustli to put this somewhere else, that is also fine because of...
• backend support: when I originally read through this PR (long, long ago!) I remember thinking that this needs the one thing that Triton extensions do not yet have — out-of-tree backends. We don't yet have a way to register the compiler.py/driver.py files with Triton so that the Python integration "just works" and it appeared your PR would need that. Now, looking at this again, I see that you are asking users to manually install the Python bits; I guess that is fine for now but if you wanted full backend support right now an in-tree plugin is the right approach. If you're fine waiting, we would wire things up when backend support is available. This question does not block merging the PR.
• testing: one early concern some of us had when discussing this on Slack was the testing story. I now see that you have lit tests (nice!) but it would be ideal if we could actually run this thing in CI... how, though?! I'm not sure if the macos runners GitHub makes available will be useful for this; what do you think about this? In any case, I'm willing to work through this over time if we can figure out...
• maintenance: adding 23k lines is a big ask not only to review now but to maintain green over time. What are your thoughts on maintaining this? This would be a good topic to discuss on Slack, but the TLDR is we need to find a balance between moving fast but also not slowing down others who may want to work on their upstreamed extensions (e.g., if this extension were to break CI, are you interested in troubleshooting the breakage?).
• confusion: another thing that contributed to me filing this PR away was a confusion over this and Add minimal Apple GPU backend #64. What is the difference? Why should we pick one over the other? Perhaps this was mentioned somewhere but I just lost context.

My sense at the moment is that if we talk through the testing and maintenance questions, this PR should be good to upstream. I'm no Apple GPU expert, though, so @plotfi may have some more review comments re: that interface and it looks like @enjustli had some thoughts about the directory layout. We can probably work those out, though, even in separate PRs, if we figure out a solid plan for the maintenance and testing angles.

[imperatormk]

@abrown No problem thank you for the response! Just to mention that in the meanwhile just today a new repo has been created for this (currently empty). Let me push a clean version there now and we can try to coin out a CI pipeline (so we get reliable testing scheme as you mentioned) + other specifics without interfering with this codebase for the time being, while trying to keep it faithful so it can be brought in here at some point as per a decision of yours.

Regarding maintenance I have no problem being around for anything that pops up and will be happy to help, however I meant to ask earlier too what is the vision for that angle for this repo in general, if a fault in one extension can affect the status of the whole collection or if say a change is made in the build system to accommodate a new extension and needs to be ensured it won't affect the existing ones, that sounds somewhat challenging. Does this confirm these 23k additions should be indeed kept in the new separate repo for the time being so as not to overload the development of triton-ext?

PS. the 'minimal' backend PR was made according to @plotfi 's request to start slow but from that point I'd have no idea how to continue building upon it given the current state.

Hope this addresses your points for now, eager to hear your opinion!

[abrown]

I am only a fraction of the way through looking at this but I thought I would leave a few comments (I got bogged down in llvm-metal-target--how necessary is that?). One high-level comment is the testing story; is there any way to emulate this?

[abrown]

😀

[imperatormk]

😁

[imperatormk]

I am only a fraction of the way through looking at this but I thought I would leave a few comments (I got bogged down in llvm-metal-target--how necessary is that?). One high-level comment is the testing story; is there any way to emulate this?

Yes the LLVM part is very needed, it is positioned as a counterpart of NVPTX for CUDA and then some, besides the common things like say BC serialization/AIR metadata attachment it also does lots of work to reconcile the IR emitted by the LLVM version that Triton uses to the version that the older LLVM that Xcode tooling uses will accept. Illustration:

• Tier 1: or disjoint, zext nneg, freeze usages are not valid so need to be stripped or canonicalized to an equivalent
• Tier 2: aggregate stores are not valid so

%buf = alloca [2 x <2 x float>]
store [2 x <2 x float>] %arr, ptr %buf      ; whole-aggregate store (illegal)
%ld  = load  [2 x <2 x float>], ptr %buf    ; whole-aggregate load  (illegal)

needs to become roughly

%p0 = getelementptr <2 x float>, ptr %buf, i64 0
store <2 x float> (extractvalue %arr, 0), ptr %p0
%p1 = getelementptr <2 x float>, ptr %buf, i64 1
store <2 x float> (extractvalue %arr, 1), ptr %p1
%e0 = load <2 x float>, ptr %p0
%e1 = load <2 x float>, ptr %p1
%ld = insertvalue [2 x <2 x float>] (insertvalue undef, %e0, 0), %e1, 1

• Tier 3: the AIR reader needs typed pointer records, so opaque ptr has to be reconstructed to a concrete pointee (up until macOS 26 I think and even there edge cases are found) so:

define void @k(ptr addrspace(1) %p) {
  store [2 x <2 x float>] %arr, ptr %p
}

needs to become

define void @k(float addrspace(1)* %p) {
  %c = bitcast [2 x <2 x float>]* %buf to <2 x float>*
  store <2 x float> %v, <2 x float>* %c
}

At scale this becomes far more delicate: pointers-to-i1 crash the JIT -> remap i8; MMA needs each device pointer pinned to its matrix element type otherwise the JIT rejects it -> keep int8/half/bf16 pointers at their element type.... the list goes on and I hope you see the problem here. I'm now working on that code to: 1. reorder everything so it's more traceable and obvious 2. strip away/simplify what is done excessively (fix in one place makes other parts moot). It doesn't help that when the AGX JIT rejects the BC it will give a short message at best and obscure 'invalid bitcode' at worst. So then it's a cycle of llvm-reduce -> fix the code up (after trying to figure out how to by lowering a .metal equivalent to AIR and see how native tooling does it) -> rerun a large test suite to make sure no side effects were introduced. I was hoping that if this manages to land in the 'raw' but best-effort comprehensive version as is now, it will become more sophisticated as it gets more users/contributors.

If you want to try some of that in isolation I can prepare some concrete runner or anything else that might be needed (needs GPU)

[plotfi]

@abrown If you've reviewed these changes, I have no problem with your approval. Sorry for being unavailable for review these days.

[robtaylor]

I am only a fraction of the way through looking at this but I thought I would leave a few comments (I got bogged down in llvm-metal-target--how necessary is that?). One high-level comment is the testing story; is there any way to emulate this?

I believe the macos-xlarge runners allow using the Apple GPU.

[robtaylor]

@imperatormk out of interest, will this backend work for gluon?

[imperatormk]

@robtaylor Should! No expected blockers; I meant to work on this but then focused on the core part because I wasn't sure it would be interesting to anyone initially. Now I'm working on a playground that will support multiple producers (I figure it can be useful for someone not necessarily wanting to go through a local build just to try it out a bit) and can try to add in Gluon there, will keep you posted!

[robtaylor]

thanks man, thats really exciting!

…

On Mon, Jul 6, 2026 at 10:13 AM Darko ***@***.***> wrote: *imperatormk* left a comment (triton-lang/triton-ext#48) <#48 (comment)> @robtaylor <https://github.com/robtaylor> Should! No expected blockers; I meant to work on this but then focused on the core part because I wasn't sure it would be interesting to anyone initially. Now I'm working on a playground that will support multiple producers (I figure it can be useful for someone not necessarily wanting to go through a local build just to try it out a bit) and can try to add in Gluon there, will keep you posted! — Reply to this email directly, view it on GitHub <#48?email_source=notifications&email_token=AAB4O4GO23P3GLDCG7MZGOD5DNUUPA5CNFSNUABFM5UWIORPF5TWS5BNNB2WEL2JONZXKZKDN5WW2ZLOOQXTIOBZGEYDKNZXGYYKM4TFMFZW63VHNVSW45DJN5XKKZLWMVXHJLDGN5XXIZLSL5RWY2LDNM#issuecomment-4891057760>, or unsubscribe <https://github.com/notifications/unsubscribe-auth/AAB4O4F5IMMLRL6A4J4SXJD5DNUUPAVCNFSNUABGKJSXA33TNF2G64TZHMYTCMJQGEZTQMBQGY5US43TOVSTWNBQHA2TINJSHAYDNILWAI> . You are receiving this because you were mentioned.Message ID: ***@***.***>

[imperatormk]

No problem, it worked out almost right off the shelf; available to test in the playground now (live url listed in https://github.com/imperatormk/applegpu-playground). Let me know how it goes!