`hipMallocFromPoolAsync` broken on ROCm 7.2

[luraess]

It seems AMDGPU.jl fails to allocate any device memory, as reported in #869 and lately while having a CI runner using ROCm 7.2 container.

Description

On ROCm 7.2 (AMD Radeon RX 7800 XT / gfx1101), any attempt to allocate device memory via
AMDGPU.jl fails immediately. Even the simplest operation crashes:

julia> using AMDGPU
julia> AMDGPU.ROCArray([1, 2, 3])
# → error / silent NULL pointer

Possible cause

AMDGPU.jl currently uses hipMallocFromPoolAsync exclusively for all HIPBuffer allocations.
On ROCm 7.2 this function appears to be broken (returns a NULL pointer or errors) for at least
some GPU/driver combinations, with no fallback path.

This should be confirmed by:

using AMDGPU, AMDGPU.HIP

dev          = AMDGPU.device()
stream       = AMDGPU.HIPStream(); # pretty printing fix in #900 
default_pool = HIP.default_memory_pool(dev)

ptr_ref = Ref{Ptr{Cvoid}}()
HIP.hipMallocFromPoolAsync(ptr_ref, 24, default_pool, stream)
HIP.wait(stream)
println("ptr: $(ptr_ref[])")   # returns NULL on affected systems

While the synchronous fallback works correctly:

ptr_ref = Ref{Ptr{Cvoid}}()
HIP.hipMalloc(ptr_ref, 24)
println("ptr: $(ptr_ref[])")   # non-NULL, works fine

Environment

ROCm version	7.2 (rocm/dev-ubuntu-24.04:7.2-complete)
GPU	AMD Radeon RX 7800 XT (DID 0x747e, gfx1101)
Julia	1.12
AMDGPU.jl	master

Possible fix

Branch lr/rocm7 adds a per-device broken flag with an automatic explicit fallback to hipMalloc:

• A global POOL_ALLOC_BROKEN::LockedObject(Set{Int64}) tracks device IDs where pool
  allocation has failed.
• On first allocation, hipMallocFromPoolAsync is attempted. If it returns NULL (or exhausts
  all retry phases), the device is added to POOL_ALLOC_BROKEN and all subsequent allocations
  on that device use synchronous hipMalloc.
• HIPBuffer gains a pool_alloc::Bool field so the correct free function is used
  (hipFreeAsync for pool buffers, hipFree otherwise) — mixing them causes memory corruption.

The change is in
src/runtime/memory/hip.jl.

I would be happy to get any insights and ideally others with access to AMD GPU to test. I will give it a try asap on MI300a, once I finalise the setup there.

Disclaimer: I iterated on the issue with Claude code. The main purpose is debugging and ideally we can work towards a proper fix after input from others.

[luraess]

EDIT: rocsparse test all fail on ROCm 7.2.2 and gfx1101.

[luraess]

Actually, looks like hipMallocAsync and hipMallocFromPoolAsync are broken in ROCm 7.2.2 (at least for gfx1101) even in plain C++. hipMalloc works fine. A ROCm 7.2 bug?

HIP C++ reproducer: debug_pool.cpp

[luraess]

Weirdly, all test pass on ROCm 7.2 when running on MI300a. This may suggest hipMallocAsync and hipMallocFromPoolAsync may be broken on RDNA3 GPUs with ROCm 7+...

[vchuravy]

Do you have the C++ reproducer?

[luraess]

There is the debug_pool.cpp code which shows async malloc failure @vchuravy (2 comments above)

[luraess]

BTW, here is a simpler MWE that shows the issue in HIP C++:

// Minimal reproducer for hipMallocAsync segfault on gfx1100/gfx1101 + ROCm 7.x
// Known broken: hipMallocAsync and hipMallocFromPoolAsync segfault or return null.
// hipMalloc (synchronous) works fine as baseline.
//
// Compile: hipcc -o debug_pool debug_pool.cpp
// Run:     ./debug_pool

#include <hip/hip_runtime.h>
#include <cstdio>
#include <cstdint>

#define CHECK(cmd) do {                                          \
    hipError_t e = (cmd);                                       \
    if (e != hipSuccess) {                                       \
        printf("FAIL [%s]: %s\n", #cmd, hipGetErrorString(e));  \
        return 1;                                                \
    }                                                            \
} while(0)

