This is an automated email from the ASF dual-hosted git repository.
spectrometerHBH pushed a commit to branch main
in repository https://gitbox.apache.org/repos/asf/tvm.git
The following commit(s) were added to refs/heads/main by this push:
new 8ddd97d106 [TIRx][Test] Gate tcgen05.mma collector certification on
nvcc >= 13.4 (#20270)
8ddd97d106 is described below
commit 8ddd97d106810a1db9238fb12e1981f23f43857e
Author: Bohan Hou <[email protected]>
AuthorDate: Fri Sep 4 19:58:48 2026 -0400
[TIRx][Test] Gate tcgen05.mma collector certification on nvcc >= 13.4
(#20270)
## Summary
`test_ptx_tcgen05_mma_block_size_form` assembles the collector-qualified
`tcgen05.mma ... block_scale` forms at `sm_107f`. Those are PTX ISA 9.4
forms (#20266) and need the CUDA 13.4 ptxas; on a CUDA 13.2 host the
test fails with `Unknown modifier '.collector::b::fill'` instead of
skipping.
This moves that half into
`test_ptx_tcgen05_mma_block_size_collector_form`, guarded by
`env.has_nvcc_version(13, 4)`, the same gate the whole-table
certification tests already use. The `sm_100a` block32 certification and
the `mxf4` block16 legality check stay in the original test and keep
running on older toolchains.
## Testing
On a CUDA 13.2 host (B200):
```
tests/python/tirx/codegen/test_ptx_dialect.py -k tcgen05_mma_block_size
3 passed, 1 skipped (collector form: need nvcc >= 13.4)
```
The full `tests/python/tirx/` suite on the same host previously had this
as its only failure (2749 passed, 321 skipped).
---
tests/python/tirx/codegen/test_ptx_dialect.py | 35 +++++++++++++++++----------
1 file changed, 22 insertions(+), 13 deletions(-)
diff --git a/tests/python/tirx/codegen/test_ptx_dialect.py
b/tests/python/tirx/codegen/test_ptx_dialect.py
index 2304824933..4290125f50 100644
--- a/tests/python/tirx/codegen/test_ptx_dialect.py
+++ b/tests/python/tirx/codegen/test_ptx_dialect.py
@@ -2354,6 +2354,28 @@ def test_ptx_tcgen05_mma_block_size_form():
assert "tcgen05.mma.cta_group::1.kind::mxf4.block_scale.block32" in src
_assert_ptxas_ok(src, arch="sm_100a")
+ with pytest.raises((ValueError, tvm.error.DiagnosticError),
match="mxf4.*block32"):
+
+ @T.prim_func
+ def invalid_mxf4_block16():
+ T.device_entry()
+ tmem = T.local_scalar("uint32")
+ desc = T.local_scalar("uint64")
+ idesc = T.local_scalar("uint32")
+ flag = T.local_scalar("uint32")
+ T.ptx["tcgen05.mma.cta_group::1.kind::mxf4.block_scale.block16"](
+ tmem, desc, desc, idesc, tmem, tmem, T.ptx.pred(flag)
+ )
+
+
[email protected](
+ not env.has_nvcc_version(13, 4),
+ reason="collector-qualified block_scale MMA is a PTX 9.4 form; need nvcc
>= 13.4",
+)
+@requires_nvcc
+def test_ptx_tcgen05_mma_block_size_collector_form():
+ """PTX 9.4 collector qualifiers on block-scaled MMA certify at their
sm_107f floor."""
+
@T.prim_func
def sm107_collector_kernel(a_ptr: T.handle):
A = T.match_buffer(a_ptr, (32,), "uint32")
@@ -2388,19 +2410,6 @@ def test_ptx_tcgen05_mma_block_size_form():
assert ts_collector_opcode in collector_src
_assert_ptxas_ok(collector_src, arch="sm_107f")
- with pytest.raises((ValueError, tvm.error.DiagnosticError),
match="mxf4.*block32"):
-
- @T.prim_func
- def invalid_mxf4_block16():
- T.device_entry()
- tmem = T.local_scalar("uint32")
- desc = T.local_scalar("uint64")
- idesc = T.local_scalar("uint32")
- flag = T.local_scalar("uint32")
- T.ptx["tcgen05.mma.cta_group::1.kind::mxf4.block_scale.block16"](
- tmem, desc, desc, idesc, tmem, tmem, T.ptx.pred(flag)
- )
-
def test_ptx_tcgen05_mma_block_size_collector_legality():
from tvm.backend.cuda.ptx.table import TABLE, tokens_for
