[Windows][Strix Halo gfx1151] HIP allocations capped at OS-visible RAM when BIOS UMA/VGM = 96 GB (ROCm 10.1, driver 26.10.41.05) #8922

"""HIP allocation ceiling test for Windows (Strix Halo gfx1151).
Allocates 4 GiB blocks with hipMalloc until 2 consecutive failures,
memsets + syncs each block, re-touches all at the end, then frees.
Usage:  set OCL_SET_SVM_SIZE=262144  (optional)
        python hip_alloc_test.py [max_gib] [dll_path]
"""
import ctypes, os, sys, time

MAX_GIB = int(sys.argv[1]) if len(sys.argv) > 1 else 120
DLL = sys.argv[2] if len(sys.argv) > 2 else "amdhip64_7.dll"
BLOCK = 4 << 30
GIB = 1 << 30

if os.path.isdir(r"C:\ROCm\10.1\bin"):
    os.add_dll_directory(r"C:\ROCm\10.1\bin")
hip = ctypes.WinDLL(DLL, winmode=0)

def err(code):
    hip.hipGetErrorString.restype = ctypes.c_char_p
    return f"{code} {hip.hipGetErrorString(code).decode()}"

def meminfo():
    free, total = ctypes.c_size_t(), ctypes.c_size_t()
    rc = hip.hipMemGetInfo(ctypes.byref(free), ctypes.byref(total))
    return f"free {free.value/GIB:.2f} / total {total.value/GIB:.2f} GiB" if rc == 0 else f"hipMemGetInfo failed: {err(rc)}"

ver = ctypes.c_int()
hip.hipRuntimeGetVersion(ctypes.byref(ver))
print(f"HIP runtime {ver.value}, OCL_SET_SVM_SIZE={os.environ.get('OCL_SET_SVM_SIZE')}")
print("before:", meminfo())

ptrs, fails = [], 0
while len(ptrs) * BLOCK < MAX_GIB * GIB and fails < 2:
    p = ctypes.c_void_p()
    t = time.time()
    rc = hip.hipMalloc(ctypes.byref(p), ctypes.c_size_t(BLOCK))
    if rc != 0:
        fails += 1
        print(f"block {len(ptrs)+1}: hipMalloc FAIL {err(rc)}")
        continue
    fails = 0
    rc = hip.hipMemset(p, 0xAB, ctypes.c_size_t(BLOCK)) or hip.hipDeviceSynchronize()
    ptrs.append(p)
    print(f"block {len(ptrs)}: OK, total {len(ptrs)*4} GiB, {time.time()-t:.1f}s" + (f", memset/sync {err(rc)}" if rc else ""))

print(f"max allocated: {len(ptrs)*4} GiB")
print("after alloc:", meminfo())
errors = sum(1 for p in ptrs if hip.hipMemset(p, 0x5A, ctypes.c_size_t(BLOCK)) != 0)
rc = hip.hipDeviceSynchronize()
print(f"final touch: {errors} memset errors, sync {err(rc)}")
for p in ptrs:
    hip.hipFree(p)
print("after free:", meminfo())