[DeviceSanitizer] Support check invalid kernel argument (#14512)

UR: oneapi-src/unified-runtime#1843

also fix parallel_for_double.cpp

---------

Co-authored-by: omarahmed1111 <[email protected]>

Author: AllanZyne

sycl/cmake/modules/FetchUnifiedRuntime.cmake

@@ -116,13 +116,13 @@ if(SYCL_PI_UR_USE_FETCH_CONTENT)
   endfunction()
 
   set(UNIFIED_RUNTIME_REPO "https://github.com/oneapi-src/unified-runtime.git")
-  # commit a89657c2e36d7152820b9ecc088422a54ca1d844
-  # Merge: 2355a7d6 be7057c4
-  # Author: Omar Ahmed <omar.ahmed@codeplay.com>
-  # Date:   Wed Aug 7 12:01:33 2024 +0100
-  #     Merge pull request #1699 from PietroGhg/pietro/usm_fixes
-  #     [NATIVECPU] Implement urUSMGetMemAllocInfo and aligned alloc
-  set(UNIFIED_RUNTIME_TAG a89657c2e36d7152820b9ecc088422a54ca1d844)
+  # commit 2d3524e74ab6a09d1ad1a96d977973fac32ca55d
+  # Merge: 6b2e678d d6e93fa1
+  # Author: Kenneth Benzie (Benie) <k.benzie@codeplay.com>
+  # Date:   Wed Aug 7 14:23:09 2024 +0100
+  #     Merge pull request #1930 from oneapi-src/benie/no-import-in-pragma-region
+  #     Make pragma region names joined by _
+  set(UNIFIED_RUNTIME_TAG 2d3524e74ab6a09d1ad1a96d977973fac32ca55d)
 
   set(UMF_BUILD_EXAMPLES OFF CACHE INTERNAL "EXAMPLES")
   # Due to the use of dependentloadflag and no installer for UMF and hwloc we need

sycl/test-e2e/AddressSanitizer/double-free/double-free.cpp

@@ -1,10 +1,10 @@
 // REQUIRES: linux
 // RUN: %{build} %device_asan_flags -O0 -g -o %t
-// RUN: %force_device_asan_rt UR_LAYER_ASAN_OPTIONS=quarantine_size_mb:1 %{run} not %t 2>&1 | FileCheck --check-prefixes CHECK,CHECK-DEVICE %s
+// RUN: %force_device_asan_rt UR_LAYER_ASAN_OPTIONS="quarantine_size_mb:1;detect_kernel_arguments:0" %{run} not %t 2>&1 | FileCheck --check-prefixes CHECK,CHECK-DEVICE %s
 // RUN: %{build} %device_asan_flags -DMALLOC_HOST -O0 -g -o %t
-// RUN: %force_device_asan_rt UR_LAYER_ASAN_OPTIONS=quarantine_size_mb:1 %{run} not %t 2>&1 | FileCheck --check-prefixes CHECK,CHECK-HOST %s
+// RUN: %force_device_asan_rt UR_LAYER_ASAN_OPTIONS="quarantine_size_mb:1;detect_kernel_arguments:0" %{run} not %t 2>&1 | FileCheck --check-prefixes CHECK,CHECK-HOST %s
 // RUN: %{build} %device_asan_flags -DMALLOC_SHARED -O0 -g -o %t
-// RUN: %force_device_asan_rt UR_LAYER_ASAN_OPTIONS=quarantine_size_mb:1 %{run} not %t 2>&1 | FileCheck --check-prefixes CHECK,CHECK-SHARED %s
+// RUN: %force_device_asan_rt UR_LAYER_ASAN_OPTIONS="quarantine_size_mb:1;detect_kernel_arguments:0" %{run} not %t 2>&1 | FileCheck --check-prefixes CHECK,CHECK-SHARED %s
 #include <sycl/usm.hpp>
 
 constexpr size_t N = 64;

sycl/test-e2e/AddressSanitizer/invalid-argument/bad-context.cpp

@@ -0,0 +1,25 @@
+// REQUIRES: linux, (gpu && level_zero), cpu
+// RUN: %{build} %device_asan_flags -DMALLOC_DEVICE -O2 -g -o %t
+// RUN: env SYCL_PREFER_UR=1 UR_LAYER_ASAN_OPTIONS="detect_kernel_arguments:1" %{run} not %t 2>&1 | FileCheck --check-prefixes CHECK,CHECK-DEVICE %s
+
+#include <sycl/detail/core.hpp>
+
+#include <sycl/usm.hpp>
+
+int main() {
+  sycl::queue gpu_queue(sycl::gpu_selector_v);
+  sycl::queue cpu_queue(sycl::cpu_selector_v);
+
+  auto data = sycl::malloc_device<int>(1, cpu_queue);
+
+  gpu_queue.submit([&](sycl::handler &h) {
+    h.single_task<class MyKernel>([=]() { *data = 0; });
+  });
+  gpu_queue.wait();
+  // CHECK: DeviceSanitizer: invalid-argument on kernel
+  // CHECK: The 1th argument {{.*}} is allocated in other context
+  // CHECK: {{.*}} is located inside of Device USM region
+
+  sycl::free(data, cpu_queue);
+  return 0;
+}

sycl/test-e2e/AddressSanitizer/invalid-argument/host-pointer.cpp

@@ -0,0 +1,31 @@
+// REQUIRES: linux
+// RUN: %{build} %device_asan_flags -O2 -g -o %t
+// RUN: env SYCL_PREFER_UR=1 UR_LAYER_ASAN_OPTIONS="detect_kernel_arguments:1" ONEAPI_DEVICE_SELECTOR=level_zero:gpu %{run-unfiltered-devices} not %t 2>&1 | FileCheck --check-prefixes CHECK-GPU %s
+// RUN: env SYCL_PREFER_UR=1 UR_LAYER_ASAN_OPTIONS="detect_kernel_arguments:1" ONEAPI_DEVICE_SELECTOR=opencl:cpu %{run-unfiltered-devices} %t 2>&1 | FileCheck --check-prefixes CHECK-CPU %s
+
+#include <sycl/detail/core.hpp>
+#include <sycl/usm.hpp>
+
+///
+/// GPU devices don't support shared system USM currently, so passing host
+/// pointer to kernel is invalid.
+/// CPU devices support shared system USM.
+///
+
+int main() {
+  sycl::queue Q;
+  auto hostPtr = new int;
+
+  Q.submit([&](sycl::handler &h) {
+    h.single_task<class MyKernel>([=]() { *hostPtr = 0; });
+  });
+  Q.wait();
+
+  // CHECK-GPU: ERROR: DeviceSanitizer: invalid-argument
+  // CHECK-GPU: The 1th argument {{.*}} is not a USM pointer
+  // CHECK-CPU-NOT: ERROR: DeviceSanitizer: invalid-argument
+
+  delete hostPtr;
+  puts("PASS\n");
+  return 0;
+}

sycl/test-e2e/AddressSanitizer/invalid-argument/out-of-bounds.cpp

@@ -0,0 +1,22 @@
+// REQUIRES: linux, gpu
+// RUN: %{build} %device_asan_flags -O2 -g -o %t
+// RUN: env SYCL_PREFER_UR=1 UR_LAYER_ASAN_OPTIONS="detect_kernel_arguments:1" %{run} not %t 2>&1 | FileCheck %s
+
+#include <sycl/detail/core.hpp>
+
+#include <sycl/usm.hpp>
+
+int main() {
+  sycl::queue Q;
+  auto Data = sycl::malloc_device<int>(1, Q);
+  ++Data;
+
+  Q.submit([&](sycl::handler &h) {
+    h.single_task<class MyKernel>([=]() { *Data = 0; });
+  });
+  Q.wait();
+  // CHECK: ERROR: DeviceSanitizer: invalid-argument
+  // CHECK: The 1th argument {{.*}} is located outside of its region
+
+  return 0;
+}

sycl/test-e2e/AddressSanitizer/invalid-argument/released-pointer.cpp

@@ -0,0 +1,25 @@
+// REQUIRES: linux, gpu
+// RUN: %{build} %device_asan_flags -O2 -g -o %t
+// RUN: env SYCL_PREFER_UR=1 UR_LAYER_ASAN_OPTIONS="quarantine_size_mb:1;detect_kernel_arguments:1" %{run} not %t 2>&1 | FileCheck %s
+
+#include <sycl/detail/core.hpp>
+
+#include <sycl/usm.hpp>
+
+int main() {
+  sycl::queue Q;
+  auto Data = sycl::malloc_device<int>(1, Q);
+  sycl::free(Data, Q);
+
+  Q.submit([&](sycl::handler &h) {
+    h.single_task<class MyKernel>([=]() { *Data = 0; });
+  });
+  Q.wait();
+  // CHECK: ERROR: DeviceSanitizer: invalid-argument
+  // CHECK: The 1th argument {{.*}} is a released USM pointer
+  // CHECK: {{.*}} is located inside of Device USM region
+  // CHECK: allocated here:
+  // CHECK: freed here:
+
+  return 0;
+}

sycl/test-e2e/AddressSanitizer/use-after-free/quarantine-no-free.cpp

@@ -1,6 +1,6 @@
 // REQUIRES: linux
 // RUN: %{build} %device_asan_flags -O0 -g -o %t
-// RUN: env UR_LAYER_ASAN_OPTIONS=quarantine_size_mb:5 UR_LOG_SANITIZER=level:info %{run} not %t 2>&1 | FileCheck %s
+// RUN: env UR_LAYER_ASAN_OPTIONS="quarantine_size_mb:5;detect_kernel_arguments:0" UR_LOG_SANITIZER=level:info %{run} not %t 2>&1 | FileCheck %s
 #include <sycl/usm.hpp>
 
 /// Quarantine Cache Test
@@ -46,7 +46,7 @@ int main() {
   // CHECK:   #0 {{.*}} {{.*quarantine-no-free.cpp}}:[[@LINE-5]]
   // CHECK: [[ADDR]] is located inside of Device USM region [{{0x.*}}, {{0x.*}})
   // CHECK: allocated here:
-  // CHECK: released here:
+  // CHECK: freed here:
 
   return 0;
 }

sycl/test-e2e/AddressSanitizer/use-after-free/use-after-free.cpp

@@ -1,6 +1,6 @@
 // REQUIRES: linux
 // RUN: %{build} %device_asan_flags -O0 -g -o %t
-// RUN: env UR_LAYER_ASAN_OPTIONS=quarantine_size_mb:1 %{run} not %t 2>&1 | FileCheck %s
+// RUN: env UR_LAYER_ASAN_OPTIONS="quarantine_size_mb:5;detect_kernel_arguments:0" %{run} not %t 2>&1 | FileCheck %s
 #include <sycl/usm.hpp>
 
 constexpr size_t N = 1024;
@@ -21,7 +21,7 @@ int main() {
   // CHECK:   #0 {{.*}} {{.*use-after-free.cpp:}}[[@LINE-5]]
   // CHECK: [[ADDR]] is located inside of Device USM region [{{0x.*}}, {{0x.*}})
   // CHECK: allocated here:
-  // CHECK: released here:
+  // CHECK: freed here:
 
   return 0;
 }