AMDGPU: scope hotswap trampolines and clamp entry prefetch by harsh-amd · Pull Request #3195 · ROCm/llvm-project

harsh-amd · 2026-07-05T01:33:30Z

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

chinmaydd · 2026-07-06T17:46:31Z

LIT tests are failing and seem genuine (?)

harsh-amd · 2026-07-07T23:39:53Z

LIT tests are failing and seem genuine (?)

this should be fixed now. please take another look.

chinmaydd · 2026-07-08T17:00:48Z

-            7u);
+            KernelEntryStubInstPrefLines);
+  EXPECT_EQ(
+      AMDHSA_BITS_GET(OutRsrc3, hsa::COMPUTE_PGM_RSRC3_GFX11_INST_PREF_SIZE),


Should test with INST_PREF_SIZE for gfx12 rather than 11

I have removed the gfx11 check. this now only checks the gfx12 inst_pref field.

chinmaydd · 2026-07-08T17:08:51Z

+  uint32_t MaxInstPrefLines = 0;
+  for (const KernelDescriptorInfo &KD : Descriptors) {
+    std::optional<uint32_t> InstPrefLines =
+        Elf.getKernelDescriptorInstPrefSize(KD.KernelName, LS.Cpu);


If I understand correctly, we have 2 different sources for the prefetch. Is that intentional ?

this is intentional. deferred trampolines use the original kernel descriptor prefetch for the trailing guard. entry stubs clamp their own descriptor value separately. I added comments to make this clearer.

chinmaydd · 2026-07-08T17:10:44Z

-    Work.push_back(KD);
+    uint32_t StubInstPrefLines =
+        std::min(*InstPrefLines, KernelEntryStubInstPrefLines);
+    MaxInstPrefLines = std::max(MaxInstPrefLines, StubInstPrefLines);


If I understand correctly, there are now 2 sources for the prefetch value ? Can they disagree?

they can disagree intentionally. the original value may be larger, but the entry stub descriptor is clamped to the 256-byte stub stride. I renamed this to make that explicit.

chinmaydd · 2026-07-08T17:11:31Z

+    if (Sled.WritePos > Sled.FunctionEnd ||
+        Needed > Sled.FunctionEnd - Sled.WritePos)


Check may be redundant ?

I have folded this into one usable-end check.

chinmaydd · 2026-07-08T17:12:02Z

+    log() << "hotswap: error: trampoline prefetch guard size " << GuardBytes
+          << " exceeds size_t.\n";
+    return false;
+  }


Probably redundant

I have removed it. the guard size is bounded by the descriptor field here.

chinmaydd · 2026-07-08T17:13:22Z

+TEST(FindNearestSled, SkipsSledsOutsideInstructionFunctionRange) {
+  std::vector<NopSled> Sleds;
+  Sleds.push_back({0, 32, 0, 0, 32});
+  Sleds.push_back({96, 128, 96, 96, 160});


// {Start, End, WritePos, FunctionStart, FunctionEnd} Sleds.push_back({/*Start=*/96, /*End=*/128, /*WritePos=*/96, /*FunctionStart=*/96, /*FunctionEnd=*/160});

I have added field comments for the sled initializers.

harsh-amd requested review from chinmaydd and lamb-j as code owners July 5, 2026 01:33

github-actions Bot added comgr Related to Code Object Manager hotswap Related to the Comgr Hotswap feature labels Jul 5, 2026

ammarwa mentioned this pull request Jul 5, 2026

fix(rocr): add gfx12 inst prefetch descriptor bits ROCm/rocm-systems#8127

Closed

harsh-amd changed the title ~~AMDGPU: clamp hotswap entry-stub inst prefetch~~ AMDGPU: scope hotswap trampolines and clamp entry prefetch Jul 5, 2026

harsh-amd force-pushed the comgr-entry-trampoline-inst-pref-hotswap branch 2 times, most recently from 6f1f253 to cff7756 Compare July 7, 2026 23:36

chinmaydd reviewed Jul 8, 2026

View reviewed changes

harsh-amd force-pushed the comgr-entry-trampoline-inst-pref-hotswap branch from cff7756 to 627fea4 Compare July 8, 2026 18:21

chinmaydd approved these changes Jul 8, 2026

View reviewed changes

harsh-amd enabled auto-merge (squash) July 8, 2026 18:25

harsh-amd added 2 commits July 8, 2026 14:43

AMDGPU: clamp hotswap entry-stub inst prefetch

cbb088f

AMDGPU: keep hotswap NOP sleds within function bounds

737d5a5

harsh-amd force-pushed the comgr-entry-trampoline-inst-pref-hotswap branch from 627fea4 to 737d5a5 Compare July 8, 2026 22:25

harsh-amd merged commit fd34dbb into ROCm:amd-staging Jul 9, 2026
28 checks passed

		if (Sled.WritePos > Sled.FunctionEnd \|\|
		Needed > Sled.FunctionEnd - Sled.WritePos)

Uh oh!

Conversation

harsh-amd commented Jul 5, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Summary

Testing

Uh oh!

chinmaydd commented Jul 6, 2026

Uh oh!

harsh-amd commented Jul 7, 2026

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

chinmaydd Jul 8, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

harsh-amd commented Jul 5, 2026 •

edited

Loading

chinmaydd Jul 8, 2026 •

edited

Loading