[Windows/HIP] Fix APU large BAR detection to enable >64GB unified memory allocation

[moonshadow-25]

Description:

Problem

On Windows, HIP unconditionally disables large BAR (largeBar_ = false) for all APU devices, regardless of their actual memory architecture.
This causes unified memory architecture (UMA) APUs like AMD Strix Halo (gfx1151) to be artificially limited to ~64GB memory allocation, even
when the system has 128GB+ RAM.

Root cause: The premature redirect to RemoteUSWC heap in palmemory.cpp triggers the Windows GART allocator's 50%-of-RAM limit,
preventing applications from utilizing the full unified memory capacity.

Solution

1. paldevice.cpp — Proper large BAR detection

Remove the HIP-specific hardcoded largeBar_ = false. Instead, detect UMA APUs by checking if GpuHeapInvisible (discrete VRAM) is absent.
On UMA systems, the entire system memory is GPU-local, so resizable BAR semantics naturally apply.

Before:

if (amd::IS_HIP) {
  info_.largeBar_ = false;  // hardcoded for all APUs
} else if (heaps[Pal::GpuHeapInvisible].logicalSize == 0) {
  info_.largeBar_ = true;
}

After:
// For APU/integrated devices (no invisible/private VRAM heap), treat as large bar.
if (heaps[Pal::GpuHeapInvisible].logicalSize == 0) {
  info_.largeBar_ = true;
}

2. palmemory.cpp — Avoid premature RemoteUSWC redirect

For large-bar APUs, skip the early redirect to RemoteUSWC. Let the natural retry loop handle heap selection (Local → Persistent →
RemoteUSWC), avoiding the Windows GART 50% cap.

Before:
if (amd::IS_HIP && dev().settings().apuSystem_) {
    // always redirects, hits GART limit on large-bar APUs
    memType = RemoteUSWC;
}

After:
if (amd::IS_HIP && dev().settings().apuSystem_ && !dev().info().largeBar_) {
    // only redirect for non-large-bar APUs
    memType = RemoteUSWC;
}

Testing

Tested on AMD Strix Halo (gfx1151), 128GB unified memory, Windows 11:

┌─────────────────────────┬───────────────────┬──────────────────┐
│       Allocation        │ Before (upstream) │ After (this fix) │
├─────────────────────────┼───────────────────┼──────────────────┤
│ hipMallocManaged(64 GB) │ ❌ Out of memory  │ ✅ Success       │
├─────────────────────────┼───────────────────┼──────────────────┤
│ hipMallocManaged(70 GB) │ ❌ Out of memory  │ ✅ Success       │
├─────────────────────────┼───────────────────┼──────────────────┤
│ hipMallocManaged(80 GB) │ ❌ Out of memory  │ ✅ Success       │
├─────────────────────────┼───────────────────┼──────────────────┤
│ Cumulative 2×40 GB      │ ❌ Out of memory  │ ✅ Success       │
└─────────────────────────┴───────────────────┴──────────────────┘

Compatibility

- ✅ Discrete GPUs: No change — invisible heap exists, largeBar_ determined by hardware as before
- ✅ Non-large-bar APUs: Behavior unchanged — early RemoteUSWC redirect still applies
- ✅ Large-bar APUs: Now correctly utilize full unified memory
- ✅ OpenCL: Benefits from the same fix — UMA detection is API-agnostic

Impact

This fix is critical for enabling large model inference (LLMs, AI workloads) on APU-based systems. Without it, users with 128GB unified
memory systems cannot allocate more than ~64GB via HIP, making modern APUs unsuitable for AI workloads despite having sufficient RAM.

[moonshadow-25]

Hi ROCm team,

I've updated the branch to sync with the latest develop.

Just wanted to mention that the companion fix in llama.cpp (ggml-org/llama.cpp#20536) has been
merged and is now in production use. This PR addresses the underlying ROCm issue that was causing APU compatibility
problems.

Could someone from the team approve the CI workflows when you get a chance? Happy to address any feedback once the
checks run.

Thanks!

[gandryey]

There was a reason why large bar was disabled at least on dgpu systems, because HIP expects CPU access to the device memory, but in Windows it's not possible. We may have regressions after this change or extra fixes in the tests.

[github-actions]

This pull request has been inactive for 25 days and will be marked as stale.

If you would like to keep this PR open, please:

• Add new commits
• Add a comment explaining why it should remain open

This PR will be automatically closed in 5 days if no further activity occurs.

[github-actions]

This pull request has been automatically closed due to inactivity (30 days with no updates).

If you'd like to continue working on this, feel free to reopen the PR or create a new one.