The GitHub Actions job "Lint" on tvm.git/phase-out-flat-device-intrins has 
succeeded.
Run started by GitHub user tlopex (triggered by tlopex).

Head commit for run:
eabcbbbd530edd6427f4b58d267e408b5c09ad52 / tlopex <[email protected]>
[TIRx] Phase out flat device-intrinsic op aliases

PR #19677 registered every CUDA / Trainium device intrinsic under two Op
names: a flat `tirx.<ns>_<name>` alias plus the canonical
`tirx.<ns>.<name>`. The flat aliases were a migration shim; passes and
codegen that match an intrinsic had to check both spellings (the
dual-name IsOp pattern). The canonical namespaced form is the intended
end state.

Remove the flat device-intrinsic aliases, keeping only the canonical
namespaced ops:

- RegisterDeviceIntrinsic (backend/cuda) and RegisterNKIIntrinsic
  (backend/trn) register only the canonical name.
- Drop the flat-only macro registrations for device intrinsics; the
  canonical op with all attrs is registered from the alias table. The
  WMMA tvm_*_sync / mma_store / mma_fill builtins and the profiling
  timer_*_cuda builtins keep their flat names (no namespace / canonical
  form, category "builtin").
- Remove the redundant flat tirx.ptx_fetch_register registration.
- C++ consumers that resolved a flat op by name string now use the
  canonical name; the ptx_elect_sync / cuda_func_call dual-name matchers
  collapse to the canonical check.
- Python: every device-intrinsic wrapper call site (backend/cuda/op.py,
  backend/trn/op.py, backend/cuda/script.py, gemm_async/tcgen05.py) now
  passes the canonical op name to call_intrin; no Python source names a
  removed flat op. _canonical_device_intrin_name is retained only as
  (a) the codegen-registration bridge from the underscore basename
  (which must stay underscore because it generates the C helper name)
  to the canonical CODEGEN_REGISTRY key, and (b) a fallback for user
  code that still passes a flat name. The remaining underscore strings
  in operator/intrinsics (cp_async.py / wgmma.py CODEGEN_REGISTRY keys)
  are codegen variant names, not TIR ops.

Generated CUDA is byte-identical: helper names are literals and codegen
dispatches by op name, with the registry resolving the canonical name to
the same helper.

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

With regards,
GitHub Actions via GitBox


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

Reply via email to