[rocjitsu][DBI][hook - 1/3] Basic Block Site Selection Layer by accauble · Pull Request #13292 · ROCm/rocm-systems

accauble · 2026-10-12T18:24:19Z

Motivation

This is a stack of 3 PRs to implement the DBI runtime hooks. The goal for this stack is to instrument every block in a kernel with a call to an empty probe when a kernel is loaded.

This PR implements changes that will eventually be used by the runtime hook, which gathers the basic blocks of a kernel, figures out the first instrumentable instruction within each block, and patches the object accordingly. In particular:

it exposes read-only BasicBlocks to a new DBI site selection layer that will run the Instrumentor via Instrumentor::kernel_blocks
Instrumentor only decodes reachable code from the kernel entry and probe analysis only happens on blocks reachable from the probe entry
creates a function to find the first instrumentable instruction in a BasicBlock, where instrumentable means that it is relocatable and not part of an s_clause
site selection automatically (for now) includes one site per kernel basic block at the first instrumentable instruction
instrumentation is still (for now) all-or-nothing
instruments with (for now) an empty probe

Other Changes

Probe calls trigger object-wide refusals for wave32, calls/unresolved edges, relative register accesses, and RDNA4 image instructions.
CDNA2/CDNA3/CDNA4 kernarg preload is refused independently of the point body.
Probe resolution rejects a missing concrete target when the destination has one.
Both patch and patch_with_debug_summaries() enforce the liveness and entry-layout refusals. Only tests that call InstrumentorTestAccess::skip_patch_gates() bypass them.
Differentiates between "kernel not supported" by instrumentation and "patching failed"
Selector Decisions:
- Only supported architecture (CDNA3, CDNA4, and RDNA4)
- Only single-kernel
- Needs a site to instrument
- Kernel instrumentation unsupported
- Patch fails
Shared DBT code is just a move. No behavioral change to DBT.

Issue Tracking

Part of Issue #8879. See this issue for a list of follow-ups suggested by this PR.

therock-pr-bot · 2026-10-16T16:02:39Z

✅ All Checks Passed — Ready for Review

Check	Status	Details
📝 PR Description	✅ Pass	—
⛔ Forbidden Files	✅ Pass	—
🧪 Unit Test	✅ Pass	—
🔎 pre-commit	✅ Pass	—
🚫 Draft PR	🔜 To Be Enabled	—
🚩 Feature Flag	🔜 To Be Enabled	—
📊 Code Coverage	🔜 To Be Enabled	—
🤖 therock-pr-bot	✅ Pass	—

🎉 All checks passed! This PR is ready for review.

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

🙋 Wish to Override Policy?

therock-pr-bot · 2026-10-17T16:11:06Z

🎉 All checks passed! This PR is ready for review.

kuhar

I think we should make the documented scope match this first hook milestone. The PR description says all basic blocks, but hooks/dbi_runtime_sites.h:58-65 limits selection to reachable blocks in a single-kernel object, skips blocks without a legal anchor, and rejects reachable calls or unresolved edges. hooks/dbi_runtime_sites.cpp:23-25 also restricts targets to CDNA3, CDNA4, and RDNA4. Please carry those limits into the PR description and the runtime-hook discussion in docs/dbi-design.md.

Also document the liveness limitation at Instrumentor::patch() in code/patch/instrumentor.h: the rationale currently appears at hooks/dbi_runtime_sites.cpp:138-142, while the Instrumentor itself constructs liveness without call/return edges and does not enforce this hook's rejection policy. Direct callers need to know the supported control-flow contract without discovering it in this caller.

accauble · 2026-10-19T18:26:16Z

I believe I have addressed all the comments. I have also updated the PR description as requested.

kuhar

I think it would help to clarify the new query API documentation in a small wording pass (paths below are under emulation/rocjitsu/):

In docs/dbi-design.md:70-71, give first_anchor(block) an explicit verb and name patch() / patch_with_debug_summaries() as the two methods sharing the attempt budget. “Both entry points” now follows the two non-consuming queries, which makes the restriction easy to misread.
In lib/rocjitsu/src/rocjitsu/code/patch/instrumentor.h:441-445, separate which blocks kernel_blocks() returns from its failure cases and non-consuming behavior; the sentence beginning “The blocks reachable_kernel_blocks() finds ... out of the ones” is difficult to follow. While editing this paragraph, change “either” at line 400 to “any of these calls” for the three listed decoding triggers.

