[DLight][CUDA] Fix undefined TX in GEMV broadcast epilogue#19970
Open
Nanmur wants to merge 1 commit into
Open
Conversation
The GEMV inner-reduction schedule handles broadcast epilogues by splitting the fused epilogue loops and binding the inner loop to threadIdx.x. This path referenced TX, but TX is not defined in the scope or passed into the inner scheduling helper. Use the existing TR tile factor, which is already used for the threadIdx.x direction in this schedule, and add a regression test that covers the broadcast epilogue path.
![gemini-code-assist[bot]](https://avatars.githubusercontent.com/in/956858?s=60&v=4){kind=small}
Contributor
There was a problem hiding this comment.
Code Review
This pull request updates the GPU GEMV schedule in python/tvm/s_tir/dlight/gpu/gemv.py by changing the split factor for the epilogue block from TX to TR. It also adds a new unit test test_gemv_broadcast_epilogue in tests/python/s_tir/dlight/test_gpu_gemv.py to validate this scheduling behavior. There are no review comments, so I have no feedback to provide.
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.
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.
This PR fixes an undefined
TXreference in the DLight GPU GEMV inner-reduction schedule.The broadcast epilogue path splits fused epilogue loops and binds the inner loop to
threadIdx.x, but it referencedTX, which is not defined in the surrounding scope or passed into the helper. This PR uses the existingTRtile factor, which is already used for thethreadIdx.xdirection in this schedule.A regression test is added to cover the GEMV broadcast epilogue path.
Fixes #19969
Tests:
python -m pytest tests/python/s_tir/dlight/test_gpu_gemv.py -qpython -m pytest tests/python/s_tir/dlight/test_gpu_low_batch_gemv.py -qpython -m pytest tests/python/s_tir/dlight/test_gpu_reduction.py -qNote: The branch is based on
apache/tvm:main. Local main-branch test execution on this Windows machine could not be completed because the available compiled TVM library is from a v0.25 build and does not match the latest Python sources; the same tests passed in the matching local v0.25 environment before rebasing the patch to main.