Merge remote-tracking branch 'origin/sycl' into review/yang/restructure_asan_msan

Author: kbenzie

llvm/lib/Transforms/Instrumentation/AddressSanitizer.cpp

@@ -1623,7 +1623,8 @@ static TargetExtType *getTargetExtType(Type *Ty) {
 // store float %1, ptr %call, align 4
 // clang-format on
 static bool isJointMatrixAccess(Value *V) {
-  if (auto *CI = dyn_cast<CallInst>(V)) {
+  auto *ActualV = V->stripInBoundsOffsets();
+  if (auto *CI = dyn_cast<CallInst>(ActualV)) {
     for (Value *Op : CI->args()) {
       if (auto *AI = dyn_cast<AllocaInst>(Op->stripInBoundsOffsets()))
         if (auto *TargetTy = getTargetExtType(AI->getAllocatedType()))

llvm/test/Instrumentation/AddressSanitizer/SPIRV/ignore_target_ext_type.ll

@@ -26,11 +26,13 @@ entry:
   %a = alloca %"struct.sycl::_V1::ext::oneapi::experimental::matrix::joint_matrix", align 8
   %0 = getelementptr inbounds %"struct.sycl::_V1::ext::oneapi::experimental::matrix::joint_matrix", ptr %a, i64 0, i32 0
   %call.i35 = call spir_func ptr @_Z19__spirv_AccessChainIfN4sycl3_V13ext6oneapi12experimental6matrix9precision4tf32ELm8ELm8ELN5__spv9MatrixUseE0ELNS8_5Scope4FlagE3EEPT_PPNS8_28__spirv_CooperativeMatrixKHRIT0_XT4_EXT1_EXT2_EXT3_EEEm(ptr %0, i64 0)
+  %1 = getelementptr inbounds { i16 }, ptr %call.i35, i64 0, i32 0
 ; CHECK-NOT: call void @__asan_load
 ; CHECK-NOT: call void @__asan_store
-  %1 = load float, ptr %call.i35, align 4
+  %2 = load i16, ptr %1, align 4
   %call.i42 = call spir_func ptr @_Z19__spirv_AccessChainIfN4sycl3_V13ext6oneapi12experimental6matrix9precision4tf32ELm8ELm8ELN5__spv9MatrixUseE0ELNS8_5Scope4FlagE3EEPT_PPNS8_28__spirv_CooperativeMatrixKHRIT0_XT4_EXT1_EXT2_EXT3_EEEm(ptr %0, i64 0)
-  store float %1, ptr %call.i42, align 4
+  %3 = getelementptr inbounds { i16 }, ptr %call.i42, i64 0, i32 0
+  store i16 %2, ptr %3, align 4
   ret void
 }
 

sycl/test-e2e/Adapters/interop-l0-direct.cpp

@@ -1,5 +1,9 @@
 // REQUIRES: level_zero, level_zero_dev_kit
 // UNSUPPORTED: ze_debug
+
+// DeviceSanitizer will report error for cross context USM usage, turn it off
+// RUN: export UR_LAYER_ASAN_OPTIONS="detect_kernel_arguments:0"
+
 // RUN: %{build} %level_zero_options -o %t.out
 // RUN: env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=0 %{run} %t.out
 // RUN: env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=1 %{run} %t.out
@@ -219,6 +223,9 @@ int main() {
     std::cout << "GPU Result from Immediate Q = {" << hostOut[0] << ", "
               << hostOut[1] << ", " << hostOut[2] << "}" << std::endl;
   }
+
+  free(deviceData, InteropContext);
+
   // Check results
   buffer<int, 1> bufDataResult{data, 3};
   host_accessor hostResult{bufDataResult, read_only};

sycl/test-e2e/AddressSanitizer/common/options-statistics.cpp

@@ -1,7 +1,8 @@
 // REQUIRES: linux, cpu || (gpu && level_zero)
 // RUN: %{build} %device_asan_flags -O2 -g -o %t
 // RUN: %{run} %t 2>&1 | FileCheck %s
-// RUN: env UR_LAYER_ASAN_OPTIONS=print_stats:1 %{run} %t 2>&1 | FileCheck --check-prefixes CHECK-STATS %s
+// RUN: env UR_LAYER_ASAN_OPTIONS="print_stats:1;quarantine_size_mb:1" %{run} %t 2>&1 | FileCheck --check-prefixes CHECK-STATS %s
+// RUN: env UR_LAYER_ASAN_OPTIONS="print_stats:1;quarantine_size_mb:0" %{run} %t 2>&1 | FileCheck --check-prefixes CHECK-STATS %s
 #include <sycl/usm.hpp>
 
 /// This test is used to check enabling/disabling memory overhead statistics
@@ -12,20 +13,26 @@ constexpr std::size_t group_size = 1;
 
 int main() {
   sycl::queue Q;
-  int *array = sycl::malloc_device<int>(1024 * 1024, Q);
+  int *array1 = sycl::malloc_device<int>(10 * 1024 * 1024, Q);
+  int *array2 = sycl::malloc_device<int>(10 * 1024 * 1024, Q);
+  int *array3 = sycl::malloc_device<int>(10 * 1024 * 1024, Q);
+
+  sycl::free(array2, Q);
+  sycl::free(array3, Q);
 
   Q.submit([&](sycl::handler &cgh) {
     auto acc = sycl::local_accessor<int>(group_size, cgh);
     cgh.parallel_for<class MyKernel>(
         sycl::nd_range<1>(N, group_size), [=](sycl::nd_item<1> item) {
-          array[item.get_global_id()] = acc[item.get_local_id()];
+          array1[item.get_global_id()] = acc[item.get_local_id()];
         });
   });
   Q.wait();
   // CHECK-STATS: Stats
   // CHECK-NOT: Stats
 
-  sycl::free(array, Q);
+  sycl::free(array1, Q);
+
   std::cout << "PASS" << std::endl;
   return 0;
 }