AMDGPU: guard hotswap entry stub idempotency decode#3182
Merged
chinmaydd merged 3 commits intoJul 7, 2026
Conversation
github-actions
Bot
added
comgr
This was referenced Jul 5, 2026
chinmaydd
reviewed
Jul 6, 2026
| Ctx.Elf.getKernelStaticLdsSize(KernelName); | ||
| // Trampoline could not be applied: the original ds_*_addtid_b32 stays | ||
| // in the code object and will silently truncate M0 to 16 bits on A0 | ||
| // (DEGFXMI400-12025) whenever the runtime LDS layout exceeds 64 KiB. |
There was a problem hiding this comment.
Any reason we are removing references to internal tickets ?
Author
There was a problem hiding this comment.
That was intentional. I replaced the internal ticket with the architecture-level failure mode (gfx1250 A0 16-bit M0 truncation) so the public code comment is self-contained.
chinmaydd
reviewed
Jul 6, 2026
Comment on lines
+20
to
+23
| /// - ds_*_addtid_b32 : compute the LDS address through the ALU and | ||
| /// issue a regular ds_*_b32, bypassing the gfx1250 A0 16-bit M0 | ||
| /// truncation. On B0 the DS unit reads 20 bits of M0; on A0 it reads only | ||
| /// 16, silently dropping bits [19:16]. |
chinmaydd
reviewed
Jul 6, 2026
| llvm::ArrayRef<uint8_t> ShortCandidate(Stub.data(), MinInstSize); | ||
| EXPECT_FALSE(hasKernelEntryTrampolinePrefix(ShortCandidate, S)); | ||
| } | ||
|
|
There was a problem hiding this comment.
Is this inspired from an actual test case ? If yes, it would be nice to add a test that trips up the decoder.
Author
There was a problem hiding this comment.
Yes. The real-entry bytes are from the gfx1250 hipBLASLt MXF8/BF16 smoke. I added that reduction plus a separate decode case so the prefix gate covers the decoder-trip path too.
chinmaydd
approved these changes
Jul 7, 2026
harsh-amd
added a commit
that referenced
this pull request
Jul 9, 2026
## Summary - clamp gfx12 hotswap entry-stub INST_PREF_SIZE to the trampoline stub span so the descriptor no longer advertises prefetch past the appended trampoline/guard layout - rewrite the descriptor in COMGR hotswap, keeping trampoline-related loader policy out of ROCr - scope NOP-sled discovery and rewriting to the owning function so COMGR does not borrow sleds from adjacent code - harden entry-trampoline idempotency decoding and add lit/unit coverage for descriptor preservation, NOP-sled function scoping, and related trampoline encodings Depends on #3182. Review commit: d49d6b2. ## Testing - cmake --build build --target HotswapMCTests hotswap-rewrite -j 32 - ./build/bin/HotswapMCTests - ./build/bin/llvm-lit -sv build/tools/comgr/test-lit --filter hotswap-kernel-entry-trampoline - ./build/bin/llvm-lit -sv build/tools/comgr/test-lit --filter hotswap-nop-sled-function-scope - git diff --check - TheRock gfx1250 final dist validation on mi400 GPU 1: - hipblaslt-test --gtest_filter=*_/matmul_test.matmul/smoke_matmul_gemm_tn_mxf8_dst_bf16_1250_smoke_* passed 16 tests - /mnt/ck-build/bin/example_gemm_multiply_multiply_xdl_fp8 passed; verbose hotswap log showed trampoline prefetch guard bytes appended - ASAN (`build-asan`, `LLVM_USE_SANITIZER=Address`; `ASAN_OPTIONS=detect_leaks=0 LSAN_OPTIONS=detect_leaks=0` due local LSan/ptrace limitation): - built HotswapMCTests, HotswapElfTests, and hotswap-rewrite - ./build-asan/bin/HotswapMCTests passed 39 tests - ./build-asan/bin/HotswapElfTests passed 14 tests - ./build-asan/bin/llvm-lit -sv build-asan/tools/comgr/test-lit --filter hotswap-kernel-entry-trampoline passed 1/1 - ./build-asan/bin/llvm-lit -sv build-asan/tools/comgr/test-lit --filter hotswap-nop-sled-function-scope passed 1/1
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
2 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Summary
Guard the COMGR HotSwap kernel-entry trampoline idempotency matcher with a cheap raw-byte prefix check before invoking LLVM's AMDGPU disassembler.
The existing matcher disassembled
KernelEntryStubStridebytes at every descriptor's current entry point to decide whether the descriptor already targeted an appended entry stub. Real gfx1250 code objects can contain valid kernel entry byte streams that still trip disassembler corner cases before HotSwap rewriting completes. The matcher only needs the full decode for plausible entry stubs, so this change first checks for the generated stub prefix (global_wb; v_nop) and treats all other entries as not already stubbed.This preserves the existing full structural matcher for already-rewritten stubs, so rewrite idempotency remains intact.
Testing
/tmp/llvm-pr3182-build/bin/HotswapMCTests --gtest_filter='BuildKernelEntryTrampoline.*'amd/comgr/test-lit/hotswap-kernel-entry-trampoline.susing the existing COMGR lit helperROCR_VISIBLE_DEVICES=2and captured the reduced kernel-entry bytes used by the new regressionhipblaslt-test --gtest_filter="*_/matmul_test.matmul/smoke_matmul_gemm_tn_mxf8_dst_bf16_1250_smoke_*"passed 16/16 tests.