[DeviceSanitizer] Support detecting misaligned access error (#14148)

UR: oneapi-src/unified-runtime#1747

---------

Co-authored-by: Kenneth Benzie (Benie) <k.benzie@codeplay.com>

libdevice/sanitizer_utils.cpp

@@ -322,6 +322,9 @@ bool MemIsZero(__SYCL_GLOBAL__ const char *beg, uptr size) {
 /// ASAN Save Report
 ///
 
+static __SYCL_CONSTANT__ const char __mem_sanitizer_report[] =
+ "[kernel] SanitizerReport (ErrorType=%d, IsRecover=%d)\n";
+
 bool __asan_internal_report_save(DeviceSanitizerErrorType error_type) {
  const int Expected = ASAN_REPORT_NONE;
  int Desired = ASAN_REPORT_START;
@@ -339,8 +342,14 @@ bool __asan_internal_report_save(DeviceSanitizerErrorType error_type) {
  if (atomicCompareAndSet(&SanitizerReport.Flag, Desired, Expected) ==
  Expected) {
  SanitizerReport.ErrorType = error_type;
+ SanitizerReport.IsRecover = false;
+
  // Show we've done copying
  atomicStore(&SanitizerReport.Flag, ASAN_REPORT_FINISH);
+
+ if (__AsanDebug)
+ __spirv_ocl_printf(__mem_sanitizer_report, SanitizerReport.ErrorType,
+ SanitizerReport.IsRecover);
  return true;
  }
  return false;
@@ -419,6 +428,10 @@ bool __asan_internal_report_save(
 
  // Show we've done copying
  atomicStore(&SanitizerReport.Flag, ASAN_REPORT_FINISH);
+
+ if (__AsanDebug)
+ __spirv_ocl_printf(__mem_sanitizer_report, SanitizerReport.ErrorType,
+ SanitizerReport.IsRecover);
  return true;
  }
  return false;
@@ -428,6 +441,32 @@ bool __asan_internal_report_save(
 /// ASAN Error Reporters
 ///
 
+DeviceSanitizerMemoryType GetMemoryTypeByShadowValue(int shadow_value) {
+ switch (shadow_value) {
+ case kUsmDeviceRedzoneMagic:
+ case kUsmDeviceDeallocatedMagic:
+ return DeviceSanitizerMemoryType::USM_DEVICE;
+ case kUsmHostRedzoneMagic:
+ case kUsmHostDeallocatedMagic:
+ return DeviceSanitizerMemoryType::USM_HOST;
+ case kUsmSharedRedzoneMagic:
+ case kUsmSharedDeallocatedMagic:
+ return DeviceSanitizerMemoryType::USM_SHARED;
+ case kPrivateLeftRedzoneMagic:
+ case kPrivateMidRedzoneMagic:
+ case kPrivateRightRedzoneMagic:
+ return DeviceSanitizerMemoryType::PRIVATE;
+ case kMemBufferRedzoneMagic:
+ return DeviceSanitizerMemoryType::MEM_BUFFER;
+ case kSharedLocalRedzoneMagic:
+ return DeviceSanitizerMemoryType::LOCAL;
+ case kDeviceGlobalRedzoneMagic:
+ return DeviceSanitizerMemoryType::DEVICE_GLOBAL;
+ default:
+ return DeviceSanitizerMemoryType::UNKNOWN;
+ }
+}
+
 void __asan_report_access_error(uptr addr, uint32_t as, size_t size,
  bool is_write, uptr poisoned_addr,
  const char __SYCL_CONSTANT__ *file,
@@ -442,61 +481,56 @@ void __asan_report_access_error(uptr addr, uint32_t as, size_t size,
  }
  // FIXME: check if shadow_address out-of-bound
 
- DeviceSanitizerMemoryType memory_type;
+ DeviceSanitizerMemoryType memory_type =
+ GetMemoryTypeByShadowValue(shadow_value);
  DeviceSanitizerErrorType error_type;
 
  switch (shadow_value) {
  case kUsmDeviceRedzoneMagic:
- memory_type = DeviceSanitizerMemoryType::USM_DEVICE;
- error_type = DeviceSanitizerErrorType::OUT_OF_BOUNDS;
- break;
  case kUsmHostRedzoneMagic:
- memory_type = DeviceSanitizerMemoryType::USM_HOST;
- error_type = DeviceSanitizerErrorType::OUT_OF_BOUNDS;
- break;
  case kUsmSharedRedzoneMagic:
- memory_type = DeviceSanitizerMemoryType::USM_SHARED;
- error_type = DeviceSanitizerErrorType::OUT_OF_BOUNDS;
- break;
- case kUsmDeviceDeallocatedMagic:
- memory_type = DeviceSanitizerMemoryType::USM_DEVICE;
- error_type = DeviceSanitizerErrorType::USE_AFTER_FREE;
- break;
- case kUsmHostDeallocatedMagic:
- memory_type = DeviceSanitizerMemoryType::USM_HOST;
- error_type = DeviceSanitizerErrorType::USE_AFTER_FREE;
- break;
- case kUsmSharedDeallocatedMagic:
- memory_type = DeviceSanitizerMemoryType::USM_SHARED;
- error_type = DeviceSanitizerErrorType::USE_AFTER_FREE;
- break;
  case kPrivateLeftRedzoneMagic:
  case kPrivateMidRedzoneMagic:
  case kPrivateRightRedzoneMagic:
- memory_type = DeviceSanitizerMemoryType::PRIVATE;
- error_type = DeviceSanitizerErrorType::OUT_OF_BOUNDS;
- break;
  case kMemBufferRedzoneMagic:
- memory_type = DeviceSanitizerMemoryType::MEM_BUFFER;
- error_type = DeviceSanitizerErrorType::OUT_OF_BOUNDS;
- break;
  case kSharedLocalRedzoneMagic:
- memory_type = DeviceSanitizerMemoryType::LOCAL;
- error_type = DeviceSanitizerErrorType::OUT_OF_BOUNDS;
- break;
  case kDeviceGlobalRedzoneMagic:
- memory_type = DeviceSanitizerMemoryType::DEVICE_GLOBAL;
  error_type = DeviceSanitizerErrorType::OUT_OF_BOUNDS;
  break;
+ case kUsmDeviceDeallocatedMagic:
+ case kUsmHostDeallocatedMagic:
+ case kUsmSharedDeallocatedMagic:
+ error_type = DeviceSanitizerErrorType::USE_AFTER_FREE;
+ break;
  default:
- memory_type = DeviceSanitizerMemoryType::UNKNOWN;
  error_type = DeviceSanitizerErrorType::UNKNOWN;
  }
 
  __asan_internal_report_save(addr, as, file, line, func, is_write, size,
  memory_type, error_type, is_recover);
 }
 
+void __asan_report_misalign_error(uptr addr, uint32_t as, size_t size,
+ bool is_write, uptr poisoned_addr,
+ const char __SYCL_CONSTANT__ *file,
+ uint32_t line,
+ const char __SYCL_CONSTANT__ *func,
+ bool is_recover = false) {
+
+ auto *shadow = (__SYCL_GLOBAL__ s8 *)MemToShadow(addr, as);
+ while (*shadow >= 0) {
+ ++shadow;
+ }
+ int shadow_value = *shadow;
+
+ DeviceSanitizerErrorType error_type = DeviceSanitizerErrorType::MISALIGNED;
+ DeviceSanitizerMemoryType memory_type =
+ GetMemoryTypeByShadowValue(shadow_value);
+
+ __asan_internal_report_save(addr, as, file, line, func, is_write, size,
+ memory_type, error_type, is_recover);
+}
+
 bool __asan_report_unknown_device() {
  return __asan_internal_report_save(DeviceSanitizerErrorType::UNKNOWN_DEVICE);
 }
@@ -589,6 +623,8 @@ inline uptr __asan_region_is_poisoned(uptr beg, uint32_t as, size_t size) {
  return 0;
 }
 
+constexpr size_t AlignMask(size_t n) { return n - 1; }
+
 } // namespace
 
 ///
@@ -599,6 +635,10 @@ inline uptr __asan_region_is_poisoned(uptr beg, uint32_t as, size_t size) {
  DEVICE_EXTERN_C_NOINLINE void __asan_##type##size( \
  uptr addr, uint32_t as, const char __SYCL_CONSTANT__ *file, \
  uint32_t line, const char __SYCL_CONSTANT__ *func) { \
+ if (addr & AlignMask(size)) { \
+ __asan_report_misalign_error(addr, as, size, is_write, addr, file, line, \
+ func); \
+ } \
  if (__asan_address_is_poisoned(addr, as, size)) { \
  __asan_report_access_error(addr, as, size, is_write, addr, file, line, \
  func); \
@@ -607,6 +647,10 @@ inline uptr __asan_region_is_poisoned(uptr beg, uint32_t as, size_t size) {
  DEVICE_EXTERN_C_NOINLINE void __asan_##type##size##_noabort( \
  uptr addr, uint32_t as, const char __SYCL_CONSTANT__ *file, \
  uint32_t line, const char __SYCL_CONSTANT__ *func) { \
+ if (addr & AlignMask(size)) { \
+ __asan_report_misalign_error(addr, as, size, is_write, addr, file, line, \
+ func, true); \
+ } \
  if (__asan_address_is_poisoned(addr, as, size)) { \
  __asan_report_access_error(addr, as, size, is_write, addr, file, line, \
  func, true); \

sycl/plugins/unified_runtime/CMakeLists.txt

@@ -100,13 +100,13 @@ if(SYCL_PI_UR_USE_FETCH_CONTENT)
  endfunction()
 
  set(UNIFIED_RUNTIME_REPO "https://github.com/oneapi-src/unified-runtime.git")
- # commit 642e3437b8a0a3e702ac3d68513d7076e7618c94
- # Merge: 342bdfff 199287b7
+ # commit 167ddf9320dba95c9324e373fb863baeb46e09f1
+ # Merge: 76c6bf9a 0cd10f63
  # Author: Kenneth Benzie (Benie) <k.benzie@codeplay.com>
- # Date: Thu Jul 11 10:37:50 2024 +0100
- # Merge pull request #1829 from callumfare/callum/bindless_workaround
- # Treat bindless images handle types as native handles
- set(UNIFIED_RUNTIME_TAG 642e3437b8a0a3e702ac3d68513d7076e7618c94)
+ # Date: Thu Jul 11 15:12:38 2024 +0100
+ # Merge pull request #1747 from AllanZyne/review/yang/misalign_access
+ # [DeviceSanitizer] Support detecting misaligned access error
+ set(UNIFIED_RUNTIME_TAG 167ddf9320dba95c9324e373fb863baeb46e09f1)
 
  fetch_adapter_source(level_zero
  ${UNIFIED_RUNTIME_REPO}

sycl/test-e2e/AddressSanitizer/misaligned/misalign-int.cpp

@@ -0,0 +1,35 @@
+// REQUIRES: linux, cpu
+// RUN: %{build} %device_asan_flags -O0 -g -o %t
+// RUN: env SYCL_PREFER_UR=1 %{run} not %t 2>&1 | FileCheck %s
+// RUN: %{build} %device_asan_flags -O1 -g -o %t
+// RUN: env SYCL_PREFER_UR=1 %{run} not %t 2>&1 | FileCheck %s
+// RUN: %{build} %device_asan_flags -O2 -g -o %t
+// RUN: env SYCL_PREFER_UR=1 %{run} not %t 2>&1 | FileCheck %s
+#include <sycl/detail/core.hpp>
+#include <sycl/usm.hpp>
+
+#include <random>
+
+int main() {
+ std::random_device rd;
+ std::mt19937 gen(rd());
+ std::uniform_int_distribution<> distrib(1, 3);
+
+ sycl::queue Q;
+ constexpr std::size_t N = 4;
+ auto *array = sycl::malloc_host<int>(N, Q);
+ auto offset = distrib(gen);
+ std::cout << "offset: " << offset << std::endl;
+ array = (int *)((char *)array + offset);
+
+ Q.submit([&](sycl::handler &h) {
+ h.parallel_for<class MyKernel>(sycl::nd_range<1>(N, 1),
+ [=](sycl::nd_item<1> item) { ++array[0]; });
+ Q.wait();
+ });
+ // CHECK: ERROR: DeviceSanitizer: misaligned-access on Host USM
+ // CHECK: READ of size 4 at kernel {{<.*MyKernel>}} LID(0, 0, 0) GID({{.*}}, 0, 0)
+ // CHECK: #0 {{.*}} {{.*misalign-int.cpp}}:[[@LINE-5]]
+
+ return 0;
+}

sycl/test-e2e/AddressSanitizer/misaligned/misalign-long.cpp

@@ -0,0 +1,35 @@
+// REQUIRES: linux, cpu
+// RUN: %{build} %device_asan_flags -O0 -g -o %t
+// RUN: env SYCL_PREFER_UR=1 %{run} not %t 2>&1 | FileCheck %s
+// RUN: %{build} %device_asan_flags -O1 -g -o %t
+// RUN: env SYCL_PREFER_UR=1 %{run} not %t 2>&1 | FileCheck %s
+// RUN: %{build} %device_asan_flags -O2 -g -o %t
+// RUN: env SYCL_PREFER_UR=1 %{run} not %t 2>&1 | FileCheck %s
+#include <sycl/detail/core.hpp>
+#include <sycl/usm.hpp>
+
+#include <random>
+
+int main() {
+ std::random_device rd;
+ std::mt19937 gen(rd());
+ std::uniform_int_distribution<> distrib(1, 7);
+
+ sycl::queue Q;
+ constexpr std::size_t N = 4;
+ auto *array = sycl::malloc_shared<long long>(N, Q);
+ auto offset = distrib(gen);
+ std::cout << "offset: " << offset << std::endl;
+ array = (long long *)((char *)array + offset);
+
+ Q.submit([&](sycl::handler &h) {
+ h.parallel_for<class MyKernel>(sycl::nd_range<1>(N, 1),
+ [=](sycl::nd_item<1> item) { ++array[0]; });
+ Q.wait();
+ });
+ // CHECK: ERROR: DeviceSanitizer: misaligned-access on Shared USM
+ // CHECK: READ of size 8 at kernel {{<.*MyKernel>}} LID(0, 0, 0) GID({{.*}}, 0, 0)
+ // CHECK: #0 {{.*}} {{.*misalign-long.cpp}}:[[@LINE-5]]
+
+ return 0;
+}

sycl/test-e2e/AddressSanitizer/misaligned/misalign-short.cpp

@@ -0,0 +1,29 @@
+// REQUIRES: linux, cpu
+// RUN: %{build} %device_asan_flags -O0 -g -o %t
+// RUN: env SYCL_PREFER_UR=1 %{run} not %t 2>&1 | FileCheck %s
+// RUN: %{build} %device_asan_flags -O1 -g -o %t
+// RUN: env SYCL_PREFER_UR=1 %{run} not %t 2>&1 | FileCheck %s
+// RUN: %{build} %device_asan_flags -O2 -g -o %t
+// RUN: env SYCL_PREFER_UR=1 %{run} not %t 2>&1 | FileCheck %s
+#include <sycl/detail/core.hpp>
+#include <sycl/usm.hpp>
+
+#include <random>
+
+int main() {
+ sycl::queue Q;
+ constexpr std::size_t N = 4;
+ auto *array = sycl::malloc_device<short>(N, Q);
+ array = (short *)((char *)array + 1);
+
+ Q.submit([&](sycl::handler &h) {
+ h.parallel_for<class MyKernel>(sycl::nd_range<1>(N, 1),
+ [=](sycl::nd_item<1> item) { ++array[0]; });
+ Q.wait();
+ });
+ // CHECK: ERROR: DeviceSanitizer: misaligned-access on Device USM
+ // CHECK: READ of size 2 at kernel {{<.*MyKernel>}} LID(0, 0, 0) GID({{.*}}, 0, 0)
+ // CHECK: #0 {{.*}} {{.*misalign-short.cpp}}:[[@LINE-5]]
+
+ return 0;
+}