README.md

Introduction

AddressSanitizer (aka ASan) is a memory error detector for C/C++, we extend it to support detecting memory errors in offload code as well.

Currently, it supports to detect out-of-bounds(buffer overflow & underflow) errors in

• host/device/shared USM
• local (static only)

It depends on the GPU driver: https://ubit-gfx.intel.com/build/18430085

Usage

Install GPU driver, icpx compiler, and setup environment

unzip gfx.zip
unzip dpcpp.zip

find ./gfx -name '*.deb' -exec dpkg-deb -x {} ~/gfx-driver \;

mkdir ~/vendors
echo $HOME/dpcpp/lib/libintelocl.so > ~/vendors/intel-cpu.icd
echo $HOME/gfx-driver/usr/lib/x86_64-linux-gnu/intel-opencl/libigdrcl.so > ~/vendors/intel.icd
unset OCL_ICD_FILENAMES
export OCL_ICD_VENDORS=~/vendors

export PATH=$HOME/dpcpp/bin:$HOME/gfx-driver/usr/bin:$PATH
export LD_LIBRARY_PATH=$HOME/gfx-driver/usr/lib/x86_64-linux-gnu:$HOME/gfx-driver/usr/lib/x86_64-linux-gnu/intel-opencl:$HOME/gfx-driver/usr/lib/x86_64-linux-gnu/dri:$HOME/gfx-driver/usr/lib/x86_64-linux-gnu/mfx:$HOME/dpcpp/lib:$LD_LIBRARY_PATH

For GPU (PVC),

icpx -fsycl -fsanitize=address -fsanitize-target=device -o demo demo.cpp

For CPU,

icpx -fsycl -fsanitize=address -fsanitize-target=both -o demo demo.cpp

Flags

-fsanitize-target=host|device|both

• host: enable sanitizer only on host code
• device: enable sanitizer only on the device code
• both: enable sanitizer on both host and device code

This flag is experimental, it may change in the future.

Limits

1. Device sanitizer doesn't support detecting out-of-bounds errors on dynamic local memory

static local memory

  Q.submit([&](sycl::handler &h) {
    h.parallel_for<class MyKernel>(
        sycl::nd_range<1>(N, M), [=](sycl::nd_item<1> item) {
          // static local memory
          auto& ref = *sycl::ext::oneapi::group_local_memory<int[M]>(item.get_group(), 1, 2, 3, ...);
          ref[item.get_local_linear_id()] = 42;
        });
  });

dynamic local memory

  Q.submit([&](sycl::handler &h) {
    h.parallel_for<class MyKernel>(
        // dynamic local memory
        sycl::local_accessor<int> ref{M, h};
        sycl::nd_range<1>(N, M), [=](sycl::nd_item<1> item) {
          ref[item.get_local_linear_id()] = 42;
        });
  });

2. The total number of workgroups is limited on GPU

We allocated shadow memory of local memory for each work group, so if the total number of workgroups is very large, it will trigger out-of-resource errors.

3. The kernel is executed in sequence, and cannot concurrently be running in one thread

4. On GPU, sanitizer cannot be enabled on both host and device; On CPU, sanitizer can only be worked on when both host and device sanitizer is enabled