Skip to content

fix(hrr): bound the symbol offset in graph memcpy to/from-symbol replay - #14350

Merged
danielbolsson merged 12 commits into
developfrom
users/danielbolsson/hrr-3093
Nov 4, 2026
Merged

danielbolsson merged 12 commits into
developfrom
users/danielbolsson/hrr-3093

Conversation

@danielbolsson

@danielbolsson danielbolsson commented Oct 27, 2026 •

Copy link
Copy Markdown
Contributor

JIRA ID : AIRUNTIME-3093

Motivation

An HRR archive is untrusted input. The six graph symbol-copy replay handlers add the record's uint64_t offset to the live address of the symbol and pass the sum to CLR's 1D memcpy node calls. None of them reads the symbol's registered size, so nothing bounds offset + count. When the address does not resolve to a memory object, CLR skips its own direction and size checks, and a host-to-host or default-kind node can reach an unchecked host memcpy at an address the symbol does not own.

On gbt350 (MI350X) with a 64-byte __device__ global, a standalone program making the same hipGraphAddMemcpyNode1D(symbol + offset, ...) call showed: offsets 0 to 64 validated, offsets 65 to about 4095 succeeded silently as an out-of-bounds write, and offset 4096 faulted at graph launch. The thresholds depend on the process layout.

Technical Details

  • projects/hrr/playback/hip_playback.cpp: new resolve_symbol_range(). It looks the symbol up with its registered size and refuses the record with hipErrorInvalidValue, before any HIP call, when offset > size or count > size - offset. The sum is never formed, so it cannot wrap.
  • It replaces the inline translate_symbol calls in hipGraphAddMemcpyNodeToSymbol and hipGraphAddMemcpyNodeFromSymbol, and the file-local symbol_or_complain used by the four hipGraph{,Exec}MemcpyNodeSetParams{To,From}Symbol handlers, so all six go through one check.
  • projects/hrr/tests/unit/hrr_symbol_copy_bounds_test.cc (new, added to HRR_HANDLER_TEST_SOURCES): three handler-level cases over all six handlers, with a fake symbol registry and stand-ins for hipGraphAddMemcpyNode1D, hipGraphMemcpyNodeSetParams1D and hipGraphExecMemcpyNodeSetParams1D that count calls. No GPU is needed.

Not changed: CLR's arena fallback reports a SIZE_MAX size, which defeats its own bound check for other producers of such nodes. That is a separate change.

Issue Tracking

JIRA ID : AIRUNTIME-3093

Test Plan

  • Handler tests, hrr-handler-tests "Unit_HRR_SymbolCopyBounds_*": an in-range copy (including one ending exactly at the symbol's end and a zero-length one at offset == size) reaches HIP once, aimed at symbol + offset; out-of-range and wrapping (UINT64_MAX - 3 plus 8) records return hipErrorInvalidValue with zero HIP calls.
  • Negative control: the same tests against hip_playback.cpp reverted to develop.
  • Full hrr-unit-tests, hrr-handler-tests and hrr-integration-tests on gbt350 (MI350X, HIP_VISIBLE_DEVICES=4,5) with a capture-enabled CLR built from the same tree.
  • End to end through hrr-playback: a program captures all six spellings against a 64-byte __device__ global; offset in each record is then patched in events.bin to 64, 4096 and 0xFFFFFFFFFFFFFFFC, and each archive is replayed with this branch and with develop.

Test Result

  • hrr-handler-tests: 35 cases, 710 assertions pass (three new). hrr-unit-tests: 82 cases, 1524 assertions pass. hrr-integration-tests: 150 pass, 1 fails, Unit_HRR_SecureExecIgnoresCaptureOutput (AIRUNTIME-3074, known and unrelated).
  • Negative control: without the fix, the out-of-range and wrapping cases fail with hip_playback.cpp reverted (andraghetti's MI350 run); the in-range control passes.
  • End to end: the unpatched archive replays PASS on both builds. All 18 patched archives (6 handlers x 3 offsets) are refused by this branch with ... does not fit symbol ... (64 bytes registered); refusing the call. With develop, 16 of the 18 replay with exit 0 and no error; the other two fail inside CLR at hipGraphLaunch (offset 4096, Exec spellings).
  • hrr_test_config.py check for integration.yaml is OK (292 cases). The same check for unit.yaml already fails on develop (42 cases not in the config, none of them from this change); the new cases are in hrr-handler-tests, which no config lists.

Submission Checklist

The six symbol-copy graph handlers added the archive's uint64_t offset to
the live symbol address and passed the sum to the 1D memcpy node calls
without reading the registered symbol size, so nothing bounded offset +
count. An offset past the symbol could write or read outside it, or fault
at graph launch.

Resolve the symbol through one helper that refuses the record with
hipErrorInvalidValue, before any HIP call, when offset + count exceeds the
registered size. The comparison never forms the sum, so it cannot wrap. A
symbol registered with size 0 (the archive size was substituted) refuses
every copy.

Add handler tests with stand-ins for the three 1D HIP calls: in-range
copies reach HIP at symbol + offset; out-of-range, wrapping and size-0
records return hipErrorInvalidValue without reaching it.

AIRUNTIME-3093
@danielbolsson
danielbolsson requested a review from a team as a code owner October 27, 2026 10:29
Copilot AI balanced review requested due to automatic review settings October 27, 2026 10:29

Copilot AI left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Copilot was unable to review this pull request because the user who requested the review has insufficient remaining quota.

@therock-pr-bot

Copy link
Copy Markdown

✅ 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?

@andraghetti andraghetti left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Gated on MI350 (gbt350): head 8ad105f green; with hip_playback.cpp reverted, OutOfRange, WrappingSum and ZeroSize fail.

Comment thread projects/hrr/playback/hip_playback.cpp Resolved
…und (AIRUNTIME-3093)

playback___hipRegisterVar never registers a size-0 symbol: hipModuleGetGlobal
fails for a size-0 variable, so the symbol is not recorded at all. The
'size == 0 ||' term in resolve_symbol_range could only fire on a registry state
a real archive cannot produce. Remove it and the two unit cases that did not
earn their place: the size-0 case, and the unregistered-symbol case, which
passes against develop's own hipErrorInvalidSymbol.
Combine the symbol-copy bounds check (resolve_symbol_range, AIRUNTIME-3093)
with the host-source refusal (to_symbol_source) in every symbol-copy handler;
keep both new unit test files.
@danielbolsson
danielbolsson merged commit 601045e into develop Nov 4, 2026
58 of 59 checks passed
@danielbolsson
danielbolsson deleted the users/danielbolsson/hrr-3093 branch November 4, 2026 08:31
@systems-assistant

Copy link
Copy Markdown
Contributor

TheRock Submodule Bump Activity

Newest first

danielbolsson added a commit that referenced this pull request Nov 6, 2026
develop's #14350 added three hrr-handler-tests cases (52 compiled vs 49
listed), so the compiled-case-list check failed. The case set was checked
against the TEST_CASE names in the handler sources, not a built binary.

Co-Authored-By: Claude Sonnet 5.5 <noreply@anthropic.com>
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Projects

None yet

Development

Successfully merging this pull request may close these issues.

3 participants