spectrometerHBH opened a new pull request, #20261:
URL: https://github.com/apache/tvm/pull/20261
## Summary
Align the TIRx CUDA backend with CUDA 13.4 / PTX ISA 9.4 and add support for
SM103a and SM107a (Rubin).
This squashes the following changes:
- Add PTX 9.3 SM103a support and architecture-aware TIRx tests
- Add PTX 9.4 SM107a support
- Support Rubin FP8 BMM requirements
- Route kernel tests by exact runtime architecture
- Align the PTX dialect with CUDA 13.4 / PTX ISA 9.4 only
## `T.ptx` table (`python/tvm/backend/cuda/ptx`)
- PTX ISA 9.3 additions certified at sm_103a: `fabric.*`, `clmad`,
`multimem.st.async` / `multimem.red.async` / `multimem.cp[.reduce].async.bulk`,
strong (`.relaxed.scope.b128`) bulk copies, mbarrier `.layout` / `.phase_type`
/ report forms / `check_layout`, `ld.mmio.acquire` / `st.mmio.release`,
`tcgen05.ld.red`, `tensormap.replace` swizzle_mode 4, `fence.proxy.*::fabric`.
- PTX ISA 9.4 / SM107 delta (`_PTX_94_ENTRIES`, sm_107f / sm_107a): mixed
precision packed add/sub/mul/fma, packed set, `ld.proxy::readonly`,
`prefetch.L1::32B.valid_addr`, cvt `.pzo` / `.rz` narrow / `.scaled::n1::ue8m0`
/ `.ue5m3x2`, `cp.async.bulk[.tensor]` multicast `::16b` / `::32b`, report
mechanisms, tensor base-address / attribute overrides, `im2col_no_offs::w`,
`applypriority.async.bulk[.tensor]`, eviction-priority bulk prefetch,
`spcompress` / `spdecompress`, `tcgen05.ld[.red].spcompress`, `tcgen05.alloc` /
`dealloc .exclusive`, `tcgen05.commit ::16b` / `::32b` and
`.sync_restrict::shared::read::mma::a`, `tcgen05.mma .kind::ti16` /
`.decompress::lut::b` / `.collector::b::*`, `atom` / `red` /
`cp.reduce.async.bulk .add.noftz.f32`, `ldmatrix .m8n16 .s8.s4`.
- Every section and table number cites the CUDA 13.4 developer-preview PTX
ISA 9.4 manual; every MEASURED clause records CUDA 13.4 ptxas behaviour
(re-measured: gaps that closed are noted, none widened). `tcgen05.ld.red` uses
the documented `.redOp{.abs}{.NaN}.type` slot order.
## Runtime and codegen
- Instruction descriptor: dense K=96 for `.kind::mxf4` / `.kind::mxf4nvf4`
(cta_group::1 M=128, cta_group::2 M=256; sm_103a / sm_107a) via Table 53 bit 31.
- `cuda_module.cc`: CUDA 13.4 oversized shared-memory mode
(`CU_SHARED_MEMORY_MODE_ALLOW_OVERSIZED_SHARED_MEMORY`) for cluster kernels
above the portable opt-in limit.
## Tests
- Tests route by exact runtime architecture (`cuda_arch` marker,
`tvm.testing.env.cuda_arch`), so the sm_100a / sm_103a / sm_107a suites run
only on their own devices.
## Docs
- `.agents/skills/tirx-ptx-dialect/SKILL.md` describes the CUDA 13.4 / PTX
ISA 9.4 toolchain model and the ISA-migration procedure.
- `tirx.pyi` and `docs/tirx/api/ptx.rst` regenerated / updated.
## Verification
Verified on CUDA 13.4 (V13.4.59), 4x sm_107a:
- `PTX_CERT=1` full-table certification (33 shards, including the sm_103a
manifest)
- sampled tier and codegen unit suites
- full `tests/python/tirx` suite: 2791 passed, 251 skipped
--
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]
