[newling]

When the race detector plugin reports a race, it prints the instructions that the reporting wave executed between the conflicting memory instruction and the racing access. To produce that listing, each wave kept a ring buffer of its last 256 PCs, and each dispatch kept a shared table that mapped every PC executed by any of its waves to disassembly text. Every wave took that table's mutex on every issued instruction, even when the PC was already present. Profiling hipblaslt-bench GEMMs under multithreaded emulation showed this lock as a bottleneck, although those runs reported no races at all.

This PR removes the shared table. Each trace entry now holds the instruction's PC and a copy of its encoding, and each pending memory event keeps a copy of the instruction that produced it:

class InstructionSnapshot {
public:
  static constexpr size_t kMaxWords = 4;  // the CU's 16-byte fetch window
  // Throws std::logic_error unless the encoding has 1 to kMaxWords words.
  InstructionSnapshot(uint64_t pc, std::span<const uint32_t> encoding);
  uint64_t pc() const;
  std::span<const uint32_t> encoding() const;
  friend bool operator==(const InstructionSnapshot &, const InstructionSnapshot &) = default;

private:
  uint64_t pc_;
  std::array<uint32_t, kMaxWords> words_{};
  uint32_t word_count_;
};

Recording an instruction is now a copy of at most 16 bytes into wave-local storage, with no lock or hash lookup. Disassembly happens only when a report is printed: the plugin passes the captured bytes to the CU's target decoder through a new Wavefront::disassemble_instruction. That call is virtual because standalone plugin modules do not link the simulator. Because reports decode the bytes that actually executed, they remain correct when the code is overwritten or unmapped after execution, and when the producer of a pending event has already left the 256-entry trace. The old table kept the first text it saw at each PC.

The report layout and elision limits do not change. Markers and trace boundaries match both the PC and captured encoding, and each printed trace row is disassembled from its own bytes. This distinguishes rewritten instructions at the same PC while preserving selection for identical encodings. The "(before trace window)" label means no matching producer appears in the reporting wave's trace; it does not establish how old the conflict is, since another wave may have executed an instruction the reporting wave never executed. Distinguishing repeated executions of identical bytes at the same PC, and giving each wave in a cross-wave race its own history, are left for follow-up work.

Related changes:

• Every event registration method in EventRegistry, RaceDetector and WaveRaceState takes an InstructionSnapshot in place of a PC, and the event's PC comes from the snapshot. The snapshot's constructor rejects an empty or oversized encoding with std::logic_error, so every recorded event can be disassembled.
• The plugin identifies the executing instruction by the newest trace entry instead of wf.pc. s_swappc_b64 updates the PC before it writes its return-address registers, so a pending scalar load into those registers was previously reported at an address derived from the branch target instead of at the s_swappc_b64.
• A missing encoding or a decode failure throws std::logic_error instead of printing placeholder text. Every recorded instruction decoded successfully when it executed, so either failure indicates a bug in recording or decoding.
• Instruction no longer caches its disassembly string, and VOPD formats its two slots when asked (generator and generated sources updated). The C API still copies disassembly into the caller's buffer.
• The race detector guide now explains how to read the instruction trace in a report.

Memory changes: Instruction shrinks from 224 to 192 bytes, while each wave's trace grows from 2,064 bytes to as much as 8,192 bytes, allocated as the wave issues its first 256 instructions, and each pending event grows by 24 bytes.

Performance

Median process wall time of three interleaved hipblaslt-bench runs on emulated gfx950, comparing develop with this PR rebased onto it (Release, race detection enabled, one GEMM iteration, no cold iterations, no async helper threads, no races reported):

GEMM	Emulation / dispatch threads	Before	After
FP16 2048 × 2048 × 2048	8 / 16	8.79 s	5.14 s
BF16 4096 × 4096 × 1024	8 / 48	16.46 s	8.16 s
FP16 2048 × 2048 × 2048	1 / 1	22.73 s	22.79 s

The parallel cases improve by 1.71× and 2.02× because the lock is gone, and process CPU time falls from 129 s to 41 s and from 531 s to 85 s. With one thread there is no contention, and the difference is within run-to-run variation. Report formatting time is not measured.

Alternatives considered

1. Keep the shared text table, but make repeated visits lock-free (#13640). Atomic tracking of seen PCs avoids the lock on repeated visits, but new entries still need synchronized publication, and each PC still maps to the first text observed there. In a separate comparison where both designs also pooled detector and wave state, their wall times were within about 2% of each other. This PR is chosen for its simpler structure and for reporting the bytes that executed.
2. Record only PCs and read the code from memory when reporting (#14172). This is the cheapest to record, but it needs memory-reader plumbing and address-space lifetime handling in the plugin, and it cannot recover the executed bytes once the code has been overwritten or unmapped.

Testing

New tests overwrite a load's code with s_nop after it executes and check that the report still shows the executed ds_read_b32, including after the load has left the trace window. Repeated-PC regressions verify each row's captured bytes, marker placement and trace boundaries when producer or consumer encodings change. Other tests cover the scalar-load/s_swappc_b64 race and its waited control, literal and maximum-width encodings, decoding with no code mapping, concurrent decoding on several CUs, event registration through every method, snapshot construction limits and equality, trace wraparound, and the failure path for missing history. The race detector, trace formatting, event registry, shared wait, disassembly and standalone plugin ABI checks pass locally, and all HIP race corpus cases on gfx950 and gfx1151 pass in local and daemon modes.

ISSUE ID: #13640, #14172 (alternative proposals)

[therock-pr-bot]

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

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

[kuhar]

Added 5 comments

[newling]

Thanks @kuhar. All addressed in last pushed, except for the comment about debug_read_sgpr which I've done in a separate PR as it's not directly related #14313.

[kuhar]

Just some comments

[kuhar]

Added some comments

[kuhar]

One thing I noticed

[kuhar]

One comment, plus an inline reply

[kuhar]

It would be useful to add a short deferred-disassembly example to the plugin guide (emulation/rocjitsu/docs/plugins.md:157, or under "Adding a new plugin"). Plugins can copy the PC and size()/4 words from raw_encoding() in onAmdgpuBeforeExecuteInstruction, then call Wavefront::disassemble_instruction on the CU execution thread when a report needs text. This would make the new service discoverable alongside the guide's existing lifetime advice.

Please distinguish borrowed encodings and mnemonic views from Instruction::disassemble(), which now returns an owned std::string that callers may safely retain. Capturing bytes lets plugins defer formatting; retaining the returned string is also valid.

[kuhar]

Left some inline comments

[kuhar]

Optional: it would help to spell out the value types at these two capture points, following the style guide's preference when the return type is not visible on the line: const std::string dis in emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/code/rj_code.cpp:341, and const InstructionSnapshot instruction in emulation/rocjitsu/lib/rocjitsu/src/rocjitsu/vm/plugins/race_detector/plugin.cpp:542. This makes the owned disassembly text and captured snapshot easier to recognize without following the helper declarations.

LGTM

[github-actions]

rocjitsu post-commit fuzz passed

The 32-worker multiplexer campaigns for commit ddbfb7d489c1 completed without recording a crash or hang.

• ASan+UBSan (6-minute): 588,611 input executions (1,667.98 inputs/s)
• TSan (6-minute): 589,458 input executions (1,680.03 inputs/s)

View the workflow run.

[systems-assistant]

TheRock Submodule Bump Activity

Newest first

• 2026-11-02 — Included in TheRock via ROCm/TheRock#9668.