Unified AddressSanitizer Framework for SYCL Kernel on CPU and GPU Device.

In response to the growing demand to incorporate LLVM sanitizers into heterogeneous computing environments, this poster presents the development of a Unified Sanitizer Framework for SYCL programming. The challenge lies in designing a unified framework capable of handling the unique features of each device, such as CPUs and GPUs, which possess different memory hierarchies and address spaces. Our ongoing work on the unified framework leverages AddressSanitizer[1] (ASan), a memory error detector for C/C++ languages, specifically designed for SYCL kernels. It includes the extension of ASan’s instrumentation pass with their function definitions in libdevice, and the addition of the Sanitizer Layer to the Unified Runtime (UR), which provides common ASan runtime support. Currently, it can detect out-of-bounds errors on Unified Shared Memory (USM) and static local memory (sycl::group_local_memory) for both OpenCL CPU and Level-Zero GPU backends, providing detailed error information for efficient debugging. The framework can be extended to other SYCL backends like CUDA and HIP, and can support OpenMP. It also lays the foundation for integration of additional sanitizers like MemorySanitizer[3] (MSan) and ThreadSanitizer[2] (TSan).