Review/yang/sanitizer private

buildbot/configure.py

@@ -181,6 +181,8 @@ def do_configure(args):
         "-DSYCL_ENABLE_KERNEL_FUSION={}".format(sycl_enable_fusion),
         "-DSYCL_ENABLE_MAJOR_RELEASE_PREVIEW_LIB={}".format(sycl_preview_lib),
         "-DBUG_REPORT_URL=https://github.com/intel/llvm/issues",
+        "-DSYCL_PI_UR_USE_FETCH_CONTENT=OFF",
+        "-DSYCL_PI_UR_SOURCE_DIR=/localdisk2/yzhao/work/sycl_workspace/unified-runtime"
     ]
 
     if args.l0_headers and args.l0_loader:

llvm/lib/Transforms/Instrumentation/AddressSanitizer.cpp

@@ -87,6 +87,7 @@
 #include <sstream>
 #include <string>
 #include <tuple>
+#include <unordered_map>
 
 using namespace llvm;
 
@@ -1184,13 +1185,96 @@ AddressSanitizerPass::AddressSanitizerPass(
 
 PreservedAnalyses AddressSanitizerPass::run(Module &M,
                                             ModuleAnalysisManager &MAM) {
+  // M.dump();
   ModuleAddressSanitizer ModuleSanitizer(
       M, Options.InsertVersionCheck, Options.CompileKernel, Options.Recover,
       UseGlobalGC, UseOdrIndicator, DestructorKind, ConstructorKind);
   bool Modified = false;
   auto &FAM = MAM.getResult<FunctionAnalysisManagerModuleProxy>(M).getManager();
   const StackSafetyGlobalInfo *const SSGI =
       ClUseStackSafety ? &MAM.getResult<StackSafetyGlobalAnalysis>(M) : nullptr;
+
+  SmallVector<Function *> SpirKernels;
+  // SmallVector<Function> SpirFuncs;
+  for (Function &F : M) {
+    if (F.getCallingConv() == CallingConv::SPIR_KERNEL) {
+      SpirKernels.emplace_back(&F);
+    }
+  }
+
+  int LongSize = M.getDataLayout().getPointerSizeInBits();
+  auto* IntptrTy = Type::getIntNTy(M.getContext(), LongSize);
+
+  for (auto* F : SpirKernels) {
+    SmallVector<Type *, 16> Types;
+    for (Function::const_arg_iterator I = F->arg_begin(), E = F->arg_end();
+         I != E; ++I) {
+      Types.push_back(I->getType());
+    }
+
+    // New Argument
+    Types.push_back(IntptrTy);
+
+    FunctionType *NewFTy = FunctionType::get(F->getReturnType(), Types, false);
+
+    std::string OrigFuncName = F->getName().str();
+    F->setName(OrigFuncName + "_del");
+
+    Function *NewF =
+        Function::Create(NewFTy, F->getLinkage(), OrigFuncName, F->getParent());
+    NewF->copyAttributesFrom(F);
+    NewF->copyMetadata(F, 0);
+    NewF->setCallingConv(F->getCallingConv());
+    NewF->setDSOLocal(F->isDSOLocal());
+
+    // Set original arguments' names.
+    Function::arg_iterator NewI = NewF->arg_begin();
+    for (Function::const_arg_iterator I = F->arg_begin(), E = F->arg_end();
+         I != E; ++I, ++NewI) {
+      NewI->setName(I->getName());
+    }
+
+    NewF->splice(NewF->begin(), F);
+    assert(F->isDeclaration() &&
+           "splice does not work, original function body is not empty!");
+
+    NewF->setSubprogram(F->getSubprogram());
+
+    NewF->setComdat(F->getComdat());
+    F->setComdat(nullptr);
+
+    F->deleteBody();
+
+    for (Function::arg_iterator I = F->arg_begin(), E = F->arg_end(),
+                                NI = NewF->arg_begin();
+         I != E; ++I, ++NI) {
+      I->replaceAllUsesWith(&*NI);
+    }
+
+    // Fixup metadata
+     IRBuilder<> Builder(M.getContext());
+
+    auto FixupMetadata = [&NewF](StringRef MDName, Constant* NewV) {
+      auto *Node = NewF->getMetadata(MDName);
+      if (!Node)
+        return;
+      SmallVector<Metadata *, 8> NewMD;
+      for (unsigned I = 0; I < Node->getNumOperands(); ++I) {
+        NewMD.emplace_back(Node->getOperand(I));
+      }
+      NewMD.emplace_back(ConstantAsMetadata::get(NewV));
+      NewF->setMetadata(MDName, llvm::MDNode::get(NewF->getContext(), NewMD));
+    };
+
+    FixupMetadata("kernel_arg_buffer_location", Builder.getInt32(-1));
+    FixupMetadata("kernel_arg_runtime_aligned", Builder.getFalse());
+    FixupMetadata("kernel_arg_exclusive_ptr", Builder.getFalse());
+
+    F->removeFromParent();
+  }
+
+  M.dump();
+
   for (Function &F : M) {
     AddressSanitizer FunctionSanitizer(
         M, SSGI, Options.InstrumentationWithCallsThreshold,
@@ -1306,6 +1390,10 @@ void AddressSanitizer::AppendDebugInfoToArgs(Instruction *InsertBefore,
   auto *FuncNameGV = GetOrCreateGlobalString(*M, "__asan_func",
                                              demangle(FuncName), ConstantAS);
   Args.push_back(ConstantExpr::getPointerCast(FuncNameGV, ConstASPtrTy));
+
+  // Launch Data
+  auto* F = InsertBefore->getFunction();
+  Args.push_back(F->getArg(F->arg_size() - 1));
 }
 
 Value *AddressSanitizer::memToShadow(Value *Shadow, IRBuilder<> &IRB) {
@@ -1368,11 +1456,13 @@ void AddressSanitizer::instrumentSyclAllocateLocalMemory(CallInst *CI) {
       IRB.CreateCall(CI->getCalledFunction(),
                      {SizeWithRedZone, ConstantInt::get(IntptrTy, Align)});
 
+  auto* F = CI->getFunction();
+
   /// __asan_set_shadow_local_memory(uptr beg, size_t size, size_t
-  /// size_with_redzone)
+  /// size_with_redzone, launch_info)
   IRB.CreateCall(
       AsanSetShadowDeviceLocalFunc,
-      {IRB.CreatePointerCast(NewCI, IntptrTy), Size, SizeWithRedZone});
+      {IRB.CreatePointerCast(NewCI, IntptrTy), Size, SizeWithRedZone, F->getArg(F->arg_size() - 1)});
 
   CI->replaceAllUsesWith(NewCI);
   CI->eraseFromParent();
@@ -2830,20 +2920,22 @@ bool ModuleAddressSanitizer::instrumentModule(Module &M) {
   // Put the constructor and destructor in comdat if both
   // (1) global instrumentation is not TU-specific
   // (2) target is ELF.
-  if (UseCtorComdat && TargetTriple.isOSBinFormatELF() && CtorComdat) {
-    if (AsanCtorFunction) {
-      AsanCtorFunction->setComdat(M.getOrInsertComdat(kAsanModuleCtorName));
-      appendToGlobalCtors(M, AsanCtorFunction, Priority, AsanCtorFunction);
-    }
-    if (AsanDtorFunction) {
-      AsanDtorFunction->setComdat(M.getOrInsertComdat(kAsanModuleDtorName));
-      appendToGlobalDtors(M, AsanDtorFunction, Priority, AsanDtorFunction);
+  if (!TargetTriple.isSPIR()) { // SPIR kernel needn't AsanCtorFunction & AsanDtorFunction
+    if (UseCtorComdat && TargetTriple.isOSBinFormatELF() && CtorComdat) {
+      if (AsanCtorFunction) {
+        AsanCtorFunction->setComdat(M.getOrInsertComdat(kAsanModuleCtorName));
+        appendToGlobalCtors(M, AsanCtorFunction, Priority, AsanCtorFunction);
+      }
+      if (AsanDtorFunction) {
+        AsanDtorFunction->setComdat(M.getOrInsertComdat(kAsanModuleDtorName));
+        appendToGlobalDtors(M, AsanDtorFunction, Priority, AsanDtorFunction);
+      }
+    } else {
+      if (AsanCtorFunction)
+        appendToGlobalCtors(M, AsanCtorFunction, Priority);
+      if (AsanDtorFunction)
+        appendToGlobalDtors(M, AsanDtorFunction, Priority);
     }
-  } else {
-    if (AsanCtorFunction)
-      appendToGlobalCtors(M, AsanCtorFunction, Priority);
-    if (AsanDtorFunction)
-      appendToGlobalDtors(M, AsanDtorFunction, Priority);
   }
 
   return true;
@@ -2873,8 +2965,7 @@ void AddressSanitizer::initializeCallbacks(Module &M, const TargetLibraryInfo *T
         }
       }
 
-      // Extend __asan_load/store arguments: unsigned int address_space, char*
-      // file, unsigned int line, char* func
+      // Extend __asan_load/store(unsigned int address_space, char* file, unsigned int line, char* func, void* launch_data)
       if (TargetTriple.isSPIR()) {
         constexpr unsigned ConstantAS = 2;
         auto *Int8PtrTy = Type::getInt8Ty(*C)->getPointerTo(ConstantAS);
@@ -2883,11 +2974,13 @@ void AddressSanitizer::initializeCallbacks(Module &M, const TargetLibraryInfo *T
         Args1.push_back(Int8PtrTy);            // file
         Args1.push_back(Type::getInt32Ty(*C)); // line
         Args1.push_back(Int8PtrTy);            // func
+        Args1.push_back(IntptrTy); // launch_data
 
         Args2.push_back(Type::getInt32Ty(*C)); // address_space
         Args2.push_back(Int8PtrTy);            // file
         Args2.push_back(Type::getInt32Ty(*C)); // line
         Args2.push_back(Int8PtrTy);            // func
+        Args2.push_back(IntptrTy); // launch_data
       }
       AsanErrorCallbackSized[AccessIsWrite][Exp] = M.getOrInsertFunction(
           kAsanReportErrorTemplate + ExpStr + TypeStr + "_n" + EndingStr,
@@ -3044,7 +3137,8 @@ bool AddressSanitizer::instrumentFunction(Function &F,
     return false;
   if (F.getLinkage() == GlobalValue::AvailableExternallyLinkage) return false;
   if (!ClDebugFunc.empty() && ClDebugFunc == F.getName()) return false;
-  if (F.getName().starts_with("__asan_")) return false;
+  if (F.getName().starts_with("__asan_"))
+    return false;
   if (F.getName().contains("__sycl_service_kernel__"))
     return false;
 

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

@@ -0,0 +1,88 @@
+// REQUIRES: linux
+// RUN: %{build} %device_sanitizer_flags -DMALLOC_DEVICE -O1 -g -o %t
+// RUN: %{run} not %t &> %t.txt ; FileCheck --check-prefixes CHECK,CHECK-DEVICE --input-file %t.txt %s
+// RUN: %{build} %device_sanitizer_flags -DMALLOC_DEVICE -O2 -g -o %t
+// RUN: %{run} not %t &> %t.txt ; FileCheck --check-prefixes CHECK,CHECK-DEVICE --input-file %t.txt %s
+// RUN: %{build} %device_sanitizer_flags -DMALLOC_HOST -O2 -g -o %t
+// RUN: %{run} not %t &> %t.txt ; FileCheck --check-prefixes CHECK,CHECK-HOST --input-file %t.txt %s
+// RUN: %{build} %device_sanitizer_flags -DMALLOC_SHARED -O2 -g -o %t
+// RUN: %{run} not %t &> %t.txt ; FileCheck --check-prefixes CHECK,CHECK-SHARED --input-file %t.txt %s
+#include <sycl/sycl.hpp>
+
+using namespace sycl;
+
+int main() {
+  sycl::queue Q;
+  constexpr std::size_t N = 16;
+  constexpr std::size_t M = 16;
+  constexpr std::size_t K = N / 4;
+#if defined(MALLOC_HOST)
+  auto matrixA = (int(*)[N])sycl::malloc_host<int>(N * M, Q);
+  auto matrixB = (int(*)[N])sycl::malloc_host<int>(N * M, Q);
+  auto matrixC = (int(*)[N])sycl::malloc_host<int>(N * M, Q);
+#elif defined(MALLOC_SHARED)
+  auto matrixA = (int(*)[N])sycl::malloc_shared<int>(N * M, Q);
+  auto matrixB = (int(*)[N])sycl::malloc_shared<int>(N * M, Q);
+  auto matrixC = (int(*)[N])sycl::malloc_shared<int>(N * M, Q);
+#elif defined(MALLOC_DEVICE)
+  auto matrixA = (int(*)[N])sycl::malloc_device<int>(N * M, Q);
+  auto matrixB = (int(*)[N])sycl::malloc_device<int>(N * M, Q);
+  auto matrixC = (int(*)[N])sycl::malloc_device<int>(N * M, Q);
+#elif defined(MALLOC_SYSTEM)
+  auto matrixA = (int(*)[N])new int[N * M];
+  auto matrixB = (int(*)[N])new int[N * M];
+  auto matrixC = (int(*)[N])new int[N * M];
+#else
+#error "Must provide malloc type to run the test"
+#endif
+
+  Q.single_task([=]() {
+    for (unsigned m = 0; m < M; ++m) {
+      for (unsigned n = 0; n < N; ++n) {
+        matrixA[m][n] = n;
+        matrixB[m][n] = n + m;
+        matrixC[m][n] = 0;
+      }
+    }
+  });
+  Q.wait();
+
+  Q.submit([&](sycl::handler &h) {
+    // Local accessor, for one matrix tile:
+    constexpr unsigned int tile_size = 16;
+    local_accessor<int> tileA{tile_size, h};
+    h.parallel_for<class MatMultiply>(
+        nd_range<2>{{M, N}, {1, tile_size}}, [=](nd_item<2> item) {
+          // Indices in the global index space:
+          int m = item.get_global_id()[0];
+          int n = item.get_global_id()[1];
+          // Index in the local index space:
+          int i = item.get_local_id()[1];
+          int sum = 0;
+          for (unsigned int kk = 0; kk < K; kk += tile_size) {
+            // Load the matrix tile from matrix A, and synchronize
+            // to ensure all work-items have a consistent view
+            // of the matrix tile in local memory.
+            tileA[i] = matrixA[m][kk + i + 1]; // <== bug add "+1" intentionally
+            // 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: {{READ of size 4 at kernel <.*MatMultiply> LID\(15, 0, 0\) GID\(15, 15, 0\)}}
+            // CHECK: {{  #0 .* .*matrix_multiply.cpp:}}[[@LINE-5]]
+            item.barrier();
+            // Perform computation using the local memory tile, and
+            // matrix B in global memory.
+            for (unsigned int k = 0; k < tile_size; k++)
+              sum += tileA[k] * matrixB[kk + k][n];
+            // After computation, synchronize again, to ensure all
+            // reads from the local memory tile are complete.
+            item.barrier();
+          }
+          // Write the final result to global memory.
+          matrixC[m][n] = sum;
+        });
+  });
+  Q.wait();
+
+  return 0;
+}

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

@@ -1,3 +1,14 @@
+<<<<<<< HEAD
+// REQUIRES: linux
+// RUN: %{build} %device_sanitizer_flags -DMALLOC_DEVICE -O1 -g -o %t
+// RUN: %{run} not %t &> %t.txt ; FileCheck --check-prefixes CHECK,CHECK-DEVICE --input-file %t.txt %s
+// RUN: %{build} %device_sanitizer_flags -DMALLOC_DEVICE -O2 -g -o %t
+// RUN: %{run} not %t &> %t.txt ; FileCheck --check-prefixes CHECK,CHECK-DEVICE --input-file %t.txt %s
+// RUN: %{build} %device_sanitizer_flags -DMALLOC_HOST -O2 -g -o %t
+// RUN: %{run} not %t &> %t.txt ; FileCheck --check-prefixes CHECK,CHECK-HOST --input-file %t.txt %s
+// RUN: %{build} %device_sanitizer_flags -DMALLOC_SHARED -O2 -g -o %t
+// RUN: %{run} not %t &> %t.txt ; FileCheck --check-prefixes CHECK,CHECK-SHARED --input-file %t.txt %s
+=======
 // REQUIRES: linux, cpu
 // RUN: %{build} %device_sanitizer_flags -DMALLOC_DEVICE -O1 -g -o %t
 // RUN: env SYCL_PREFER_UR=1 ONEAPI_DEVICE_SELECTOR=opencl:cpu %{run-unfiltered-devices} not %t &> %t.txt ; FileCheck --check-prefixes CHECK,CHECK-DEVICE --input-file %t.txt %s
@@ -7,6 +18,7 @@
 // RUN: env SYCL_PREFER_UR=1 ONEAPI_DEVICE_SELECTOR=opencl:cpu %{run-unfiltered-devices} not %t &> %t.txt ; FileCheck --check-prefixes CHECK,CHECK-HOST --input-file %t.txt %s
 // RUN: %{build} %device_sanitizer_flags -DMALLOC_SHARED -O2 -g -o %t
 // RUN: env SYCL_PREFER_UR=1 ONEAPI_DEVICE_SELECTOR=opencl:cpu %{run-unfiltered-devices} not %t &> %t.txt ; FileCheck --check-prefixes CHECK,CHECK-SHARED --input-file %t.txt %s
+>>>>>>> sycl
 #include <sycl/sycl.hpp>
 
 int main() {
@@ -30,11 +42,19 @@ int main() {
         [=](sycl::nd_item<1> item) { ++array[item.get_global_id(0)]; });
   });
   Q.wait();
+<<<<<<< HEAD
+  // 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: {{READ of size 1 at kernel <.*MyKernelR_4> LID\(0, 0, 0\) GID\(12345, 0, 0\)}}
+  // CHECK: {{  #0 .* .*parallel_for_char.cpp:}}[[@LINE-8]]
+=======
   // 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
 
   return 0;
 }

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

@@ -1,3 +1,14 @@
+<<<<<<< HEAD
+// REQUIRES: linux
+// RUN: %{build} %device_sanitizer_flags -DMALLOC_DEVICE -O1 -g -o %t
+// RUN: %{run} not %t &> %t.txt ; FileCheck --check-prefixes CHECK,CHECK-DEVICE --input-file %t.txt %s
+// RUN: %{build} %device_sanitizer_flags -DMALLOC_DEVICE -O2 -g -o %t
+// RUN: %{run} not %t &> %t.txt ; FileCheck --check-prefixes CHECK,CHECK-DEVICE --input-file %t.txt %s
+// RUN: %{build} %device_sanitizer_flags -DMALLOC_HOST -O2 -g -o %t
+// RUN: %{run} not %t &> %t.txt ; FileCheck --check-prefixes CHECK,CHECK-HOST --input-file %t.txt %s
+// RUN: %{build} %device_sanitizer_flags -DMALLOC_SHARED -O2 -g -o %t
+// RUN: %{run} not %t &> %t.txt ; FileCheck --check-prefixes CHECK,CHECK-SHARED --input-file %t.txt %s
+=======
 // REQUIRES: linux, cpu, aspect-fp64
 // RUN: %{build} %device_sanitizer_flags -DMALLOC_DEVICE -O1 -g -o %t
 // RUN: env SYCL_PREFER_UR=1 ONEAPI_DEVICE_SELECTOR=opencl:cpu %{run-unfiltered-devices} not %t &> %t.txt ; FileCheck --check-prefixes CHECK,CHECK-DEVICE --input-file %t.txt %s
@@ -7,6 +18,7 @@
 // RUN: env SYCL_PREFER_UR=1 ONEAPI_DEVICE_SELECTOR=opencl:cpu %{run-unfiltered-devices} not %t &> %t.txt ; FileCheck --check-prefixes CHECK,CHECK-HOST --input-file %t.txt %s
 // RUN: %{build} %device_sanitizer_flags -DMALLOC_SHARED -O2 -g -o %t
 // RUN: env SYCL_PREFER_UR=1 ONEAPI_DEVICE_SELECTOR=opencl:cpu %{run-unfiltered-devices} not %t &> %t.txt ; FileCheck --check-prefixes CHECK,CHECK-SHARED --input-file %t.txt %s
+>>>>>>> sycl
 #include <sycl/sycl.hpp>
 
 int main() {
@@ -30,11 +42,19 @@ int main() {
         [=](sycl::nd_item<1> item) { ++array[item.get_global_id(0)]; });
   });
   Q.wait();
+<<<<<<< HEAD
+  // 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: {{READ of size 8 at kernel <.*MyKernelR_4> LID\(0, 0, 0\) GID\(123456, 0, 0\)}}
+  // CHECK: {{  #0 .* .*parallel_for_double.cpp:}}[[@LINE-8]]
+=======
   // 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
 
   return 0;
 }