Skip to content

Commit c1c22cd

Browse files
[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
1 parent c555522 commit c1c22cd

File tree

1 file changed

+17
-0
lines changed

1 file changed

+17
-0
lines changed

clang/lib/Headers/__clang_hip_libdevice_declares.h

Lines changed: 17 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -338,6 +338,23 @@ __device__ __attribute__((const)) __2f16 __ocml_sqrt_2f16(__2f16);
338338
__device__ __attribute__((const)) __2f16 __ocml_trunc_2f16(__2f16);
339339
__device__ __attribute__((const)) __2f16 __ocml_pown_2f16(__2f16, __2i16);
340340

341+
__device__ void __asan_poison_memory_region(const void *addr,
342+
__SIZE_TYPE__ size);
343+
__device__ void __asan_unpoison_memory_region(const void *addr,
344+
__SIZE_TYPE__ size);
345+
__device__ int __asan_address_is_poisoned(const void *addr);
346+
__device__ void *__asan_region_is_poisoned(void *beg, __SIZE_TYPE__ size);
347+
348+
#if __has_feature(address_sanitizer)
349+
#define ASAN_POISON_MEMORY_REGION(addr, size) \
350+
__asan_poison_memory_region((addr), (size))
351+
#define ASAN_UNPOISON_MEMORY_REGION(addr, size) \
352+
__asan_unpoison_memory_region((addr), (size))
353+
#else
354+
#define ASAN_POISON_MEMORY_REGION(addr, size) ((void)(addr), (void)(size))
355+
#define ASAN_UNPOISON_MEMORY_REGION(addr, size) ((void)(addr), (void)(size))
356+
#endif
357+
341358
#ifdef __cplusplus
342359
} // extern "C"
343360
#endif

0 commit comments

Comments
 (0)