Enable tcgen05 blockscaled ops on Thor SM110#3283
Open
xiangg-nv wants to merge 1 commit into
Open
Conversation
Collaborator
brandon-yujie-sun
left a comment
brandon-yujie-sun
left a comment
There was a problem hiding this comment.
Thanks for contributing the codes! Would it be possible to also add some test cases?
Edge-LLM NvFP4 MoE CuTeDSL kernels on Thor use tcgen05 blockscaled MMA and SMEM-to-TMEM scale-factor copies. The existing checks only admitted the SM100/SM103 paths, so source-built CuTeDSL rejected SM110. Admit Thor's blockscaled MMA arch aliases sm_101a and sm_110a, and allow the SM110f family for S2T tcgen05 copy ops. Validation: - git diff --check - python3 -m py_compile python/CuTeDSL/cutlass/cute/nvgpu/tcgen05/mma.py python/CuTeDSL/cutlass/cute/nvgpu/tcgen05/copy.py - DKG grouped_blockscaled_gemm.py documented 4-group example on Thor SM110: PASS - Edge-LLM nvfp4_moe AOT for sm_110/aarch64: 12/12 variants PASS
xiangg-nv
force-pushed
the
feature/cutedsl-thor-sm110
branch
from
e19678b to
da87ebf
Compare
Author
Not sure which kind of test case should I add, it seems that we don't have any SM110 test cases, can you help to resolve this? |
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
2 participants
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.
Edge-LLM NvFP4 MoE CuTeDSL kernels on Thor use tcgen05 blockscaled MMA and SMEM-to-TMEM scale-factor copies. (Currently use the patch to WR)
The existing checks only admitted the SM100/SM103 paths, so source-built CuTeDSL rejected SM110.
Admit Thor's blockscaled MMA arch aliases sm_101a and sm_110a, and allow the SM110f family for S2T tcgen05 copy ops.
Validation:
git diff --check
python3 -m py_compile python/CuTeDSL/cutlass/cute/nvgpu/tcgen05/mma.py python/CuTeDSL/cutlass/cute/nvgpu/tcgen05/copy.py
grouped_blockscaled_gemm.py documented 4-group example on Thor SM110: PASS
Edge-LLM nvfp4_moe AOT for sm_110/aarch64: 12/12 variants PASS