[nihalpasham]

What This PR Does

This PR makes three changes to shared-memory addressing and CTA-local TMA loads (following up on #1336, #1341, #1238, #1337, and #1249):

1. Keep .add() and .cast() on shared pointers in addrspace(3) (mir-importer)
2. Use 32-bit i32 shared-memory offsets in cuda-oxide-mlir-export (+ #[grid_constant] in blockscale_gemm_cute)
3. Remove SharedPtr<T> (-3,141 lines) and take dst: *mut u8 directly in cp_async_bulk_tensor_{1..5}d_g2s_cta (+ route ctm-rs to .shared::cta)

---

1. Keep .add() and .cast() on Shared Pointers in addrspace(3) (mir-importer)

Previously, SlotAddrSpaceMap only preserved addrspace(3) across plain _y = _x copies. Calling .add(offset) or .cast::<T>() on a shared pointer demoted it to a 64-bit generic pointer (addrspace(0)), forcing an addrspace(3) -> addrspace(0) -> addrspace(3) round-trip:

Rust kernel:
  let base = DynamicSharedArray::<u8, 1024>::get_raw();
  let tile = base.add(slot * TILE_BYTES).cast::<u32>();

BEFORE:
  base (ptr addrspace(3))
    ──► addrspacecast to generic (ptr)       <-- demoted on .add() / .cast()!
    ──► 64-bit generic GEP
    ──► addrspacecast back to shared (ptr addrspace(3))

AFTER:
  base (ptr addrspace(3))
    ──► GEP in ptr addrspace(3)              <-- stays in shared memory (0 casts)!

---

2. Use 32-bit i32 Shared-Memory Offsets in cuda-oxide-mlir-export

Shared memory (!llvm.ptr<3>) uses 32-bit addresses, but MirPtrOffsetOp and SM120 mma / epilogue lowering in cuda-oxide-mlir-export were extending 32-bit lane/warp indices to 64-bit i64 before indexing !llvm.ptr<3>.

BEFORE (64-bit i64 index math on 32-bit !llvm.ptr<3>):
  lane (i32) ──► zext to i64 ──► 64-bit swizzle math ──► GEP(!llvm.ptr<3>, i64)
  * Blinded LLVM's LoadStoreVectorizer (emitted 4 scalar 32-bit LDS instructions)
  * Left 20 64-bit SHF.R.U64 / IADD.64 instructions in the K-loop

AFTER (32-bit i32 index math on 32-bit !llvm.ptr<3>):
  lane (i32) ──────────────────► 32-bit swizzle math ──► GEP(!llvm.ptr<3>, i32)
  * LLVM's LoadStoreVectorizer automatically merges the 4 scalar loads into LDS.128!
  * Eliminates the 64-bit shift/add instructions in SASS

---

3. Remove SharedPtr<T> and Take dst: *mut u8 in .shared::cta TMA (#1336, #1341, #1238, #1337, #1249)

Why #1238 Happened on RTX 5090 (sm_120a)

Consumer Blackwell (sm_120a) has no hardware support for cross-block TMA loads. When a kernel emits a Cluster TMA load (.shared::cluster), ptxas generates an if / else with a software emulation CALL in the else branch, which allocates an 8-byte stack frame and causes ptxas to ignore setmaxnreg (warning C7506):

Cluster TMA (`.shared::cluster`) on sm_120a:
  if (ptr.block_id == my_block_id) {
      UTMALDG.2D ;                                                 // fast hardware TMA
  } else {
      CALL __cuda_syscall_cp_async_bulk_tensor_2d_tile_unicast ;   // <-- kills setmaxnreg!
  }

Local Block TMA (`.shared::cta`) on sm_120a:
  UTMALDG.2D ;   // no branch, no CALL, 0 stack frame, keeps setmaxnreg!

Why We Don't Need SharedPtr<T> (Revising #1336 / #1341)

PR #1336 (and #1341) introduced SharedPtr<T> and lowered it to ptr addrspace(3) across function signatures, return values, structs, and arrays (+3,141 lines of plumbing). That created a fundamental layout mismatch:

• rustc always sizes SharedPtr<T> at 8 bytes (*mut T).
• Under p3:32:32:32 (libNVVM / --nvptx-short-ptr), LLVM sizes ptr addrspace(3) at 4 bytes, causing return-type mismatches in llc and misaligned address crashes in arrays/structs.

With Change 1 (SlotAddrSpaceMap) + convert_g2s_cta (which already accepts both addrspace(3) and addrspace(0) pointers, just like barrier: *mut Barrier), cp_async_bulk_tensor_{1..5}d_g2s_cta works directly with dst: *mut u8:

// Direct CTA-local TMA load with plain *mut u8 (0 cvta.to.shared, 0 stack frame):
cp_async_bulk_tensor_2d_g2s_cta(
    &raw mut TILE as *mut u8,
    tensor_map,
    tile_x,
    tile_y,
    &raw mut BAR,
);

We also updated ctm_rs::tma::cp_async_bulk_tensor_shared_cta_global to call cp_async_bulk_tensor_2d_g2s_cta (.shared::cta) instead of cp_async_bulk_tensor_2d_g2s (.shared::cluster), eliminating the 8-byte stack frame and runtime CTA-id check in blockscale_gemm_ctm.

---

Verification

• All 238 examples pass scripts/smoketest.sh on RTX 5090 (sm_120a), including tma_copy/verify-code-shape.sh (.shared::cta emitted, 0 cvta.to.shared, 0 shared::cluster).
• just check (fmt-check, clippy, test, test-cuda, check-guards, check-intrinsics main, doc-check): all PASS.
• scripts/check-primitives-api-compile.py (api, fp16, blockscale): all PASS.
• setmaxnreg + .shared::cta on sm_120a (ptxas -v + nvdisasm): 168 registers, 0 stack frame, USETMAXREG.DEALLOC.CTAPOOL 0x28, USETMAXREG.TRY_ALLOC.CTAPOOL UP0, 0xe8, UTMALDG.2D, no CALL.
• 2048x2048x1024 MXFP4 Block-Scaled GEMM on RTX 5090 (Nsight nsys profile kernel-active median of 31 launches):

Kernel	Backend	Before This PR	After This PR	vs. Python CUTLASS CTM (10.880 us)
Python CUTLASS CTM (reference/benchmark.py)	cutlass-mlir	10.880 us	10.880 us	baseline
Rust blockscale_gemm_cute	cutlass-mlir	11.360 us	10.912 us	-0.032 us (-0.3%)
Rust blockscale_gemm_ctm	Native LLVM/PTX	11.616 us	10.976 us	-0.096 us (-0.9%)