Repository navigation
[Fix] Register the tcgen05.mma emitter for the sm_100 family - #169
Merged
Merged
Conversation
The MMA emitter was registered only for nvgpu_sm100a. On a B300 the current target is nvgpu_sm103a, which does not support sm100a, so every script using tcgen05.mma (all blackwell_matmul examples) failed with "Failed to find emitter ... Tcgen05MmaSSInst". .kind::f16, .kind::tf32 and .kind::f8f6f4 are family features (ptxas accepts them for sm_100a/103a/100f/103f), so register the emitter for nvgpu_sm100f. .kind::i8 is sm_100a only, so reject it on other targets with a clear error. Add requires.nvgpu_sm100f and use it for the tcgen05 instruction tests (mma, alloc/free, load/store, copy). With requires.nvgpu_sm100a they fell back to compile-only on 10.3 and never ran there. Fixes NVIDIA#168 Co-authored-by: Claude Opus 5.5 <noreply@anthropic.com> Signed-off-by: 0z5a <dezhen.lu@student.uni-tuebingen.de>
Member
|
/ok to test b6d2085 |
yaoyaoding
enabled auto-merge (squash)
November 4, 2026 17:23
Member
|
@0z5a thank you! |
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.
Summary
Fixes #168.
TMemoryMmaSSEmitterwas registered only fornvgpu_sm100a. On a B300 the current target isnvgpu_sm103a, andsm103a.supports(sm100a)is False, so every script that usestcgen05.mmafailed withFailed to find emitter ... Tcgen05MmaSSInst (registered targets: nvgpu/sm100a). That includes allexamples/blackwell_matmulscripts.Changes:
backends/emitters/cuda/tcgen05/mma.py: register the emitter fornvgpu_sm100f.sm100a,sm103a,sm100fandsm103fall support it..kind::f16,.kind::tf32and.kind::f8f6f4are family features; ptxas (CUDA 13.3) accepts them for all four targets..kind::i8is the exception. ptxas accepts it only forsm_100a, soget_mma_kindnow raisesCodeGenerationFailedfor int8 operands unless the target supportssm100a. Todayget_inst_mnkdoes not implement I8 yet, so this only changes the error message. It keeps the family registration from claiming i8 on 10.3 later.testing/_requires.py: addrequires.nvgpu_sm100f.tests/instructions/test_tcgen05_{mma,alloc_free,load_store,copy}.py: userequires.nvgpu_sm100f. Withrequires.nvgpu_sm100a, a 10.3 GPU falls back to compile-only underscope(nvgpu_sm100a), so these tests never ran there. The alloc/load/store/copy emitters are registered for basesm100and already work on 10.3. Only the test gate kept them from running.Testing
B300 (CC 10.3), CUDA 13.3, PyTorch 2.14.1+cu130, tilus
mainat 2047456.Build for each target (
scope(...)+.compile(...)): the 7blackwell_matmulexamples (1024³ fp16) andTcgen05MmaExamplefrom the test with f16 and f8 operands.sm103asm100asm100fint8 operands: before, the error is
NotImplementedErrorfromget_inst_mnkonsm100aand "Failed to find emitter" onsm103a. After, it is unchanged onsm100a, and onsm103a/sm100fit istcgen05.mma .kind::i8 is not supported on nvgpu/sm103a.Run on the B300:
One B300 GPU,
mainvs. this PR:mainexamples/blackwell_matmul/matmul_v0.py…v6.pymain(bench=False)(8192³ fp16,assert_closevs. torch)tests/instructions/test_tcgen05_{mma,alloc_free,load_store,copy}.py+tests/ir/test_targets.pyThe 4 xfails are existing
xfailcases intest_tcgen05_copy.py, the same on both sides.ruff checkandruff format --check(0.11.0) pass on the changed files.AI assistance
Claude (Anthropic) helped find this, write the change, and run the tests; the commit carries a
Co-authored-bytrailer for it. I reviewed the change and the results.