[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
Thu May 14 03:42:38 PDT 2026
https://github.com/ampandey-1995 updated https://github.com/llvm/llvm-project/pull/192240
>From 287801ac9fabbf286c064c1d69c91c6fc8662ca6 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/21] [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 e88321d822f84..a88a13a8c3be2 100644
--- a/compiler-rt/CMakeLists.txt
+++ b/compiler-rt/CMakeLists.txt
@@ -540,6 +540,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 2c70fb7712597..70c27eb3e548c 100644
--- a/compiler-rt/lib/sanitizer_common/sanitizer_platform.h
+++ b/compiler-rt/lib/sanitizer_common/sanitizer_platform.h
@@ -303,7 +303,7 @@
# define SANITIZER_LOONGARCH64 0
#endif
-#if defined(__AMDGPU__)
+#if SANITIZER_AMDGPU || defined(__AMDGPU__)
# define SANITIZER_AMDGPU 1
#else
# define SANITIZER_AMDGPU 0
>From 95b881c704c38701125b7c27c918feabbcd89827 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/21] [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 | 261 +++++++++++++++++-
compiler-rt/lib/asan/asan_allocator.h | 50 ++++
.../sanitizer_common/sanitizer_allocator.h | 27 +-
.../sanitizer_allocator_combined.h | 96 ++++++-
4 files changed, 413 insertions(+), 21 deletions(-)
diff --git a/compiler-rt/lib/asan/asan_allocator.cpp b/compiler-rt/lib/asan/asan_allocator.cpp
index 46ba7e16da9b2..a12c1689cd3a4 100644
--- a/compiler-rt/lib/asan/asan_allocator.cpp
+++ b/compiler-rt/lib/asan/asan_allocator.cpp
@@ -395,7 +395,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
@@ -536,8 +543,9 @@ struct Allocator {
}
// -------------------- Allocation/Deallocation routines ---------------
- void *Allocate(uptr size, uptr alignment, BufferedStackTrace *stack,
- AllocType alloc_type, bool can_fill) {
+ void* Allocate(uptr size, uptr alignment, BufferedStackTrace* stack,
+ AllocType alloc_type, bool can_fill,
+ DeviceAllocationInfo* da_info = nullptr) {
if (UNLIKELY(!AsanInited()))
AsanInitFromRtl();
if (UNLIKELY(IsRssLimitExceeded())) {
@@ -592,11 +600,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);
+ allocated = allocator.Allocate(cache, needed_size, 8, da_info);
}
if (UNLIKELY(!allocated)) {
SetAllocatorOutOfMemory();
@@ -1463,3 +1471,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 a02d1434a273d..068808e2bb4e3 100644
--- a/compiler-rt/lib/asan/asan_allocator.h
+++ b/compiler-rt/lib/asan/asan_allocator.h
@@ -323,4 +323,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 cb09aa2905a5759b2d05334eb9cdd71384a0176f 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/21] [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 c1401ee2ef3c56352230c1d9a968eaf8c6bd49ab 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/21] [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 a1cda7f93382f..9a46d5ae1d145 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 fc836f0d46a7fec71be83ceeeb833d61f1aa85a1 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/21] [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 18abd9af108b39af7513d3093e91171ebab7bb80 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/21] 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 8f40492d8c023a6d823d801261c274769fb69deb 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/21] 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 60542b3fbcedb8fdde6be34b8861a7dceaa3c634 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/21] 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 a12c1689cd3a4..976653f1288a2 100644
--- a/compiler-rt/lib/asan/asan_allocator.cpp
+++ b/compiler-rt/lib/asan/asan_allocator.cpp
@@ -1714,5 +1714,6 @@ hsa_status_t asan_hsa_init() {
}
return status;
}
-#endif
} // namespace __asan
+
+#endif // SANITIZER_AMDGPU
>From fdc192c687685cb08683c9d62d94dcc2bd2ae293 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/21] [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 cc56662b8a630857c3cc67f06792e539bdc31589 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/21] [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 f8d88d7ba8e82ed7ad1e39a6efa02e44fb5959d5 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/21] 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 5ea925e5d4ff573855a30bd0bfac2fd79062512f 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/21] [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 a88a13a8c3be2..29bc8e05f1f0e 100644
--- a/compiler-rt/CMakeLists.txt
+++ b/compiler-rt/CMakeLists.txt
@@ -542,22 +542,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 796c11643b6ef02e5d220e79016377e1b27ff926 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/21] [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 29bc8e05f1f0e..515d9529e7a84 100644
--- a/compiler-rt/CMakeLists.txt
+++ b/compiler-rt/CMakeLists.txt
@@ -542,8 +542,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 e9fc966ad2fa3cd948b54732baf996dda21a8ed7 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/21] [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 db0443ddaede95bb24c6c4589e75fc1b4b037447 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/21] [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 515d9529e7a84..aec41e03ac265 100644
--- a/compiler-rt/CMakeLists.txt
+++ b/compiler-rt/CMakeLists.txt
@@ -543,7 +543,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 e460d294c09d5e82e2299c5d122308d6bab2e0ae 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/21] [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 4dd2187df2272..35c1f374fbdb8 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 bbe876828217f..157ddb3dfc344 100644
--- a/compiler-rt/lib/sanitizer_common/sanitizer_linux.cpp
+++ b/compiler-rt/lib/sanitizer_common/sanitizer_linux.cpp
@@ -2882,6 +2882,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 f86fe4b40e0e31bac966fd38e04b6c50b8f1f696 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/21] [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 878377900ebd3276c2c292f3feef4c61c404714f 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/21] [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 afcffc053328e24172b5d4537c1b4f5863c863f4 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/21] [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 f7830b34a0ba295dabbed03700e222028b5ac8c8 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/21] [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 9f6bf43d0ca8a3ae86359edf349c200214a88408 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/21] 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 976653f1288a2..8d04f6234ec6f 100644
--- a/compiler-rt/lib/asan/asan_allocator.cpp
+++ b/compiler-rt/lib/asan/asan_allocator.cpp
@@ -1604,16 +1604,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(
@@ -1664,6 +1673,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);
More information about the llvm-commits
mailing list