[DeviceSanitizer] Support nullpointer detection & enable GPU tests

libdevice/sanitizer_utils.cpp

@@ -52,6 +52,8 @@ extern SYCL_EXTERNAL __SYCL_LOCAL__ void *
 __spirv_GenericCastToPtrExplicit_ToLocal(void *, int);
 extern SYCL_EXTERNAL __SYCL_PRIVATE__ void *
 __spirv_GenericCastToPtrExplicit_ToPrivate(void *, int);
+
+extern "C" SYCL_EXTERNAL void __devicelib_exit();
 #endif // __USE_SPIR_BUILTIN__
 
 static const __SYCL_CONSTANT__ char __asan_shadow_value_start[] =
@@ -422,6 +424,7 @@ bool __asan_internal_report_save(DeviceSanitizerErrorType error_type) {
                          SanitizerReport.IsRecover);
     return true;
   }
+  __devicelib_exit();
   return false;
 }
 
@@ -504,6 +507,7 @@ bool __asan_internal_report_save(
                          SanitizerReport.IsRecover);
     return true;
   }
+  __devicelib_exit();
   return false;
 }
 
@@ -572,6 +576,9 @@ void __asan_report_access_error(uptr addr, uint32_t as, size_t size,
   case kUsmSharedDeallocatedMagic:
     error_type = DeviceSanitizerErrorType::USE_AFTER_FREE;
     break;
+  case kNullPointerRedzoneMagic:
+    error_type = DeviceSanitizerErrorType::NULL_POINTER;
+    break;
   default:
     error_type = DeviceSanitizerErrorType::UNKNOWN;
   }

sycl/cmake/modules/FetchUnifiedRuntime.cmake

@@ -115,13 +115,14 @@ if(SYCL_PI_UR_USE_FETCH_CONTENT)
       CACHE PATH "Path to external '${name}' adapter source dir" FORCE)
   endfunction()
 
-  set(UNIFIED_RUNTIME_REPO "https://github.com/oneapi-src/unified-runtime.git")
-  # commit bc1a28ede0df7f837047b632e00437587672c134
-  # Author: Omar Ahmed <omar.ahmed@codeplay.com>
-  # Date:   Mon Jul 29 16:44:58 2024 +0100
-  #     Merge pull request #1819 from DBDuncan/sean/rename-interop-to-external
-  #    [Bindless][Exp] Rename interop related structs/funcs with "external"
-  set(UNIFIED_RUNTIME_TAG bc1a28ede0df7f837047b632e00437587672c134)
+  set(UNIFIED_RUNTIME_REPO "https://github.com/AllanZyne/unified-runtime.git")
+  # commit a985a81dc9ba8adfcc8b54e35ad287e97766fb3e
+  # Merge: b7b0c8b3 f772f907
+  # Author: Piotr Balcer <piotr.balcer@intel.com>
+  # Date:   Mon Jul 29 09:11:29 2024 +0200
+  #     Merge pull request #1905 from igchor/umf_hwloc_disable
+  #     Bump UMF version to allow disabling hwloc
+  set(UNIFIED_RUNTIME_TAG review/yang/dsan_nullpointer)
 
   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/nullpointer/nullpointer.cpp

@@ -0,0 +1,26 @@
+// REQUIRES: linux
+// RUN-NOT: %{build} %device_asan_flags -O0 -g -o %t
+// RUN-NOT: %{run} not %t 2>&1 | FileCheck %s
+// RUN: %{build} %device_asan_flags -O1 -g -o %t
+// RUN: %{run} not %t 2>&1 | FileCheck %s
+// RUN: %{build} %device_asan_flags -O2 -g -o %t
+// RUN: %{run} not %t 2>&1 | FileCheck %s
+
+#include <sycl/detail/core.hpp>
+
+int main() {
+  sycl::queue Q;
+  constexpr std::size_t N = 4;
+  int *array = 0;
+
+  Q.submit([&](sycl::handler &h) {
+    h.parallel_for<class MyKernel>(
+        sycl::nd_range<1>(N, 1), [=](sycl::nd_item<1> item) { array[0] = 0; });
+    Q.wait();
+  });
+  // CHECK: ERROR: DeviceSanitizer: null-pointer-access on Unknown Memory
+  // CHECK: WRITE of size 4 at kernel {{<.*MyKernel>}} LID(0, 0, 0) GID({{.*}}, 0, 0)
+  // CHECK:   #0 {{.*}} {{.*null-pointer.cpp}}:[[@LINE-5]]
+
+  return 0;
+}