[clang] c1c22cd - [ASan][HIP] Add ASan declarations and macros. (#167522)

[via cfe-commits]

Author: Amit Kumar Pandey
Date: 2025-11-17T11:13:09-08:00
New Revision: c1c22cd3e16beb3937eb0d11da014451397be5d6

URL: https://github.com/llvm/llvm-project/commit/c1c22cd3e16beb3937eb0d11da014451397be5d6
DIFF: https://github.com/llvm/llvm-project/commit/c1c22cd3e16beb3937eb0d11da014451397be5d6.diff

LOG: [ASan][HIP] Add ASan declarations and macros. (#167522)

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

Added: 
    

Modified: 
    clang/lib/Headers/__clang_hip_libdevice_declares.h

Removed: 
    


################################################################################
diff  --git a/clang/lib/Headers/__clang_hip_libdevice_declares.h b/clang/lib/Headers/__clang_hip_libdevice_declares.h
index fa8d918248dd0..fad9c6ca7ffc5 100644
--- a/clang/lib/Headers/__clang_hip_libdevice_declares.h
+++ b/clang/lib/Headers/__clang_hip_libdevice_declares.h
@@ -338,6 +338,23 @@ __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_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_TYPE__ 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