feat: isolate cutlass._mlir imports behind compat gateway (#118)

[bikrammajhi]

Summary

Closes #118 — isolates all direct cutlass._mlir usage behind a single compatibility gateway.

CuTeDSL's generated MLIR/NVVM bindings (cutlass._mlir.{ir,arith,llvm,vector,nvvm,cute}) are private implementation detail with no cross-version stability contract — the tcgen05_ld/st breakage between CutDSL 4.5.2 and 4.5.3 was the first confirmed incident. cuLA imported these directly from 10 kernel modules, so a patch release could break kernel emission silently, delayed until import/JIT time.

Changes

1. cula/ops/_mlir_compat.py — the only module that may touch cutlass._mlir:
   • Lazy dialect loading (plain import cula never touches private bindings)
   • Version contract mirroring pyproject.toml (>=4.4.2,<4.7,!=4.5.0), enforced with a fail-fast RuntimeError naming the offending version
   • Canary probes on every dialect's entry points, so a missing/renamed binding fails with an actionable message instead of surfacing mid-JIT
   • vector_extract_element() helper with cross-version dispatch (see below)
2. Migration — all 10 consumer files bind their dialect aliases from the gateway (one import-line swap per file, zero behavioral changes).
3. Bug fix discovered during prep — CutDSL renamed vector.extractelement to vector.extract in the 4.6 line; against nvidia-cutlass-dsl==4.6.2 (resolved by our own pyproject.toml range) store_256b raises AttributeError at JIT time. store_256b is used by the KDA SM100 backward path. The gateway's vector_extract_element dispatches to whichever binding exists, restoring 4.6 compatibility.
4. Tests — 15 headless tests for the gateway: version parsing, contract enforcement, missing dialect/canary fault injection, vector dispatch on both 4.5/4.6 shapes, and a real-wheel smoke test. No GPU required.
5. scripts/modal_validate.py — H100/CUDA 12.9 validation harness (runs the headless suite + SM90 prefill/decode tests in one modal run).

Validation status

• Headless: 15/15 tests/test_cutedsl_compat.py green against nvidia-cutlass-dsl==4.6.2
• All 10 migrated consumer modules import cleanly against the installed wheel
• SM90 suite on Hopper (H100) — running via modal run scripts/modal_validate.py; will attach results

Notes for reviewers

• The old extractelement call in store_256b passed the position as a keyword (position=...); the new vector.extract binding takes it positionally ((source, dynamic_position, static_position)). The dispatch helper deliberately keeps the exact call shapes per branch.
• Canary choices reflect what the migrated consumers actually call; extend _CANARIES/_ANY_OF_CANARIES when new private-API usage lands.

[bikrammajhi]

Aligned with #119 (merged) — this PR builds on top of the tcgen05 signature-detection work; no overlap: #119 owns _cutedsl_compat.py capability probing, this PR owns the remaining cutlass.isolated-mode isolation.

Scope recap

• cula/ops/_mlir_compat.py: new gateway — dialect-presence canaries (arith/cute/ir/llvm/nvvm/vector), lazy cutlass._mlir.\*aliases, version-parse helper, module-level __getattr__ shim, mlir_context_dir().
• 10 op files migrated off direct cutlass._mlir imports.
• cula/ops/ptx.py + cula/ops/sm100/ptx.py: store_256b is now version-agnostic — probes the real API (fixes the CuTeDSL 4.6.2 extractelement→extract break).
• Tests: tests/test_cutedsl_compat.py extends the Detect tcgen05 load/store API by signature #119 file with gateway + JIT-probe coverage; scripts/modal_validate.py runs the compat + SM90 prefill/decode suites headless on an H100 (image pull/build + pytest; JSON summary returned).

Healthy on the minimal local smoke suite; the automated H100 run is in flight — I'll post the summary when it's back.