## Motivation Fixes a bug found in PR #13292 (see [the comment](#13292 (comment))). On GFX9, there must be 5 wait states between a VALU instruction writing an SGPR and a VMEM instruction. Additionally, there must be 1 wait state between an M0 write and read. The DBI envelope uses `v_readlane` to restore spilled SGPRs and `s_mov_b32` to restore M0, but did not verify the wait. This PR makes the fix by counting the number of instructions emitted after the last restore and if that value is less than 1 or 5 and the target is a gfx9, it emits s_nops to create the space.

accauble · 2026-10-20T16:14:53Z

Addressed round 2 comments.

kuhar

I think it would help to update these test comments to match the cases they describe (paths under emulation/rocjitsu/):

tests/code/kernel_descriptor_scan_test.cpp:55-58: move the paragraph about rewriting an entry into .rodata back above EntryOutsideTextIsDropped. It currently introduces ObjectScanFindsEveryKernel, which does no such rewrite.
tests/dbi/hsa_dbi_nop_asm_test.cpp:70: describe choosing one legal anchor from each of the first two eligible kernel-reachable blocks. The fixture no longer takes the first two relocatable instructions across .text; that distinction is what keeps leading stubs out of this test.

accauble · 2026-10-21T15:21:19Z

Made those quick updates and moved a rogue comment

kuhar

Added some comments

kuhar

It would help to include the new direct-Instrumentor::patch() refusal contract in the PR description, alongside the selector changes. The October 23 update describes most of it, but the description still omits it:

Probe calls trigger object-wide refusals for wave32, calls/unresolved edges, relative register accesses, and RDNA4 image instructions.
CDNA3/CDNA4 kernarg preload is refused independently of the point body.
Probe resolution rejects a missing concrete target when the destination has one, while patch_with_debug_summaries() bypasses the liveness and entry-layout refusals.

These changes affect direct Instrumentor callers as well as the hook. Stating the object-wide scope and debug bypass here would make the API change discoverable without reconstructing it from the discussion.

kuhar

One thing I noticed

## Motivation `patch` is a wrapper for `patch_with_debug_summaries`. In PR #13292, it started to become a convenient place to put gates around instrumentation but have tests avoid them, but this was never the intended design. This PR adds a test-only setting for `patch_with_debug_summaries` so that the instrumentation gates can fall in one place and the tests can still by-pass the gates.

kuhar

We should update the PR description's patch_with_debug_summaries() bullet: both patch entry points now enforce the liveness and kernarg-preload refusals. Only tests explicitly using InstrumentorTestAccess::skip_patch_gates() bypass them. instrumentor.h:509-514 and the regressions for both entry points match this contract, so the current claim that the debug API bypasses the checks describes behavior it no longer provides.

## Motivation `patch` is a wrapper for `patch_with_debug_summaries`. In PR #13292, it started to become a convenient place to put gates around instrumentation but have tests avoid them, but this was never the intended design. This PR adds a test-only setting for `patch_with_debug_summaries` so that the instrumentation gates can fall in one place and the tests can still by-pass the gates.

kuhar

One thing I noticed

kuhar · 2026-10-28T18:31:56Z

+    return out;
+  }
+
+  InstrumentedCodeObject patched = instrumentor.patch();


We should make the Instrumentor refuse or repair PC-relative references to data that moves when .text grows, and cover this through instrument_each_block(). At code/patch/instrumentor.cpp:1659, replace_text(new_text) receives no PcRelativeDataRelocations, so it shifts allocated sections after .text without updating the kernel's s_getpc_b64 + add address calculation.

A single-kernel, wave64 gfx1201 HIP object reading an initialized __device__ int[64] returns Instrumented with one site and the empty probe. Its .data/g_table moves from 0x3970 to 0x4970, but the add literal remains 0x2268. On hardware, the original produces the expected 303, 306, ...; the patched object returns zero in all 64 lanes. Changing only that literal to 0x3268 restores all 64 results. The original/failure/control sequence was reproduced on gfx1201; CDNA was not tested.

Until these references are relocated, a conservative refusal in the Instrumentor could surface as UnsupportedKernel. Please add a selector regression with a PC-relative reference to .data after .text. The underlying patcher usage predates this PR, but the new automatic selector accepts this object and returns silently incorrect code.

That refusal is in the next PR in this stack. If there is nothing else, I'd rather keep it there and stop the back and forth going on here.

When adding the refusal in #13293, I think we should put it in the Instrumentor and map it to UnsupportedKernel here. Instrumentor::patch_with_debug_summaries() still calls replace_text(new_text) without data relocations, so a selector-only check would leave direct patch() callers exposed to the same stale reference. A shared patch gate would protect both entry points. Deferring runtime activation to #13293 is consistent with this PR: instrument_each_block() currently has only a unit-test caller.

