fix crash

Author: AllanZyne

clang/lib/Driver/ToolChains/SYCL.cpp

@@ -691,6 +691,10 @@ SYCL::getDeviceLibraries(const Compilation &C, const llvm::Triple &TargetTriple,
         SanitizeVal = "address";
         break;
       }
+      if (Arg.find("-fsanitize=memory") != std::string::npos) {
+        SanitizeVal = "memory";
+        break;
+      }
     }
   }
 

libdevice/sanitizer/msan_rtl.cpp

@@ -68,11 +68,6 @@ void __msan_internal_report_save(const uint32_t size,
   const int Expected = MSAN_REPORT_NONE;
   int Desired = MSAN_REPORT_START;
 
-  if (UNLIKELY(!__MsanLaunchInfo)) {
-    __spirv_ocl_printf(__msan_print_warning_nolaunchinfo);
-    return;
-  }
-
   auto &SanitizerReport =
       ((__SYCL_GLOBAL__ MsanLaunchInfo *)__MsanLaunchInfo.get())->Report;
 
@@ -116,8 +111,6 @@ void __msan_internal_report_save(const uint32_t size,
 
     // Show we've done copying
     atomicStore(&SanitizerReport.Flag, MSAN_REPORT_FINISH);
-
-    MSAN_DEBUG(__spirv_ocl_printf(__msan_print_report, size, func));
   }
 }
 
@@ -162,18 +155,18 @@ MSAN_MAYBE_WARNING(u32, 4)
 MSAN_MAYBE_WARNING(u64, 8)
 
 DEVICE_EXTERN_C_NOINLINE uptr __msan_get_shadow(uptr addr, uint32_t as) {
+  // Return clean shadow (0s) by default
+  uptr shadow_ptr = (uptr)CleanShadow;
+
   if (UNLIKELY(!__MsanLaunchInfo)) {
     __spirv_ocl_printf(__msan_print_warning_nolaunchinfo);
-    return 0;
+    return shadow_ptr;
   }
 
   auto launch_info = (__SYCL_GLOBAL__ MsanLaunchInfo *)__MsanLaunchInfo.get();
   MSAN_DEBUG(__spirv_ocl_printf(__msan_print_launchinfo, (void *)launch_info,
                                 launch_info->GlobalShadowOffset));
 
-  // Return clean shadow (0s) by default
-  uptr shadow_ptr = (uptr)CleanShadow;
-
   if (LIKELY(launch_info->DeviceTy == DeviceType::CPU)) {
     shadow_ptr = __msan_get_shadow_cpu(addr);
   } else if (launch_info->DeviceTy == DeviceType::GPU_PVC) {

llvm/lib/Transforms/Instrumentation/MemorySanitizer.cpp

@@ -161,6 +161,7 @@
 #include "llvm/Analysis/GlobalsModRef.h"
 #include "llvm/Analysis/TargetLibraryInfo.h"
 #include "llvm/Analysis/ValueTracking.h"
+#include "llvm/Demangle/Demangle.h"
 #include "llvm/IR/Argument.h"
 #include "llvm/IR/AttributeMask.h"
 #include "llvm/IR/Attributes.h"
@@ -196,6 +197,7 @@
 #include "llvm/Support/DebugCounter.h"
 #include "llvm/Support/ErrorHandling.h"
 #include "llvm/Support/MathExtras.h"
+#include "llvm/Support/Path.h"
 #include "llvm/Support/raw_ostream.h"
 #include "llvm/TargetParser/Triple.h"
 #include "llvm/Transforms/Utils/BasicBlockUtils.h"
@@ -528,6 +530,19 @@ static const PlatformMemoryMapParams NetBSD_X86_MemoryMapParams = {
     &NetBSD_X86_64_MemoryMapParams,
 };
 
+// SPIR Linux
+static const MemoryMapParams Intel_SPIR_MemoryMapParams = {
+    0, // AndMask
+    0, // XorMask (not used)
+    0, // ShadowBase (not used)
+    0, // OriginBase
+};
+
+static const PlatformMemoryMapParams Intel_GFX_MemoryMapParams = {
+    nullptr,
+    &Intel_SPIR_MemoryMapParams,
+};
+
 // Spir memory address space
 static constexpr unsigned kSpirOffloadPrivateAS = 0;
 static constexpr unsigned kSpirOffloadGlobalAS = 1;
@@ -1069,6 +1084,7 @@ void MemorySanitizer::initializeModule(Module &M) {
       // NOTE: Support SPIR or SPIRV only, without MapParams
       if (!TargetTriple.isSPIROrSPIRV())
         report_fatal_error("unsupported architecture");
+      MapParams = Intel_GFX_MemoryMapParams.bits64;
       break;
     default:
       report_fatal_error("unsupported operating system");

sycl/test-e2e/MemorySanitizer/check_call.cpp

@@ -13,19 +13,17 @@ __attribute__((noinline)) long long foo(int data1, long long data2) {
 
 int main() {
   sycl::queue Q;
-  auto *array1 = sycl::malloc_device<int>(2, Q);
-  auto *array2 = sycl::malloc_device<long long>(2, Q);
+  auto *array = sycl::malloc_device<int>(2, Q);
 
   Q.submit([&](sycl::handler &h) {
     h.single_task<class MyKernel>(
-        [=]() { array2[0] = foo(array1[0], array2[1]); });
+        [=]() { array[0] = foo(array[0], array[1]); });
   });
   Q.wait();
   // CHECK: use-of-uninitialized-value
   // CHECK: kernel <{{.*MyKernel}}>
   // CHECK: #0 {{.*}} {{.*check_call.cpp}}:[[@LINE-5]]
 
-  sycl::free(array1, Q);
-  sycl::free(array2, Q);
+  sycl::free(array, Q);
   return 0;
 }