[randomizedcoder]

Closes #8801

Summary

We moved a set of AMD GPUs (MI50/gfx906, W5700/gfx1010, WX 2100/gfx803) to a machine running Wayland while waiting for the amdgpu firmware reset issue to be resolved on our primary system. To our surprise, rocminfo and all ROCm/HSA workloads (including llama.cpp with ROCm) failed completely — no GPUs detected at all.

The root cause: the WX 2100 (gfx803/Polaris12) has DoorbellType=1, and the GpuAgent constructor throws HSA_STATUS_ERROR for any doorbell type other than 2. The DiscoverGpu handler in amd_topology.cpp only catches HSA_STATUS_ERROR_INVALID_ISA, so the uncaught exception propagates up and kills hsa_init() for all devices — including the perfectly supported MI50 and W5700.

How this happened

Git blame shows this was introduced across three commits by different authors:

1. Nov 2020 — ba667661c5 (Konstantin Zhuravlyov): "Implement Target ID Proposal"

   • Added the try/catch in DiscoverGpu that catches HSA_STATUS_ERROR_INVALID_ISA to gracefully skip GPUs with unrecognized ISA
   • Added throw hsa_exception(HSA_STATUS_ERROR_INVALID_ISA, ...) in GpuAgent for unknown ISA
   • The catch handler was designed specifically for ISA mismatches — nothing else

2. May 2025 — 0d70045817 (David Yat Sin): "rocr: Remove deprecated doorbell type 1 support"

   • Correctly removed ~95 lines of type-1 doorbell handling from amd_aql_queue.cpp
   • Added a 3-line guard in amd_gpu_agent.cpp: if (DoorbellType != 2) throw HSA_STATUS_ERROR
   • The bug: used HSA_STATUS_ERROR (generic) instead of HSA_STATUS_ERROR_INVALID_ISA, so the existing catch handler doesn't catch it

3. Mar 2026 — bc17c5424c (Sunday Clement): "rocr: Skip Unsupported GPUs from being enumerated (rocr: Skip Unsupported GPUs from being enumerated #4198)"

   • Refactored the catch handler with better formatting and GpuNodeDescription()
   • Still only catches HSA_STATUS_ERROR_INVALID_ISA — the doorbell guard's HSA_STATUS_ERROR remained uncaught

The removal of type-1 support was correct. The issue is just that the rejection path used the wrong exception status code for the existing catch handler.

Changes

The removal of type-1 support was correct; the only defect was the rejection path using HSA_STATUS_ERROR where the existing DiscoverGpu catch expects HSA_STATUS_ERROR_INVALID_ISA. The fix aligns the two.

1. amd_gpu_agent.cpp — the fix

