Conversation
…calculation on gfx1250 TensorDataMoverLoad.calculateStartAddr() / calculateStartAddrWaveSeparated() (Components/TensorDataMover.py) had two independent address-calc bugs for TDM (TDMInst=3) sparse (2:4) kernels, neither caught by existing spmm_tdm_*.yaml regression configs (all GlobalSplitU=[1], Batch=1): - GSU offset: gsuOffsetBytes used the full (uncompressed) DepthU for the sparse-tracked A/B operand and for Metadata, instead of the halved DepthU used everywhere else in this file for those tensors. Each GSU group beyond group 0 advanced by double the correct stride, corrupting ~98% of output elements for GlobalSplitU>1. - Batch offset: Metadata's per-batch pointer advance scaled the host-provided StrideXX by bpe=0.25 (same as A/B), but per ContractionProblem::normalizeSparse, Metadata's stride is already in metadata-native Int8 byte units (bpe=1). Scaling by 0.25 again quartered the offset, so batches beyond 0 read metadata from the wrong location. Fixed both by halving depthU for the sparse-tracked operand/Metadata before computing gsuOffsetBytes, and by using batchBpe=1 for Metadata when scaling the batch stride, in both address-calc methods. Extended spmm_tdm_all.yaml with two new BenchmarkProblems blocks (Sparse=1 and Sparse=2, NN, fp16) covering GlobalSplitU=[1,2] combined with Batch=[1,2] -- the previously-untested combination that exposed both bugs. Verified on gfx1250 hardware: full sparse/gfx1250 TDM regression suite (spmm_tdm_all, spmm_tdm_multicast, spmm_tdm_gl2prefetch, spmm_tdm_ldspadding, spmm_tdm_f16_transposes, tdm_general_batched_gfx1250, tdm_gfx1250, tdm_multicast_gfx1250, tdmpluslds_tdm_gfx1250) passes with no regressions; new GSU/batch blocks pass where they previously failed. Co-authored-by: Cursor <cursoragent@cursor.com>
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. |
Contributor
There was a problem hiding this comment.
Copilot review overview
🔵 Needs a closer look
One or more issues must be addressed before approval.
Review effort: Lite
Findings: None
What changed in this PR
Fixes gfx1250 sparse TDM address calculations for GSU and strided batches.
Changes:
- Corrects compressed
DepthUscaling for sparse operands and metadata. - Uses byte-native metadata batch strides.
- Adds Sparse A/B GSU-and-batch regression configurations.
| File | Description |
|---|---|
| projects/hipblaslt/tensilelite/Tensile/Tests/common/sparse/gfx1250/spmm_tdm_all.yaml | Updated as part of this pull request. |
| projects/hipblaslt/tensilelite/Tensile/Components/TensorDataMover.py | Updated as part of this pull request. |
💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.
Codecov Report❌ Patch coverage is
Additional details and impacted files@@ Coverage Diff @@
## develop #12339 +/- ##
===========================================
+ Coverage 70.71% 70.75% +0.04%
===========================================
Files 2847 2849 +2
Lines 466776 468153 +1377
Branches 68835 69160 +325
===========================================
+ Hits 330060 331234 +1174
- Misses 112771 112925 +154
- Partials 23945 23994 +49
*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
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 - AIHPSPAR-278
AIHPSPAR-278
Motivation
TDM (
TDMInst=3) sparse (2:4) GEMM kernels on gfx1250 produced numerically wrong results in two independent, previously-untested scenarios:TDMInst=3+GlobalSplitU > 1: ~98% of output elements incorrect.TDMInst=3(or any TDM) +StridedBatchedcount > 1: metadata read from the wrong batch offset, corrupting every batch beyond batch 0.Both
TensorDataMoverLoad.calculateStartAddr()andcalculateStartAddrWaveSeparated()(Components/TensorDataMover.py) treat the sparse-tracked operand's/Metadata'sDepthU/bpescaling inconsistently between the tile-address path and the GSU/batch-offset path. Existingspmm_tdm_*.yamlregression configs never caught this because they all usedGlobalSplitU=[1]andBatch=1.Technical Details
gsuOffsetByteswas computed asdepthU * bpeusing the full (uncompressed)DepthUfor the sparse-tracked A/B operand and forMetadata, while every other address computation in this file (andinitTDMDescriptor/tdmSplitLdsBoundary) already halvesDepthUto account for the 2:4 compression. Each GSU group beyond group 0 therefore advanced by double the correct stride, reading from the wrong K range. Fixed by halvingdepthUbefore computinggsuOffsetByteswhenever(Sparse==1 and isA) or (Sparse==2 and isB) or isM, in bothcalculateStartAddr()andcalculateStartAddrWaveSeparated().StrideXX(from the host) bybpe=0.25for theMetadatatensor, same as for A/B. But perContractionProblem::normalizeSparse, Metadata's host-provided stride is already expressed in metadata-native Int8 byte units (bpe=1), unlike A/B whose strides are in units of their own element type. Scaling by0.25again quartered the per-batch metadata offset, so batches beyond 0 read metadata from the wrong location. Fixed by usingbatchBpe=1forMetadata(bpe for A/B unchanged), in both address-calc methods.spmm_tdm_all.yaml(gfx1250) with two newBenchmarkProblemsblocks (NN, fp16,Sparse=1andSparse=2) coveringGlobalSplitU=[1,2]combined withBatch=[1,2]— the previously-untested combination that exposed both bugs.Test Plan
spmm_tdm_all.yamlon real gfx1250 hardware.spmm_tdm_all,spmm_tdm_multicast,spmm_tdm_gl2prefetch,spmm_tdm_ldspadding,spmm_tdm_f16_transposes,tdm_general_batched_gfx1250,tdm_gfx1250,tdm_multicast_gfx1250,tdmpluslds_tdm_gfx1250.Tensile/Tests -m common -k "gfx1250 and not streamk").Test Result
Sparse=1andSparse=2blocks PASSED forGlobalSplitU=[1,2]xBatch=[1,2].-k "gfx1250 and not streamk"): in progress, 0 failures so far — will update this section with final counts once complete.JIRA: https://amd-hub.atlassian.net/browse/AIHPSPAR-278
Submission Checklist