Repository navigation
fix(tensilelite): key the solution cache on every field its predicates read - #12861
Merged
jaopaulolc merged 10 commits intoOct 7, 2026
Merged
Conversation
…s read CachingLibrary answers repeated solution queries from a cache keyed on ContractionProblemGemm, but the key left out fields that solution predicates and the workspace check read. A problem that differed only in one of them was served the solutions cached for another. After an f16-bias heuristic query, a bf16-bias query of the same shape returned the f16 kernels, and hipblasLtMatmul with one of them wrote 0 instead of the bias. Compare and hash the bias data type, bias enum, factorDim, gate-residual type, activation enum, compute type, amaxD type, C == D, beta, alpha and beta restrictions, sparsity, device user arguments, GSU, fallback status, persistent-kernel eligibility and grouped GEMM count, and also compare the bias sizes. The new fields go in a second compare and hash call, which keeps cache hits fast. The grouped GEMM hash now reuses the per-problem hash instead of its own shorter field list. The new CachingLibraryCollision cases cache a solution for a problem that a real predicate accepts and check that a problem it rejects is not served from the cache.
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 #12861 +/- ##
===========================================
- Coverage 71.93% 70.92% -1.01%
===========================================
Files 2863 2946 +83
Lines 472701 485740 +13039
Branches 69827 70150 +323
===========================================
+ Hits 340000 344476 +4476
- Misses 109118 117306 +8188
- Partials 23583 23958 +375
*This pull request uses carry forward flags. Click here to find out more. 🚀 New features to boost your workflow:
|
jaopaulolc
added
the
rocm:gemm algos
…public API Every f16 GEMM is looked up through the one CachingLibrary at the root of the master library, and BiasDataTypeWhiteList is a solution predicate below that cache. On develop, after an f16-bias heuristic query, a bf16-bias query of the same shape returns the cached f16-bias solutions. No shipped f16 GEMM kernel reads a bf16 bias, so matmulIsAlgoSupported rejects them and hipblasLtMatmul with one writes D without the bias. The new case runs the two queries through hipblasLtMatmulAlgoGetHeuristic and checks every bf16-bias result with matmulIsAlgoSupported and with a GEMM whose D must equal the bias.
The solution predicates read beta only as beta == 0 (BetaZero, and BufferLoadOffsetLimitCheck_Beta in 311 of 314 shipped gfx950 library files) and beta == 1 (BetaOne). Keying the raw value gave every other beta its own cache entry, so the key now holds that category instead. betaRestriction does not cover beta == 1: for complex types hipBLASLt stores |beta| in beta() but derives the restriction from the complex value. Compute type is read only by isGemmTypeSame in the GEMM_TYPE_ONLY search of findAllSolutions, which bypasses the cache, so it leaves the key. The beta test now holds the restriction fixed and separates 0, 1 and other values. A new test checks that beta 0.5 and 2.0 share one entry.
A field stays in the cache key only if a predicate in the shipped gfx950 libraries reads it and hipBLASLt can set it apart from the rest of the key. - Beta is keyed as beta == 0, which BufferLoadOffsetLimitCheck_Beta reads in 311 of 314 library files. BetaZero and BetaOne ship in none. - The alpha and beta restrictions, C == D, the grouped GEMM count and persistent-kernel eligibility are dropped: AlphaValue, BetaValue, CEqualsD, SynchronizerSizeCheck and PersistentKernelCheck ship in none. - Sparsity, the gate-residual type and factorDim are dropped because hipBLASLt never sets them. Bias length, biasEnum and device user arguments are dropped because hipBLASLt derives them from D's sizes, the epilogue and grouped_gemm. - The amaxD type stays. hipBLASLt sets it to the compute type, so it also keys computeTypeElementSize(), which WorkspaceCheck reads. The synthetic tests for dropped fields are removed. The beta tests now check that zero and non-zero beta are kept apart and that 0.5, 2.0 and 1.0 share one entry.
AlexBrownAMD
approved these changes
Oct 5, 2026
jaopaulolc
enabled auto-merge (squash)
October 5, 2026 21:52
jaopaulolc
deleted the
users/jolabega/fix-tensilelite-caching-library-key
branch
October 7, 2026 21:44
Contributor
TheRock Submodule Bump ActivityNewest first
|
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
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 : AIHPBLAS-4984
Motivation
hipBLASLt asks TensileLite for the GPU kernels, called solutions, that suit each matrix multiplication (GEMM). TensileLite's
CachingLibrarycaches the answer at the root of the loaded library, keyed on the problem, but the key leaves out fields that per-solution checks read. On gfx950, a bf16-bias heuristic query of an f16 GEMM finds no solutions, but after an f16-bias query of the same shape it gets the cached f16-bias ones.hipblasLtMatmulwith one returns success and writes D without the bias.Technical Details
hipblasLtMatmulAlgoGetHeuristicreachesCachingLibrary::findTopSolutionsthroughgetBestSolutionsandMasterSolutionLibrary::findTopSolutions. A cache hit skips each solution'sBiasDataTypeWhiteListcheck, andhipblasLtMatmulruns a solution by index without it. Bias type does not split libraries; it is checked per solution, below the cache. TheGemmInstanceand grouped GEMM heuristics share the cache;getAllAlgosandmatmulIsAlgoSupporteddo not.Host code only.
Comparisonandstd::hashforContractionProblemGemmalso cover these fields, each with the check that reads it and the number of the 314 shipped gfx950 library files that use it:BiasDataTypeWhiteList, 311;ActivationEnumWhiteList, 266;BufferLoadOffsetLimitCheck_Beta, 311;WorkspaceCheck, 304;AmaxDCheck, 311, andWorkspaceCheck;WorkgroupMappingXCCCheck, 311.Other fields that checks read stay out of the key. hipBLASLt never sets sparsity, the gate-residual type or
factorDim. It derives bias length,biasEnumand device user arguments from D's sizes, the epilogue andgrouped_gemm, which the key already covers. No shipped library checks the alpha and beta restrictions,C == D, the grouped GEMM count or persistent-kernel eligibility.The new fields go in a second compare and hash call to keep hits fast. The grouped GEMM hash now reuses the per-problem hash.
Test Plan
Run
hipblaslt-test --gtest_filter='CachingLibraryCollision.*'.smoke_Bf16BiasHeuristicIsNotServedF16BiasSolutionsruns the f16-then-bf16 queries through the public API. Every bf16 result must passmatmulIsAlgoSupportedand give D equal to the bias. Fivesmoke_*PartOfKeytests check that a problem a real predicate rejects is not served a solution cached for one it accepts. Two tests check that equal keys still hit: identical problems, and beta 0.5, 2.0 and 1.0.Test Result
On an MI355X (gfx950) with a TheRock gfx950 device library, the bf16 query on
developreturned 4 solutions.matmulIsAlgoSupportedrejected all 4, and each left every element of D wrong. Here it returns none. The public-API test and the 5smoke_*PartOfKeytests fail ondevelop; all 12 suite tests pass here. Median heuristic latency over 3 rounds, small library: warm 1.70–1.71 µs before, 1.71–1.72 µs after; cold about 94–99 µs both. Automated GitHub checks are pending.Submission Checklist
Risk level
Low (2/5): only cache lookups change. Problems differing in a predicate-read field get separate entries, so results match an uncached search.