From ef75e2133d046223c4a21a6702660524901cf612 Mon Sep 17 00:00:00 2001 From: Satyanvesh Dittakavi Date: Wed, 30 Sep 2026 13:02:24 +0000 Subject: [PATCH 1/2] fix(clr): match SVM frees by range so hipFree accepts offset pointers --- projects/clr/rocclr/platform/memory.cpp | 20 ++++++++++++++++---- projects/clr/rocclr/platform/memory.hpp | 2 +- 2 files changed, 17 insertions(+), 5 deletions(-) diff --git a/projects/clr/rocclr/platform/memory.cpp b/projects/clr/rocclr/platform/memory.cpp index bcdb52d4d220..13f64b3ee6ba 100644 --- a/projects/clr/rocclr/platform/memory.cpp +++ b/projects/clr/rocclr/platform/memory.cpp @@ -1510,9 +1510,20 @@ void SvmBuffer::Add(uintptr_t k, uintptr_t v) { Allocated_.insert(std::pair(k, v)); } -bool SvmBuffer::Remove(uintptr_t k) { +bool SvmBuffer::Remove(uintptr_t ptr, uintptr_t* base) { std::scoped_lock lock(AllocatedLock_); - return Allocated_.erase(k) != 0; + // ptr may not be base pointer, so look up the allocation that contains it. + auto it = Allocated_.upper_bound(ptr); + if (it == Allocated_.begin()) { + return false; + } + --it; + if (ptr < it->first || ptr >= it->second) { + return false; + } + *base = it->first; + Allocated_.erase(it); + return true; } bool SvmBuffer::Contains(uintptr_t ptr) { @@ -1539,10 +1550,11 @@ void* SvmBuffer::malloc(Context& context, cl_svm_mem_flags flags, size_t size, s } bool SvmBuffer::free(const Context& context, void* ptr) { - if (!Remove(reinterpret_cast(ptr))) { + uintptr_t base_ptr = 0; + if (!Remove(reinterpret_cast(ptr), &base_ptr)) { return false; } - context.svmFree(ptr); + context.svmFree(reinterpret_cast(base_ptr)); return true; } diff --git a/projects/clr/rocclr/platform/memory.hpp b/projects/clr/rocclr/platform/memory.hpp index df21af6bcf8e..e74f93b41e56 100644 --- a/projects/clr/rocclr/platform/memory.hpp +++ b/projects/clr/rocclr/platform/memory.hpp @@ -706,7 +706,7 @@ class SvmBuffer : AllStatic { private: static void Add(uintptr_t k, uintptr_t v); - static bool Remove(uintptr_t k); + static bool Remove(uintptr_t ptr, uintptr_t* base_ptr); static bool Contains(uintptr_t ptr); static std::map Allocated_; // !< Allocated buffers From a9acc8baae9a372a2686d70f57eb1bdba5d83de7 Mon Sep 17 00:00:00 2001 From: Satyanvesh Dittakavi Date: Wed, 30 Sep 2026 15:51:40 +0000 Subject: [PATCH 2/2] Add catch tests verifying the hipFree behavior --- .../catch/config/configs/unit/memory.yaml | 10 +++ .../hip-tests/catch/unit/memory/hipFree.cc | 62 +++++++++++++++++++ 2 files changed, 72 insertions(+) diff --git a/projects/hip-tests/catch/config/configs/unit/memory.yaml b/projects/hip-tests/catch/config/configs/unit/memory.yaml index 2a9f7008af42..071764cdd272 100644 --- a/projects/hip-tests/catch/config/configs/unit/memory.yaml +++ b/projects/hip-tests/catch/config/configs/unit/memory.yaml @@ -2146,6 +2146,16 @@ memory: # Tests double free which throws off the sanitizer disabled: [asan] tags: [alloc, P1] + Unit_hipFreeInteriorPtrDev: + <<: *level_2 + # Frees the same allocation twice which throws off the sanitizer + disabled: [asan] + tags: [alloc, P1] + Unit_hipFreeInteriorPtrHost: + <<: *level_2 + # Frees the same allocation twice which throws off the sanitizer + disabled: [asan] + tags: [alloc, P1] Unit_hipFreeDoubleArrayFree: <<: *level_2 tags: [alloc] diff --git a/projects/hip-tests/catch/unit/memory/hipFree.cc b/projects/hip-tests/catch/unit/memory/hipFree.cc index 8336f502b960..9a4cd5a708bb 100644 --- a/projects/hip-tests/catch/unit/memory/hipFree.cc +++ b/projects/hip-tests/catch/unit/memory/hipFree.cc @@ -216,6 +216,68 @@ HIP_TEST_CASE(Unit_hipFreeDoubleHost) { HIP_CHECK_ERROR(hipHostFree(ptr), hipErrorInvalidValue); } +// Freeing device memory through a pointer that is not the base address +HIP_TEST_CASE(Unit_hipFreeInteriorPtrDev) { + constexpr size_t size = 1024; + char* ptr{}; + HIP_CHECK(hipMalloc(&ptr, size)); + + SECTION("InteriorPtr") { +#if HT_NVIDIA + HIP_CHECK_ERROR(hipFree(ptr + size / 2), hipErrorInvalidValue); + HIP_CHECK(hipFree(ptr)); +#else + // The whole allocation is released, so the base address is no longer valid + HIP_CHECK(hipFree(ptr + size / 2)); + HIP_CHECK_ERROR(hipFree(ptr), hipErrorInvalidValue); +#endif + } + SECTION("LastBytePtr") { +#if HT_NVIDIA + HIP_CHECK_ERROR(hipFree(ptr + size - 1), hipErrorInvalidValue); + HIP_CHECK(hipFree(ptr)); +#else + HIP_CHECK(hipFree(ptr + size - 1)); + HIP_CHECK_ERROR(hipFree(ptr), hipErrorInvalidValue); +#endif + } + SECTION("PtrPastEnd") { + HIP_CHECK_ERROR(hipFree(ptr + size), hipErrorInvalidValue); + HIP_CHECK(hipFree(ptr)); + } +} + +// Freeing pinned host memory through a pointer that is not the base address +HIP_TEST_CASE(Unit_hipFreeInteriorPtrHost) { + constexpr size_t size = 1024; + char* ptr{}; + HIP_CHECK(hipHostMalloc(&ptr, size)); + + SECTION("InteriorPtr") { +#if HT_NVIDIA + HIP_CHECK_ERROR(hipHostFree(ptr + size / 2), hipErrorInvalidValue); + HIP_CHECK(hipHostFree(ptr)); +#else + // The whole allocation is released, so the base address is no longer valid + HIP_CHECK(hipHostFree(ptr + size / 2)); + HIP_CHECK_ERROR(hipHostFree(ptr), hipErrorInvalidValue); +#endif + } + SECTION("LastBytePtr") { +#if HT_NVIDIA + HIP_CHECK_ERROR(hipHostFree(ptr + size - 1), hipErrorInvalidValue); + HIP_CHECK(hipHostFree(ptr)); +#else + HIP_CHECK(hipHostFree(ptr + size - 1)); + HIP_CHECK_ERROR(hipHostFree(ptr), hipErrorInvalidValue); +#endif + } + SECTION("PtrPastEnd") { + HIP_CHECK_ERROR(hipHostFree(ptr + size), hipErrorInvalidValue); + HIP_CHECK(hipHostFree(ptr)); + } +} + #if HT_NVIDIA HIP_TEST_CASE(Unit_hipFreeDoubleArrayFree) { HIP_SKIP_TEST("tracked issue EXSWCPHIPT-120.");