[DeviceSanitizer] Support nullpointer detection & enable GPU tests

libdevice/sanitizer_utils.cpp

@@ -55,6 +55,8 @@ __spirv_GenericCastToPtrExplicit_ToPrivate(void *, int);
 
 extern SYCL_EXTERNAL __attribute__((convergent)) void
 __spirv_ControlBarrier(uint32_t Execution, uint32_t Memory, uint32_t Semantics);
+
+extern "C" SYCL_EXTERNAL void __devicelib_exit();
 #endif // __USE_SPIR_BUILTIN__
 
 static const __SYCL_CONSTANT__ char __asan_shadow_value_start[] =
@@ -104,8 +106,8 @@ enum ADDRESS_SPACE : uint32_t {
 
 namespace {
 
-bool __asan_report_unknown_device();
-bool __asan_report_out_of_shadow_bounds();
+void __asan_report_unknown_device();
+void __asan_report_out_of_shadow_bounds();
 void __asan_print_shadow_memory(uptr addr, uptr shadow_address, uint32_t as);
 
 __SYCL_GLOBAL__ void *ToGlobal(void *ptr) {
@@ -182,10 +184,11 @@ inline uptr MemToShadow_DG2(uptr addr, uint32_t as) {
                       ((addr & (slm_size - 1)) >> ASAN_SHADOW_SCALE);
 
     if (shadow_ptr > shadow_offset_end) {
-      if (__asan_report_out_of_shadow_bounds()) {
+      if (__AsanDebug) {
         __spirv_ocl_printf(__local_shadow_out_of_bound, addr, shadow_ptr,
                            wg_lid, (uptr)shadow_offset);
       }
+      __asan_report_out_of_shadow_bounds();
       return 0;
     }
     return shadow_ptr;
@@ -215,10 +218,11 @@ inline uptr MemToShadow_DG2(uptr addr, uint32_t as) {
                       ((addr & (ASAN_PRIVATE_SIZE - 1)) >> ASAN_SHADOW_SCALE);
 
     if (shadow_ptr > shadow_offset_end) {
-      if (__asan_report_out_of_shadow_bounds()) {
+      if (__AsanDebug) {
         __spirv_ocl_printf(__private_shadow_out_of_bound, addr, shadow_ptr,
                            WG_LID, (uptr)shadow_offset);
       }
+      __asan_report_out_of_shadow_bounds();
       return 0;
     }
     return shadow_ptr;
@@ -245,10 +249,11 @@ inline uptr MemToShadow_PVC(uptr addr, uint32_t as) {
     }
 
     if (shadow_ptr > __AsanShadowMemoryGlobalEnd) {
-      if (__asan_report_out_of_shadow_bounds()) {
+      if (__AsanDebug) {
         __spirv_ocl_printf(__global_shadow_out_of_bound, addr, shadow_ptr,
                            (uptr)__AsanShadowMemoryGlobalStart);
       }
+      __asan_report_out_of_shadow_bounds();
       return 0;
     }
     return shadow_ptr;
@@ -281,10 +286,11 @@ inline uptr MemToShadow_PVC(uptr addr, uint32_t as) {
                       ((addr & (SLM_SIZE - 1)) >> ASAN_SHADOW_SCALE);
 
     if (shadow_ptr > shadow_offset_end) {
-      if (__asan_report_out_of_shadow_bounds()) {
+      if (__AsanDebug) {
         __spirv_ocl_printf(__local_shadow_out_of_bound, addr, shadow_ptr,
                            wg_lid, (uptr)shadow_offset);
       }
+      __asan_report_out_of_shadow_bounds();
       return 0;
     }
     return shadow_ptr;
@@ -314,10 +320,11 @@ inline uptr MemToShadow_PVC(uptr addr, uint32_t as) {
                       ((addr & (ASAN_PRIVATE_SIZE - 1)) >> ASAN_SHADOW_SCALE);
 
     if (shadow_ptr > shadow_offset_end) {
-      if (__asan_report_out_of_shadow_bounds()) {
+      if (__AsanDebug) {
         __spirv_ocl_printf(__private_shadow_out_of_bound, addr, shadow_ptr,
                            WG_LID, (uptr)shadow_offset);
       }
+      __asan_report_out_of_shadow_bounds();
       return 0;
     }
     return shadow_ptr;
@@ -336,14 +343,13 @@ inline uptr MemToShadow(uptr addr, uint32_t as) {
   } else if (__DeviceType == DeviceType::GPU_DG2) {
     shadow_ptr = MemToShadow_DG2(addr, as);
   } else {
-    if (__asan_report_unknown_device() && __AsanDebug) {
+    if (__AsanDebug) {
       __spirv_ocl_printf(__asan_print_unsupport_device_type, (int)__DeviceType);
     }
-    return shadow_ptr;
+    __asan_report_unknown_device();
+    return 0;
   }
 
-// FIXME: OCL "O2" optimizer doesn't work well with following code
-#if 0
   if (__AsanDebug) {
     if (shadow_ptr) {
       if (as == ADDRESS_SPACE_PRIVATE)
@@ -355,7 +361,6 @@ inline uptr MemToShadow(uptr addr, uint32_t as) {
       __spirv_ocl_printf(__asan_print_shadow_value2, addr, as, shadow_ptr);
     }
   }
-#endif
 
   return shadow_ptr;
 }
@@ -398,7 +403,7 @@ bool MemIsZero(__SYCL_GLOBAL__ const char *beg, uptr size) {
 static __SYCL_CONSTANT__ const char __mem_sanitizer_report[] =
     "[kernel] SanitizerReport (ErrorType=%d, IsRecover=%d)\n";
 
-bool __asan_internal_report_save(DeviceSanitizerErrorType error_type) {
+void __asan_internal_report_save(DeviceSanitizerErrorType error_type) {
   const int Expected = ASAN_REPORT_NONE;
   int Desired = ASAN_REPORT_START;
 
@@ -423,12 +428,11 @@ bool __asan_internal_report_save(DeviceSanitizerErrorType error_type) {
     if (__AsanDebug)
       __spirv_ocl_printf(__mem_sanitizer_report, SanitizerReport.ErrorType,
                          SanitizerReport.IsRecover);
-    return true;
   }
-  return false;
+  __devicelib_exit();
 }
 
-bool __asan_internal_report_save(
+void __asan_internal_report_save(
     uptr ptr, uint32_t as, const char __SYCL_CONSTANT__ *file, uint32_t line,
     const char __SYCL_CONSTANT__ *func, bool is_write, uint32_t access_size,
     DeviceSanitizerMemoryType memory_type, DeviceSanitizerErrorType error_type,
@@ -505,9 +509,8 @@ bool __asan_internal_report_save(
     if (__AsanDebug)
       __spirv_ocl_printf(__mem_sanitizer_report, SanitizerReport.ErrorType,
                          SanitizerReport.IsRecover);
-    return true;
   }
-  return false;
+  __devicelib_exit();
 }
 
 ///
@@ -575,6 +578,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;
   }
@@ -604,13 +610,12 @@ void __asan_report_misalign_error(uptr addr, uint32_t as, size_t size,
                               memory_type, error_type, is_recover);
 }
 
-bool __asan_report_unknown_device() {
-  return __asan_internal_report_save(DeviceSanitizerErrorType::UNKNOWN_DEVICE);
+void __asan_report_unknown_device() {
+  __asan_internal_report_save(DeviceSanitizerErrorType::UNKNOWN_DEVICE);
 }
 
-bool __asan_report_out_of_shadow_bounds() {
-  return __asan_internal_report_save(
-      DeviceSanitizerErrorType::OUT_OF_SHADOW_BOUNDS);
+void __asan_report_out_of_shadow_bounds() {
+  __asan_internal_report_save(DeviceSanitizerErrorType::OUT_OF_SHADOW_BOUNDS);
 }
 
 ///

sycl/cmake/modules/FetchUnifiedRuntime.cmake

@@ -117,13 +117,13 @@ if(SYCL_UR_USE_FETCH_CONTENT)
   endfunction()
 
   set(UNIFIED_RUNTIME_REPO "https://github.com/oneapi-src/unified-runtime.git")
-  # commit cfecab08e6e6dbb694f614b4f6271a258a41fc10
-  # Merge: 10fd78c1 5bebef5d
+  # commit 4517290650a9938537666e6409fb8e0db73ff4d8
+  # Merge: 6298474e 3dbb7a2a
   # Author: Omar Ahmed <[email protected]>
-  # Date:   Tue Sep 17 12:26:35 2024 +0100
-  #   Merge pull request #1874 from PietroGhg/pietro/membarrier
-  #   [NATIVECPU] Support atomic fence queries
-  set(UNIFIED_RUNTIME_TAG cfecab08e6e6dbb694f614b4f6271a258a41fc10)
+  # Date:   Wed Sep 18 08:48:03 2024 +0100
+  #     Merge pull request #1914 from AllanZyne/review/yang/dsan_nullpointer
+  #     [DeviceSanitizer] Support nullpointer detection
+  set(UNIFIED_RUNTIME_TAG 4517290650a9938537666e6409fb8e0db73ff4d8)
 
   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/lit.local.cfg

@@ -10,5 +10,5 @@ config.substitutions.append(
 
 config.unsupported_features += ['cuda', 'hip']
 
-# FIXME: Skip gen devices, waiting for gfx driver uplifting
-config.unsupported_features += ['gpu-intel-gen9', 'gpu-intel-gen11', 'gpu-intel-gen12', 'gpu-intel-dg2', 'gpu-intel-pvc']
+# FIXME: Skip some of gpu devices, waiting for gfx driver uplifting
+config.unsupported_features += ['gpu-intel-gen9', 'gpu-intel-gen11', 'gpu-intel-gen12', 'gpu-intel-pvc']

sycl/test-e2e/AddressSanitizer/nullpointer/global_nullptr.cpp

@@ -0,0 +1,32 @@
+// REQUIRES: linux
+// RUN: %{build} %device_asan_flags -O0 -g -o %t
+// RUN: %{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 = nullptr;
+
+  Q.submit([&](sycl::handler &h) {
+    h.parallel_for<class MyKernel>(
+        sycl::nd_range<1>(N, 1), [=](sycl::nd_item<1> item) {
+          auto private_array =
+              sycl::ext::oneapi::experimental::static_address_cast<
+                  sycl::access::address_space::private_space,
+                  sycl::access::decorated::no>(array);
+          private_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: {{.*global_nullptr.cpp}}:[[@LINE-5]]
+
+  return 0;
+}

sycl/test-e2e/AddressSanitizer/nullpointer/private_nullptr.cpp

@@ -0,0 +1,36 @@
+// REQUIRES: linux
+// RUN: %{build} %device_asan_flags -O0 -g -o %t
+// RUN: %{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
+
+// FIXME: There's an issue in gfx driver, so this test pending here.
+// XFAIL: *
+
+#include <sycl/detail/core.hpp>
+#include <sycl/ext/oneapi/experimental/address_cast.hpp>
+
+int main() {
+  sycl::queue Q;
+  constexpr std::size_t N = 4;
+  int *array = nullptr;
+
+  Q.submit([&](sycl::handler &h) {
+    h.parallel_for<class MyKernel>(
+        sycl::nd_range<1>(N, 1), [=](sycl::nd_item<1> item) {
+          auto private_array =
+              sycl::ext::oneapi::experimental::static_address_cast<
+                  sycl::access::address_space::private_space,
+                  sycl::access::decorated::no>(array);
+          private_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: {{.*private_nullptr.cpp}}:[[@LINE-5]]
+
+  return 0;
+}