From fad918fa72e31ac390ea5186cb60b9acbb7fcb12 Mon Sep 17 00:00:00 2001 From: Amit Kumar Pandey Date: Tue, 11 Nov 2025 20:56:03 +0530 Subject: [PATCH 1/2] [ASan][HIP] Add ASan declarations and macros. This patch adds the following device ASan hooks and guarded macros in __clang_hip_libdevice_declares.h - Function Declarations - __asan_poison_memory_region - __asan_unpoison_memory_region - __asan_address_is_poisoned - __asan_region_is_poisoned - Macros - ASAN_POISON_MEMORY_REGION - ASAN_UNPOISON_MEMORY_REGION --- .../lib/Headers/__clang_hip_libdevice_declares.h | 15 +++++++++++++++ 1 file changed, 15 insertions(+) diff --git a/clang/lib/Headers/__clang_hip_libdevice_declares.h b/clang/lib/Headers/__clang_hip_libdevice_declares.h index fa8d918248dd0..9861e674333a7 100644 --- a/clang/lib/Headers/__clang_hip_libdevice_declares.h +++ b/clang/lib/Headers/__clang_hip_libdevice_declares.h @@ -338,6 +338,21 @@ __device__ __attribute__((const)) __2f16 __ocml_sqrt_2f16(__2f16); __device__ __attribute__((const)) __2f16 __ocml_trunc_2f16(__2f16); __device__ __attribute__((const)) __2f16 __ocml_pown_2f16(__2f16, __2i16); +__device__ void __asan_poison_memory_region(const void *addr, size_t size); +__device__ void __asan_unpoison_memory_region(const void *addr, size_t size); +__device__ int __asan_address_is_poisoned(const void *addr); +__device__ void *__asan_region_is_poisoned(void *beg, size_t size); + +#if __has_feature(address_sanitizer) +#define ASAN_POISON_MEMORY_REGION(addr, size) \ + __asan_poison_memory_region((addr), (size)) +#define ASAN_UNPOISON_MEMORY_REGION(addr, size) \ + __asan_unpoison_memory_region((addr), (size)) +#else +#define ASAN_POISON_MEMORY_REGION(addr, size) ((void)(addr), (void)(size)) +#define ASAN_UNPOISON_MEMORY_REGION(addr, size) ((void)(addr), (void)(size)) +#endif + #ifdef __cplusplus } // extern "C" #endif From 5d7b6ad691c6d03bb3f0820d605c84ff57405fe3 Mon Sep 17 00:00:00 2001 From: Amit Kumar Pandey Date: Tue, 11 Nov 2025 21:54:24 +0530 Subject: [PATCH 2/2] Fix Build error. Switch from 'size_t' to '__SIZE_TYPE__' --- clang/lib/Headers/__clang_hip_libdevice_declares.h | 8 +++++--- 1 file changed, 5 insertions(+), 3 deletions(-) diff --git a/clang/lib/Headers/__clang_hip_libdevice_declares.h b/clang/lib/Headers/__clang_hip_libdevice_declares.h index 9861e674333a7..fad9c6ca7ffc5 100644 --- a/clang/lib/Headers/__clang_hip_libdevice_declares.h +++ b/clang/lib/Headers/__clang_hip_libdevice_declares.h @@ -338,10 +338,12 @@ __device__ __attribute__((const)) __2f16 __ocml_sqrt_2f16(__2f16); __device__ __attribute__((const)) __2f16 __ocml_trunc_2f16(__2f16); __device__ __attribute__((const)) __2f16 __ocml_pown_2f16(__2f16, __2i16); -__device__ void __asan_poison_memory_region(const void *addr, size_t size); -__device__ void __asan_unpoison_memory_region(const void *addr, size_t size); +__device__ void __asan_poison_memory_region(const void *addr, + __SIZE_TYPE__ size); +__device__ void __asan_unpoison_memory_region(const void *addr, + __SIZE_TYPE__ size); __device__ int __asan_address_is_poisoned(const void *addr); -__device__ void *__asan_region_is_poisoned(void *beg, size_t size); +__device__ void *__asan_region_is_poisoned(void *beg, __SIZE_TYPE__ size); #if __has_feature(address_sanitizer) #define ASAN_POISON_MEMORY_REGION(addr, size) \