Repository navigation
[S-TIR] Fix software pipeline offsets for legacy MMA intrinsics - #19742
Merged
Merged
Conversation
Contributor
There was a problem hiding this comment.
Code Review
This pull request adds support for legacy operations (such as legacy PTX ldmatrix and MMA ops) in the software pipeline transformation by introducing a helper function IsOp and updating the buffer access rewriter. The feedback suggests declaring the IsOp helper function as static to ensure internal linkage and prevent potential linker collisions.
Important
The consumer version of Gemini Code Assist on GitHub is being sunset. Starting June 18, 2026, new organization installations will be blocked, and all code review activity will officially cease on July 17, 2026.
For more details on the timeline and next steps, please review the Help Documentation.
tlopex
force-pushed
the
investigate-s-tir-nested-pipeline-mma
branch
2 times, most recently
from
June 12, 2026 03:03
adbd870 to
3705f39
Compare
tlopex
force-pushed
the
investigate-s-tir-nested-pipeline-mma
branch
from
June 12, 2026 03:06
3705f39 to
e56ca66
Compare
spectrometerHBH
approved these changes
Jun 12, 2026
MasterJH5574
pushed a commit
to MasterJH5574/tvm
that referenced
this pull request
Jun 15, 2026
…he#19742) This pr fixes `InjectSoftwarePipeline` to rewrite opaque buffer offsets for legacy PTX MMA intrinsics, including `ptx_ldmatrix_legacy`, `ptx_mma_legacy`, `mma_store_legacy`, and `mma_fill_legacy`. ### Failure `test_async_nested_pipeline_mma_gemm_ideal_annotation` failed with a real numerical mismatch during the final GEMM result check: ```text Mismatched elements: 6744298 / 16777216 (40.2%) Max absolute difference: 6.706421 Max relative difference: 0.00653978 ``` A diagnostic run with a fixed seed reproduced the same class of error: ```text nested_full bad 6826830 / 16777216 max_abs 7.0405273 mean_abs 0.9874781 ``` The adjacent simple pipeline test passed, and disabling async copy did not change the mismatch, so this was not a cp.async lowering issue, tolerance noise, or GPU flakiness. ### Root Cause `InjectSoftwarePipeline` may add a leading version dimension to pipeline buffers. Normal buffer loads/stores and newer opaque PTX intrinsics already had their offsets rewritten to include the pipeline version slot. However, the legacy MMA intrinsics were not covered. In this test, the warp buffers became multi-versioned, but the legacy `ldmatrix`/`mma` offsets still pointed to the original slot. As a result, the second `ldmatrix` stage overwrote the first warp fragment, and both MMA stages read the same fragment. That skipped one K fragment and duplicated another, producing the numerical mismatch above. (cherry picked from commit 2a16d98)
MasterJH5574
pushed a commit
to MasterJH5574/tvm
that referenced
this pull request
Jun 15, 2026
…he#19742) This pr fixes `InjectSoftwarePipeline` to rewrite opaque buffer offsets for legacy PTX MMA intrinsics, including `ptx_ldmatrix_legacy`, `ptx_mma_legacy`, `mma_store_legacy`, and `mma_fill_legacy`. ### Failure `test_async_nested_pipeline_mma_gemm_ideal_annotation` failed with a real numerical mismatch during the final GEMM result check: ```text Mismatched elements: 6744298 / 16777216 (40.2%) Max absolute difference: 6.706421 Max relative difference: 0.00653978 ``` A diagnostic run with a fixed seed reproduced the same class of error: ```text nested_full bad 6826830 / 16777216 max_abs 7.0405273 mean_abs 0.9874781 ``` The adjacent simple pipeline test passed, and disabling async copy did not change the mismatch, so this was not a cp.async lowering issue, tolerance noise, or GPU flakiness. ### Root Cause `InjectSoftwarePipeline` may add a leading version dimension to pipeline buffers. Normal buffer loads/stores and newer opaque PTX intrinsics already had their offsets rewritten to include the pipeline version slot. However, the legacy MMA intrinsics were not covered. In this test, the warp buffers became multi-versioned, but the legacy `ldmatrix`/`mma` offsets still pointed to the original slot. As a result, the second `ldmatrix` stage overwrote the first warp fragment, and both MMA stages read the same fragment. That skipped one K fragment and duplicated another, producing the numerical mismatch above. (cherry picked from commit 2a16d98)
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.
This pr fixes
InjectSoftwarePipelineto rewrite opaque buffer offsets for legacy PTX MMA intrinsics, includingptx_ldmatrix_legacy,ptx_mma_legacy,mma_store_legacy, andmma_fill_legacy.Failure
test_async_nested_pipeline_mma_gemm_ideal_annotationfailed with a real numerical mismatch during the final GEMM result check:A diagnostic run with a fixed seed reproduced the same class of error:
The adjacent simple pipeline test passed, and disabling async copy did not change the mismatch, so this was not a cp.async lowering issue, tolerance noise, or GPU flakiness.
Root Cause
InjectSoftwarePipelinemay add a leading version dimension to pipeline buffers. Normal buffer loads/stores and newer opaque PTX intrinsics already had their offsets rewritten to include the pipeline version slot.However, the legacy MMA intrinsics were not covered. In this test, the warp buffers became multi-versioned, but the legacy
ldmatrix/mmaoffsets still pointed to the original slot. As a result, the secondldmatrixstage overwrote the first warp fragment, and both MMA stages read the same fragment. That skipped one K fragment and duplicated another, producing the numerical mismatch above.