[multi-gpu] Phases 4 + 6 + AIR-on-GPU realigned to PE=wavefront model#1618
Merged
Conversation
…lti-launch fixes Replaces the in-flight phase 4 PR (Xilinx#1579), reframing symmetric-heap allocation as a memref memory_space attribute rather than an op-attribute on memref.alloc. Foundation for the future channel-to-cacheline lowering (see docs/MultiGPUPhase56Redesign.md). Why memory_space (not op-attribute): - Travels with SSA values automatically; no need to trace defining-op chain when checking from a memref.load/store site. - The AIR herd verifier can reject/accept based purely on the memref type, not on op metadata of a possibly-distant alloc. - Lets memrefs constructed via memref.view / wrap_bytes / etc. (not just memref.alloc) carry the symmetric-heap tag — required for the in-flight cacheline tests where buffers are wrapped from runtime pointers. Changes: * `#air.symmetric_heap` custom memref memory_space attribute, defined via TableGen AttrDef in mlir/include/air/Dialect/AIR/AIROpBase.td. Dialect plumbing in AIROpBase.td (let useDefaultAttributePrinterParser = 1), CMakeLists.txt (mlir_tablegen for AIRAttrs.{h,cpp}.inc), AIRDialect.h (#include the generated header), AIRDialect.cpp (addAttributes<> in initialize()). * AIR herd verifier (verifyComputeMemoryAccess in AIRDialect.cpp) skips the L1-or-better-only check for memrefs whose memory_space is #air.symmetric_heap. Same for verifyAllocMemorySpace (segment-level alloc check). Direct memref.load/store on symmetric-heap memrefs inside air.herd bodies is now legal — required for the kernel-driven cross-rank cacheline pattern on GPU. * AIRSymmetricAllocToMgpu pass dispatches on the memref result type's memory_space, replacing the op-attribute check `op->hasAttr( "air.symmetric")`. FileCheck unit tests rewritten to use `memref.alloc() : memref<..., #air.symmetric_heap>` instead of `memref.alloc() {air.symmetric}`. All 6 cases pass. * AIR-to-ROCDL pass: two fixes uncovered by the air_hierarchy/ cacheline e2e baseline (kept in-tree as the target shape but currently blocked on a separate multi-launch use-after-free in this same pass — see docs/MultiGPUPhase56Redesign.md). 1. 1D / N-D launch + herd shapes are now handled (previously assumed 2D, would crash on getSizeOperands()[1] OOB for 1D). The new `sizeOrOne` helper materializes a `1` constant for missing dims. 2. Pattern set is now frozen once and reused across multiple launches (previously moved into applyPatternsGreedily on first iteration, FrozenRewritePatternSet ctor crashed on the second). * Plan doc (docs/MultiGPUPhase56Redesign.md) updated with the infrastructure landed today, what's still blocked (air-to-rocdl multi-launch use-after-free), and recommended next steps. * test/gpu/multi_gpu/air_hierarchy/cacheline.mlir + Makefile: the target IR shape that phase 6's redesigned lowering should emit. Kept as a draft baseline; will run e2e once the air-to-rocdl multi-launch issue is resolved separately. Tested: - lit unit test for AIRSymmetricAllocToMgpu: 6/6 pass - air_hierarchy e2e: blocked at air-to-rocdl multi-launch crash; the 1D and pattern-reuse fixes here unblock the first two errors but the third (use-after-free during block deletion) needs separate work Supersedes: - PR Xilinx#1579 (phase 4 op-attribute approach) — close as superseded Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
…wrong pairing, missing arg replace)
Multi-launch programs (e.g., scf.if %is_producer { air.launch P } else
{ air.launch C }) crashed AIRToROCDLPass with "Cannot destroy a value
that still has uses!" during block destruction. Three compounding bugs
identified via a minimal reproducer (two empty air.launch ops with
air.segment + air.herd):
1. blkIdx / gridIdx are class-level vectors that accumulated across
iterations of the launch walk. The second launch ended up using the
first launch's herd-size operands, dangling values into the new
gpu.launch op. Fixed by clearing them at the top of each (launch,
segment) pair (along with gridXVal / gridYVal).
2. The post-conversion body-move loop used a *nested* walk:
module.walk(gpuLaunchOp) { module.walk(airLaunchOp) { move } }
which pairwise-matched every gpu.launch with every air.launch in the
module, folding multi-launch programs into the first gpu.launch and
leaving the others empty. Fixed by performing the body move *inside*
the first walk where the 1:1 pairing is established.
3. The body-move did not replace the air.launch's own block args
(kernel operands) with their outer values before moving the body.
After air.launch was later erased, the moved ops dangled into the
destroyed block args (the actual use-after-free). Fixed with the
same replaceAllUsesWith pattern that deleteAirHerd /
deleteAirSegment use.
The 1D-launch and pattern-reuse fixes from the prior commit on this
branch are preserved.
Tested:
- Minimal reproducer (/tmp/multi_launch_repro.mlir): two air.launches
with memref.store inside each herd. Now lowers cleanly to two
distinct gpu.launch ops with correct operand routing.
- air_hierarchy/cacheline.mlir: now progresses past air-to-rocdl all
the way through to a GPU binary (after sed-stripping
#air.symmetric_heap from the pre-mlir-opt IR; see Makefile + plan
doc). Runtime hang remains; likely the herd's tile-id → block-dim
mapping doesn't preserve the cooperative cache-line semantics the
handwritten kernel relies on. Documented as next step in
docs/MultiGPUPhase56Redesign.md.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
…2.3 doc The air_hierarchy/cacheline.mlir test was hanging at runtime because the herd was declared (1,1) PE but the body used gpu.thread_id x to address 32 lanes — incompatible with the standard AIR→GPU mapping where the herd iteration space becomes the GPU blockDim (PE → thread). Per AIRComputeModel.md §4.1/§4.5/§5 — and the convention used by test/gpu/4k_4k_mul/air_sync.mlir — a 1-PE herd lowers to 1 GPU thread. Rewrite producer + consumer herds as (64, 1) so blockDim = (64, 1, 1) (one full MI3xx wavefront), and use the herd tile id %tx in place of gpu.thread_id x. Lanes 0..31 do real work, lanes 32..63 stay idle but keep the wavefront full so the consumer's gpu.shuffle width=64 can see the producer's flag lane. Also fix the §2.3 statement that claimed "PE instances mapped to individual warps" on GPU — that contradicted the rest of the doc and the actual lowering in air-to-rocdl, which is PE→thread. Replace with a pointer to §4.1 for the precise mapping, plus guidance that any herd needing wave-cooperative ops (e.g. gpu.shuffle) should pick a herd size that's a multiple of the target wavefront width. Verified on MI350X (gfx950, 8 GPUs): producer + consumer both report PASS, with data[0]=100 and flag=1 transmitted across the symmetric heap. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Adds the redesigned phase-6 pass per docs/MultiGPUPhase56Redesign.md.
Lowers air.channel.put/get ops on channels of type "gpu_symmetric_heap"
into the kernel-driven cacheline pattern that
test/gpu/multi_gpu/handwritten/cacheline.mlir writes by hand:
put -> air.translate %src, %from, %to, %bases + cooperative memref.store
(lanes 0..30 publish payload, lane 31 stores sync flag = 1)
get -> scf.while spin loop + gpu.shuffle idx broadcast of lane 31's
observation until the flag arrives; sink store keeps the spin
alive past DCE in subsequent passes
The pass:
- Pairs puts/gets via existing air::getTheOtherChannelOpThroughSymbol util
- Infers producer/consumer ranks from enclosing scf.if (cmpi eq %rid, %k)
rank-dispatch context
- Finds %bases by type-matching memref<?xindex, #air.symmetric_heap> in
the put/get's enclosing herd kernel args (semantically-unique combo
for the symmetric-heap base table; clear error on 0 or >1 matches)
- Errors loudly if any precondition isn't met (herd scope, rank
dispatch, bases arg, memref shape)
- Erases the channel symbol after expansion
Initial scope: 1-put / 1-get cacheline pattern with
memref<32xi32, #air.symmetric_heap> source/destination. Allgather +
multi-wire topologies are deferred.
Verified end-to-end on MI350X (gfx950, 8 GPUs): produces
"=== ALL 2 RANKS PASSED ===" with data[0]=100 and flag=1 published
cross-rank, functionally equivalent to the handwritten reference and
to air_hierarchy/cacheline.mlir.
Files added:
- mlir/include/air/Conversion/AIRGpuChannelToCachelinePass.h
- mlir/lib/Conversion/AIRGpuChannelToCachelinePass.cpp
- mlir/test/Conversion/AIRGpuChannelToCacheline/cacheline.mlir
- test/gpu/multi_gpu/air_channel/{cacheline.mlir,Makefile}
Wired into existing pass registration (GPUPasses.td, GPUPassDetail.h,
CMakeLists.txt, Passes.cpp).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Switch the cacheline spin loop in all three multi_gpu cacheline tests
(handwritten/, air_hierarchy/, air_channel/) and in the
air-gpu-channel-to-cacheline pass's get-side expansion from
%final_v = scf.while (%dummy = %c0) : (i32) -> i32 {
%v = scf.if %active -> i32 { memref.load %dst[%tx] : ... } ...
scf.condition(%not_ready) %v : i32
} do { ^bb0(%vi : i32): scf.yield %vi : i32 }
scf.if %active { memref.store %final_v, %dst[%tx] : ... } // sink
to the upstream-idiomatic shape from
mlir/test/Integration/GPU/CUDA/concurrent-kernels.mlir:
scf.while : () -> () {
%v = scf.if %active -> i32 {
%loaded = memref.atomic_rmw addi %c0_i32, %dst[%tx]
: (i32, memref<32xi32, ...>) -> i32
scf.yield %loaded : i32
} else { scf.yield %c0_i32 : i32 }
%flag, _ = gpu.shuffle idx %v, %c31, %c64 : i32
%not_ready = arith.cmpi ne, %flag, %c1 : i32
scf.condition(%not_ready)
} do { scf.yield }
`memref.atomic_rmw addi %c0` is functionally a load (adds 0, returns the
prior value) but carries both Read and Write effects in its
MemoryEffectOpInterface. This:
1. Survives the DCE inside air-to-rocdl's `applyPatternsGreedily`. With
the previous plain memref.load, MLIR's `wouldOpBeTriviallyDead`
considered the spin's body to have only Read effects and the
greedy driver killed the entire scf.while (verified in
GreedyPatternRewriteDriver.cpp:483-490 + SideEffectInterfaces.cpp:84-95).
The previous workaround was a sink store after the loop to make
%final_v used; this commit removes the workaround.
2. Encodes "this read must be observable across producers" as an
IR-level fact instead of relying on a plain memref.load happening to
stay observable through the lowering chain.
3. Lets us drop the iter-arg/result plumbing (%dummy, %final_v) and the
trailing scf.if/store sink — the spin loop becomes a self-contained
zero-result scf.while.
Verified on MI350X (gfx950, 8 GPUs): all three tiers
(handwritten/cacheline, air_hierarchy/cacheline, air_channel/cacheline)
still produce "=== ALL 2 RANKS PASSED ===" with data[0]=100 and flag=1
cross-rank.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
…gfx950
All four multi_gpu test Makefiles (handwritten, air_rank, air_hierarchy,
air_channel) previously hardcoded the AMDGPU chip target in their
mlir-opt pass pipeline (some at gfx942, the newer ones at gfx950),
which meant running them on a node with a different MI3xx variant
silently produced an unloadable binary (hipErrorNoBinaryForGpu) at
runtime — and the tests reported MISMATCH instead of a clean
precondition failure.
Replace the hardcoded chip with a GFX_TARGET make variable, auto-
detected in two layers:
1. amdgpu-arch (ROCm's clang tool; ships at /opt/rocm/llvm/bin/).
Most direct — outputs the gfx name for the visible GPU(s) verbatim.
2. Fallback: parse /sys/class/kfd/kfd/topology/nodes/*/properties for
the first GPU node's `gfx_target_version`, which encodes the chip
as major*10000 + minor*100 + step. Decode to `gfx<major><minor><step>`
(e.g., 90500 -> gfx950, 90402 -> gfx942). Holds for the MI3xx family
these tests target. Doesn't depend on ROCm tooling, only the kernel
module exposing the topology files.
Users can override via `make GFX_TARGET=gfx<NNN>` for cross-compilation
or on heterogeneous systems. check-preconditions now fails loudly with
a clear message if neither detection layer produces a value.
Verified on MI350X (gfx950, 8 GPUs): all four e2e tests detect gfx950
correctly and produce "=== ALL 2 RANKS PASSED ===". Manual override
test (`make GFX_TARGET=gfx906 ...`) confirmed the variable propagates
through the pipeline pass arguments.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
erwei-xilinx
force-pushed
the
multigpu-air-hierarchy
branch
from
de86582 to
25560b2
Compare
…ng scratchpad This file was an internal planning document — in-flight PR statuses, "drop phase 5", "what is currently wrong", a TL;DR aimed at the author rather than the project. It got included in commit 4a6cd86 by mistake and should not have landed. The actual normative documentation lives in docs/AIRComputeModel.md (GPU §4 explains the AIR-on-GPU mapping that this PR's phase 6 pass lowers into). PR descriptions cover the rationale for the redesign. Removal only, no behavior change. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
…efront)
The AIR-on-GPU lowering had drifted to PE → thread over time, contradicting
docs/AIRComputeModel.md §2.3 which states the herd PE maps to a wavefront
(not a thread). The implementation drift, plus an earlier silent "fix" of
§2.3 to match the drift, made the model inconsistent across §2.3 vs §4 vs
the air-to-rocdl lowering vs the GPU tests.
This commit restores PE → wavefront as the canonical model and aligns every
GPU-side artifact to it:
Doc (docs/AIRComputeModel.md):
* §2.3: restored original PE → wavefront wording; added pointer to
gpu.lane_id for lane access inside the herd body.
* §4.1: rewrote the mapping table — blockDim = (Nx * wave_size, Ny, 1),
tile_x → thread_id_x / wave_size, lane within PE → gpu.lane_id.
Clarified that the GPU kernel boundary is at the (innermost) segment,
not the herd; herd is an inline parallel block inside the kernel.
* §4.3: clarified L1 (private) semantics under PE = wavefront.
* §4.5: rewrote the matmul example to herd (4,1) — 4 PEs × 64 lanes =
256 effective work items — fitting the §2.3 wavefront-slot budget.
* §5: summary table: "air.herd tile" → "Single GPU wavefront",
"L1" → "Per-PE (per-warp) VGPRs / private scratch".
Pass (mlir/lib/Conversion/AIRToROCDLPass.cpp,
mlir/include/air/Conversion/GPUPasses.td):
* Added `wave-size` option (default 64; configurable for other arches).
* convertLaunchToGPULaunch: blockDim.x = herd.Nx * wave_size.
* deleteAirHerd: herd block args remapped to warp-id within block:
tile_x → thread_id_x / wave_size
tile_y → thread_id_y
size_x → block_dim_x / wave_size
size_y → block_dim_y
Pass (mlir/lib/Conversion/AIRGpuChannelToCachelinePass.cpp):
* put / get expansions emit `gpu.lane_id` for the per-lane index
instead of using the herd's tile_x (which is now a warp id, not a
thread id).
* Updated FileCheck unit test to expect gpu.lane_id.
Tests:
* test/gpu/multi_gpu/air_hierarchy/cacheline.mlir:
herd (64, 1) → herd (1, 1) — one PE = one wave; lane index from
gpu.lane_id (lanes 0..30 publish payload, lane 31 publishes flag).
* test/gpu/multi_gpu/air_channel/cacheline.mlir: same.
* test/gpu/4k_4k_mul/air_sync.mlir matmul:
herd (256, 1) → herd (4, 1) with %gtid = %tx * 64 + gpu.lane_id;
same 256-thread workgroup, expressed as 4 wavefronts × 64 lanes.
Honors §2.3's ≤32 PE per herd budget.
* test/gpu/4k_4k_mul/run.sh: GFX_TARGET auto-detected (amdgpu-arch
primary, KFD-topology fallback). Fixed a pre-existing bug where
LLVM_LIB_DIR was set to LLVM_INSTALL_DIR (missing /lib) when
LLVM_INSTALL_DIR was set in the environment.
Verified on MI350X (gfx950, 8 GPUs):
* FileCheck mlir/test/Conversion/AIRGpuChannelToCacheline/cacheline.mlir
* multi_gpu/handwritten/cacheline → === ALL 2 RANKS PASSED ===
* multi_gpu/air_hierarchy/cacheline → === ALL 2 RANKS PASSED ===
* multi_gpu/air_channel/cacheline → === ALL 2 RANKS PASSED ===
* 4k_4k_mul matmul → "Output Matched!"
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
§2.3 ("Platform-specific iteration space semantics") is about herd shape
and placement budgets — not how to access state inside a PE. The
gpu.lane_id / gpu.shuffle sentence I added in the previous commit (about
PE = wavefront) belonged in §4, not §2.3; §4.1 already carries that
guidance in the mapping table. Removing the bolted-on sentence restores
§2.3 to its original wording (three claims: herd ⊆ CU, PE → warp,
≤32 PE budget) with no dialect-op leak.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Previous commits expanded §4 with new intro paragraphs, an extra
mapping-table row, an L1 description rewrite, and matmul-example
elaboration — all of which were elaboration beyond what reverting
§2.3 strictly required.
Restore everything to the original wording except the cells that are
*literally wrong* under PE = wavefront:
§4.1 mapping table:
- "blockDim = (bx, by, 1)" → "blockDim = (bx * wave_size, by, 1),
with each PE materialised as one wavefront"
- "Herd tile index → (threadIdx.x, threadIdx.y)" → "warp-id within
block: (threadIdx.x / wave_size, threadIdx.y)"
§4.1 prose: "number of threads per block" → "number of wavefronts per
block"; "per-thread air.herd body" → "per-PE air.herd body". Two-word
tweaks; no new paragraphs.
§4.5 matmul example:
- herd (256, 1) → herd (4, 1) (256 PEs violated §2.3's ≤32 budget;
4 PEs × 64 lanes = same 256 work items)
- mapping summary updated to match
§5 summary table:
- "air.herd tile → Single GPU thread" → "Single GPU wavefront"
§4.3 (memory space mapping) reverted entirely — original "Per-thread
private (VGPRs/scratch)" is technically still accurate under PE = warp
(MLIR's "private" space is per-thread regardless of how a PE is
defined).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
erwei-xilinx
requested review from
eddierichter-amd,
fifield and
jgmelber
as code owners
erwei-xilinx
changed the title
Contributor
There was a problem hiding this comment.
Pull request overview
This PR advances the multi-GPU stack by introducing a first-class #air.symmetric_heap memref memory_space attribute, adding/reshaping GPU-oriented lowering passes (symmetric alloc lowering + channel put/get → cacheline pattern), updating the AIR→ROCDL lowering to follow a wavefront-based compute model, and expanding end-to-end GPU test coverage (including more robust GFX target selection in harnesses).
Changes:
- Add
#air.symmetric_heapas an AIR dialect attribute and use it to drive multi-GPU symmetric-heap behavior (verifiers + alloc lowering). - Introduce
air-symmetric-alloc-to-mgpuandair-gpu-channel-to-cachelineconversion passes plus FileCheck coverage. - Update
air-to-rocdllowering to treatair.herdtiles as wavefronts (warp-id mapping + blockDim scaling) and refresh GPU tests/Makefiles accordingly (incl. auto-detectGFX_TARGET).
Reviewed changes
Copilot reviewed 25 out of 25 changed files in this pull request and generated 8 comments.
Show a summary per file
| File | Description |
|---|---|
| test/gpu/multi_gpu/handwritten/Makefile | Auto-detect GFX_TARGET; plumb into ROCm codegen pipeline. |
| test/gpu/multi_gpu/handwritten/cacheline.mlir | Switch consumer spin loop to zero-result scf.while + memref.atomic_rmw. |
| test/gpu/multi_gpu/air_rank/Makefile | Auto-detect GFX_TARGET; plumb into ROCm codegen pipeline. |
| test/gpu/multi_gpu/air_hierarchy/Makefile | New e2e harness for AIR hierarchy variant; uses GFX_TARGET auto-detect and tag-stripping. |
| test/gpu/multi_gpu/air_hierarchy/cacheline.mlir | New AIR-hierarchy “handwritten-equivalent” cacheline test. |
| test/gpu/multi_gpu/air_channel/Makefile | New e2e harness for channel-based variant; includes air-gpu-channel-to-cacheline. |
| test/gpu/multi_gpu/air_channel/cacheline.mlir | New channel.put/get cacheline test intended to lower to the handwritten pattern. |
| test/gpu/4k_4k_mul/run.sh | Auto-detect GFX_TARGET; improve LLVM/AIR lib dir detection. |
| test/gpu/4k_4k_mul/air_sync.mlir | Update herd sizing to wavefront granularity using gpu.lane_id. |
| mlir/test/Conversion/AIRSymmetricAllocToMgpu/symmetric_alloc.mlir | New FileCheck coverage for symmetric alloc lowering. |
| mlir/test/Conversion/AIRGpuChannelToCacheline/cacheline.mlir | New FileCheck coverage for channel→cacheline expansion. |
| mlir/lib/Dialect/AIR/IR/AIRDialect.cpp | Register AIR attrs; exempt symmetric-heap memrefs from NPU memory-space checks. |
| mlir/lib/Conversion/Passes.cpp | Wire new GPU passes into registration when AIR_ENABLE_GPU. |
| mlir/lib/Conversion/CMakeLists.txt | Build new conversion pass implementations. |
| mlir/lib/Conversion/AIRToROCDLPass.cpp | Fix multi-launch correctness, add 1D support, reuse frozen patterns, and remap herd tiles to wavefronts (wave-size). |
| mlir/lib/Conversion/AIRSymmetricAllocToMgpuPass.cpp | New pass lowering memref.alloc/dealloc tagged #air.symmetric_heap to mgpu runtime calls. |
| mlir/lib/Conversion/AIRGpuChannelToCachelinePass.cpp | New pass expanding gpu_symmetric_heap channel put/get into cacheline pattern. |
| mlir/include/air/Dialect/AIR/CMakeLists.txt | TableGen for AIR attribute defs. |
| mlir/include/air/Dialect/AIR/AIROpBase.td | Define #air.symmetric_heap attribute as memref memory_space marker. |
| mlir/include/air/Dialect/AIR/AIRDialect.h | Include generated AIRAttrs defs. |
| mlir/include/air/Conversion/GPUPasses.td | Add new pass defs + wave-size option for air-to-rocdl. |
| mlir/include/air/Conversion/GPUPassDetail.h | Add SCF include; add GEN_PASS defs for new GPU passes. |
| mlir/include/air/Conversion/AIRSymmetricAllocToMgpuPass.h | New pass header. |
| mlir/include/air/Conversion/AIRGpuChannelToCachelinePass.h | New pass header. |
| docs/AIRComputeModel.md | Update GPU mapping text to wavefront-based herd semantics. |
💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.
…ANY RUN line The second RUN line in AIRGpuChannelToCacheline/cacheline.mlir invoked FileCheck with --check-prefix=ANY but the file has no ANY: patterns, so FileCheck errors with "no check strings found with prefix 'ANY'". Drop the RUN line; the remaining (default-prefix) RUN line covers the test. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
…) null in convertLaunchToGPULaunch `blkIdx[i]` is allowed to be a BlockArgument — happens when an air.herd/air.launch size operand is passed in as an SSA value rather than declared as a constant in scope. In that case getDefiningOp() returns null and the unconditional `blockXValOp->moveBefore(launchOp)` would crash with a null deref. Add the obvious null guard: only move when there's a defining op (a BlockArgument already dominates launchOp from the enclosing scope, so no move is needed). Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
…air-symmetric-alloc-to-mgpu buildMemrefDescriptor hard-codes row-major strides from the shape and offset=0; it can't faithfully represent strided / affine layouts. Without an explicit check, a non-identity layout (`memref<NxT, strided<...>, #air.symmetric_heap>`) would be silently miscompiled — the runtime allocation would be sized correctly but the descriptor strides would not match the user-declared layout. Add an `isIdentity()` precondition check in the AllocOp loop, before calling buildMemrefDescriptor. Emit a clear pass-failure diagnostic naming the requirement. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
…n GPUPasses.td dependentDialects ConvertAIRToROCDL's dependentDialects mixed `gpu::GPUDialect` / `LLVM::LLVMDialect` (no `mlir::` prefix) with `mlir::arith::ArithDialect` (with prefix). The rest of this file (and the repo's convention) uses the bare style. Drop the `mlir::` prefix from the arith entry for consistency. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
…cacheline (rank inference, bases shape, wave-size) Three related fixes to the channel-to-cacheline pass surfaced by Copilot review on Xilinx#1618. 1. inferRankFromEnclosingIf (Copilot #2): the pass walked up scf.if chains and accepted the first `cmpi eq %v, %const` as a rank dispatch. It never checked %v derives from %rid of the enclosing air.rank, so any other `eq %v, %const` (e.g. a lane predicate `cmpi eq %lane, %c0` inside the herd body) would mis-infer "rank 0". Add a helper `isRankIdOfEnclosingAirRank(v)` and require %v be a rank-id block arg of an enclosing air.rank before extracting the constant. Also handle the (rid, const) and (const, rid) orderings symmetrically via a small lambda instead of two near-identical branches. 2. findUniqueBasesArg (Copilot #3): the heap_bases arg search filtered only on element type (index) + memory_space (#air.symmetric_heap) but missed `memrefTy.getRank() == 1`. A 2-D index memref with the symmetric_heap memory_space would silently match, and downstream air.translate creation would fail with a poor message. Add the rank check. 3. Wave-size option (Copilot Xilinx#4): the spin loop's gpu.shuffle hard- coded `width = 64`, which is wrong on wave32 targets (NVIDIA). The sister pass air-to-rocdl already has a `wave-size` option; mirror it here. Default 64 (matching air-to-rocdl); user must set both options consistently. Updated td description spells out the must- match-air-to-rocdl + must-be->=32 (cacheline memref shape) constraints. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
…hannelToCachelinePass.cpp
CI's clang-format-17 wants slightly different line breaks than what
my local clang-format produced:
- Break after `(Value rid,` in the lambda signature in
inferRankFromEnclosingIf (line was 1 char too long without wrap).
- Break after `gpu::ShuffleOp::create(` instead of after the second
arg in expandGetToCachelineSpin (after the waveWidthI32 rename
pushed total width over the limit).
Pure formatting; no semantic change.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
erwei-xilinx
added a commit
to erwei-xilinx/mlir-air-erwei
that referenced
this pull request
May 19, 2026
The default-path RUN line in opt_shim_dma_bds.mlir (`device=npu1`, empty `shim-dma-tile-sizes`) hit the 30s lit timeout on Assert builds after the merge of Xilinx#1618. The combined time of the three RUN lines (default + NPUTILED + AIE1) over 18 functions exceeded the budget. The default path now invokes `tilePerfectlyNested` with tile size 1 per loop level, which fully unrolls shim loop nests. The pre-PR detailed CHECK lines (asserting one fully-folded `air.channel.put` per put) no longer match this output, and the smoke-style CHECKs they were replaced with provided no real verification. Move the default RUN into a sibling `opt_shim_dma_bds_default.mlir` (same IR body) so each file gets its own 30s lit budget. Replace smoke CHECKs with per-function op counts (`CHECK-COUNT-N: air.channel` or `air.dma_memcpy_nd`) and `CHECK-NOT: scf.for` to validate the unrolling invariant. NPUTILED + AIE1 RUNs stay in the main file. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
fifield
pushed a commit
to fifield/mlir-air
that referenced
this pull request
May 21, 2026
* [Transform] Auto-derive shim DMA tile size from BD-queue cost model
`runtime_loop_tiling_sizes` (CLI flag `--air-runtime-loop-tiling-sizes`,
pass option `shim-dma-tile-sizes`) is a user-supplied magic number used in
~40 example call sites. Choosing it wrong silently overflows the per-channel
shim DMA start-queue (XAIE_DMA_MAX_QUEUE_SIZE = 4) at runtime, causing
ERT_CMD_STATE_TIMEOUT. Choosing one value for a multi-launch module is also
impossible because each launch has different BD/iter counts — recently
demonstrated by the L-D drop-in into o_gemv_ffn where the module-wide
[16,16] hung the cascade launch's compile.
The compiler has all the inputs needed to pick the right tile per loop:
N = min(T, ⌊K / B⌋)
T = trip count of the surviving outer shim scf.for (already used)
K = XAIE_DMA_MAX_QUEUE_SIZE = 4 (aie-rt/.../xaie_dma.c:45)
B = max distinct BD configurations per shim channel per iter
(counted via the existing chansMappedToEquivalentBDs predicate
from AIRToAIESchedulingUtils.cpp::getRepeatCounts)
Verified against two reference configs:
- K=2048 matvec at user-set [16,16]: surviving loop fully absorbed by
wrap-and-stride → tile setting has no effect → unbounded
- LD cascade at user-set [2,2]: B=2 (alternating R + A_bulk patterns
on the same shim channel) → predicted ⌊4/2⌋ = 2 ✓
- Hardware sweep on cascade: tile in {1,2} runs; tile=4 compiles but
runtime-hangs with ERT_CMD_STATE_TIMEOUT — confirms K=4 binds at
runtime, not at compile time
- End-to-end o_gemv_ffn 7-launch with no runtime_loop_tiling_sizes set
anywhere: compiles, runs, correlation 0.9988 vs CPU reference
Surface:
- New pass option `auto-derive-tile-sizes` (default false); when true the
per-loop tile size is computed from the formula above. Per-loop user
override via `shim-dma-tile-sizes=N,M` still wins.
- aircc sets auto-derive=true by default; an explicit
`--air-runtime-loop-tiling-sizes=N,M` still overrides. Passing the flag
with no values disables tiling (preserves existing behavior).
- No IR attribute, no per-launch knob — the compiler decides per loop.
Replaces draft PR Xilinx#1616 (per-launch attribute). The attribute approach added
IR surface that rots; this version eliminates the surface entirely.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
* [Transform] Address PR review: hoist predicate, centralize queue depth, add aircc opt-out
Review-driven fixes on top of the auto-derive shim DMA tile size pass:
- Hoist chansMappedToEquivalentBDs from a private lambda inside
air::getRepeatCounts to a public helper in AIRToAIESchedulingUtils;
both the existing repeat-count canonicalization and the new BD-queue
cost model now call the single shared predicate.
- Centralize the per-channel shim DMA start-queue depth K=4 as
air::getShimDmaStartQueueDepth(targetModel), parameterized by the
AIE target model. AIETargetModel does not currently expose this
(only getNumBDs / getNumLocks / switchbox counts); TODO marker on
the helper for when an mlir-aie accessor lands.
- Hoist the mode check (user override vs auto-derive vs neither) out
of the per-loop tile-resolver loop; replace the prior break-on-empty
with an explicit assert so per-loop emptiness can't silently abandon
subsequent shim loops.
- Diagnostic when the cost model clamps tile to 1 (B > K): emit a
remark on the offending scf.for so queue-budget exhaustion is
visible in -mlir-print-ir-after-all and -v logs.
- aircc: refresh --air-runtime-loop-tiling-sizes help text (the old
"omit to disable tiling" is stale post-PR); add a
--no-air-auto-derive-tile-sizes opt-out so users can bisect default-
flip regressions without inventing tiling values; drop the dead
runtimeLoopTilingSizesPresent variable and unreachable "flag present
but empty" branch (cl::list<unsigned> rejects empty values at the
parser, so that arm cannot be reached from the CLI).
Test additions on top of opt_shim_dma_bds_auto_tile.mlir:
- Three RUN lines (auto, user override, default-off) sharing CHECK
prefix — wrap-and-stride folding converges all three on the same
final IR for these inputs, so this is a regression guard against
mode-specific divergence.
- Fourth RUN line capturing stderr with REMARK prefix; verifies the
B>K diagnostic fires on the new b5_above_queue_depth case.
- b5_above_queue_depth: B=5 > K=4, tile clamped to 1, 40 puts after
unroll.
- two_shim_loops_per_launch: two sibling shim for loops with
different B per loop — exercises per-loop tile resolution
(regression guard against the prior module-wide tiling behavior).
New mlir/test/aircc/runtime_loop_tiling.mlir: aircc smoke test
covering all three CLI surfaces — auto-derive default, explicit
user override (`--air-runtime-loop-tiling-sizes=2
--air-runtime-loop-tiling-sizes=2`, the correct list syntax used by
the Python backend), and the new --no-air-auto-derive-tile-sizes
opt-out. Each mode is verified to produce placed.*.mlir.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
* Sweep runtime_loop_tiling_sizes from tests, examples, and Python backend
Now that air-opt-shim-dma-bds derives the per-loop tile size from the
BD-queue cost model (auto-derive-tile-sizes=true, set by aircc when no
explicit --air-runtime-loop-tiling-sizes is given), the user-supplied
runtime_loop_tiling_sizes magic numbers are noise: they override a
working cost model with values that are at best redundant and at worst
the original footgun (a too-large module-wide tile overflows the
per-channel shim DMA start queue at runtime).
This commit removes the kwarg from:
* 53 xrt e2e tests
* 102 programming_examples scripts
* 4 dead `tiling = ...` locals in flash_attention examples that
were only used to feed the kwarg
* the XRTBackend / XRTRunner Python constructors and their
aircc forwarding path
Validated locally on NPU2 (amdhx370/Strix) post-removal — the cost
model now drives tile selection on every previously-overridden test:
* 33_triton_matmul_ver2: PASS
* 52_dma_pad_passthrough: PASS
* 53_matmul_padding_bf16 (M=N=500 K=784): PASS
* 54_matmul_padding_f32_bf16_emulation: PASS
This is the first time the cost model is exercised by the bulk of
the e2e suite — prior CI only validated the user-override path.
The aircc CLI flag `--air-runtime-loop-tiling-sizes` is kept as an
escape hatch (a tuning knob for cases where the cost model picks
badly), but the Python convenience parameter is removed — Python
callers wanting an override can drop to aircc directly.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
* Fix black formatting on two swept files
Removing runtime_loop_tiling_sizes from these two left residue that
black wanted to reflow:
- gen.py: single-element list `["L1",]` left on one line
- run.py: single-kwarg call collapsed onto one line
Caught by Lint and Format CI on the sweep commit.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
* Restore runtime_loop_tiling_sizes on 6 NPU2 tests that fail under cost model on CI
Strix CI (amdhx370) on the prior sweep commit showed 6 swept tests
producing all-zero output:
xrt/01_air_to_npu
xrt/11_gemm_bias_fusion
xrt/14_conv2d_i8_extern_vec
xrt/17_gemm_8x16_transform_vec_4x4
xrt/25_batch_matmul_bf16
xrt/40_triton_vec_add (bf16_emulation variant)
Local NPU2 (also Strix) passes the same tests with the cost model
driving. The CI environment uses pinned mlir-aie (b37dc33d4) +
llvm-aie (2026051501+f4933ef7) — both newer than my local — so the
interaction between the cost model's chosen tile and the CI
toolchain's DMA lowering is breaking the data path on those six
specific workloads.
Rather than block this PR while that's investigated, restore each
test's prior explicit runtime_loop_tiling_sizes and tag with a TODO
so the sweep can be reapplied once root cause is fixed. The bulk
~47 NPU2 swept tests continue to exercise the cost model in CI.
Also re-add the runtime_loop_tiling_sizes parameter to the Python
XRTBackend / XRTRunner constructors so these six tests can pass it
again. Updated docstrings to reflect post-PR semantics: empty list
means use the cost model.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
* [Transform] Cost model: short-circuit tile=1 when B<=1
The auto-derive cost model was emitting tile = floor(K/B) for every
surviving outer shim scf.for, including B=1 cases where the per-loop
BD count is already at the minimum. With B=1 the queue is not the
bottleneck, so tiling can only reorder the schedule (strip-mine +
unroll) without shrinking BD pressure.
On CI Strix HW (mlir-aie b37dc33d4 + llvm-aie 2026051501+f4933ef7),
the resulting schedule perturbation broke the data path on six
previously-overridden tests:
Diffing pass_045_after_air-opt-shim-dma-bds.mlir between explicit
runtime_loop_tiling_sizes=[1,1] and auto-derive on test 14
(conv2d_i8_extern_vec):
explicit: for(3) { for(8) { ... }} 24 segment dispatches
auto: for(3) { for(3) { for(8) {...}}} 72 segment dispatches
BD count on @channel_5: 18 in both
Tiling added a loop level without reducing BD count — pure
schedule perturbation. On the CI toolchain this produces all-zero
output; on my local (older llvm-aie 2026051001) it still works,
hence the CI-only failure on the prior commit.
Fix: in computeAutoTileSize, return {1} early when
maxDistinctPerChannel <= 1. The queue constraint N*B <= K is
trivially satisfied at tile=1 for B=1, and skipping the strip-mine
avoids the downstream perturbation.
This restores correctness on all 53 previously-overridden xrt tests
without the per-test runtime_loop_tiling_sizes=[N,M] revert from the
prior commit — those restores are now undone, putting the cost model
back in charge of every NPU2 e2e test in CI.
Lit test counts unchanged: B=1 inputs (b1_one_put_per_iter,
trip_clamp_b1_trip2, two_shim_loops_per_launch first loop) still
produce 1 surviving put because wrap-and-stride absorbs the
untouched scf.for. Comments updated.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
* Make cost-model auto-derive aircc opt-in; restore sweep
The two prior CI cycles showed that flipping the BD-queue cost model
to be aircc's default tile-size source has interactions with downstream
Strix/Phoenix HW lowering that I can't reproduce locally (CI uses
newer mlir-aie b37dc33d4 + llvm-aie 2026051501; my local pins are
older). Without local repro, deeper diagnosis is impractical in this
session.
Pragmatic resolution:
1. tools/aircc/aircc.cpp: rename `--no-air-auto-derive-tile-sizes`
(opt-out) to `--air-auto-derive-tile-sizes` (opt-in), default OFF.
Pre-PR behavior is now the aircc default — empty list means
"no shim DMA tiling", matching what main expected.
2. Cost-model code itself is unchanged from the original PR (no
B<=1 short-circuit) — the feature is still implemented and
invokable, just not on by default. The lit tests still exercise
it because they pass `auto-derive-tile-sizes=true` explicitly
to the pass.
3. Restored runtime_loop_tiling_sizes=[N,M] on all 53 xrt e2e tests
and 102 programming_examples by checking out the pre-sweep
versions from c2c85f8~1. The Python XRTBackend / XRTRunner
`runtime_loop_tiling_sizes` parameter remains in place (re-added
in the prior commit b0c8e94).
4. Updated mlir/test/aircc/runtime_loop_tiling.mlir smoke test to
cover the three real surfaces: default (no tiling), explicit
override, and the new --air-auto-derive-tile-sizes opt-in.
Net effect on the PR:
- The cost model exists as a feature (lit-tested, opt-in via aircc
CLI flag, opt-in via pass option).
- Default behavior matches main — no breakage.
- Sweep is undone; the magic-number values stay until a future PR
can drive a cost-model-default flip safely.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
* [Transform] Cost model returns tile=1 (BD-pool is binding, not start-queue)
Root cause for the prior CI cycles: the original `tile = floor(K/B)`
formula budgeted only against the per-channel start-queue depth K=4
(XAIE_DMA_MAX_QUEUE_SIZE). Empirical evidence on the CI-pinned
toolchain (mlir-aie b37dc33d4 + llvm-aie 2026051501+f4933ef7), built
locally and reproduced on real Strix NPU2 via test/xrt/14_conv2d_i8_
extern_vec:
tile=1 -> compile PASS + run PASS
tile=2 -> compile PASS, runtime output wrong
tile=4 -> compile FAIL: 'aiex.dma_configure_task' op
Allocator exhausted available buffer descriptor IDs
The binding constraint is the per-channel BD-descriptor pool, halved
again by ping-pong. K/B says "tile=4 is fine" on a B=1 loop, but the
unrolled body's distinct BDs (when they don't fold via wrap-and-
stride) exhaust the BD allocator. Without a reliable foldability
predictor, tile=1 is the only universally safe choice.
The cost model now returns a vector of 1s matching the perfectly-
nested loop depth. tile=1 strip-mines each level by 1 -- an iter-
count no-op -- but still invokes tilePerfectlyNested and the post-
tile fixup, which downstream lowering relies on (pre-PR no-tile
path skipped this and exhausted the BD allocator on test 14).
Verified locally on Strix amdhx370 with the CI-pinned mlir-aie +
llvm-aie. Tests that previously needed an explicit
`runtime_loop_tiling_sizes=[N,M]` now pass through the cost model
with no override:
* test/xrt/14_conv2d_i8_extern_vec : PASS
* test/xrt/33_triton_matmul_ver2 : PASS
* test/xrt/53_matmul_padding_bf16 : PASS
Changes:
* mlir/lib/Transform/AIRDependencyScheduleOpt.cpp:
computeAutoTileSize returns SmallVector<unsigned>(depth, 1).
Dead K-based logic + emitRemark removed.
* mlir/lib/Conversion/AIRToAIESchedulingUtils.cpp,
mlir/include/air/Conversion/AIRToAIESchedulingUtils.h:
drop getShimDmaStartQueueDepth (no longer used; useful future
re-add when a smarter cost model lands).
* tools/aircc/aircc.cpp: auto-derive is the aircc default again
(cl::init(true) on --air-auto-derive-tile-sizes), since the
cost model is now safe to drive every test.
* mlir/test/Transform/AIRDependencyScheduleOpt/opt_shim_dma_bds_
auto_tile.mlir: REMARK RUN removed (no diagnostic to emit when
tile is uniformly 1); user-override RUN switched to [1,1] to
match the cost-model output.
* Re-applied the sweep removing runtime_loop_tiling_sizes=[N,M]
from 53 xrt tests and 102 programming_examples (the values are
noise once the cost model is correct).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
* Fix black formatting on swept test/xrt/11_gemm_bias_fusion/gen.py
The sweep removed runtime_loop_tiling_sizes from a single-element
list `channel_multiplexing=["L1",]` that black wants expanded.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
* Fix black formatting on two more swept files
Black 26.5.1 (CI version) flags an extra blank line in these two
swept files after runtime_loop_tiling_sizes removal. Local black
25.1.0 was lenient about it; the newer CI version isn't.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
* Apply clang-format-17 to AIRDependencyScheduleOpt.cpp
Tiny re-flow on the computeAutoTileSize lambda signature that
clang-format-17 (CI version) wants.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
* Restore runtime_loop_tiling_sizes=[4,4] on ffn_swiglu/prefill
CI Strix amdhx370 segfaulted at runtime on this single test with the
cost model driving (tile=1). Reproduced locally 3/3 PASS with the
exact CI-pinned mlir-aie + llvm-aie + same Strix silicon family, so
this is a runner-specific environment issue I can't diagnose
remotely. Pin the explicit value back with a TODO for follow-up;
the other 154 swept tests / examples remain on the cost model.
27 of 28 CI checks pass; this restores the last one without
reverting the rest of the sweep.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
* Restore runtime_loop_tiling_sizes=[16,16] on llama32_1b GEMV preset
CI Strix amdhx370 timed out at 600s during `make compile` on
llama32_1b/run_npu2_makefile_peano_synthetic_verify after the sweep.
This multi-launch model has many GEMV launches; with cost-model
tile=1 the per-launch compile time scales poorly and the aggregate
exceeds the 10-min budget. The pre-PR value [16,16] gave acceptable
compile time.
Two of 156 swept tests now have explicit overrides
(ffn_swiglu/prefill, llama32_1b GEMV preset); the rest stay on the
cost model. Both are tagged with TODOs.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
* Drop auto-derive-tile-sizes option; just default empty tile to 1
The auto-derive-tile-sizes pass option + matching --air-auto-derive-tile-
sizes aircc CLI flag were scaffolding for a "smart" cost model that
doesn't exist. The cost model returns tile=1 unconditionally, so the
option only ever gated whether the pass falls back to tile=1 vs skips
tiling.
Collapsed to: empty `shim-dma-tile-sizes` defaults to tile=1 per level.
No new flag, no aircc CLI surface, no opt-in/opt-out toggle.
- Passes.td: drop clAutoDeriveTileSizes option
- AIRDependencyScheduleOpt.cpp: no flag check; empty clTileSizes
triggers the same per-level-1 vector that auto-derive returned
- aircc.cpp: drop --air-auto-derive-tile-sizes flag and the
conditional forwarding
- opt_shim_dma_bds_auto_tile.mlir: deleted (it tested the flag)
- aircc/runtime_loop_tiling.mlir: deleted (it tested the CLI flag)
- opt_shim_dma_bds.mlir: dropped the default-no-tile RUN line and
its CHECK lines; that behavior no longer exists. NPUTILED and
AIE1 RUNs unchanged.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
* Revert chansMappedToEquivalentBDs public hoist
The hoist was for the cost model to call the predicate; the cost
model no longer does any analysis (just returns tile=1), so the
only caller is back to air::getRepeatCounts itself. Move the
predicate back to a private lambda inside that function.
Net effect: AIRToAIESchedulingUtils.h diff vs main is zero on this
predicate.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
* Restore default-RUN line in opt_shim_dma_bds.mlir with smoke-test CHECKs
The deleted RUN exercised the no-tile-sizes default path, which is
now distinct from the explicit shim-dma-tile-sizes=2,2 path. Restored
with a new DEFAULT prefix and per-function smoke checks (LABEL +
"contains air.channel"). Doesn't validate exact IR shape, but proves
the pass produces output for each function with the new default.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
* Try removing both TODO restores; see if CI still needs them
If CI is green: the explicit values weren't needed (possibly transient
runner state when the original failures hit), and the cost model is
sufficient.
If CI fails: the TODOs were accurate; revert this commit.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
* Restore llama32_1b GEMV preset; split default-path test out
CI on b3b6b5a confirmed two TODOs were not just paranoid:
1. llama32_1b GEMV preset really does need runtime_loop_tiling_sizes=
[16, 16]. Without it, multi-launch compile times out at 600s.
Restored with a tightened TODO citing the CI evidence.
2. opt_shim_dma_bds.mlir default-RUN line I restored in 8379c02
timed out at 30s in CI Assert builds. The pass got slower with
tile=1-by-default (tilePerfectlyNested + fixup runs where the
pre-PR no-tile path skipped); on a 18-function input that added
~8s to lit and pushed past the 30s per-test timeout.
Compromise: keep the user-override RUN lines (NPUTILED + AIE1)
on the big file, move the default-path smoke check to a new
small file (opt_shim_dma_bds_default.mlir, 1 function). Both
complete well within the timeout.
ffn_swiglu/prefill remains on the cost-model default; CI's earlier
segfault wasn't re-exercised this round because llama32_1b's timeout
aborted the Strix run. Next CI cycle tells us.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
* Revert the sweep: restore original runtime_loop_tiling_sizes everywhere
The sweep removed runtime_loop_tiling_sizes=[N,M] from 153 callers on
the assumption that the cost model's tile=1 default could replace them.
The intervening CI cycles showed two examples needed to be restored
back to their original values for compile-time and runtime reasons.
For the rest, the values were performance-tuned, and tile=1 by default
likely regresses runtime perf even when it compiles and runs correctly.
This commit restores all 153 swept files to their pre-PR state by
checking out origin/main for test/xrt and programming_examples.
Side effect: my earlier TODO-comment restores on ffn_swiglu/prefill
and llama32_1b backend_presets become redundant (those files were
restored to main, which already contains the explicit values without
my TODO).
Net effect of the PR on test/example files:
- 3 black-format reflows on files black 26.5.1 (CI version) wants
to reformat (test/xrt/11_gemm_bias_fusion, test/xrt/18_matmul_*,
test/xrt/19_matmul_*); unrelated to runtime_loop_tiling_sizes.
- No removals of runtime_loop_tiling_sizes; performance-tuned
values stay where the authors put them.
The pass-level change (empty shim-dma-tile-sizes defaults to tile=1)
stays, so callers that DON'T set a value go through the cost model.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
* Drop now-dead AIRToAIESchedulingUtils.h include
Added when the cost model called chansMappedToEquivalentBDs +
getShimDmaStartQueueDepth from that header. Both are reverted (the
predicate is back as a private lambda; the queue-depth helper was
deleted). AIRDependencyScheduleOpt.cpp uses no symbols from that
header anymore.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
* Restore default-no-tile RUN in opt_shim_dma_bds.mlir; drop duplicate file
Restoring the RUN line that exercises the empty-shim-dma-tile-sizes
default path. Added per-function smoke CHECKs (LABEL + "contains
air.channel"), since the original detailed CHECKs matched old skip-
tiling output and don't apply to the new tile=1 default IR.
Deletes the now-duplicate opt_shim_dma_bds_default.mlir.
Risk: CI Assert build timed out at 30s on this configuration on a
prior commit (8379c02). If it times out again, follow-up commits
will split the file or trim functions to fit the budget.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
* [test] Add cascade lit test: minimal reproducer of the bug this PR fixes
Two distinct shim BD patterns per iter on one channel (R + A_bulk on
different memrefs, not collapsible via repeat_count). Under the previous
[16,16] preset this shape exhausted the per-tile BD allocator during
aircc lowering. With default tile=1 it compiles and emits 8*2 = 16
channel.puts on @cascade.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
* [test] Split default-path RUN into sibling file with op-count CHECKs
The default-path RUN line in opt_shim_dma_bds.mlir (`device=npu1`,
empty `shim-dma-tile-sizes`) hit the 30s lit timeout on Assert builds
after the merge of Xilinx#1618. The combined time of the three RUN lines
(default + NPUTILED + AIE1) over 18 functions exceeded the budget.
The default path now invokes `tilePerfectlyNested` with tile size 1
per loop level, which fully unrolls shim loop nests. The pre-PR
detailed CHECK lines (asserting one fully-folded `air.channel.put`
per put) no longer match this output, and the smoke-style CHECKs
they were replaced with provided no real verification.
Move the default RUN into a sibling `opt_shim_dma_bds_default.mlir`
(same IR body) so each file gets its own 30s lit budget. Replace
smoke CHECKs with per-function op counts (`CHECK-COUNT-N: air.channel`
or `air.dma_memcpy_nd`) and `CHECK-NOT: scf.for` to validate the
unrolling invariant. NPUTILED + AIE1 RUNs stay in the main file.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
* [test] Fix file header comment to match actual filename
Copilot review noted the cascade test's header line still referenced
the pre-rename `opt_shim_dma_bds_cascade_default.mlir`.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
* [test] Add multi-launch case to cascade reproducer
The original o_gemv_ffn failure shape was multiple heterogeneous
air.launch ops in one func — a cascade-style launch (B=2, wants small
tile) alongside absorbable launches (B=1, indifferent to tile). The
prior cascade test used only a single launch, missing the structural
condition that motivated the PR.
Add a second func with two launches in one module — one cascade, one
absorbable — and assert per-launch op counts to confirm the pass
handles each surviving shim loop independently.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
* [test] Rename default-path test, drop func3 (4096-dma unroll blew lit budget)
Rename opt_shim_dma_bds_default.mlir → opt_shim_dma_bds_empty_tile_sizes.mlir
to name the input condition under test (empty `shim-dma-tile-sizes`)
rather than the vague "default".
Drop func3: its 3-deep shim loop nest (32x8x4) wrapping an air.herd
with 4 air.dma_memcpy_nd ops unrolls to 4096 dmas under the empty-tile
path, which alone exceeds the 30s Assert-build lit budget. The other
17 functions exercise the same unrolling on air.channel.put/get in
<100ms total. NPUTILED + AIE1 RUNs in the sibling file still cover
func3 with bounded tile sizes.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
---------
Co-authored-by: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
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.
Summary
Consolidates and supersedes #1579 (Phase 4), #1580 (Phase 5), and #1581 (Phase 6). Builds on already-merged #1577 (Phase 2) and #1578 (Phase 3).
What this PR delivers
Phase 4 (redesigned) —
#air.symmetric_heapmemref memory_space attribute +AIRSymmetricAllocToMgpupass dispatching on the new attribute. Replaces #1579's{air.symmetric}op-attribute approach. The memref-type carries the tag, so verifiers /air.translate/ channel pass can answer "is this on the symmetric heap?" by inspecting the type — no defining-op chain walk, andmemref.view-derived memrefs also carry the tag.Phase 6 (redesigned) —
air-gpu-channel-to-cachelinepass (replacing*-to-mgpu). Walksair.channelops withchannel_type = "gpu_symmetric_heap", pairs put/get via existingair::getTheOtherChannelOpThroughSymbol, infers producer/consumer ranks from the enclosingscf.if (cmpi eq %rid, %const)rank-dispatch, finds%basesby type-matchingmemref<?xindex, #air.symmetric_heap>in the herd's kernel args, and expands:air.translate+ cooperativememref.store(flag at lane 31)scf.whilespin +memref.atomic_rmw addi 0+gpu.shuffle idxbroadcast of lane 31 (upstream idiom frommlir/test/Integration/GPU/CUDA/concurrent-kernels.mlir)AIR-on-GPU model alignment to AIRComputeModel.md §2.3 (PE = wavefront) — the existing
air-to-rocdllowering had drifted to PE → thread, contradicting §2.3's normative claim that "PE instances [are] mapped to individual warps". This PR restores PE → wavefront across:air-to-rocdl:blockDim.x = herd.Nx * wave_size(default 64, configurable via newwave-sizepass option); herd tile_x →threadIdx.x / wave_size(warp-id-within-block); the lane index inside a PE comes fromgpu.lane_id.air-gpu-channel-to-cacheline: per-lane index fromgpu.lane_id, not from the herd's tile_x.AIRComputeModel.md) to align cells that contradicted §2.3.4k_4k_mul/air_sync.mlirmatmul: rewritten fromherd in (256, 1)(violated §2.3's ≤32-PE budget) toherd in (4, 1)with%gtid = %tx * 64 + gpu.lane_id. Same 256-thread workgroup, expressed as 4 wavefronts × 64 lanes.herd (1, 1)(= one wavefront) +gpu.lane_id.Drive-bys (collateral cleanup unblocked by the above):
air-to-rocdlcorrectness fixes: 1D launch support (was hardcoded 2D),FrozenRewritePatternSetreuse, three multi-launch bugs.multi_gpu/Makefiles +4k_4k_mul/run.shauto-detectGFX_TARGET(amdgpu-arch primary, KFD-topology fallback). Drops hardcodedgfx942/gfx950. Also fixes a pre-existingLLVM_LIB_DIRderivation bug inrun.sh.Phase 5 dropped
#1580 (
air-cross-rank-dma-to-mgpu) closed without replacement.air.dma_memcpy_nddoesn't naturally express the producer/consumer asymmetry of the cacheline pattern; the channel-based phase 6 covers the same use case correctly.Deferred
memref<32xi32, #air.symmetric_heap>forchannel.put/get.Test plan
mlir/test/Conversion/AIRGpuChannelToCacheline/cacheline.mlirmlir/test/Conversion/AIRSymmetricAllocToMgpu/symmetric_alloc.mlir(6 cases)test/gpu/multi_gpu/handwritten/cacheline→=== ALL 2 RANKS PASSED ===test/gpu/multi_gpu/air_rank/cacheline→=== ALL 2 RANKS PASSED ===test/gpu/multi_gpu/air_hierarchy/cacheline→=== ALL 2 RANKS PASSED ===test/gpu/multi_gpu/air_channel/cacheline→=== ALL 2 RANKS PASSED ===(data[0]=100, flag=1)test/gpu/4k_4k_mul/air_sync.mlirmatmul →Output Matched!GFX_TARGET=gfx906override confirmed to propagate.🤖 Generated with Claude Code