Repository navigation
fix(rocr/aie): key a full-ELF kernel's PDI placements by section, not content - #13908
Merged
Merged
Conversation
ypapadop-amd
requested review from
a team,
atgutier,
cfreeamd,
dayatsin-amd,
kentrussell and
shwetagkhatri
as code owners
October 21, 2026 14:20
github-actions
Bot
added
documentation
✅ All Policy Checks Passed
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
|
🚫 Please fix the failed policies before requesting reviews. The following policy checks failed:
The |
Contributor
There was a problem hiding this comment.
🟢 Approval recommended
The focused change preserves same-section sharing, includes targeted regression coverage, and has no unresolved findings.
No open issues 🎉
What changed in this PR
Fixes repeated full-ELF kernel dispatches by keeping distinct PDI sections in separate device placements, even when their bytes match.
Changes:
- Keys PDI entries by ELF section index while preserving same-section sharing.
- Adds regression coverage for identical-content sections.
- Updates documentation to describe placement semantics and the observed fix.
| File | Description |
|---|---|
| projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_aie_elf.cpp | Groups PDI placements by section index. |
| projects/rocr-runtime/runtime/hsa-runtime/core/inc/amd_aie_elf.h | Documents section-based placement and observed behavior. |
| projects/rocr-runtime/rocrtst/suites/aie/elf_parse_test.cc | Tests separate entries for identical-content sections. |
| projects/rocr-runtime/docs/aie-hsaco-format.md | Updates the documented PDI placement rules. |
Review effort: Balanced
ypapadop-amd
removed
documentation
ypapadop-amd
force-pushed
the
users/ypapadop-amd/fix-full-elf-pdi-dedup
branch
from
October 21, 2026 14:36
dd5ba9c to
80b49a8
Compare
github-actions
Bot
added
the
documentation
dayatsin-amd
approved these changes
Oct 21, 2026
cfreeamd
approved these changes
Oct 21, 2026
ypapadop-amd
force-pushed
the
users/ypapadop-amd/fix-full-elf-pdi-dedup
branch
from
October 21, 2026 22:44
80b49a8 to
a87909e
Compare
1 task done
… content The full-ELF reader deduplicated a kernel's PDI patch sites by comparing the referenced section's raw bytes: two sites whose sections held equal bytes were placed in device memory once and shared one address, however many sites named them. MLIR-AIR's control code loads an empty PDI both before and after the design runs, as two separate ELF sections with equal (empty-configuration) bytes. Under content-based dedup those two sections share one device placement. A kernel built that way dispatched correctly once and incorrectly on every dispatch after. Keying placements by section index instead of content fixes it: two sites naming the same section still share one placement (unchanged, and still covered by ParsesTwoPatchSitesForOnePdi), but two different sections are now placed separately even when their bytes are identical. Nothing in this reader explains the mechanism -- it only decides which bytes get a device placement and writes the address into the control code; it does not touch the firmware or driver layer that loads a PDI at that address. This is recorded as an observed fix, not a diagnosed root cause. Adds KeepsIdenticalPdiSectionsSeparate, which aliases a spare section onto .pdi.1's own bytes (same offset and size) so its content is identical by construction, then checks both sections get distinct Pdi entries and patch offsets. Fails without this change (k.pdis.size() == 1 instead of 2). Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
ypapadop-amd
force-pushed
the
users/ypapadop-amd/fix-full-elf-pdi-dedup
branch
from
October 23, 2026 17:09
a87909e to
f17d29c
Compare
ypapadop-amd
added
the
ci:platform:linux
erwei-xilinx
added a commit
that referenced
this pull request
Oct 25, 2026
… variables (#14004) ## Motivation A full ELF can relocate the address of a buffer it does not carry. The common case is the control scratchpad (`scratch-pad-ctrl`), which holds the runtime parameters that a design reads in its preamble. MLIR-AIR's LLM decode uses it to pass the context length to one ELF for every token. The AIE hsaco loader refused any such relocation ("unsupported relocation symbol"). So a design with scratchpad parameters could not be loaded. Applications worked around this: they allocated the scratchpad, wrote its device address into the ELF, pointed the relocation at symbol 0 so that the loader skipped it, and repacked the hsaco for each allocation (for example amd/Triton-XDNA#158). ## Technical Details A relocation whose symbol is neither a PDI nor an argument index is now an **external variable**. It is resolved the same way as a GPU code object's undefined symbols: the application allocates the buffer and defines it before the load. ```cpp hsa_amd_memory_pool_allocate(dev_pool, size, 0, &scratchpad); hsa_executable_agent_global_variable_define(exe, aie_agent, "scratch-pad-ctrl", scratchpad); hsa_executable_load_agent_code_object(exe, aie_agent, reader, nullptr, nullptr); ``` - **Parse** (`amd_aie_elf.cpp`): the parser records each external symbol's patch sites and `st_size` in `Kernel::external_symbols`. Only `Address64` is accepted, as for PDIs. A site must hold 8 bytes inside the control code. All the symbols that name one variable must have the same `st_size`. - **Load** (`executable.cpp`): the loader looks up each name in `agent_symbols_`, the table that `hsa_executable_agent_global_variable_define` fills. If the variable is not defined, the load fails with `HSA_STATUS_ERROR_VARIABLE_UNDEFINED`. If its `st_size` bytes are not inside one allocation that the AIE agent can reach, the load fails with `HSA_STATUS_ERROR_INVALID_ALLOCATION`. Otherwise the loader writes the variable's agent address plus the addend into the control code, next to the PDI addresses, and records its BO handle. - **Dispatch** (`amd_xdna_driver.cpp`): `BuildFullElfCommand` flushes each external variable from the CPU cache and adds its BO to the command. So the host can change a parameter between dispatches with an ordinary store, and the application does not flush. The packet still lists only the kernel's arguments. - **Docs**: `docs/aie-hsaco-format.md` sections 7.4, 7.5, 10 and 11, and a new section 9.3. No change to the hsaco format, the packer, or the packet ABI. Two points for reviewers: 1. The command processor copies the scratchpad itself. On aie2p it does so only from the device-SVM pool (dev heap). A scratchpad in the data pool loads, but the dispatch fails. This PR documents the requirement but does not enforce it, because enforcement needs an allocation-to-region lookup that `Runtime` does not have. Should the loader refuse the load? 2. The new test design follows MLIR-AIR's sequence: load the empty device, then the design, then `sync_parameters()`, run, and load the empty device again. Without the final empty load, aie2p firmware gives each dispatch the scratchpad value of the dispatch before it. XRT gives the same result for one ELF run through a new `xrt::run` for each dispatch, so this comes from the firmware, not from this change. It is noted in the design file. ## Issue Tracking Follow-up to #12846 (AIE hsaco support). Independent of #13908. ## Test Plan - `elf_parse_test`: 3 new tests on a real scratchpad artifact (`vector_scalar_add_param.py`, a new IRON design that is built only as a full ELF): - The scratchpad is parsed as an external symbol, with the correct sites and size. - Any scheme other than `Address64` is refused. - An out-of-range site is refused. - `test_dispatch`: new `ExternalVariableTest` cases on hardware: - The load refuses an undefined scratchpad, a scratchpad in host memory, and a scratchpad that runs past the end of its allocation or of a VMM mapping. - One loaded kernel is dispatched 5 times, and the host changes the parameter with a plain store before each dispatch. Each output is checked against that dispatch's value. - The full AIE suite, as a regression check. - `TryLoad`/`LoadAndFreeze` take an optional list of variables to define before the load. `add_aie_kernel`'s full-ELF half is split out as `add_aie_full_elf` so that a full-ELF-only design can use it. ## Test Result On an aie2p (Krackan) NPU with `hsa-runtime64` built from this branch: - `elf_parse_test`: 16/16 (13 before this PR). - `aie_code_test`: 12/12. - `test_dispatch`: 56/56, run several times. - `test_memory`: 40/40. - `test_error_callback`: 2/2. - `ExternalVariableTest*` and the existing `FullElf*` tests: 20 consecutive runs each, with no failures. During development, one run of `ExternalVariableTest*` hung, right after the artifacts were rebuilt. It did not happen again in the 30 runs after it. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/rocm-systems/blob/develop/CONTRIBUTING.md. 🤖 Generated with [Claude Code](https://claude.com/claude-code)
Contributor
TheRock Submodule Bump ActivityNewest first
|
1 task done
3 of 4 tasks
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.
Motivation
The full-ELF reader placed a kernel's PDI patch sites in device memory keyed
by byte content: two sites whose sections held equal bytes shared one device
placement and one address. MLIR-AIR's control code loads an empty PDI both
before and after the design runs, as two separate ELF sections holding equal
(empty-configuration) bytes. Under content-based dedup those two sections
collapsed to one placement. A kernel built that way dispatched correctly once
and incorrectly on every dispatch after.
Technical Details
Key a kernel's PDI placements by ELF section index instead of by content.
Two sites naming the same section still share one placement, unchanged
from before (
ParsesTwoPatchSitesForOnePdistill covers this). Twodifferent sections are now placed separately even when their bytes are
identical.
Nothing in this reader explains the mechanism — it only decides which bytes
get a device placement and writes the address into the control code; it
does not touch the firmware or driver layer that loads a PDI at that
address. This is recorded as an observed fix, not a diagnosed root cause
(see the updated
Kernel::pdisdoc comment).Also updates
docs/aie-hsaco-format.mdsection 7.4, which described the oldcontent-based dedup as the authoritative behavior.
Issue Tracking
Part of #12846
Test Plan
elf_parse_test: addedAieElfParse.KeepsIdenticalPdiSectionsSeparate,which aliases a spare section onto
.pdi.1's own bytes (same offset andsize, so the content match is by construction) and checks both sections
get distinct
Pdientries and patch offsets. Confirmed it fails withoutthis change (
k.pdis.size() == 1, expected 2) and passes with it.hsa-runtime64from this branch and, on an aie2p (Strix) NPU, ran afull-ELF kernel back to back (repeated dispatch of the same prepared
kernel) and a full-ELF kernel with a control scratchpad dispatched many
times in a row. Both produced correct results on every dispatch, where the
unpatched build gave a correct first dispatch and wrong results after.
Test Result
elf_parse_test: 14/14 pass (13 before this PR's new test).Hardware: repeated full-ELF dispatches correct on every run, where they were
previously correct only on the first.
Submission Checklist