Fix hipErrorNotFound for helper kernels under I-cache rotation.#9266
Open
geotseng-amd wants to merge 1 commit into
Open
Fix hipErrorNotFound for helper kernels under I-cache rotation.#9266geotseng-amd wants to merge 1 commit into
geotseng-amd wants to merge 1 commit into
Conversation
github-actions
Bot
added
project: hipblaslt
project: hipsparselt
ci:hipsparselt-fast
labels
❌ PR Check — Action Required
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
|
🚫 Please fix the failed policies before requesting reviews. The following policy checks failed:
The |
Codecov Report✅ All modified and coverable lines are covered by tests. ❌ Your project status has failed because the head coverage (76.92%) is below the target coverage (80.00%). You can increase the head coverage or adjust the target coverage. Additional details and impacted files@@ Coverage Diff @@
## develop #9266 +/- ##
===========================================
- Coverage 64.74% 64.73% -0.00%
===========================================
Files 2675 2675
Lines 420248 420303 +55
Branches 62446 62456 +10
===========================================
+ Hits 272051 272080 +29
- Misses 127699 127722 +23
- Partials 20498 20501 +3
*This pull request uses carry forward flags. Click here to find out more. 🚀 New features to boost your workflow:
|
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
3 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.
This PR is credited to @Jinp800125
Motivation
Benchmarking with I-cache rotation enabled (
--icache-rotate-copies) crashed at kernel launch withhipErrorNotFound("named symbol not found"), e.g. for the helper kernelCijk_SB_PostGSU2_VW4. I-cache rotation duplicates only the main--code-object(.co) modules into rotation slots; the helper kernel bundle (Kernels.so-000-<arch>.hsaco, containing PostGSU/betaOnly/reduction/conversion kernels) is loaded only once. When a launch landed on a rotation slot other than copy 0, the helper symbol could not be resolved and the run aborted.Technical Details
HipSolutionAdapter::getKernel: added a fallback that searches the basem_moduleslist when a symbol is not found in the selected rotation copy's modules. This prevents launch failures for any module that isn't duplicated across rotation slots (safety net).HipSolutionAdapter::loadHelperKernelExtraCopies(int extraCopies): new method (declared inHipSolutionAdapter.hpp) that mirrors the helper bundle (Kernels.so-000-<arch>.hsaco) into every rotation slot, analogous to howloadCodeObjectFileExtraCopiesduplicates main.cofiles. It is idempotent (guarded bym_extraModuleCopies[0].size() >= m_modules.size()) and a safe no-op when no helper bundle is present.client/main.cpp: wiredloadHelperKernelExtraCopiesinto both extra-copy paths, each placed after the helper is present inm_modules:--icache-rotate-copies=-1): called right after the main-copy loop inside the per-solution setup.--icache-rotate-copies=N, N>0): the main copies load before lazy-load init, so the helper is mirrored in a new block immediately afterinitializeLazyLoading.Net effect: helper kernels now truly rotate across I-cache slots instead of resolving only from copy 0, with the
getKernelfallback remaining as a defensive backstop.Test Plan
tensilelite-host(adapter changes) andtensilelite-client(main.cppchanges).--icache-rotate-copies=-1and an explicit--icache-rotate-copies=N) to confirm the launch no longer fails.Test Result
main.cpprecompiled andtensilelite-clientlinked successfully; no linter errors.Submission Checklist