spectrometerHBH opened a new pull request, #20271: URL: https://github.com/apache/tvm/pull/20271
## Summary Two `T.ptx` table follow-ups to #20266 (`python/tvm/backend/cuda/ptx/table.py`). - **Collector-A-only SM107 block-scale MMA.** On the SM107 `tcgen05.mma` block-scale entries `collector_b` becomes optional, so the activation-stationary FP8 form with collector A alone is accepted alongside the A+B form. - **Open `cp.async` wait-group immediates.** `cp.async.wait_group N` and `cp.async.bulk.wait_group[.read] N` no longer restrict `N` to `0..7`. PTX ISA 9.7.10.28.3.3 / 9.7.10.28.6.2 define `N` only as an integer constant with no value domain, and CUDA 13.4 ptxas at sm_107a accepts values beyond 7 (measured: 8, 9, 16, 255; the bulk `.read` form also 2147483647 and -1). Each call-site constant still becomes its own helper; enumeration and full-table certification sample the open operand at 0 and certify the instruction shape. ## Tests (`tests/python/tirx/codegen/test_ptx_dialect.py`) - New `test_ptx_tcgen05_mma_block_scale_collector_a_without_block_size`: SM107 activation-stationary FP8 accepts collector A without `.block*`. - Wait-group tests updated for the open immediate. ## Downstream Needed by the Rubin masked grouped GEMM and block-scaled fused-MoE SwiGLU ports in mlc-ai/tirx-kernels. -- 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]
