Add standalone C++/HIP examples investigating hipMemAllocationTypeUncached for P2P atomics

[Copilot]

hipMemAllocationTypeUncached (0x40000000) is accepted by hipMemCreate but CLR still hardcodes the coarse-grained GPU pool, so P2P atomics crash identically to hipMemAllocationTypePinned. These examples make that reproducible and debuggable on any machine.

iris/hip.py

• Add hipMemAllocationTypeUncached = 0x40000000 constant
• Add alloc_type parameter to mem_create() (default: hipMemAllocationTypePinned, backward-compatible)
• Thread alloc_type through _HIPVMemP2P and _run_p2p_atomics_hip

tests/unittests/test_hsa_vmem_reproducible.py

• test_hip_vmem_uncached_alloc_type: probes hipMemCreate with uncached type; skips if driver rejects it, otherwise tests local allocation + single-rank atomic only. Cross-rank access not tested — empirically it still causes SIGSEGV (same as pinned), confirming CLR ignores the type.
• Replace magic tolerances 0.01 / 0.5 with named constants _ATOMIC_EXACT_TOL / _ATOMIC_COUNT_TOL

csrc/standalone/ (new)

Two minimal, self-contained C++ programs — no MPI, no torchrun. Each spawns two ranks via fork + execl("/proc/self/exe") (self-exec avoids HSA internal thread corruption from plain fork), coordinating over a socketpair with SCM_RIGHTS FD passing.

p2p_atomics_hsa.cpp — Path 3 (HSA fine-grained, correct)

// Allocate from the fine-grained pool directly — CLR never reaches this path
hsa_amd_vmem_handle_create(fine_grained_pool, size, MEMORY_TYPE_NONE, 0, &handle);
hsa_amd_vmem_export_shareable_handle(&dmabuf_fd, handle, 0);
// ... SCM_RIGHTS exchange ...
hsa_amd_vmem_import_shareable_handle(peer_fd, &peer_handle);
// P2P sys-scope atomics → 0/N failures

p2p_atomics_hip.cpp — Path 2 (HIP VMem, always coarse-grained)

./p2p_atomics_hip                # safe: P2P non-atomic read only
./p2p_atomics_hip --uncached     # same setup with hipMemAllocationTypeUncached
./p2p_atomics_hip --atomics      # WARNING: GPU page fault (SIGSEGV) — demonstrates the bug

CLR routes both --pinned and --uncached through the coarse-grained pool (CoarseGrain=1 in KFD). P2P atomics crash identically for both.

Implementation notes:

• hipMemset/hipMemcpy(D2H) silently fail on HSA VMem VAs (not registered in HIP's pointer tables); init and readback use k_zero / k_copy kernels with a hipMalloc bounce buffer
• k_copy calls __threadfence_system() before the read for cross-GPU write visibility
• Agent-scope atomic results are informational only; sys-scope drives PASS/FAIL

Original prompt

---

This section details on the original issue you should resolve

<issue_title>VMem allocator and atomics</issue_title>
<issue_description>From previous tests, only the runner machine had issues running some atomics (all reduce for example) when using the vmem allocator. Could be a bad initialization flag. I am not able to reproduce the bug on other GPUs. Please use the vmem allocator, run all unittests and figure out a fix. For reference, consider shallow cloning https://github.com/ROCm/rocm-systems to understand how the low level apis. Your outcomes should be 2 things: 1: minimal repro that doesn't include any iris code that show the bug, 2: ideally a bug fix patch (I suspect the code had some wrong flag when allocating memory or something similar)</issue_description>

Comments on the Issue (you are @copilot in this section)

• Fixes VMem allocator and atomics #418

---

💡 You can make Copilot smarter by setting up custom instructions, customizing its development environment and configuring Model Context Protocol (MCP) servers. Learn more Copilot coding agent tips in the docs.