[DeviceSanitizer] Support detecting out-of-bounds error on DeviceGlobals

libdevice/include/device-sanitizer-report.hpp

@@ -8,9 +8,6 @@
 //===----------------------------------------------------------------------===//
 #pragma once
 
-// Treat this header as system one to workaround frontend's restriction
-#pragma clang system_header
-
 #include <cinttypes>
 
 enum class DeviceSanitizerErrorType : int32_t {
@@ -29,6 +26,7 @@ enum class DeviceSanitizerMemoryType : int32_t {
   LOCAL,
   PRIVATE,
   MEM_BUFFER,
+  DEVICE_GLOBAL,
 };
 
 // NOTE Layout of this structure should be aligned with the one in

libdevice/include/sanitizer_device_utils.hpp

@@ -10,16 +10,10 @@
 #include "spir_global_var.hpp"
 #include <cstdint>
 
-// Treat this header as system one to workaround frontend's restriction
-#pragma clang system_header
-
 template <typename T>
 class
 #ifdef __SYCL_DEVICE_ONLY__
-    [[__sycl_detail__::global_variable_allowed, __sycl_detail__::device_global,
-      __sycl_detail__::add_ir_attributes_global_variable(
-          "sycl-device-global-size", "sycl-device-image-scope", sizeof(T),
-          nullptr)]]
+    [[__sycl_detail__::global_variable_allowed, __sycl_detail__::device_global]]
 #endif
     DeviceGlobal {
 public:

libdevice/sanitizer_utils.cpp

@@ -54,6 +54,7 @@ const int kUsmDeviceRedzoneMagic = (char)0x81;
 const int kUsmHostRedzoneMagic = (char)0x82;
 const int kUsmSharedRedzoneMagic = (char)0x83;
 const int kMemBufferRedzoneMagic = (char)0x84;
+const int kDeviceGlobalRedZoneMagic = (char)0x85;
 
 const int kUsmDeviceDeallocatedMagic = (char)0x91;
 const int kUsmHostDeallocatedMagic = (char)0x92;
@@ -366,6 +367,10 @@ void __asan_report_access_error(uptr addr, int32_t as, size_t size,
     memory_type = DeviceSanitizerMemoryType::LOCAL;
     error_type = DeviceSanitizerErrorType::OUT_OF_BOUND;
     break;
+  case kDeviceGlobalRedZoneMagic:
+    memory_type = DeviceSanitizerMemoryType::DEVICE_GLOBAL;
+    error_type = DeviceSanitizerErrorType::OUT_OF_BOUND;
+    break;
   default:
     memory_type = DeviceSanitizerMemoryType::UNKNOWN;
     error_type = DeviceSanitizerErrorType::UNKNOWN;

llvm/tools/sycl-post-link/sycl-post-link.cpp

@@ -883,6 +883,114 @@ static bool removeDeviceGlobalFromCompilerUsed(Module &M) {
   return true;
 }
 
+// Add extra red zone to each image scope device globals if the module has been
+// instrumented by sanitizer pass. And record their infomation like size, red
+// zone size, beginning address.
+static bool instrumentDeviceGlobal(Module &M) {
+  if (!M.getNamedGlobal("__DeviceSanitizerReportMem"))
+    return false;
+
+  auto &DL = M.getDataLayout();
+  IRBuilder<> IRB(M.getContext());
+  SmallVector<GlobalVariable *, 8> GlobalsToRemove;
+  SmallVector<GlobalVariable *, 8> NewDeviceGlobals;
+  SmallVector<Constant *, 8> DeviceGlobalMetadata;
+
+
+  constexpr uint64_t MaxRZ = 1 << 18;
+  const uint64_t MinRZ = 32;
+
+  Type *IntptrTy =
+      Type::getIntNTy(M.getContext(), M.getDataLayout().getPointerSizeInBits());
+
+  // Device global meta data is described by a structure
+  //  size_t device_global_size
+  //  size_t device_global_size_with_red_zone
+  //  size_t beginning address of the device global
+  StructType *StructTy = StructType::get(IntptrTy, IntptrTy, IntptrTy);
+
+  for (auto &G : M.globals()) {
+    // Non image scope device globals are implemented by device USM, and the
+    // out-of-bounds check for them will be done by sanitizer USM part. So we
+    // exclude them here.
+    if (isDeviceGlobalVariable(G) && hasDeviceImageScopeProperty(G)) {
+      Type *Ty = G.getValueType();
+      const uint64_t SizeInBytes = DL.getTypeAllocSize(Ty);
+      const uint64_t RightRedzoneSize = [&] {
+        uint64_t RZ = 0;
+        if (SizeInBytes <= MinRZ / 2) {
+          // Reduce redzone size for small size objects, e.g. int, char[1].
+          // Optimize when SizeInBytes is less than or equal to half of MinRZ.
+          RZ = MinRZ - SizeInBytes;
+        } else {
+          // Calculate RZ, where MinRZ <= RZ <= MaxRZ, and RZ ~ 1/4 *
+          // SizeInBytes.
+          RZ = std::clamp((SizeInBytes / MinRZ / 4) * MinRZ, MinRZ, MaxRZ);
+
+          // Round up to multiple of MinRZ.
+          if (SizeInBytes % MinRZ)
+            RZ += MinRZ - (SizeInBytes % MinRZ);
+        }
+
+        assert((RZ + SizeInBytes) % MinRZ == 0);
+        return RZ;
+      }();
+      Type *RightRedZoneTy = ArrayType::get(IRB.getInt8Ty(), RightRedzoneSize);
+      StructType *NewTy = StructType::get(Ty, RightRedZoneTy);
+      Constant *NewInitializer = ConstantStruct::get(
+          NewTy, G.getInitializer(), Constant::getNullValue(RightRedZoneTy));
+
+      // Create a new global variable with enough space for a redzone.
+      GlobalVariable *NewGlobal = new GlobalVariable(
+          M, NewTy, G.isConstant(), G.getLinkage(), NewInitializer, "", &G,
+          G.getThreadLocalMode(), G.getAddressSpace());
+      NewGlobal->copyAttributesFrom(&G);
+      NewGlobal->setComdat(G.getComdat());
+      NewGlobal->setAlignment(Align(MinRZ));
+      NewGlobal->copyMetadata(&G, 0);
+
+      Value *Indices2[2];
+      Indices2[0] = IRB.getInt32(0);
+      Indices2[1] = IRB.getInt32(0);
+
+      G.replaceAllUsesWith(
+          ConstantExpr::getGetElementPtr(NewTy, NewGlobal, Indices2, true));
+      NewGlobal->takeName(&G);
+      GlobalsToRemove.push_back(&G);
+      NewDeviceGlobals.push_back(NewGlobal);
+      DeviceGlobalMetadata.push_back(ConstantStruct::get(
+          StructTy, ConstantInt::get(IntptrTy, SizeInBytes),
+          ConstantInt::get(IntptrTy, SizeInBytes + RightRedzoneSize),
+          ConstantExpr::getPointerCast(NewGlobal, IntptrTy)));
+    }
+  }
+
+  if (GlobalsToRemove.empty())
+    return false;
+
+  // Create global to record number of device globals
+  GlobalVariable *NumOfDeviceGlobals = new GlobalVariable(
+      M, IntptrTy, false, GlobalValue::ExternalLinkage,
+      ConstantInt::get(IntptrTy, NewDeviceGlobals.size()),
+      "__AsanDeviceGlobalCount", nullptr, GlobalValue::NotThreadLocal, 1);
+  NumOfDeviceGlobals->setUnnamedAddr(GlobalValue::UnnamedAddr::Local);
+
+  // Create meta data global to record device globals' information
+  ArrayType *ArrayTy = ArrayType::get(StructTy, NewDeviceGlobals.size());
+  Constant *MetadataInitializer =
+      ConstantArray::get(ArrayTy, DeviceGlobalMetadata);
+  GlobalVariable *AsanDeviceGlobalMetadata = new GlobalVariable(
+      M, MetadataInitializer->getType(), false, GlobalValue::ExternalLinkage,
+      MetadataInitializer, "__AsanDeviceGlobalMetadata", nullptr,
+      GlobalValue::NotThreadLocal, 1);
+  AsanDeviceGlobalMetadata->setUnnamedAddr(GlobalValue::UnnamedAddr::Local);
+
+  for (auto *G : GlobalsToRemove)
+    G->eraseFromParent();
+
+  return true;
+}
+
 SmallVector<module_split::ModuleDesc, 2>
 handleESIMD(module_split::ModuleDesc &&MDesc, bool &Modified,
             bool &SplitOccurred) {
@@ -980,6 +1088,10 @@ processInputModule(std::unique_ptr<Module> M) {
   // "llvm.compiler.used" they can be erased safely.
   Modified |= removeDeviceGlobalFromCompilerUsed(*M.get());
 
+  // Instrument each image scope device globals if the module has been
+  // instrumented by sanitizer pass.
+  Modified |= instrumentDeviceGlobal(*M.get());
+
   // Do invoke_simd processing before splitting because this:
   // - saves processing time (the pass is run once, even though on larger IR)
   // - doing it before SYCL/ESIMD splitting is required for correctness

sycl/test-e2e/AddressSanitizer/out-of-bounds/DeviceGlobal/device_global.cpp

@@ -0,0 +1,24 @@
+// REQUIRES: linux, cpu
+// RUN: %{build} %device_sanitizer_flags -O2 -g -o %t.out
+// RUN: env SYCL_PREFER_UR=1 ONEAPI_DEVICE_SELECTOR=opencl:cpu %{run-unfiltered-devices} not %t.out 2>&1 | FileCheck %s
+
+#include <sycl/sycl.hpp>
+
+using namespace sycl;
+using namespace sycl::ext::oneapi;
+using namespace sycl::ext::oneapi::experimental;
+
+sycl::ext::oneapi::experimental::device_global<char[5]> dev_global;
+
+int main() {
+  sycl::queue Q;
+
+  Q.submit([&](sycl::handler &h) {
+     h.single_task<class Test>([=]() { dev_global[8] = 42; });
+     // CHECK: ERROR: DeviceSanitizer: out-of-bounds-access on Device USM
+     // CHECK: {{WRITE of size 1 at kernel <.*Test> LID\(0, 0, 0\) GID\(0, 0, 0\)}}
+     // CHECK: {{#0 .* .*device_global.cpp:}}[[@LINE-3]]
+   }).wait();
+
+  return 0;
+}

sycl/test-e2e/AddressSanitizer/out-of-bounds/DeviceGlobal/device_global_image_scope.cpp

@@ -0,0 +1,31 @@
+// REQUIRES: linux, cpu
+// RUN: %{build} %device_sanitizer_flags -O2 -g -o %t.out
+// RUN: env SYCL_PREFER_UR=1 ONEAPI_DEVICE_SELECTOR=opencl:cpu %{run-unfiltered-devices} not %t.out 2>&1 | FileCheck %s
+
+#include <sycl/sycl.hpp>
+
+using namespace sycl;
+using namespace sycl::ext::oneapi;
+using namespace sycl::ext::oneapi::experimental;
+
+sycl::ext::oneapi::experimental::device_global<
+    int[4], decltype(properties(device_image_scope, host_access_read_write))>
+    dev_global;
+
+int main() {
+  sycl::queue Q;
+
+  Q.submit([&](sycl::handler &h) {
+    h.single_task<class Test>([=]() {
+      dev_global[4] = 42;
+      // CHECK: ERROR: DeviceSanitizer: out-of-bounds-access on Device Global
+      // CHECK: {{WRITE of size 4 at kernel <.*Test> LID\(0, 0, 0\) GID\(0, 0, 0\)}}
+      // CHECK: {{#0 .* .*device_global_image_scope.cpp:}}[[@LINE-3]]
+    });
+  }).wait();
+
+  int val;
+  Q.copy(dev_global, &val).wait();
+
+  return 0;
+}

sycl/test-e2e/AddressSanitizer/out-of-bounds/DeviceGlobal/device_global_image_scope_unaligned.cpp

@@ -0,0 +1,29 @@
+// REQUIRES: linux, cpu
+// RUN: %{build} %device_sanitizer_flags -O2 -g -o %t.out
+// RUN: env SYCL_PREFER_UR=1 ONEAPI_DEVICE_SELECTOR=opencl:cpu %{run-unfiltered-devices} not %t.out 2>&1 | FileCheck %s
+
+#include <sycl/sycl.hpp>
+
+using namespace sycl;
+using namespace sycl::ext::oneapi;
+using namespace sycl::ext::oneapi::experimental;
+
+sycl::ext::oneapi::experimental::device_global<
+    char[5], decltype(properties(device_image_scope, host_access_read_write))>
+    dev_global;
+
+int main() {
+  sycl::queue Q;
+
+  Q.submit([&](sycl::handler &h) {
+     h.single_task<class Test>([=]() { dev_global[8] = 42; });
+     // CHECK: ERROR: DeviceSanitizer: out-of-bounds-access on Device Global
+     // CHECK: {{WRITE of size 1 at kernel <.*Test> LID\(0, 0, 0\) GID\(0, 0, 0\)}}
+     // CHECK: {{#0 .* .*device_global_image_scope_unaligned.cpp:}}[[@LINE-3]]
+   }).wait();
+
+  char val;
+  Q.copy(dev_global, &val).wait();
+
+  return 0;
+}

sycl/test-e2e/AddressSanitizer/out-of-bounds/DeviceGlobal/multi_device_images.cpp

@@ -0,0 +1,54 @@
+// REQUIRES: linux, cpu
+// RUN: %{build} %device_sanitizer_flags -O2 -g -DUSER_CODE_1 -c -o %t1.o
+// RUN: %{build} %device_sanitizer_flags -O2 -g -DUSER_CODE_2 -c -o %t2.o
+// RUN: %clangxx -fsycl %device_sanitizer_flags -O2 -g %t1.o %t2.o -o %t.out
+// RUN: env SYCL_PREFER_UR=1 ONEAPI_DEVICE_SELECTOR=opencl:cpu %{run-unfiltered-devices} not %t.out 2>&1 | FileCheck %s
+
+#include <sycl/sycl.hpp>
+
+using namespace sycl;
+using namespace sycl::ext::oneapi;
+using namespace sycl::ext::oneapi::experimental;
+
+#ifdef USER_CODE_1
+
+sycl::ext::oneapi::experimental::device_global<
+    int[4], decltype(properties(device_image_scope, host_access_read_write))>
+    dev_global2;
+
+void foo() {
+  sycl::queue Q;
+
+  Q.submit([&](sycl::handler &h) {
+    h.single_task<class Test2>([=]() {
+      dev_global2[4] = 42;
+      // CHECK: ERROR: DeviceSanitizer: out-of-bounds-access on Device Global
+      // CHECK: {{WRITE of size 4 at kernel <.*Test2> LID\(0, 0, 0\) GID\(0, 0, 0\)}}
+      // CHECK: {{#0 .* .*multi_device_images.cpp:}}[[@LINE-3]]
+    });
+  }).wait();
+}
+
+#else
+
+sycl::ext::oneapi::experimental::device_global<
+    int, decltype(properties(device_image_scope, host_access_read_write))>
+    dev_global;
+
+extern void foo();
+
+int main() {
+  sycl::queue Q;
+
+  Q.submit([&](sycl::handler &h) {
+    h.single_task<class Test1>([=]() {
+      dev_global = 42;
+    });
+  }).wait();
+
+  foo();
+
+  return 0;
+}
+
+#endif

sycl/test-e2e/AddressSanitizer/out-of-bounds/USM/parallel_for_char.cpp

@@ -30,9 +30,9 @@ int main() {
         [=](sycl::nd_item<1> item) { ++array[item.get_global_id(0)]; });
   });
   Q.wait();
-  // CHECK-DEVICE: ERROR: DeviceSanitizer: out-of-bounds-access on USM Device Memory
-  // CHECK-HOST:   ERROR: DeviceSanitizer: out-of-bounds-access on USM Host Memory
-  // CHECK-SHARED: ERROR: DeviceSanitizer: out-of-bounds-access on USM Shared Memory
+  // CHECK-DEVICE: ERROR: DeviceSanitizer: out-of-bounds-access on Device USM
+  // CHECK-HOST:   ERROR: DeviceSanitizer: out-of-bounds-access on Host USM
+  // CHECK-SHARED: ERROR: DeviceSanitizer: out-of-bounds-access on Shared USM
   // CHECK: {{READ of size 1 at kernel <.*MyKernelR_4> LID\(0, 0, 0\) GID\(12345, 0, 0\)}}
   // CHECK: {{  #0 .* .*parallel_for_char.cpp:}}[[@LINE-7]]
 

sycl/test-e2e/AddressSanitizer/out-of-bounds/USM/parallel_for_double.cpp

@@ -30,9 +30,9 @@ int main() {
         [=](sycl::nd_item<1> item) { ++array[item.get_global_id(0)]; });
   });
   Q.wait();
-  // CHECK-DEVICE: ERROR: DeviceSanitizer: out-of-bounds-access on USM Device Memory
-  // CHECK-HOST:   ERROR: DeviceSanitizer: out-of-bounds-access on USM Host Memory
-  // CHECK-SHARED: ERROR: DeviceSanitizer: out-of-bounds-access on USM Shared Memory
+  // CHECK-DEVICE: ERROR: DeviceSanitizer: out-of-bounds-access on Device USM
+  // CHECK-HOST:   ERROR: DeviceSanitizer: out-of-bounds-access on Host USM
+  // CHECK-SHARED: ERROR: DeviceSanitizer: out-of-bounds-access on Shared USM
   // CHECK: {{READ of size 8 at kernel <.*MyKernelR_4> LID\(0, 0, 0\) GID\(123456, 0, 0\)}}
   // CHECK: {{  #0 .* .*parallel_for_double.cpp:}}[[@LINE-7]]
 

sycl/test-e2e/AddressSanitizer/out-of-bounds/USM/parallel_for_func.cpp

@@ -10,9 +10,9 @@
 #include <sycl/sycl.hpp>
 
 __attribute__((noinline)) void foo(int *array, size_t i) { array[i] = 1; }
-// CHECK-DEVICE: ERROR: DeviceSanitizer: out-of-bounds-access on USM Device Memory
-// CHECK-HOST:   ERROR: DeviceSanitizer: out-of-bounds-access on USM Host Memory
-// CHECK-SHARED: ERROR: DeviceSanitizer: out-of-bounds-access on USM Shared Memory
+// CHECK-DEVICE: ERROR: DeviceSanitizer: out-of-bounds-access on Device USM
+// CHECK-HOST:   ERROR: DeviceSanitizer: out-of-bounds-access on Host USM
+// CHECK-SHARED: ERROR: DeviceSanitizer: out-of-bounds-access on Shared USM
 // CHECK: {{WRITE of size 4 at kernel <.*MyKernel> LID\(0, 0, 0\) GID\(123, 0, 0\)}}
 // CHECK: {{  #0 foo\(int\*, unsigned long\) .*parallel_for_func.cpp:}}[[@LINE-5]]
 

sycl/test-e2e/AddressSanitizer/out-of-bounds/USM/parallel_for_int.cpp

@@ -30,9 +30,9 @@ int main() {
         [=](sycl::nd_item<1> item) { ++array[item.get_global_id(0)]; });
   });
   Q.wait();
