diff --git a/clang/lib/Driver/ToolChains/SYCL.cpp b/clang/lib/Driver/ToolChains/SYCL.cpp index 782a084ff3187..63e138090a421 100644 --- a/clang/lib/Driver/ToolChains/SYCL.cpp +++ b/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; + } } } diff --git a/libdevice/sanitizer/msan_rtl.cpp b/libdevice/sanitizer/msan_rtl.cpp index 81196b4f9f5da..ab02f4d0662e5 100644 --- a/libdevice/sanitizer/msan_rtl.cpp +++ b/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) { diff --git a/llvm/lib/Transforms/Instrumentation/MemorySanitizer.cpp b/llvm/lib/Transforms/Instrumentation/MemorySanitizer.cpp index a5e746cff0b3f..acf0a4f84f359 100644 --- a/llvm/lib/Transforms/Instrumentation/MemorySanitizer.cpp +++ b/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"); diff --git a/sycl/test-e2e/MemorySanitizer/check_call.cpp b/sycl/test-e2e/MemorySanitizer/check_call.cpp index 53804e86aae19..985fa9080aae6 100644 --- a/sycl/test-e2e/MemorySanitizer/check_call.cpp +++ b/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(2, Q); - auto *array2 = sycl::malloc_device(2, Q); + auto *array = sycl::malloc_device(2, Q); Q.submit([&](sycl::handler &h) { h.single_task( - [=]() { 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; }