The GitHub Actions job "Lint" on tvm.git/pr/ptx-isa-9.4-cuda-13.4-sm103a-sm107a 
has failed.
Run started by GitHub user spectrometerHBH (triggered by spectrometerHBH).

Head commit for run:
1b3a011575f81ea34bf8ad33667a0914b2d5425a / Bohan Hou <[email protected]>
[CUDA][TIRx] PTX ISA 9.4 / CUDA 13.4 support for SM103a and SM107a (Rubin)

Squash of the dev-branch commits above the mirror:
- Add PTX 9.3 SM103a support and architecture-aware TIRx tests
- [CUDA][TIRx] Add PTX 9.4 SM107a support
- [CUDA][TIRx] Support Rubin FP8 BMM requirements
- [CUDA][TIRx] Route kernel tests by exact runtime architecture
- [CUDA][TIRx] 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 route by exact runtime architecture (`cuda_arch` marker,
  `tvm.testing.env.cuda_arch`), so 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.

Verified on CUDA 13.4 (V13.4.59), 4x sm_107a: PTX_CERT=1 full-table
certification (33 shards, incl. the sm_103a manifest), sampled tier, codegen
unit suites, and the full tests/python/tirx suite (2791 passed, 251 skipped).

Report URL: https://github.com/apache/tvm/actions/runs/33597966860

With regards,
GitHub Actions via GitBox


---------------------------------------------------------------------
To unsubscribe, e-mail: [email protected]
For additional commands, e-mail: [email protected]

Reply via email to