tlopex opened a new pull request, #19742: URL: https://github.com/apache/tvm/pull/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. -- This is an automated message from the Apache Git Service. To respond to the message, please log on to GitHub and use the URL above to go to the specific comment. To unsubscribe, e-mail: [email protected] For queries about this service, please contact Infrastructure at: [email protected] --------------------------------------------------------------------- To unsubscribe, e-mail: [email protected] For additional commands, e-mail: [email protected]
