Skip to content

fix(clr): match SVM frees by range so hipFree accepts offset pointers #12525

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
satyanveshd merged 2 commits into from
Oct 8, 2026
Merged

fix(clr): match SVM frees by range so hipFree accepts offset pointers #12525

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
ef75e21
fix(clr): match SVM frees by range so hipFree accepts offset pointers
satyanveshd Sep 30, 2026
a9acc8b
Add catch tests verifying the hipFree behavior
satyanveshd Sep 30, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
20 changes: 16 additions & 4 deletions projects/clr/rocclr/platform/memory.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1510,9 +1510,20 @@ void SvmBuffer::Add(uintptr_t k, uintptr_t v) {
Allocated_.insert(std::pair<uintptr_t, uintptr_t>(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);
Comment thread
satyanveshd marked this conversation as resolved.
Resolved
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) {
Expand All @@ -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<uintptr_t>(ptr))) {
uintptr_t base_ptr = 0;
if (!Remove(reinterpret_cast<uintptr_t>(ptr), &base_ptr)) {
return false;
}
context.svmFree(ptr);
context.svmFree(reinterpret_cast<void*>(base_ptr));
return true;
}

Expand Down
2 changes: 1 addition & 1 deletion projects/clr/rocclr/platform/memory.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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<uintptr_t, uintptr_t> Allocated_; // !< Allocated buffers
Expand Down
10 changes: 10 additions & 0 deletions projects/hip-tests/catch/config/configs/unit/memory.yaml
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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]
Expand Down
62 changes: 62 additions & 0 deletions projects/hip-tests/catch/unit/memory/hipFree.cc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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
Comment thread
chrispaquot marked this conversation as resolved.
Resolved
HIP_CHECK_ERROR(hipFree(ptr + size / 2), hipErrorInvalidValue);
HIP_CHECK(hipFree(ptr));
#else
Comment thread
satyanveshd marked this conversation as resolved.
Resolved
// 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
Comment thread
chrispaquot marked this conversation as resolved.
Resolved
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.");
Expand Down
Loading