• Replace bare if (DoorbellType != 2) throw HSA_STATUS_ERROR with a documented switch statement.
• Throw HSA_STATUS_ERROR_INVALID_ISA instead of HSA_STATUS_ERROR, so the existing DiscoverGpu handler catches it and skips just that device.
• Include the actual doorbell type value in the error message.
• Document all 4 doorbell types with their GPU families (from the kernel's kfd_topology.c).
• Include <sstream> directly (previously only transitively available) and bind the message to a named std::string before .c_str() to avoid a dangling-pointer smell.

2. amd_topology.cpp — skip path

• Keep the DiscoverGpu catch scoped to HSA_STATUS_ERROR_INVALID_ISA only; any other status is a genuine failure (e.g. out-of-memory, GetWallclockFrequency) and is still rethrown.
• Always print the skip message to stderr (was previously gated behind ifdebug) so users can see why a device was skipped.

Affected hardware

Any multi-GPU system containing a pre-Vega AMD GPU alongside supported GPUs:

DoorbellType	GPU families	Supported
0 (PRE_1_0)	Kaveri, Hawaii, Tonga	No
1 (1_0)	Carrizo, Fiji, Polaris10/11/12, Vegam	No
2 (2_0)	Vega and newer (GCN 5.0+, GC IP >= 9.0.1)	Yes
3 (Reserved)	Future use	No

Testing

Unit tests (standalone, no hardware required):

13 tests in doorbell_type_test.cc covering exhaustive combinatorial validation of the 2-bit doorbell type field:

• Positive/negative classification for each known doorbell type
• Exhaustive: exactly 1 of 4 possible values is supported
• Bit extraction: correct isolation of bits 12-13 from capability word
• Real hardware regression data: MI50 (0xac73a280), WX 2100 (0x00001280), W5700 (0x2883a280)
• Future-proofing: field width validation, synthetic values (666, 0xDEAD, 0xFFFF) anticipate the field being widened beyond 2 bits

All 13 tests pass:

[==========] 13 tests from 1 test suite.
[  PASSED  ] 13 tests.

Hardware integration test (MI50 + W5700 + WX 2100):

Validated on NixOS (kernel 6.19, ROCm 7.2) with all three GPUs present:

• rocminfo detects MI50 (gfx906) and W5700 (gfx1010), skips WX 2100 gracefully
• llama-cpp service on W5700 starts and runs successfully (Qwen3-8B, 5G VRAM)
• MI50 HSA initialization succeeds (gfx906 agent enumerated)
• stderr shows informative skip message:

  ROCm/HSA: Skipping unsupported GPU (node 2): Agent creation failed.
  The GPU node uses unsupported doorbell type 1 (only type 2 is currently supported).
  

rocminfo output confirms all supported agents:

Agent 1: AMD Ryzen Threadripper PRO 3945WX 12-Cores (CPU)
Agent 2: gfx906 — AMD Radeon Graphics (MI50)
Agent 3: gfx1010 — AMD Radeon Pro W5700

Build notes: Validating required multiple full ROCm stack rebuilds on NixOS. The dependency tree is deep — e.g.:

[1/6/22 built] building hipblaslt-7.2.0 (buildPhase): Generating assembly kernels done in 4922.3s!  260680/260680

This is part of why we invested in the standalone unit tests — they compile in seconds and catch the combinatorial logic errors without requiring a full ROCm rebuild.

Also submitted

A corresponding patch has been submitted to nixpkgs for the packaged ROCm runtime.

[randomizedcoder]

@Sundance636

[randomizedcoder]

@dayatsin-amd

[randomizedcoder]

How to trigger the bug?

[dayatsin-amd]

Thank you for the changes. They look ok to me.

Will monitor CI tests to confirm they pass before merging.

[randomizedcoder]

@dayatsin-amd

Great! Thanks.

There are more PRs filed, I think probably the most serious is the container ID truncation
https://github.com/ROCm/rocm-systems/pulls/randomizedcoder

[randomizedcoder]

@dayatsin-amd G'day! Heads up — the CI failure in GpuDiscoveryDeprecatedDoorbellTest was a test environment issue, not a problem with the runtime fix. I've pushed a fix in 5ddcf8abce.

What happened: The test reads KFD sysfs topology and sees all 8 GPUs on the CI host, but HSA only reports 1 because the container only has 1 GPU passed through. The strict EXPECT_EQ(8, 1) assertion failed. KFD sysfs exposes the full host topology regardless of container boundaries, while HSA respects the GPU visibility restrictions.

The fix: Replaced the strict equality with a bounded check — HSA must expose at least 1 GPU (when KFD has supported nodes) and must never expose more than KFD reports as supported. This preserves the important invariant (deprecated GPUs never leak through) while working correctly in restricted environments.

One thing this surfaced though — the integration test for this fix is designed to catch bugs on mixed GPU systems (deprecated + supported GPUs together), but with CI only passing 1 GPU per container, that scenario can never actually occur in CI. The standalone unit tests in doorbell_type_test.cc still validate the classification logic without hardware, so regressions in the doorbell handling code are caught. But the full integration path (KFD scan → HSA init → agent filtering) only gets a real workout on multi-GPU or mixed-GPU hardware.

Would it be feasible to pass through multiple GPUs to the rocrtst CI containers? Even 2 would make the test more meaningful, and it would probably benefit other multi-GPU tests in the suite too. Alternatively, a KFD topology mock/override mechanism (where the test injects a synthetic DoorbellType=1 node) would let us simulate deprecated hardware without needing the physical card — that might be the most robust long-term approach.

Also — if it would help to have the exact card that triggered this bug for CI testing, the Radeon Pro WX 2100 (gfx803/Polaris12, DoorbellType=1) is available on eBay for ~$32: https://www.ebay.com/itm/287179806374. We'd be happy to purchase and ship one to AMD for the CI system. Even though this card isn't "supported" by ROCm, its mere presence in a system caused complete ROCm failure for all other GPUs — which is exactly the kind of thing CI should be catching.

No urgency on the CI question — just flagging it while it's fresh. The fix should let CI pass now so this can merge. Thanks!

[randomizedcoder]

G'day! Just a heads up — I've submitted #5135 to fix the flaky rocprofv3-test-roctx-pause-resume-tracing test that's been blocking CI here. Hopefully once that lands, re-runs should pass cleanly.

[randomizedcoder]

@dayatsin-amd

Sorry to bug you, but I think if you approve this again, the CI will work now.

[randomizedcoder]

@dayatsin-amd

Any ideas? Not a lot of detail in the log :(

[randomizedcoder]

There is good progress by the AMD firmware team, so hopefully the cards will be a lot better soon, but without this PR getting merged, NONE of my AMD cards will work.
https://gitlab.freedesktop.org/drm/amd/-/work_items/5125

[randomizedcoder]

@cfreeamd @dayatsin-amd — thanks for the review. Rebased onto current develop and addressed the feedback.

1. hsa_shut_down() at the end of Run() — you're right, TestBase::Close() calls it via CommonCleanUp(). Removed the redundant call. While I was at it I also removed the matching hsa_init() call earlier in Run() — TestBase::SetUp() already initialises HSA via InitAndSetupHSA(), so calling hsa_init() again was creating a 2-init/2-shutdown refcount imbalance. The regression test is still meaningful: if hsa_init() aborted under deprecated-GPU presence, TestBase::SetUp() would fail before Run() ever executes. Convention check: debug_basic.cc, agent_props.cc, etc. all leave the lifecycle to TestBase; only reference_count.cc owns it explicitly and that test deliberately overrides SetUp() to a no-op.

2. doorbell_type_test.cc location — moved into suites/functional/ next to gpu_discovery_deprecated.cc. The build-instruction comment in the file header is updated for the new location and the actual gtest layout in this tree (the bundled gtest-all.cpp has stale src/*.cc includes that don't match the .cpp files actually present, so the comment now lists the individual sources to compile, mirroring rocrtst/gtest/CMakeLists.txt).

3. switch vs if-else in amd_gpu_agent.cpp — left as a switch, if you don't mind. The four documented doorbell-type values (0/1/2/3) are easier to follow as labelled cases than as a single negated if predicate, and a future maintainer adding a new supported type adds one case rather than mutating a condition. The fall-through default: already throws HSA_STATUS_ERROR_INVALID_ISA so DiscoverGpu's existing catch handles it, and the message includes the actual doorbell-type integer. Happy to revisit if you'd prefer otherwise.

Branch was rebased onto upstream/develop (clean — no conflicts after ~210 upstream commits). Standalone unit tests still pass from the new location:

[==========] 13 tests from 1 test case ran. (1 ms total)
[  PASSED  ] 13 tests.

[randomizedcoder]

@cfreeamd @dayatsin-amd

This bug is still active. Look, I don't really care how we get this fixed. We can keep the code, but ultimately we need all the 2 bits covered please.

This is beyond a joke now.

[266-750Balloons]

Also running into this issue. I hope this can be sorted out soon. I mostly just want ROCm for the occasional Blender render (emphasis on occasional; otherwise I would have gone NVidia, but since I don't depend on Blender for anything serious, the easier desktop usability of AMD is much more important to me).

Here's my system info just in case:
OS: Debian Testing Forky
ROCm version: TheRock7.13 (was also running into the problem on 7.2.3)
CPU: AMD Ryzen 5 2600 Six-Core Processor
GPU 1: RX 9060 XT 16 GB (gfx1200)
GPU 2: RX 580 8 GB (gfx803); I literally just use this one for passing through to VMs.

[dayatsin-amd]

@randomizedcoder It seems the CI is complaining about these errors:

[rocrtst] ld.lld: error: duplicate symbol: DoorbellTypeValidation_Supported_Type2_Vega_And_Newer_Test::test_info_
[rocrtst] >>> defined at doorbell_type_test.cc
[rocrtst] >>> CMakeFiles/rocrtst64.dir/_w/rocm-systems/rocm-systems/projects/rocr-runtime/rocrtst/suites/functional/doorbell_type_test.cc.o:(DoorbellTypeValidation_Supported_Type2_Vega_And_Newer_Test::test_info)
[rocrtst] >>> defined at gpu_discovery_deprecated.cc
[rocrtst] >>> CMakeFiles/rocrtst64.dir/__w/rocm-systems/rocm-systems/projects/rocr-runtime/rocrtst/suites/functional/gpu_discovery_deprecated.cc.o:(.bss+0x0)
[rocrtst]

[randomizedcoder]

@dayatsin-amd — quick update: build passes now (Linux 2h27m, Windows 2h15m). The original ld.lld duplicate-symbol blocker is resolved.

The one red mark left is Test rocrtst (shard 1 of 1), but it's an unrelated flake — not from anything this PR touches:

ROCr Virtual Memory Test - InterProcess
projects/rocr-runtime/rocrtst/suites/functional/virtual_memory.cc:1950: Failure
  hsa_amd_vmem_import_shareable_handle(dmabuf_fd, &imported_handle)
  Actual: 4109   Expected: HSA_STATUS_SUCCESS
##[error]The action 'Test' has timed out after 15 minutes.

That's the inter-process dma-buf shareable-handle test — different file, different feature. The 15-min timeout suggests the child process hangs after the failure. All other Linux test shards (hip-tests x4, rocgdb, rocprofiler-sdk, rocr-debug-agent, sanity) pass.

Would you mind re-running just that job, or is there a path to merge from here? Happy to do whatever's most useful.

[shwetagkhatri]

Now new tests need to be added to test_categories.yaml. Otherwsie the test will not run in CI tier pipelines. Please add the test to all (quick/standard/comprehensive/full) tiers.
In test_categories.yaml under standard.test_patterns, add:
- "rocrtstFunc.GpuDiscoveryDeprecatedDoorbellTest"

[randomizedcoder]

Thanks for the review, @shwetagkhatri — all three points are addressed in e9ddb991eb:

1. test_categories.yaml — registered rocrtstFunc.GpuDiscoveryDeprecatedDoorbellTest in all four tiers (quick, standard, comprehensive, full), each grouped next to rocrtstFunc.AgentPropertiesTests.
2. amd_topology.cpp — collapsed the two redundant comment lines into the single line you suggested: // - HSA_STATUS_ERROR_INVALID_ISA: unrecognized GPU or deprecated doorbell type (pre-Vega GPUs).
3. doorbell_type_test.cc — added the standard AMD MIT SPDX header to match the sibling files in suites/functional/.

I'd also like to share some friendly feedback on the process. This PR has now been open for 84 days (since April 12), for a change that's small and low-risk — it's purely about how the 2-bit DoorbellType field (bits 12–13 of the capability word) is handled: rejecting an unsupported value with HSA_STATUS_ERROR_INVALID_ISA so DiscoverGpu skips that one GPU, instead of the current HSA_STATUS_ERROR that escapes the catch and aborts hsa_init() for every GPU in the system. The impact isn't limited to us — @266-750Balloons independently hit the exact same failure (comment above, May 21): a single pre-Vega card in the chassis silently takes down the entire ROCm stack for all other GPUs.

More broadly, I've found it discouraging that a number of the PRs and issues I've filed against this repo haven't received much traction — several PRs were closed without review, and I've had fixes go unlooked-at to the point where I opened a separate issue to escalate (e.g. #7081, the amdsmi container-ID truncation that's affecting a production MI300X deployment). I really want to be a constructive contributor here, and I'm very happy to keep addressing feedback — it would just help a lot to know these contributions are being seriously considered.

This one already has an approval from @dayatsin-amd. I'd genuinely appreciate help getting it prioritized and merged, and I'm glad to make any further tweaks needed. Thanks for your time.

[randomizedcoder]

Thanks so much for approving and running the CI — really appreciate it! 🙏

The results look great: all the substantive checks are green — rocr-runtime Code Restrictions, ROCR Runtime WSL Test, and HIP NVIDIA CI all pass, and on the Windows side the full package build plus all 12 hip-tests shards and the sanity check passed too.

The only two red marks look like unrelated CI infrastructure hiccups rather than anything in this change:

• The Linux "Build Linux Packages" jobs (gfx94X + gfx950) failed while checking out a rocprofiler-sdk submodule, before the build ever reached rocr-runtime/rocrtst:

  CMake Error at cmake/rocprofiler_utilities.cmake:233 (message):
    Command: "git submodule update --init external/gotcha"
  ninja: build stopped: cannot make progress due to previous errors.
  

• therock-pr-bot is failing on a missing script (tools/systems_pr_bot/policy_check.py).

Neither touches the doorbell fix. @shwetagkhatri — given the green results everywhere the code is actually exercised, and the existing approval from @dayatsin-amd, I think this is in good shape to merge. A quick re-run of the Linux job should clear it to all-green, and I'm happy to rebase if that helps trigger a fresh run. Thanks again for the help getting this moving!

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

[therock-pr-bot]

🚫 Please fix the failed policies before requesting reviews.

The following policy checks failed:

• ❌ PR Title/Description

The Not ready to Review label has been added to this PR.
Once all policies pass, the label will be removed automatically.

[randomizedcoder]

Filed #8801 to document the root cause, and added Closes #8801 to the description so the therock-pr-bot PR-description policy passes.

I'll be direct: I'm extremely disappointed by how this has gone. This is an ~3-line fix for a pre-existing runtime bug — root cause traced in #8801 to commit 0d70045817, which throws HSA_STATUS_ERROR (a status the DiscoverGpu catch handler was never written to catch), so a single pre-Vega card aborts hsa_init() for every GPU in the system. It was approved on 2026-04-14, has now been open 97 days, cleared four successive rounds of feedback, and independently hit by a second reporter (@266-750Balloons). The latest blocker was a bot requiring me — the reporter — to file an issue number for a bug I did not introduce.

That issue is now filed and referenced. Requesting a CI re-run and a path to merge. @dayatsin-amd's approval already stands; @shwetagkhatri's 2026-06-16 test_categories.yaml request is addressed in e9ddb991eb (all four tiers) — could that review be re-run or dismissed?

[dayatsin-amd]

@shwetagkhatri can you review rocrtst changes. Thanks

[saadrahim]

@cfreeamd Can you please help merge this fix in? I only see some minor changes needed. The ROCr team should make the changes needed or propose them in the PR.

[shwetagkhatri]

Hi @randomizedcoder ,

Thank you for your patience and for the care you put into this PR. You carried it through several rounds of successful feedbacks. The experience you have had getting a small, correct fix landed has not been good, and that is on us. I am sorry it has taken this long, and I want to get it merged.

To get this unblocked I have rebased your branch cleanly onto the current develop. Your commits and authorship are preserved exactly, and I have added the two small items from the last review round on top: registering the DoorbellTypeValidation unit tests in all four CI tiers in test_categories.yaml, and narrowing the DiscoverGpu catch to HSA_STATUS_ERROR_INVALID_ISA so unrelated construction failures surface instead of being silently dropped. You can see the clean, conflict-free result on tip of develop here for reference: #10136 (opened as a draft. it is only a reference, the fix is still meant to land through this PR).

There are two ways we can bring this in, and I am happy with whichever you prefer:

1. With your OK, I force-push the rebased branch to your rocminfo-bug branch (using --force-with-lease so nothing of yours is lost if you have pushed since). PR fix(rocr): fix multi-GPU init failure from deprecated doorbell type #4936 then updates in place, with your commits and authorship.
2. If you would rather keep control of your own branch, the rebased branch is already in fix(rocr): multi-GPU init failure from deprecated doorbell type (reference rebase of #4936) #10136, so you can adopt it directly:
   git fetch https://github.com/ROCm/rocm-systems.git users/shweta/randomizedcoder-rocminfo-bug
   git checkout rocminfo-bug
   git reset --hard FETCH_HEAD
   git push --force-with-lease origin rocminfo-bug
   That points rocminfo-bug at the clean history and updates this PR in place under your name.

Let me know which you prefer and I will take care of my side promptly, and I will help shepherd it through the remaining CI and to merge. Thanks again for sticking with this.

[randomizedcoder]

@shwetagkhatri - I really don't mind how we do it. This is a small fix for a 2 bit field. - I'd love to do a little debrief about this once we get this landed

[saadrahim]

@shwetagkhatri please force push to this PR and let's get this merged.

[randomizedcoder]

@shwetagkhatri @saadrahim

Rebased this branch onto current develop (it was ~3300 commits behind) and squashed everything into a single commit so the history is clean — then addressed the outstanding review points.

Changes since the last review:

• amd_gpu_agent.cpp: the doorbell guard now throws HSA_STATUS_ERROR_INVALID_ISA (matching what DiscoverGpu already catches) via a documented switch, added <sstream> directly, and bound the error message to a named std::string before .c_str().
• amd_topology.cpp: kept the DiscoverGpu catch scoped to HSA_STATUS_ERROR_INVALID_ISA only — any other status is a genuine failure (OOM, GetWallclockFrequency, etc.) and is rethrown rather than masked.
• test_categories.yaml: registered the standalone DoorbellTypeValidation.* unit tests across all four CI tiers (no hardware needed; 13/13 pass locally).

I also ran the project's static-analysis tools (semgrep / flawfinder / cpplint) scoped to the changed lines — no new findings remain.

This supersedes the reference rebase in #10136 (same content plus the fixes above, authorship preserved). The red CI on #10136 was unrelated infrastructure — TheRock source-fetch failures and a test runner with no GPU
(HSA_STATUS_ERROR_OUT_OF_RESOURCES); the same jobs are currently red across other open PRs.

[shwetagkhatri]

Ok. All approvals are done. Now waiting on CI to complete. If you see Rebase and merge button Green, please go ahead and merge @randomizedcoder . I will circle back in couple of hours.

[randomizedcoder]

Follow-up: verified on hardware that this fix works as merged.

Setup: Lenovo ThinkStation P620, kernel 7.2.7, two AMD GPUs — a Radeon Instinct MI50 (gfx906, KFD doorbell type 2) and a Radeon PRO WX 2100 (Polaris12 / gfx803, KFD doorbell type 1). That is exactly the multi-GPU trigger from #8801.

With ROCm 10.1.0 (ROCr from the therock-10.1 tag, which contains 8210339, built via the nixpkgs rocmPackages: 7.2.3 -> 10.1.0 update, NixOS/nixpkgs#573885):

$ rocminfo
ROCm/HSA: Skipping unsupported GPU: Polaris12 (node 2).
  Reason: Agent creation failed.
The GPU node uses unsupported doorbell type 1 (only type 2 is currently supported).

  Use ROCR_VISIBLE_DEVICES to limit to supported GPU(s) if needed.

exit 0, and the MI50 is enumerated (gfx906:sramecc+:xnack-). llama.cpp (HIP backend) on the same box reports found 1 ROCm devices and runs Qwen2.5-3B Q4_K_M at ~158 tok/s tg64 on the MI50, with the same skip line on stderr.

Control: the previous ROCm 7.2.3 rocminfo on the same machine, same cards, fails at hsa_init() with HSA_STATUS_ERROR: A generic error has occurred and lists no agents — the behaviour this PR fixed.

Thanks for getting it merged.