[compiler-rt] [compiler-rt][asan] Add AMDGPU ASan support via HSA API interceptors. (PR #192240)
Amit Kumar Pandey via llvm-commits
llvm-commits at lists.llvm.org
Mon Jun 1 03:21:10 PDT 2026
https://github.com/ampandey-1995 updated https://github.com/llvm/llvm-project/pull/192240
>From 7d0e7d6111b1d8adde95a21b47c6b30861393866 Mon Sep 17 00:00:00 2001
From: Amit Pandey <pandey.kumaramit2023 at gmail.com>
Date: Wed, 15 Apr 2026 16:00:46 +0530
Subject: [PATCH 01/27] [compiler-rt][asan] Add AMDGPU ASan support via HSA API
interceptors.
MIME-Version: 1.0
Content-Type: text/plain; charset=UTF-8
Content-Transfer-Encoding: 8bit
Add initial AddressSanitizer support for AMDGPU by intercepting HSA APIs
used for host-side initialization and CPU↔GPU memory operations
(allocation, copies, IPC, vmem, and other related queries).
This support is guarded by `SANITIZER_AMDGPU`, which enables inclusion of
the ROCm HSA and COMgr headers required to build the interceptors.
---
compiler-rt/CMakeLists.txt | 20 +++
compiler-rt/lib/asan/asan_interceptors.cpp | 164 ++++++++++++++++++
.../lib/sanitizer_common/sanitizer_platform.h | 2 +-
3 files changed, 185 insertions(+), 1 deletion(-)
diff --git a/compiler-rt/CMakeLists.txt b/compiler-rt/CMakeLists.txt
index 39034fd9ba67d..8118255041a3d 100644
--- a/compiler-rt/CMakeLists.txt
+++ b/compiler-rt/CMakeLists.txt
@@ -549,6 +549,26 @@ elseif(COMPILER_RT_HAS_G_FLAG)
list(APPEND SANITIZER_COMMON_CFLAGS -g)
endif()
+if(SANITIZER_AMDGPU)
+ list(APPEND SANITIZER_COMMON_CFLAGS -DSANITIZER_AMDGPU=1)
+ message(STATUS "Looking 'hsa' and 'amd_comgr' header")
+ find_path(HSA_INCLUDE NAMES hsa.h HINTS ${SANITIZER_HSA_INCLUDE_PATH} /opt/rocm/include PATH_SUFFIXES hsa)
+ if(NOT HSA_INCLUDE)
+ message(FATAL_ERROR "Required header 'hsa.h' not found in path ${HSA_INCLUDE}. Aborting SANITIZER_AMDGPU build")
+ endif()
+ message(STATUS "Found 'hsa.h' in ${HSA_INCLUDE}")
+ include_directories(${HSA_INCLUDE})
+ find_path(COMgr_INCLUDE NAMES amd_comgr.h.in HINTS ${SANITIZER_COMGR_INCLUDE_PATH} PATH_SUFFIXES amd_comgr)
+ if(NOT COMgr_INCLUDE)
+ find_path(COMgr_INCLUDE NAMES amd_comgr.h HINTS /opt/rocm/include PATH_SUFFIXES amd_comgr)
+ if(NOT COMgr_INCLUDE)
+ message(FATAL_ERROR "Required header 'amd_comgr.h/amd_comgr.h.in' not found in path ${COMgr_INCLUDE}. Aborting SANITIZER_AMDGPU build")
+ endif()
+ endif()
+ message(STATUS "Found 'amd_comgr.h.in/amd_comgr.h' in ${COMgr_INCLUDE}")
+ include_directories(${COMgr_INCLUDE})
+endif()
+
if(LLVM_ENABLE_MODULES)
# Sanitizers cannot be built with -fmodules. The interceptors intentionally
# don't include system headers, which is incompatible with modules.
diff --git a/compiler-rt/lib/asan/asan_interceptors.cpp b/compiler-rt/lib/asan/asan_interceptors.cpp
index 6d024e58b27cf..924d36afc925b 100644
--- a/compiler-rt/lib/asan/asan_interceptors.cpp
+++ b/compiler-rt/lib/asan/asan_interceptors.cpp
@@ -895,6 +895,166 @@ DEFINE_REAL(int, vfork, )
DECLARE_EXTERN_INTERCEPTOR_AND_WRAPPER(int, vfork, )
# endif
+# if SANITIZER_AMDGPU
+void ENSURE_HSA_INITED();
+
+INTERCEPTOR(hsa_status_t, hsa_amd_memory_pool_allocate,
+ hsa_amd_memory_pool_t memory_pool, size_t size, uint32_t flags,
+ void** ptr) {
+ AsanInitFromRtl();
+ ENSURE_HSA_INITED();
+ GET_STACK_TRACE_MALLOC;
+ return asan_hsa_amd_memory_pool_allocate(memory_pool, size, flags, ptr,
+ &stack);
+}
+
+INTERCEPTOR(hsa_status_t, hsa_amd_memory_pool_free, void* ptr) {
+ AsanInitFromRtl();
+ ENSURE_HSA_INITED();
+ GET_STACK_TRACE_FREE;
+ return asan_hsa_amd_memory_pool_free(ptr, &stack);
+}
+
+INTERCEPTOR(hsa_status_t, hsa_amd_agents_allow_access, uint32_t num_agents,
+ const hsa_agent_t* agents, const uint32_t* flags, const void* ptr) {
+ AsanInitFromRtl();
+ ENSURE_HSA_INITED();
+ GET_STACK_TRACE_FREE;
+ return asan_hsa_amd_agents_allow_access(num_agents, agents, flags, ptr,
+ &stack);
+}
+
+INTERCEPTOR(hsa_status_t, hsa_memory_copy, void* dst, const void* src,
+ size_t size) {
+ AsanInitFromRtl();
+ ENSURE_HSA_INITED();
+ if (flags()->replace_intrin) {
+ if (dst != src) {
+ CHECK_RANGES_OVERLAP("hsa_memory_copy", dst, size, src, size);
+ }
+ ASAN_READ_RANGE(nullptr, src, size);
+ ASAN_WRITE_RANGE(nullptr, dst, size);
+ }
+ return REAL(hsa_memory_copy)(dst, src, size);
+}
+
+INTERCEPTOR(hsa_status_t, hsa_amd_memory_async_copy, void* dst,
+ hsa_agent_t dst_agent, const void* src, hsa_agent_t src_agent,
+ size_t size, uint32_t num_dep_signals,
+ const hsa_signal_t* dep_signals, hsa_signal_t completion_signal) {
+ AsanInitFromRtl();
+ ENSURE_HSA_INITED();
+ if (flags()->replace_intrin) {
+ if (dst != src) {
+ CHECK_RANGES_OVERLAP("hsa_amd_memory_async_copy", dst, size, src, size);
+ }
+ ASAN_READ_RANGE(nullptr, src, size);
+ ASAN_WRITE_RANGE(nullptr, dst, size);
+ }
+ return REAL(hsa_amd_memory_async_copy)(dst, dst_agent, src, src_agent, size,
+ num_dep_signals, dep_signals,
+ completion_signal);
+}
+
+# if HSA_AMD_INTERFACE_VERSION_MINOR >= 1
+INTERCEPTOR(hsa_status_t, hsa_amd_memory_async_copy_on_engine, void* dst,
+ hsa_agent_t dst_agent, const void* src, hsa_agent_t src_agent,
+ size_t size, uint32_t num_dep_signals,
+ const hsa_signal_t* dep_signals, hsa_signal_t completion_signal,
+ hsa_amd_sdma_engine_id_t engine_id, bool force_copy_on_sdma) {
+ AsanInitFromRtl();
+ ENSURE_HSA_INITED();
+ if (flags()->replace_intrin) {
+ if (dst != src) {
+ CHECK_RANGES_OVERLAP("hsa_amd_memory_async_copy_on_engine", dst, size,
+ src, size);
+ }
+ ASAN_READ_RANGE(nullptr, src, size);
+ ASAN_WRITE_RANGE(nullptr, dst, size);
+ }
+ return REAL(hsa_amd_memory_async_copy_on_engine)(
+ dst, dst_agent, src, src_agent, size, num_dep_signals, dep_signals,
+ completion_signal, engine_id, force_copy_on_sdma);
+}
+# endif
+
+INTERCEPTOR(hsa_status_t, hsa_amd_ipc_memory_create, void* ptr, size_t len,
+ hsa_amd_ipc_memory_t* handle) {
+ AsanInitFromRtl();
+ ENSURE_HSA_INITED();
+ return asan_hsa_amd_ipc_memory_create(ptr, len, handle);
+}
+
+INTERCEPTOR(hsa_status_t, hsa_amd_ipc_memory_attach,
+ const hsa_amd_ipc_memory_t* handle, size_t len, uint32_t num_agents,
+ const hsa_agent_t* mapping_agents, void** mapped_ptr) {
+ AsanInitFromRtl();
+ ENSURE_HSA_INITED();
+ return asan_hsa_amd_ipc_memory_attach(handle, len, num_agents, mapping_agents,
+ mapped_ptr);
+}
+
+INTERCEPTOR(hsa_status_t, hsa_amd_ipc_memory_detach, void* mapped_ptr) {
+ AsanInitFromRtl();
+ ENSURE_HSA_INITED();
+ return asan_hsa_amd_ipc_memory_detach(mapped_ptr);
+}
+
+INTERCEPTOR(hsa_status_t, hsa_amd_vmem_address_reserve_align, void** ptr,
+ size_t size, uint64_t address, uint64_t alignment, uint64_t flags) {
+ AsanInitFromRtl();
+ ENSURE_HSA_INITED();
+ GET_STACK_TRACE_MALLOC;
+ return asan_hsa_amd_vmem_address_reserve_align(ptr, size, address, alignment,
+ flags, &stack);
+}
+
+INTERCEPTOR(hsa_status_t, hsa_amd_vmem_address_free, void* ptr, size_t size) {
+ AsanInitFromRtl();
+ ENSURE_HSA_INITED();
+ GET_STACK_TRACE_FREE;
+ return asan_hsa_amd_vmem_address_free(ptr, size, &stack);
+}
+
+INTERCEPTOR(hsa_status_t, hsa_amd_pointer_info, const void* ptr,
+ hsa_amd_pointer_info_t* info, void* (*alloc)(size_t),
+ uint32_t* num_agents_accessible, hsa_agent_t** accessible) {
+ AsanInitFromRtl();
+ ENSURE_HSA_INITED();
+ return asan_hsa_amd_pointer_info(ptr, info, alloc, num_agents_accessible,
+ accessible);
+}
+
+INTERCEPTOR(hsa_status_t, hsa_init) {
+ AsanInitFromRtl();
+ ENSURE_HSA_INITED();
+ return asan_hsa_init();
+}
+
+void InitializeAmdgpuInterceptors() {
+ ASAN_INTERCEPT_FUNC(hsa_init);
+ ASAN_INTERCEPT_FUNC(hsa_memory_copy);
+ ASAN_INTERCEPT_FUNC(hsa_amd_memory_pool_allocate);
+ ASAN_INTERCEPT_FUNC(hsa_amd_memory_pool_free);
+ ASAN_INTERCEPT_FUNC(hsa_amd_agents_allow_access);
+ ASAN_INTERCEPT_FUNC(hsa_amd_memory_async_copy);
+# if HSA_AMD_INTERFACE_VERSION_MINOR >= 1
+ ASAN_INTERCEPT_FUNC(hsa_amd_memory_async_copy_on_engine);
+# endif
+ ASAN_INTERCEPT_FUNC(hsa_amd_ipc_memory_create);
+ ASAN_INTERCEPT_FUNC(hsa_amd_ipc_memory_attach);
+ ASAN_INTERCEPT_FUNC(hsa_amd_ipc_memory_detach);
+ ASAN_INTERCEPT_FUNC(hsa_amd_vmem_address_reserve_align);
+ ASAN_INTERCEPT_FUNC(hsa_amd_vmem_address_free);
+ ASAN_INTERCEPT_FUNC(hsa_amd_pointer_info);
+}
+
+void ENSURE_HSA_INITED() {
+ if (!REAL(hsa_init))
+ InitializeAmdgpuInterceptors();
+}
+# endif
+
// ---------------------- InitializeAsanInterceptors ---------------- {{{1
namespace __asan {
void InitializeAsanInterceptors() {
@@ -1011,6 +1171,10 @@ void InitializeAsanInterceptors() {
ASAN_INTERCEPT_FUNC(vfork);
# endif
+# if SANITIZER_AMDGPU
+ InitializeAmdgpuInterceptors();
+# endif
+
VReport(1, "AddressSanitizer: libc interceptors initialized\n");
}
diff --git a/compiler-rt/lib/sanitizer_common/sanitizer_platform.h b/compiler-rt/lib/sanitizer_common/sanitizer_platform.h
index 3d85959ac94a6..6b5ffd0af9595 100644
--- a/compiler-rt/lib/sanitizer_common/sanitizer_platform.h
+++ b/compiler-rt/lib/sanitizer_common/sanitizer_platform.h
@@ -315,7 +315,7 @@
# define SANITIZER_ALPHA 0
#endif
-#if defined(__AMDGPU__)
+#if SANITIZER_AMDGPU || defined(__AMDGPU__)
# define SANITIZER_AMDGPU 1
#else
# define SANITIZER_AMDGPU 0
>From 41adbe2ad94e5f20d5165106fa5bc7a3a13728b2 Mon Sep 17 00:00:00 2001
From: Amit Pandey <pandey.kumaramit2023 at gmail.com>
Date: Wed, 22 Apr 2026 15:10:01 +0530
Subject: [PATCH 02/27] [compiler-rt][HSA] DeviceAllocator route through
CombinedAllocator and HSA Allocator wrappers.
For SANITIZER_AMDGPU, the sanitizer `CombinedAllocator` gains an
optional `DeviceAllocator` (DeviceAllocatorT): init can turn it on,
Allocate can take a DeviceAllocationInfo* and send those requests to the
device path, and free / pointer queries / stats / locks also consider
device-owned chunks.
`asan_allocator` passes `da_info` into the underlying combined
allocator, enables the device side at InitLinkerInitialized on AMDGPU
builds, and adds HSA-related wrappers (hsa_init, pool alloc/free, agents
allow access, IPC, vmem, pointer_info) plus declarations/includes in
asan_allocator.h.
sanitizer_allocator.h pulls in HSA headers (when AMDGPU) and
sanitizer_allocator_device.h before the combined allocator.
---
compiler-rt/lib/asan/asan_allocator.cpp | 260 +++++++++++++++++-
compiler-rt/lib/asan/asan_allocator.h | 50 ++++
.../sanitizer_common/sanitizer_allocator.h | 27 +-
.../sanitizer_allocator_combined.h | 96 ++++++-
4 files changed, 412 insertions(+), 21 deletions(-)
diff --git a/compiler-rt/lib/asan/asan_allocator.cpp b/compiler-rt/lib/asan/asan_allocator.cpp
index 6af197ab4cb1c..7477f9e52e533 100644
--- a/compiler-rt/lib/asan/asan_allocator.cpp
+++ b/compiler-rt/lib/asan/asan_allocator.cpp
@@ -393,7 +393,14 @@ struct Allocator {
void InitLinkerInitialized(const AllocatorOptions& options) {
SetAllocatorMayReturnNull(options.may_return_null);
+#if SANITIZER_AMDGPU
+ // Device-backed HSA allocations (e.g. hsa_amd_memory_pool_allocate) use
+ // CombinedAllocator's device path; it must be enabled for InitMemFuncs/
+ // AmdgpuMemFuncs::Init to run at startup.
+ allocator.InitLinkerInitialized(options.release_to_os_interval_ms, 0, true);
+#else
allocator.InitLinkerInitialized(options.release_to_os_interval_ms);
+#endif
SharedInitCode(options);
max_user_defined_malloc_size = common_flags()->max_allocation_size_mb
? common_flags()->max_allocation_size_mb
@@ -540,7 +547,7 @@ struct Allocator {
// (true) or a fatal Report*+Die() (false).
void* AllocateImpl(uptr size, uptr alignment, BufferedStackTrace* stack,
AllocType alloc_type, bool can_fill,
- bool may_return_null) {
+ bool may_return_null,DeviceAllocationInfo *da_info = nullptr) {
if (UNLIKELY(!AsanInited()))
AsanInitFromRtl();
if (UNLIKELY(IsRssLimitExceeded())) {
@@ -595,11 +602,11 @@ struct Allocator {
void* allocated;
if (t) {
AllocatorCache* cache = GetAllocatorCache(&t->malloc_storage());
- allocated = allocator.Allocate(cache, needed_size, 8);
+ allocated = allocator.Allocate(cache, needed_size, 8, da_info);
} else {
SpinMutexLock l(&fallback_mutex);
- AllocatorCache* cache = &fallback_allocator_cache;
- allocated = allocator.Allocate(cache, needed_size, 8);
+ AllocatorCache *cache = &fallback_allocator_cache;
+ allocated = allocator.Allocate(cache, needed_size, 8, da_info);
}
if (UNLIKELY(!allocated)) {
SetAllocatorOutOfMemory();
@@ -1475,3 +1482,248 @@ int __asan_update_allocation_context(void* addr) {
GET_STACK_TRACE_MALLOC;
return instance.UpdateAllocationStack((uptr)addr, &stack);
}
+
+#if SANITIZER_AMDGPU
+
+DECLARE_REAL(hsa_status_t, hsa_init);
+DECLARE_REAL(hsa_status_t, hsa_amd_agents_allow_access, uint32_t num_agents,
+ const hsa_agent_t* agents, const uint32_t* flags, const void* ptr)
+DECLARE_REAL(hsa_status_t, hsa_amd_memory_pool_allocate,
+ hsa_amd_memory_pool_t memory_pool, size_t size, uint32_t flags,
+ void** ptr)
+DECLARE_REAL(hsa_status_t, hsa_amd_memory_pool_free, void* ptr)
+DECLARE_REAL(hsa_status_t, hsa_amd_ipc_memory_create, void* ptr, size_t len,
+ hsa_amd_ipc_memory_t* handle)
+DECLARE_REAL(hsa_status_t, hsa_amd_ipc_memory_attach,
+ const hsa_amd_ipc_memory_t* handle, size_t len,
+ uint32_t num_agents, const hsa_agent_t* mapping_agents,
+ void** mapped_ptr)
+DECLARE_REAL(hsa_status_t, hsa_amd_ipc_memory_detach, void* mapped_ptr)
+DECLARE_REAL(hsa_status_t, hsa_amd_vmem_address_reserve_align, void** ptr,
+ size_t size, uint64_t address, uint64_t alignment, uint64_t flags)
+DECLARE_REAL(hsa_status_t, hsa_amd_vmem_address_free, void* ptr, size_t size)
+DECLARE_REAL(hsa_status_t, hsa_amd_pointer_info, const void* ptr,
+ hsa_amd_pointer_info_t* info, void* (*alloc)(size_t),
+ uint32_t* num_agents_accessible, hsa_agent_t** accessible)
+DECLARE_REAL(hsa_status_t, hsa_amd_register_system_event_handler,
+ hsa_amd_system_event_callback_t, void*)
+
+namespace __asan {
+// Always align to page boundary to match current ROCr behavior
+static const size_t kPageSize_ = 4096;
+
+hsa_status_t asan_hsa_amd_memory_pool_allocate(
+ hsa_amd_memory_pool_t memory_pool, size_t size, uint32_t flags, void** ptr,
+ BufferedStackTrace* stack) {
+ AmdgpuAllocationInfo aa_info;
+ aa_info.alloc_func =
+ reinterpret_cast<void*>(asan_hsa_amd_memory_pool_allocate);
+ aa_info.memory_pool = memory_pool;
+ aa_info.size = size;
+ aa_info.flags = flags;
+ aa_info.ptr = nullptr;
+ SetErrnoOnNull(*ptr = instance.Allocate(size, kPageSize_, stack, FROM_MALLOC,
+ false, &aa_info));
+ return aa_info.status;
+}
+
+hsa_status_t asan_hsa_amd_memory_pool_free(void* ptr,
+ BufferedStackTrace* stack) {
+ void* p = get_allocator().GetBlockBegin(ptr);
+ if (p) {
+ instance.Deallocate(ptr, 0, 0, stack, FROM_MALLOC);
+ return HSA_STATUS_SUCCESS;
+ }
+ return REAL(hsa_amd_memory_pool_free)(ptr);
+}
+
+hsa_status_t asan_hsa_amd_agents_allow_access(uint32_t num_agents,
+ const hsa_agent_t* agents,
+ const uint32_t* flags,
+ const void* ptr,
+ BufferedStackTrace* stack) {
+ void* p = get_allocator().GetBlockBegin(ptr);
+ return REAL(hsa_amd_agents_allow_access)(num_agents, agents, flags,
+ p ? p : ptr);
+}
+
+// For asan allocator, kMetadataSize is 0 and maximum redzone size is 2048. This
+// implies for device allocation, the gap between user_beg and GetBlockBegin()
+// is always one kPageSize_
+// IPC calls use static_assert to make sure kMetadataSize = 0
+//
+# if SANITIZER_CAN_USE_ALLOCATOR64
+static struct AP64<LocalAddressSpaceView> AP_;
+# else
+static struct AP32<LocalAddressSpaceView> AP_;
+# endif
+
+hsa_status_t asan_hsa_amd_ipc_memory_create(void* ptr, size_t len,
+ hsa_amd_ipc_memory_t* handle) {
+ void* ptr_ = get_allocator().GetBlockBegin(ptr);
+ AsanChunk* m = ptr_
+ ? instance.GetAsanChunkByAddr(reinterpret_cast<uptr>(ptr_))
+ : nullptr;
+ if (ptr_ && m) {
+ static_assert(AP_.kMetadataSize == 0, "Expression below requires this");
+ uptr p = reinterpret_cast<uptr>(ptr);
+ uptr p_ = reinterpret_cast<uptr>(ptr_);
+ if (p == p_ + kPageSize_ && len == m->UsedSize()) {
+ size_t len_ = get_allocator().GetActuallyAllocatedSize(ptr_);
+ return REAL(hsa_amd_ipc_memory_create)(ptr_, len_, handle);
+ }
+ }
+ return REAL(hsa_amd_ipc_memory_create)(ptr, len, handle);
+}
+
+hsa_status_t asan_hsa_amd_ipc_memory_attach(const hsa_amd_ipc_memory_t* handle,
+ size_t len, uint32_t num_agents,
+ const hsa_agent_t* mapping_agents,
+ void** mapped_ptr) {
+ static_assert(AP_.kMetadataSize == 0, "Expression below requires this");
+ size_t len_ = len + kPageSize_;
+ hsa_status_t status = REAL(hsa_amd_ipc_memory_attach)(
+ handle, len_, num_agents, mapping_agents, mapped_ptr);
+ if (status == HSA_STATUS_SUCCESS && mapped_ptr) {
+ uptr mapped_base = reinterpret_cast<uptr>(*mapped_ptr);
+ uptr user_beg = mapped_base + kPageSize_;
+ uptr tail_beg = RoundUpTo(user_beg + len, ASAN_SHADOW_GRANULARITY);
+ uptr mapped_end = mapped_base + kPageSize_ + RoundUpTo(len, kPageSize_);
+
+ PoisonShadow(mapped_base, kPageSize_, kAsanHeapLeftRedzoneMagic);
+
+ if (mapped_end > tail_beg)
+ PoisonShadow(tail_beg, mapped_end - tail_beg, kAsanHeapLeftRedzoneMagic);
+
+ uptr size_rounded_down = RoundDownTo(len, ASAN_SHADOW_GRANULARITY);
+ if (size_rounded_down)
+ PoisonShadow(user_beg, size_rounded_down, 0);
+
+ if (len != size_rounded_down && CanPoisonMemory()) {
+ u8* shadow = (u8*)MemToShadow(user_beg + size_rounded_down);
+ *shadow = flags()->poison_partial
+ ? static_cast<u8>(len & (ASAN_SHADOW_GRANULARITY - 1))
+ : 0;
+ }
+
+ *mapped_ptr = reinterpret_cast<void*>(user_beg);
+ }
+ return status;
+}
+
+hsa_status_t asan_hsa_amd_ipc_memory_detach(void* mapped_ptr) {
+ static_assert(AP_.kMetadataSize == 0, "Expression below requires this");
+ uptr mapped_base = reinterpret_cast<uptr>(mapped_ptr) - kPageSize_;
+
+ hsa_amd_pointer_info_t info;
+ info.size = sizeof(hsa_amd_pointer_info_t);
+ if (REAL(hsa_amd_pointer_info)(reinterpret_cast<void*>(mapped_base), &info,
+ nullptr, nullptr,
+ nullptr) == HSA_STATUS_SUCCESS) {
+ PoisonShadow(mapped_base, info.sizeInBytes, 0);
+ FlushUnneededASanShadowMemory(mapped_base, info.sizeInBytes);
+ }
+
+ return REAL(hsa_amd_ipc_memory_detach)(reinterpret_cast<void*>(mapped_base));
+}
+
+hsa_status_t asan_hsa_amd_vmem_address_reserve_align(
+ void** ptr, size_t size, uint64_t address, uint64_t alignment,
+ uint64_t flags, BufferedStackTrace* stack) {
+ // Bypass the tracking for a fixed address since it cannot be supported.
+ // Reasons:
+ // 1. Address may not meet the alignment/page-size requirement.
+ // 2. Requested range overlaps an existing reserved/mapped range.
+ // 3. Insufficient VA space to honor that exact placement.
+ if (address)
+ return REAL(hsa_amd_vmem_address_reserve_align)(ptr, size, address,
+ alignment, flags);
+
+ if (alignment < kPageSize_)
+ alignment = kPageSize_;
+
+ if (UNLIKELY(!IsPowerOfTwo(alignment))) {
+ errno = errno_EINVAL;
+ return HSA_STATUS_ERROR_INVALID_ARGUMENT;
+ }
+
+ AmdgpuAllocationInfo aa_info;
+ aa_info.alloc_func =
+ reinterpret_cast<void*>(asan_hsa_amd_vmem_address_reserve_align);
+ aa_info.memory_pool = {0};
+ aa_info.size = size;
+ aa_info.flags64 = flags;
+ aa_info.address = 0;
+ aa_info.alignment = alignment;
+ aa_info.ptr = nullptr;
+ SetErrnoOnNull(*ptr = instance.Allocate(size, alignment, stack, FROM_MALLOC,
+ false, &aa_info));
+
+ return aa_info.status;
+}
+
+hsa_status_t asan_hsa_amd_vmem_address_free(void* ptr, size_t size,
+ BufferedStackTrace* stack) {
+ if (UNLIKELY(!IsAligned(reinterpret_cast<uptr>(ptr), kPageSize_))) {
+ errno = errno_EINVAL;
+ return HSA_STATUS_ERROR_INVALID_ARGUMENT;
+ }
+ if (size == 0) {
+ errno = errno_EINVAL;
+ return HSA_STATUS_ERROR_INVALID_ARGUMENT;
+ }
+
+ void* p = get_allocator().GetBlockBegin(ptr);
+ if (p) {
+ instance.Deallocate(ptr, 0, 0, stack, FROM_MALLOC);
+ return HSA_STATUS_SUCCESS;
+ }
+ return REAL(hsa_amd_vmem_address_free)(ptr, size);
+}
+
+hsa_status_t asan_hsa_amd_pointer_info(const void* ptr,
+ hsa_amd_pointer_info_t* info,
+ void* (*alloc)(size_t),
+ uint32_t* num_agents_accessible,
+ hsa_agent_t** accessible) {
+ void* ptr_ = get_allocator().GetBlockBegin(ptr);
+ AsanChunk* m = ptr_
+ ? instance.GetAsanChunkByAddr(reinterpret_cast<uptr>(ptr_))
+ : nullptr;
+ if (ptr_ && m) {
+ hsa_status_t status = REAL(hsa_amd_pointer_info)(
+ ptr_, info, alloc, num_agents_accessible, accessible);
+ if (status == HSA_STATUS_SUCCESS && info) {
+ static_assert(AP_.kMetadataSize == 0, "Expression below requires this");
+ // Adjust base address of agent,host and sizeInBytes so as to return
+ // the actual pointer information of user allocation rather than asan
+ // allocation. Asan allocation pointer info can be acquired using internal
+ // 'GetPointerInfo'
+ info->agentBaseAddress = reinterpret_cast<void*>(
+ reinterpret_cast<uptr>(info->agentBaseAddress) + kPageSize_);
+ info->hostBaseAddress = reinterpret_cast<void*>(
+ reinterpret_cast<uptr>(info->hostBaseAddress) + kPageSize_);
+ info->sizeInBytes = m->UsedSize();
+ }
+ return status;
+ }
+ return REAL(hsa_amd_pointer_info)(ptr, info, alloc, num_agents_accessible,
+ accessible);
+}
+
+hsa_status_t asan_hsa_init() {
+ hsa_status_t status = REAL(hsa_init)();
+ if (status == HSA_STATUS_SUCCESS) {
+ // Only clear state when recovering from a prior shutdown (avoids clearing
+ // amdgpu_event_registered on every refcount bump and re-registering).
+ if (__sanitizer::AmdgpuMemFuncs::IsAmdgpuRuntimeShutdown())
+ __sanitizer::AmdgpuMemFuncs::ClearAmdgpuRuntimeShutdownState();
+ // Load HSA entry points once the runtime is up; device allocator may stay
+ // disabled, but interceptors and RegisterSystemEventHandlers need them.
+ if (__sanitizer::AmdgpuMemFuncs::Init())
+ __sanitizer::AmdgpuMemFuncs::RegisterSystemEventHandlers();
+ }
+ return status;
+}
+#endif
+} // namespace __asan
diff --git a/compiler-rt/lib/asan/asan_allocator.h b/compiler-rt/lib/asan/asan_allocator.h
index 0e185dfa31bd1..785544959e0f9 100644
--- a/compiler-rt/lib/asan/asan_allocator.h
+++ b/compiler-rt/lib/asan/asan_allocator.h
@@ -330,4 +330,54 @@ void PrintInternalAllocatorStats();
void AsanSoftRssLimitExceededCallback(bool exceeded);
} // namespace __asan
+
+#if SANITIZER_AMDGPU
+
+# if defined(__has_include)
+# if __has_include("hsa.h")
+# include "hsa.h"
+# include "hsa_ext_amd.h"
+# elif __has_include("hsa/hsa.h")
+# include "hsa/hsa.h"
+# include "hsa/hsa_ext_amd.h"
+# endif
+# else
+# include "hsa/hsa.h"
+# include "hsa/hsa_ext_amd.h"
+# endif
+
+namespace __asan {
+hsa_status_t asan_hsa_amd_memory_pool_allocate(
+ hsa_amd_memory_pool_t memory_pool, size_t size, uint32_t flags, void** ptr,
+ BufferedStackTrace* stack);
+hsa_status_t asan_hsa_amd_memory_pool_free(void* ptr,
+ BufferedStackTrace* stack);
+hsa_status_t asan_hsa_amd_agents_allow_access(uint32_t num_agents,
+ const hsa_agent_t* agents,
+ const uint32_t* flags,
+ const void* ptr,
+ BufferedStackTrace* stack);
+hsa_status_t asan_hsa_amd_ipc_memory_create(void* ptr, size_t len,
+ hsa_amd_ipc_memory_t* handle);
+hsa_status_t asan_hsa_amd_ipc_memory_attach(const hsa_amd_ipc_memory_t* handle,
+ size_t len, uint32_t num_agents,
+ const hsa_agent_t* mapping_agents,
+ void** mapped_ptr);
+hsa_status_t asan_hsa_amd_ipc_memory_detach(void* mapped_ptr);
+hsa_status_t asan_hsa_amd_vmem_address_reserve_align(void** ptr, size_t size,
+ uint64_t address,
+ uint64_t alignment,
+ uint64_t flags,
+ BufferedStackTrace* stack);
+hsa_status_t asan_hsa_amd_vmem_address_free(void* ptr, size_t size,
+ BufferedStackTrace* stack);
+hsa_status_t asan_hsa_amd_pointer_info(const void* ptr,
+ hsa_amd_pointer_info_t* info,
+ void* (*alloc)(size_t),
+ uint32_t* num_agents_accessible,
+ hsa_agent_t** accessible);
+hsa_status_t asan_hsa_init();
+} // namespace __asan
+#endif
+
#endif // ASAN_ALLOCATOR_H
diff --git a/compiler-rt/lib/sanitizer_common/sanitizer_allocator.h b/compiler-rt/lib/sanitizer_common/sanitizer_allocator.h
index 6154f7810334b..dfaa50eabab13 100644
--- a/compiler-rt/lib/sanitizer_common/sanitizer_allocator.h
+++ b/compiler-rt/lib/sanitizer_common/sanitizer_allocator.h
@@ -24,6 +24,22 @@
#include "sanitizer_procmaps.h"
#include "sanitizer_type_traits.h"
+#if SANITIZER_AMDGPU
+
+# if defined(__has_include)
+# if __has_include("hsa.h")
+# include "hsa.h"
+# include "hsa_ext_amd.h"
+# elif __has_include("hsa/hsa.h")
+# include "hsa/hsa.h"
+# include "hsa/hsa_ext_amd.h"
+# endif
+# else
+# include "hsa/hsa.h"
+# include "hsa/hsa_ext_amd.h"
+# endif
+#endif
+
namespace __sanitizer {
// Allows the tools to name their allocations appropriately.
@@ -63,13 +79,14 @@ struct NoOpMapUnmapCallback {
void OnUnmap(uptr p, uptr size) const {}
};
-#include "sanitizer_allocator_size_class_map.h"
-#include "sanitizer_allocator_stats.h"
-#include "sanitizer_allocator_primary64.h"
-#include "sanitizer_allocator_primary32.h"
+#include "sanitizer_allocator_combined.h"
+#include "sanitizer_allocator_device.h"
#include "sanitizer_allocator_local_cache.h"
+#include "sanitizer_allocator_primary32.h"
+#include "sanitizer_allocator_primary64.h"
#include "sanitizer_allocator_secondary.h"
-#include "sanitizer_allocator_combined.h"
+#include "sanitizer_allocator_size_class_map.h"
+#include "sanitizer_allocator_stats.h"
bool IsRssLimitExceeded();
void SetRssLimitExceeded(bool limit_exceeded);
diff --git a/compiler-rt/lib/sanitizer_common/sanitizer_allocator_combined.h b/compiler-rt/lib/sanitizer_common/sanitizer_allocator_combined.h
index 49940d9b5d505..b3ec606f61cb1 100644
--- a/compiler-rt/lib/sanitizer_common/sanitizer_allocator_combined.h
+++ b/compiler-rt/lib/sanitizer_common/sanitizer_allocator_combined.h
@@ -29,19 +29,32 @@ class CombinedAllocator {
LargeMmapAllocatorPtrArray,
typename PrimaryAllocator::AddressSpaceView>;
- void InitLinkerInitialized(s32 release_to_os_interval_ms,
- uptr heap_start = 0) {
+#if SANITIZER_AMDGPU
+ using DeviceAllocator =
+ DeviceAllocatorT<typename PrimaryAllocator::MapUnmapCallback>;
+#endif
+
+ void InitLinkerInitialized(s32 release_to_os_interval_ms, uptr heap_start = 0,
+ bool enable_device_allocator = false) {
primary_.Init(release_to_os_interval_ms, heap_start);
secondary_.InitLinkerInitialized();
+#if SANITIZER_AMDGPU
+ device_.Init(enable_device_allocator, primary_.kMetadataSize);
+#endif
}
- void Init(s32 release_to_os_interval_ms, uptr heap_start = 0) {
+ void Init(s32 release_to_os_interval_ms, uptr heap_start = 0,
+ bool enable_device_allocator = false) {
stats_.Init();
primary_.Init(release_to_os_interval_ms, heap_start);
secondary_.Init();
+#if SANITIZER_AMDGPU
+ device_.Init(enable_device_allocator, primary_.kMetadataSize);
+#endif
}
- void *Allocate(AllocatorCache *cache, uptr size, uptr alignment) {
+ void* Allocate(AllocatorCache* cache, uptr size, uptr alignment,
+ DeviceAllocationInfo* da_info = nullptr) {
// Returning 0 on malloc(0) may break a lot of code.
if (size == 0)
size = 1;
@@ -65,7 +78,12 @@ class CombinedAllocator {
// alignment without such requirement, and allocating 'size' would use
// extraneous memory, so we employ 'original_size'.
void *res;
- if (primary_.CanAllocate(size, alignment))
+#if SANITIZER_AMDGPU
+ if (da_info)
+ res = device_.Allocate(&stats_, original_size, alignment, da_info);
+ else
+#endif
+ if (primary_.CanAllocate(size, alignment))
res = cache->Allocate(&primary_, primary_.ClassID(size));
else
res = secondary_.Allocate(&stats_, original_size, alignment);
@@ -90,8 +108,12 @@ class CombinedAllocator {
if (!p) return;
if (primary_.PointerIsMine(p))
cache->Deallocate(&primary_, primary_.GetSizeClass(p), p);
- else
+ else if (secondary_.PointerIsMine(p))
secondary_.Deallocate(&stats_, p);
+#if SANITIZER_AMDGPU
+ else if (device_.PointerIsMine(p))
+ device_.Deallocate(&stats_, p);
+#endif
}
void *Reallocate(AllocatorCache *cache, void *p, uptr new_size,
@@ -115,7 +137,13 @@ class CombinedAllocator {
bool PointerIsMine(const void *p) const {
if (primary_.PointerIsMine(p))
return true;
- return secondary_.PointerIsMine(p);
+ if (secondary_.PointerIsMine(p))
+ return true;
+#if SANITIZER_AMDGPU
+ if (device_.PointerIsMine(p))
+ return true;
+#endif
+ return false;
}
bool FromPrimary(const void *p) const { return primary_.PointerIsMine(p); }
@@ -123,31 +151,60 @@ class CombinedAllocator {
void *GetMetaData(const void *p) {
if (primary_.PointerIsMine(p))
return primary_.GetMetaData(p);
- return secondary_.GetMetaData(p);
+ if (secondary_.PointerIsMine(p))
+ return secondary_.GetMetaData(p);
+#if SANITIZER_AMDGPU
+ if (device_.PointerIsMine(p))
+ return device_.GetMetaData(p);
+#endif
+ return nullptr;
}
void *GetBlockBegin(const void *p) {
if (primary_.PointerIsMine(p))
return primary_.GetBlockBegin(p);
- return secondary_.GetBlockBegin(p);
+ if (secondary_.PointerIsMine(p))
+ return secondary_.GetBlockBegin(p);
+#if SANITIZER_AMDGPU
+ if (device_.PointerIsMine(p))
+ return device_.GetBlockBegin(p);
+#endif
+ return nullptr;
}
// This function does the same as GetBlockBegin, but is much faster.
// Must be called with the allocator locked.
void *GetBlockBeginFastLocked(const void *p) {
+ void* beg;
if (primary_.PointerIsMine(p))
return primary_.GetBlockBegin(p);
- return secondary_.GetBlockBeginFastLocked(p);
+ if ((beg = secondary_.GetBlockBeginFastLocked(p)))
+ return beg;
+#if SANITIZER_AMDGPU
+ if ((beg = device_.GetBlockBeginFastLocked(p)))
+ return beg;
+#endif
+ return nullptr;
}
uptr GetActuallyAllocatedSize(void *p) {
if (primary_.PointerIsMine(p))
return primary_.GetActuallyAllocatedSize(p);
- return secondary_.GetActuallyAllocatedSize(p);
+ if (secondary_.PointerIsMine(p))
+ return secondary_.GetActuallyAllocatedSize(p);
+#if SANITIZER_AMDGPU
+ if (device_.PointerIsMine(p))
+ return device_.GetActuallyAllocatedSize(p);
+#endif
+ return 0;
}
uptr TotalMemoryUsed() {
- return primary_.TotalMemoryUsed() + secondary_.TotalMemoryUsed();
+ return primary_.TotalMemoryUsed() + secondary_.TotalMemoryUsed()
+#if SANITIZER_AMDGPU
+ + device_.TotalMemoryUsed()
+#endif
+ ;
}
void TestOnlyUnmap() { primary_.TestOnlyUnmap(); }
@@ -171,11 +228,17 @@ class CombinedAllocator {
void PrintStats() {
primary_.PrintStats();
secondary_.PrintStats();
+#if SANITIZER_AMDGPU
+ device_.PrintStats();
+#endif
}
// ForceLock() and ForceUnlock() are needed to implement Darwin malloc zone
// introspection API.
void ForceLock() SANITIZER_NO_THREAD_SAFETY_ANALYSIS {
+#if SANITIZER_AMDGPU
+ device_.ForceLock();
+#endif
primary_.ForceLock();
secondary_.ForceLock();
}
@@ -183,6 +246,9 @@ class CombinedAllocator {
void ForceUnlock() SANITIZER_NO_THREAD_SAFETY_ANALYSIS {
secondary_.ForceUnlock();
primary_.ForceUnlock();
+#if SANITIZER_AMDGPU
+ device_.ForceUnlock();
+#endif
}
// Iterate over all existing chunks.
@@ -190,10 +256,16 @@ class CombinedAllocator {
void ForEachChunk(ForEachChunkCallback callback, void *arg) {
primary_.ForEachChunk(callback, arg);
secondary_.ForEachChunk(callback, arg);
+#if SANITIZER_AMDGPU
+ device_.ForEachChunk(callback, arg);
+#endif
}
private:
PrimaryAllocator primary_;
SecondaryAllocator secondary_;
+#if SANITIZER_AMDGPU
+ DeviceAllocator device_;
+#endif
AllocatorGlobalStats stats_;
};
>From 13c9082c1686bce47c71880635b7d35cf17cc6d7 Mon Sep 17 00:00:00 2001
From: Amit Pandey <pandey.kumaramit2023 at gmail.com>
Date: Wed, 22 Apr 2026 15:20:46 +0530
Subject: [PATCH 03/27] [compiler-rt][sanitizer_common] AMDGPU
`DeviceAllocatorT` and AmdgpuMemFuncs(HSA dlsym, shutdown hook)
Under `SANITIZER_AMDGPU`, this adds `DeviceAllocatorT`
(sanitizer_allocator_device.h) to track HSA-backed slabs (allocate/free,
chunk lookup, stats, shutdown fallbacks) and AmdgpuMemFuncs
(sanitizer_allocator_amdgpu.{h,cpp}) to dlsym ROCr hsa_amd_* entry
points, perform pool or vmem allocations, free by pointer type, query
hsa_amd_pointer_info, and register a one-shot shutdown callback so the
allocator can stop using HSA after runtime teardown.
---
.../lib/sanitizer_common/CMakeLists.txt | 1 +
.../sanitizer_allocator_amdgpu.cpp | 209 +++++++++++
.../sanitizer_allocator_amdgpu.h | 48 +++
.../sanitizer_allocator_device.h | 354 ++++++++++++++++++
4 files changed, 612 insertions(+)
create mode 100644 compiler-rt/lib/sanitizer_common/sanitizer_allocator_amdgpu.cpp
create mode 100644 compiler-rt/lib/sanitizer_common/sanitizer_allocator_amdgpu.h
create mode 100644 compiler-rt/lib/sanitizer_common/sanitizer_allocator_device.h
diff --git a/compiler-rt/lib/sanitizer_common/CMakeLists.txt b/compiler-rt/lib/sanitizer_common/CMakeLists.txt
index 96c23c6d8ab82..5e13b5b18cc56 100644
--- a/compiler-rt/lib/sanitizer_common/CMakeLists.txt
+++ b/compiler-rt/lib/sanitizer_common/CMakeLists.txt
@@ -3,6 +3,7 @@
set(SANITIZER_SOURCES_NOTERMINATION
sanitizer_allocator.cpp
+ sanitizer_allocator_amdgpu.cpp
sanitizer_common.cpp
sanitizer_deadlock_detector1.cpp
sanitizer_deadlock_detector2.cpp
diff --git a/compiler-rt/lib/sanitizer_common/sanitizer_allocator_amdgpu.cpp b/compiler-rt/lib/sanitizer_common/sanitizer_allocator_amdgpu.cpp
new file mode 100644
index 0000000000000..3597db7c9d1bf
--- /dev/null
+++ b/compiler-rt/lib/sanitizer_common/sanitizer_allocator_amdgpu.cpp
@@ -0,0 +1,209 @@
+//===-- sanitizer_allocator_amdgpu.cpp --------------------------*- C++ -*-===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+//
+// Part of the Sanitizer Allocator.
+//
+//===----------------------------------------------------------------------===//
+#if SANITIZER_AMDGPU
+# include <dlfcn.h> // For dlsym
+
+# include "sanitizer_allocator.h"
+# include "sanitizer_atomic.h"
+
+namespace __sanitizer {
+struct HsaFunctions {
+ // -------------- Memory Allocate/Deallocate Functions ----------------
+ hsa_status_t (*memory_pool_allocate)(hsa_amd_memory_pool_t memory_pool,
+ size_t size, uint32_t flags, void** ptr);
+ hsa_status_t (*memory_pool_free)(void* ptr);
+ hsa_status_t (*pointer_info)(void* ptr, hsa_amd_pointer_info_t* info,
+ void* (*alloc)(size_t),
+ uint32_t* num_agents_accessible,
+ hsa_agent_t** accessible);
+ hsa_status_t (*vmem_address_reserve_align)(void** ptr, size_t size,
+ uint64_t address,
+ uint64_t alignment,
+ uint64_t flags);
+ hsa_status_t (*vmem_address_free)(void* ptr, size_t size);
+
+ // ----------------- System Event Register Function -------------------
+ hsa_status_t (*register_system_event_handler)(
+ hsa_amd_system_event_callback_t callback, void* data);
+};
+
+static HsaFunctions hsa_amd;
+
+// Always align to page boundary to match current ROCr behavior
+static const size_t kPageSize_ = 4096;
+
+static atomic_uint8_t amdgpu_runtime_shutdown{0};
+static atomic_uint8_t amdgpu_event_registered{0};
+
+# define LOAD_HSA_FUNC_WITH_ERROR_CHECK(func, name, success) \
+ func = (decltype(func))dlsym(RTLD_NEXT, name); \
+ if (!func) { \
+ VReport(2, "Amdgpu Init: Failed to load " #name " function\n"); \
+ success = false; \
+ }
+
+// Check AMDGPU runtime shutdown state
+bool AmdgpuMemFuncs::IsAmdgpuRuntimeShutdown() {
+ return static_cast<bool>(
+ atomic_load(&amdgpu_runtime_shutdown, memory_order_acquire));
+}
+
+// Notify AMDGPU runtime shutdown to allocator
+void AmdgpuMemFuncs::NotifyAmdgpuRuntimeShutdown() {
+ uint8_t shutdown = 0;
+ if (atomic_compare_exchange_strong(&amdgpu_runtime_shutdown, &shutdown, 1,
+ memory_order_acq_rel)) {
+ VReport(2, "Amdgpu Allocator: AMDGPU runtime shutdown detected\n");
+ }
+}
+
+// Clear shutdown state when hsa_init() succeeds again (re-init after shutdown).
+// Resets amdgpu_runtime_shutdown so allocator operations are enabled, and
+// amdgpu_event_registered so RegisterSystemEventHandlers() will register the
+// shutdown callback for the new runtime instance.
+void AmdgpuMemFuncs::ClearAmdgpuRuntimeShutdownState() {
+ atomic_store(&amdgpu_runtime_shutdown, 0, memory_order_release);
+ atomic_store(&amdgpu_event_registered, 0, memory_order_release);
+}
+
+bool AmdgpuMemFuncs::Init() {
+ bool success = true;
+ LOAD_HSA_FUNC_WITH_ERROR_CHECK(hsa_amd.memory_pool_allocate,
+ "hsa_amd_memory_pool_allocate", success);
+ LOAD_HSA_FUNC_WITH_ERROR_CHECK(hsa_amd.memory_pool_free,
+ "hsa_amd_memory_pool_free", success);
+ LOAD_HSA_FUNC_WITH_ERROR_CHECK(hsa_amd.pointer_info, "hsa_amd_pointer_info",
+ success);
+ LOAD_HSA_FUNC_WITH_ERROR_CHECK(hsa_amd.vmem_address_reserve_align,
+ "hsa_amd_vmem_address_reserve_align", success);
+ LOAD_HSA_FUNC_WITH_ERROR_CHECK(hsa_amd.vmem_address_free,
+ "hsa_amd_vmem_address_free", success);
+ LOAD_HSA_FUNC_WITH_ERROR_CHECK(hsa_amd.register_system_event_handler,
+ "hsa_amd_register_system_event_handler",
+ success);
+ if (!success) {
+ VReport(1, "Amdgpu Init: Failed to load AMDGPU runtime functions\n");
+ return false;
+ }
+ return true;
+}
+
+void* AmdgpuMemFuncs::Allocate(uptr size, uptr alignment,
+ DeviceAllocationInfo* da_info) {
+ // Do not allocate if AMDGPU runtime is shutdown
+ if (UNLIKELY(IsAmdgpuRuntimeShutdown())) {
+ VReport(1,
+ "Amdgpu Allocate: Runtime shutdown, skipping allocation for size "
+ "%zu alignment %zu\n",
+ size, alignment);
+ return nullptr;
+ }
+
+ AmdgpuAllocationInfo* aa_info =
+ reinterpret_cast<AmdgpuAllocationInfo*>(da_info);
+ if (!aa_info->memory_pool.handle) {
+ aa_info->status = hsa_amd.vmem_address_reserve_align(
+ &aa_info->ptr, size, aa_info->address, aa_info->alignment,
+ aa_info->flags64);
+ } else {
+ aa_info->status = hsa_amd.memory_pool_allocate(
+ aa_info->memory_pool, size, aa_info->flags, &aa_info->ptr);
+ }
+ if (aa_info->status != HSA_STATUS_SUCCESS)
+ return nullptr;
+
+ return aa_info->ptr;
+}
+
+void AmdgpuMemFuncs::Deallocate(void* p) {
+ // Deallocate does nothing after AMDGPU runtime shutdown
+ if (UNLIKELY(IsAmdgpuRuntimeShutdown())) {
+ VReport(
+ 1,
+ "Amdgpu Deallocate: Runtime shutdown, skipping deallocation for %p\n",
+ reinterpret_cast<void*>(p));
+ return;
+ }
+
+ DevicePointerInfo DevPtrInfo;
+ if (AmdgpuMemFuncs::GetPointerInfo(reinterpret_cast<uptr>(p), &DevPtrInfo)) {
+ if (DevPtrInfo.type == HSA_EXT_POINTER_TYPE_HSA) {
+ UNUSED hsa_status_t status = hsa_amd.memory_pool_free(p);
+ } else if (DevPtrInfo.type == HSA_EXT_POINTER_TYPE_RESERVED_ADDR) {
+ UNUSED hsa_status_t status =
+ hsa_amd.vmem_address_free(p, DevPtrInfo.map_size);
+ }
+ }
+}
+
+bool AmdgpuMemFuncs::GetPointerInfo(uptr ptr, DevicePointerInfo* ptr_info) {
+ // GetPointerInfo returns false after AMDGPU runtime shutdown
+ if (UNLIKELY(IsAmdgpuRuntimeShutdown())) {
+ VReport(1,
+ "Amdgpu GetPointerInfo: Runtime shutdown, skipping query for %p\n",
+ reinterpret_cast<void*>(ptr));
+ return false;
+ }
+
+ hsa_amd_pointer_info_t info;
+ info.size = sizeof(hsa_amd_pointer_info_t);
+ hsa_status_t status =
+ hsa_amd.pointer_info(reinterpret_cast<void*>(ptr), &info, 0, 0, 0);
+
+ if (status != HSA_STATUS_SUCCESS)
+ return false;
+
+ if (info.type == HSA_EXT_POINTER_TYPE_RESERVED_ADDR)
+ ptr_info->map_beg = reinterpret_cast<uptr>(info.hostBaseAddress);
+ else if (info.type == HSA_EXT_POINTER_TYPE_HSA)
+ ptr_info->map_beg = reinterpret_cast<uptr>(info.agentBaseAddress);
+ ptr_info->map_size = info.sizeInBytes;
+ ptr_info->type = reinterpret_cast<hsa_amd_pointer_type_t>(info.type);
+
+ return true;
+}
+// Register shutdown system event handler only once
+// TODO: Register multiple event handlers if needed in future
+void AmdgpuMemFuncs::RegisterSystemEventHandlers() {
+ uint8_t registered = 0;
+ // Check if shutdown event handler is already registered
+ if (atomic_compare_exchange_strong(&amdgpu_event_registered, ®istered, 1,
+ memory_order_acq_rel)) {
+ // Callback to detect and notify AMDGPU runtime shutdown
+ hsa_amd_system_event_callback_t callback = [](const hsa_amd_event_t* event,
+ void* data) {
+ if (!event)
+ return HSA_STATUS_ERROR_INVALID_ARGUMENT;
+ if (event->event_type == HSA_AMD_SYSTEM_SHUTDOWN_EVENT)
+ AmdgpuMemFuncs::NotifyAmdgpuRuntimeShutdown();
+ return HSA_STATUS_SUCCESS;
+ };
+ // Register the event callback
+ hsa_status_t status =
+ hsa_amd.register_system_event_handler(callback, nullptr);
+ // Check as registered if successful
+ if (status == HSA_STATUS_SUCCESS)
+ VReport(
+ 1,
+ "Amdgpu RegisterSystemEventHandlers: Registered shutdown event \n");
+ else {
+ VReport(1,
+ "Amdgpu RegisterSystemEventHandlers: Failed to register shutdown "
+ "event \n");
+ atomic_store(&amdgpu_event_registered, 0, memory_order_release);
+ }
+ }
+}
+
+uptr AmdgpuMemFuncs::GetPageSize() { return kPageSize_; }
+} // namespace __sanitizer
+#endif // SANITIZER_AMDGPU
diff --git a/compiler-rt/lib/sanitizer_common/sanitizer_allocator_amdgpu.h b/compiler-rt/lib/sanitizer_common/sanitizer_allocator_amdgpu.h
new file mode 100644
index 0000000000000..ff77911efb546
--- /dev/null
+++ b/compiler-rt/lib/sanitizer_common/sanitizer_allocator_amdgpu.h
@@ -0,0 +1,48 @@
+//===-- sanitizer_allocator_amdgpu.h ----------------------------*- C++ -*-===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+//
+// Part of the Sanitizer Allocator.
+//
+//===----------------------------------------------------------------------===//
+#ifndef SANITIZER_ALLOCATOR_H
+# error This file must be included inside sanitizer_allocator_device.h
+#endif
+
+#if SANITIZER_AMDGPU
+class AmdgpuMemFuncs {
+ public:
+ static bool Init();
+ static void* Allocate(uptr size, uptr alignment,
+ DeviceAllocationInfo* da_info);
+ static void Deallocate(void* p);
+ static bool GetPointerInfo(uptr ptr, DevicePointerInfo* ptr_info);
+ static uptr GetPageSize();
+ static void RegisterSystemEventHandlers();
+ static bool IsAmdgpuRuntimeShutdown();
+ static void ClearAmdgpuRuntimeShutdownState();
+
+ private:
+ static void NotifyAmdgpuRuntimeShutdown();
+};
+
+struct AmdgpuAllocationInfo : public DeviceAllocationInfo {
+ AmdgpuAllocationInfo() : DeviceAllocationInfo(DAT_AMDGPU) {
+ status = HSA_STATUS_SUCCESS;
+ alloc_func = nullptr;
+ }
+ hsa_status_t status;
+ void* alloc_func;
+ hsa_amd_memory_pool_t memory_pool;
+ u64 alignment;
+ u64 address;
+ u64 flags64;
+ usize size;
+ u32 flags;
+ void* ptr;
+};
+#endif // SANITIZER_AMDGPU
diff --git a/compiler-rt/lib/sanitizer_common/sanitizer_allocator_device.h b/compiler-rt/lib/sanitizer_common/sanitizer_allocator_device.h
new file mode 100644
index 0000000000000..c6a79760e902a
--- /dev/null
+++ b/compiler-rt/lib/sanitizer_common/sanitizer_allocator_device.h
@@ -0,0 +1,354 @@
+//===-- sanitizer_allocator_device.h ----------------------------*- C++ -*-===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+//
+// Part of the Sanitizer Allocator.
+//
+//===----------------------------------------------------------------------===//
+#ifndef SANITIZER_ALLOCATOR_H
+# error This file must be included inside sanitizer_allocator.h
+#endif
+
+struct DeviceAllocationInfo;
+#if SANITIZER_AMDGPU
+// Device memory allocation usually requires additional information, we can put
+// all the additional information into a data structure DeviceAllocationInfo.
+// This is only a parent structure since different vendors may require
+// different allocation info.
+typedef enum {
+ DAT_UNKNOWN = 0,
+ DAT_AMDGPU = 1,
+} DeviceAllocationType;
+
+struct DeviceAllocationInfo {
+ DeviceAllocationInfo(DeviceAllocationType type = DAT_UNKNOWN) {
+ type_ = type;
+ }
+ DeviceAllocationType type_;
+};
+
+struct DevicePointerInfo {
+ u64 type;
+ uptr map_beg;
+ uptr map_size;
+};
+
+# include "sanitizer_allocator_amdgpu.h"
+
+template <class MapUnmapCallback = NoOpMapUnmapCallback>
+class DeviceAllocatorT {
+ public:
+ using PtrArrayT = DefaultLargeMmapAllocatorPtrArray;
+ using DeviceMemFuncs = AmdgpuMemFuncs;
+
+ void Init(bool enable, uptr kMetadataSize) {
+ internal_memset(this, 0, sizeof(*this));
+ enabled_ = enable;
+ if (!enable)
+ return;
+ kMetadataSize_ = kMetadataSize;
+ chunks_ = reinterpret_cast<uptr*>(ptr_array_.Init());
+ InitMemFuncs();
+ }
+
+ void* Allocate(AllocatorStats* stat, uptr size, uptr alignment,
+ DeviceAllocationInfo* da_info) {
+ if (!da_info || !InitMemFuncs())
+ return nullptr;
+
+ // Allocate an extra page for Metadata
+ if (kMetadataSize_ + (size % page_size_) > page_size_) {
+ size += page_size_;
+ }
+ CHECK(IsPowerOfTwo(alignment));
+ uptr map_size = RoundUpMapSize(size);
+ if (alignment > page_size_)
+ map_size += alignment;
+ // Overflow.
+ if (map_size < size) {
+ Report(
+ "WARNING: %s: DeviceAllocator allocation overflow: "
+ "0x%zx bytes with 0x%zx alignment requested\n",
+ SanitizerToolName, map_size, alignment);
+ return nullptr;
+ }
+ void* ptr = DeviceMemFuncs::Allocate(map_size, alignment, da_info);
+ if (!ptr)
+ return nullptr;
+ uptr map_beg = reinterpret_cast<uptr>(ptr);
+ CHECK(IsAligned(map_beg, page_size_));
+ MapUnmapCallback().OnMap(map_beg, map_size);
+ uptr map_end = map_beg + map_size;
+ uptr res = map_beg;
+ if (res & (alignment - 1)) // Align.
+ res += alignment - (res & (alignment - 1));
+ CHECK(IsAligned(res, alignment));
+ CHECK(IsAligned(res, page_size_));
+ CHECK_GE(res + size, map_beg);
+ CHECK_LE(res + size, map_end);
+ uptr size_log = MostSignificantSetBitIndex(map_size);
+ CHECK_LT(size_log, ARRAY_SIZE(stats.by_size_log));
+ {
+ SpinMutexLock l(&mutex_);
+ ptr_array_.EnsureSpace(n_chunks_);
+ uptr idx = n_chunks_++;
+ chunks_[idx] = map_beg;
+ chunks_sorted_ = false;
+ stats.n_allocs++;
+ stats.currently_allocated += map_size;
+ stats.max_allocated = Max(stats.max_allocated, stats.currently_allocated);
+ stats.by_size_log[size_log]++;
+ stat->Add(AllocatorStatAllocated, map_size);
+ stat->Add(AllocatorStatMapped, map_size);
+ }
+ return reinterpret_cast<void*>(res);
+ }
+
+ void Deallocate(AllocatorStats* stat, void* p) {
+ Header header, *h;
+ {
+ SpinMutexLock l(&mutex_);
+ uptr idx;
+ uptr p_ = reinterpret_cast<uptr>(p);
+ EnsureSortedChunks(); // Avoid doing the sort while iterating.
+ for (idx = 0; idx < n_chunks_; idx++) {
+ if (chunks_[idx] >= p_)
+ break;
+ }
+ CHECK_EQ(chunks_[idx], p_);
+ CHECK_LT(idx, n_chunks_);
+ h = GetHeader(chunks_[idx], &header);
+ chunks_[idx] = chunks_[--n_chunks_];
+ chunks_sorted_ = false;
+ stats.n_frees++;
+ stats.currently_allocated -= h->map_size;
+ stat->Sub(AllocatorStatAllocated, h->map_size);
+ stat->Sub(AllocatorStatMapped, h->map_size);
+ }
+ MapUnmapCallback().OnUnmap(h->map_beg, h->map_size);
+ DeviceMemFuncs::Deallocate(p);
+ }
+
+ uptr TotalMemoryUsed() {
+ Header header;
+ SpinMutexLock l(&mutex_);
+ uptr res = 0;
+ for (uptr i = 0; i < n_chunks_; i++) {
+ Header* h = GetHeader(chunks_[i], &header);
+ res += RoundUpMapSize(h->map_size);
+ }
+ return res;
+ }
+
+ bool PointerIsMine(const void* p) const {
+ return GetBlockBegin(p) != nullptr;
+ }
+
+ uptr GetActuallyAllocatedSize(void* p) {
+ Header header;
+ uptr p_ = reinterpret_cast<uptr>(p);
+ Header* h = GetHeaderAnyPointer(p_, &header);
+ return h ? h->map_size : 0;
+ }
+
+ void* GetMetaData(const void* p) {
+ Header header;
+ uptr p_ = reinterpret_cast<uptr>(p);
+ Header* h = GetHeaderAnyPointer(p_, &header);
+ return h ? reinterpret_cast<void*>(h->map_beg + h->map_size -
+ kMetadataSize_)
+ : nullptr;
+ }
+
+ void* GetBlockBegin(const void* ptr) const {
+ Header header;
+ if (!mem_funcs_inited_)
+ return nullptr;
+ uptr p = reinterpret_cast<uptr>(ptr);
+ SpinMutexLock l(&mutex_);
+ uptr nearest_chunk = 0;
+ // Cache-friendly linear search.
+ for (uptr i = 0; i < n_chunks_; i++) {
+ uptr ch = chunks_[i];
+ if (p < ch)
+ continue; // p is at left to this chunk, skip it.
+ if (p - ch < p - nearest_chunk)
+ nearest_chunk = ch;
+ }
+ if (!nearest_chunk)
+ return nullptr;
+ if (p != nearest_chunk) {
+ Header* h = GetHeader(nearest_chunk, &header);
+ CHECK_GE(nearest_chunk, h->map_beg);
+ CHECK_LT(nearest_chunk, h->map_beg + h->map_size);
+ CHECK_LE(nearest_chunk, p);
+ if (h->map_beg + h->map_size <= p) {
+ return nullptr;
+ }
+ }
+ return GetUser(nearest_chunk);
+ }
+
+ void EnsureSortedChunks() {
+ if (chunks_sorted_)
+ return;
+ Sort(reinterpret_cast<uptr*>(chunks_), n_chunks_);
+ chunks_sorted_ = true;
+ }
+
+ // This function does the same as GetBlockBegin, but is much faster.
+ // Must be called with the allocator locked.
+ void* GetBlockBeginFastLocked(const void* ptr) {
+ if (!mem_funcs_inited_)
+ return nullptr;
+ mutex_.CheckLocked();
+ uptr p = reinterpret_cast<uptr>(ptr);
+ uptr n = n_chunks_;
+ if (!n)
+ return nullptr;
+ EnsureSortedChunks();
+ Header header, *h;
+ h = GetHeader(chunks_[n - 1], &header);
+ uptr min_mmap_ = chunks_[0];
+ uptr max_mmap_ = chunks_[n - 1] + h->map_size;
+ if (p < min_mmap_)
+ return nullptr;
+ if (p >= max_mmap_) {
+ // TODO (bingma): If dev_runtime_unloaded_ = true, map_size is limited
+ // to one page and we might miss a valid 'ptr'. If we hit cases where
+ // this kind of miss is unacceptable, we will need to implement a full
+ // solution with higher cost
+ return nullptr;
+ }
+ uptr beg = 0, end = n - 1;
+ // This loop is a log(n) lower_bound. It does not check for the exact match
+ // to avoid expensive cache-thrashing loads.
+ while (end - beg >= 2) {
+ uptr mid = (beg + end) / 2; // Invariant: mid >= beg + 1
+ if (p < chunks_[mid])
+ end = mid - 1; // We are not interested in chunks[mid].
+ else
+ beg = mid; // chunks[mid] may still be what we want.
+ }
+
+ if (beg < end) {
+ CHECK_EQ(beg + 1, end);
+ // There are 2 chunks left, choose one.
+ if (p >= chunks_[end])
+ beg = end;
+ }
+
+ if (p != chunks_[beg]) {
+ h = GetHeader(chunks_[beg], &header);
+ CHECK_NE(h, nullptr);
+ if (p < h->map_beg)
+ return nullptr;
+ if (h->map_beg + h->map_size <= p) {
+ // TODO (bingma): See above TODO in this function
+ return nullptr;
+ }
+ }
+ return GetUser(chunks_[beg]);
+ }
+
+ void PrintStats() {
+ Printf(
+ "Stats: DeviceAllocator: allocated %zd times, "
+ "remains %zd (%zd K) max %zd M; by size logs: ",
+ stats.n_allocs, stats.n_allocs - stats.n_frees,
+ stats.currently_allocated >> 10, stats.max_allocated >> 20);
+ for (uptr i = 0; i < ARRAY_SIZE(stats.by_size_log); i++) {
+ uptr c = stats.by_size_log[i];
+ if (!c)
+ continue;
+ Printf("%zd:%zd; ", i, c);
+ }
+ Printf("\n");
+ }
+
+ // ForceLock() and ForceUnlock() are needed to implement Darwin malloc zone
+ // introspection API.
+ void ForceLock() SANITIZER_ACQUIRE(mutex_) { mutex_.Lock(); }
+
+ void ForceUnlock() SANITIZER_RELEASE(mutex_) { mutex_.Unlock(); }
+
+ // Iterate over all existing chunks.
+ // The allocator must be locked when calling this function.
+ void ForEachChunk(ForEachChunkCallback callback, void* arg) {
+ EnsureSortedChunks(); // Avoid doing the sort while iterating.
+ for (uptr i = 0; i < n_chunks_; i++) {
+ const uptr t = chunks_[i];
+ callback(t, arg);
+ // Consistency check: verify that the array did not change.
+ CHECK_EQ(chunks_[i], t);
+ }
+ }
+
+ private:
+ bool InitMemFuncs() {
+ if (!enabled_ || mem_funcs_inited_ || mem_funcs_init_count_ >= 2) {
+ return mem_funcs_inited_;
+ }
+ mem_funcs_inited_ = DeviceMemFuncs::Init();
+ mem_funcs_init_count_++;
+ if (mem_funcs_inited_)
+ page_size_ = DeviceMemFuncs::GetPageSize();
+ return mem_funcs_inited_;
+ }
+
+ typedef DevicePointerInfo Header;
+
+ Header* GetHeaderAnyPointer(uptr p, Header* h) const {
+ CHECK(IsAligned(p, page_size_));
+ return DeviceMemFuncs::GetPointerInfo(p, h) ? h : nullptr;
+ }
+
+ Header* GetHeader(uptr chunk, Header* h) const {
+ // Device allocator has dependency on device runtime. If device runtime
+ // is unloaded, GetPointerInfo() will fail. For such case, we can still
+ // return a valid value for map_beg, map_size will be limited to one page
+ if (LIKELY(!dev_runtime_unloaded_)) {
+ if (DeviceMemFuncs::GetPointerInfo(chunk, h))
+ return h;
+ // If GetPointerInfo() fails, we don't assume the runtime is unloaded yet.
+ // We just return a conservative single-page header. Here mark/check the
+ // runtime shutdown state
+ dev_runtime_unloaded_ = DeviceMemFuncs::IsAmdgpuRuntimeShutdown();
+ }
+ // If we reach here, device runtime is unloaded.
+ // Fallback: conservative single-page header
+ h->map_beg = chunk;
+ h->map_size = page_size_;
+ return h;
+ }
+
+ void* GetUser(const uptr ptr) const { return reinterpret_cast<void*>(ptr); }
+
+ uptr RoundUpMapSize(uptr size) {
+ return RoundUpTo(size, page_size_) + page_size_;
+ }
+
+ bool enabled_;
+ bool mem_funcs_inited_;
+ mutable bool dev_runtime_unloaded_;
+ // Maximum of mem_funcs_init_count_ is 2:
+ // 1. The initial init called from Init(...), it could fail if
+ // libhsa-runtime64.so is dynamically loaded with dlopen()
+ // 2. A potential deferred init called by Allocate(...)
+ u32 mem_funcs_init_count_;
+ uptr kMetadataSize_;
+ uptr page_size_;
+ uptr* chunks_;
+ PtrArrayT ptr_array_;
+ uptr n_chunks_;
+ bool chunks_sorted_;
+ struct Stats {
+ uptr n_allocs, n_frees, currently_allocated, max_allocated, by_size_log[64];
+ } stats;
+ mutable StaticSpinMutex mutex_;
+};
+#endif // SANITIZER_AMDGPU
>From 0b6e351278435244879e58ed001fe363629aca5d Mon Sep 17 00:00:00 2001
From: Amit Pandey <pandey.kumaramit2023 at gmail.com>
Date: Wed, 22 Apr 2026 15:33:28 +0530
Subject: [PATCH 04/27] [compiler-rt][HSA-Test] Add AMDGPU ASan lit feature and
ROCm/HSA regression tests.
Wire the compiler-rt test suite for optional SANITIZER_AMDGPU builds and
add AddressSanitizer AMDGPU integration tests that link against ROCm's
libhsa-runtime64.
Lit/CMake
- pythonize_bool(SANITIZER_AMDGPU) in compiler-rt/test/CMakeLists.txt
and pass sanitizer_amdgpu into lit via lit.common.configured.in.
- When sanitizer_amdgpu is true, register the sanitizer-amdgpu lit
feature so tests can REQUIRES: sanitizer-amdgpu.
ASan AMDGPU tests (test/asan/TestCases/AMDGPU)
- lit.local.cfg.py: gate tests on Linux, dynamic ASan, ROCm discovery
(ROCM_PATH or /opt/rocm), and define %rocm_include, %rocm_lib, and
rpath substitutions for the shared ASan runtime and HSA DSO.
- hsa_memory_copy_overlap.cpp: overlapping hsa_memory_copy ranges
trigger the ASan overlap diagnostic.
- hsa_amd_memory_pool_allocate_double_free.cpp: allocate from a
runtime-alloc HSA memory pool then double hsa_amd_memory_pool_free;
expect ASan double-free reporting.
---
compiler-rt/test/CMakeLists.txt | 2 +
...a_amd_memory_pool_allocate_double_free.cpp | 89 +++++++++++++++++++
.../AMDGPU/hsa_memory_copy_overlap.cpp | 39 ++++++++
.../asan/TestCases/AMDGPU/lit.local.cfg.py | 58 ++++++++++++
compiler-rt/test/lit.common.cfg.py | 3 +
compiler-rt/test/lit.common.configured.in | 1 +
6 files changed, 192 insertions(+)
create mode 100644 compiler-rt/test/asan/TestCases/AMDGPU/hsa_amd_memory_pool_allocate_double_free.cpp
create mode 100644 compiler-rt/test/asan/TestCases/AMDGPU/hsa_memory_copy_overlap.cpp
create mode 100644 compiler-rt/test/asan/TestCases/AMDGPU/lit.local.cfg.py
diff --git a/compiler-rt/test/CMakeLists.txt b/compiler-rt/test/CMakeLists.txt
index 3fab82518e75f..0b28a532f8e22 100644
--- a/compiler-rt/test/CMakeLists.txt
+++ b/compiler-rt/test/CMakeLists.txt
@@ -16,6 +16,8 @@ pythonize_bool(COMPILER_RT_HAS_AARCH64_SME)
pythonize_bool(COMPILER_RT_HAS_NO_DEFAULT_CONFIG_FLAG)
+pythonize_bool(SANITIZER_AMDGPU)
+
if(LLVM_TREE_AVAILABLE OR NOT COMPILER_RT_STANDALONE_BUILD)
set(COMPILER_RT_BUILT_WITH_LLVM TRUE)
else()
diff --git a/compiler-rt/test/asan/TestCases/AMDGPU/hsa_amd_memory_pool_allocate_double_free.cpp b/compiler-rt/test/asan/TestCases/AMDGPU/hsa_amd_memory_pool_allocate_double_free.cpp
new file mode 100644
index 0000000000000..76899716e2398
--- /dev/null
+++ b/compiler-rt/test/asan/TestCases/AMDGPU/hsa_amd_memory_pool_allocate_double_free.cpp
@@ -0,0 +1,89 @@
+// RUN: %clangxx_asan -O0 -isystem %rocm_include %s -o %t -L%rocm_lib -lhsa-runtime64 \
+// RUN: -Wl,-rpath,%rocm_lib -Wl,-rpath,%compiler_rt_libdir
+// RUN: not %run %t 2>&1 | FileCheck %s
+//
+// Regression test for the AddressSanitizer hsa_amd_memory_pool_allocate /
+// hsa_amd_memory_pool_free interceptors: freeing the same pool allocation
+// twice is diagnosed (same family of checks as double-free on malloc).
+//
+// Link against ROCm's HSA runtime. Tests under TestCases/AMDGPU run only when
+// lit finds a ROCm install (see lit.local.cfg.py): $ROCM_PATH or /opt/rocm,
+// with include/hsa/hsa.h and libhsa-runtime64. Compiler-rt must be built with
+// SANITIZER_AMDGPU enabled. The suite uses the dynamic ASan runtime only.
+//
+// REQUIRES: sanitizer-amdgpu, linux, stable-runtime, rocm
+// UNSUPPORTED: android
+
+#include <hsa/hsa.h>
+#include <hsa/hsa_ext_amd.h>
+
+#include <stdio.h>
+
+struct PoolSearch {
+ hsa_amd_memory_pool_t pool;
+ bool found;
+};
+
+static hsa_status_t find_alloc_pool(hsa_amd_memory_pool_t pool, void *data) {
+ auto *ps = static_cast<PoolSearch *>(data);
+ bool allow = false;
+ if (hsa_amd_memory_pool_get_info(
+ pool, HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_ALLOWED, &allow) !=
+ HSA_STATUS_SUCCESS ||
+ !allow)
+ return HSA_STATUS_SUCCESS;
+ ps->pool = pool;
+ ps->found = true;
+ return HSA_STATUS_INFO_BREAK;
+}
+
+static hsa_status_t find_agent_with_pool(hsa_agent_t agent, void *data) {
+ (void)agent;
+ auto *ps = static_cast<PoolSearch *>(data);
+ ps->found = false;
+ hsa_status_t st =
+ hsa_amd_agent_iterate_memory_pools(agent, find_alloc_pool, ps);
+ if (st != HSA_STATUS_SUCCESS && st != HSA_STATUS_INFO_BREAK)
+ return st;
+ if (ps->found)
+ return HSA_STATUS_INFO_BREAK;
+ return HSA_STATUS_SUCCESS;
+}
+
+int main() {
+ if (hsa_init() != HSA_STATUS_SUCCESS) {
+ fprintf(stderr, "hsa_init failed\n");
+ return 1;
+ }
+
+ PoolSearch ps = {};
+ ps.pool.handle = 0;
+ ps.found = false;
+
+ hsa_status_t it = hsa_iterate_agents(find_agent_with_pool, &ps);
+ if (it != HSA_STATUS_SUCCESS && it != HSA_STATUS_INFO_BREAK) {
+ fprintf(stderr, "hsa_iterate_agents failed\n");
+ return 1;
+ }
+ if (!ps.found) {
+ fprintf(stderr, "no runtime-alloc HSA memory pool found\n");
+ return 1;
+ }
+
+ void *mem = nullptr;
+ if (hsa_amd_memory_pool_allocate(ps.pool, 64, 0, &mem) !=
+ HSA_STATUS_SUCCESS ||
+ !mem) {
+ fprintf(stderr, "hsa_amd_memory_pool_allocate failed\n");
+ return 1;
+ }
+
+ (void)hsa_amd_memory_pool_free(mem);
+ (void)hsa_amd_memory_pool_free(mem);
+
+ fprintf(stderr, "expected double-free report\n");
+ return 0;
+}
+
+// CHECK: ERROR: AddressSanitizer: attempting double-free
+// CHECK: SUMMARY: AddressSanitizer: double-free
diff --git a/compiler-rt/test/asan/TestCases/AMDGPU/hsa_memory_copy_overlap.cpp b/compiler-rt/test/asan/TestCases/AMDGPU/hsa_memory_copy_overlap.cpp
new file mode 100644
index 0000000000000..7895899a87dd5
--- /dev/null
+++ b/compiler-rt/test/asan/TestCases/AMDGPU/hsa_memory_copy_overlap.cpp
@@ -0,0 +1,39 @@
+// RUN: %clangxx_asan -O0 -isystem %rocm_include %s -o %t -L%rocm_lib -lhsa-runtime64 \
+// RUN: -Wl,-rpath,%rocm_lib -Wl,-rpath,%compiler_rt_libdir
+// RUN: not %run %t 2>&1 | FileCheck %s
+//
+// Regression test for the AddressSanitizer hsa_memory_copy interceptor: invalid
+// overlapping ranges are diagnosed (same family of checks as memcpy).
+//
+// Link against ROCm's HSA runtime. Tests under TestCases/AMDGPU run only when
+// lit finds a ROCm install (see lit.local.cfg.py): $ROCM_PATH or /opt/rocm,
+// with include/hsa/hsa.h and libhsa-runtime64. Compiler-rt must be built with
+// SANITIZER_AMDGPU enabled.
+//
+// REQUIRES: sanitizer-amdgpu, linux, stable-runtime, rocm
+// UNSUPPORTED: android
+
+#include <hsa/hsa.h>
+
+#include <stdio.h>
+#include <stdlib.h>
+
+int main() {
+ if (hsa_init() != HSA_STATUS_SUCCESS) {
+ fprintf(stderr, "hsa_init failed\n");
+ return 1;
+ }
+
+ char buf[128];
+ char *dst = buf;
+ char *src = buf + 40;
+ // Ranges [buf, buf+64) and [buf+40, buf+104) overlap; dst != src so the
+ // interceptor runs CHECK_RANGES_OVERLAP.
+ (void)hsa_memory_copy(dst, src, 64);
+ fprintf(stderr, "expected hsa_memory_copy overlap report\n");
+ return 0;
+}
+
+// CHECK: hsa_memory_copy-param-overlap: memory ranges
+// CHECK: [{{0x.*,[ ]*0x.*}}) and [{{0x.*,[ ]*0x.*}}) overlap
+// CHECK: SUMMARY: AddressSanitizer: hsa_memory_copy-param-overlap
diff --git a/compiler-rt/test/asan/TestCases/AMDGPU/lit.local.cfg.py b/compiler-rt/test/asan/TestCases/AMDGPU/lit.local.cfg.py
new file mode 100644
index 0000000000000..b825c27eacf21
--- /dev/null
+++ b/compiler-rt/test/asan/TestCases/AMDGPU/lit.local.cfg.py
@@ -0,0 +1,58 @@
+import glob
+import os
+
+
+def getRoot(config):
+ if not config.parent:
+ return config
+ return getRoot(config.parent)
+
+
+def rocm_lib_dir(rocm_root):
+ """Return lib or lib64 under rocm_root that provides libhsa-runtime64."""
+ for libname in ("lib", "lib64"):
+ libdir = os.path.join(rocm_root, libname)
+ if not os.path.isdir(libdir):
+ continue
+ if glob.glob(os.path.join(libdir, "libhsa-runtime64.so*")):
+ return libdir
+ return None
+
+
+def rocm_is_available(rocm_root):
+ if not rocm_root or not os.path.isdir(rocm_root):
+ return False
+ hsa_h = os.path.join(rocm_root, "include", "hsa", "hsa.h")
+ if not os.path.isfile(hsa_h):
+ return False
+ return rocm_lib_dir(rocm_root) is not None
+
+
+root = getRoot(config)
+
+# AMDGPU ASan tests are only run with the dynamic ASan runtime (-shared-libasan).
+if "asan-static-runtime" in root.available_features:
+ config.unsupported = True
+elif root.target_os != "Linux":
+ config.unsupported = True
+else:
+ rocm_root = os.environ.get("ROCM_PATH", "/opt/rocm")
+ if not rocm_is_available(rocm_root):
+ config.unsupported = True
+ else:
+ # Dynamic ASan (-shared-libasan) adds libclang_rt.asan*.so as DT_NEEDED; embed
+ # the host compiler-rt lib dir in RUNPATH so the loader finds it (same path as
+ # LD_LIBRARY_PATH in lit.common.cfg.py, but explicit in the linked binary).
+ rt_libdir = getattr(root, "compiler_rt_libdir", None)
+ if not rt_libdir or not os.path.isdir(rt_libdir):
+ config.unsupported = True
+ elif not glob.glob(os.path.join(rt_libdir, "libclang_rt.asan*.so")):
+ config.unsupported = True
+ else:
+ config.available_features.add("rocm")
+ rocm_lib = rocm_lib_dir(rocm_root)
+ rocm_include = os.path.join(rocm_root, "include")
+ config.substitutions.append(("%rocm_root", rocm_root))
+ config.substitutions.append(("%rocm_include", rocm_include))
+ config.substitutions.append(("%rocm_lib", rocm_lib))
+ config.substitutions.append(("%compiler_rt_libdir", rt_libdir))
diff --git a/compiler-rt/test/lit.common.cfg.py b/compiler-rt/test/lit.common.cfg.py
index 12446b4708d20..4821178500f49 100644
--- a/compiler-rt/test/lit.common.cfg.py
+++ b/compiler-rt/test/lit.common.cfg.py
@@ -562,6 +562,9 @@ def get_ios_commands_dir():
if config.gwp_asan:
config.available_features.add("gwp_asan")
+if getattr(config, "sanitizer_amdgpu", False):
+ config.available_features.add("sanitizer-amdgpu")
+
lit.util.usePlatformSdkOnDarwin(config, lit_config)
min_macos_deployment_target_substitutions = [
diff --git a/compiler-rt/test/lit.common.configured.in b/compiler-rt/test/lit.common.configured.in
index cad956aedf94a..0dde837fa52df 100644
--- a/compiler-rt/test/lit.common.configured.in
+++ b/compiler-rt/test/lit.common.configured.in
@@ -26,6 +26,7 @@ set_default("clang", "@COMPILER_RT_RESOLVED_TEST_COMPILER@")
set_default("compiler_id", "@COMPILER_RT_TEST_COMPILER_ID@")
set_default("python_executable", "@Python3_EXECUTABLE@")
set_default("compiler_rt_debug", @COMPILER_RT_DEBUG_PYBOOL@)
+set_default("sanitizer_amdgpu", @SANITIZER_AMDGPU_PYBOOL@)
set_default("compiler_rt_intercept_libdispatch", @COMPILER_RT_INTERCEPT_LIBDISPATCH_PYBOOL@)
set_default("compiler_rt_output_dir", "@COMPILER_RT_RESOLVED_OUTPUT_DIR@")
set_default("compiler_rt_bindir", "@COMPILER_RT_RESOLVED_EXEC_OUTPUT_DIR@")
>From 0288d411469b55fe784048c9bc86af9af10598e1 Mon Sep 17 00:00:00 2001
From: Amit Pandey <pandey.kumaramit2023 at gmail.com>
Date: Wed, 22 Apr 2026 15:37:45 +0530
Subject: [PATCH 05/27] [compiler-rt][sanitizer_common] Skip nolibc unittest
subprocess for SANITIZER_AMDGPU
MIME-Version: 1.0
Content-Type: text/plain; charset=UTF-8
Content-Transfer-Encoding: 8bit
For SANITIZER_AMDGPU builds, the sanitizer_common -nostdlib / nolibc
helper binary and NolibcMain subprocess test are turned off in CMake
(they clash with AMDGPU’s libdl / libc needs and with compiling every
unittest with HSA).
A new compile flag COMPILER_RT_SKIP_NOLIBC_SUBPROCESS_TEST gates the
test in sanitizer_nolibc_test.cpp so normal configs are unchanged.
---
.../lib/sanitizer_common/tests/CMakeLists.txt | 18 +++++++++++++++---
.../tests/sanitizer_nolibc_test.cpp | 5 ++++-
2 files changed, 19 insertions(+), 4 deletions(-)
diff --git a/compiler-rt/lib/sanitizer_common/tests/CMakeLists.txt b/compiler-rt/lib/sanitizer_common/tests/CMakeLists.txt
index 55c7d665e639f..b8c58041d31a4 100644
--- a/compiler-rt/lib/sanitizer_common/tests/CMakeLists.txt
+++ b/compiler-rt/lib/sanitizer_common/tests/CMakeLists.txt
@@ -80,6 +80,15 @@ set(SANITIZER_TEST_CFLAGS_COMMON
-Wno-gnu-zero-variadic-macro-arguments
)
+# Do not pass -DSANITIZER_AMDGPU=1 to all unittests: that would pull <hsa.h>
+# and AMDGPU-only types into every TU without ROCm include paths. When
+# SANITIZER_AMDGPU is enabled, skip the NolibcMain subprocess test instead
+# (see NOT SANITIZER_AMDGPU targets below and sanitizer_nolibc_test.cpp).
+if(SANITIZER_AMDGPU)
+ list(APPEND SANITIZER_TEST_CFLAGS_COMMON
+ -DCOMPILER_RT_SKIP_NOLIBC_SUBPROCESS_TEST=1)
+endif()
+
set(SANITIZER_TEST_LINK_FLAGS_COMMON
${COMPILER_RT_UNITTEST_LINK_FLAGS}
${COMPILER_RT_UNWINDER_LINK_LIBS}
@@ -183,10 +192,13 @@ macro(add_sanitizer_tests_for_arch arch)
CFLAGS ${SANITIZER_TEST_CFLAGS_COMMON} ${extra_flags}
LINK_FLAGS ${SANITIZER_TEST_LINK_FLAGS_COMMON} ${TARGET_LINK_FLAGS} ${extra_flags})
- if("${CMAKE_SYSTEM_NAME}" STREQUAL "Linux" AND "${arch}" STREQUAL "x86_64")
+ if("${CMAKE_SYSTEM_NAME}" STREQUAL "Linux" AND "${arch}" STREQUAL "x86_64"
+ AND NOT SANITIZER_AMDGPU)
# Test that the libc-independent part of sanitizer_common is indeed
# independent of libc, by linking this binary without libc (here) and
- # executing it (unit test in sanitizer_nolibc_test.cpp).
+ # executing it (unit test in sanitizer_nolibc_test.cpp). Omitted when
+ # SANITIZER_AMDGPU: AMDGPU runtime code needs libdl / libcdep and cannot
+ # link with -nostdlib.
get_target_flags_for_arch(${arch} TARGET_FLAGS)
clang_compile(sanitizer_nolibc_test_main.${arch}.o
sanitizer_nolibc_test_main.cpp
@@ -212,7 +224,7 @@ if(COMPILER_RT_CAN_EXECUTE_TESTS AND NOT ANDROID)
$<TARGET_OBJECTS:RTSanitizerCommonLibc.osx>
$<TARGET_OBJECTS:RTSanitizerCommonSymbolizer.osx>)
else()
- if(CAN_TARGET_x86_64)
+ if(CAN_TARGET_x86_64 AND NOT SANITIZER_AMDGPU)
add_sanitizer_common_lib("RTSanitizerCommon.test.nolibc.x86_64"
$<TARGET_OBJECTS:RTSanitizerCommon.x86_64>
$<TARGET_OBJECTS:RTSanitizerCommonNoLibc.x86_64>)
diff --git a/compiler-rt/lib/sanitizer_common/tests/sanitizer_nolibc_test.cpp b/compiler-rt/lib/sanitizer_common/tests/sanitizer_nolibc_test.cpp
index 41376ee094307..b17cc731aaa0a 100644
--- a/compiler-rt/lib/sanitizer_common/tests/sanitizer_nolibc_test.cpp
+++ b/compiler-rt/lib/sanitizer_common/tests/sanitizer_nolibc_test.cpp
@@ -19,7 +19,10 @@
extern const char *argv0;
-#if SANITIZER_LINUX && defined(__x86_64__)
+// When SANITIZER_AMDGPU is enabled, CMake defines this macro and does not build
+// Sanitizer-*-Test-Nolibc (see tests/CMakeLists.txt).
+#if SANITIZER_LINUX && defined(__x86_64__) && \
+ !defined(COMPILER_RT_SKIP_NOLIBC_SUBPROCESS_TEST)
TEST(SanitizerCommon, NolibcMain) {
std::string NolibcTestPath = argv0;
NolibcTestPath += "-Nolibc";
>From 9d46bca9eba0643cf3ff1d3292b8c3e29c028b83 Mon Sep 17 00:00:00 2001
From: Amit Pandey <pandey.kumaramit2023 at gmail.com>
Date: Wed, 22 Apr 2026 16:02:12 +0530
Subject: [PATCH 06/27] Fix the header inclusion order for device allocator.
---
.../lib/sanitizer_common/sanitizer_allocator.h | 13 ++++++-------
1 file changed, 6 insertions(+), 7 deletions(-)
diff --git a/compiler-rt/lib/sanitizer_common/sanitizer_allocator.h b/compiler-rt/lib/sanitizer_common/sanitizer_allocator.h
index dfaa50eabab13..1b058fd155d04 100644
--- a/compiler-rt/lib/sanitizer_common/sanitizer_allocator.h
+++ b/compiler-rt/lib/sanitizer_common/sanitizer_allocator.h
@@ -25,7 +25,6 @@
#include "sanitizer_type_traits.h"
#if SANITIZER_AMDGPU
-
# if defined(__has_include)
# if __has_include("hsa.h")
# include "hsa.h"
@@ -79,14 +78,14 @@ struct NoOpMapUnmapCallback {
void OnUnmap(uptr p, uptr size) const {}
};
-#include "sanitizer_allocator_combined.h"
-#include "sanitizer_allocator_device.h"
-#include "sanitizer_allocator_local_cache.h"
-#include "sanitizer_allocator_primary32.h"
-#include "sanitizer_allocator_primary64.h"
-#include "sanitizer_allocator_secondary.h"
#include "sanitizer_allocator_size_class_map.h"
#include "sanitizer_allocator_stats.h"
+#include "sanitizer_allocator_primary64.h"
+#include "sanitizer_allocator_primary32.h"
+#include "sanitizer_allocator_local_cache.h"
+#include "sanitizer_allocator_secondary.h"
+#include "sanitizer_allocator_device.h"
+#include "sanitizer_allocator_combined.h"
bool IsRssLimitExceeded();
void SetRssLimitExceeded(bool limit_exceeded);
>From 8bb1dde4985db69cdb12111a339446949014bf5b Mon Sep 17 00:00:00 2001
From: Amit Pandey <pandey.kumaramit2023 at gmail.com>
Date: Wed, 22 Apr 2026 16:59:28 +0530
Subject: [PATCH 07/27] Fix clang-format issue
---
compiler-rt/lib/sanitizer_common/sanitizer_allocator.h | 4 ++++
1 file changed, 4 insertions(+)
diff --git a/compiler-rt/lib/sanitizer_common/sanitizer_allocator.h b/compiler-rt/lib/sanitizer_common/sanitizer_allocator.h
index 1b058fd155d04..d803444f69dd6 100644
--- a/compiler-rt/lib/sanitizer_common/sanitizer_allocator.h
+++ b/compiler-rt/lib/sanitizer_common/sanitizer_allocator.h
@@ -78,6 +78,9 @@ struct NoOpMapUnmapCallback {
void OnUnmap(uptr p, uptr size) const {}
};
+// clang-format off
+// Include order is load-bearing (MemoryMapper, AllocatorStats, secondary
+// typedefs, then DeviceAllocatorT, then CombinedAllocator). Do not sort.
#include "sanitizer_allocator_size_class_map.h"
#include "sanitizer_allocator_stats.h"
#include "sanitizer_allocator_primary64.h"
@@ -86,6 +89,7 @@ struct NoOpMapUnmapCallback {
#include "sanitizer_allocator_secondary.h"
#include "sanitizer_allocator_device.h"
#include "sanitizer_allocator_combined.h"
+// clang-format on
bool IsRssLimitExceeded();
void SetRssLimitExceeded(bool limit_exceeded);
>From 530e3796591fd1db378fe0a2ee5204e3a2a502f8 Mon Sep 17 00:00:00 2001
From: Amit Pandey <pandey.kumaramit2023 at gmail.com>
Date: Wed, 22 Apr 2026 17:05:40 +0530
Subject: [PATCH 08/27] Fix extra closing brace error.
```
error: extraneous closing brace ('}')
1718 | } // namespace __asan
| ^
1 error generated.
```
---
compiler-rt/lib/asan/asan_allocator.cpp | 3 ++-
1 file changed, 2 insertions(+), 1 deletion(-)
diff --git a/compiler-rt/lib/asan/asan_allocator.cpp b/compiler-rt/lib/asan/asan_allocator.cpp
index 7477f9e52e533..4acc456a244ca 100644
--- a/compiler-rt/lib/asan/asan_allocator.cpp
+++ b/compiler-rt/lib/asan/asan_allocator.cpp
@@ -1725,5 +1725,6 @@ hsa_status_t asan_hsa_init() {
}
return status;
}
-#endif
} // namespace __asan
+
+#endif // SANITIZER_AMDGPU
>From 88762ebfbba6f4fd7092f7d560468c76c82c9948 Mon Sep 17 00:00:00 2001
From: Amit Pandey <pandey.kumaramit2023 at gmail.com>
Date: Thu, 23 Apr 2026 15:17:13 +0530
Subject: [PATCH 09/27] [compiler-rt][HSA-Test] Skip AMDGPU ROCm lit tests on
32-bit x86.
ROCm ships libhsa-runtime64 as a 64-bit DSO only, so i386 ASan runs fail
at link time (elf32-i386 vs 64-bit library). Mark TestCases/AMDGPU as
unsupported when the suite targets 32-bit x86.
---
...a_amd_memory_pool_allocate_double_free.cpp | 5 ----
.../AMDGPU/hsa_memory_copy_overlap.cpp | 5 ----
.../asan/TestCases/AMDGPU/lit.local.cfg.py | 23 +++++++++++++++----
3 files changed, 19 insertions(+), 14 deletions(-)
diff --git a/compiler-rt/test/asan/TestCases/AMDGPU/hsa_amd_memory_pool_allocate_double_free.cpp b/compiler-rt/test/asan/TestCases/AMDGPU/hsa_amd_memory_pool_allocate_double_free.cpp
index 76899716e2398..702ec5c100d25 100644
--- a/compiler-rt/test/asan/TestCases/AMDGPU/hsa_amd_memory_pool_allocate_double_free.cpp
+++ b/compiler-rt/test/asan/TestCases/AMDGPU/hsa_amd_memory_pool_allocate_double_free.cpp
@@ -6,11 +6,6 @@
// hsa_amd_memory_pool_free interceptors: freeing the same pool allocation
// twice is diagnosed (same family of checks as double-free on malloc).
//
-// Link against ROCm's HSA runtime. Tests under TestCases/AMDGPU run only when
-// lit finds a ROCm install (see lit.local.cfg.py): $ROCM_PATH or /opt/rocm,
-// with include/hsa/hsa.h and libhsa-runtime64. Compiler-rt must be built with
-// SANITIZER_AMDGPU enabled. The suite uses the dynamic ASan runtime only.
-//
// REQUIRES: sanitizer-amdgpu, linux, stable-runtime, rocm
// UNSUPPORTED: android
diff --git a/compiler-rt/test/asan/TestCases/AMDGPU/hsa_memory_copy_overlap.cpp b/compiler-rt/test/asan/TestCases/AMDGPU/hsa_memory_copy_overlap.cpp
index 7895899a87dd5..51f90366d40b2 100644
--- a/compiler-rt/test/asan/TestCases/AMDGPU/hsa_memory_copy_overlap.cpp
+++ b/compiler-rt/test/asan/TestCases/AMDGPU/hsa_memory_copy_overlap.cpp
@@ -5,11 +5,6 @@
// Regression test for the AddressSanitizer hsa_memory_copy interceptor: invalid
// overlapping ranges are diagnosed (same family of checks as memcpy).
//
-// Link against ROCm's HSA runtime. Tests under TestCases/AMDGPU run only when
-// lit finds a ROCm install (see lit.local.cfg.py): $ROCM_PATH or /opt/rocm,
-// with include/hsa/hsa.h and libhsa-runtime64. Compiler-rt must be built with
-// SANITIZER_AMDGPU enabled.
-//
// REQUIRES: sanitizer-amdgpu, linux, stable-runtime, rocm
// UNSUPPORTED: android
diff --git a/compiler-rt/test/asan/TestCases/AMDGPU/lit.local.cfg.py b/compiler-rt/test/asan/TestCases/AMDGPU/lit.local.cfg.py
index b825c27eacf21..7807c20f3a4cd 100644
--- a/compiler-rt/test/asan/TestCases/AMDGPU/lit.local.cfg.py
+++ b/compiler-rt/test/asan/TestCases/AMDGPU/lit.local.cfg.py
@@ -1,12 +1,25 @@
+# Link against ROCm's HSA runtime. Tests under TestCases/AMDGPU run only when
+# lit finds a ROCm install (see lit.local.cfg.py): $ROCM_PATH or /opt/rocm,
+# with include/hsa/hsa.h and libhsa-runtime64. Compiler-rt must be built with
+# SANITIZER_AMDGPU enabled. The suite uses the dynamic ASan runtime only.
+
import glob
import os
-
def getRoot(config):
if not config.parent:
return config
return getRoot(config.parent)
+def walk_config_attr(cfg, name):
+ """Return the first defined attribute `name` walking cfg -> parents."""
+ while cfg is not None:
+ if hasattr(cfg, name):
+ val = getattr(cfg, name)
+ if val is not None:
+ return val
+ cfg = cfg.parent
+ return None
def rocm_lib_dir(rocm_root):
"""Return lib or lib64 under rocm_root that provides libhsa-runtime64."""
@@ -18,7 +31,6 @@ def rocm_lib_dir(rocm_root):
return libdir
return None
-
def rocm_is_available(rocm_root):
if not rocm_root or not os.path.isdir(rocm_root):
return False
@@ -27,14 +39,17 @@ def rocm_is_available(rocm_root):
return False
return rocm_lib_dir(rocm_root) is not None
-
root = getRoot(config)
-
# AMDGPU ASan tests are only run with the dynamic ASan runtime (-shared-libasan).
if "asan-static-runtime" in root.available_features:
config.unsupported = True
elif root.target_os != "Linux":
config.unsupported = True
+elif walk_config_attr(config, "bits") == "32" or walk_config_attr(
+ config, "target_arch"
+) in ("i386", "i686"):
+ # ROCm libhsa-runtime64.so is 64-bit only (link fails: incompatible with elf32-i386).
+ config.unsupported = True
else:
rocm_root = os.environ.get("ROCM_PATH", "/opt/rocm")
if not rocm_is_available(rocm_root):
>From c69dd17a23b6fa928c2e726fea26e78f670bc110 Mon Sep 17 00:00:00 2001
From: Amit Pandey <pandey.kumaramit2023 at gmail.com>
Date: Thu, 23 Apr 2026 20:53:07 +0530
Subject: [PATCH 10/27] [sanitizer_common] Add GetMetaDataFastLocked for locked
allocator paths
Introduce CombinedAllocator::GetMetaDataFastLocked, mirroring
GetMetaData but using GetBlockBeginFastLocked for secondary (and AMDGPU
device) lookups so callers that already hold the allocator lock do not
take the secondary mutex again.
Use the fast path from LSAN (MetadataFastLocked) and HWASAN where
metadata is read under the same locking discipline as chunk iteration.
---
compiler-rt/lib/hwasan/hwasan_allocator.cpp | 15 ++++++++-------
compiler-rt/lib/lsan/lsan_allocator.cpp | 10 ++++++++--
.../sanitizer_allocator_combined.h | 15 +++++++++++++++
3 files changed, 31 insertions(+), 9 deletions(-)
diff --git a/compiler-rt/lib/hwasan/hwasan_allocator.cpp b/compiler-rt/lib/hwasan/hwasan_allocator.cpp
index 80cc8e1b69a23..7a03caac80778 100644
--- a/compiler-rt/lib/hwasan/hwasan_allocator.cpp
+++ b/compiler-rt/lib/hwasan/hwasan_allocator.cpp
@@ -571,8 +571,9 @@ uptr PointsIntoChunk(void *p) {
reinterpret_cast<uptr>(__hwasan::allocator.GetBlockBeginFastLocked(p));
if (!chunk)
return 0;
- __hwasan::Metadata *metadata = reinterpret_cast<__hwasan::Metadata *>(
- __hwasan::allocator.GetMetaData(reinterpret_cast<void *>(chunk)));
+ __hwasan::Metadata* metadata = reinterpret_cast<__hwasan::Metadata*>(
+ __hwasan::allocator.GetMetaDataFastLocked(
+ reinterpret_cast<void*>(chunk)));
if (!metadata || !metadata->IsAllocated())
return 0;
if (addr < chunk + metadata->GetRequestedSize())
@@ -588,8 +589,8 @@ uptr GetUserBegin(uptr chunk) {
reinterpret_cast<void *>(chunk));
if (!block)
return 0;
- __hwasan::Metadata *metadata = reinterpret_cast<__hwasan::Metadata *>(
- __hwasan::allocator.GetMetaData(block));
+ __hwasan::Metadata* metadata = reinterpret_cast<__hwasan::Metadata*>(
+ __hwasan::allocator.GetMetaDataFastLocked(block));
if (!metadata || !metadata->IsAllocated())
return 0;
@@ -605,9 +606,9 @@ uptr GetUserAddr(uptr chunk) {
LsanMetadata::LsanMetadata(uptr chunk) {
CHECK_EQ(UntagAddr(chunk), chunk);
- metadata_ =
- chunk ? __hwasan::allocator.GetMetaData(reinterpret_cast<void *>(chunk))
- : nullptr;
+ metadata_ = chunk ? __hwasan::allocator.GetMetaDataFastLocked(
+ reinterpret_cast<void*>(chunk))
+ : nullptr;
}
bool LsanMetadata::allocated() const {
diff --git a/compiler-rt/lib/lsan/lsan_allocator.cpp b/compiler-rt/lib/lsan/lsan_allocator.cpp
index a436d9c07ac6c..e690b86013a5d 100644
--- a/compiler-rt/lib/lsan/lsan_allocator.cpp
+++ b/compiler-rt/lib/lsan/lsan_allocator.cpp
@@ -60,6 +60,12 @@ static ChunkMetadata *Metadata(const void *p) {
return reinterpret_cast<ChunkMetadata *>(allocator.GetMetaData(p));
}
+// Same as Metadata, but must be called with the allocator locked (via
+// ForceLock). Avoids re-acquiring the secondary allocator mutex.
+static ChunkMetadata* MetadataFastLocked(const void* p) {
+ return reinterpret_cast<ChunkMetadata*>(allocator.GetMetaDataFastLocked(p));
+}
+
static void RegisterAllocation(const StackTrace &stack, void *p, uptr size) {
if (!p) return;
ChunkMetadata *m = Metadata(p);
@@ -287,7 +293,7 @@ uptr PointsIntoChunk(void* p) {
// LargeMmapAllocator considers pointers to the meta-region of a chunk to be
// valid, but we don't want that.
if (addr < chunk) return 0;
- ChunkMetadata *m = Metadata(reinterpret_cast<void *>(chunk));
+ ChunkMetadata* m = MetadataFastLocked(reinterpret_cast<void*>(chunk));
CHECK(m);
if (!m->allocated)
return 0;
@@ -307,7 +313,7 @@ uptr GetUserAddr(uptr chunk) {
}
LsanMetadata::LsanMetadata(uptr chunk) {
- metadata_ = Metadata(reinterpret_cast<void *>(chunk));
+ metadata_ = MetadataFastLocked(reinterpret_cast<void*>(chunk));
CHECK(metadata_);
}
diff --git a/compiler-rt/lib/sanitizer_common/sanitizer_allocator_combined.h b/compiler-rt/lib/sanitizer_common/sanitizer_allocator_combined.h
index b3ec606f61cb1..1397959b66679 100644
--- a/compiler-rt/lib/sanitizer_common/sanitizer_allocator_combined.h
+++ b/compiler-rt/lib/sanitizer_common/sanitizer_allocator_combined.h
@@ -160,6 +160,21 @@ class CombinedAllocator {
return nullptr;
}
+ // Same as GetMetaData, but must be called with the allocator locked
+ // (via ForceLock). Uses GetBlockBeginFastLocked for secondary/device checks
+ // to avoid re-acquiring the mutex already held by the caller.
+ void* GetMetaDataFastLocked(const void* p) {
+ if (primary_.PointerIsMine(p))
+ return primary_.GetMetaData(p);
+ if (secondary_.GetBlockBeginFastLocked(p))
+ return secondary_.GetMetaData(p);
+#if SANITIZER_AMDGPU
+ if (device_.GetBlockBeginFastLocked(p))
+ return device_.GetMetaData(p);
+#endif
+ return nullptr;
+ }
+
void *GetBlockBegin(const void *p) {
if (primary_.PointerIsMine(p))
return primary_.GetBlockBegin(p);
>From a65c271a383cb8a0463048695d99dd676c17ea84 Mon Sep 17 00:00:00 2001
From: Amit Pandey <pandey.kumaramit2023 at gmail.com>
Date: Thu, 23 Apr 2026 21:08:32 +0530
Subject: [PATCH 11/27] Fix python code formatting issues for lit.local.cfg.py
---
compiler-rt/test/asan/TestCases/AMDGPU/lit.local.cfg.py | 5 +++++
1 file changed, 5 insertions(+)
diff --git a/compiler-rt/test/asan/TestCases/AMDGPU/lit.local.cfg.py b/compiler-rt/test/asan/TestCases/AMDGPU/lit.local.cfg.py
index 7807c20f3a4cd..a544e8d3d0b9d 100644
--- a/compiler-rt/test/asan/TestCases/AMDGPU/lit.local.cfg.py
+++ b/compiler-rt/test/asan/TestCases/AMDGPU/lit.local.cfg.py
@@ -6,11 +6,13 @@
import glob
import os
+
def getRoot(config):
if not config.parent:
return config
return getRoot(config.parent)
+
def walk_config_attr(cfg, name):
"""Return the first defined attribute `name` walking cfg -> parents."""
while cfg is not None:
@@ -21,6 +23,7 @@ def walk_config_attr(cfg, name):
cfg = cfg.parent
return None
+
def rocm_lib_dir(rocm_root):
"""Return lib or lib64 under rocm_root that provides libhsa-runtime64."""
for libname in ("lib", "lib64"):
@@ -31,6 +34,7 @@ def rocm_lib_dir(rocm_root):
return libdir
return None
+
def rocm_is_available(rocm_root):
if not rocm_root or not os.path.isdir(rocm_root):
return False
@@ -39,6 +43,7 @@ def rocm_is_available(rocm_root):
return False
return rocm_lib_dir(rocm_root) is not None
+
root = getRoot(config)
# AMDGPU ASan tests are only run with the dynamic ASan runtime (-shared-libasan).
if "asan-static-runtime" in root.available_features:
>From 601e2e3338fc11e73103da88efac908275c6f1b3 Mon Sep 17 00:00:00 2001
From: Amit Pandey <pandey.kumaramit2023 at gmail.com>
Date: Fri, 24 Apr 2026 14:12:10 +0530
Subject: [PATCH 12/27] [compiler-rt][AMDASAN-cmake] Add Find HSA/AMDComgr
CMake modules
Add `FindHSA.cmake` and `FindAMDComgr.cmake` modules invoked via
find_package(... MODULE).
Discovery honors HSA_ROOT / AMDComgr_ROOT, ROCM_PATH (and ROCM_PATH
env), legacy SANITIZER_*_INCLUDE_PATH hints (including leaf include
dirs), and default legacy search path `/opt/rocm/include`, then sets
HSA_INCLUDE_DIR and AMDComgr_INCLUDE_DIR for include_directories.
---
compiler-rt/CMakeLists.txt | 19 +-----
compiler-rt/cmake/Modules/FindAMDComgr.cmake | 68 ++++++++++++++++++++
compiler-rt/cmake/Modules/FindHSA.cmake | 59 +++++++++++++++++
3 files changed, 130 insertions(+), 16 deletions(-)
create mode 100644 compiler-rt/cmake/Modules/FindAMDComgr.cmake
create mode 100644 compiler-rt/cmake/Modules/FindHSA.cmake
diff --git a/compiler-rt/CMakeLists.txt b/compiler-rt/CMakeLists.txt
index 8118255041a3d..d6612c6730d9c 100644
--- a/compiler-rt/CMakeLists.txt
+++ b/compiler-rt/CMakeLists.txt
@@ -551,22 +551,9 @@ endif()
if(SANITIZER_AMDGPU)
list(APPEND SANITIZER_COMMON_CFLAGS -DSANITIZER_AMDGPU=1)
- message(STATUS "Looking 'hsa' and 'amd_comgr' header")
- find_path(HSA_INCLUDE NAMES hsa.h HINTS ${SANITIZER_HSA_INCLUDE_PATH} /opt/rocm/include PATH_SUFFIXES hsa)
- if(NOT HSA_INCLUDE)
- message(FATAL_ERROR "Required header 'hsa.h' not found in path ${HSA_INCLUDE}. Aborting SANITIZER_AMDGPU build")
- endif()
- message(STATUS "Found 'hsa.h' in ${HSA_INCLUDE}")
- include_directories(${HSA_INCLUDE})
- find_path(COMgr_INCLUDE NAMES amd_comgr.h.in HINTS ${SANITIZER_COMGR_INCLUDE_PATH} PATH_SUFFIXES amd_comgr)
- if(NOT COMgr_INCLUDE)
- find_path(COMgr_INCLUDE NAMES amd_comgr.h HINTS /opt/rocm/include PATH_SUFFIXES amd_comgr)
- if(NOT COMgr_INCLUDE)
- message(FATAL_ERROR "Required header 'amd_comgr.h/amd_comgr.h.in' not found in path ${COMgr_INCLUDE}. Aborting SANITIZER_AMDGPU build")
- endif()
- endif()
- message(STATUS "Found 'amd_comgr.h.in/amd_comgr.h' in ${COMgr_INCLUDE}")
- include_directories(${COMgr_INCLUDE})
+ find_package(HSA REQUIRED MODULE)
+ find_package(AMDComgr REQUIRED MODULE)
+ include_directories(${HSA_INCLUDE_DIR} ${AMDComgr_INCLUDE_DIR})
endif()
if(LLVM_ENABLE_MODULES)
diff --git a/compiler-rt/cmake/Modules/FindAMDComgr.cmake b/compiler-rt/cmake/Modules/FindAMDComgr.cmake
new file mode 100644
index 0000000000000..2303cbd5d89f3
--- /dev/null
+++ b/compiler-rt/cmake/Modules/FindAMDComgr.cmake
@@ -0,0 +1,68 @@
+# Find AMD COMGR headers (amd_comgr.h or amd_comgr.h.in).
+#
+# This module is used by compiler-rt when SANITIZER_AMDGPU is enabled.
+#
+# The following variables may be set by the user:
+# AMDComgr_ROOT ROCm / amd_comgr install prefix
+# ROCM_PATH Typical ROCm install (expects include/amd_comgr/...)
+# SANITIZER_COMGR_INCLUDE_PATH
+# Legacy hint: base directory searched with
+# PATH_SUFFIXES amd_comgr (same as original CMake).
+#
+# The following cache variables may be set by this module:
+# AMDComgr_INCLUDE_DIR Directory containing amd_comgr.h(.in)
+#
+# This module defines:
+# AMDComgr_FOUND TRUE if a supported amd_comgr header was found
+#
+# Example:
+# find_package(AMDComgr REQUIRED)
+
+include(FindPackageHandleStandardArgs)
+
+set(_amdcomgr_search_paths "")
+foreach(_root IN ITEMS "${AMDComgr_ROOT}" "${ROCM_PATH}" "$ENV{ROCM_PATH}")
+ if(_root)
+ list(APPEND _amdcomgr_search_paths "${_root}/include")
+ endif()
+endforeach()
+if(SANITIZER_COMGR_INCLUDE_PATH)
+ list(APPEND _amdcomgr_search_paths "${SANITIZER_COMGR_INCLUDE_PATH}")
+endif()
+# Default Legacy ROCm include path
+list(APPEND _amdcomgr_search_paths "/opt/rocm/include")
+
+find_path(
+ AMDComgr_INCLUDE_DIR
+ NAMES amd_comgr.h.in
+ HINTS ${_amdcomgr_search_paths}
+ PATH_SUFFIXES amd_comgr
+)
+
+if(NOT AMDComgr_INCLUDE_DIR)
+ find_path(
+ AMDComgr_INCLUDE_DIR
+ NAMES amd_comgr.h
+ HINTS ${_amdcomgr_search_paths}
+ PATH_SUFFIXES amd_comgr
+ )
+endif()
+
+# If SANITIZER_COMGR_INCLUDE_PATH points at the leaf `amd_comgr` directory,
+# retry without PATH_SUFFIXES.
+if(NOT AMDComgr_INCLUDE_DIR AND SANITIZER_COMGR_INCLUDE_PATH)
+ find_path(
+ AMDComgr_INCLUDE_DIR
+ NAMES amd_comgr.h.in amd_comgr.h
+ HINTS "${SANITIZER_COMGR_INCLUDE_PATH}"
+ NO_DEFAULT_PATH
+ NO_CMAKE_PATH
+ NO_CMAKE_ENVIRONMENT_PATH
+ NO_SYSTEM_ENVIRONMENT_PATH
+ NO_CMAKE_SYSTEM_PATH
+ )
+endif()
+
+find_package_handle_standard_args(AMDComgr REQUIRED_VARS AMDComgr_INCLUDE_DIR)
+
+mark_as_advanced(AMDComgr_INCLUDE_DIR)
diff --git a/compiler-rt/cmake/Modules/FindHSA.cmake b/compiler-rt/cmake/Modules/FindHSA.cmake
new file mode 100644
index 0000000000000..bd0f57a31c47f
--- /dev/null
+++ b/compiler-rt/cmake/Modules/FindHSA.cmake
@@ -0,0 +1,59 @@
+# Find ROCm HSA runtime headers (hsa.h/hsa_ext_amd.h).
+#
+# This module is used by compiler-rt when SANITIZER_AMDGPU is enabled.
+#
+# The following variables may be set by the user:
+# HSA_ROOT ROCm / HSA install prefix (expects include/hsa/hsa.h)
+# ROCM_PATH Same as typical ROCm install (expects include/hsa/hsa.h)
+# SANITIZER_HSA_INCLUDE_PATH
+# Legacy hint: directory searched like the original
+# find_path(HINTS ... PATH_SUFFIXES hsa) entry.
+#
+# The following cache variables may be set by this module:
+# HSA_INCLUDE_DIR Directory containing hsa.h (typically .../include/hsa)
+#
+# This module defines:
+# HSA_FOUND TRUE if hsa.h was found
+#
+# Example:
+# find_package(HSA REQUIRED)
+
+include(FindPackageHandleStandardArgs)
+
+set(_hsa_search_paths "")
+foreach(_root IN ITEMS "${HSA_ROOT}" "${ROCM_PATH}" "$ENV{ROCM_PATH}")
+ if(_root)
+ list(APPEND _hsa_search_paths "${_root}/include")
+ endif()
+endforeach()
+if(SANITIZER_HSA_INCLUDE_PATH)
+ list(APPEND _hsa_search_paths "${SANITIZER_HSA_INCLUDE_PATH}")
+endif()
+# Default Legacy ROCm include path
+list(APPEND _hsa_search_paths "/opt/rocm/include")
+
+find_path(
+ HSA_INCLUDE_DIR
+ NAMES hsa.h
+ HINTS ${_hsa_search_paths}
+ PATH_SUFFIXES hsa
+)
+
+# If SANITIZER_HSA_INCLUDE_PATH points at the leaf `hsa` directory, retry
+# without PATH_SUFFIXES.
+if(NOT HSA_INCLUDE_DIR AND SANITIZER_HSA_INCLUDE_PATH)
+ find_path(
+ HSA_INCLUDE_DIR
+ NAMES hsa.h
+ HINTS "${SANITIZER_HSA_INCLUDE_PATH}"
+ NO_DEFAULT_PATH
+ NO_CMAKE_PATH
+ NO_CMAKE_ENVIRONMENT_PATH
+ NO_SYSTEM_ENVIRONMENT_PATH
+ NO_CMAKE_SYSTEM_PATH
+ )
+endif()
+
+find_package_handle_standard_args(HSA REQUIRED_VARS HSA_INCLUDE_DIR)
+
+mark_as_advanced(HSA_INCLUDE_DIR)
>From e5347e9dbe192055d634fafddbdfb1bb785cd1ce Mon Sep 17 00:00:00 2001
From: Amit Pandey <pandey.kumaramit2023 at gmail.com>
Date: Fri, 24 Apr 2026 16:44:44 +0530
Subject: [PATCH 13/27] [compiler-rt][AMDGPU-Cmake] Fold ROCm HSA/COMGr find
logic into CompilerRTAMDGPUUtils.
Replace FindHSA.cmake and FindAMDComgr.cmake with a single
CompilerRTAMDGPUUtils.cmake module that defines
compiler_rt_find_amdgpu_sanitizer_dependencies().
---
compiler-rt/CMakeLists.txt | 4 +-
.../cmake/Modules/CompilerRTAMDGPUUtils.cmake | 101 ++++++++++++++++++
compiler-rt/cmake/Modules/FindAMDComgr.cmake | 68 ------------
compiler-rt/cmake/Modules/FindHSA.cmake | 59 ----------
4 files changed, 103 insertions(+), 129 deletions(-)
create mode 100644 compiler-rt/cmake/Modules/CompilerRTAMDGPUUtils.cmake
delete mode 100644 compiler-rt/cmake/Modules/FindAMDComgr.cmake
delete mode 100644 compiler-rt/cmake/Modules/FindHSA.cmake
diff --git a/compiler-rt/CMakeLists.txt b/compiler-rt/CMakeLists.txt
index d6612c6730d9c..a174d30dc1dac 100644
--- a/compiler-rt/CMakeLists.txt
+++ b/compiler-rt/CMakeLists.txt
@@ -551,8 +551,8 @@ endif()
if(SANITIZER_AMDGPU)
list(APPEND SANITIZER_COMMON_CFLAGS -DSANITIZER_AMDGPU=1)
- find_package(HSA REQUIRED MODULE)
- find_package(AMDComgr REQUIRED MODULE)
+ include(CompilerRTAMDGPUUtils)
+ compiler_rt_find_amdgpu_sanitizer_dependencies()
include_directories(${HSA_INCLUDE_DIR} ${AMDComgr_INCLUDE_DIR})
endif()
diff --git a/compiler-rt/cmake/Modules/CompilerRTAMDGPUUtils.cmake b/compiler-rt/cmake/Modules/CompilerRTAMDGPUUtils.cmake
new file mode 100644
index 0000000000000..e9d3af7ca0518
--- /dev/null
+++ b/compiler-rt/cmake/Modules/CompilerRTAMDGPUUtils.cmake
@@ -0,0 +1,101 @@
+# ROCm header discovery for compiler-rt when SANITIZER_AMDGPU is enabled.
+#
+# Include this module and call:
+# compiler_rt_find_amdgpu_sanitizer_dependencies()
+#
+# User-settable hints (optional):
+# HSA_ROOT / AMDComgr_ROOT Install prefix (expects include/hsa/... or include/amd_comgr/...)
+# ROCM_PATH Typical ROCm layout; also honors $ENV{ROCM_PATH}
+# SANITIZER_HSA_INCLUDE_PATH Legacy: same search as original find_path(... PATH_SUFFIXES hsa)
+# SANITIZER_COMGR_INCLUDE_PATH
+# Legacy: same search as original find_path(... PATH_SUFFIXES amd_comgr)
+#
+# Output (same as the former FindHSA / FindAMDComgr modules):
+# HSA_INCLUDE_DIR, HSA_FOUND
+# AMDComgr_INCLUDE_DIR, AMDComgr_FOUND
+#
+# This call is REQUIRED-style: missing headers trigger a fatal error from
+# find_package_handle_standard_args.
+
+include(FindPackageHandleStandardArgs)
+
+macro(compiler_rt_find_amdgpu_sanitizer_dependencies)
+ # --- HSA (hsa.h) ---
+ set(_hsa_search_paths "")
+ foreach(_root IN ITEMS "${HSA_ROOT}" "${ROCM_PATH}" "$ENV{ROCM_PATH}")
+ if(_root)
+ list(APPEND _hsa_search_paths "${_root}/include")
+ endif()
+ endforeach()
+ if(SANITIZER_HSA_INCLUDE_PATH)
+ list(APPEND _hsa_search_paths "${SANITIZER_HSA_INCLUDE_PATH}")
+ endif()
+ list(APPEND _hsa_search_paths "/opt/rocm/include")
+
+ find_path(
+ HSA_INCLUDE_DIR
+ NAMES hsa.h
+ HINTS ${_hsa_search_paths}
+ PATH_SUFFIXES hsa
+ )
+
+ if(NOT HSA_INCLUDE_DIR AND SANITIZER_HSA_INCLUDE_PATH)
+ find_path(
+ HSA_INCLUDE_DIR
+ NAMES hsa.h
+ HINTS "${SANITIZER_HSA_INCLUDE_PATH}"
+ NO_DEFAULT_PATH
+ NO_CMAKE_PATH
+ NO_CMAKE_ENVIRONMENT_PATH
+ NO_SYSTEM_ENVIRONMENT_PATH
+ NO_CMAKE_SYSTEM_PATH
+ )
+ endif()
+
+ find_package_handle_standard_args(HSA REQUIRED_VARS HSA_INCLUDE_DIR)
+ mark_as_advanced(HSA_INCLUDE_DIR)
+
+ # --- AMD COMGR (amd_comgr.h.in / amd_comgr.h) ---
+ set(_amdcomgr_search_paths "")
+ foreach(_root IN ITEMS "${AMDComgr_ROOT}" "${ROCM_PATH}" "$ENV{ROCM_PATH}")
+ if(_root)
+ list(APPEND _amdcomgr_search_paths "${_root}/include")
+ endif()
+ endforeach()
+ if(SANITIZER_COMGR_INCLUDE_PATH)
+ list(APPEND _amdcomgr_search_paths "${SANITIZER_COMGR_INCLUDE_PATH}")
+ endif()
+ list(APPEND _amdcomgr_search_paths "/opt/rocm/include")
+
+ find_path(
+ AMDComgr_INCLUDE_DIR
+ NAMES amd_comgr.h.in
+ HINTS ${_amdcomgr_search_paths}
+ PATH_SUFFIXES amd_comgr
+ )
+
+ if(NOT AMDComgr_INCLUDE_DIR)
+ find_path(
+ AMDComgr_INCLUDE_DIR
+ NAMES amd_comgr.h
+ HINTS ${_amdcomgr_search_paths}
+ PATH_SUFFIXES amd_comgr
+ )
+ endif()
+
+ if(NOT AMDComgr_INCLUDE_DIR AND SANITIZER_COMGR_INCLUDE_PATH)
+ find_path(
+ AMDComgr_INCLUDE_DIR
+ NAMES amd_comgr.h.in amd_comgr.h
+ HINTS "${SANITIZER_COMGR_INCLUDE_PATH}"
+ NO_DEFAULT_PATH
+ NO_CMAKE_PATH
+ NO_CMAKE_ENVIRONMENT_PATH
+ NO_SYSTEM_ENVIRONMENT_PATH
+ NO_CMAKE_SYSTEM_PATH
+ )
+ endif()
+
+ find_package_handle_standard_args(AMDComgr REQUIRED_VARS AMDComgr_INCLUDE_DIR)
+ mark_as_advanced(AMDComgr_INCLUDE_DIR)
+endmacro()
diff --git a/compiler-rt/cmake/Modules/FindAMDComgr.cmake b/compiler-rt/cmake/Modules/FindAMDComgr.cmake
deleted file mode 100644
index 2303cbd5d89f3..0000000000000
--- a/compiler-rt/cmake/Modules/FindAMDComgr.cmake
+++ /dev/null
@@ -1,68 +0,0 @@
-# Find AMD COMGR headers (amd_comgr.h or amd_comgr.h.in).
-#
-# This module is used by compiler-rt when SANITIZER_AMDGPU is enabled.
-#
-# The following variables may be set by the user:
-# AMDComgr_ROOT ROCm / amd_comgr install prefix
-# ROCM_PATH Typical ROCm install (expects include/amd_comgr/...)
-# SANITIZER_COMGR_INCLUDE_PATH
-# Legacy hint: base directory searched with
-# PATH_SUFFIXES amd_comgr (same as original CMake).
-#
-# The following cache variables may be set by this module:
-# AMDComgr_INCLUDE_DIR Directory containing amd_comgr.h(.in)
-#
-# This module defines:
-# AMDComgr_FOUND TRUE if a supported amd_comgr header was found
-#
-# Example:
-# find_package(AMDComgr REQUIRED)
-
-include(FindPackageHandleStandardArgs)
-
-set(_amdcomgr_search_paths "")
-foreach(_root IN ITEMS "${AMDComgr_ROOT}" "${ROCM_PATH}" "$ENV{ROCM_PATH}")
- if(_root)
- list(APPEND _amdcomgr_search_paths "${_root}/include")
- endif()
-endforeach()
-if(SANITIZER_COMGR_INCLUDE_PATH)
- list(APPEND _amdcomgr_search_paths "${SANITIZER_COMGR_INCLUDE_PATH}")
-endif()
-# Default Legacy ROCm include path
-list(APPEND _amdcomgr_search_paths "/opt/rocm/include")
-
-find_path(
- AMDComgr_INCLUDE_DIR
- NAMES amd_comgr.h.in
- HINTS ${_amdcomgr_search_paths}
- PATH_SUFFIXES amd_comgr
-)
-
-if(NOT AMDComgr_INCLUDE_DIR)
- find_path(
- AMDComgr_INCLUDE_DIR
- NAMES amd_comgr.h
- HINTS ${_amdcomgr_search_paths}
- PATH_SUFFIXES amd_comgr
- )
-endif()
-
-# If SANITIZER_COMGR_INCLUDE_PATH points at the leaf `amd_comgr` directory,
-# retry without PATH_SUFFIXES.
-if(NOT AMDComgr_INCLUDE_DIR AND SANITIZER_COMGR_INCLUDE_PATH)
- find_path(
- AMDComgr_INCLUDE_DIR
- NAMES amd_comgr.h.in amd_comgr.h
- HINTS "${SANITIZER_COMGR_INCLUDE_PATH}"
- NO_DEFAULT_PATH
- NO_CMAKE_PATH
- NO_CMAKE_ENVIRONMENT_PATH
- NO_SYSTEM_ENVIRONMENT_PATH
- NO_CMAKE_SYSTEM_PATH
- )
-endif()
-
-find_package_handle_standard_args(AMDComgr REQUIRED_VARS AMDComgr_INCLUDE_DIR)
-
-mark_as_advanced(AMDComgr_INCLUDE_DIR)
diff --git a/compiler-rt/cmake/Modules/FindHSA.cmake b/compiler-rt/cmake/Modules/FindHSA.cmake
deleted file mode 100644
index bd0f57a31c47f..0000000000000
--- a/compiler-rt/cmake/Modules/FindHSA.cmake
+++ /dev/null
@@ -1,59 +0,0 @@
-# Find ROCm HSA runtime headers (hsa.h/hsa_ext_amd.h).
-#
-# This module is used by compiler-rt when SANITIZER_AMDGPU is enabled.
-#
-# The following variables may be set by the user:
-# HSA_ROOT ROCm / HSA install prefix (expects include/hsa/hsa.h)
-# ROCM_PATH Same as typical ROCm install (expects include/hsa/hsa.h)
-# SANITIZER_HSA_INCLUDE_PATH
-# Legacy hint: directory searched like the original
-# find_path(HINTS ... PATH_SUFFIXES hsa) entry.
-#
-# The following cache variables may be set by this module:
-# HSA_INCLUDE_DIR Directory containing hsa.h (typically .../include/hsa)
-#
-# This module defines:
-# HSA_FOUND TRUE if hsa.h was found
-#
-# Example:
-# find_package(HSA REQUIRED)
-
-include(FindPackageHandleStandardArgs)
-
-set(_hsa_search_paths "")
-foreach(_root IN ITEMS "${HSA_ROOT}" "${ROCM_PATH}" "$ENV{ROCM_PATH}")
- if(_root)
- list(APPEND _hsa_search_paths "${_root}/include")
- endif()
-endforeach()
-if(SANITIZER_HSA_INCLUDE_PATH)
- list(APPEND _hsa_search_paths "${SANITIZER_HSA_INCLUDE_PATH}")
-endif()
-# Default Legacy ROCm include path
-list(APPEND _hsa_search_paths "/opt/rocm/include")
-
-find_path(
- HSA_INCLUDE_DIR
- NAMES hsa.h
- HINTS ${_hsa_search_paths}
- PATH_SUFFIXES hsa
-)
-
-# If SANITIZER_HSA_INCLUDE_PATH points at the leaf `hsa` directory, retry
-# without PATH_SUFFIXES.
-if(NOT HSA_INCLUDE_DIR AND SANITIZER_HSA_INCLUDE_PATH)
- find_path(
- HSA_INCLUDE_DIR
- NAMES hsa.h
- HINTS "${SANITIZER_HSA_INCLUDE_PATH}"
- NO_DEFAULT_PATH
- NO_CMAKE_PATH
- NO_CMAKE_ENVIRONMENT_PATH
- NO_SYSTEM_ENVIRONMENT_PATH
- NO_CMAKE_SYSTEM_PATH
- )
-endif()
-
-find_package_handle_standard_args(HSA REQUIRED_VARS HSA_INCLUDE_DIR)
-
-mark_as_advanced(HSA_INCLUDE_DIR)
>From d18264194a4347793b1907b022f4c8f8a2c132a4 Mon Sep 17 00:00:00 2001
From: Amit Pandey <pandey.kumaramit2023 at gmail.com>
Date: Mon, 27 Apr 2026 16:31:13 +0530
Subject: [PATCH 14/27] [compiler-rt][HSA-Test] Add
hsa_amd_pointer_info_memory_pool test.
---
.../hsa_amd_pointer_info_memory_pool.cpp | 101 ++++++++++++++++++
1 file changed, 101 insertions(+)
create mode 100644 compiler-rt/test/asan/TestCases/AMDGPU/hsa_amd_pointer_info_memory_pool.cpp
diff --git a/compiler-rt/test/asan/TestCases/AMDGPU/hsa_amd_pointer_info_memory_pool.cpp b/compiler-rt/test/asan/TestCases/AMDGPU/hsa_amd_pointer_info_memory_pool.cpp
new file mode 100644
index 0000000000000..13a05b8b35237
--- /dev/null
+++ b/compiler-rt/test/asan/TestCases/AMDGPU/hsa_amd_pointer_info_memory_pool.cpp
@@ -0,0 +1,101 @@
+// RUN: %clangxx_asan -O0 -isystem %rocm_include %s -o %t -L%rocm_lib -lhsa-runtime64 \
+// RUN: -Wl,-rpath,%rocm_lib -Wl,-rpath,%compiler_rt_libdir
+// RUN: %run %t 2>&1 | FileCheck %s
+//
+// Regression test for the AddressSanitizer hsa_amd_pointer_info interceptor on
+// hsa_amd_memory_pool_allocate pointers: reported sizeInBytes matches the user
+// request (ASan unwraps the page-sized host wrapper from pointer metadata).
+//
+// REQUIRES: sanitizer-amdgpu, linux, stable-runtime, rocm
+// UNSUPPORTED: android
+
+#include <hsa/hsa.h>
+#include <hsa/hsa_ext_amd.h>
+
+#include <stdio.h>
+
+struct PoolSearch {
+ hsa_amd_memory_pool_t pool;
+ bool found;
+};
+
+static hsa_status_t find_alloc_pool(hsa_amd_memory_pool_t pool, void *data) {
+ auto *ps = static_cast<PoolSearch *>(data);
+ bool allow = false;
+ if (hsa_amd_memory_pool_get_info(
+ pool, HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_ALLOWED, &allow) !=
+ HSA_STATUS_SUCCESS ||
+ !allow)
+ return HSA_STATUS_SUCCESS;
+ ps->pool = pool;
+ ps->found = true;
+ return HSA_STATUS_INFO_BREAK;
+}
+
+static hsa_status_t find_agent_with_pool(hsa_agent_t agent, void *data) {
+ (void)agent;
+ auto *ps = static_cast<PoolSearch *>(data);
+ ps->found = false;
+ hsa_status_t st =
+ hsa_amd_agent_iterate_memory_pools(agent, find_alloc_pool, ps);
+ if (st != HSA_STATUS_SUCCESS && st != HSA_STATUS_INFO_BREAK)
+ return st;
+ if (ps->found)
+ return HSA_STATUS_INFO_BREAK;
+ return HSA_STATUS_SUCCESS;
+}
+
+int main() {
+ if (hsa_init() != HSA_STATUS_SUCCESS) {
+ fprintf(stderr, "hsa_init failed\n");
+ return 1;
+ }
+
+ PoolSearch ps = {};
+ ps.pool.handle = 0;
+ ps.found = false;
+
+ hsa_status_t it = hsa_iterate_agents(find_agent_with_pool, &ps);
+ if (it != HSA_STATUS_SUCCESS && it != HSA_STATUS_INFO_BREAK) {
+ fprintf(stderr, "hsa_iterate_agents failed\n");
+ return 1;
+ }
+ if (!ps.found) {
+ fprintf(stderr, "no runtime-alloc HSA memory pool found\n");
+ return 1;
+ }
+
+ void *mem = nullptr;
+ if (hsa_amd_memory_pool_allocate(ps.pool, 64, 0, &mem) !=
+ HSA_STATUS_SUCCESS ||
+ !mem) {
+ fprintf(stderr, "hsa_amd_memory_pool_allocate failed\n");
+ return 1;
+ }
+
+ hsa_amd_pointer_info_t info = {};
+ info.size = sizeof(hsa_amd_pointer_info_t);
+
+ if (hsa_amd_pointer_info(mem, &info, nullptr, nullptr, nullptr) !=
+ HSA_STATUS_SUCCESS) {
+ fprintf(stderr, "hsa_amd_pointer_info failed\n");
+ return 1;
+ }
+
+ printf("pointer_info_pool type: %d\n", info.type);
+ printf("pointer_info_pool sizeInBytes: %zu\n", info.sizeInBytes);
+ printf("pointer_info_pool begin: %p\n", info.agentBaseAddress);
+ printf("pointer_info_pool end: %p\n",
+ (void *)((uintptr_t)info.agentBaseAddress + info.sizeInBytes));
+
+ if (hsa_amd_memory_pool_free(mem) != HSA_STATUS_SUCCESS) {
+ fprintf(stderr, "hsa_amd_memory_pool_free failed\n");
+ return 1;
+ }
+ return 0;
+}
+
+// CHECK: pointer_info_pool type: 1
+// CHECK-NEXT: pointer_info_pool sizeInBytes: 64
+// CHECK-NEXT: pointer_info_pool begin: 0x{{[0-9a-f]+}}
+// CHECK-NEXT: pointer_info_pool end: 0x{{[0-9a-f]+}}
>From c17efb3438ea89337f25c0e3dc5bb09fbb2803eb Mon Sep 17 00:00:00 2001
From: Amit Pandey <pandey.kumaramit2023 at gmail.com>
Date: Mon, 27 Apr 2026 16:57:11 +0530
Subject: [PATCH 15/27] [compiler-rt][Cmake] Rename macro.
Rename macro `compiler_rt_find_amdgpu_sanitizer_dependencies` to
`compiler_rt_find_amdgpu_runtime_headers`.
---
compiler-rt/CMakeLists.txt | 2 +-
compiler-rt/cmake/Modules/CompilerRTAMDGPUUtils.cmake | 2 +-
2 files changed, 2 insertions(+), 2 deletions(-)
diff --git a/compiler-rt/CMakeLists.txt b/compiler-rt/CMakeLists.txt
index a174d30dc1dac..340fd194a1390 100644
--- a/compiler-rt/CMakeLists.txt
+++ b/compiler-rt/CMakeLists.txt
@@ -552,7 +552,7 @@ endif()
if(SANITIZER_AMDGPU)
list(APPEND SANITIZER_COMMON_CFLAGS -DSANITIZER_AMDGPU=1)
include(CompilerRTAMDGPUUtils)
- compiler_rt_find_amdgpu_sanitizer_dependencies()
+ compiler_rt_find_amdgpu_runtime_headers()
include_directories(${HSA_INCLUDE_DIR} ${AMDComgr_INCLUDE_DIR})
endif()
diff --git a/compiler-rt/cmake/Modules/CompilerRTAMDGPUUtils.cmake b/compiler-rt/cmake/Modules/CompilerRTAMDGPUUtils.cmake
index e9d3af7ca0518..0ce46b220d565 100644
--- a/compiler-rt/cmake/Modules/CompilerRTAMDGPUUtils.cmake
+++ b/compiler-rt/cmake/Modules/CompilerRTAMDGPUUtils.cmake
@@ -19,7 +19,7 @@
include(FindPackageHandleStandardArgs)
-macro(compiler_rt_find_amdgpu_sanitizer_dependencies)
+macro(compiler_rt_find_amdgpu_runtime_headers)
# --- HSA (hsa.h) ---
set(_hsa_search_paths "")
foreach(_root IN ITEMS "${HSA_ROOT}" "${ROCM_PATH}" "$ENV{ROCM_PATH}")
>From fcf2f30fa36b22448d347ce1cd4469e20f7c3934 Mon Sep 17 00:00:00 2001
From: Amit Pandey <pandey.kumaramit2023 at gmail.com>
Date: Tue, 28 Apr 2026 09:48:36 +0530
Subject: [PATCH 16/27] [compiler-rt][asan] Force RTLD_GLOBAL for
HSA/HIP/OpenCL runtime dlopen.
- When ASan intercepts dlopen of libamdhip64.so, libhsa-runtime64.so,
or libamdocl64.so, add RTLD_GLOBAL if it was omitted, so symbols are
visible as required by the AMDGPU device sanitizer runtime.
- Add PatchHsaRuntimeDlopenFlag in sanitizer_common/sanitizer_linux
and call it from the dlopen interceptor.
- Update CompilerRTAMDGPUUtils.cmake comments for the AMDGPU runtime
header search.
---
.../cmake/Modules/CompilerRTAMDGPUUtils.cmake | 25 ++++++++++---------
compiler-rt/lib/asan/asan_interceptors.cpp | 1 +
.../lib/sanitizer_common/sanitizer_common.h | 6 +++++
.../lib/sanitizer_common/sanitizer_linux.cpp | 18 +++++++++++++
4 files changed, 38 insertions(+), 12 deletions(-)
diff --git a/compiler-rt/cmake/Modules/CompilerRTAMDGPUUtils.cmake b/compiler-rt/cmake/Modules/CompilerRTAMDGPUUtils.cmake
index 0ce46b220d565..58a81ffdffd68 100644
--- a/compiler-rt/cmake/Modules/CompilerRTAMDGPUUtils.cmake
+++ b/compiler-rt/cmake/Modules/CompilerRTAMDGPUUtils.cmake
@@ -1,21 +1,20 @@
-# ROCm header discovery for compiler-rt when SANITIZER_AMDGPU is enabled.
+# ADMGPU runtime headers discovery for compiler-rt when SANITIZER_AMDGPU is enabled.
#
-# Include this module and call:
-# compiler_rt_find_amdgpu_sanitizer_dependencies()
+# Usage: Include this module and call
+# `compiler_rt_find_amdgpu_runtime_headers()`
#
# User-settable hints (optional):
-# HSA_ROOT / AMDComgr_ROOT Install prefix (expects include/hsa/... or include/amd_comgr/...)
-# ROCM_PATH Typical ROCm layout; also honors $ENV{ROCM_PATH}
-# SANITIZER_HSA_INCLUDE_PATH Legacy: same search as original find_path(... PATH_SUFFIXES hsa)
-# SANITIZER_COMGR_INCLUDE_PATH
-# Legacy: same search as original find_path(... PATH_SUFFIXES amd_comgr)
-#
-# Output (same as the former FindHSA / FindAMDComgr modules):
+# HSA_ROOT / AMDComgr_ROOT Install prefix (expects include/hsa/... or include/amd_comgr/...)
+# ROCM_PATH Typical ROCm layout; also honors $ENV{ROCM_PATH}
+# SANITIZER_HSA_INCLUDE_PATH Custom HSA Include Path from source tree
+# SANITIZER_COMGR_INCLUDE_PATH Custom COMGR Include Path from source tree
+#
+# Output CMake variables:
# HSA_INCLUDE_DIR, HSA_FOUND
# AMDComgr_INCLUDE_DIR, AMDComgr_FOUND
#
-# This call is REQUIRED-style: missing headers trigger a fatal error from
-# find_package_handle_standard_args.
+# This call is REQUIRED-style: missing headers triggers a fatal error from
+# `find_package_handle_standard_args`.
include(FindPackageHandleStandardArgs)
@@ -30,6 +29,7 @@ macro(compiler_rt_find_amdgpu_runtime_headers)
if(SANITIZER_HSA_INCLUDE_PATH)
list(APPEND _hsa_search_paths "${SANITIZER_HSA_INCLUDE_PATH}")
endif()
+ # Default Search Fallback: ROCm include path.
list(APPEND _hsa_search_paths "/opt/rocm/include")
find_path(
@@ -65,6 +65,7 @@ macro(compiler_rt_find_amdgpu_runtime_headers)
if(SANITIZER_COMGR_INCLUDE_PATH)
list(APPEND _amdcomgr_search_paths "${SANITIZER_COMGR_INCLUDE_PATH}")
endif()
+ # Default Search Fallback: ROCm include path.
list(APPEND _amdcomgr_search_paths "/opt/rocm/include")
find_path(
diff --git a/compiler-rt/lib/asan/asan_interceptors.cpp b/compiler-rt/lib/asan/asan_interceptors.cpp
index 924d36afc925b..16f22bbcc2bc5 100644
--- a/compiler-rt/lib/asan/asan_interceptors.cpp
+++ b/compiler-rt/lib/asan/asan_interceptors.cpp
@@ -164,6 +164,7 @@ DECLARE_REAL_AND_INTERCEPTOR(void, free, void*)
if (flags()->strict_init_order) \
StopInitOrderChecking(); \
CheckNoDeepBind(filename, flag); \
+ PatchHsaRuntimeDlopenFlag(filename, flag); \
REAL(dlopen)(filename, flag); \
})
# define COMMON_INTERCEPTOR_ON_EXIT(ctx) OnExit()
diff --git a/compiler-rt/lib/sanitizer_common/sanitizer_common.h b/compiler-rt/lib/sanitizer_common/sanitizer_common.h
index 2aacec1cd8994..d556d2cac2be4 100644
--- a/compiler-rt/lib/sanitizer_common/sanitizer_common.h
+++ b/compiler-rt/lib/sanitizer_common/sanitizer_common.h
@@ -1094,6 +1094,12 @@ const s32 kReleaseToOSIntervalNever = -1;
void CheckNoDeepBind(const char *filename, int flag);
+#if SANITIZER_AMDGPU
+void PatchHsaRuntimeDlopenFlag(const char* filename, int& flag);
+#else
+inline void PatchHsaRuntimeDlopenFlag(const char* filename, int& flag) {}
+#endif
+
// Returns the requested amount of random data (up to 256 bytes) that can then
// be used to seed a PRNG. Defaults to blocking like the underlying syscall.
bool GetRandom(void *buffer, uptr length, bool blocking = true);
diff --git a/compiler-rt/lib/sanitizer_common/sanitizer_linux.cpp b/compiler-rt/lib/sanitizer_common/sanitizer_linux.cpp
index 1bdbccc0a03ec..7378c1ad6c7dd 100644
--- a/compiler-rt/lib/sanitizer_common/sanitizer_linux.cpp
+++ b/compiler-rt/lib/sanitizer_common/sanitizer_linux.cpp
@@ -2901,6 +2901,24 @@ void CheckNoDeepBind(const char *filename, int flag) {
# endif
}
+# if SANITIZER_AMDGPU
+void PatchHsaRuntimeDlopenFlag(const char* filename, int& flag) {
+ if (filename &&
+ (internal_strstr(filename, "libamdhip64.so") ||
+ internal_strstr(filename, "libhsa-runtime64.so") ||
+ internal_strstr(filename, "libamdocl64.so")) &&
+ !(flag & RTLD_GLOBAL)) {
+ flag |= RTLD_GLOBAL;
+ if (Verbosity() >= 2) {
+ Printf(
+ "RTLD_GLOBAL flag on dlopen call forced on for %s due to AMDGPU "
+ "device sanitizer runtime requirements.\n",
+ filename);
+ }
+ }
+}
+# endif
+
uptr FindAvailableMemoryRange(uptr size, uptr alignment, uptr left_padding,
uptr *largest_gap_found,
uptr *max_occupied_addr) {
>From 46e021c6f8e7a9a7c68f7325426896b35560f3bf Mon Sep 17 00:00:00 2001
From: Amit Pandey <pandey.kumaramit2023 at gmail.com>
Date: Mon, 4 May 2026 13:51:15 +0530
Subject: [PATCH 17/27] [compiler-rt] Only use hidden sanitizer interface on
true AMDGPU TUs
SANITIZER_AMDGPU=1 is also used for host (e.g. x86_64) runtimes that
include the HSA stack. In that configuration, mapping
SANITIZER_INTERFACE_ATTRIBUTE to visibility("hidden") clashes with
replaceable operator new / delete, which must match the
default-visibility ABI.
Require defined(__AMDGPU__) in addition to SANITIZER_AMDGPU for the
hidden/weak branch, so only AMDGPU-target translation units get the
stricter attributes; NVPTX behavior is unchanged.
---
compiler-rt/lib/sanitizer_common/sanitizer_internal_defs.h | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/compiler-rt/lib/sanitizer_common/sanitizer_internal_defs.h b/compiler-rt/lib/sanitizer_common/sanitizer_internal_defs.h
index c694897b6556b..b0a61686aa8d4 100644
--- a/compiler-rt/lib/sanitizer_common/sanitizer_internal_defs.h
+++ b/compiler-rt/lib/sanitizer_common/sanitizer_internal_defs.h
@@ -40,7 +40,7 @@
# if SANITIZER_GO
# define SANITIZER_INTERFACE_ATTRIBUTE
# define SANITIZER_WEAK_ATTRIBUTE
-# elif SANITIZER_AMDGPU || SANITIZER_NVPTX
+# elif (SANITIZER_AMDGPU && defined(__AMDGPU__)) || SANITIZER_NVPTX
# define SANITIZER_INTERFACE_ATTRIBUTE __attribute__((visibility("hidden")))
# define SANITIZER_WEAK_ATTRIBUTE __attribute__((weak))
# else
>From 7cdc873d2ecd9938187a07354bba8f568df05706 Mon Sep 17 00:00:00 2001
From: Amit Pandey <pandey.kumaramit2023 at gmail.com>
Date: Wed, 6 May 2026 13:29:39 +0530
Subject: [PATCH 18/27] [compiler-rt][HSA-Test] Add more AMDGPU ASan HSA Lit
Tests.
- Add test hsa_amd_ipc_memory_roundtrip.cpp
- Add test hsa_amd_memory_async_copy_overlap.cpp
- Add test hsa_amd_ipc_memory_attach_heap_oob.cpp
---
.../hsa_amd_ipc_memory_attach_heap_oob.cpp | 159 ++++++++++++++++++
.../AMDGPU/hsa_amd_ipc_memory_roundtrip.cpp | 149 ++++++++++++++++
.../hsa_amd_memory_async_copy_overlap.cpp | 62 +++++++
...a_amd_memory_pool_allocate_double_free.cpp | 2 +-
4 files changed, 371 insertions(+), 1 deletion(-)
create mode 100644 compiler-rt/test/asan/TestCases/AMDGPU/hsa_amd_ipc_memory_attach_heap_oob.cpp
create mode 100644 compiler-rt/test/asan/TestCases/AMDGPU/hsa_amd_ipc_memory_roundtrip.cpp
create mode 100644 compiler-rt/test/asan/TestCases/AMDGPU/hsa_amd_memory_async_copy_overlap.cpp
diff --git a/compiler-rt/test/asan/TestCases/AMDGPU/hsa_amd_ipc_memory_attach_heap_oob.cpp b/compiler-rt/test/asan/TestCases/AMDGPU/hsa_amd_ipc_memory_attach_heap_oob.cpp
new file mode 100644
index 0000000000000..ec1f38b2038f9
--- /dev/null
+++ b/compiler-rt/test/asan/TestCases/AMDGPU/hsa_amd_ipc_memory_attach_heap_oob.cpp
@@ -0,0 +1,159 @@
+// RUN: %clangxx_asan -O0 -isystem %rocm_include %s -o %t -L%rocm_lib -lhsa-runtime64 \
+// RUN: -Wl,-rpath,%rocm_lib -Wl,-rpath,%compiler_rt_libdir
+// RUN: not %run %t 2>&1 | FileCheck %s
+//
+// After hsa_amd_ipc_memory_attach, AddressSanitizer poisons a trailing redzone
+// matching the pool allocation layout. A one-past-end store must be reported as
+// a heap-buffer-overflow.
+// hsa_amd_ipc_memory_create only supports coarse-grained GPU allocations; skip
+// fine-grained pools and non-GPU agents.
+//
+// The bad store is instrumented host code; VRAM imports may need CPU access
+// enabled (best-effort hsa_amd_agents_allow_access) or the fault can be SIGSEGV
+// instead of AddressSanitizer.
+//
+// REQUIRES: sanitizer-amdgpu, linux, stable-runtime, rocm
+// UNSUPPORTED: android
+
+#include <hsa/hsa.h>
+#include <hsa/hsa_ext_amd.h>
+
+#include <stdio.h>
+
+struct PoolSearch {
+ hsa_amd_memory_pool_t pool;
+ bool found;
+};
+
+static hsa_status_t find_coarse_gpu_ipc_pool(hsa_amd_memory_pool_t pool,
+ void *data) {
+ auto *ps = static_cast<PoolSearch *>(data);
+
+ hsa_amd_segment_t segment = HSA_AMD_SEGMENT_PRIVATE;
+ if (hsa_amd_memory_pool_get_info(pool, HSA_AMD_MEMORY_POOL_INFO_SEGMENT,
+ &segment) != HSA_STATUS_SUCCESS)
+ return HSA_STATUS_SUCCESS;
+ if (segment != HSA_AMD_SEGMENT_GLOBAL)
+ return HSA_STATUS_SUCCESS;
+
+ uint32_t global_flags = 0;
+ if (hsa_amd_memory_pool_get_info(pool, HSA_AMD_MEMORY_POOL_INFO_GLOBAL_FLAGS,
+ &global_flags) != HSA_STATUS_SUCCESS)
+ return HSA_STATUS_SUCCESS;
+ if ((global_flags & HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_COARSE_GRAINED) == 0)
+ return HSA_STATUS_SUCCESS;
+
+ bool allow = false;
+ if (hsa_amd_memory_pool_get_info(
+ pool, HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_ALLOWED, &allow) !=
+ HSA_STATUS_SUCCESS ||
+ !allow)
+ return HSA_STATUS_SUCCESS;
+
+ ps->pool = pool;
+ ps->found = true;
+ return HSA_STATUS_INFO_BREAK;
+}
+
+static hsa_agent_t g_cpu_agent = {};
+
+static hsa_status_t pick_first_cpu_agent(hsa_agent_t agent, void * /*data*/) {
+ hsa_device_type_t dev = HSA_DEVICE_TYPE_GPU;
+ if (hsa_agent_get_info(agent, HSA_AGENT_INFO_DEVICE, &dev) !=
+ HSA_STATUS_SUCCESS)
+ return HSA_STATUS_SUCCESS;
+ if (dev != HSA_DEVICE_TYPE_CPU)
+ return HSA_STATUS_SUCCESS;
+ g_cpu_agent = agent;
+ return HSA_STATUS_INFO_BREAK;
+}
+
+static hsa_status_t find_gpu_agent_with_ipc_pool(hsa_agent_t agent,
+ void *data) {
+ hsa_device_type_t dev = HSA_DEVICE_TYPE_CPU;
+ if (hsa_agent_get_info(agent, HSA_AGENT_INFO_DEVICE, &dev) !=
+ HSA_STATUS_SUCCESS)
+ return HSA_STATUS_SUCCESS;
+ if (dev != HSA_DEVICE_TYPE_GPU)
+ return HSA_STATUS_SUCCESS;
+
+ auto *ps = static_cast<PoolSearch *>(data);
+ ps->found = false;
+ hsa_status_t st =
+ hsa_amd_agent_iterate_memory_pools(agent, find_coarse_gpu_ipc_pool, ps);
+ if (st != HSA_STATUS_SUCCESS && st != HSA_STATUS_INFO_BREAK)
+ return st;
+ if (ps->found)
+ return HSA_STATUS_INFO_BREAK;
+ return HSA_STATUS_SUCCESS;
+}
+
+int main() {
+ if (hsa_init() != HSA_STATUS_SUCCESS) {
+ fprintf(stderr, "hsa_init failed\n");
+ return 1;
+ }
+
+ PoolSearch ps = {};
+ ps.pool.handle = 0;
+ ps.found = false;
+
+ hsa_status_t it = hsa_iterate_agents(find_gpu_agent_with_ipc_pool, &ps);
+ if (it != HSA_STATUS_SUCCESS && it != HSA_STATUS_INFO_BREAK) {
+ fprintf(stderr, "hsa_iterate_agents failed\n");
+ return 1;
+ }
+ if (!ps.found) {
+ fprintf(stderr,
+ "no coarse-grained GPU runtime-alloc HSA memory pool found\n");
+ return 1;
+ }
+
+ constexpr size_t kBytes = 64;
+ void *mem = nullptr;
+ if (hsa_amd_memory_pool_allocate(ps.pool, kBytes, 0, &mem) !=
+ HSA_STATUS_SUCCESS ||
+ !mem) {
+ fprintf(stderr, "hsa_amd_memory_pool_allocate failed\n");
+ return 1;
+ }
+
+ hsa_amd_ipc_memory_t ipc = {};
+ if (hsa_amd_ipc_memory_create(mem, kBytes, &ipc) != HSA_STATUS_SUCCESS) {
+ fprintf(stderr, "hsa_amd_ipc_memory_create failed\n");
+ (void)hsa_amd_memory_pool_free(mem);
+ return 1;
+ }
+
+ void *mapped = nullptr;
+ if (hsa_amd_ipc_memory_attach(&ipc, kBytes, /*num_agents=*/0,
+ /*mapping_agents=*/nullptr,
+ &mapped) != HSA_STATUS_SUCCESS ||
+ !mapped) {
+ fprintf(stderr, "hsa_amd_ipc_memory_attach failed\n");
+ (void)hsa_amd_memory_pool_free(mem);
+ return 1;
+ }
+
+ g_cpu_agent.handle = 0;
+ (void)hsa_iterate_agents(pick_first_cpu_agent, nullptr);
+ if (g_cpu_agent.handle != 0) {
+ /* Best-effort: allow the host CPU to access the imported range so the store
+ below is a normal fault checked by ASan, not an unmapped-device SIGSEGV. */
+ (void)hsa_amd_agents_allow_access(/*num_agents=*/1, &g_cpu_agent,
+ /*flags=*/nullptr, mapped);
+ }
+
+ auto *p = reinterpret_cast<volatile char *>(mapped);
+ // One byte past the 64-byte imported region; should land in ASan's tail redzone.
+ p[kBytes] = 1;
+
+ fprintf(stderr, "expected heap-buffer-overflow after ipc attach\n");
+ (void)hsa_amd_ipc_memory_detach(mapped);
+ (void)hsa_amd_memory_pool_free(mem);
+ return 0;
+}
+
+// CHECK: ERROR: AddressSanitizer: heap-buffer-overflow
+// CHECK-NEXT: WRITE of size 1 at {{0x[0-9a-f]+}} thread T0
+// CHECK: SUMMARY: AddressSanitizer: heap-buffer-overflow {{.*}}hsa_amd_ipc_memory_attach_heap_oob.cpp:149:{{[0-9]+}} in main
diff --git a/compiler-rt/test/asan/TestCases/AMDGPU/hsa_amd_ipc_memory_roundtrip.cpp b/compiler-rt/test/asan/TestCases/AMDGPU/hsa_amd_ipc_memory_roundtrip.cpp
new file mode 100644
index 0000000000000..2994a9dc47460
--- /dev/null
+++ b/compiler-rt/test/asan/TestCases/AMDGPU/hsa_amd_ipc_memory_roundtrip.cpp
@@ -0,0 +1,149 @@
+// RUN: %clangxx_asan -O0 -isystem %rocm_include %s -o %t -L%rocm_lib -lhsa-runtime64 \
+// RUN: -Wl,-rpath,%rocm_lib -Wl,-rpath,%compiler_rt_libdir
+// RUN: %run %t 2>&1 | FileCheck %s
+//
+// Regression test for AddressSanitizer hsa_amd_ipc_memory_{create,attach,detach}:
+// same-process IPC round-trip on a pool allocation (unwrap on create, user pointer
+// adjustment and shadow on attach, base adjustment on detach).
+// hsa_amd_ipc_memory_create only supports coarse-grained GPU allocations; skip
+// fine-grained pools and non-GPU agents.
+//
+// Coarse-grained device memory is often not mapped for CPU stores; do not
+// read/write *mapped from the host. Validate with hsa_amd_pointer_info instead.
+//
+// REQUIRES: sanitizer-amdgpu, linux, stable-runtime, rocm
+// UNSUPPORTED: android
+
+#include <hsa/hsa.h>
+#include <hsa/hsa_ext_amd.h>
+
+#include <stdio.h>
+
+struct PoolSearch {
+ hsa_amd_memory_pool_t pool;
+ bool found;
+};
+
+static hsa_status_t find_coarse_gpu_ipc_pool(hsa_amd_memory_pool_t pool,
+ void *data) {
+ auto *ps = static_cast<PoolSearch *>(data);
+
+ hsa_amd_segment_t segment = HSA_AMD_SEGMENT_PRIVATE;
+ if (hsa_amd_memory_pool_get_info(pool, HSA_AMD_MEMORY_POOL_INFO_SEGMENT,
+ &segment) != HSA_STATUS_SUCCESS)
+ return HSA_STATUS_SUCCESS;
+ if (segment != HSA_AMD_SEGMENT_GLOBAL)
+ return HSA_STATUS_SUCCESS;
+
+ uint32_t global_flags = 0;
+ if (hsa_amd_memory_pool_get_info(pool, HSA_AMD_MEMORY_POOL_INFO_GLOBAL_FLAGS,
+ &global_flags) != HSA_STATUS_SUCCESS)
+ return HSA_STATUS_SUCCESS;
+ if ((global_flags & HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_COARSE_GRAINED) == 0)
+ return HSA_STATUS_SUCCESS;
+
+ bool allow = false;
+ if (hsa_amd_memory_pool_get_info(
+ pool, HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_ALLOWED, &allow) !=
+ HSA_STATUS_SUCCESS ||
+ !allow)
+ return HSA_STATUS_SUCCESS;
+
+ ps->pool = pool;
+ ps->found = true;
+ return HSA_STATUS_INFO_BREAK;
+}
+
+static hsa_status_t find_gpu_agent_with_ipc_pool(hsa_agent_t agent,
+ void *data) {
+ hsa_device_type_t dev = HSA_DEVICE_TYPE_CPU;
+ if (hsa_agent_get_info(agent, HSA_AGENT_INFO_DEVICE, &dev) !=
+ HSA_STATUS_SUCCESS)
+ return HSA_STATUS_SUCCESS;
+ if (dev != HSA_DEVICE_TYPE_GPU)
+ return HSA_STATUS_SUCCESS;
+
+ auto *ps = static_cast<PoolSearch *>(data);
+ ps->found = false;
+ hsa_status_t st =
+ hsa_amd_agent_iterate_memory_pools(agent, find_coarse_gpu_ipc_pool, ps);
+ if (st != HSA_STATUS_SUCCESS && st != HSA_STATUS_INFO_BREAK)
+ return st;
+ if (ps->found)
+ return HSA_STATUS_INFO_BREAK;
+ return HSA_STATUS_SUCCESS;
+}
+
+int main() {
+ if (hsa_init() != HSA_STATUS_SUCCESS) {
+ fprintf(stderr, "hsa_init failed\n");
+ return 1;
+ }
+
+ PoolSearch ps = {};
+ ps.pool.handle = 0;
+ ps.found = false;
+
+ hsa_status_t it = hsa_iterate_agents(find_gpu_agent_with_ipc_pool, &ps);
+ if (it != HSA_STATUS_SUCCESS && it != HSA_STATUS_INFO_BREAK) {
+ fprintf(stderr, "hsa_iterate_agents failed\n");
+ return 1;
+ }
+ if (!ps.found) {
+ fprintf(stderr,
+ "no coarse-grained GPU runtime-alloc HSA memory pool found\n");
+ return 1;
+ }
+
+ constexpr size_t kBytes = 64;
+ void *mem = nullptr;
+ if (hsa_amd_memory_pool_allocate(ps.pool, kBytes, 0, &mem) !=
+ HSA_STATUS_SUCCESS ||
+ !mem) {
+ fprintf(stderr, "hsa_amd_memory_pool_allocate failed\n");
+ return 1;
+ }
+
+ hsa_amd_ipc_memory_t ipc = {};
+ if (hsa_amd_ipc_memory_create(mem, kBytes, &ipc) != HSA_STATUS_SUCCESS) {
+ fprintf(stderr, "hsa_amd_ipc_memory_create failed\n");
+ (void)hsa_amd_memory_pool_free(mem);
+ return 1;
+ }
+
+ void *mapped = nullptr;
+ if (hsa_amd_ipc_memory_attach(&ipc, kBytes, /*num_agents=*/0,
+ /*mapping_agents=*/nullptr,
+ &mapped) != HSA_STATUS_SUCCESS ||
+ !mapped) {
+ fprintf(stderr, "hsa_amd_ipc_memory_attach failed\n");
+ (void)hsa_amd_memory_pool_free(mem);
+ return 1;
+ }
+
+ hsa_amd_pointer_info_t info = {};
+ info.size = sizeof(hsa_amd_pointer_info_t);
+ if (hsa_amd_pointer_info(mapped, &info, nullptr, nullptr, nullptr) !=
+ HSA_STATUS_SUCCESS) {
+ fprintf(stderr, "hsa_amd_pointer_info on imported mapping failed\n");
+ (void)hsa_amd_ipc_memory_detach(mapped);
+ (void)hsa_amd_memory_pool_free(mem);
+ return 1;
+ }
+
+ if (hsa_amd_ipc_memory_detach(mapped) != HSA_STATUS_SUCCESS) {
+ fprintf(stderr, "hsa_amd_ipc_memory_detach failed\n");
+ (void)hsa_amd_memory_pool_free(mem);
+ return 1;
+ }
+
+ if (hsa_amd_memory_pool_free(mem) != HSA_STATUS_SUCCESS) {
+ fprintf(stderr, "hsa_amd_memory_pool_free failed\n");
+ return 1;
+ }
+
+ printf("ipc roundtrip ok\n");
+ return 0;
+}
+
+// CHECK: ipc roundtrip ok
diff --git a/compiler-rt/test/asan/TestCases/AMDGPU/hsa_amd_memory_async_copy_overlap.cpp b/compiler-rt/test/asan/TestCases/AMDGPU/hsa_amd_memory_async_copy_overlap.cpp
new file mode 100644
index 0000000000000..7d4cc994d9fa1
--- /dev/null
+++ b/compiler-rt/test/asan/TestCases/AMDGPU/hsa_amd_memory_async_copy_overlap.cpp
@@ -0,0 +1,62 @@
+// RUN: %clangxx_asan -O0 -isystem %rocm_include %s -o %t -L%rocm_lib -lhsa-runtime64 \
+// RUN: -Wl,-rpath,%rocm_lib -Wl,-rpath,%compiler_rt_libdir
+// RUN: not %run %t 2>&1 | FileCheck %s
+//
+// Regression test for the AddressSanitizer hsa_amd_memory_async_copy interceptor:
+// invalid overlapping ranges are diagnosed (same family of checks as memcpy).
+//
+// REQUIRES: sanitizer-amdgpu, linux, stable-runtime, rocm
+// UNSUPPORTED: android
+
+#include <hsa/hsa.h>
+#include <hsa/hsa_ext_amd.h>
+
+#include <stdio.h>
+#include <stdlib.h>
+
+static hsa_agent_t g_agent = {};
+
+static hsa_status_t pick_first_agent(hsa_agent_t agent, void * /*data*/) {
+ g_agent = agent;
+ return HSA_STATUS_INFO_BREAK;
+}
+
+int main() {
+ if (hsa_init() != HSA_STATUS_SUCCESS) {
+ fprintf(stderr, "hsa_init failed\n");
+ return 1;
+ }
+
+ hsa_status_t it = hsa_iterate_agents(pick_first_agent, nullptr);
+ if (it != HSA_STATUS_SUCCESS && it != HSA_STATUS_INFO_BREAK) {
+ fprintf(stderr, "hsa_iterate_agents failed\n");
+ return 1;
+ }
+ if (g_agent.handle == 0) {
+ fprintf(stderr, "no HSA agent found\n");
+ return 1;
+ }
+
+ hsa_signal_t completion = {};
+ if (hsa_signal_create(/*initial_value=*/0, /*num_consumers=*/0,
+ /*consumers=*/nullptr,
+ &completion) != HSA_STATUS_SUCCESS) {
+ fprintf(stderr, "hsa_signal_create failed\n");
+ return 1;
+ }
+
+ char buf[128];
+ char *dst = buf;
+ char *src = buf + 40;
+ // Ranges [buf, buf+64) and [buf+40, buf+104) overlap; dst != src so the
+ // interceptor runs CHECK_RANGES_OVERLAP before scheduling the async copy.
+ (void)hsa_amd_memory_async_copy(dst, g_agent, src, g_agent, 64,
+ /*num_dep_signals=*/0,
+ /*dep_signals=*/nullptr, completion);
+ fprintf(stderr, "expected hsa_amd_memory_async_copy overlap report\n");
+ return 0;
+}
+
+// CHECK: hsa_amd_memory_async_copy-param-overlap: memory ranges
+// CHECK: [{{0x.*,[ ]*0x.*}}) and [{{0x.*,[ ]*0x.*}}) overlap
+// CHECK: SUMMARY: AddressSanitizer: hsa_amd_memory_async_copy-param-overlap
diff --git a/compiler-rt/test/asan/TestCases/AMDGPU/hsa_amd_memory_pool_allocate_double_free.cpp b/compiler-rt/test/asan/TestCases/AMDGPU/hsa_amd_memory_pool_allocate_double_free.cpp
index 702ec5c100d25..119b463af89e5 100644
--- a/compiler-rt/test/asan/TestCases/AMDGPU/hsa_amd_memory_pool_allocate_double_free.cpp
+++ b/compiler-rt/test/asan/TestCases/AMDGPU/hsa_amd_memory_pool_allocate_double_free.cpp
@@ -81,4 +81,4 @@ int main() {
}
// CHECK: ERROR: AddressSanitizer: attempting double-free
-// CHECK: SUMMARY: AddressSanitizer: double-free
+// CHECK: SUMMARY: AddressSanitizer: double-free {{.*}}hsa_amd_memory_pool_allocate_double_free.cpp:77:{{[0-9]+}} in main
>From 746f16b139f50d2a8bfef85822445463ff6305fd Mon Sep 17 00:00:00 2001
From: Amit Pandey <pandey.kumaramit2023 at gmail.com>
Date: Mon, 11 May 2026 11:31:02 +0530
Subject: [PATCH 19/27] [compiler-rt][asan-test][AMDGPU] Disable LeakSanitizer
for HSA lit tests.
When the AMDGPU ASan TestCases/AMDGPU suite is enabled, append
`detect_leaks=0` to ASAN_OPTIONS.
Also extend the %env_asan_opts= and %export_asan_opts= substitution
prefixes so RUN lines that override ASAN_OPTIONS stay consistent with
environment passed to plain %run invocations.
---
.../asan/TestCases/AMDGPU/lit.local.cfg.py | 30 +++++++++++++++++++
1 file changed, 30 insertions(+)
diff --git a/compiler-rt/test/asan/TestCases/AMDGPU/lit.local.cfg.py b/compiler-rt/test/asan/TestCases/AMDGPU/lit.local.cfg.py
index a544e8d3d0b9d..fc2aa6def57e4 100644
--- a/compiler-rt/test/asan/TestCases/AMDGPU/lit.local.cfg.py
+++ b/compiler-rt/test/asan/TestCases/AMDGPU/lit.local.cfg.py
@@ -44,6 +44,29 @@ def rocm_is_available(rocm_root):
return rocm_lib_dir(rocm_root) is not None
+def append_detect_leaks_off_asan_subst(substitutions):
+ """Append detect_leaks=0 to %env_asan_opts= / %export_asan_opts= default prefixes."""
+
+ def patch_body(body):
+ if body:
+ return (
+ body + "detect_leaks=0:"
+ if body.endswith(":")
+ else body + ":detect_leaks=0:"
+ )
+ return "detect_leaks=0:"
+
+ pairs = (
+ ("%env_asan_opts=", "env ASAN_OPTIONS="),
+ ("%export_asan_opts=", "export ASAN_OPTIONS="),
+ )
+ for i, (pat, repl) in enumerate(substitutions):
+ for sub_pat, opt_prefix in pairs:
+ if pat == sub_pat and repl.startswith(opt_prefix):
+ body = repl[len(opt_prefix) :]
+ substitutions[i] = (pat, opt_prefix + patch_body(body))
+
+
root = getRoot(config)
# AMDGPU ASan tests are only run with the dynamic ASan runtime (-shared-libasan).
if "asan-static-runtime" in root.available_features:
@@ -70,6 +93,13 @@ def rocm_is_available(rocm_root):
config.unsupported = True
else:
config.available_features.add("rocm")
+ # Linux ASan defaults to leak detection; disable for ROCm/HSA tests.
+ _asan = config.environment.get("ASAN_OPTIONS", "")
+ if _asan:
+ config.environment["ASAN_OPTIONS"] = _asan + ":detect_leaks=0"
+ else:
+ config.environment["ASAN_OPTIONS"] = "detect_leaks=0"
+ append_detect_leaks_off_asan_subst(config.substitutions)
rocm_lib = rocm_lib_dir(rocm_root)
rocm_include = os.path.join(rocm_root, "include")
config.substitutions.append(("%rocm_root", rocm_root))
>From 96d4fe8d00c504fcf649c9ff7315e300dce44eac Mon Sep 17 00:00:00 2001
From: Amit Pandey <pandey.kumaramit2023 at gmail.com>
Date: Mon, 11 May 2026 11:39:44 +0530
Subject: [PATCH 20/27] [compiler-rt][asan][AMDGPU] Refactor shared common HSA
utility functions into 'hsa_amd_test_helpers.h'
Add 'hsa_amd_test_helpers.h' with inline helpers for hsa_init,iterating
agents, finding a runtime-alloc pool, finding a coarse-grained GPU IPC
pool, picking CPU/first agents without globals and remove duplicated
callbacks invocation.
---
.../cmake/Modules/CompilerRTAMDGPUUtils.cmake | 6 +-
.../hsa_amd_ipc_memory_attach_heap_oob.cpp | 100 ++--------
.../AMDGPU/hsa_amd_ipc_memory_roundtrip.cpp | 76 +-------
.../hsa_amd_memory_async_copy_overlap.cpp | 25 +--
...a_amd_memory_pool_allocate_double_free.cpp | 53 +----
.../hsa_amd_pointer_info_memory_pool.cpp | 51 +----
.../TestCases/AMDGPU/hsa_amd_test_helpers.h | 184 ++++++++++++++++++
7 files changed, 225 insertions(+), 270 deletions(-)
create mode 100644 compiler-rt/test/asan/TestCases/AMDGPU/hsa_amd_test_helpers.h
diff --git a/compiler-rt/cmake/Modules/CompilerRTAMDGPUUtils.cmake b/compiler-rt/cmake/Modules/CompilerRTAMDGPUUtils.cmake
index 58a81ffdffd68..9ad77d4efb66e 100644
--- a/compiler-rt/cmake/Modules/CompilerRTAMDGPUUtils.cmake
+++ b/compiler-rt/cmake/Modules/CompilerRTAMDGPUUtils.cmake
@@ -1,4 +1,4 @@
-# ADMGPU runtime headers discovery for compiler-rt when SANITIZER_AMDGPU is enabled.
+# AMDGPU runtime headers discovery for compiler-rt when SANITIZER_AMDGPU is enabled.
#
# Usage: Include this module and call
# `compiler_rt_find_amdgpu_runtime_headers()`
@@ -39,6 +39,8 @@ macro(compiler_rt_find_amdgpu_runtime_headers)
PATH_SUFFIXES hsa
)
+ # If SANITIZER_HSA_INCLUDE_PATH points at the leaf `hsa` directory,
+ # retry without PATH_SUFFIXES.
if(NOT HSA_INCLUDE_DIR AND SANITIZER_HSA_INCLUDE_PATH)
find_path(
HSA_INCLUDE_DIR
@@ -84,6 +86,8 @@ macro(compiler_rt_find_amdgpu_runtime_headers)
)
endif()
+ # If SANITIZER_COMGR_INCLUDE_PATH points at the leaf `amd_comgr` directory,
+ # retry without PATH_SUFFIXES.
if(NOT AMDComgr_INCLUDE_DIR AND SANITIZER_COMGR_INCLUDE_PATH)
find_path(
AMDComgr_INCLUDE_DIR
diff --git a/compiler-rt/test/asan/TestCases/AMDGPU/hsa_amd_ipc_memory_attach_heap_oob.cpp b/compiler-rt/test/asan/TestCases/AMDGPU/hsa_amd_ipc_memory_attach_heap_oob.cpp
index ec1f38b2038f9..f39b86e6fe333 100644
--- a/compiler-rt/test/asan/TestCases/AMDGPU/hsa_amd_ipc_memory_attach_heap_oob.cpp
+++ b/compiler-rt/test/asan/TestCases/AMDGPU/hsa_amd_ipc_memory_attach_heap_oob.cpp
@@ -15,99 +15,20 @@
// REQUIRES: sanitizer-amdgpu, linux, stable-runtime, rocm
// UNSUPPORTED: android
+#include "hsa_amd_test_helpers.h"
+
#include <hsa/hsa.h>
#include <hsa/hsa_ext_amd.h>
#include <stdio.h>
-struct PoolSearch {
- hsa_amd_memory_pool_t pool;
- bool found;
-};
-
-static hsa_status_t find_coarse_gpu_ipc_pool(hsa_amd_memory_pool_t pool,
- void *data) {
- auto *ps = static_cast<PoolSearch *>(data);
-
- hsa_amd_segment_t segment = HSA_AMD_SEGMENT_PRIVATE;
- if (hsa_amd_memory_pool_get_info(pool, HSA_AMD_MEMORY_POOL_INFO_SEGMENT,
- &segment) != HSA_STATUS_SUCCESS)
- return HSA_STATUS_SUCCESS;
- if (segment != HSA_AMD_SEGMENT_GLOBAL)
- return HSA_STATUS_SUCCESS;
-
- uint32_t global_flags = 0;
- if (hsa_amd_memory_pool_get_info(pool, HSA_AMD_MEMORY_POOL_INFO_GLOBAL_FLAGS,
- &global_flags) != HSA_STATUS_SUCCESS)
- return HSA_STATUS_SUCCESS;
- if ((global_flags & HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_COARSE_GRAINED) == 0)
- return HSA_STATUS_SUCCESS;
-
- bool allow = false;
- if (hsa_amd_memory_pool_get_info(
- pool, HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_ALLOWED, &allow) !=
- HSA_STATUS_SUCCESS ||
- !allow)
- return HSA_STATUS_SUCCESS;
-
- ps->pool = pool;
- ps->found = true;
- return HSA_STATUS_INFO_BREAK;
-}
-
-static hsa_agent_t g_cpu_agent = {};
-
-static hsa_status_t pick_first_cpu_agent(hsa_agent_t agent, void * /*data*/) {
- hsa_device_type_t dev = HSA_DEVICE_TYPE_GPU;
- if (hsa_agent_get_info(agent, HSA_AGENT_INFO_DEVICE, &dev) !=
- HSA_STATUS_SUCCESS)
- return HSA_STATUS_SUCCESS;
- if (dev != HSA_DEVICE_TYPE_CPU)
- return HSA_STATUS_SUCCESS;
- g_cpu_agent = agent;
- return HSA_STATUS_INFO_BREAK;
-}
-
-static hsa_status_t find_gpu_agent_with_ipc_pool(hsa_agent_t agent,
- void *data) {
- hsa_device_type_t dev = HSA_DEVICE_TYPE_CPU;
- if (hsa_agent_get_info(agent, HSA_AGENT_INFO_DEVICE, &dev) !=
- HSA_STATUS_SUCCESS)
- return HSA_STATUS_SUCCESS;
- if (dev != HSA_DEVICE_TYPE_GPU)
- return HSA_STATUS_SUCCESS;
-
- auto *ps = static_cast<PoolSearch *>(data);
- ps->found = false;
- hsa_status_t st =
- hsa_amd_agent_iterate_memory_pools(agent, find_coarse_gpu_ipc_pool, ps);
- if (st != HSA_STATUS_SUCCESS && st != HSA_STATUS_INFO_BREAK)
- return st;
- if (ps->found)
- return HSA_STATUS_INFO_BREAK;
- return HSA_STATUS_SUCCESS;
-}
-
int main() {
- if (hsa_init() != HSA_STATUS_SUCCESS) {
- fprintf(stderr, "hsa_init failed\n");
+ if (hsa_amd_test_require_init())
return 1;
- }
-
- PoolSearch ps = {};
- ps.pool.handle = 0;
- ps.found = false;
- hsa_status_t it = hsa_iterate_agents(find_gpu_agent_with_ipc_pool, &ps);
- if (it != HSA_STATUS_SUCCESS && it != HSA_STATUS_INFO_BREAK) {
- fprintf(stderr, "hsa_iterate_agents failed\n");
+ HsaAmdPoolSearch ps;
+ if (hsa_amd_test_find_first_coarse_gpu_ipc_pool(&ps))
return 1;
- }
- if (!ps.found) {
- fprintf(stderr,
- "no coarse-grained GPU runtime-alloc HSA memory pool found\n");
- return 1;
- }
constexpr size_t kBytes = 64;
void *mem = nullptr;
@@ -135,12 +56,13 @@ int main() {
return 1;
}
- g_cpu_agent.handle = 0;
- (void)hsa_iterate_agents(pick_first_cpu_agent, nullptr);
- if (g_cpu_agent.handle != 0) {
+ HsaAmdCpuAgentPick cpu;
+ hsa_amd_test_cpu_agent_pick_init(&cpu);
+ (void)hsa_iterate_agents(hsa_amd_test_pick_first_cpu_agent_cb, &cpu);
+ if (cpu.agent.handle != 0) {
/* Best-effort: allow the host CPU to access the imported range so the store
below is a normal fault checked by ASan, not an unmapped-device SIGSEGV. */
- (void)hsa_amd_agents_allow_access(/*num_agents=*/1, &g_cpu_agent,
+ (void)hsa_amd_agents_allow_access(/*num_agents=*/1, &cpu.agent,
/*flags=*/nullptr, mapped);
}
@@ -156,4 +78,4 @@ int main() {
// CHECK: ERROR: AddressSanitizer: heap-buffer-overflow
// CHECK-NEXT: WRITE of size 1 at {{0x[0-9a-f]+}} thread T0
-// CHECK: SUMMARY: AddressSanitizer: heap-buffer-overflow {{.*}}hsa_amd_ipc_memory_attach_heap_oob.cpp:149:{{[0-9]+}} in main
+// CHECK: SUMMARY: AddressSanitizer: heap-buffer-overflow {{.*}}hsa_amd_ipc_memory_attach_heap_oob.cpp:71:{{[0-9]+}} in main
diff --git a/compiler-rt/test/asan/TestCases/AMDGPU/hsa_amd_ipc_memory_roundtrip.cpp b/compiler-rt/test/asan/TestCases/AMDGPU/hsa_amd_ipc_memory_roundtrip.cpp
index 2994a9dc47460..f84b1bbdde2a2 100644
--- a/compiler-rt/test/asan/TestCases/AMDGPU/hsa_amd_ipc_memory_roundtrip.cpp
+++ b/compiler-rt/test/asan/TestCases/AMDGPU/hsa_amd_ipc_memory_roundtrip.cpp
@@ -14,86 +14,20 @@
// REQUIRES: sanitizer-amdgpu, linux, stable-runtime, rocm
// UNSUPPORTED: android
+#include "hsa_amd_test_helpers.h"
+
#include <hsa/hsa.h>
#include <hsa/hsa_ext_amd.h>
#include <stdio.h>
-struct PoolSearch {
- hsa_amd_memory_pool_t pool;
- bool found;
-};
-
-static hsa_status_t find_coarse_gpu_ipc_pool(hsa_amd_memory_pool_t pool,
- void *data) {
- auto *ps = static_cast<PoolSearch *>(data);
-
- hsa_amd_segment_t segment = HSA_AMD_SEGMENT_PRIVATE;
- if (hsa_amd_memory_pool_get_info(pool, HSA_AMD_MEMORY_POOL_INFO_SEGMENT,
- &segment) != HSA_STATUS_SUCCESS)
- return HSA_STATUS_SUCCESS;
- if (segment != HSA_AMD_SEGMENT_GLOBAL)
- return HSA_STATUS_SUCCESS;
-
- uint32_t global_flags = 0;
- if (hsa_amd_memory_pool_get_info(pool, HSA_AMD_MEMORY_POOL_INFO_GLOBAL_FLAGS,
- &global_flags) != HSA_STATUS_SUCCESS)
- return HSA_STATUS_SUCCESS;
- if ((global_flags & HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_COARSE_GRAINED) == 0)
- return HSA_STATUS_SUCCESS;
-
- bool allow = false;
- if (hsa_amd_memory_pool_get_info(
- pool, HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_ALLOWED, &allow) !=
- HSA_STATUS_SUCCESS ||
- !allow)
- return HSA_STATUS_SUCCESS;
-
- ps->pool = pool;
- ps->found = true;
- return HSA_STATUS_INFO_BREAK;
-}
-
-static hsa_status_t find_gpu_agent_with_ipc_pool(hsa_agent_t agent,
- void *data) {
- hsa_device_type_t dev = HSA_DEVICE_TYPE_CPU;
- if (hsa_agent_get_info(agent, HSA_AGENT_INFO_DEVICE, &dev) !=
- HSA_STATUS_SUCCESS)
- return HSA_STATUS_SUCCESS;
- if (dev != HSA_DEVICE_TYPE_GPU)
- return HSA_STATUS_SUCCESS;
-
- auto *ps = static_cast<PoolSearch *>(data);
- ps->found = false;
- hsa_status_t st =
- hsa_amd_agent_iterate_memory_pools(agent, find_coarse_gpu_ipc_pool, ps);
- if (st != HSA_STATUS_SUCCESS && st != HSA_STATUS_INFO_BREAK)
- return st;
- if (ps->found)
- return HSA_STATUS_INFO_BREAK;
- return HSA_STATUS_SUCCESS;
-}
-
int main() {
- if (hsa_init() != HSA_STATUS_SUCCESS) {
- fprintf(stderr, "hsa_init failed\n");
+ if (hsa_amd_test_require_init())
return 1;
- }
- PoolSearch ps = {};
- ps.pool.handle = 0;
- ps.found = false;
-
- hsa_status_t it = hsa_iterate_agents(find_gpu_agent_with_ipc_pool, &ps);
- if (it != HSA_STATUS_SUCCESS && it != HSA_STATUS_INFO_BREAK) {
- fprintf(stderr, "hsa_iterate_agents failed\n");
- return 1;
- }
- if (!ps.found) {
- fprintf(stderr,
- "no coarse-grained GPU runtime-alloc HSA memory pool found\n");
+ HsaAmdPoolSearch ps;
+ if (hsa_amd_test_find_first_coarse_gpu_ipc_pool(&ps))
return 1;
- }
constexpr size_t kBytes = 64;
void *mem = nullptr;
diff --git a/compiler-rt/test/asan/TestCases/AMDGPU/hsa_amd_memory_async_copy_overlap.cpp b/compiler-rt/test/asan/TestCases/AMDGPU/hsa_amd_memory_async_copy_overlap.cpp
index 7d4cc994d9fa1..c773b55b8e866 100644
--- a/compiler-rt/test/asan/TestCases/AMDGPU/hsa_amd_memory_async_copy_overlap.cpp
+++ b/compiler-rt/test/asan/TestCases/AMDGPU/hsa_amd_memory_async_copy_overlap.cpp
@@ -8,31 +8,24 @@
// REQUIRES: sanitizer-amdgpu, linux, stable-runtime, rocm
// UNSUPPORTED: android
+#include "hsa_amd_test_helpers.h"
+
#include <hsa/hsa.h>
#include <hsa/hsa_ext_amd.h>
#include <stdio.h>
#include <stdlib.h>
-static hsa_agent_t g_agent = {};
-
-static hsa_status_t pick_first_agent(hsa_agent_t agent, void * /*data*/) {
- g_agent = agent;
- return HSA_STATUS_INFO_BREAK;
-}
-
int main() {
- if (hsa_init() != HSA_STATUS_SUCCESS) {
- fprintf(stderr, "hsa_init failed\n");
+ if (hsa_amd_test_require_init())
return 1;
- }
- hsa_status_t it = hsa_iterate_agents(pick_first_agent, nullptr);
- if (it != HSA_STATUS_SUCCESS && it != HSA_STATUS_INFO_BREAK) {
- fprintf(stderr, "hsa_iterate_agents failed\n");
+ HsaAmdAgentPick pick;
+ hsa_amd_test_agent_pick_init(&pick);
+ hsa_status_t it = hsa_iterate_agents(hsa_amd_test_pick_first_agent_cb, &pick);
+ if (hsa_amd_test_iterate_agents_ok(it))
return 1;
- }
- if (g_agent.handle == 0) {
+ if (pick.agent.handle == 0) {
fprintf(stderr, "no HSA agent found\n");
return 1;
}
@@ -50,7 +43,7 @@ int main() {
char *src = buf + 40;
// Ranges [buf, buf+64) and [buf+40, buf+104) overlap; dst != src so the
// interceptor runs CHECK_RANGES_OVERLAP before scheduling the async copy.
- (void)hsa_amd_memory_async_copy(dst, g_agent, src, g_agent, 64,
+ (void)hsa_amd_memory_async_copy(dst, pick.agent, src, pick.agent, 64,
/*num_dep_signals=*/0,
/*dep_signals=*/nullptr, completion);
fprintf(stderr, "expected hsa_amd_memory_async_copy overlap report\n");
diff --git a/compiler-rt/test/asan/TestCases/AMDGPU/hsa_amd_memory_pool_allocate_double_free.cpp b/compiler-rt/test/asan/TestCases/AMDGPU/hsa_amd_memory_pool_allocate_double_free.cpp
index 119b463af89e5..82b21962b0948 100644
--- a/compiler-rt/test/asan/TestCases/AMDGPU/hsa_amd_memory_pool_allocate_double_free.cpp
+++ b/compiler-rt/test/asan/TestCases/AMDGPU/hsa_amd_memory_pool_allocate_double_free.cpp
@@ -9,61 +9,20 @@
// REQUIRES: sanitizer-amdgpu, linux, stable-runtime, rocm
// UNSUPPORTED: android
+#include "hsa_amd_test_helpers.h"
+
#include <hsa/hsa.h>
#include <hsa/hsa_ext_amd.h>
#include <stdio.h>
-struct PoolSearch {
- hsa_amd_memory_pool_t pool;
- bool found;
-};
-
-static hsa_status_t find_alloc_pool(hsa_amd_memory_pool_t pool, void *data) {
- auto *ps = static_cast<PoolSearch *>(data);
- bool allow = false;
- if (hsa_amd_memory_pool_get_info(
- pool, HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_ALLOWED, &allow) !=
- HSA_STATUS_SUCCESS ||
- !allow)
- return HSA_STATUS_SUCCESS;
- ps->pool = pool;
- ps->found = true;
- return HSA_STATUS_INFO_BREAK;
-}
-
-static hsa_status_t find_agent_with_pool(hsa_agent_t agent, void *data) {
- (void)agent;
- auto *ps = static_cast<PoolSearch *>(data);
- ps->found = false;
- hsa_status_t st =
- hsa_amd_agent_iterate_memory_pools(agent, find_alloc_pool, ps);
- if (st != HSA_STATUS_SUCCESS && st != HSA_STATUS_INFO_BREAK)
- return st;
- if (ps->found)
- return HSA_STATUS_INFO_BREAK;
- return HSA_STATUS_SUCCESS;
-}
-
int main() {
- if (hsa_init() != HSA_STATUS_SUCCESS) {
- fprintf(stderr, "hsa_init failed\n");
+ if (hsa_amd_test_require_init())
return 1;
- }
-
- PoolSearch ps = {};
- ps.pool.handle = 0;
- ps.found = false;
- hsa_status_t it = hsa_iterate_agents(find_agent_with_pool, &ps);
- if (it != HSA_STATUS_SUCCESS && it != HSA_STATUS_INFO_BREAK) {
- fprintf(stderr, "hsa_iterate_agents failed\n");
+ HsaAmdPoolSearch ps;
+ if (hsa_amd_test_find_first_runtime_alloc_pool(&ps))
return 1;
- }
- if (!ps.found) {
- fprintf(stderr, "no runtime-alloc HSA memory pool found\n");
- return 1;
- }
void *mem = nullptr;
if (hsa_amd_memory_pool_allocate(ps.pool, 64, 0, &mem) !=
@@ -81,4 +40,4 @@ int main() {
}
// CHECK: ERROR: AddressSanitizer: attempting double-free
-// CHECK: SUMMARY: AddressSanitizer: double-free {{.*}}hsa_amd_memory_pool_allocate_double_free.cpp:77:{{[0-9]+}} in main
+// CHECK: SUMMARY: AddressSanitizer: double-free {{.*}}hsa_amd_memory_pool_allocate_double_free.cpp:36:{{[0-9]+}} in main
diff --git a/compiler-rt/test/asan/TestCases/AMDGPU/hsa_amd_pointer_info_memory_pool.cpp b/compiler-rt/test/asan/TestCases/AMDGPU/hsa_amd_pointer_info_memory_pool.cpp
index 13a05b8b35237..c6c54373210e5 100644
--- a/compiler-rt/test/asan/TestCases/AMDGPU/hsa_amd_pointer_info_memory_pool.cpp
+++ b/compiler-rt/test/asan/TestCases/AMDGPU/hsa_amd_pointer_info_memory_pool.cpp
@@ -9,61 +9,20 @@
// REQUIRES: sanitizer-amdgpu, linux, stable-runtime, rocm
// UNSUPPORTED: android
+#include "hsa_amd_test_helpers.h"
+
#include <hsa/hsa.h>
#include <hsa/hsa_ext_amd.h>
#include <stdio.h>
-struct PoolSearch {
- hsa_amd_memory_pool_t pool;
- bool found;
-};
-
-static hsa_status_t find_alloc_pool(hsa_amd_memory_pool_t pool, void *data) {
- auto *ps = static_cast<PoolSearch *>(data);
- bool allow = false;
- if (hsa_amd_memory_pool_get_info(
- pool, HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_ALLOWED, &allow) !=
- HSA_STATUS_SUCCESS ||
- !allow)
- return HSA_STATUS_SUCCESS;
- ps->pool = pool;
- ps->found = true;
- return HSA_STATUS_INFO_BREAK;
-}
-
-static hsa_status_t find_agent_with_pool(hsa_agent_t agent, void *data) {
- (void)agent;
- auto *ps = static_cast<PoolSearch *>(data);
- ps->found = false;
- hsa_status_t st =
- hsa_amd_agent_iterate_memory_pools(agent, find_alloc_pool, ps);
- if (st != HSA_STATUS_SUCCESS && st != HSA_STATUS_INFO_BREAK)
- return st;
- if (ps->found)
- return HSA_STATUS_INFO_BREAK;
- return HSA_STATUS_SUCCESS;
-}
-
int main() {
- if (hsa_init() != HSA_STATUS_SUCCESS) {
- fprintf(stderr, "hsa_init failed\n");
+ if (hsa_amd_test_require_init())
return 1;
- }
-
- PoolSearch ps = {};
- ps.pool.handle = 0;
- ps.found = false;
- hsa_status_t it = hsa_iterate_agents(find_agent_with_pool, &ps);
- if (it != HSA_STATUS_SUCCESS && it != HSA_STATUS_INFO_BREAK) {
- fprintf(stderr, "hsa_iterate_agents failed\n");
+ HsaAmdPoolSearch ps;
+ if (hsa_amd_test_find_first_runtime_alloc_pool(&ps))
return 1;
- }
- if (!ps.found) {
- fprintf(stderr, "no runtime-alloc HSA memory pool found\n");
- return 1;
- }
void *mem = nullptr;
if (hsa_amd_memory_pool_allocate(ps.pool, 64, 0, &mem) !=
diff --git a/compiler-rt/test/asan/TestCases/AMDGPU/hsa_amd_test_helpers.h b/compiler-rt/test/asan/TestCases/AMDGPU/hsa_amd_test_helpers.h
new file mode 100644
index 0000000000000..53b7f571fcbcf
--- /dev/null
+++ b/compiler-rt/test/asan/TestCases/AMDGPU/hsa_amd_test_helpers.h
@@ -0,0 +1,184 @@
+//===-- hsa_amd_test_helpers.h - shared helpers for AMDGPU ASan HSA tests --===//
+//
+// Common ROCm/HSA discovery and init helpers for hsa_amd*.cpp tests in this
+// directory. Intended only for compiler-rt lit tests.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef COMPILER_RT_ASAN_AMDGPU_HSA_AMD_TEST_HELPERS_H
+#define COMPILER_RT_ASAN_AMDGPU_HSA_AMD_TEST_HELPERS_H
+
+#include <hsa/hsa.h>
+#include <hsa/hsa_ext_amd.h>
+
+#include <stdio.h>
+
+inline int hsa_amd_test_require_init() {
+ if (hsa_init() != HSA_STATUS_SUCCESS) {
+ fprintf(stderr, "hsa_init failed\n");
+ return 1;
+ }
+ return 0;
+}
+
+/// Return 1 if `st` indicates a hard failure from hsa_iterate_agents.
+inline int hsa_amd_test_iterate_agents_ok(hsa_status_t st) {
+ if (st != HSA_STATUS_SUCCESS && st != HSA_STATUS_INFO_BREAK) {
+ fprintf(stderr, "hsa_iterate_agents failed\n");
+ return 1;
+ }
+ return 0;
+}
+
+struct HsaAmdPoolSearch {
+ hsa_amd_memory_pool_t pool;
+ bool found;
+};
+
+inline hsa_status_t
+hsa_amd_test_find_runtime_alloc_pool_cb(hsa_amd_memory_pool_t pool,
+ void *data) {
+ auto *ps = static_cast<HsaAmdPoolSearch *>(data);
+ bool allow = false;
+ if (hsa_amd_memory_pool_get_info(
+ pool, HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_ALLOWED, &allow) !=
+ HSA_STATUS_SUCCESS ||
+ !allow)
+ return HSA_STATUS_SUCCESS;
+ ps->pool = pool;
+ ps->found = true;
+ return HSA_STATUS_INFO_BREAK;
+}
+
+inline hsa_status_t
+hsa_amd_test_find_agent_with_runtime_alloc_pool_cb(hsa_agent_t agent,
+ void *data) {
+ (void)agent;
+ auto *ps = static_cast<HsaAmdPoolSearch *>(data);
+ ps->found = false;
+ hsa_status_t st = hsa_amd_agent_iterate_memory_pools(
+ agent, hsa_amd_test_find_runtime_alloc_pool_cb, ps);
+ if (st != HSA_STATUS_SUCCESS && st != HSA_STATUS_INFO_BREAK)
+ return st;
+ if (ps->found)
+ return HSA_STATUS_INFO_BREAK;
+ return HSA_STATUS_SUCCESS;
+}
+
+inline int hsa_amd_test_find_first_runtime_alloc_pool(HsaAmdPoolSearch *ps) {
+ ps->pool.handle = 0;
+ ps->found = false;
+ hsa_status_t it = hsa_iterate_agents(
+ hsa_amd_test_find_agent_with_runtime_alloc_pool_cb, ps);
+ if (hsa_amd_test_iterate_agents_ok(it))
+ return 1;
+ if (!ps->found) {
+ fprintf(stderr, "no runtime-alloc HSA memory pool found\n");
+ return 1;
+ }
+ return 0;
+}
+
+inline hsa_status_t
+hsa_amd_test_find_coarse_gpu_ipc_pool_cb(hsa_amd_memory_pool_t pool,
+ void *data) {
+ auto *ps = static_cast<HsaAmdPoolSearch *>(data);
+
+ hsa_amd_segment_t segment = HSA_AMD_SEGMENT_PRIVATE;
+ if (hsa_amd_memory_pool_get_info(pool, HSA_AMD_MEMORY_POOL_INFO_SEGMENT,
+ &segment) != HSA_STATUS_SUCCESS)
+ return HSA_STATUS_SUCCESS;
+ if (segment != HSA_AMD_SEGMENT_GLOBAL)
+ return HSA_STATUS_SUCCESS;
+
+ uint32_t global_flags = 0;
+ if (hsa_amd_memory_pool_get_info(pool, HSA_AMD_MEMORY_POOL_INFO_GLOBAL_FLAGS,
+ &global_flags) != HSA_STATUS_SUCCESS)
+ return HSA_STATUS_SUCCESS;
+ if ((global_flags & HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_COARSE_GRAINED) == 0)
+ return HSA_STATUS_SUCCESS;
+
+ bool allow = false;
+ if (hsa_amd_memory_pool_get_info(
+ pool, HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_ALLOWED, &allow) !=
+ HSA_STATUS_SUCCESS ||
+ !allow)
+ return HSA_STATUS_SUCCESS;
+
+ ps->pool = pool;
+ ps->found = true;
+ return HSA_STATUS_INFO_BREAK;
+}
+
+inline hsa_status_t
+hsa_amd_test_find_gpu_agent_with_ipc_pool_cb(hsa_agent_t agent, void *data) {
+ hsa_device_type_t dev = HSA_DEVICE_TYPE_CPU;
+ if (hsa_agent_get_info(agent, HSA_AGENT_INFO_DEVICE, &dev) !=
+ HSA_STATUS_SUCCESS)
+ return HSA_STATUS_SUCCESS;
+ if (dev != HSA_DEVICE_TYPE_GPU)
+ return HSA_STATUS_SUCCESS;
+
+ auto *ps = static_cast<HsaAmdPoolSearch *>(data);
+ ps->found = false;
+ hsa_status_t st = hsa_amd_agent_iterate_memory_pools(
+ agent, hsa_amd_test_find_coarse_gpu_ipc_pool_cb, ps);
+ if (st != HSA_STATUS_SUCCESS && st != HSA_STATUS_INFO_BREAK)
+ return st;
+ if (ps->found)
+ return HSA_STATUS_INFO_BREAK;
+ return HSA_STATUS_SUCCESS;
+}
+
+inline int hsa_amd_test_find_first_coarse_gpu_ipc_pool(HsaAmdPoolSearch *ps) {
+ ps->pool.handle = 0;
+ ps->found = false;
+ hsa_status_t it =
+ hsa_iterate_agents(hsa_amd_test_find_gpu_agent_with_ipc_pool_cb, ps);
+ if (hsa_amd_test_iterate_agents_ok(it))
+ return 1;
+ if (!ps->found) {
+ fprintf(stderr,
+ "no coarse-grained GPU runtime-alloc HSA memory pool found\n");
+ return 1;
+ }
+ return 0;
+}
+
+struct HsaAmdCpuAgentPick {
+ hsa_agent_t agent;
+};
+
+inline void hsa_amd_test_cpu_agent_pick_init(HsaAmdCpuAgentPick *out) {
+ out->agent.handle = 0;
+}
+
+inline hsa_status_t hsa_amd_test_pick_first_cpu_agent_cb(hsa_agent_t agent,
+ void *data) {
+ auto *out = static_cast<HsaAmdCpuAgentPick *>(data);
+ hsa_device_type_t dev = HSA_DEVICE_TYPE_GPU;
+ if (hsa_agent_get_info(agent, HSA_AGENT_INFO_DEVICE, &dev) !=
+ HSA_STATUS_SUCCESS)
+ return HSA_STATUS_SUCCESS;
+ if (dev != HSA_DEVICE_TYPE_CPU)
+ return HSA_STATUS_SUCCESS;
+ out->agent = agent;
+ return HSA_STATUS_INFO_BREAK;
+}
+
+struct HsaAmdAgentPick {
+ hsa_agent_t agent;
+};
+
+inline void hsa_amd_test_agent_pick_init(HsaAmdAgentPick *out) {
+ out->agent.handle = 0;
+}
+
+inline hsa_status_t hsa_amd_test_pick_first_agent_cb(hsa_agent_t agent,
+ void *data) {
+ auto *out = static_cast<HsaAmdAgentPick *>(data);
+ out->agent = agent;
+ return HSA_STATUS_INFO_BREAK;
+}
+
+#endif // COMPILER_RT_ASAN_AMDGPU_HSA_AMD_TEST_HELPERS_H
>From 536af5327b4cf42972b0f6dbea509d31b4733f7f Mon Sep 17 00:00:00 2001
From: Amit Pandey <pandey.kumaramit2023 at gmail.com>
Date: Thu, 14 May 2026 14:19:47 +0530
Subject: [PATCH 21/27] Address Comments of @bing-ma.
- Add same reserved va/size check validation logic at the time of
freeing VMEM memory.
- Ensure ASan HSA interceptors never return `HSA_STATUS_SUCCESS` when
allocation yields a null pointer through following cases like early
device runtime shutdown, failed device runtime initialization,
device allocator size overflow, or unsuccessful device allocation
without an updated HSA status).
- ROCr can return from hsa_amd_ipc_memory_detach before the import is
fully torn down (e.g. thunk_bo when hsaKmtMemoryVaUnmap fails).
Defer PoisonShadow and FlushUnneededASanShadowMemory calls until
detach succeeds so synthetic redzones remain while the mapping may
still be live.
---
compiler-rt/lib/asan/asan_allocator.cpp | 26 +++++++++++++++----
.../sanitizer_allocator_amdgpu.cpp | 6 +++--
.../sanitizer_allocator_amdgpu.h | 6 +++++
.../sanitizer_allocator_device.h | 15 +++++++++--
4 files changed, 44 insertions(+), 9 deletions(-)
diff --git a/compiler-rt/lib/asan/asan_allocator.cpp b/compiler-rt/lib/asan/asan_allocator.cpp
index 4acc456a244ca..1fbbbb0fded2d 100644
--- a/compiler-rt/lib/asan/asan_allocator.cpp
+++ b/compiler-rt/lib/asan/asan_allocator.cpp
@@ -1615,16 +1615,25 @@ hsa_status_t asan_hsa_amd_ipc_memory_detach(void* mapped_ptr) {
static_assert(AP_.kMetadataSize == 0, "Expression below requires this");
uptr mapped_base = reinterpret_cast<uptr>(mapped_ptr) - kPageSize_;
+ // Snapshot mapping size before detach: ROCr may return an error while the
+ // import is still live (e.g. thunk_bo path if hsaKmtMemoryVaUnmap fails).
+ // Only clear ASan shadow after a successful detach so redzones stay valid
+ // if teardown did not complete.
hsa_amd_pointer_info_t info;
info.size = sizeof(hsa_amd_pointer_info_t);
+ uptr mapped_sz = 0;
if (REAL(hsa_amd_pointer_info)(reinterpret_cast<void*>(mapped_base), &info,
nullptr, nullptr,
- nullptr) == HSA_STATUS_SUCCESS) {
- PoisonShadow(mapped_base, info.sizeInBytes, 0);
- FlushUnneededASanShadowMemory(mapped_base, info.sizeInBytes);
- }
+ nullptr) == HSA_STATUS_SUCCESS)
+ mapped_sz = static_cast<uptr>(info.sizeInBytes);
- return REAL(hsa_amd_ipc_memory_detach)(reinterpret_cast<void*>(mapped_base));
+ const hsa_status_t status =
+ REAL(hsa_amd_ipc_memory_detach)(reinterpret_cast<void*>(mapped_base));
+ if (status == HSA_STATUS_SUCCESS && mapped_sz) {
+ PoisonShadow(mapped_base, mapped_sz, 0);
+ FlushUnneededASanShadowMemory(mapped_base, mapped_sz);
+ }
+ return status;
}
hsa_status_t asan_hsa_amd_vmem_address_reserve_align(
@@ -1675,6 +1684,13 @@ hsa_status_t asan_hsa_amd_vmem_address_free(void* ptr, size_t size,
void* p = get_allocator().GetBlockBegin(ptr);
if (p) {
+ // Match ROCr: require the same va/size as reserve before freeing metadata.
+ uptr ptr_ = reinterpret_cast<uptr>(ptr);
+ AsanChunk* m = reinterpret_cast<AsanChunk*>(ptr_ - kChunkHeaderSize);
+ if (m->Beg() != ptr_ || size != m->UsedSize()) {
+ errno = errno_EINVAL;
+ return HSA_STATUS_ERROR_INVALID_ARGUMENT;
+ }
instance.Deallocate(ptr, 0, 0, stack, FROM_MALLOC);
return HSA_STATUS_SUCCESS;
}
diff --git a/compiler-rt/lib/sanitizer_common/sanitizer_allocator_amdgpu.cpp b/compiler-rt/lib/sanitizer_common/sanitizer_allocator_amdgpu.cpp
index 3597db7c9d1bf..87db693476abd 100644
--- a/compiler-rt/lib/sanitizer_common/sanitizer_allocator_amdgpu.cpp
+++ b/compiler-rt/lib/sanitizer_common/sanitizer_allocator_amdgpu.cpp
@@ -99,17 +99,19 @@ bool AmdgpuMemFuncs::Init() {
void* AmdgpuMemFuncs::Allocate(uptr size, uptr alignment,
DeviceAllocationInfo* da_info) {
+ AmdgpuAllocationInfo* aa_info =
+ reinterpret_cast<AmdgpuAllocationInfo*>(da_info);
+
// Do not allocate if AMDGPU runtime is shutdown
if (UNLIKELY(IsAmdgpuRuntimeShutdown())) {
VReport(1,
"Amdgpu Allocate: Runtime shutdown, skipping allocation for size "
"%zu alignment %zu\n",
size, alignment);
+ aa_info->EnsureFailureStatus(HSA_STATUS_ERROR_INVALID_RUNTIME_STATE);
return nullptr;
}
- AmdgpuAllocationInfo* aa_info =
- reinterpret_cast<AmdgpuAllocationInfo*>(da_info);
if (!aa_info->memory_pool.handle) {
aa_info->status = hsa_amd.vmem_address_reserve_align(
&aa_info->ptr, size, aa_info->address, aa_info->alignment,
diff --git a/compiler-rt/lib/sanitizer_common/sanitizer_allocator_amdgpu.h b/compiler-rt/lib/sanitizer_common/sanitizer_allocator_amdgpu.h
index ff77911efb546..0f11b7a343d0a 100644
--- a/compiler-rt/lib/sanitizer_common/sanitizer_allocator_amdgpu.h
+++ b/compiler-rt/lib/sanitizer_common/sanitizer_allocator_amdgpu.h
@@ -35,6 +35,12 @@ struct AmdgpuAllocationInfo : public DeviceAllocationInfo {
status = HSA_STATUS_SUCCESS;
alloc_func = nullptr;
}
+ // If allocation fails without an HSA API status, record one so callers never
+ // see *ptr == nullptr with status still SUCCESS.
+ void EnsureFailureStatus(hsa_status_t err) {
+ if (status == HSA_STATUS_SUCCESS)
+ status = err;
+ }
hsa_status_t status;
void* alloc_func;
hsa_amd_memory_pool_t memory_pool;
diff --git a/compiler-rt/lib/sanitizer_common/sanitizer_allocator_device.h b/compiler-rt/lib/sanitizer_common/sanitizer_allocator_device.h
index c6a79760e902a..eb29ed7bde0f1 100644
--- a/compiler-rt/lib/sanitizer_common/sanitizer_allocator_device.h
+++ b/compiler-rt/lib/sanitizer_common/sanitizer_allocator_device.h
@@ -57,8 +57,12 @@ class DeviceAllocatorT {
void* Allocate(AllocatorStats* stat, uptr size, uptr alignment,
DeviceAllocationInfo* da_info) {
- if (!da_info || !InitMemFuncs())
+ if (!da_info || !InitMemFuncs()) {
+ if (da_info && da_info->type_ == DAT_AMDGPU)
+ reinterpret_cast<AmdgpuAllocationInfo*>(da_info)->EnsureFailureStatus(
+ HSA_STATUS_ERROR_NOT_INITIALIZED);
return nullptr;
+ }
// Allocate an extra page for Metadata
if (kMetadataSize_ + (size % page_size_) > page_size_) {
@@ -74,11 +78,18 @@ class DeviceAllocatorT {
"WARNING: %s: DeviceAllocator allocation overflow: "
"0x%zx bytes with 0x%zx alignment requested\n",
SanitizerToolName, map_size, alignment);
+ if (da_info->type_ == DAT_AMDGPU)
+ reinterpret_cast<AmdgpuAllocationInfo*>(da_info)->EnsureFailureStatus(
+ HSA_STATUS_ERROR_OUT_OF_RESOURCES);
return nullptr;
}
void* ptr = DeviceMemFuncs::Allocate(map_size, alignment, da_info);
- if (!ptr)
+ if (!ptr) {
+ if (da_info->type_ == DAT_AMDGPU)
+ reinterpret_cast<AmdgpuAllocationInfo*>(da_info)->EnsureFailureStatus(
+ HSA_STATUS_ERROR_OUT_OF_RESOURCES);
return nullptr;
+ }
uptr map_beg = reinterpret_cast<uptr>(ptr);
CHECK(IsAligned(map_beg, page_size_));
MapUnmapCallback().OnMap(map_beg, map_size);
>From 18f8d2773478fd1259e0c5a4a699033442db5340 Mon Sep 17 00:00:00 2001
From: Amit Pandey <pandey.kumaramit2023 at gmail.com>
Date: Mon, 18 May 2026 13:58:48 +0530
Subject: [PATCH 22/27] [compiler-rt][asan] Fix asan_hsa_amd_pointer_info query
for asan tracked vmem memory.
`hsa_amd_pointer_info()` used a fixed +page shift from the device map
base to the user-visible pointer. That only matches the IPC-attach
layout, not intercepted pool/vmem allocations where the gap can include
alignment padding and ASan redzones (e.g. reserve_align with alignment >
page size).
Query ROCr on the reservation base (map_beg), then adjust reported bases
by the actual offset from the live user chunk. Passthrough to the real
API when the pointer is not tracked or is not the allocated user base.
Add lit tests for pointer_info after 64 KiB-aligned vmem reserve and for
vmem double-free diagnosis.
---
compiler-rt/lib/asan/asan_allocator.cpp | 56 ++++++++-------
...sa_amd_pointer_info_vmem_reserve_align.cpp | 72 +++++++++++++++++++
...vmem_address_reserve_align_double_free.cpp | 47 ++++++++++++
3 files changed, 150 insertions(+), 25 deletions(-)
create mode 100644 compiler-rt/test/asan/TestCases/AMDGPU/hsa_amd_pointer_info_vmem_reserve_align.cpp
create mode 100644 compiler-rt/test/asan/TestCases/AMDGPU/hsa_amd_vmem_address_reserve_align_double_free.cpp
diff --git a/compiler-rt/lib/asan/asan_allocator.cpp b/compiler-rt/lib/asan/asan_allocator.cpp
index 1fbbbb0fded2d..66dcb65890bb3 100644
--- a/compiler-rt/lib/asan/asan_allocator.cpp
+++ b/compiler-rt/lib/asan/asan_allocator.cpp
@@ -1547,9 +1547,6 @@ hsa_status_t asan_hsa_amd_agents_allow_access(uint32_t num_agents,
p ? p : ptr);
}
-// For asan allocator, kMetadataSize is 0 and maximum redzone size is 2048. This
-// implies for device allocation, the gap between user_beg and GetBlockBegin()
-// is always one kPageSize_
// IPC calls use static_assert to make sure kMetadataSize = 0
//
# if SANITIZER_CAN_USE_ALLOCATOR64
@@ -1702,29 +1699,38 @@ hsa_status_t asan_hsa_amd_pointer_info(const void* ptr,
void* (*alloc)(size_t),
uint32_t* num_agents_accessible,
hsa_agent_t** accessible) {
- void* ptr_ = get_allocator().GetBlockBegin(ptr);
- AsanChunk* m = ptr_
- ? instance.GetAsanChunkByAddr(reinterpret_cast<uptr>(ptr_))
- : nullptr;
- if (ptr_ && m) {
- hsa_status_t status = REAL(hsa_amd_pointer_info)(
- ptr_, info, alloc, num_agents_accessible, accessible);
- if (status == HSA_STATUS_SUCCESS && info) {
- static_assert(AP_.kMetadataSize == 0, "Expression below requires this");
- // Adjust base address of agent,host and sizeInBytes so as to return
- // the actual pointer information of user allocation rather than asan
- // allocation. Asan allocation pointer info can be acquired using internal
- // 'GetPointerInfo'
- info->agentBaseAddress = reinterpret_cast<void*>(
- reinterpret_cast<uptr>(info->agentBaseAddress) + kPageSize_);
- info->hostBaseAddress = reinterpret_cast<void*>(
- reinterpret_cast<uptr>(info->hostBaseAddress) + kPageSize_);
- info->sizeInBytes = m->UsedSize();
- }
- return status;
+ // Device/pool mappings are keyed by the ROCr reservation base (map_beg), not
+ // the user pointer returned from intercepted allocate/reserve.
+ void* hsa_map_base = get_allocator().GetBlockBegin(ptr);
+ if (!hsa_map_base)
+ // Not tracked by ASan; query ROCr on the caller's pointer unchanged.
+ return REAL(hsa_amd_pointer_info)(ptr, info, alloc, num_agents_accessible,
+ accessible);
+
+ uptr user = reinterpret_cast<uptr>(ptr);
+ AsanChunk* m = reinterpret_cast<AsanChunk*>(user - kChunkHeaderSize);
+ if (atomic_load(&m->chunk_state, memory_order_acquire) != CHUNK_ALLOCATED ||
+ m->Beg() != user)
+ // Inside a tracked mapping but not the live user base (redzone, interior,
+ // or freed); do not apply the user-visible pointer_info adjustment.
+ return REAL(hsa_amd_pointer_info)(ptr, info, alloc, num_agents_accessible,
+ accessible);
+
+ hsa_status_t status = REAL(hsa_amd_pointer_info)(
+ hsa_map_base, info, alloc, num_agents_accessible, accessible);
+ if (status == HSA_STATUS_SUCCESS && info) {
+ static_assert(AP_.kMetadataSize == 0, "Expression below requires this");
+ // User VA may be above hsa_map_base when VMem API's(i.e
+ // hsa_amd_vmem_address_reserve_align()) uses alignment > page size (device
+ // allocator padding and ASan redzones/headers).
+ const uptr offset = user - reinterpret_cast<uptr>(hsa_map_base);
+ info->agentBaseAddress = reinterpret_cast<void*>(
+ reinterpret_cast<uptr>(info->agentBaseAddress) + offset);
+ info->hostBaseAddress = reinterpret_cast<void*>(
+ reinterpret_cast<uptr>(info->hostBaseAddress) + offset);
+ info->sizeInBytes = m->UsedSize();
}
- return REAL(hsa_amd_pointer_info)(ptr, info, alloc, num_agents_accessible,
- accessible);
+ return status;
}
hsa_status_t asan_hsa_init() {
diff --git a/compiler-rt/test/asan/TestCases/AMDGPU/hsa_amd_pointer_info_vmem_reserve_align.cpp b/compiler-rt/test/asan/TestCases/AMDGPU/hsa_amd_pointer_info_vmem_reserve_align.cpp
new file mode 100644
index 0000000000000..06c09d041dd67
--- /dev/null
+++ b/compiler-rt/test/asan/TestCases/AMDGPU/hsa_amd_pointer_info_vmem_reserve_align.cpp
@@ -0,0 +1,72 @@
+// XFAIL: *
+//
+//
+// RUN: %clangxx_asan -O0 -isystem %rocm_include %s -o %t -L%rocm_lib -lhsa-runtime64 \
+// RUN: -Wl,-rpath,%rocm_lib -Wl,-rpath,%compiler_rt_libdir
+// RUN: %run %t 2>&1 | FileCheck %s
+//
+// hsa_amd_pointer_info on vmem reserved via reserve_align() must report the
+// user-visible base and size, not the internal HSA backing mapping. A fixed
+// +page offset is wrong when alignment is larger than the page size.
+//
+// REQUIRES: sanitizer-amdgpu, linux, stable-runtime, rocm
+// UNSUPPORTED: android
+
+#include "hsa_amd_test_helpers.h"
+
+#include <hsa/hsa.h>
+#include <hsa/hsa_ext_amd.h>
+
+#include <inttypes.h>
+#include <stdint.h>
+#include <stdio.h>
+
+int main() {
+ if (hsa_amd_test_require_init())
+ return 1;
+
+ const size_t kSize = 4096;
+ const uint64_t kAlign = 65536;
+ void *mem = nullptr;
+ if (hsa_amd_vmem_address_reserve_align(&mem, kSize, /*address=*/0, kAlign,
+ /*flags=*/0) != HSA_STATUS_SUCCESS ||
+ !mem) {
+ fprintf(stderr, "hsa_amd_vmem_address_reserve_align failed\n");
+ return 1;
+ }
+
+ const uintptr_t user = reinterpret_cast<uintptr_t>(mem);
+ if (user % kAlign != 0) {
+ fprintf(stderr, "reserved address not %" PRIu64 "-byte aligned\n", kAlign);
+ return 1;
+ }
+
+ hsa_amd_pointer_info_t info = {};
+ info.size = sizeof(hsa_amd_pointer_info_t);
+ if (hsa_amd_pointer_info(mem, &info, nullptr, nullptr, nullptr) !=
+ HSA_STATUS_SUCCESS) {
+ fprintf(stderr, "hsa_amd_pointer_info failed\n");
+ return 1;
+ }
+
+ if (info.agentBaseAddress != mem ||
+ reinterpret_cast<uintptr_t>(info.agentBaseAddress) + info.sizeInBytes !=
+ user + kSize) {
+ fprintf(stderr,
+ "pointer_info mismatch: user=%p begin=%p size=%zu (expected %zu)\n",
+ mem, info.agentBaseAddress, info.sizeInBytes, kSize);
+ return 1;
+ }
+
+ printf("pointer_info_vmem sizeInBytes: %zu\n", info.sizeInBytes);
+ printf("pointer_info_vmem begin: %p\n", info.agentBaseAddress);
+
+ if (hsa_amd_vmem_address_free(mem, kSize) != HSA_STATUS_SUCCESS) {
+ fprintf(stderr, "hsa_amd_vmem_address_free failed\n");
+ return 1;
+ }
+ return 0;
+}
+
+// CHECK: pointer_info_vmem sizeInBytes: 4096
+// CHECK-NEXT: pointer_info_vmem begin: 0x{{[0-9a-f]+}}
diff --git a/compiler-rt/test/asan/TestCases/AMDGPU/hsa_amd_vmem_address_reserve_align_double_free.cpp b/compiler-rt/test/asan/TestCases/AMDGPU/hsa_amd_vmem_address_reserve_align_double_free.cpp
new file mode 100644
index 0000000000000..038fec02f7be3
--- /dev/null
+++ b/compiler-rt/test/asan/TestCases/AMDGPU/hsa_amd_vmem_address_reserve_align_double_free.cpp
@@ -0,0 +1,47 @@
+// XFAIL: *
+//
+//
+// RUN: %clangxx_asan -O0 -isystem %rocm_include %s -o %t -L%rocm_lib -lhsa-runtime64 \
+// RUN: -Wl,-rpath,%rocm_lib -Wl,-rpath,%compiler_rt_libdir
+// RUN: not %run %t 2>&1 | FileCheck %s
+//
+// Regression test for the AddressSanitizer hsa_amd_vmem_address_reserve_align /
+// hsa_amd_vmem_address_free interceptors: freeing the same reserved range twice
+// is diagnosed.
+//
+// REQUIRES: sanitizer-amdgpu, linux, stable-runtime, rocm
+// UNSUPPORTED: android
+
+#include "hsa_amd_test_helpers.h"
+
+#include <hsa/hsa.h>
+#include <hsa/hsa_ext_amd.h>
+
+#include <stdio.h>
+
+int main() {
+ if (hsa_amd_test_require_init())
+ return 1;
+
+ // Size must be a non-zero multiple of the page size; address must be 0 so
+ // the interceptor records the reservation for double-free diagnosis (see
+ // asan_hsa_amd_vmem_address_reserve_align).
+ const size_t kSize = 4096;
+ void *mem = nullptr;
+ if (hsa_amd_vmem_address_reserve_align(&mem, kSize, /*address=*/0,
+ /*alignment=*/4096,
+ /*flags=*/0) != HSA_STATUS_SUCCESS ||
+ !mem) {
+ fprintf(stderr, "hsa_amd_vmem_address_reserve_align failed\n");
+ return 1;
+ }
+
+ (void)hsa_amd_vmem_address_free(mem, kSize);
+ (void)hsa_amd_vmem_address_free(mem, kSize);
+
+ fprintf(stderr, "expected double-free report\n");
+ return 0;
+}
+
+// CHECK: ERROR: AddressSanitizer: attempting double-free
+// CHECK: SUMMARY: AddressSanitizer: double-free
>From ee2097824229e29b3fa1e4f4b987f7f4f43874c1 Mon Sep 17 00:00:00 2001
From: Amit Pandey <pandey.kumaramit2023 at gmail.com>
Date: Mon, 18 May 2026 16:34:06 +0530
Subject: [PATCH 23/27] [compiler-rt] Move device allocator HSA failure status
into AMDGPU layer
DeviceAllocatorT failure paths in sanitizer_allocator_device.h recorded
HSA_STATUS_* via AmdgpuAllocationInfo casts, which tied the generic
device allocator header to ROCr/HSA types.
Add AmdgpuMemFuncs::NoteDeviceAllocatorFailure() and call it from the
device allocator on init failure, size overflow, and null allocation.
Map failure kinds to hsa_status_t in sanitizer_allocator_amdgpu.cpp
only.
---
.../sanitizer_allocator_amdgpu.cpp | 16 ++++++++++++++++
.../sanitizer_allocator_amdgpu.h | 9 +++++++++
.../sanitizer_allocator_device.h | 16 +++++++---------
3 files changed, 32 insertions(+), 9 deletions(-)
diff --git a/compiler-rt/lib/sanitizer_common/sanitizer_allocator_amdgpu.cpp b/compiler-rt/lib/sanitizer_common/sanitizer_allocator_amdgpu.cpp
index 87db693476abd..f1aab70cf189b 100644
--- a/compiler-rt/lib/sanitizer_common/sanitizer_allocator_amdgpu.cpp
+++ b/compiler-rt/lib/sanitizer_common/sanitizer_allocator_amdgpu.cpp
@@ -75,6 +75,22 @@ void AmdgpuMemFuncs::ClearAmdgpuRuntimeShutdownState() {
atomic_store(&amdgpu_event_registered, 0, memory_order_release);
}
+void AmdgpuMemFuncs::NoteDeviceAllocatorFailure(DeviceAllocationInfo* da_info,
+ DeviceAllocFailure failure) {
+ if (!da_info || da_info->type_ != DAT_AMDGPU)
+ return;
+ AmdgpuAllocationInfo* aa_info =
+ reinterpret_cast<AmdgpuAllocationInfo*>(da_info);
+ switch (failure) {
+ case kNotInitialized:
+ aa_info->EnsureFailureStatus(HSA_STATUS_ERROR_NOT_INITIALIZED);
+ break;
+ case kOutOfResources:
+ aa_info->EnsureFailureStatus(HSA_STATUS_ERROR_OUT_OF_RESOURCES);
+ break;
+ }
+}
+
bool AmdgpuMemFuncs::Init() {
bool success = true;
LOAD_HSA_FUNC_WITH_ERROR_CHECK(hsa_amd.memory_pool_allocate,
diff --git a/compiler-rt/lib/sanitizer_common/sanitizer_allocator_amdgpu.h b/compiler-rt/lib/sanitizer_common/sanitizer_allocator_amdgpu.h
index 0f11b7a343d0a..884b6b30d7459 100644
--- a/compiler-rt/lib/sanitizer_common/sanitizer_allocator_amdgpu.h
+++ b/compiler-rt/lib/sanitizer_common/sanitizer_allocator_amdgpu.h
@@ -16,6 +16,11 @@
#if SANITIZER_AMDGPU
class AmdgpuMemFuncs {
public:
+ enum DeviceAllocFailure {
+ kNotInitialized,
+ kOutOfResources,
+ };
+
static bool Init();
static void* Allocate(uptr size, uptr alignment,
DeviceAllocationInfo* da_info);
@@ -25,6 +30,10 @@ class AmdgpuMemFuncs {
static void RegisterSystemEventHandlers();
static bool IsAmdgpuRuntimeShutdown();
static void ClearAmdgpuRuntimeShutdownState();
+ // Record an HSA error on da_info when DeviceAllocatorT fails before/without a
+ // successful AmdgpuMemFuncs::Allocate (keeps HSA types out of device.h).
+ static void NoteDeviceAllocatorFailure(DeviceAllocationInfo* da_info,
+ DeviceAllocFailure failure);
private:
static void NotifyAmdgpuRuntimeShutdown();
diff --git a/compiler-rt/lib/sanitizer_common/sanitizer_allocator_device.h b/compiler-rt/lib/sanitizer_common/sanitizer_allocator_device.h
index eb29ed7bde0f1..d1a0981e32fa1 100644
--- a/compiler-rt/lib/sanitizer_common/sanitizer_allocator_device.h
+++ b/compiler-rt/lib/sanitizer_common/sanitizer_allocator_device.h
@@ -58,9 +58,9 @@ class DeviceAllocatorT {
void* Allocate(AllocatorStats* stat, uptr size, uptr alignment,
DeviceAllocationInfo* da_info) {
if (!da_info || !InitMemFuncs()) {
- if (da_info && da_info->type_ == DAT_AMDGPU)
- reinterpret_cast<AmdgpuAllocationInfo*>(da_info)->EnsureFailureStatus(
- HSA_STATUS_ERROR_NOT_INITIALIZED);
+ if (da_info)
+ DeviceMemFuncs::NoteDeviceAllocatorFailure(
+ da_info, DeviceMemFuncs::kNotInitialized);
return nullptr;
}
@@ -78,16 +78,14 @@ class DeviceAllocatorT {
"WARNING: %s: DeviceAllocator allocation overflow: "
"0x%zx bytes with 0x%zx alignment requested\n",
SanitizerToolName, map_size, alignment);
- if (da_info->type_ == DAT_AMDGPU)
- reinterpret_cast<AmdgpuAllocationInfo*>(da_info)->EnsureFailureStatus(
- HSA_STATUS_ERROR_OUT_OF_RESOURCES);
+ DeviceMemFuncs::NoteDeviceAllocatorFailure(
+ da_info, DeviceMemFuncs::kOutOfResources);
return nullptr;
}
void* ptr = DeviceMemFuncs::Allocate(map_size, alignment, da_info);
if (!ptr) {
- if (da_info->type_ == DAT_AMDGPU)
- reinterpret_cast<AmdgpuAllocationInfo*>(da_info)->EnsureFailureStatus(
- HSA_STATUS_ERROR_OUT_OF_RESOURCES);
+ DeviceMemFuncs::NoteDeviceAllocatorFailure(
+ da_info, DeviceMemFuncs::kOutOfResources);
return nullptr;
}
uptr map_beg = reinterpret_cast<uptr>(ptr);
>From 728d434d668e777de09ea2ed73e032dc6ca0d586 Mon Sep 17 00:00:00 2001
From: Amit Pandey <pandey.kumaramit2023 at gmail.com>
Date: Thu, 21 May 2026 12:16:51 +0530
Subject: [PATCH 24/27] [compiler-rt] Decouple DeviceAllocatorT from AMDGPU/HSA
backend.
Make DeviceAllocatorT vendor-neutral with a pluggable DeviceAllocator
backend (NoOpDeviceAllocator / AmdgpuDeviceAllocator). Always wire the
device layer through CombinedAllocator; keep HSA types and ROCr dlsym in
the AMDGPU backend only.
---
compiler-rt/lib/asan/asan_allocator.cpp | 21 ++---
.../sanitizer_common/sanitizer_allocator.h | 3 +
.../sanitizer_allocator_amdgpu.cpp | 42 ++++-----
.../sanitizer_allocator_amdgpu.h | 28 +++---
.../sanitizer_allocator_combined.h | 47 +---------
.../sanitizer_allocator_device.h | 94 +++++++++++++------
6 files changed, 117 insertions(+), 118 deletions(-)
diff --git a/compiler-rt/lib/asan/asan_allocator.cpp b/compiler-rt/lib/asan/asan_allocator.cpp
index 66dcb65890bb3..24b22035a10fb 100644
--- a/compiler-rt/lib/asan/asan_allocator.cpp
+++ b/compiler-rt/lib/asan/asan_allocator.cpp
@@ -393,14 +393,11 @@ struct Allocator {
void InitLinkerInitialized(const AllocatorOptions& options) {
SetAllocatorMayReturnNull(options.may_return_null);
-#if SANITIZER_AMDGPU
- // Device-backed HSA allocations (e.g. hsa_amd_memory_pool_allocate) use
- // CombinedAllocator's device path; it must be enabled for InitMemFuncs/
- // AmdgpuMemFuncs::Init to run at startup.
- allocator.InitLinkerInitialized(options.release_to_os_interval_ms, 0, true);
-#else
- allocator.InitLinkerInitialized(options.release_to_os_interval_ms);
-#endif
+ // Device-backed allocations use CombinedAllocator's device path. On
+ // SANITIZER_AMDGPU builds, enable it so InitMemFuncs() /
+ // AmdgpuDeviceAllocator::Init() run at startup.
+ allocator.InitLinkerInitialized(options.release_to_os_interval_ms, 0,
+ SANITIZER_AMDGPU);
SharedInitCode(options);
max_user_defined_malloc_size = common_flags()->max_allocation_size_mb
? common_flags()->max_allocation_size_mb
@@ -1738,12 +1735,12 @@ hsa_status_t asan_hsa_init() {
if (status == HSA_STATUS_SUCCESS) {
// Only clear state when recovering from a prior shutdown (avoids clearing
// amdgpu_event_registered on every refcount bump and re-registering).
- if (__sanitizer::AmdgpuMemFuncs::IsAmdgpuRuntimeShutdown())
- __sanitizer::AmdgpuMemFuncs::ClearAmdgpuRuntimeShutdownState();
+ if (__sanitizer::AmdgpuDeviceAllocator::IsRuntimeShutdown())
+ __sanitizer::AmdgpuDeviceAllocator::ClearRuntimeShutdownState();
// Load HSA entry points once the runtime is up; device allocator may stay
// disabled, but interceptors and RegisterSystemEventHandlers need them.
- if (__sanitizer::AmdgpuMemFuncs::Init())
- __sanitizer::AmdgpuMemFuncs::RegisterSystemEventHandlers();
+ if (__sanitizer::AmdgpuDeviceAllocator::Init())
+ __sanitizer::AmdgpuDeviceAllocator::RegisterSystemEventHandlers();
}
return status;
}
diff --git a/compiler-rt/lib/sanitizer_common/sanitizer_allocator.h b/compiler-rt/lib/sanitizer_common/sanitizer_allocator.h
index d803444f69dd6..726ef132d0e40 100644
--- a/compiler-rt/lib/sanitizer_common/sanitizer_allocator.h
+++ b/compiler-rt/lib/sanitizer_common/sanitizer_allocator.h
@@ -88,6 +88,9 @@ struct NoOpMapUnmapCallback {
#include "sanitizer_allocator_local_cache.h"
#include "sanitizer_allocator_secondary.h"
#include "sanitizer_allocator_device.h"
+#if SANITIZER_AMDGPU
+# include "sanitizer_allocator_amdgpu.h"
+#endif
#include "sanitizer_allocator_combined.h"
// clang-format on
diff --git a/compiler-rt/lib/sanitizer_common/sanitizer_allocator_amdgpu.cpp b/compiler-rt/lib/sanitizer_common/sanitizer_allocator_amdgpu.cpp
index f1aab70cf189b..461ad033cac9f 100644
--- a/compiler-rt/lib/sanitizer_common/sanitizer_allocator_amdgpu.cpp
+++ b/compiler-rt/lib/sanitizer_common/sanitizer_allocator_amdgpu.cpp
@@ -51,14 +51,12 @@ static atomic_uint8_t amdgpu_event_registered{0};
success = false; \
}
-// Check AMDGPU runtime shutdown state
-bool AmdgpuMemFuncs::IsAmdgpuRuntimeShutdown() {
+bool AmdgpuDeviceAllocator::IsRuntimeShutdown() {
return static_cast<bool>(
atomic_load(&amdgpu_runtime_shutdown, memory_order_acquire));
}
-// Notify AMDGPU runtime shutdown to allocator
-void AmdgpuMemFuncs::NotifyAmdgpuRuntimeShutdown() {
+void AmdgpuDeviceAllocator::NotifyRuntimeShutdown() {
uint8_t shutdown = 0;
if (atomic_compare_exchange_strong(&amdgpu_runtime_shutdown, &shutdown, 1,
memory_order_acq_rel)) {
@@ -70,28 +68,28 @@ void AmdgpuMemFuncs::NotifyAmdgpuRuntimeShutdown() {
// Resets amdgpu_runtime_shutdown so allocator operations are enabled, and
// amdgpu_event_registered so RegisterSystemEventHandlers() will register the
// shutdown callback for the new runtime instance.
-void AmdgpuMemFuncs::ClearAmdgpuRuntimeShutdownState() {
+void AmdgpuDeviceAllocator::ClearRuntimeShutdownState() {
atomic_store(&amdgpu_runtime_shutdown, 0, memory_order_release);
atomic_store(&amdgpu_event_registered, 0, memory_order_release);
}
-void AmdgpuMemFuncs::NoteDeviceAllocatorFailure(DeviceAllocationInfo* da_info,
- DeviceAllocFailure failure) {
+void AmdgpuDeviceAllocator::NoteDeviceAllocatorFailure(
+ DeviceAllocationInfo* da_info, DeviceAllocFailure failure) {
if (!da_info || da_info->type_ != DAT_AMDGPU)
return;
AmdgpuAllocationInfo* aa_info =
reinterpret_cast<AmdgpuAllocationInfo*>(da_info);
switch (failure) {
- case kNotInitialized:
+ case DEV_ALLOC_FAILURE_NOT_INITIALIZED:
aa_info->EnsureFailureStatus(HSA_STATUS_ERROR_NOT_INITIALIZED);
break;
- case kOutOfResources:
+ case DEV_ALLOC_FAILURE_OUT_OF_RESOURCES:
aa_info->EnsureFailureStatus(HSA_STATUS_ERROR_OUT_OF_RESOURCES);
break;
}
}
-bool AmdgpuMemFuncs::Init() {
+bool AmdgpuDeviceAllocator::Init() {
bool success = true;
LOAD_HSA_FUNC_WITH_ERROR_CHECK(hsa_amd.memory_pool_allocate,
"hsa_amd_memory_pool_allocate", success);
@@ -113,13 +111,13 @@ bool AmdgpuMemFuncs::Init() {
return true;
}
-void* AmdgpuMemFuncs::Allocate(uptr size, uptr alignment,
- DeviceAllocationInfo* da_info) {
+void* AmdgpuDeviceAllocator::Allocate(uptr size, uptr alignment,
+ DeviceAllocationInfo* da_info) {
AmdgpuAllocationInfo* aa_info =
reinterpret_cast<AmdgpuAllocationInfo*>(da_info);
// Do not allocate if AMDGPU runtime is shutdown
- if (UNLIKELY(IsAmdgpuRuntimeShutdown())) {
+ if (UNLIKELY(IsRuntimeShutdown())) {
VReport(1,
"Amdgpu Allocate: Runtime shutdown, skipping allocation for size "
"%zu alignment %zu\n",
@@ -142,9 +140,9 @@ void* AmdgpuMemFuncs::Allocate(uptr size, uptr alignment,
return aa_info->ptr;
}
-void AmdgpuMemFuncs::Deallocate(void* p) {
+void AmdgpuDeviceAllocator::Deallocate(void* p) {
// Deallocate does nothing after AMDGPU runtime shutdown
- if (UNLIKELY(IsAmdgpuRuntimeShutdown())) {
+ if (UNLIKELY(IsRuntimeShutdown())) {
VReport(
1,
"Amdgpu Deallocate: Runtime shutdown, skipping deallocation for %p\n",
@@ -153,7 +151,8 @@ void AmdgpuMemFuncs::Deallocate(void* p) {
}
DevicePointerInfo DevPtrInfo;
- if (AmdgpuMemFuncs::GetPointerInfo(reinterpret_cast<uptr>(p), &DevPtrInfo)) {
+ if (AmdgpuDeviceAllocator::GetPointerInfo(reinterpret_cast<uptr>(p),
+ &DevPtrInfo)) {
if (DevPtrInfo.type == HSA_EXT_POINTER_TYPE_HSA) {
UNUSED hsa_status_t status = hsa_amd.memory_pool_free(p);
} else if (DevPtrInfo.type == HSA_EXT_POINTER_TYPE_RESERVED_ADDR) {
@@ -163,9 +162,10 @@ void AmdgpuMemFuncs::Deallocate(void* p) {
}
}
-bool AmdgpuMemFuncs::GetPointerInfo(uptr ptr, DevicePointerInfo* ptr_info) {
+bool AmdgpuDeviceAllocator::GetPointerInfo(uptr ptr,
+ DevicePointerInfo* ptr_info) {
// GetPointerInfo returns false after AMDGPU runtime shutdown
- if (UNLIKELY(IsAmdgpuRuntimeShutdown())) {
+ if (UNLIKELY(IsRuntimeShutdown())) {
VReport(1,
"Amdgpu GetPointerInfo: Runtime shutdown, skipping query for %p\n",
reinterpret_cast<void*>(ptr));
@@ -191,7 +191,7 @@ bool AmdgpuMemFuncs::GetPointerInfo(uptr ptr, DevicePointerInfo* ptr_info) {
}
// Register shutdown system event handler only once
// TODO: Register multiple event handlers if needed in future
-void AmdgpuMemFuncs::RegisterSystemEventHandlers() {
+void AmdgpuDeviceAllocator::RegisterSystemEventHandlers() {
uint8_t registered = 0;
// Check if shutdown event handler is already registered
if (atomic_compare_exchange_strong(&amdgpu_event_registered, ®istered, 1,
@@ -202,7 +202,7 @@ void AmdgpuMemFuncs::RegisterSystemEventHandlers() {
if (!event)
return HSA_STATUS_ERROR_INVALID_ARGUMENT;
if (event->event_type == HSA_AMD_SYSTEM_SHUTDOWN_EVENT)
- AmdgpuMemFuncs::NotifyAmdgpuRuntimeShutdown();
+ AmdgpuDeviceAllocator::NotifyRuntimeShutdown();
return HSA_STATUS_SUCCESS;
};
// Register the event callback
@@ -222,6 +222,6 @@ void AmdgpuMemFuncs::RegisterSystemEventHandlers() {
}
}
-uptr AmdgpuMemFuncs::GetPageSize() { return kPageSize_; }
+uptr AmdgpuDeviceAllocator::GetPageSize() { return kPageSize_; }
} // namespace __sanitizer
#endif // SANITIZER_AMDGPU
diff --git a/compiler-rt/lib/sanitizer_common/sanitizer_allocator_amdgpu.h b/compiler-rt/lib/sanitizer_common/sanitizer_allocator_amdgpu.h
index 884b6b30d7459..efc6c8f869e3b 100644
--- a/compiler-rt/lib/sanitizer_common/sanitizer_allocator_amdgpu.h
+++ b/compiler-rt/lib/sanitizer_common/sanitizer_allocator_amdgpu.h
@@ -6,21 +6,20 @@
//
//===----------------------------------------------------------------------===//
//
-// Part of the Sanitizer Allocator.
+// AMDGPU / HSA device memory backend for DeviceAllocatorT.
//
//===----------------------------------------------------------------------===//
#ifndef SANITIZER_ALLOCATOR_H
-# error This file must be included inside sanitizer_allocator_device.h
+# error This file must be included inside sanitizer_allocator.h
#endif
#if SANITIZER_AMDGPU
-class AmdgpuMemFuncs {
- public:
- enum DeviceAllocFailure {
- kNotInitialized,
- kOutOfResources,
- };
+static const DeviceAllocationType DAT_AMDGPU =
+ static_cast<DeviceAllocationType>(1);
+
+class AmdgpuDeviceAllocator {
+ public:
static bool Init();
static void* Allocate(uptr size, uptr alignment,
DeviceAllocationInfo* da_info);
@@ -28,15 +27,13 @@ class AmdgpuMemFuncs {
static bool GetPointerInfo(uptr ptr, DevicePointerInfo* ptr_info);
static uptr GetPageSize();
static void RegisterSystemEventHandlers();
- static bool IsAmdgpuRuntimeShutdown();
- static void ClearAmdgpuRuntimeShutdownState();
- // Record an HSA error on da_info when DeviceAllocatorT fails before/without a
- // successful AmdgpuMemFuncs::Allocate (keeps HSA types out of device.h).
+ static bool IsRuntimeShutdown();
+ static void ClearRuntimeShutdownState();
static void NoteDeviceAllocatorFailure(DeviceAllocationInfo* da_info,
DeviceAllocFailure failure);
private:
- static void NotifyAmdgpuRuntimeShutdown();
+ static void NotifyRuntimeShutdown();
};
struct AmdgpuAllocationInfo : public DeviceAllocationInfo {
@@ -60,4 +57,9 @@ struct AmdgpuAllocationInfo : public DeviceAllocationInfo {
u32 flags;
void* ptr;
};
+
+template <class MapUnmapCallback>
+using DefaultDeviceAllocator =
+ DeviceAllocatorT<MapUnmapCallback, AmdgpuDeviceAllocator>;
+
#endif // SANITIZER_AMDGPU
diff --git a/compiler-rt/lib/sanitizer_common/sanitizer_allocator_combined.h b/compiler-rt/lib/sanitizer_common/sanitizer_allocator_combined.h
index 1397959b66679..fc2a2e5b50c8f 100644
--- a/compiler-rt/lib/sanitizer_common/sanitizer_allocator_combined.h
+++ b/compiler-rt/lib/sanitizer_common/sanitizer_allocator_combined.h
@@ -28,19 +28,14 @@ class CombinedAllocator {
LargeMmapAllocator<typename PrimaryAllocator::MapUnmapCallback,
LargeMmapAllocatorPtrArray,
typename PrimaryAllocator::AddressSpaceView>;
-
-#if SANITIZER_AMDGPU
using DeviceAllocator =
- DeviceAllocatorT<typename PrimaryAllocator::MapUnmapCallback>;
-#endif
+ DefaultDeviceAllocator<typename PrimaryAllocator::MapUnmapCallback>;
void InitLinkerInitialized(s32 release_to_os_interval_ms, uptr heap_start = 0,
bool enable_device_allocator = false) {
primary_.Init(release_to_os_interval_ms, heap_start);
secondary_.InitLinkerInitialized();
-#if SANITIZER_AMDGPU
device_.Init(enable_device_allocator, primary_.kMetadataSize);
-#endif
}
void Init(s32 release_to_os_interval_ms, uptr heap_start = 0,
@@ -48,9 +43,7 @@ class CombinedAllocator {
stats_.Init();
primary_.Init(release_to_os_interval_ms, heap_start);
secondary_.Init();
-#if SANITIZER_AMDGPU
device_.Init(enable_device_allocator, primary_.kMetadataSize);
-#endif
}
void* Allocate(AllocatorCache* cache, uptr size, uptr alignment,
@@ -77,13 +70,10 @@ class CombinedAllocator {
// using a non-fixed base address). The secondary takes care of the
// alignment without such requirement, and allocating 'size' would use
// extraneous memory, so we employ 'original_size'.
- void *res;
-#if SANITIZER_AMDGPU
+ void* res;
if (da_info)
res = device_.Allocate(&stats_, original_size, alignment, da_info);
- else
-#endif
- if (primary_.CanAllocate(size, alignment))
+ else if (primary_.CanAllocate(size, alignment))
res = cache->Allocate(&primary_, primary_.ClassID(size));
else
res = secondary_.Allocate(&stats_, original_size, alignment);
@@ -110,10 +100,8 @@ class CombinedAllocator {
cache->Deallocate(&primary_, primary_.GetSizeClass(p), p);
else if (secondary_.PointerIsMine(p))
secondary_.Deallocate(&stats_, p);
-#if SANITIZER_AMDGPU
else if (device_.PointerIsMine(p))
device_.Deallocate(&stats_, p);
-#endif
}
void *Reallocate(AllocatorCache *cache, void *p, uptr new_size,
@@ -139,10 +127,8 @@ class CombinedAllocator {
return true;
if (secondary_.PointerIsMine(p))
return true;
-#if SANITIZER_AMDGPU
if (device_.PointerIsMine(p))
return true;
-#endif
return false;
}
@@ -153,10 +139,8 @@ class CombinedAllocator {
return primary_.GetMetaData(p);
if (secondary_.PointerIsMine(p))
return secondary_.GetMetaData(p);
-#if SANITIZER_AMDGPU
if (device_.PointerIsMine(p))
return device_.GetMetaData(p);
-#endif
return nullptr;
}
@@ -168,10 +152,8 @@ class CombinedAllocator {
return primary_.GetMetaData(p);
if (secondary_.GetBlockBeginFastLocked(p))
return secondary_.GetMetaData(p);
-#if SANITIZER_AMDGPU
if (device_.GetBlockBeginFastLocked(p))
return device_.GetMetaData(p);
-#endif
return nullptr;
}
@@ -180,10 +162,8 @@ class CombinedAllocator {
return primary_.GetBlockBegin(p);
if (secondary_.PointerIsMine(p))
return secondary_.GetBlockBegin(p);
-#if SANITIZER_AMDGPU
if (device_.PointerIsMine(p))
return device_.GetBlockBegin(p);
-#endif
return nullptr;
}
@@ -195,10 +175,8 @@ class CombinedAllocator {
return primary_.GetBlockBegin(p);
if ((beg = secondary_.GetBlockBeginFastLocked(p)))
return beg;
-#if SANITIZER_AMDGPU
if ((beg = device_.GetBlockBeginFastLocked(p)))
return beg;
-#endif
return nullptr;
}
@@ -207,19 +185,14 @@ class CombinedAllocator {
return primary_.GetActuallyAllocatedSize(p);
if (secondary_.PointerIsMine(p))
return secondary_.GetActuallyAllocatedSize(p);
-#if SANITIZER_AMDGPU
if (device_.PointerIsMine(p))
return device_.GetActuallyAllocatedSize(p);
-#endif
return 0;
}
uptr TotalMemoryUsed() {
- return primary_.TotalMemoryUsed() + secondary_.TotalMemoryUsed()
-#if SANITIZER_AMDGPU
- + device_.TotalMemoryUsed()
-#endif
- ;
+ return primary_.TotalMemoryUsed() + secondary_.TotalMemoryUsed() +
+ device_.TotalMemoryUsed();
}
void TestOnlyUnmap() { primary_.TestOnlyUnmap(); }
@@ -243,17 +216,13 @@ class CombinedAllocator {
void PrintStats() {
primary_.PrintStats();
secondary_.PrintStats();
-#if SANITIZER_AMDGPU
device_.PrintStats();
-#endif
}
// ForceLock() and ForceUnlock() are needed to implement Darwin malloc zone
// introspection API.
void ForceLock() SANITIZER_NO_THREAD_SAFETY_ANALYSIS {
-#if SANITIZER_AMDGPU
device_.ForceLock();
-#endif
primary_.ForceLock();
secondary_.ForceLock();
}
@@ -261,9 +230,7 @@ class CombinedAllocator {
void ForceUnlock() SANITIZER_NO_THREAD_SAFETY_ANALYSIS {
secondary_.ForceUnlock();
primary_.ForceUnlock();
-#if SANITIZER_AMDGPU
device_.ForceUnlock();
-#endif
}
// Iterate over all existing chunks.
@@ -271,16 +238,12 @@ class CombinedAllocator {
void ForEachChunk(ForEachChunkCallback callback, void *arg) {
primary_.ForEachChunk(callback, arg);
secondary_.ForEachChunk(callback, arg);
-#if SANITIZER_AMDGPU
device_.ForEachChunk(callback, arg);
-#endif
}
private:
PrimaryAllocator primary_;
SecondaryAllocator secondary_;
-#if SANITIZER_AMDGPU
DeviceAllocator device_;
-#endif
AllocatorGlobalStats stats_;
};
diff --git a/compiler-rt/lib/sanitizer_common/sanitizer_allocator_device.h b/compiler-rt/lib/sanitizer_common/sanitizer_allocator_device.h
index d1a0981e32fa1..00d8a5385ce8d 100644
--- a/compiler-rt/lib/sanitizer_common/sanitizer_allocator_device.h
+++ b/compiler-rt/lib/sanitizer_common/sanitizer_allocator_device.h
@@ -13,21 +13,22 @@
# error This file must be included inside sanitizer_allocator.h
#endif
-struct DeviceAllocationInfo;
-#if SANITIZER_AMDGPU
-// Device memory allocation usually requires additional information, we can put
-// all the additional information into a data structure DeviceAllocationInfo.
-// This is only a parent structure since different vendors may require
-// different allocation info.
-typedef enum {
+// Vendor-neutral device heap book-keeping. Runtime access (allocate, pointer
+// queries, deallocate etc.) is provided by a DeviceAllocator template
+// parameter.
+
+enum DeviceAllocationType {
DAT_UNKNOWN = 0,
- DAT_AMDGPU = 1,
-} DeviceAllocationType;
+};
+
+enum DeviceAllocFailure {
+ DEV_ALLOC_FAILURE_NOT_INITIALIZED,
+ DEV_ALLOC_FAILURE_OUT_OF_RESOURCES,
+};
struct DeviceAllocationInfo {
- DeviceAllocationInfo(DeviceAllocationType type = DAT_UNKNOWN) {
- type_ = type;
- }
+ explicit DeviceAllocationInfo(DeviceAllocationType type = DAT_UNKNOWN)
+ : type_(type) {}
DeviceAllocationType type_;
};
@@ -37,13 +38,41 @@ struct DevicePointerInfo {
uptr map_size;
};
-# include "sanitizer_allocator_amdgpu.h"
+struct NoOpDeviceAllocator {
+ static bool Init() { return false; }
-template <class MapUnmapCallback = NoOpMapUnmapCallback>
+ static void* Allocate(uptr size, uptr alignment,
+ DeviceAllocationInfo* da_info) {
+ (void)size;
+ (void)alignment;
+ (void)da_info;
+ return nullptr;
+ }
+
+ static void Deallocate(void* p) { (void)p; }
+
+ static bool GetPointerInfo(uptr ptr, DevicePointerInfo* ptr_info) {
+ (void)ptr;
+ (void)ptr_info;
+ return false;
+ }
+
+ static uptr GetPageSize() { return 4096; }
+
+ static bool IsRuntimeShutdown() { return false; }
+
+ static void NoteDeviceAllocatorFailure(DeviceAllocationInfo* da_info,
+ DeviceAllocFailure failure) {
+ (void)da_info;
+ (void)failure;
+ }
+};
+
+template <class MapUnmapCallback = NoOpMapUnmapCallback,
+ class DeviceAllocator = NoOpDeviceAllocator>
class DeviceAllocatorT {
public:
using PtrArrayT = DefaultLargeMmapAllocatorPtrArray;
- using DeviceMemFuncs = AmdgpuMemFuncs;
void Init(bool enable, uptr kMetadataSize) {
internal_memset(this, 0, sizeof(*this));
@@ -59,8 +88,8 @@ class DeviceAllocatorT {
DeviceAllocationInfo* da_info) {
if (!da_info || !InitMemFuncs()) {
if (da_info)
- DeviceMemFuncs::NoteDeviceAllocatorFailure(
- da_info, DeviceMemFuncs::kNotInitialized);
+ DeviceAllocator::NoteDeviceAllocatorFailure(
+ da_info, DEV_ALLOC_FAILURE_NOT_INITIALIZED);
return nullptr;
}
@@ -78,14 +107,14 @@ class DeviceAllocatorT {
"WARNING: %s: DeviceAllocator allocation overflow: "
"0x%zx bytes with 0x%zx alignment requested\n",
SanitizerToolName, map_size, alignment);
- DeviceMemFuncs::NoteDeviceAllocatorFailure(
- da_info, DeviceMemFuncs::kOutOfResources);
+ DeviceAllocator::NoteDeviceAllocatorFailure(
+ da_info, DEV_ALLOC_FAILURE_OUT_OF_RESOURCES);
return nullptr;
}
- void* ptr = DeviceMemFuncs::Allocate(map_size, alignment, da_info);
+ void* ptr = DeviceAllocator::Allocate(map_size, alignment, da_info);
if (!ptr) {
- DeviceMemFuncs::NoteDeviceAllocatorFailure(
- da_info, DeviceMemFuncs::kOutOfResources);
+ DeviceAllocator::NoteDeviceAllocatorFailure(
+ da_info, DEV_ALLOC_FAILURE_OUT_OF_RESOURCES);
return nullptr;
}
uptr map_beg = reinterpret_cast<uptr>(ptr);
@@ -139,7 +168,7 @@ class DeviceAllocatorT {
stat->Sub(AllocatorStatMapped, h->map_size);
}
MapUnmapCallback().OnUnmap(h->map_beg, h->map_size);
- DeviceMemFuncs::Deallocate(p);
+ DeviceAllocator::Deallocate(p);
}
uptr TotalMemoryUsed() {
@@ -302,10 +331,10 @@ class DeviceAllocatorT {
if (!enabled_ || mem_funcs_inited_ || mem_funcs_init_count_ >= 2) {
return mem_funcs_inited_;
}
- mem_funcs_inited_ = DeviceMemFuncs::Init();
+ mem_funcs_inited_ = DeviceAllocator::Init();
mem_funcs_init_count_++;
if (mem_funcs_inited_)
- page_size_ = DeviceMemFuncs::GetPageSize();
+ page_size_ = DeviceAllocator::GetPageSize();
return mem_funcs_inited_;
}
@@ -313,7 +342,7 @@ class DeviceAllocatorT {
Header* GetHeaderAnyPointer(uptr p, Header* h) const {
CHECK(IsAligned(p, page_size_));
- return DeviceMemFuncs::GetPointerInfo(p, h) ? h : nullptr;
+ return DeviceAllocator::GetPointerInfo(p, h) ? h : nullptr;
}
Header* GetHeader(uptr chunk, Header* h) const {
@@ -321,12 +350,12 @@ class DeviceAllocatorT {
// is unloaded, GetPointerInfo() will fail. For such case, we can still
// return a valid value for map_beg, map_size will be limited to one page
if (LIKELY(!dev_runtime_unloaded_)) {
- if (DeviceMemFuncs::GetPointerInfo(chunk, h))
+ if (DeviceAllocator::GetPointerInfo(chunk, h))
return h;
// If GetPointerInfo() fails, we don't assume the runtime is unloaded yet.
// We just return a conservative single-page header. Here mark/check the
// runtime shutdown state
- dev_runtime_unloaded_ = DeviceMemFuncs::IsAmdgpuRuntimeShutdown();
+ dev_runtime_unloaded_ = DeviceAllocator::IsRuntimeShutdown();
}
// If we reach here, device runtime is unloaded.
// Fallback: conservative single-page header
@@ -346,7 +375,7 @@ class DeviceAllocatorT {
mutable bool dev_runtime_unloaded_;
// Maximum of mem_funcs_init_count_ is 2:
// 1. The initial init called from Init(...), it could fail if
- // libhsa-runtime64.so is dynamically loaded with dlopen()
+ // the device runtime is dynamically loaded with dlopen()
// 2. A potential deferred init called by Allocate(...)
u32 mem_funcs_init_count_;
uptr kMetadataSize_;
@@ -360,4 +389,9 @@ class DeviceAllocatorT {
} stats;
mutable StaticSpinMutex mutex_;
};
-#endif // SANITIZER_AMDGPU
+
+#if !SANITIZER_AMDGPU
+template <class MapUnmapCallback>
+using DefaultDeviceAllocator =
+ DeviceAllocatorT<MapUnmapCallback, NoOpDeviceAllocator>;
+#endif
>From 591ee2fb27d2eb7524d4a3921e90368f1f2b86d2 Mon Sep 17 00:00:00 2001
From: Amit Pandey <pandey.kumaramit2023 at gmail.com>
Date: Tue, 26 May 2026 16:41:36 +0530
Subject: [PATCH 25/27] [ASan][AMDGPU] Fix GPU-only vmem reserve tracking and
double-free reporting
hsa_amd_vmem_address_reserve_align without HSA_AMD_VMEM_ADDRESS_NO_REGISTER
reserves KFD address-only VA that is not host-writable. Sending those calls
through instance.Allocate() faulted when writing AsanChunk metadata before the
user pointer.
For that path, call the real HSA reserve/free APIs and record reservations in
VmemGpuReserveTracker. Report double-free from asan_hsa_amd_vmem_address_free
when the same range is freed twice. Reservations with NO_REGISTER still use the
existing device-heap path (e.g. HIP managed memory).
Drop XFAIL from the vmem reserve_align lit tests now that the interceptors
behave correctly.
---
compiler-rt/lib/asan/asan_allocator.cpp | 24 +++++++++++++
compiler-rt/lib/asan/asan_descriptions.cpp | 3 ++
compiler-rt/lib/asan/asan_errors.h | 5 ++-
compiler-rt/lib/asan/asan_report.cpp | 16 +++++++++
.../sanitizer_allocator_amdgpu.cpp | 32 +++++++++++++++++
.../sanitizer_allocator_amdgpu.h | 34 +++++++++++++++++++
...sa_amd_pointer_info_vmem_reserve_align.cpp | 3 --
...vmem_address_reserve_align_double_free.cpp | 13 +++----
8 files changed, 120 insertions(+), 10 deletions(-)
diff --git a/compiler-rt/lib/asan/asan_allocator.cpp b/compiler-rt/lib/asan/asan_allocator.cpp
index 24b22035a10fb..a479241605a26 100644
--- a/compiler-rt/lib/asan/asan_allocator.cpp
+++ b/compiler-rt/lib/asan/asan_allocator.cpp
@@ -1650,6 +1650,15 @@ hsa_status_t asan_hsa_amd_vmem_address_reserve_align(
return HSA_STATUS_ERROR_INVALID_ARGUMENT;
}
+ if (!__sanitizer::AmdgpuVmemReserveUsesHostMapping(flags)) {
+ const hsa_status_t status = REAL(hsa_amd_vmem_address_reserve_align)(
+ ptr, size, address, alignment, flags);
+ if (status == HSA_STATUS_SUCCESS && ptr && *ptr)
+ __sanitizer::VmemGpuReserveTracker::Get().OnReserve(
+ reinterpret_cast<uptr>(*ptr), size);
+ return status;
+ }
+
AmdgpuAllocationInfo aa_info;
aa_info.alloc_func =
reinterpret_cast<void*>(asan_hsa_amd_vmem_address_reserve_align);
@@ -1688,6 +1697,21 @@ hsa_status_t asan_hsa_amd_vmem_address_free(void* ptr, size_t size,
instance.Deallocate(ptr, 0, 0, stack, FROM_MALLOC);
return HSA_STATUS_SUCCESS;
}
+
+ using VmemGpuReserveTracker = __sanitizer::VmemGpuReserveTracker;
+ switch (
+ VmemGpuReserveTracker::Get().OnFree(reinterpret_cast<uptr>(ptr), size)) {
+ case VmemGpuReserveTracker::kNotTracked:
+ break;
+ case VmemGpuReserveTracker::kFirstFree:
+ return REAL(hsa_amd_vmem_address_free)(ptr, size);
+ case VmemGpuReserveTracker::kSizeMismatch:
+ errno = errno_EINVAL;
+ return HSA_STATUS_ERROR_INVALID_ARGUMENT;
+ case VmemGpuReserveTracker::kDoubleFree:
+ ReportDoubleFree(reinterpret_cast<uptr>(ptr), stack);
+ return HSA_STATUS_SUCCESS;
+ }
return REAL(hsa_amd_vmem_address_free)(ptr, size);
}
diff --git a/compiler-rt/lib/asan/asan_descriptions.cpp b/compiler-rt/lib/asan/asan_descriptions.cpp
index 551b819a4c436..e6e815591f4f6 100644
--- a/compiler-rt/lib/asan/asan_descriptions.cpp
+++ b/compiler-rt/lib/asan/asan_descriptions.cpp
@@ -418,6 +418,9 @@ void StackAddressDescription::Print() const {
}
void HeapAddressDescription::Print() const {
+ if (alloc_tid == kInvalidTid)
+ return;
+
PrintHeapChunkAccess(addr, chunk_access);
asanThreadRegistry().CheckLocked();
diff --git a/compiler-rt/lib/asan/asan_errors.h b/compiler-rt/lib/asan/asan_errors.h
index d9a626e711282..3c5872ac2bef1 100644
--- a/compiler-rt/lib/asan/asan_errors.h
+++ b/compiler-rt/lib/asan/asan_errors.h
@@ -73,7 +73,10 @@ struct ErrorDoubleFree : ErrorBase {
: ErrorBase(tid, 42, "double-free"),
second_free_stack(stack) {
CHECK_GT(second_free_stack->size, 0);
- GetHeapAddressInformation(addr, 1, &addr_description);
+ if (!GetHeapAddressInformation(addr, 1, &addr_description)) {
+ internal_memset(&addr_description, 0, sizeof(addr_description));
+ addr_description.addr = addr;
+ }
}
void Print();
};
diff --git a/compiler-rt/lib/asan/asan_report.cpp b/compiler-rt/lib/asan/asan_report.cpp
index df797deaa5dbd..331ce8e321c11 100644
--- a/compiler-rt/lib/asan/asan_report.cpp
+++ b/compiler-rt/lib/asan/asan_report.cpp
@@ -249,6 +249,22 @@ void ReportDeadlySignal(const SignalContext &sig) {
}
void ReportDoubleFree(uptr addr, BufferedStackTrace *free_stack) {
+ // GPU-only vmem reservations (VmemGpuReserveTracker) are not backed by
+ // AsanChunk metadata. Avoid reading chunk headers at such addresses to
+ // prevent potential crashes in StackDepot (e.g. 0xaa poison patterns).
+ if (!FindHeapChunkByAddress(addr).IsValid()) {
+ ScopedInErrorReport in_report;
+ Decorator d;
+ Printf("%s", d.Error());
+ Report(
+ "ERROR: AddressSanitizer: attempting double-free on %p in thread %s:\n",
+ (void*)addr, AsanThreadIdAndName(GetCurrentTidOrInvalid()).c_str());
+ Printf("%s", d.Default());
+ CHECK_GT(free_stack->size, 0);
+ free_stack->Print();
+ ReportErrorSummary("double-free", free_stack);
+ return;
+ }
ScopedInErrorReport in_report;
ErrorDoubleFree error(GetCurrentTidOrInvalid(), free_stack, addr);
in_report.ReportError(error);
diff --git a/compiler-rt/lib/sanitizer_common/sanitizer_allocator_amdgpu.cpp b/compiler-rt/lib/sanitizer_common/sanitizer_allocator_amdgpu.cpp
index 461ad033cac9f..879f892499922 100644
--- a/compiler-rt/lib/sanitizer_common/sanitizer_allocator_amdgpu.cpp
+++ b/compiler-rt/lib/sanitizer_common/sanitizer_allocator_amdgpu.cpp
@@ -223,5 +223,37 @@ void AmdgpuDeviceAllocator::RegisterSystemEventHandlers() {
}
uptr AmdgpuDeviceAllocator::GetPageSize() { return kPageSize_; }
+
+VmemGpuReserveTracker& VmemGpuReserveTracker::Get() {
+ static VmemGpuReserveTracker tracker;
+ return tracker;
+}
+
+void VmemGpuReserveTracker::OnReserve(uptr ptr, uptr size) {
+ SpinMutexLock l(&mu_);
+ VmemGpuReservation entry;
+ entry.ptr = ptr;
+ entry.size = size;
+ entry.freed = false;
+ reservations_.push_back(entry);
+}
+
+VmemGpuReserveTracker::FreeResult VmemGpuReserveTracker::OnFree(uptr ptr,
+ uptr size) {
+ SpinMutexLock l(&mu_);
+ for (uptr i = 0; i < reservations_.size(); ++i) {
+ VmemGpuReservation& r = reservations_[i];
+ if (r.ptr != ptr)
+ continue;
+ if (r.size != size)
+ return kSizeMismatch;
+ if (r.freed)
+ return kDoubleFree;
+ r.freed = true;
+ return kFirstFree;
+ }
+ return kNotTracked;
+}
+
} // namespace __sanitizer
#endif // SANITIZER_AMDGPU
diff --git a/compiler-rt/lib/sanitizer_common/sanitizer_allocator_amdgpu.h b/compiler-rt/lib/sanitizer_common/sanitizer_allocator_amdgpu.h
index efc6c8f869e3b..c33d8ace22fec 100644
--- a/compiler-rt/lib/sanitizer_common/sanitizer_allocator_amdgpu.h
+++ b/compiler-rt/lib/sanitizer_common/sanitizer_allocator_amdgpu.h
@@ -36,6 +36,40 @@ class AmdgpuDeviceAllocator {
static void NotifyRuntimeShutdown();
};
+// GPU-only hsa_amd_vmem_address_reserve_align (without
+// HSA_AMD_VMEM_ADDRESS_NO_REGISTER) returns KFD "address only" VA that is not
+// host-writable. Sanitizer runtimes that cannot place metadata in the reserved
+// range track such reservations here for double-free detection on free.
+class VmemGpuReserveTracker {
+ public:
+ enum FreeResult {
+ kNotTracked,
+ kFirstFree,
+ kSizeMismatch,
+ kDoubleFree,
+ };
+
+ static VmemGpuReserveTracker& Get();
+
+ void OnReserve(uptr ptr, uptr size);
+ FreeResult OnFree(uptr ptr, uptr size);
+
+ private:
+ struct VmemGpuReservation {
+ uptr ptr;
+ uptr size;
+ bool freed;
+ };
+
+ StaticSpinMutex mu_;
+ InternalMmapVector<VmemGpuReservation> reservations_;
+};
+
+// True when ROCr reserves host-accessible VA (e.g. HIP managed / SVM).
+inline bool AmdgpuVmemReserveUsesHostMapping(u64 flags) {
+ return (flags & HSA_AMD_VMEM_ADDRESS_NO_REGISTER) != 0;
+}
+
struct AmdgpuAllocationInfo : public DeviceAllocationInfo {
AmdgpuAllocationInfo() : DeviceAllocationInfo(DAT_AMDGPU) {
status = HSA_STATUS_SUCCESS;
diff --git a/compiler-rt/test/asan/TestCases/AMDGPU/hsa_amd_pointer_info_vmem_reserve_align.cpp b/compiler-rt/test/asan/TestCases/AMDGPU/hsa_amd_pointer_info_vmem_reserve_align.cpp
index 06c09d041dd67..5c22d85b23f21 100644
--- a/compiler-rt/test/asan/TestCases/AMDGPU/hsa_amd_pointer_info_vmem_reserve_align.cpp
+++ b/compiler-rt/test/asan/TestCases/AMDGPU/hsa_amd_pointer_info_vmem_reserve_align.cpp
@@ -1,6 +1,3 @@
-// XFAIL: *
-//
-//
// RUN: %clangxx_asan -O0 -isystem %rocm_include %s -o %t -L%rocm_lib -lhsa-runtime64 \
// RUN: -Wl,-rpath,%rocm_lib -Wl,-rpath,%compiler_rt_libdir
// RUN: %run %t 2>&1 | FileCheck %s
diff --git a/compiler-rt/test/asan/TestCases/AMDGPU/hsa_amd_vmem_address_reserve_align_double_free.cpp b/compiler-rt/test/asan/TestCases/AMDGPU/hsa_amd_vmem_address_reserve_align_double_free.cpp
index 038fec02f7be3..1828e6e930bbd 100644
--- a/compiler-rt/test/asan/TestCases/AMDGPU/hsa_amd_vmem_address_reserve_align_double_free.cpp
+++ b/compiler-rt/test/asan/TestCases/AMDGPU/hsa_amd_vmem_address_reserve_align_double_free.cpp
@@ -1,6 +1,3 @@
-// XFAIL: *
-//
-//
// RUN: %clangxx_asan -O0 -isystem %rocm_include %s -o %t -L%rocm_lib -lhsa-runtime64 \
// RUN: -Wl,-rpath,%rocm_lib -Wl,-rpath,%compiler_rt_libdir
// RUN: not %run %t 2>&1 | FileCheck %s
@@ -28,9 +25,13 @@ int main() {
// asan_hsa_amd_vmem_address_reserve_align).
const size_t kSize = 4096;
void *mem = nullptr;
- if (hsa_amd_vmem_address_reserve_align(&mem, kSize, /*address=*/0,
- /*alignment=*/4096,
- /*flags=*/0) != HSA_STATUS_SUCCESS ||
+
+ // NOTE: To use `hipMallocManaged` way of reserving memory,
+ // use either `HSA_AMD_VMEM_ADDRESS_NO_REGISTER` in flags.
+ if (hsa_amd_vmem_address_reserve_align(
+ &mem, kSize, /*address=*/0,
+ /*alignment=*/4096,
+ /*flags=*/HSA_AMD_VMEM_ADDRESS_NO_REGISTER) != HSA_STATUS_SUCCESS ||
!mem) {
fprintf(stderr, "hsa_amd_vmem_address_reserve_align failed\n");
return 1;
>From d08a2d726a4d37f5d93a6aa9dbfb382232ed6164 Mon Sep 17 00:00:00 2001
From: Amit Pandey <pandey.kumaramit2023 at gmail.com>
Date: Mon, 1 Jun 2026 12:36:35 +0530
Subject: [PATCH 26/27] Fix merge confilct in asan_allocator.cpp
---
compiler-rt/lib/asan/asan_allocator.cpp | 6 +++---
1 file changed, 3 insertions(+), 3 deletions(-)
diff --git a/compiler-rt/lib/asan/asan_allocator.cpp b/compiler-rt/lib/asan/asan_allocator.cpp
index a479241605a26..f9da22f7f1902 100644
--- a/compiler-rt/lib/asan/asan_allocator.cpp
+++ b/compiler-rt/lib/asan/asan_allocator.cpp
@@ -544,7 +544,7 @@ struct Allocator {
// (true) or a fatal Report*+Die() (false).
void* AllocateImpl(uptr size, uptr alignment, BufferedStackTrace* stack,
AllocType alloc_type, bool can_fill,
- bool may_return_null,DeviceAllocationInfo *da_info = nullptr) {
+ bool may_return_null,DeviceAllocationInfo *da_info) {
if (UNLIKELY(!AsanInited()))
AsanInitFromRtl();
if (UNLIKELY(IsRssLimitExceeded())) {
@@ -688,9 +688,9 @@ struct Allocator {
// Defer to the global, flag controlled, OOM policy.
void* Allocate(uptr size, uptr alignment, BufferedStackTrace* stack,
- AllocType alloc_type, bool can_fill) {
+ AllocType alloc_type, bool can_fill, DeviceAllocationInfo *da_info = nullptr) {
return AllocateImpl(size, alignment, stack, alloc_type, can_fill,
- AllocatorMayReturnNull());
+ AllocatorMayReturnNull(),da_info);
}
// Set quarantine flag if chunk is allocated, issue ASan error report on
>From 2509db44351535871a5dac804d8b6830728c084b Mon Sep 17 00:00:00 2001
From: Amit Pandey <pandey.kumaramit2023 at gmail.com>
Date: Mon, 1 Jun 2026 15:05:56 +0530
Subject: [PATCH 27/27] [compiler-rt][asan][AMDGPU] Fix pointer_info bases for
vmem and device ASan internal GetPointerInfo bases.
Only apply hsa_amd_pointer_info offset fixups to agentBaseAddress and
hostBaseAddress when ROCr returned a non-null value, so unmapped
RESERVED_ADDR reservations do not get a fabricated agent base from NULL
+ offset. Still shift hostBaseAddress for NO_REGISTER reserves where
ROCr reports os_addr below the aligned user VA.
Improve AmdgpuDeviceAllocator::GetPointerInfo to derive map_beg for all
relevant HSA pointer types, handle aligned NO_REGISTER vmem
reservations, and reject UNKNOWN or zero-size mappings. Update the vmem
pointer_info lit test to use NO_REGISTER host-assisted reserves and
assert on hostBaseAddress instead of agentBaseAddress.
---
compiler-rt/lib/asan/asan_allocator.cpp | 23 ++++++----
.../sanitizer_allocator_amdgpu.cpp | 42 +++++++++++++++++--
...sa_amd_pointer_info_vmem_reserve_align.cpp | 19 +++++----
3 files changed, 65 insertions(+), 19 deletions(-)
diff --git a/compiler-rt/lib/asan/asan_allocator.cpp b/compiler-rt/lib/asan/asan_allocator.cpp
index f9da22f7f1902..f6973eef30e86 100644
--- a/compiler-rt/lib/asan/asan_allocator.cpp
+++ b/compiler-rt/lib/asan/asan_allocator.cpp
@@ -543,8 +543,8 @@ struct Allocator {
// may_return_null tells AllocateImpl() whether OOM should produce a nullptr
// (true) or a fatal Report*+Die() (false).
void* AllocateImpl(uptr size, uptr alignment, BufferedStackTrace* stack,
- AllocType alloc_type, bool can_fill,
- bool may_return_null,DeviceAllocationInfo *da_info) {
+ AllocType alloc_type, bool can_fill, bool may_return_null,
+ DeviceAllocationInfo* da_info) {
if (UNLIKELY(!AsanInited()))
AsanInitFromRtl();
if (UNLIKELY(IsRssLimitExceeded())) {
@@ -688,9 +688,10 @@ struct Allocator {
// Defer to the global, flag controlled, OOM policy.
void* Allocate(uptr size, uptr alignment, BufferedStackTrace* stack,
- AllocType alloc_type, bool can_fill, DeviceAllocationInfo *da_info = nullptr) {
+ AllocType alloc_type, bool can_fill,
+ DeviceAllocationInfo* da_info = nullptr) {
return AllocateImpl(size, alignment, stack, alloc_type, can_fill,
- AllocatorMayReturnNull(),da_info);
+ AllocatorMayReturnNull(), da_info);
}
// Set quarantine flag if chunk is allocated, issue ASan error report on
@@ -1745,10 +1746,16 @@ hsa_status_t asan_hsa_amd_pointer_info(const void* ptr,
// hsa_amd_vmem_address_reserve_align()) uses alignment > page size (device
// allocator padding and ASan redzones/headers).
const uptr offset = user - reinterpret_cast<uptr>(hsa_map_base);
- info->agentBaseAddress = reinterpret_cast<void*>(
- reinterpret_cast<uptr>(info->agentBaseAddress) + offset);
- info->hostBaseAddress = reinterpret_cast<void*>(
- reinterpret_cast<uptr>(info->hostBaseAddress) + offset);
+ // hostBaseAddress must reflect the user-visible reservation base (ROCr may
+ // report os_addr below an aligned user VA for NO_REGISTER reserves).
+ if (info->hostBaseAddress)
+ info->hostBaseAddress = reinterpret_cast<void*>(
+ reinterpret_cast<uptr>(info->hostBaseAddress) + offset);
+ // agentBaseAddress is NULL for unmapped RESERVED_ADDR (no GPU backing).
+ // Do not turn NULL + offset into a fake agent address.
+ if (info->agentBaseAddress)
+ info->agentBaseAddress = reinterpret_cast<void*>(
+ reinterpret_cast<uptr>(info->agentBaseAddress) + offset);
info->sizeInBytes = m->UsedSize();
}
return status;
diff --git a/compiler-rt/lib/sanitizer_common/sanitizer_allocator_amdgpu.cpp b/compiler-rt/lib/sanitizer_common/sanitizer_allocator_amdgpu.cpp
index 879f892499922..5844fffaba386 100644
--- a/compiler-rt/lib/sanitizer_common/sanitizer_allocator_amdgpu.cpp
+++ b/compiler-rt/lib/sanitizer_common/sanitizer_allocator_amdgpu.cpp
@@ -162,8 +162,38 @@ void AmdgpuDeviceAllocator::Deallocate(void* p) {
}
}
+static uptr AmdgpuPointerInfoMapBase(uptr ptr,
+ const hsa_amd_pointer_info_t& info) {
+ switch (info.type) {
+ case HSA_EXT_POINTER_TYPE_RESERVED_ADDR: {
+ uptr map_beg = reinterpret_cast<uptr>(info.hostBaseAddress);
+ // ROCr sets hostBaseAddress to the OS/KFD reservation base (os_addr).
+ // With HSA_AMD_VMEM_ADDRESS_NO_REGISTER the user VA from reserve_align
+ // is AlignUp(os_addr, alignment) and is the handle for vmem_address_free;
+ // it can differ from hostBaseAddress when alignment > page size.
+ if (map_beg && ptr >= map_beg && ptr < map_beg + info.sizeInBytes)
+ return map_beg;
+ return ptr;
+ }
+ case HSA_EXT_POINTER_TYPE_HSA:
+ case HSA_EXT_POINTER_TYPE_HSA_VMEM:
+ case HSA_EXT_POINTER_TYPE_IPC:
+ // Match ROCr PtrInfo block_info: prefer host mapping when present.
+ if (info.hostBaseAddress)
+ return reinterpret_cast<uptr>(info.hostBaseAddress);
+ return reinterpret_cast<uptr>(info.agentBaseAddress);
+ case HSA_EXT_POINTER_TYPE_LOCKED:
+ return reinterpret_cast<uptr>(info.hostBaseAddress);
+ default:
+ return 0;
+ }
+}
+
bool AmdgpuDeviceAllocator::GetPointerInfo(uptr ptr,
DevicePointerInfo* ptr_info) {
+ if (!ptr_info)
+ return false;
+
// GetPointerInfo returns false after AMDGPU runtime shutdown
if (UNLIKELY(IsRuntimeShutdown())) {
VReport(1,
@@ -180,10 +210,14 @@ bool AmdgpuDeviceAllocator::GetPointerInfo(uptr ptr,
if (status != HSA_STATUS_SUCCESS)
return false;
- if (info.type == HSA_EXT_POINTER_TYPE_RESERVED_ADDR)
- ptr_info->map_beg = reinterpret_cast<uptr>(info.hostBaseAddress);
- else if (info.type == HSA_EXT_POINTER_TYPE_HSA)
- ptr_info->map_beg = reinterpret_cast<uptr>(info.agentBaseAddress);
+ if (info.type == HSA_EXT_POINTER_TYPE_UNKNOWN || !info.sizeInBytes)
+ return false;
+
+ const uptr map_beg = AmdgpuPointerInfoMapBase(ptr, info);
+ if (!map_beg)
+ return false;
+
+ ptr_info->map_beg = map_beg;
ptr_info->map_size = info.sizeInBytes;
ptr_info->type = reinterpret_cast<hsa_amd_pointer_type_t>(info.type);
diff --git a/compiler-rt/test/asan/TestCases/AMDGPU/hsa_amd_pointer_info_vmem_reserve_align.cpp b/compiler-rt/test/asan/TestCases/AMDGPU/hsa_amd_pointer_info_vmem_reserve_align.cpp
index 5c22d85b23f21..dceb8eca54d41 100644
--- a/compiler-rt/test/asan/TestCases/AMDGPU/hsa_amd_pointer_info_vmem_reserve_align.cpp
+++ b/compiler-rt/test/asan/TestCases/AMDGPU/hsa_amd_pointer_info_vmem_reserve_align.cpp
@@ -22,11 +22,14 @@ int main() {
if (hsa_amd_test_require_init())
return 1;
- const size_t kSize = 4096;
- const uint64_t kAlign = 65536;
+ const size_t kSize = 4096; //4KiB size
+ // Here alignment is larger than a page size, so the reserved address will be aligned to 64KiB.
+ const uint64_t kAlign = 65536; //64KiB alignment
void *mem = nullptr;
- if (hsa_amd_vmem_address_reserve_align(&mem, kSize, /*address=*/0, kAlign,
- /*flags=*/0) != HSA_STATUS_SUCCESS ||
+ // Host-accessible vmem reserve via NO_REGISTER (host mmap, not KFD VA).
+ if (hsa_amd_vmem_address_reserve_align(
+ &mem, kSize, /*address=*/0, kAlign,
+ /*flags=*/HSA_AMD_VMEM_ADDRESS_NO_REGISTER) != HSA_STATUS_SUCCESS ||
!mem) {
fprintf(stderr, "hsa_amd_vmem_address_reserve_align failed\n");
return 1;
@@ -40,14 +43,16 @@ int main() {
hsa_amd_pointer_info_t info = {};
info.size = sizeof(hsa_amd_pointer_info_t);
+ // Intercepted by ASan, where hostBaseAddress is the user pointer.
if (hsa_amd_pointer_info(mem, &info, nullptr, nullptr, nullptr) !=
HSA_STATUS_SUCCESS) {
fprintf(stderr, "hsa_amd_pointer_info failed\n");
return 1;
}
- if (info.agentBaseAddress != mem ||
- reinterpret_cast<uintptr_t>(info.agentBaseAddress) + info.sizeInBytes !=
+ // Verify: User Pointer(mem) returned is hostBaseAddress as expected.
+ if (info.hostBaseAddress != mem ||
+ reinterpret_cast<uintptr_t>(info.hostBaseAddress) + info.sizeInBytes !=
user + kSize) {
fprintf(stderr,
"pointer_info mismatch: user=%p begin=%p size=%zu (expected %zu)\n",
@@ -56,7 +61,7 @@ int main() {
}
printf("pointer_info_vmem sizeInBytes: %zu\n", info.sizeInBytes);
- printf("pointer_info_vmem begin: %p\n", info.agentBaseAddress);
+ printf("pointer_info_vmem begin: %p\n", info.hostBaseAddress);
if (hsa_amd_vmem_address_free(mem, kSize) != HSA_STATUS_SUCCESS) {
fprintf(stderr, "hsa_amd_vmem_address_free failed\n");
More information about the llvm-commits
mailing list