int main() {
    // Print device info
    hipDeviceProp_t prop;
    CHECK(hipGetDeviceProperties(&prop, 0));
    printf("Device: %s  arch: %s\n\n", prop.name, prop.gcnArchName);

    hipStream_t stream;
    CHECK(hipStreamCreate(&stream));

    // TEST 1: hipMalloc (synchronous) — expected to work
    {
        void* ptr = nullptr;
        hipError_t e = hipMalloc(&ptr, 64);
        printf("[hipMalloc]             %s\n", (e == hipSuccess && ptr) ? "OK" : hipGetErrorString(e));
        if (ptr) (void)hipFree(ptr);
    }

    // TEST 2: hipMallocAsync — OOM or fails on gfx1100/gfx1101 + ROCm 7.x
    {
        void* ptr = nullptr;
        hipError_t e = hipMallocAsync(&ptr, 64, stream);
        (void)hipStreamSynchronize(stream);
        printf("[hipMallocAsync]        %s\n", (e == hipSuccess && ptr) ? "OK" : hipGetErrorString(e));
        if (ptr) { (void)hipFreeAsync(ptr, stream); (void)hipStreamSynchronize(stream); }
    }

    // TEST 3: hipMallocFromPoolAsync — OOM or fails on gfx1100/gfx1101 + ROCm 7.x
    {
        hipMemPoolProps pool_props{};
        pool_props.allocType       = hipMemAllocationTypePinned;
        pool_props.location.type   = hipMemLocationTypeDevice;
        pool_props.location.id     = 0;
        hipMemPool_t pool;
        CHECK(hipMemPoolCreate(&pool, &pool_props));

        void* ptr = nullptr;
        hipError_t e = hipMallocFromPoolAsync(&ptr, 64, pool, stream);
        (void)hipStreamSynchronize(stream);
        printf("[hipMallocFromPoolAsync] %s\n", (e == hipSuccess && ptr) ? "OK" : hipGetErrorString(e));
        if (ptr) { (void)hipFreeAsync(ptr, stream); (void)hipStreamSynchronize(stream); }
        (void)hipMemPoolDestroy(pool);
    }

    (void)hipStreamDestroy(stream);
    return 0;
}

Running on ROCm 7.2 on gfx1101:

lraess@gpuci:~/data/AMDGPU.jl$ hipcc -o debug_rocm debug_rocm.cpp
lraess@gpuci:~/data/AMDGPU.jl$ ./debug_rocm
Device: AMD Radeon RX 7800 XT  arch: gfx1101

[hipMalloc]             OK
[hipMallocAsync]        out of memory
[hipMallocFromPoolAsync] out of memory

According to https://rocm.docs.amd.com/en/latest/compatibility/compatibility-matrix.html, gfx1101 should be supported by ROCm 7.2.1-2.

[luraess]

Interestingly, all works fine on MI300a + ROCm 7.2 (and all tests pass)

[vchuravy]

Seems to work okay on a gfx1030

vchuravy@odin ~> ./debug_pool
Device: AMD Radeon RX 6800 XT  arch: gfx1030

[hipMalloc]             OK
[hipMallocAsync]        OK
[hipMallocFromPoolAsync] OK

vchuravy@odin ~> hipcc --version
HIP version: 7.2.53211-9999
AMD clang version 22.0.0git (/srcdest/rocm-llvm f58b06dce1f9c15707c5f808fd002e18c2accf7e)
Target: x86_64-pc-linux-gnu
Thread model: posix
InstalledDir: /opt/rocm/lib/llvm/bin

[luraess]

Uh, interesting... So may be specific to gfx1101

lraess@gpuci:~/data/AMDGPU.jl$ hipcc --version
HIP version: 7.2.53211-671d39a71e
AMD clang version 22.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-7.2.2 26084 f58b06dce1f9c15707c5f808fd002e18c2accf7e)
Target: x86_64-unknown-linux-gnu
Thread model: posix
InstalledDir: /opt/rocm-7.2.2/lib/llvm/bin

[luraess]

Can confirm it also fails on RX 7900 XTX gfx1100

Device: AMD Radeon RX 7900 XTX  arch: gfx1100

[hipMalloc]             OK
[hipMallocAsync]        out of memory
[hipMallocFromPoolAsync] out of memory
root@amdgpu1-debug:/mnt/host/data/AMDGPU.jl# 

and

root@amdgpu1-debug:/mnt/host/data/AMDGPU.jl# hipcc --version
HIP version: 7.2.26015-fc0010cf6a
AMD clang version 22.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-7.2.0 26014 7b800a19466229b8479a78de19143dc33c3ab9b5)
Target: x86_64-unknown-linux-gnu
Thread model: posix
InstalledDir: /opt/rocm-7.2.0/lib/llvm/bin