From 23dc6e12f79411a56d7e2f43c222db9f7d4aebc9 Mon Sep 17 00:00:00 2001 From: David Schneller Date: Mon, 30 Mar 2026 02:54:08 +0200 Subject: [PATCH] feat: switch to coarse-grained managed HIP allocations by default --- interfaces/hip/Memory.cpp | 5 +++++ 1 file changed, 5 insertions(+) diff --git a/interfaces/hip/Memory.cpp b/interfaces/hip/Memory.cpp index 6fff772..b4ef78a 100644 --- a/interfaces/hip/Memory.cpp +++ b/interfaces/hip/Memory.cpp @@ -26,11 +26,16 @@ void* ConcreteAPI::allocUnifiedMem(size_t size, bool compress, Destination hint) isFlagSet(status); void* devPtr; APIWRAP(hipMallocManaged(&devPtr, size, hipMemAttachGlobal)); + + // make coarse-grained memory access behavior the default (match with allocGlobMem) + APIWRAP(hipMemAdvise(devPtr, size, hipMemAdviseSetCoarseGrain, 1)); + if (hint == Destination::Host) { APIWRAP(hipMemAdvise(devPtr, size, hipMemAdviseSetPreferredLocation, hipCpuDeviceId)); } else { APIWRAP(hipMemAdvise(devPtr, size, hipMemAdviseSetPreferredLocation, getDeviceId())); } + statistics.allocatedMemBytes += size; statistics.allocatedUnifiedMemBytes += size; memToSizeMap[devPtr] = size;