[nihalpasham]

What This PR Does

This PR aligns our native opt + llc PTX backend with nvcc and libNVVM so shared-memory pointers (ptr addrspace(3)) use 32-bit (%r) registers instead of 64-bit (%rd) registers:

1. Set NVPTX_DATALAYOUT_PTX = NVPTX_DATALAYOUT_FULL (llvm-export) so opt -O2 sees p3:32:32:32.
2. Pass --nvptx-short-ptr to llc (cuda-oxide-codegen) so llc emits 32-bit %r shared-memory addresses.
3. Return ptr addrspace(3) as a 64-bit generic ptr across .func return slots (mir-lower) so non-inlined device functions returning shared pointers work cleanly under --nvptx-short-ptr.

---

1. Why PtxExportConfig Was Emitting 64-Bit Shared Pointers

Previously, NvvmExportConfig (libNVVM) and PtxExportConfig (opt + llc) used two different data layouts:

1. NvvmExportConfig (`NVPTX_DATALAYOUT_FULL`, matching `nvcc`):
   "e-p:64:64:64-p3:32:32:32-..."
                 ^^^^^^^^^^^
                 Tells LLVM that shared memory (`addrspace(3)`) is 32-bit!

2. PtxExportConfig (`NVPTX_DATALAYOUT_PTX`, before this PR):
   "e-i64:64-i128:128-v16:16-v32:32-n16:32:64"
                 (Omitted `p3`, so `addrspace(3)` defaulted to 64-bit!)

With p3:32:32:32 in NVPTX_DATALAYOUT_PTX and --nvptx-short-ptr in base_llc_args:

• opt -O2 narrows addrspace(3) GEP indices from i64 to i32.
• llc emits 32-bit %r registers for all shared-memory operations:

BEFORE (64-bit %rd shared-memory pointers in PTX):
  mbarrier.arrive.shared.b64                 %rd147, [%rd42];
  ld.shared.v2.b64                           {%rd127, %rd128}, [%rd126+50688];
  ldmatrix.sync.aligned.m8n8.x4.shared.b16   {...}, [%rd129];
  stmatrix.sync.aligned.m8n8.x2.shared.b16   [%rd87], {...};

AFTER (32-bit %r shared-memory pointers in PTX, matching `nvcc`):
  mbarrier.arrive.shared.b64                 %rd64, [%r48];
  ld.shared.v2.b64                           {%rd56, %rd57}, [%r281+50688];
  ldmatrix.sync.aligned.m8n8.x4.shared.b16   {...}, [%r290];
  stmatrix.sync.aligned.m8n8.x2.shared.b16   [%r13], {...};

---

2. Why mir-lower Coerces ptr addrspace(3) Returns Across .func Boundaries

In PTX, a function returns its value through a named .param box (func_retval0 in the definition, retval0 at the call site), and both boxes must have the same bit width:

1. FUNCTION DEFINITION                         2. CALLER SITE
──────────────────────                         ──────────────
.func (.param .b64 func_retval0) my_fn(...)    {
{                                                .param .b64 retval0;
    ...                                          call.uni (retval0), my_fn, (...);
    st.param.b64 [func_retval0], %rd2; ◄───────► ld.param.b64 %rd_out, [retval0];
    ret;                                       }
}

When --nvptx-short-ptr is enabled, LLVM 23's NVPTXAsmPrinter has a bug when a non-inlined .func returns a 32-bit ptr addrspace(3) directly (for example, #[inline(never)] fn bits_to_pointer(bits: usize) -> *mut SharedArray<u32, 32> in shared_ptr_transmute):

BEFORE (returning `ptr addrspace(3)` directly under `--nvptx-short-ptr`):
  1. Definition box (`func_retval0`):   .param .b64   (8 bytes - LLVM NVPTXAsmPrinter bug!)
  2. Caller box     (`retval0`):        .param .b32   (4 bytes)
                                        --> ptxas error: Type of argument does not match
                                            formal parameter 'func_retval0'!

AFTER (`shared_return_abi_info` in `mir-lower`, matching `nvcc`):
  1. Inside `my_fn`:                    Casts `ptr addrspace(3)` -> 8-byte generic `ptr` at `ret`
  2. Definition box (`func_retval0`):   .param .b64   (8 bytes)
  3. Caller box     (`retval0`):        .param .b64   (8 bytes)
  4. At caller after `call`:            Casts 8-byte generic `ptr` -> `ptr addrspace(3)`
                                        (and when inlined, `opt -O2` folds both casts away!)

---

Verification

• scripts/smoketest.sh on RTX 5090 (sm_120a): 238 / 238 examples PASS (including shared_ptr_transmute, shared_memory_helper, and tma_copy).
• ptxas -v -arch=sm_120a on blockscale_gemm_ctm:
  • Register usage drops from 162 to 160 registers (0 bytes stack frame, 0 spill bytes).
• Unit & Integration Tests:
  • cargo test --all-targets -p mir-lower (331 unit + 150 integration tests): PASS
  • cargo test --all-targets -p llvm-export (168 tests): PASS
  • cargo test --all-targets -p cuda-oxide-codegen (1,294 tests): PASS
  • just fmt-check, just clippy, just check-guards: all PASS