feat(fused-mlp-megakernels): portable SIMT GeGLU fallback for SM110 (Thor) - #5
Open
shideqin wants to merge 1 commit into
Open
Conversation
…or SM110 The SM100-family CUTLASS 4.0 megakernel uses tcgen05/TMA descriptor paths that assert at runtime on sm_110a (Thor). Add a block-tiled pure-SIMT FMA reference (32x32 tile, shared-memory K chunks) computing the same fusion and route sm_110a / FLASHRT_FORCE_SIMT to it in torch_binding.cpp; sm_100/sm_103 keep the CUTLASS megakernel. Record the installed-artifact Thor validation (7/7 rows, cos>=0.99999976) and benchmark (85.6ms on M768,N16384,K2048, ~22x faster than the one-thread-per-output reference).
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.
Summary
The SM100-family CUTLASS 4.0 megakernel (
fp16_geglu_fused_out) uses tcgen05/TMA descriptor paths that assert at runtime onsm_110a(Thor). Add a block-tiled pure-SIMT FMA reference (portable_geglu_simt.cu, 32x32 tile, shared-memory K chunks) computing the same fusion and routesm_110a/FLASHRT_FORCE_SIMTto it intorch_binding.cpp;sm_100/sm_103keep the CUTLASS megakernel.Validation (real hardware, NVIDIA Thor sm_110a)
torch.compile(fullgraph=True)+ CUDA Graph, worst rowM768,N16384,K2048max=0.00146 cos=0.99999988.M768,N16384,K2048). Remains a compatibility path (~52x slower than eager) since no native Thor backend exists.Per AGENTS.md the fallback is explicit (capability-gated +
FLASHRT_FORCE_SIMT) and does not replace the qualified SM100/103 megakernel.