[SYCL][DeviceSanitizer] Checking out-of-bounds error on sycl::local_accessor

[AllanZyne]

UR: oneapi-src/unified-runtime#1532

To check sycl::local_accessor(aka, dynamic local memory), we need to extend a new argument in spir kernel, this is because:

• ASan needs to know some size information of local buffer, like its size and size with redzone, so that it can poison its shadow memory
• By using this new argument, we can also pass some per-launch information (that is, it is different in each launch of kernel). One obvious example is SanitizerReport, which saves the error message, so that we can store and print multiple error reports for one kernel with different arguments. Another example is the shadow memory of local memory, this should be different per-launch as well, since one kernel can be launched multiple times and executed in parallel.

I named this argument as "__asan_launch", which is a pointer pointed to "LaunchInfo" structure and allocated it in shared USM. To make this pointer can be used in spir_func w/o extending their argument, I created a global external local memory (external, so that it can be shared with other translation units, and its instance is defined in libdevice), and save the "__asan_launch" into this local memory immediately at the entry of kernel.

UR can't check the name of kernel arguments, so it can't know if the kernel has "__asan_launch". So I assume the "__asan_launch" is always there, and added a check to prevent DAE pass from removing it.

[AllanZyne]

Hi @intel/compute-runtime-maintain, can you help to review?
Thanks very much!

[kbenzie]

UR LGTM

[kbenzie]

There are failing tests on this job.

This is likely to hold up the UR merge pipeline so I may need to revert the UR change in oneapi-src/unified-runtime#1532

********************
Failed Tests (18):
  SYCL :: AddressSanitizer/common/demangle-kernel-name.cpp
  SYCL :: AddressSanitizer/common/kernel-debug.cpp
  SYCL :: AddressSanitizer/out-of-bounds/DeviceGlobal/device_global.cpp
  SYCL :: AddressSanitizer/out-of-bounds/DeviceGlobal/device_global_image_scope.cpp
  SYCL :: AddressSanitizer/out-of-bounds/DeviceGlobal/device_global_image_scope_unaligned.cpp
  SYCL :: AddressSanitizer/out-of-bounds/DeviceGlobal/multi_device_images.cpp
  SYCL :: AddressSanitizer/out-of-bounds/USM/parallel_for_char.cpp
  SYCL :: AddressSanitizer/out-of-bounds/USM/parallel_for_double.cpp
  SYCL :: AddressSanitizer/out-of-bounds/USM/parallel_for_func.cpp
  SYCL :: AddressSanitizer/out-of-bounds/USM/parallel_for_int.cpp
  SYCL :: AddressSanitizer/out-of-bounds/USM/parallel_for_short.cpp
  SYCL :: AddressSanitizer/out-of-bounds/local/group_local_memory.cpp
  SYCL :: AddressSanitizer/out-of-bounds/local/local_accessor_basic.cpp
  SYCL :: AddressSanitizer/out-of-bounds/local/local_accessor_function.cpp
  SYCL :: AddressSanitizer/out-of-bounds/local/local_accessor_multiargs.cpp
  SYCL :: AddressSanitizer/out-of-bounds/local/multiple_source.cpp
  SYCL :: AddressSanitizer/use-after-free/quarantine-no-free.cpp
  SYCL :: AddressSanitizer/use-after-free/use-after-free.cpp

[kbenzie]

@AllanZyne please attempt to fix these test failures when you start today, if they are not fixed when I'm back at work tomorrow morning in Europe I'll go ahead with the revert of oneapi-src/unified-runtime#1532.

[AllanZyne]

@AllanZyne please attempt to fix these test failures when you start today, if they are not fixed when I'm back at work tomorrow morning in Europe I'll go ahead with the revert of oneapi-src/unified-runtime#1532.

@kbenzie It seems like the main tag of UR is not correct.

[kbenzie]

@kbenzie It seems like the main tag of UR is not correct.

🤦 thanks for fixing it

[kbenzie]

@intel/llvm-gatekeepers please merge

[steffenlarsen]

@intel/dpcpp-tools-reviewers approval is needed.

[kbenzie]

@intel/dpcpp-tools-reviewers approval is needed.

I'm struggling to see where this is visible in the UI. Isn't there usually a sheild for a required reviewer teams?

Edit: I see it now. I'll keep this in mind in future.

[kbenzie]

@AlexeySachkov have your review comments been addressed?

[AllanZyne]

Thank @kbenzie for your help!

[AllanZyne]

@intel/llvm-gatekeepers could you please merge this PR? Thanks!