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

[AllanZyne]

UR: https://github.com/oneapi-src/unified-runtime/pull/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/dpcpp-tools-reviewers @intel/unified-runtime-reviewers @intel/llvm-reviewers-runtime, please review. Thanks very much!

[AllanZyne]

I think we also should check whether the device supports aspect::usm_device_allocations. I think you can do that by adding aspect-usm_device_allocations to the REQUIRES lines for each test.

The devices (CPU/PVC/DG2) which I specify to be tested on must already support USM. This is an implicit aspect of test devices. So I don't think I need to add this aspect.

[maarquitos14]

I think we also should check whether the device supports aspect::usm_device_allocations. I think you can do that by adding aspect-usm_device_allocations to the REQUIRES lines for each test.

The devices (CPU/PVC/DG2) which I specify to be tested on must already support USM. This is an implicit aspect of test devices. So I don't think I need to add this aspect.

Sorry, I don't see the requirement that these tests need to be CPU/PVC/DG2 to run, but even with that, I don't think it hurts to check that it supports usm_device_allocations, so I would add it to be on the safe side.

[AllanZyne]

I think we also should check whether the device supports aspect::usm_device_allocations. I think you can do that by adding aspect-usm_device_allocations to the REQUIRES lines for each test.

The devices (CPU/PVC/DG2) which I specify to be tested on must already support USM. This is an implicit aspect of test devices. So I don't think I need to add this aspect.

Sorry, I don't see the requirement that these tests need to be CPU/PVC/DG2 to run, but even with that, I don't think it hurts to check that it supports usm_device_allocations, so I would add it to be on the safe side.

After reconsidering I still don't think I should add them. Adding "usm_device_allocations" will make these tests look like they support any device that has "usm_device_allocations" aspect. But the truth is it only supports limited types of device.

If you're interested, I added more specific device type check on this PR https://github.com/intel/llvm/pull/13450 https://github.com/intel/llvm/pull/13450/files#diff-a48a43565418aed9195e4ebaa034fe1d617eb835a4969ea0222b99b1185e79db.

[maarquitos14]

I think we also should check whether the device supports aspect::usm_device_allocations. I think you can do that by adding aspect-usm_device_allocations to the REQUIRES lines for each test.

The devices (CPU/PVC/DG2) which I specify to be tested on must already support USM. This is an implicit aspect of test devices. So I don't think I need to add this aspect.

Sorry, I don't see the requirement that these tests need to be CPU/PVC/DG2 to run, but even with that, I don't think it hurts to check that it supports usm_device_allocations, so I would add it to be on the safe side.

After reconsidering I still don't think I should add them. Adding "usm_device_allocations" will make these tests look like they support any device that has "usm_device_allocations" aspect. But the truth is it only supports limited types of device.

If you're interested, I added more specific device type check on this PR #13450 https://github.com/intel/llvm/pull/13450/files#diff-a48a43565418aed9195e4ebaa034fe1d617eb835a4969ea0222b99b1185e79db.

Okay, I see your point. I'm just slightly concerned that if we ever add a new device to the config file that doesn't support this aspect, it will fail. Could we at least add a comment either in the config file saying something along the lines of Only add devices that support usm_device_allocations aspect or in the test files saying This test assumes it can only run in CPU/PVC/DG2 devices, which support usm_device_allocations aspect, so that if anybody ever adds a new device, they will know?

[AllanZyne]

I think we also should check whether the device supports aspect::usm_device_allocations. I think you can do that by adding aspect-usm_device_allocations to the REQUIRES lines for each test.

The devices (CPU/PVC/DG2) which I specify to be tested on must already support USM. This is an implicit aspect of test devices. So I don't think I need to add this aspect.

Sorry, I don't see the requirement that these tests need to be CPU/PVC/DG2 to run, but even with that, I don't think it hurts to check that it supports usm_device_allocations, so I would add it to be on the safe side.

After reconsidering I still don't think I should add them. Adding "usm_device_allocations" will make these tests look like they support any device that has "usm_device_allocations" aspect. But the truth is it only supports limited types of device. If you're interested, I added more specific device type check on this PR #13450 https://github.com/intel/llvm/pull/13450/files#diff-a48a43565418aed9195e4ebaa034fe1d617eb835a4969ea0222b99b1185e79db.

Okay, I see your point. I'm just slightly concerned that if we ever add a new device to the config file that doesn't support this aspect, it will fail. Could we at least add a comment either in the config file saying something along the lines of Only add devices that support usm_device_allocations aspect or in the test files saying This test assumes it can only run in CPU/PVC/DG2 devices, which support usm_device_allocations aspect, so that if anybody ever adds a new device, they will know?

Only add devices that support usm_device_allocations aspect

To enable device sanitizer, the device which supports "usm_device_allocations" is not enough. Currently, we develop dsan device by device.

Ok, I'll add This test assumes it can only run in CPU/PVC/DG2 devices, which support usm_device_allocations aspect this message.

[AllanZyne]

Hi @intel/dpcpp-tools-reviewers @intel/unified-runtime-reviewers Please review this PR, thanks very much!

[AllanZyne]

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

[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 https://github.com/oneapi-src/unified-runtime/pull/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 https://github.com/oneapi-src/unified-runtime/pull/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!