Skip to content

fix(rocr): Select I$ and GL2 invalidation for gfx125x code caches - #13924

Merged
renjithravindrankannath merged 1 commit into
developfrom
users/dayatsin/gfx125x-code-cache-inv
Nov 3, 2026
Merged

renjithravindrankannath merged 1 commit into
developfrom
users/dayatsin/gfx125x-code-cache-inv

Conversation

@dayatsin-amd

@dayatsin-amd dayatsin-amd commented Oct 21, 2026 •

Copy link
Copy Markdown
Contributor

Summary

On gfx125x, code-cache invalidation selects only the instruction cache and GL2. Scalar and vector caches are maintained by dispatch, and GCR_CNTL[9] is reserved. A non-null pointer uses RANGE, a null pointer uses ALL, and the two operations are sequenced in reverse. The shared base/size encoding is unchanged, and every other ISA keeps its existing acquire-mem packet.

JIRA ID

The internal change did not name a ticket.

Test plan

  • On gfx125x, load a code object and launch a kernel so the non-null RANGE invalidation runs. c07-1 (gfx1250, official amdgpu 7.1.0-2424461), kernel result 1. Two RANGE packets: isa=12.5, 8 dwords, header c0065800, DW7 00025002. Size granules matched (14449 → 0x39, 36864 → 0x90) and the base fields matched the pointers.
  • On gfx125x, exercise the null-pointer ALL invalidation and confirm it completes. Same run submitted one null packet and waited for it: header c0065800, DW2/DW3 ffffffff/ff, DW7 00024001. Product callers pass the code-object pointer, so this packet was submitted once from the test build.
  • gfx1100 on dayatsin-test-01 (gfx10+, not gfx125x) submits the existing acquire-mem packet.
  • gfx950 on e3-2b (pre-gfx10) submits the existing acquire-mem packet.

Made with Cursor

K$ and V$ are maintained by dispatch, and GCR_CNTL[9] is reserved on
gfx125x. Use RANGE for a non-null pointer and ALL for a null pointer,
sequenced in reverse. Keep the common base/size encoding and leave
other ISA paths unchanged.

Signed-off-by: Shane Xiao <shane.xiao@amd.com>
Signed-off-by: David Yat Sin <David.YatSin@amd.com>
Co-authored-by: Cursor <cursoragent@cursor.com>
@therock-pr-bot

Copy link
Copy Markdown

❌ PR Check — Action Required

