[hipBLASLt] Make generated code-object linking deterministic - #11129
Open
newling wants to merge 1 commit into
Open
[hipBLASLt] Make generated code-object linking deterministic#11129newling wants to merge 1 commit into
newling wants to merge 1 commit into
Conversation
github-actions
Bot
added
project: hipblaslt
project: hipsparselt
ci:hipsparselt-fast
labels
✅ All Checks Passed — Ready for Review
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
|
🎉 All checks passed! This PR is ready for review. |
Codecov Report✅ All modified and coverable lines are covered by tests. Additional details and impacted files@@ Coverage Diff @@
## develop #11129 +/- ##
===========================================
- Coverage 69.65% 69.65% -0.00%
===========================================
Files 2789 2789
Lines 459421 459422 +1
Branches 67739 67740 +1
===========================================
- Hits 319988 319987 -1
- Misses 116264 116266 +2
Partials 23169 23169
*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
1 participant
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.
JIRA ID: ROCM-28430
What this changes
hipBLASLt includes a large collection of precompiled GPU kernels. During the
Tensile library-generation step, each kernel is first assembled into an object
file (
.o). The generator then combines—or “links”—many of those object filesinto a GPU code object (
.co).A
.cois the binary container that hipBLASLt loads at runtime to execute itsGPU kernels. In the usual compressed build, it contains a GPU-specific ELF
binary together with the information needed to load it for a particular GPU
architecture.
The order in which object files are given to the linker matters. It determines
where the individual kernel functions are placed inside the resulting binary,
and therefore affects the final
.cobytes and build ID.For explicitly grouped code objects, Tensile currently collects the object-file
paths in a Python
setand passes that set directly to the linker. Python setsdo not have a stable iteration order: their order can change with the Python
hash seed, the object paths, or other details of the build environment.
That means two builds can start from the same source and the same collection of
kernel object files, but link those files in different orders and produce
different
.cobinaries.This PR sorts the object-file paths immediately before invoking the linker.
With a fixed set of object files, the linker now receives the same ordered list
on every build.
Why this is useful
Deterministic linking makes generated hipBLASLt libraries easier to:
set iteration.
This does intentionally establish one canonical physical order for the kernel
functions inside each generated
.co. It does not change which kernels areincluded or their assembled instruction bytes.
The build-time cost is negligible: sorting 1,919 real object paths took
approximately 0.09 ms, compared with 0.52–0.55 seconds for linking and
4.32–4.36 seconds for compression.
Testing
The new characterization test constructs the same explicit code-object group
using two different kernel input orders. It verifies that both cases give the
linker the same lexicographically sorted list and that the linker receives a
concrete
list, not a set iterator.Validation performed:
origin/develop.ToolchainAssembly/test_assembly_char.pyfile passes:PYTHONHASHSEED=0,1, and42.git diff --checkpasses.Scope
This is a build-reproducibility fix. It makes the physical layout of generated
code objects deterministic, but it does not claim that one particular layout
changes or fixes numerical behavior.