Repository navigation
Conversation
Metal.jl emits these in their newest form and rewrote them for older AIR versions in its own `finish_ir!`. Like the barrier and atomic ABIs, that is a property of the AIR version being targeted, so do it here during AIR lowering: before AIR 2.8, the simdgroup matrix load/store calls get the legacy elements-per-row, origin and transpose operands, and before AIR 2.9 the operands AIR 2.9 appended to the tensor intrinsics are dropped. Calls already in the older form, as emitted by released versions of Metal.jl, are left alone. The element type of the simdgroup matrix pointer operand is now also derived here, for the downgrader.
Member
Author
|
Closing: the AIR legalization of intrinsics Metal.jl emits stays next to the emitter in Metal.jl, so the two packages don't have to release in lockstep, and GPUCompiler doesn't need to recognize every form an older Metal.jl may emit. |
Codecov Report❌ Patch coverage is
Additional details and impacted files@@ Coverage Diff @@
## main #1046 +/- ##
==========================================
- Coverage 92.34% 91.51% -0.83%
==========================================
Files 31 31
Lines 7341 7413 +72
==========================================
+ Hits 6779 6784 +5
- Misses 562 629 +67 ☔ View full report in Codecov by Harness. 🚀 New features to boost your workflow:
|
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
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.
Metal.jl emits the simdgroup matrix load/store and tensor intrinsics in their newest AIR form, and rewrote them for older AIR versions in its own
finish_ir!override. That rewrite depends only on the targeted AIR version, like the barrier and atomic legalization this package already does, and is the kind of thing an LLVM back-end handles for the other GPU targets. So this moves it intolower_air!, next tolegalize_barrier_abi!:air.simdgroup_matrix_8x8_{load,store}: before AIR 2.8, the dims/strides/origin vectors become the legacy elements-per-row, origin and transpose operands. The transpose flag has to be an immediate, so one of the strides must be the literal 1 (as Metal.jl emits them); other layouts are an error rather than an assertion.air.init_strided_private_tensor.i32.{global,local}andair.get_extent_private_tensor.i32: before AIR 2.9, the operands AIR 2.9 appended (alwaysnull, 0, -1and0in Apple's front-end) are dropped.Calls already in the older form are left alone. Released versions of Metal.jl still rewrite them in
finish_ir!, which runs beforelower_air!, so they keep working with this version. Older forms aren't upgraded for newer targets.Tests cover both families at the AIR versions on either side of the change, with constant and run-time operands, and that the legacy forms pass through. Validated on Julia 1.10 and 1.13. Metal.jl's simdgroup, linalg and bfloat16 tests pass with JuliaGPU/Metal.jl#1096, which removes its copy of the rewrites, including execution at
macos=v"15"(AIR 2.7) on macOS 27.