spectrometerHBH opened a new pull request, #20266:
URL: https://github.com/apache/tvm/pull/20266
## Summary
Follow-up to #20261 for two `tcgen05` forms in the `T.ptx` table
(`python/tvm/backend/cuda/ptx/table.py`), per the CUDA 13.4 / PTX ISA 9.4
manual.
- **`tcgen05.mma` block-scaled with explicit block size**
(`tcgen05_mma_block_scale_block_{ss,ts}`): the entries move from the base table
into `_PTX_94_ENTRIES` and gain optional `collector_a` / `collector_b` slots,
so each entry owns both the no-collector (ISA 9.7.18.10.10.1 syntax form 2) and
collector-qualified (form 4) variants, which share the same operand shape per A
location. Table 68 block sizes are unchanged (block32 for `kind::mxf8f6f4` and
`kind::mxf4`, block16/block32 for `kind::mxf4nvf4`). The check additionally
rejects collector B without collector A. The entries certify at sm_107f because
collector B is SM107-only; the no-collector and collector-A-only forms keep
their sm_100f floor.
- **`tcgen05.ld.red`**: `redval` is now bound as an output operand
(`rw="w"`). ISA 9.7.18.8.3 writes the reduction result into it; with an input
binding the kernel compiles but never observes the hardware max (measured on
GB300).
## Tests (`tests/python/tirx/codegen/test_ptx_dialect.py`)
- `test_ptx_tcgen05_mma_block_size_collector_legality`: collector A/B
legality and ptxas certification of the collector forms at sm_107f.
- `test_ptx_tcgen05_mma_block_size_no_b_certifies_at_sm100f`: no-collector
and collector-A-only forms still certify at sm_100f.
- `test_ptx_94_sm107_arch_floors_and_delta` (renamed from
`test_ptx_94_family_specific_arch_floors_and_delta`): entries owning SM107
variants certify at sm_107f and remain in the 9.4 delta.
- `test_ptx_tcgen05_ld_red_binds_redval_as_output`: `redval` is rendered
with an output constraint.
## Downstream
The FP4 FlashAttention-4 port in mlc-ai/tirx-kernels uses
`tcgen05.ld.red...max.f32` and depends on the `redval` binding fix.
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: [email protected]
For queries about this service, please contact Infrastructure at:
[email protected]
---------------------------------------------------------------------
To unsubscribe, e-mail: [email protected]
For additional commands, e-mail: [email protected]