-  // CHECK-DEVICE: ERROR: DeviceSanitizer: out-of-bounds-access on USM Device Memory
-  // CHECK-HOST:   ERROR: DeviceSanitizer: out-of-bounds-access on USM Host Memory
-  // CHECK-SHARED: ERROR: DeviceSanitizer: out-of-bounds-access on USM Shared Memory
+  // CHECK-DEVICE: ERROR: DeviceSanitizer: out-of-bounds-access on Device USM
+  // CHECK-HOST:   ERROR: DeviceSanitizer: out-of-bounds-access on Host USM
+  // CHECK-SHARED: ERROR: DeviceSanitizer: out-of-bounds-access on Shared USM
   // CHECK: {{READ of size 4 at kernel <.*MyKernelR_4> LID\(0, 0, 0\) GID\(1234567, 0, 0\)}}
   // CHECK: {{  #0 .* .*parallel_for_int.cpp:}}[[@LINE-7]]
 

sycl/test-e2e/AddressSanitizer/out-of-bounds/USM/parallel_for_short.cpp

@@ -30,9 +30,9 @@ int main() {
         [=](sycl::nd_item<1> item) { ++array[item.get_global_id(0)]; });
   });
   Q.wait();
-  // CHECK-DEVICE: ERROR: DeviceSanitizer: out-of-bounds-access on USM Device Memory
-  // CHECK-HOST:   ERROR: DeviceSanitizer: out-of-bounds-access on USM Host Memory
-  // CHECK-SHARED: ERROR: DeviceSanitizer: out-of-bounds-access on USM Shared Memory
+  // CHECK-DEVICE: ERROR: DeviceSanitizer: out-of-bounds-access on Device USM
+  // CHECK-HOST:   ERROR: DeviceSanitizer: out-of-bounds-access on Host USM
+  // CHECK-SHARED: ERROR: DeviceSanitizer: out-of-bounds-access on Shared USM
   // CHECK: {{READ of size 2 at kernel <.*MyKernelR_4> LID\(0, 0, 0\) GID\(123456789, 0, 0\)}}
   // CHECK: {{  #0 .* .*parallel_for_short.cpp:}}[[@LINE-7]]