* Expose blocks for site selection * Add first_anchor, which finds the first instrumentable insn in the given block. Instrumentable means it is relocatable and that it isn't covered by an s_clause

* For now, instrument each block. Decides whether to instrument a given code object and then instruments with a probe call. * Patching is still all-or-nothing

* Instrumentor now decodes only kernel-reachable code * Probe analysis only decodes from entry * Add DbiLoadDecision Count

* Update comments/documentation * Create and share scan_object_kernel_descriptors helper * Clearer errors for anchor offsets * Updated inlined nop and empty probe test fixtures

* `patch` will not instrument objects with wave32 kernel, a call, or unresolved successor (rather than Site Selector) * `patch` will not use a probe with no target when the object has one * Add new liveness gate for relative reg access * Add new gate to prevent instrumentation of a kernel with a kernarg preload * Document wave32 liveness issue (Issue #14022) * Stale doc fixes

* Instrumentor reports errors as "unsupported" and "patch failure" * Share `uses_kernarg_preload_firmware_skip` and add CDNA2 to it * Revert decision to include inlined assembly in liveness check since it is only allowed to contain nops at this point

* develop has a test-only parameter for patch_with_debug_summaries so all the gates in `patch` can be moved down.

* develop now handles the wave size in liveness, so undo the gates added to prevent its instrumentation

kuhar

We should update the PR description's refusal list: uniform wave32 kernels are now supported. DbiRuntimeSites.InstrumentsAWave32Kernel expects Instrumented, and the Instrumentor passes the kernels' shared wave size to liveness; only mixed Wave32/Wave64 objects remain refused on wave-size grounds. Please drop wave32 from the object-wide refusal bullet and mention that distinction.

kuhar · 2026-10-30T00:20:48Z

+- RDNA4 image address VGPRs. The VIMAGE and VSAMPLE decoders do not report them, so `InstDefUse` misses those reads (`is_unmodeled_image_access()`).
 - Memory dependencies. Liveness is purely register-based.

+For the call, unresolved-edge, relative-access and image limits, both `Instrumentor` patch entry points refuse a code object in which a point calls a probe if it has a call, an unresolved edge (code that runs off the end of `.text` included), a relative register access or an RDNA4 image instruction, and sets `InstrumentedCodeObject::kernel_unsupported` to tell the refusal from a failure to patch the queued points. A test can skip these patch gates through `InstrumentorTestAccess`.


Nit: It would help to name InstrumentorTestAccess::skip_patch_gates() in this sentence and the kernarg-wrapper paragraph near line 500. That gives test authors the actual opt-out call, matching instrumentor.h, without having to look up the helper type first.

github-actions Bot added documentation Improvements or additions to documentation project: rocjitsu labels Oct 12, 2026

accauble added this pull request to stack #13294 October 12, 2026 18:25

systems-assistant Bot added the organization: ROCm label Oct 12, 2026

Base automatically changed from users/accauble/rj-dbi-log-record-probe to develop October 16, 2026 16:01

accauble force-pushed the users/accauble/rj-dbi-runtime-site-selection branch from afdfd83 to 03ba35f Compare October 17, 2026 16:09

accauble marked this pull request as ready for review October 17, 2026 16:39

accauble requested review from a team and atgutier as code owners October 17, 2026 16:39

kuhar reviewed Oct 18, 2026

View reviewed changes

accauble requested a review from kuhar October 19, 2026 18:26

accauble mentioned this pull request Oct 19, 2026

[rocjitsu][DBI] Space gfx9 sgpr and M0 restores in envelope #13765

Merged

kuhar reviewed Oct 20, 2026

View reviewed changes

Comment thread emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/code/patch/instrumentor.cpp Resolved

Comment thread emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/hooks/dbi_runtime_sites.cpp Resolved

accauble force-pushed the users/accauble/rj-dbi-runtime-site-selection branch from 9ad197d to f5a7bc7 Compare October 20, 2026 16:13

accauble requested a review from kuhar October 20, 2026 16:15

kuhar reviewed Oct 20, 2026

View reviewed changes

Comment thread emulation/rocjitsu/tests/CMakeLists.txt Resolved

accauble mentioned this pull request Oct 20, 2026

[rocjitsu][DBIT][refactor] Refactor shared hotswap and dbi runtime #13842

Open

accauble requested a review from kuhar October 21, 2026 15:20

accauble force-pushed the users/accauble/rj-dbi-runtime-site-selection branch from 796a183 to c7474df Compare October 21, 2026 16:06

kuhar reviewed Oct 24, 2026

View reviewed changes

Comment thread emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/code/kernel_descriptor_scan.cpp Resolved

Comment thread emulation/rocjitsu/tests/code/kernel_descriptor_scan_test.cpp Resolved

Comment thread emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/code/patch/instrumentor.h Resolved

accauble requested a review from kuhar October 26, 2026 13:36

kuhar reviewed Oct 26, 2026

View reviewed changes

Comment thread emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/code/patch/instrumentor.cpp Resolved

accauble requested a review from kuhar October 27, 2026 15:54

accauble mentioned this pull request Oct 27, 2026

[rocjitsu][DBI] Move patch gate out and make it test-only #14388

Merged

kuhar reviewed Oct 27, 2026

View reviewed changes

Comment thread emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/code/kernel_descriptor_scan.h Resolved

accauble force-pushed the users/accauble/rj-dbi-runtime-site-selection branch from 6d79c45 to 240b770 Compare October 27, 2026 20:17

accauble force-pushed the users/accauble/rj-dbi-runtime-site-selection branch from 240b770 to 0492276 Compare October 27, 2026 20:53

kuhar reviewed Oct 28, 2026

View reviewed changes

Comment thread emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/code/patch/instrumentor.h Resolved

accauble requested a review from kuhar October 28, 2026 14:20

kuhar reviewed Oct 28, 2026

View reviewed changes

accauble requested a review from kuhar October 28, 2026 20:38

accauble added 13 commits October 29, 2026 14:14

[rocjitsu][DBI] Expose Instrumentor's blocks and first anchors

2539b88

* Expose blocks for site selection * Add first_anchor, which finds the first instrumentable insn in the given block. Instrumentable means it is relocatable and that it isn't covered by an s_clause

[rocjitsu][DBI] Site selector for DBI runtime

51ac90c

* For now, instrument each block. Decides whether to instrument a given code object and then instruments with a probe call. * Patching is still all-or-nothing

[rocjitsu][DBI] Address reviewer comments (round 1)

cdbf0f1

* Instrumentor now decodes only kernel-reachable code * Probe analysis only decodes from entry * Add DbiLoadDecision Count

[rocjitsu][DBI] Address reviewer comments (round 2)

3c43a99

* Update comments/documentation * Create and share scan_object_kernel_descriptors helper * Clearer errors for anchor offsets * Updated inlined nop and empty probe test fixtures

[rocjitsu][DBI] Fixed comments and hook platform req

040fa59

[rocjitsu][DBI] liveness bug gate + rename

66add16

[rocjitsu][DBI] Remove CDNA2 from uses_kernarg_preload_firmware_skip

59f1428

[rocjitsu][DBI] Refuse CDNA2 at patch too

806f74f

[rocjitsu][DBI] Move gates to patch_with_debug_summaries

d29090b

* develop has a test-only parameter for patch_with_debug_summaries so all the gates in `patch` can be moved down.

[rocjitsu][DBI] Comment tweaks

dc4dbee

[rocjitsu][DBI] Remove wave32 liveness gate

373b4ba

* develop now handles the wave size in liveness, so undo the gates added to prevent its instrumentation

accauble force-pushed the users/accauble/rj-dbi-runtime-site-selection branch from ff3780a to 373b4ba Compare October 29, 2026 20:15

kuhar reviewed Oct 30, 2026

View reviewed changes

accauble closed this Nov 6, 2026

Conversation

accauble commented Oct 12, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Motivation

Other Changes

Issue Tracking

Uh oh!

therock-pr-bot Bot commented Oct 16, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

✅ All Checks Passed — Ready for Review

Uh oh!

therock-pr-bot Bot commented Oct 17, 2026

Uh oh!

kuhar left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

accauble commented Oct 19, 2026

Uh oh!

kuhar left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

accauble commented Oct 20, 2026

Uh oh!

kuhar left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

accauble commented Oct 21, 2026

Uh oh!

kuhar left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

kuhar left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

kuhar left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

kuhar left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

kuhar left a comment

Choose a reason for hiding this comment

Uh oh!

kuhar Oct 28, 2026

Choose a reason for hiding this comment

Uh oh!

accauble Oct 28, 2026

Choose a reason for hiding this comment

Uh oh!

kuhar Oct 30, 2026

Choose a reason for hiding this comment

Uh oh!

kuhar left a comment

Choose a reason for hiding this comment

Uh oh!

kuhar Oct 30, 2026

Choose a reason for hiding this comment

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

accauble commented Oct 12, 2026 •

edited

Loading

therock-pr-bot Bot commented Oct 16, 2026 •

edited

Loading