spectrometerHBH opened a new pull request, #20224:
URL: https://github.com/apache/tvm/pull/20224
## Summary
This PR audits the table-driven `T.ptx` dialect against PTX ISA 9.2, the ISA
version shipped by CUDA 13.2, and brings the registered surface, operand
contracts, generated stubs, and certification tests into alignment.
The main changes are:
- add missing PTX 9.2 forms for vector `red`, `cp.reduce.async.bulk`, atomic
bit-bucket destinations, mbarrier bit buckets, `mapa` carrier shapes, and the
`lop3` predicate-destination sibling;
- model relaxed operand carriers for `ld`, `st`, `ldu`, scalar `cvt`,
`slct`, and `st.bulk` while preserving existing canonical helper names and
signatures;
- correct `wgmma.mma_async` `scale-d` from a text immediate to a runtime
predicate operand;
- distinguish generic and explicit shared address carriers for bare versus
`.shared::*` `tcgen05.alloc` and `tcgen05.commit` forms;
- add an instruction-local `lop3 immLut` check for the documented 0..255
domain;
- add a private byte-register bridge for the PTX 9.2
`st.async.release.{b8,u8,s8}` forms, without changing the existing 16/32/64-bit
helpers;
- extend generated-stub, invariant, round-trip, exact-rendering, and
production-shaped nvcc/ptxas certification coverage.
## CUDA 13.2 boundaries
The table exposes callable inline-assembly helpers, so documented PTX forms
that CUDA 13.2 ptxas does not compile are reported explicitly instead of being
advertised as usable helpers. This applies narrowly to bf16 atomic bit-bucket
forms, widened floating source registers for scalar `cvt`, and widened carriers
on 256-bit `ld`/`st`. Exact-width forms remain registered and certified.
## Testing
- `pre-commit run --files` on all changed files
- `python -m pytest -q tests/python/tirx/codegen/test_ptx_dialect.py
tests/python/tirx/codegen/test_ptx_cvt.py
tests/python/tirx/codegen/test_ptx_addr.py
tests/python/tirx/codegen/test_codegen_hopper.py` (`212 passed, 34 skipped`)
- full TIRX suite: `2877 passed, 91 skipped, 3 xpassed`; seven GPU OOM cases
caused by unrelated concurrent GPU processes passed when rerun serially on a
clean GPU
--
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]
