feat(rocm): unified memory true by construction on integrated APUs (issue #41 F6, approach (b))

[localai-bot]

What this is

The ratified approach (b) from issue #41's F6: on a ROCm device probing
hipDeviceAttributeIntegrated=1 + ManagedMemory=1 + ConcurrentManagedAccess=1,
Backend::Alloc uses hipMallocManaged(hipMemAttachGlobal) and UnifiedMemory()
returns true exactly then — host access becomes API-guaranteed rather than
architecturally incidental, unblocking the zero-kernel reference-tier M2 path on
Strix Halo (gfx1151) and the 780M (gfx1103), where XNACK-less RDNA3 probes
PageableMemoryAccess=0 and the W0 CUDA-shaped conjunction read false.

Draft opened at start per the helper protocol; full work lands on this branch.

Honesty statement — read before reviewing

Written blind. No AMD GPU and no hipcc exist on the authoring machine; the
.hip delta in this PR has never been compiled. That is this lane's standing
policy (the W0 skeleton landed the same way and was then verified by four
community boards on #41). CPU-side gates are green here; the HIP-side evidence
is owed by board owners, itemized below. A compile error in this PR is useful
data, not a review failure.

The change

• src/vt/rocm/rocm_backend.hip — probe the two managed attributes (failure
  defaults to 0, never unregisters the device); UseManagedAlloc() predicate;
  managed branch in Alloc (bytes==0 stays on hipMalloc); UnifiedMemory() = managed_branch || (integrated && pageable) (W0 ground kept intact); Free
  stays hipFree — the HIP API documents it as the release call for both
  hipMalloc and hipMallocManaged (mirrors cudaFree; no hipFreeManaged
  exists), certainty stated in the code.
• Allocation-site audit (all of them): Alloc branches; Free single-path;
  inherited AllocPinned/FreePinned delegate to Alloc/Free
  (src/vt/backend.cpp:19-20) and so ride the same branch, coherent with the
  pinned contract (include/vt/backend.h:76-78); rocm_rmsnorm.hip has zero
  allocation sites; no pool paths exist in the skeleton. Discrete devices are
  byte-identical to W0: the branch is provably dead at Integrated=0, and a
  runtime test asserts it.
• include/vt/rocm/rocm_runtime.h — HIP-free introspection probes
  ManagedAllocActive(index) / IntegratedDevice(index) (the testable seam,
  also what board owners report).
• tests/vt/test_rocm_backend.cpp — two runtime-gated cases: (1) alloc path
  and UnifiedMemory() move together (discrete: both dead/false; integrated:
  managed active and unified true, loud failure with the probe triple
  otherwise); (2) issue ROCm (AMD GPU) backend #41 F6's decisive experiment as a standing gate —
  host-writes inputs with no Copy, the native RmsNorm kernel reads them,
  host-reads the kernel-written output with no Copy. Both compile on every
  build via the -Werror syntax-check object; both no-op without a device.
• CMakeLists.txt — F1/F3 absorption: when ROCM_PATH exists, derive
  CMAKE_HIP_COMPILER_ROCM_ROOT, seed --rocm-path into CMAKE_HIP_FLAGS,
  export ROCM_PATH into the environment — each ONLY when unset, before
  check_language(HIP). Arch/TheRock should now configure with no manual
  flags. F2 was downstream of the unidentified compiler and disappears with
  F1/F3 (documented, no separate change).
• Docs + records: docs/ROCM.md §3.1 (decision + behavior table) and §5.2
  (board-owner sequence), docs/BUILD.md, spec
  .agents/specs/rocm-unified-memory-b.md (decision quoted verbatim +
  community verification table), backend-matrix row, porting-inventory §9.14
  (ADDITIVE deviation: upstream vLLM allocates via torch and has no analog —
  rocm.py:75-77, 909-910 only name the APUs), STATUS/BENCHMARKS/FEATURES
  rows, state + NOW.

Gates run here (CPU only)

• scripts/agent-preflight.sh EXIT=0 (all record gates + mutation suites).
• CPU configure + -Werror object-compile of the platform and test TUs
  (vllm_rocm_platform_syntax_check): clean; compile-mutation of the new
  probe API went red (2 errors) and green on restore.
• The .hip files remain excluded from non-HIP builds — byte-identical CPU
  builds.

Evidence the community owes (PENDING — the M0/M1 table shape from #41)

Gate	Board	Report
Configure with NO manual flags	gfx1151, gfx1103 (Arch/TheRock)	configure log
.hip compile of this delta	any	first error text, or clean
ctest -R 'rocm|cross_device' incl. the 2 new cases	all four boards	pass/fail + the printed integrated/managed-alloc/UnifiedMemory triple
Discrete branch provably dead	gfx1100, gfx1201	triple = 0/0/false
M2: small dense model, greedy parity vs --device cpu	gfx1151, gfx1103	tokens + VT_OP_PROVIDER_STATS=1 fallback list

Sequencing note: composes with #140 (discrete gfx1201 lane) — this PR touches
the integrated-path branch only; will rebase onto main if #140 lands first.

Closes the F6 blocker on #41 (task #286).

[localai-bot]

Landed on main as e9f3561 (mudler-authored squash, lease-guarded push). Basis: PR CI 12/12 green on the head (full CPU build + both sanitizers), local -Werror syntax-check target green, nine record gates green; the mutation-review round was deliberately skipped as disproportionate for a blind-HIP lane whose real gate is community hardware (disclosed in the commit). @jimmykarily @arch-btw @tbrasser — approach (b) is now on main: pull main (or keep testing the branch, identical content), run the flag-free configure + ctest -R 'rocm|cross_device' + the M2 attempt per the issue-41 handoff comment, and post the table. A compile error remains the most useful thing you can post.