Check Status Details
📝 PR Description ❌ Fail Error: PR description must reference a JIRA ID, ISSUE ID, or a GitHub closing keyword.
Expected: include a JIRA ID / ISSUE ID line (separator : or -, or omitted; value may be a JIRA key, a number with/without #, or a link), OR a closing keyword + issue reference. Accepted examples:
• JIRA ID : TESTAUTO-6039
• JIRA ID - #330
• JIRA ID #330
• JIRA ID (on separate line)
ROCM-25757
• ISSUE ID : TESTUTO-3334
• ISSUE ID #3334
• ISSUE ID - TESTAUTO-3433
• ISSUE ID (on separate line)
AIRUNTIME-2352
• ISSUE ID : https://github.com/<org_name>/<repo_name>/issues/1234
• Closes #10
• Fixes octo-org/octo-repo#100
• Resolves: #123
• #123
• https://github.com/<org_name>/<repo_name>/issues/123
Current: no valid JIRA/ISSUE/closing-keyword reference found
⛔ Forbidden Files ✅ Pass —
🧪 Unit Test ⚠️ Warning Error: Source/code files changed without an accompanying unit test.
Expected: add at least one test file named like test_<name>.py / test_<name>.cpp (or <name>_test.*).
Current: code file(s) changed: projects/rocr-runtime/runtime/hsa-runtime/core/inc/amd_gpu_pm4.h, projects/rocr-runtime/runtime/hsa-runtime/core/runtime/amd_gpu_agent.cpp; no test file found
🔎 pre-commit ⏳ Pending ⏳ Still running…
🚫 Draft PR 🔜 To Be Enabled —
🚩 Feature Flag 🔜 To Be Enabled —
📊 Code Coverage 🔜 To Be Enabled —

⚠️ 1 policy check(s) failed. Please address the issues above before this PR can be Reviewed.

🚫 Please fix the failed policies

  • ❌ PR Description

The Not ready to Review label was added to this PR. Once all policies pass, the label is removed automatically.

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

🙋 Wish to Override Policy?

@therock-pr-bot

Copy link
Copy Markdown

🚫 Please fix the failed policies before requesting reviews.

The following policy checks failed:

  • ❌ PR Description

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

@dayatsin-amd

Copy link
Copy Markdown
Contributor Author

CI test failures are unrelated to this change

This PR only changes InvalidateCodeCaches for gfx12.5, selecting I$ and GL2 invalidation. gfx94X and gfx110X stay on the existing acquire-mem packet. On the shared gfx10+ path the only edit is replacing the literal 1 with PM4_ACQUIRE_MEM_GCR_CNTL_GLI_INV_ALL (1U), so the encoded dword is the same.

Completed test failures:

  • Linux gfx94X rocprofiler-sdk (job 119396966509): 2 of 2179 tests failed, both in tests.integration.validate.app-coordinated-shutdown-fork (test_json_is_valid, test_flushed_markers). pytest raises UnicodeDecodeError: 'utf-8' codec can't decode byte 0xbb while reading a trace. gfx94X does not take the new path.
  • Windows gfx110X hipdnn-samples (job 119413754358): hipdnn_sample_autotune failed because MIOpen kernels were missing (kpack: kernel not found for miopen_lib_gfx1101.kpack / miopen_lib_gfx11-generic.kpack). gfx110X does not take the new path.

The two Multi-Arch CI Summary failures come from earlier runs that were cancelled (setup, linux_build_and_test, windows_build_and_test), superseded by the run above. gfx125X test jobs were skipped, so CI did not exercise the new path. gfx94X rccl and rocshmem were still queued when this was checked.

therock-pr-bot is the separate description-policy failure (no JIRA / issue id), already noted on the PR.

@dayatsin-amd

Copy link
Copy Markdown
Contributor Author

Re-running the failed CI tests on TheRock Multi-Arch CI run 39637381994:

  • Linux::release / Test gfx94X-dcgpu / Test rocprofiler-sdk (shard 1/1)
  • Windows::release / Test gfx110X-all / Test hipdnn-samples (shard 1/1)
  • Multi-Arch CI Summary

@dayatsin-amd

Copy link
Copy Markdown
Contributor Author

Re-running the failed CI tests on https://github.com/ROCm/rocm-systems/actions/runs/39637381994 (attempt 3):

  • Test rocprofiler-sdk (shard 1/1) (gfx94X-dcgpu)

The Test step timed out after 20 minutes. This PR only changes code-cache invalidation for gfx125x, so gfx94X stays on the existing acquire-mem packet. Multi-Arch CI Summary will run again after that job finishes.

@dayatsin-amd dayatsin-amd self-assigned this Oct 30, 2026
@renjithravindrankannath

Copy link
Copy Markdown
Contributor

Bypassing the rule due to known CI job failure.

@renjithravindrankannath
renjithravindrankannath merged commit a9c3b00 into develop Nov 3, 2026
845 of 872 checks passed
@renjithravindrankannath
renjithravindrankannath deleted the users/dayatsin/gfx125x-code-cache-inv branch November 3, 2026 16:42
systems-assistant Bot pushed a commit to ROCm/ROCR-Runtime that referenced this pull request Nov 3, 2026
 caches (#13924)
MIME-Version: 1.0
Content-Type: text/plain; charset=UTF-8
Content-Transfer-Encoding: 8bit

## Summary

On gfx125x, code-cache invalidation selects only the instruction cache
and GL2. Scalar and vector caches are maintained by dispatch, and
GCR_CNTL[9] is reserved. A non-null pointer uses RANGE, a null pointer
uses ALL, and the two operations are sequenced in reverse. The shared
base/size encoding is unchanged, and every other ISA keeps its existing
acquire-mem packet.

## JIRA ID

The internal change did not name a ticket.

## Test plan

- [x] On gfx125x, load a code object and launch a kernel so the non-null
RANGE invalidation runs. c07-1 (gfx1250, official amdgpu 7.1.0-2424461),
kernel result 1. Two RANGE packets: isa=12.5, 8 dwords, header
`c0065800`, DW7 `00025002`. Size granules matched (14449 → `0x39`, 36864
→ `0x90`) and the base fields matched the pointers.
- [x] On gfx125x, exercise the null-pointer ALL invalidation and confirm
it completes. Same run submitted one null packet and waited for it:
header `c0065800`, DW2/DW3 `ffffffff`/`ff`, DW7 `00024001`. Product
callers pass the code-object pointer, so this packet was submitted once
from the test build.
- [x] gfx1100 on dayatsin-test-01 (gfx10+, not gfx125x) submits the
existing acquire-mem packet.
- [x] gfx950 on e3-2b (pre-gfx10) submits the existing acquire-mem
packet.

Made with [Cursor](https://cursor.com)

Signed-off-by: Shane Xiao <shane.xiao@amd.com>
Signed-off-by: David Yat Sin <David.YatSin@amd.com>
Co-authored-by: Shane Xiao <shane.xiao@amd.com>
Co-authored-by: Cursor <cursoragent@cursor.com>
[rocm-systems] ROCm/rocm-systems#13924 (commit a9c3b00)
@systems-assistant

Copy link
Copy Markdown
Contributor

TheRock Submodule Bump Activity

Newest first

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.

5 participants