[erwei-xilinx]

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.

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

• Look over the contributing guidelines at https://github.com/ROCm/rocm-systems/blob/develop/CONTRIBUTING.md.

🤖 Generated with Claude Code

[therock-pr-bot]

✅ All Policy Checks Passed

Check	Status	Details
📝 PR Description	✅ Pass	—
⛔ Forbidden Files	✅ Pass	—
🧪 Unit Test	✅ Pass	—
🚫 Draft PR	🔜 To Be Enabled	—
🚩 Feature Flag	🔜 To Be Enabled	—
📊 Code Coverage	🔜 To Be Enabled	—

🎉 All policy checks passed!

📖 Need help? See the Policy FAQ for details on every check and how to fix failures.

🙋 Wish to Override Policy?

[Copilot]

🟡 Changes recommended

The allocation check accepts external-variable ranges that extend beyond virtual-memory mappings, risking invalid memory access during dispatch.

1 open issue

• Validate VMM pointer extents against the actual mapping base · New

What changed in this PR

Adds external-variable support to the AIE full-ELF loading path, allowing applications to define buffers such as the control scratchpad without rewriting and repackaging code objects.

Changes:

• Parses external-symbol relocations and resolves them against application-defined variables.
• Records external buffers for per-dispatch cache flushing and residency.
• Adds scratchpad tests, build integration, and usage documentation without changing the hsaco format or packet ABI.

File	Description
projects/​rocr-runtime/​runtime/​hsa-runtime/​loader/​executable.cpp	Resolves external variables, validates allocations, and patches addresses.
projects/​rocr-runtime/​runtime/​hsa-runtime/​core/​runtime/​amd_aie_elf.cpp	Parses and validates external-symbol relocation sites.
projects/​rocr-runtime/​runtime/​hsa-runtime/​core/​inc/​amd_aie_section.h	Records external buffers in kernel descriptors.
projects/​rocr-runtime/​runtime/​hsa-runtime/​core/​inc/​amd_aie_elf.h	Defines parsed external-symbol metadata.
projects/​rocr-runtime/​runtime/​hsa-runtime/​core/​driver/​xdna/​amd_xdna_driver.cpp	Flushes external buffers and includes their handles in commands.
projects/​rocr-runtime/​rocrtst/​suites/​aie/​vector_scalar_add_param.py	Adds a scratchpad-driven vector-add design.
projects/​rocr-runtime/​rocrtst/​suites/​aie/​test_dispatch.cc	Tests variable validation and parameter changes between dispatches.
projects/​rocr-runtime/​rocrtst/​suites/​aie/​elf_parse_test.cc	Tests external-symbol parsing and invalid relocation rejection.
projects/​rocr-runtime/​rocrtst/​suites/​aie/​CMakeLists.txt	Extracts full-ELF build support and integrates scratchpad artifacts.
projects/​rocr-runtime/​docs/​aie-hsaco-format.md	Documents external-variable loading, memory requirements, and dispatch behavior.

_{Review effort: Balanced}

[cfreeamd]

🤖 AI-assisted pre-review (for @cfreeamd)

Core loader (executable.cpp), ELF parser (amd_aie_elf.cpp), and driver dispatch path (amd_xdna_driver.cpp) look correct. Variable resolution follows the same agent_symbols_ lookup path as GPU undefined symbols; cleanup on failure destroys the executable and reader; size consistency across multiple relocations naming the same symbol is enforced. New tests cover the three refusal paths (undefined, non-AIE memory, out-of-range) and the parameter-change-between-dispatches scenario. rocrtst_yaml_missing flag is a false positive — elf_parse_test and test_dispatch are standalone GTest binaries, not entries in the rocrtst main runner, so test_categories.yaml / platform_config.yaml updates are not required.

[systems-assistant]

TheRock Submodule Bump Activity

Newest first

• 2026-10-26 — Included in TheRock via ROCm/TheRock#9486.