[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
Wed May 6 01:04:48 PDT 2026
https://github.com/ampandey-1995 updated https://github.com/llvm/llvm-project/pull/192240
>From 8d4a95b16951b03306b3bbcdae6a7b4d999b1f57 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/18] [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 6e8045631b218..f2576a39350fb 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 a1f3cf799e344..d088fe4b96863 100644
--- a/compiler-rt/lib/sanitizer_common/sanitizer_platform.h
+++ b/compiler-rt/lib/sanitizer_common/sanitizer_platform.h
@@ -302,7 +302,7 @@
# define SANITIZER_LOONGARCH64 0
#endif
-#if defined(__AMDGPU__)
+#if SANITIZER_AMDGPU || defined(__AMDGPU__)
# define SANITIZER_AMDGPU 1
#else
# define SANITIZER_AMDGPU 0
>From deb260c851752854c5f11beef0cb7b41259a6250 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/18] [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 7a706f523172520832d9c0f30cf66236da4fb9d0 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/18] [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 e723301d04940d5bb7327842ef990d1d092d0e56 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/18] [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 fef8f7ab297cc..6cfe268e15b80 100644
--- a/compiler-rt/test/lit.common.cfg.py
+++ b/compiler-rt/test/lit.common.cfg.py
@@ -574,6 +574,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 a62a8dbaa98be9eec8e8371bf8e6d6fb2fd17d0d 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/18] [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 673dd2cb21eb586284b8c1182b8dbfc06b2ea7b1 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/18] 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 982c9187a0a19db0c34e1963e4db0f2e6b83323c 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/18] 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 65982d89b58d6499772933259ca073304471c099 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/18] 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 fe74cc0ae1d3cecd218e91703d09d13128e31a6e 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/18] [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 8663ae8a9ab020caa01b653ee6b4b470bd10a449 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/18] [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 18f5a110f4ce10052b555c3abd4af205143f4d06 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/18] 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 ebe669cf018da5c620d93ef6543a55afa4f7d8a7 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/18] [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 f2576a39350fb..3ff98abda2ff1 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 c1ba87f865af2a25881f8bb01a2f89a5dcbdc269 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/18] [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 3ff98abda2ff1..d1356cf9eaea3 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 97d696217961234494b347e66c5ddad200b29412 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/18] [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 38d90cb63ab740b0f3979dd86a3e5f27f13c2991 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/18] [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 d1356cf9eaea3..4dd47fb4a78c7 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 ecf7a30ae06f43a4f189a21012a4455d0cd3e10b 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/18] [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 700a240c5b0ea1d78ad029263c9670b08a1ba842 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/18] [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 c16dfb1540aae5f1d6d67a2a27db5a8b86eedc25 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/18] [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 | 158 ++++++++++++++++++
.../AMDGPU/hsa_amd_ipc_memory_roundtrip.cpp | 149 +++++++++++++++++
.../hsa_amd_memory_async_copy_overlap.cpp | 61 +++++++
3 files changed, 368 insertions(+)
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..85bc29f2cdff1
--- /dev/null
+++ b/compiler-rt/test/asan/TestCases/AMDGPU/hsa_amd_ipc_memory_attach_heap_oob.cpp
@@ -0,0 +1,158 @@
+// 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: SUMMARY: AddressSanitizer: heap-buffer-overflow
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..a258473b6ce92
--- /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 -g
+// 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..deb31e29f7875
--- /dev/null
+++ b/compiler-rt/test/asan/TestCases/AMDGPU/hsa_amd_memory_async_copy_overlap.cpp
@@ -0,0 +1,61 @@
+// 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
More information about the llvm-commits
mailing list