diff --git a/clang/lib/Driver/OffloadBundler.cpp b/clang/lib/Driver/OffloadBundler.cpp
index 1839d52de7a89..d2149c86c8a28 100644
--- a/clang/lib/Driver/OffloadBundler.cpp
+++ b/clang/lib/Driver/OffloadBundler.cpp
@@ -692,7 +692,7 @@ class ObjectFileHandler final : public FileHandler {
         if (SF->isIR() &&
             (Name == "llvm.used" || Name == "llvm.compiler.used" ||
              Name == "__AsanDeviceGlobalMetadata" ||
-             Name == "__AsanKernelMetadata"))
+             Name == "__AsanKernelMetadata" || Name == "__MsanKernelMetadata"))
           continue;
 
         // Add symbol name with the target prefix to the buffer.
diff --git a/clang/lib/Driver/SanitizerArgs.cpp b/clang/lib/Driver/SanitizerArgs.cpp
index 788d6682e26c3..90cdbfec6f4e2 100644
--- a/clang/lib/Driver/SanitizerArgs.cpp
+++ b/clang/lib/Driver/SanitizerArgs.cpp
@@ -1169,6 +1169,7 @@ void SanitizerArgs::addArgs(const ToolChain &TC, const llvm::opt::ArgList &Args,
   // SPIR/SPIRV sanitizer support is experimental and will pass a fixed set of
   // flags
   if (TC.getTriple().isSPIROrSPIRV()) {
+#if !defined(_WIN32)
     if (Sanitizers.has(SanitizerKind::Address)) {
       CmdArgs.push_back("-fsanitize=address");
       CmdArgs.push_back("-fsanitize-address-use-after-return=never");
@@ -1200,7 +1201,26 @@ void SanitizerArgs::addArgs(const ToolChain &TC, const llvm::opt::ArgList &Args,
 
       addSpecialCaseListOpt(Args, CmdArgs,
                             "-fsanitize-ignorelist=", UserIgnorelistFiles);
+    } else if (Sanitizers.has(SanitizerKind::Memory)) {
+      CmdArgs.push_back("-fsanitize=memory");
+
+      CmdArgs.push_back("-mllvm");
+      CmdArgs.push_back("-msan-instrumentation-with-call-threshold=0");
+
+      CmdArgs.push_back("-mllvm");
+      CmdArgs.push_back("-msan-eager-checks=1");
     }
+#else // _WIN32
+    std::string SanitizeArg;
+    if (Sanitizers.has(SanitizerKind::Address))
+      SanitizeArg = "-fsanitize=address";
+    else if (Sanitizers.has(SanitizerKind::Memory))
+      SanitizeArg = "-fsanitize=memory";
+
+    if (!SanitizeArg.empty())
+      TC.getDriver().Diag(diag::warn_drv_unsupported_option_for_target)
+          << SanitizeArg << TC.getTripleString();
+#endif
     return;
   }
 
diff --git a/clang/lib/Driver/ToolChains/SYCL.cpp b/clang/lib/Driver/ToolChains/SYCL.cpp
index 41231e50f83e6..ed254e01bd41e 100644
--- a/clang/lib/Driver/ToolChains/SYCL.cpp
+++ b/clang/lib/Driver/ToolChains/SYCL.cpp
@@ -555,6 +555,7 @@ SYCL::getDeviceLibraries(const Compilation &C, const llvm::Triple &TargetTriple,
       {"libsycl-asan-cpu", "internal"},
       {"libsycl-asan-dg2", "internal"},
       {"libsycl-asan-pvc", "internal"}};
+  const SYCLDeviceLibsList SYCLDeviceMsanLibs = {{"libsycl-msan", "internal"}};
 #endif
 
   const SYCLDeviceLibsList SYCLNativeCpuDeviceLibs = {
@@ -670,12 +671,15 @@ SYCL::getDeviceLibraries(const Compilation &C, const llvm::Triple &TargetTriple,
   };
 
   std::string SanitizeVal;
+  std::string SanitizeArg;
   size_t sanitizer_lib_idx = getSingleBuildTarget();
   if (Arg *A = Args.getLastArg(options::OPT_fsanitize_EQ,
                                options::OPT_fno_sanitize_EQ)) {
     if (A->getOption().matches(options::OPT_fsanitize_EQ) &&
-        A->getValues().size() == 1)
+        A->getValues().size() == 1) {
       SanitizeVal = A->getValue();
+      SanitizeArg = A->getAsString(Args);
+    }
   } else {
     // User can pass -fsanitize=address to device compiler via
     // -Xsycl-target-frontend, sanitize device library must be
@@ -699,6 +703,12 @@ SYCL::getDeviceLibraries(const Compilation &C, const llvm::Triple &TargetTriple,
     for (const std::string &Arg : ArgVals) {
       if (Arg.find("-fsanitize=address") != std::string::npos) {
         SanitizeVal = "address";
+        SanitizeArg = Arg;
+        break;
+      }
+      if (Arg.find("-fsanitize=memory") != std::string::npos) {
+        SanitizeVal = "memory";
+        SanitizeArg = Arg;
         break;
       }
     }
@@ -706,7 +716,8 @@ SYCL::getDeviceLibraries(const Compilation &C, const llvm::Triple &TargetTriple,
 
   if (SanitizeVal == "address")
     addSingleLibrary(SYCLDeviceAsanLibs[sanitizer_lib_idx]);
-
+  else if (SanitizeVal == "memory")
+    addLibraries(SYCLDeviceMsanLibs);
 #endif
 
   if (isNativeCPU)
@@ -826,6 +837,7 @@ static llvm::SmallVector<StringRef, 16> SYCLDeviceLibList{
     "asan-pvc",
     "asan-cpu",
     "asan-dg2",
+    "msan",
 #endif
     "imf",
     "imf-fp64",
@@ -1665,11 +1677,11 @@ SYCLToolChain::SYCLToolChain(const Driver &D, const llvm::Triple &Triple,
       if (SupportedByNativeCPU(*this, Opt))
         continue;
       // All sanitizer options are not currently supported, except
-      // AddressSanitizer
+      // AddressSanitizer and MemorySanitizer
       if (A->getOption().getID() == options::OPT_fsanitize_EQ &&
           A->getValues().size() == 1) {
         std::string SanitizeVal = A->getValue();
-        if (SanitizeVal == "address")
+        if (SanitizeVal == "address" || SanitizeVal == "memory")
           continue;
       }
       D.Diag(clang::diag::warn_drv_unsupported_option_for_target)
@@ -1710,7 +1722,7 @@ SYCLToolChain::TranslateArgs(const llvm::opt::DerivedArgList &Args,
         if (Opt.getID() == options::OPT_fsanitize_EQ &&
             A->getValues().size() == 1) {
           std::string SanitizeVal = A->getValue();
-          if (SanitizeVal == "address") {
+          if (SanitizeVal == "address" || SanitizeVal == "memory") {
             if (IsNewDAL)
               DAL->append(A);
             continue;
@@ -2119,5 +2131,5 @@ void SYCLToolChain::AddClangCXXStdlibIncludeArgs(const ArgList &Args,
 }
 
 SanitizerMask SYCLToolChain::getSupportedSanitizers() const {
-  return SanitizerKind::Address;
+  return SanitizerKind::Address | SanitizerKind::Memory;
 }
diff --git a/clang/test/Driver/sycl-device-lib-old-model.cpp b/clang/test/Driver/sycl-device-lib-old-model.cpp
index ab8f900f6710d..b4e3a3f9cf164 100644
--- a/clang/test/Driver/sycl-device-lib-old-model.cpp
+++ b/clang/test/Driver/sycl-device-lib-old-model.cpp
@@ -198,35 +198,35 @@
 /// ###########################################################################
 /// test behavior of libsycl-asan.bc linking when -fsanitize=address is available
 // RUN: %clangxx -fsycl --no-offload-new-driver %s --sysroot=%S/Inputs/SYCL -fsanitize=address -### 2>&1 \
-// RUN:   | FileCheck %s -check-prefix=SYCL_DEVICE_LIB_SANITIZER
+// RUN:   | FileCheck %s -check-prefix=SYCL_DEVICE_LIB_ASAN
 // RUN: %clangxx -fsycl --no-offload-new-driver %s --sysroot=%S/Inputs/SYCL -Xsycl-target-frontend -fsanitize=address -### 2>&1 \
-// RUN:   | FileCheck %s -check-prefix=SYCL_DEVICE_LIB_SANITIZER
+// RUN:   | FileCheck %s -check-prefix=SYCL_DEVICE_LIB_ASAN
 // RUN: %clangxx -fsycl --no-offload-new-driver %s --sysroot=%S/Inputs/SYCL -Xsycl-target-frontend=spir64 -fsanitize=address -### 2>&1 \
-// RUN:   | FileCheck %s -check-prefix=SYCL_DEVICE_LIB_SANITIZER
+// RUN:   | FileCheck %s -check-prefix=SYCL_DEVICE_LIB_ASAN
 // RUN: %clangxx -fsycl --no-offload-new-driver %s --sysroot=%S/Inputs/SYCL -Xarch_device -fsanitize=address -### 2>&1 \
-// RUN:   | FileCheck %s -check-prefix=SYCL_DEVICE_LIB_SANITIZER
+// RUN:   | FileCheck %s -check-prefix=SYCL_DEVICE_LIB_ASAN
 // RUN: %clangxx -fsycl --no-offload-new-driver %s --sysroot=%S/Inputs/SYCL -Xarch_device "-fsanitize=address -DUSE_SYCL_DEVICE_ASAN" -### 2>&1 \
-// RUN:   | FileCheck %s -check-prefix=SYCL_DEVICE_LIB_SANITIZER
+// RUN:   | FileCheck %s -check-prefix=SYCL_DEVICE_LIB_ASAN
 // RUN: %clangxx -fsycl --no-offload-new-driver %s --sysroot=%S/Inputs/SYCL -Xarch_device "-fsanitize=address -DUSE_SYCL_DEVICE_ASAN" -### 2>&1 \
 // RUN:   | FileCheck %s -check-prefix=SYCL_DEVICE_ASAN_MACRO
-// SYCL_DEVICE_LIB_SANITIZER: llvm-link{{.*}} "{{.*}}libsycl-crt.bc"
-// SYCL_DEVICE_LIB_SANITIZER-SAME: "{{.*}}libsycl-complex.bc"
-// SYCL_DEVICE_LIB_SANITIZER-SAME: "{{.*}}libsycl-complex-fp64.bc"
-// SYCL_DEVICE_LIB_SANITIZER-SAME: "{{.*}}libsycl-cmath.bc"
-// SYCL_DEVICE_LIB_SANITIZER-SAME: "{{.*}}libsycl-cmath-fp64.bc"
-// SYCL_DEVICE_LIB_SANITIZER-SAME: "{{.*}}libsycl-imf.bc"
-// SYCL_DEVICE_LIB_SANITIZER-SAME: "{{.*}}libsycl-imf-fp64.bc"
-// SYCL_DEVICE_LIB_SANITIZER-SAME: "{{.*}}libsycl-imf-bf16.bc"
-// SYCL_DEVICE_LIB_SANITIZER-SAME: "{{.*}}libsycl-fallback-cassert.bc"
-// SYCL_DEVICE_LIB_SANITIZER-SAME: "{{.*}}libsycl-fallback-cstring.bc"
-// SYCL_DEVICE_LIB_SANITIZER-SAME: "{{.*}}libsycl-fallback-complex.bc"
-// SYCL_DEVICE_LIB_SANITIZER-SAME: "{{.*}}libsycl-fallback-complex-fp64.bc"
-// SYCL_DEVICE_LIB_SANITIZER-SAME: "{{.*}}libsycl-fallback-cmath.bc"
-// SYCL_DEVICE_LIB_SANITIZER-SAME: "{{.*}}libsycl-fallback-cmath-fp64.bc"
-// SYCL_DEVICE_LIB_SANITIZER-SAME: "{{.*}}libsycl-fallback-imf.bc"
-// SYCL_DEVICE_LIB_SANITIZER-SAME: "{{.*}}libsycl-fallback-imf-fp64.bc"
-// SYCL_DEVICE_LIB_SANITIZER-SAME: "{{.*}}libsycl-fallback-imf-bf16.bc"
-// SYCL_DEVICE_LIB_SANITIZER-SAME: "{{.*}}libsycl-asan.bc"
+// SYCL_DEVICE_LIB_ASAN: llvm-link{{.*}} "{{.*}}libsycl-crt.bc"
+// SYCL_DEVICE_LIB_ASAN-SAME: "{{.*}}libsycl-complex.bc"
+// SYCL_DEVICE_LIB_ASAN-SAME: "{{.*}}libsycl-complex-fp64.bc"
+// SYCL_DEVICE_LIB_ASAN-SAME: "{{.*}}libsycl-cmath.bc"
+// SYCL_DEVICE_LIB_ASAN-SAME: "{{.*}}libsycl-cmath-fp64.bc"
+// SYCL_DEVICE_LIB_ASAN-SAME: "{{.*}}libsycl-imf.bc"
+// SYCL_DEVICE_LIB_ASAN-SAME: "{{.*}}libsycl-imf-fp64.bc"
+// SYCL_DEVICE_LIB_ASAN-SAME: "{{.*}}libsycl-imf-bf16.bc"
+// SYCL_DEVICE_LIB_ASAN-SAME: "{{.*}}libsycl-fallback-cassert.bc"
+// SYCL_DEVICE_LIB_ASAN-SAME: "{{.*}}libsycl-fallback-cstring.bc"
+// SYCL_DEVICE_LIB_ASAN-SAME: "{{.*}}libsycl-fallback-complex.bc"
+// SYCL_DEVICE_LIB_ASAN-SAME: "{{.*}}libsycl-fallback-complex-fp64.bc"
+// SYCL_DEVICE_LIB_ASAN-SAME: "{{.*}}libsycl-fallback-cmath.bc"
+// SYCL_DEVICE_LIB_ASAN-SAME: "{{.*}}libsycl-fallback-cmath-fp64.bc"
+// SYCL_DEVICE_LIB_ASAN-SAME: "{{.*}}libsycl-fallback-imf.bc"
+// SYCL_DEVICE_LIB_ASAN-SAME: "{{.*}}libsycl-fallback-imf-fp64.bc"
+// SYCL_DEVICE_LIB_ASAN-SAME: "{{.*}}libsycl-fallback-imf-bf16.bc"
+// SYCL_DEVICE_LIB_ASAN-SAME: "{{.*}}libsycl-asan.bc"
 // SYCL_DEVICE_ASAN_MACRO: "-cc1"
 // SYCL_DEVICE_ASAN_MACRO-SAME: "USE_SYCL_DEVICE_ASAN"
 // SYCL_DEVICE_ASAN_MACRO: llvm-link{{.*}} "-only-needed"
@@ -317,3 +317,41 @@
 // SYCL_DEVICE_LIB_ASAN_DG2-SAME: "{{.*}}libsycl-fallback-imf-fp64.bc"
 // SYCL_DEVICE_LIB_ASAN_DG2-SAME: "{{.*}}libsycl-fallback-imf-bf16.bc"
 // SYCL_DEVICE_LIB_ASAN_DG2-SAME: "{{.*}}libsycl-asan-dg2.bc"
+
+
+/// ###########################################################################
+/// test behavior of libsycl-msan.bc linking when -fsanitize=memory is available
+// RUN: %clangxx -fsycl --no-offload-new-driver %s --sysroot=%S/Inputs/SYCL -fsanitize=memory -### 2>&1 \
+// RUN:   | FileCheck %s -check-prefix=SYCL_DEVICE_LIB_MSAN
+// RUN: %clangxx -fsycl --no-offload-new-driver %s --sysroot=%S/Inputs/SYCL -Xsycl-target-frontend -fsanitize=memory -### 2>&1 \
+// RUN:   | FileCheck %s -check-prefix=SYCL_DEVICE_LIB_MSAN
+// RUN: %clangxx -fsycl --no-offload-new-driver %s --sysroot=%S/Inputs/SYCL -Xsycl-target-frontend=spir64 -fsanitize=memory -### 2>&1 \
+// RUN:   | FileCheck %s -check-prefix=SYCL_DEVICE_LIB_MSAN
+// RUN: %clangxx -fsycl --no-offload-new-driver %s --sysroot=%S/Inputs/SYCL -Xarch_device -fsanitize=memory -### 2>&1 \
+// RUN:   | FileCheck %s -check-prefix=SYCL_DEVICE_LIB_MSAN
+// RUN: %clangxx -fsycl --no-offload-new-driver %s --sysroot=%S/Inputs/SYCL -Xarch_device "-fsanitize=memory -DUSE_SYCL_DEVICE_MSAN" -### 2>&1 \
+// RUN:   | FileCheck %s -check-prefix=SYCL_DEVICE_LIB_MSAN
+// RUN: %clangxx -fsycl --no-offload-new-driver %s --sysroot=%S/Inputs/SYCL -Xarch_device "-fsanitize=memory -DUSE_SYCL_DEVICE_MSAN" -### 2>&1 \
+// RUN:   | FileCheck %s -check-prefix=SYCL_DEVICE_MSAN_MACRO
+// SYCL_DEVICE_LIB_MSAN: llvm-link{{.*}} "{{.*}}libsycl-crt.bc"
+// SYCL_DEVICE_LIB_MSAN-SAME: "{{.*}}libsycl-complex.bc"
+// SYCL_DEVICE_LIB_MSAN-SAME: "{{.*}}libsycl-complex-fp64.bc"
+// SYCL_DEVICE_LIB_MSAN-SAME: "{{.*}}libsycl-cmath.bc"
+// SYCL_DEVICE_LIB_MSAN-SAME: "{{.*}}libsycl-cmath-fp64.bc"
+// SYCL_DEVICE_LIB_MSAN-SAME: "{{.*}}libsycl-imf.bc"
+// SYCL_DEVICE_LIB_MSAN-SAME: "{{.*}}libsycl-imf-fp64.bc"
+// SYCL_DEVICE_LIB_MSAN-SAME: "{{.*}}libsycl-imf-bf16.bc"
+// SYCL_DEVICE_LIB_MSAN-SAME: "{{.*}}libsycl-fallback-cassert.bc"
+// SYCL_DEVICE_LIB_MSAN-SAME: "{{.*}}libsycl-fallback-cstring.bc"
+// SYCL_DEVICE_LIB_MSAN-SAME: "{{.*}}libsycl-fallback-complex.bc"
+// SYCL_DEVICE_LIB_MSAN-SAME: "{{.*}}libsycl-fallback-complex-fp64.bc"
+// SYCL_DEVICE_LIB_MSAN-SAME: "{{.*}}libsycl-fallback-cmath.bc"
+// SYCL_DEVICE_LIB_MSAN-SAME: "{{.*}}libsycl-fallback-cmath-fp64.bc"
+// SYCL_DEVICE_LIB_MSAN-SAME: "{{.*}}libsycl-fallback-imf.bc"
+// SYCL_DEVICE_LIB_MSAN-SAME: "{{.*}}libsycl-fallback-imf-fp64.bc"
+// SYCL_DEVICE_LIB_MSAN-SAME: "{{.*}}libsycl-fallback-imf-bf16.bc"
+// SYCL_DEVICE_LIB_MSAN-SAME: "{{.*}}libsycl-msan.bc"
+// SYCL_DEVICE_MSAN_MACRO: "-cc1"
+// SYCL_DEVICE_MSAN_MACRO-SAME: "USE_SYCL_DEVICE_MSAN"
+// SYCL_DEVICE_MSAN_MACRO: llvm-link{{.*}} "-only-needed"
+// SYCL_DEVICE_MSAN_MACRO-SAME: "{{.*}}libsycl-msan.bc"
diff --git a/clang/test/Driver/sycl-device-lib.cpp b/clang/test/Driver/sycl-device-lib.cpp
index e84eaadc5405a..9e07edf2287fa 100644
--- a/clang/test/Driver/sycl-device-lib.cpp
+++ b/clang/test/Driver/sycl-device-lib.cpp
@@ -155,36 +155,36 @@
 /// ###########################################################################
 /// test behavior of libsycl-asan.o linking when -fsanitize=address is available
 // RUN: %clangxx -fsycl --offload-new-driver %s --sysroot=%S/Inputs/SYCL -fsanitize=address -### 2>&1 \
-// RUN:   | FileCheck %s -check-prefix=SYCL_DEVICE_LIB_SANITIZER
+// RUN:   | FileCheck %s -check-prefix=SYCL_DEVICE_LIB_ASAN
 // RUN: %clangxx -fsycl --offload-new-driver %s --sysroot=%S/Inputs/SYCL -Xsycl-target-frontend -fsanitize=address -### 2>&1 \
-// RUN:   | FileCheck %s -check-prefix=SYCL_DEVICE_LIB_SANITIZER
+// RUN:   | FileCheck %s -check-prefix=SYCL_DEVICE_LIB_ASAN
 // RUN: %clangxx -fsycl --offload-new-driver %s --sysroot=%S/Inputs/SYCL -Xsycl-target-frontend=spir64 -fsanitize=address -### 2>&1 \
-// RUN:   | FileCheck %s -check-prefix=SYCL_DEVICE_LIB_SANITIZER
+// RUN:   | FileCheck %s -check-prefix=SYCL_DEVICE_LIB_ASAN
 // RUN: %clangxx -fsycl --offload-new-driver %s --sysroot=%S/Inputs/SYCL -Xarch_device -fsanitize=address -### 2>&1 \
-// RUN:   | FileCheck %s -check-prefix=SYCL_DEVICE_LIB_SANITIZER
+// RUN:   | FileCheck %s -check-prefix=SYCL_DEVICE_LIB_ASAN
 // RUN: %clangxx -fsycl --offload-new-driver %s --sysroot=%S/Inputs/SYCL -Xarch_device "-fsanitize=address -DUSE_SYCL_DEVICE_ASAN" -### 2>&1 \
-// RUN:   | FileCheck %s -check-prefix=SYCL_DEVICE_LIB_SANITIZER
+// RUN:   | FileCheck %s -check-prefix=SYCL_DEVICE_LIB_ASAN
 // RUN: %clangxx -fsycl --offload-new-driver %s --sysroot=%S/Inputs/SYCL -Xarch_device "-fsanitize=address -DUSE_SYCL_DEVICE_ASAN" -### 2>&1 \
 // RUN:   | FileCheck %s -check-prefix=SYCL_DEVICE_ASAN_MACRO
-// SYCL_DEVICE_LIB_SANITIZER: clang-linker-wrapper{{.*}} "-sycl-device-libraries
-// SYCL_DEVICE_LIB_SANITIZER: {{.*}}libsycl-crt.new.o
-// SYCL_DEVICE_LIB_SANITIZER-SAME: {{.*}}libsycl-complex.
-// SYCL_DEVICE_LIB_SANITIZER-SAME: {{.*}}libsycl-complex-fp64.
-// SYCL_DEVICE_LIB_SANITIZER-SAME: {{.*}}libsycl-cmath.new.o
-// SYCL_DEVICE_LIB_SANITIZER-SAME: {{.*}}libsycl-cmath-fp64.new.o
-// SYCL_DEVICE_LIB_SANITIZER-SAME: {{.*}}libsycl-imf.new.o
-// SYCL_DEVICE_LIB_SANITIZER-SAME: {{.*}}libsycl-imf-fp64.new.o
-// SYCL_DEVICE_LIB_SANITIZER-SAME: {{.*}}libsycl-imf-bf16.new.o
-// SYCL_DEVICE_LIB_SANITIZER-SAME: {{.*}}libsycl-fallback-cassert.new.o
-// SYCL_DEVICE_LIB_SANITIZER-SAME: {{.*}}libsycl-fallback-cstring.new.o
-// SYCL_DEVICE_LIB_SANITIZER-SAME: {{.*}}libsycl-fallback-complex.new.o
-// SYCL_DEVICE_LIB_SANITIZER-SAME: {{.*}}libsycl-fallback-complex-fp64.new.o
-// SYCL_DEVICE_LIB_SANITIZER-SAME: {{.*}}libsycl-fallback-cmath.new.o
-// SYCL_DEVICE_LIB_SANITIZER-SAME: {{.*}}libsycl-fallback-cmath-fp64.new.o
-// SYCL_DEVICE_LIB_SANITIZER-SAME: {{.*}}libsycl-fallback-imf.new.o
-// SYCL_DEVICE_LIB_SANITIZER-SAME: {{.*}}libsycl-fallback-imf-fp64.new.o
-// SYCL_DEVICE_LIB_SANITIZER-SAME: {{.*}}libsycl-fallback-imf-bf16.new.o
-// SYCL_DEVICE_LIB_SANITIZER-SAME: {{.*}}libsycl-asan.new.o
+// SYCL_DEVICE_LIB_ASAN: clang-linker-wrapper{{.*}} "-sycl-device-libraries
+// SYCL_DEVICE_LIB_ASAN: {{.*}}libsycl-crt.new.o
+// SYCL_DEVICE_LIB_ASAN-SAME: {{.*}}libsycl-complex.
+// SYCL_DEVICE_LIB_ASAN-SAME: {{.*}}libsycl-complex-fp64.
+// SYCL_DEVICE_LIB_ASAN-SAME: {{.*}}libsycl-cmath.new.o
+// SYCL_DEVICE_LIB_ASAN-SAME: {{.*}}libsycl-cmath-fp64.new.o
+// SYCL_DEVICE_LIB_ASAN-SAME: {{.*}}libsycl-imf.new.o
+// SYCL_DEVICE_LIB_ASAN-SAME: {{.*}}libsycl-imf-fp64.new.o
+// SYCL_DEVICE_LIB_ASAN-SAME: {{.*}}libsycl-imf-bf16.new.o
+// SYCL_DEVICE_LIB_ASAN-SAME: {{.*}}libsycl-fallback-cassert.new.o
+// SYCL_DEVICE_LIB_ASAN-SAME: {{.*}}libsycl-fallback-cstring.new.o
+// SYCL_DEVICE_LIB_ASAN-SAME: {{.*}}libsycl-fallback-complex.new.o
+// SYCL_DEVICE_LIB_ASAN-SAME: {{.*}}libsycl-fallback-complex-fp64.new.o
+// SYCL_DEVICE_LIB_ASAN-SAME: {{.*}}libsycl-fallback-cmath.new.o
+// SYCL_DEVICE_LIB_ASAN-SAME: {{.*}}libsycl-fallback-cmath-fp64.new.o
+// SYCL_DEVICE_LIB_ASAN-SAME: {{.*}}libsycl-fallback-imf.new.o
+// SYCL_DEVICE_LIB_ASAN-SAME: {{.*}}libsycl-fallback-imf-fp64.new.o
+// SYCL_DEVICE_LIB_ASAN-SAME: {{.*}}libsycl-fallback-imf-bf16.new.o
+// SYCL_DEVICE_LIB_ASAN-SAME: {{.*}}libsycl-asan.new.o
 // SYCL_DEVICE_ASAN_MACRO: "-cc1"
 // SYCL_DEVICE_ASAN_MACRO-SAME: "USE_SYCL_DEVICE_ASAN"
 // SYCL_DEVICE_ASAN_MACRO: libsycl-asan.new.o
@@ -314,3 +314,41 @@
 // SYCL_DEVICE_LIB_ASAN_MUL-SAME: {{.*}}libsycl-fallback-imf-fp64.new.o
 // SYCL_DEVICE_LIB_ASAN_MUL-SAME: {{.*}}libsycl-fallback-imf-bf16.new.o
 // SYCL_DEVICE_LIB_ASAN_MUL-SAME: {{.*}}libsycl-asan.new.o
+
+
+/// ###########################################################################
+/// test behavior of libsycl-msan.o linking when -fsanitize=memory is available
+// RUN: %clangxx -fsycl --offload-new-driver %s --sysroot=%S/Inputs/SYCL -fsanitize=memory -### 2>&1 \
+// RUN:   | FileCheck %s -check-prefix=SYCL_DEVICE_LIB_MSAN
+// RUN: %clangxx -fsycl --offload-new-driver %s --sysroot=%S/Inputs/SYCL -Xsycl-target-frontend -fsanitize=memory -### 2>&1 \
+// RUN:   | FileCheck %s -check-prefix=SYCL_DEVICE_LIB_MSAN
+// RUN: %clangxx -fsycl --offload-new-driver %s --sysroot=%S/Inputs/SYCL -Xsycl-target-frontend=spir64 -fsanitize=memory -### 2>&1 \
+// RUN:   | FileCheck %s -check-prefix=SYCL_DEVICE_LIB_MSAN
+// RUN: %clangxx -fsycl --offload-new-driver %s --sysroot=%S/Inputs/SYCL -Xarch_device -fsanitize=memory -### 2>&1 \
+// RUN:   | FileCheck %s -check-prefix=SYCL_DEVICE_LIB_MSAN
+// RUN: %clangxx -fsycl --offload-new-driver %s --sysroot=%S/Inputs/SYCL -Xarch_device "-fsanitize=memory -DUSE_SYCL_DEVICE_MSAN" -### 2>&1 \
+// RUN:   | FileCheck %s -check-prefix=SYCL_DEVICE_LIB_MSAN
+// RUN: %clangxx -fsycl --offload-new-driver %s --sysroot=%S/Inputs/SYCL -Xarch_device "-fsanitize=memory -DUSE_SYCL_DEVICE_MSAN" -### 2>&1 \
+// RUN:   | FileCheck %s -check-prefix=SYCL_DEVICE_MSAN_MACRO
+// SYCL_DEVICE_LIB_MSAN: clang-linker-wrapper{{.*}} "-sycl-device-libraries
+// SYCL_DEVICE_LIB_MSAN: {{.*}}libsycl-crt.new.o
+// SYCL_DEVICE_LIB_MSAN-SAME: {{.*}}libsycl-complex.
+// SYCL_DEVICE_LIB_MSAN-SAME: {{.*}}libsycl-complex-fp64.
+// SYCL_DEVICE_LIB_MSAN-SAME: {{.*}}libsycl-cmath.new.o
+// SYCL_DEVICE_LIB_MSAN-SAME: {{.*}}libsycl-cmath-fp64.new.o
+// SYCL_DEVICE_LIB_MSAN-SAME: {{.*}}libsycl-imf.new.o
+// SYCL_DEVICE_LIB_MSAN-SAME: {{.*}}libsycl-imf-fp64.new.o
+// SYCL_DEVICE_LIB_MSAN-SAME: {{.*}}libsycl-imf-bf16.new.o
+// SYCL_DEVICE_LIB_MSAN-SAME: {{.*}}libsycl-fallback-cassert.new.o
+// SYCL_DEVICE_LIB_MSAN-SAME: {{.*}}libsycl-fallback-cstring.new.o
+// SYCL_DEVICE_LIB_MSAN-SAME: {{.*}}libsycl-fallback-complex.new.o
+// SYCL_DEVICE_LIB_MSAN-SAME: {{.*}}libsycl-fallback-complex-fp64.new.o
+// SYCL_DEVICE_LIB_MSAN-SAME: {{.*}}libsycl-fallback-cmath.new.o
+// SYCL_DEVICE_LIB_MSAN-SAME: {{.*}}libsycl-fallback-cmath-fp64.new.o
+// SYCL_DEVICE_LIB_MSAN-SAME: {{.*}}libsycl-fallback-imf.new.o
+// SYCL_DEVICE_LIB_MSAN-SAME: {{.*}}libsycl-fallback-imf-fp64.new.o
+// SYCL_DEVICE_LIB_MSAN-SAME: {{.*}}libsycl-fallback-imf-bf16.new.o
+// SYCL_DEVICE_LIB_MSAN-SAME: {{.*}}libsycl-msan.new.o
+// SYCL_DEVICE_MSAN_MACRO: "-cc1"
+// SYCL_DEVICE_MSAN_MACRO-SAME: "USE_SYCL_DEVICE_MSAN"
+// SYCL_DEVICE_MSAN_MACRO: libsycl-msan.new.o
diff --git a/clang/test/Driver/sycl-device-sanitizer-win.cpp b/clang/test/Driver/sycl-device-sanitizer-win.cpp
new file mode 100644
index 0000000000000..563d36e381df8
--- /dev/null
+++ b/clang/test/Driver/sycl-device-sanitizer-win.cpp
@@ -0,0 +1,24 @@
+///
+/// Perform several driver tests for SYCL device side sanitizers on Windows
+///
+
+// REQUIRES: system-windows
+
+/// ###########################################################################
+
+// RUN: %clangxx -fsycl -fsanitize=address -c %s -### 2>&1 \
+// RUN:   | FileCheck --check-prefix=SYCL-ASAN %s
+// RUN: %clangxx -fsycl -Xarch_device -fsanitize=address -c %s -### 2>&1 \
+// RUN:   | FileCheck --check-prefix=SYCL-ASAN %s
+
+// SYCL-ASAN: ignoring '-fsanitize=address' option as it is not currently supported for target 'spir64-unknown-unknown'
+
+/// ###########################################################################
+
+// We need to add "not" here since "error: unsupported option '-fsanitize=memory' for target 'x86_64-pc-windows-msvc'"
+// RUN: not %clangxx -fsycl -fsanitize=memory -c %s -### 2>&1 \
+// RUN:   | FileCheck --check-prefix=SYCL-MSAN %s
+// RUN: %clangxx -fsycl -Xarch_device -fsanitize=memory -c %s -### 2>&1 \
+// RUN:   | FileCheck --check-prefix=SYCL-MSAN %s
+
+// SYCL-MSAN: ignoring '-fsanitize=memory' option as it is not currently supported for target 'spir64-unknown-unknown'
diff --git a/clang/test/Driver/sycl-device-sanitizer.cpp b/clang/test/Driver/sycl-device-sanitizer.cpp
index e9ddcf091c268..b25bb3c524955 100644
--- a/clang/test/Driver/sycl-device-sanitizer.cpp
+++ b/clang/test/Driver/sycl-device-sanitizer.cpp
@@ -1,3 +1,11 @@
+///
+/// Perform several driver tests for SYCL device side sanitizers on Linux
+///
+
+// UNSUPPORTED: system-windows
+
+/// ###########################################################################
+
 // RUN: %clangxx -fsycl -fsanitize=address -c %s -### 2>&1 \
 // RUN:   | FileCheck --check-prefix=SYCL-ASAN %s
 // SYCL-ASAN: clang{{.*}} "-fsycl-is-device"
@@ -13,9 +21,9 @@
 // SYCL-ASAN-SAME: "-mllvm" "-asan-mapping-scale=4"
 
 // RUN: %clangxx -fsycl -Xarch_device -fsanitize=address -c %s -### 2>&1 \
-// RUN:   | FileCheck --check-prefix=SYCL-XARCH-DEVICE %s
-// SYCL-XARCH-DEVICE: clang{{.*}} "-fsycl-is-device"
-// SYCL-XARCH-DEVICE-SAME: -fsanitize=address
+// RUN:   | FileCheck --check-prefix=SYCL-ASAN-XARCH-DEVICE %s
+// SYCL-ASAN-XARCH-DEVICE: clang{{.*}} "-fsycl-is-device"
+// SYCL-ASAN-XARCH-DEVICE-SAME: -fsanitize=address
 
 // RUN: %clangxx -fsycl -Xarch_device -fsanitize=address -Xarch_device -fsanitize-recover=address -c %s -### 2>&1 \
 // RUN:   | FileCheck --check-prefix=SYCL-ASAN-RECOVER %s
@@ -30,3 +38,17 @@
 // SYCL-ASAN-FILTER: clang{{.*}} "-fsycl-is-device"
 // SYCL-ASAN-FILTER-SAME: -fsanitize=address
 // SYCL-ASAN-FILTER-SAME: "-mllvm" "-asan-stack=0"
+
+/// ###########################################################################
+
+// RUN: %clangxx -fsycl -fsanitize=memory -c %s -### 2>&1 \
+// RUN:   | FileCheck --check-prefix=SYCL-MSAN %s
+// SYCL-MSAN: clang{{.*}} "-fsycl-is-device"
+// SYCL-MSAN-SAME: -fsanitize=memory
+// SYCL-MSAN-SAME: "-mllvm" "-msan-instrumentation-with-call-threshold=0"
+// SYCL-MSAN-SAME: "-mllvm" "-msan-eager-checks=1"
+
+// RUN: %clangxx -fsycl -Xarch_device -fsanitize=memory -c %s -### 2>&1 \
+// RUN:   | FileCheck --check-prefix=SYCL-MSAN-XARCH-DEVICE %s
+// SYCL-MSAN-XARCH-DEVICE: clang{{.*}} "-fsycl-is-device"
+// SYCL-MSAN-XARCH-DEVICE-SAME: -fsanitize=memory
diff --git a/libdevice/cmake/modules/SYCLLibdevice.cmake b/libdevice/cmake/modules/SYCLLibdevice.cmake
index a2ca6a03f50f1..0b2c1780a4756 100644
--- a/libdevice/cmake/modules/SYCLLibdevice.cmake
+++ b/libdevice/cmake/modules/SYCLLibdevice.cmake
@@ -288,6 +288,13 @@ if (NOT MSVC AND UR_SANITIZER_INCLUDE_DIR)
                                             ${sanitizer_generic_compile_opts}
                                             ${sycl_dg2_target_opt}
                                             -D__LIBDEVICE_DG2__)
+  
+  set(msan_obj_deps
+    device.h atomic.hpp spirv_vars.h
+    ${UR_SANITIZER_INCLUDE_DIR}/msan/msan_libdevice.hpp
+    include/msan_rtl.hpp
+    include/spir_global_var.hpp
+    sycl-compiler)
 endif()
 
 if("native_cpu" IN_LIST SYCL_ENABLE_BACKENDS)
@@ -373,6 +380,14 @@ else()
                         OPTS ${asan_${asan_device}_compile_opts_${asan_ft}})
       endforeach()
     endforeach()
+
+    # msan jit
+    add_devicelibs(libsycl-msan
+      SRC sanitizer/msan_rtl.cpp
+      DEPENDENCIES ${msan_obj_deps}
+      EXTRA_OPTS -fno-sycl-instrument-device-code
+                 -I${UR_SANITIZER_INCLUDE_DIR}
+                 -I${CMAKE_CURRENT_SOURCE_DIR})
   endif()
 endif()
 
diff --git a/libdevice/include/msan_rtl.hpp b/libdevice/include/msan_rtl.hpp
new file mode 100644
index 0000000000000..1c160c3f65a46
--- /dev/null
+++ b/libdevice/include/msan_rtl.hpp
@@ -0,0 +1,14 @@
+//==-- msan_rtl.hpp - Declaration for sanitizer global var ---==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+#pragma once
+
+#include "sanitizer_defs.hpp"
+#include "spir_global_var.hpp"
+
+// Treat this header as system one to workaround frontend's restriction
+#pragma clang system_header
diff --git a/libdevice/sanitizer/msan_rtl.cpp b/libdevice/sanitizer/msan_rtl.cpp
new file mode 100644
index 0000000000000..ab02f4d0662e5
--- /dev/null
+++ b/libdevice/sanitizer/msan_rtl.cpp
@@ -0,0 +1,185 @@
+//==--- msan_rtl.cpp - device memory sanitizer runtime library -------------==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#include "include/msan_rtl.hpp"
+#include "atomic.hpp"
+#include "device.h"
+#include "msan/msan_libdevice.hpp"
+#include "spirv_vars.h"
+
+DeviceGlobal<void *> __MsanLaunchInfo;
+
+constexpr int MSAN_REPORT_NONE = 0;
+constexpr int MSAN_REPORT_START = 1;
+constexpr int MSAN_REPORT_FINISH = 2;
+
+static const uint64_t CleanShadow[16] = {};
+
+static const __SYCL_CONSTANT__ char __msan_print_warning_return[] =
+    "[kernel] !!! msan warning return\n";
+
+static const __SYCL_CONSTANT__ char __msan_print_shadow[] =
+    "[kernel] __msan_get_shadow(addr=%p, as=%d) = %p: %02X\n";
+
+static const __SYCL_CONSTANT__ char __msan_print_warning_nolaunchinfo[] =
+    "[kernel] !!! __mem_warning_nolaunchinfo\n";
+
+static const __SYCL_CONSTANT__ char __msan_print_launchinfo[] =
+    "[kernel] !!! launchinfo %p (GlobalShadow=%p)\n";
+
+static const __SYCL_CONSTANT__ char __msan_print_report[] =
+    "[kernel] %d bytes uninitialized at kernel %s\n";
+
+static const __SYCL_CONSTANT__ char __msan_print_unsupport_device_type[] =
+    "[kernel] Unsupport device type: %d\n";
+
+#if defined(__SPIR__) || defined(__SPIRV__)
+
+#if defined(__SYCL_DEVICE_ONLY__)
+#define __USE_SPIR_BUILTIN__ 1
+#endif
+
+#if __USE_SPIR_BUILTIN__
+extern SYCL_EXTERNAL int
+__spirv_ocl_printf(const __SYCL_CONSTANT__ char *Format, ...);
+extern "C" SYCL_EXTERNAL void __devicelib_exit();
+#endif
+
+#define MSAN_DEBUG(X)                                                          \
+  do {                                                                         \
+    auto launch_info =                                                         \
+        (__SYCL_GLOBAL__ const MsanLaunchInfo *)__MsanLaunchInfo.get();        \
+    if (launch_info->Debug) {                                                  \
+      X;                                                                       \
+    }                                                                          \
+  } while (false)
+
+namespace {
+
+void __msan_internal_report_save(const uint32_t size,
+                                 const char __SYCL_CONSTANT__ *file,
+                                 const uint32_t line,
+                                 const char __SYCL_CONSTANT__ *func) {
+  const int Expected = MSAN_REPORT_NONE;
+  int Desired = MSAN_REPORT_START;
+
+  auto &SanitizerReport =
+      ((__SYCL_GLOBAL__ MsanLaunchInfo *)__MsanLaunchInfo.get())->Report;
+
+  if (atomicCompareAndSet(&SanitizerReport.Flag, Desired, Expected) ==
+      Expected) {
+
+    int FileLength = 0;
+    int FuncLength = 0;
+
+    if (file)
+      for (auto *C = file; *C != '\0'; ++C, ++FileLength)
+        ;
+    if (func)
+      for (auto *C = func; *C != '\0'; ++C, ++FuncLength)
+        ;
+
+    int MaxFileIdx = sizeof(SanitizerReport.File) - 1;
+    int MaxFuncIdx = sizeof(SanitizerReport.Func) - 1;
+
+    if (FileLength < MaxFileIdx)
+      MaxFileIdx = FileLength;
+    if (FuncLength < MaxFuncIdx)
+      MaxFuncIdx = FuncLength;
+
+    for (int Idx = 0; Idx < MaxFileIdx; ++Idx)
+      SanitizerReport.File[Idx] = file[Idx];
+    SanitizerReport.File[MaxFileIdx] = '\0';
+
+    for (int Idx = 0; Idx < MaxFuncIdx; ++Idx)
+      SanitizerReport.Func[Idx] = func[Idx];
+    SanitizerReport.Func[MaxFuncIdx] = '\0';
+
+    SanitizerReport.AccessSize = size;
+    SanitizerReport.Line = line;
+    SanitizerReport.GID0 = __spirv_GlobalInvocationId_x();
+    SanitizerReport.GID1 = __spirv_GlobalInvocationId_y();
+    SanitizerReport.GID2 = __spirv_GlobalInvocationId_z();
+    SanitizerReport.LID0 = __spirv_LocalInvocationId_x();
+    SanitizerReport.LID1 = __spirv_LocalInvocationId_y();
+    SanitizerReport.LID2 = __spirv_LocalInvocationId_z();
+
+    // Show we've done copying
+    atomicStore(&SanitizerReport.Flag, MSAN_REPORT_FINISH);
+  }
+}
+
+void __msan_report_error(const uint32_t size,
+                         const char __SYCL_CONSTANT__ *file,
+                         const uint32_t line,
+                         const char __SYCL_CONSTANT__ *func) {
+  __msan_internal_report_save(size, file, line, func);
+
+  auto launch = (__SYCL_GLOBAL__ MsanLaunchInfo *)__MsanLaunchInfo.get();
+  if (!launch->IsRecover) {
+    __devicelib_exit();
+  }
+}
+
+inline uptr __msan_get_shadow_cpu(uptr addr) {
+  return addr ^ 0x500000000000ULL;
+}
+
+inline uptr __msan_get_shadow_pvc(uptr addr, uint32_t as) {
+  // Device USM only
+  uptr shadow_ptr = ((__SYCL_GLOBAL__ MsanLaunchInfo *)__MsanLaunchInfo.get())
+                        ->GlobalShadowOffset +
+                    (addr & 0x3FFF'FFFF'FFFFULL);
+  return shadow_ptr;
+}
+
+} // namespace
+
+#define MSAN_MAYBE_WARNING(type, size)                                         \
+  DEVICE_EXTERN_C_NOINLINE void __msan_maybe_warning_##size(                   \
+      type s, u32 o, const char __SYCL_CONSTANT__ *file, uint32_t line,        \
+      const char __SYCL_CONSTANT__ *func) {                                    \
+    if (UNLIKELY(s)) {                                                         \
+      __msan_report_error(size, file, line, func);                             \
+    }                                                                          \
+  }
+
+MSAN_MAYBE_WARNING(u8, 1)
+MSAN_MAYBE_WARNING(u16, 2)
+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 shadow_ptr;
+  }
+
+  auto launch_info = (__SYCL_GLOBAL__ MsanLaunchInfo *)__MsanLaunchInfo.get();
+  MSAN_DEBUG(__spirv_ocl_printf(__msan_print_launchinfo, (void *)launch_info,
+                                launch_info->GlobalShadowOffset));
+
+  if (LIKELY(launch_info->DeviceTy == DeviceType::CPU)) {
+    shadow_ptr = __msan_get_shadow_cpu(addr);
+  } else if (launch_info->DeviceTy == DeviceType::GPU_PVC) {
+    shadow_ptr = __msan_get_shadow_pvc(addr, as);
+  } else {
+    MSAN_DEBUG(__spirv_ocl_printf(__msan_print_unsupport_device_type,
+                                  launch_info->DeviceTy));
+  }
+
+  MSAN_DEBUG(__spirv_ocl_printf(__msan_print_shadow, (void *)addr, as,
+                                (void *)shadow_ptr, *(u8 *)shadow_ptr));
+
+  return shadow_ptr;
+}
+
+#endif // __SPIR__ || __SPIRV__
diff --git a/llvm/include/llvm/SYCLLowerIR/ComputeModuleRuntimeInfo.h b/llvm/include/llvm/SYCLLowerIR/ComputeModuleRuntimeInfo.h
index e7cff6c730051..ba8471965aa76 100644
--- a/llvm/include/llvm/SYCLLowerIR/ComputeModuleRuntimeInfo.h
+++ b/llvm/include/llvm/SYCLLowerIR/ComputeModuleRuntimeInfo.h
@@ -29,6 +29,7 @@ struct GlobalBinImageProps {
   bool EmitDeviceGlobalPropSet;
 };
 bool isModuleUsingAsan(const Module &M);
+bool isModuleUsingMsan(const Module &M);
 using PropSetRegTy = llvm::util::PropertySetRegistry;
 using EntryPointSet = SetVector<Function *>;
 
diff --git a/llvm/include/llvm/SYCLLowerIR/AsanKernelMetadata.h b/llvm/include/llvm/SYCLLowerIR/SanitizerKernelMetadata.h
similarity index 75%
rename from llvm/include/llvm/SYCLLowerIR/AsanKernelMetadata.h
rename to llvm/include/llvm/SYCLLowerIR/SanitizerKernelMetadata.h
index 56286927395b9..1b30ea7b05b25 100644
--- a/llvm/include/llvm/SYCLLowerIR/AsanKernelMetadata.h
+++ b/llvm/include/llvm/SYCLLowerIR/SanitizerKernelMetadata.h
@@ -1,4 +1,4 @@
-//===-- AsanKernelMetadata.h - fix kernel medatadata for sanitizer ---===//
+//===-- SanitizerKernelMetadata.h - fix kernel medatadata for sanitizer ---===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
@@ -6,7 +6,7 @@
 //
 //===----------------------------------------------------------------------===//
 // This pass fixes attributes and metadata of the global variable
-// "__AsanKernelMetadata"
+// "__XsanKernelMetadata"
 //===----------------------------------------------------------------------===//
 
 #pragma once
@@ -15,7 +15,8 @@
 
 namespace llvm {
 
-class AsanKernelMetadataPass : public PassInfoMixin<AsanKernelMetadataPass> {
+class SanitizerKernelMetadataPass
+    : public PassInfoMixin<SanitizerKernelMetadataPass> {
 public:
   PreservedAnalyses run(Module &M, ModuleAnalysisManager &MAM);
 };
diff --git a/llvm/lib/SYCLLowerIR/CMakeLists.txt b/llvm/lib/SYCLLowerIR/CMakeLists.txt
index fc254a4809bb5..bc221b912977c 100644
--- a/llvm/lib/SYCLLowerIR/CMakeLists.txt
+++ b/llvm/lib/SYCLLowerIR/CMakeLists.txt
@@ -69,12 +69,13 @@ add_llvm_component_library(LLVMSYCLLowerIR
   SYCLPropagateJointMatrixUsage.cpp
   SYCLVirtualFunctionsAnalysis.cpp
   SYCLUtils.cpp
-  AsanKernelMetadata.cpp
 
   LocalAccessorToSharedMemory.cpp
   GlobalOffset.cpp
   TargetHelpers.cpp
 
+  SanitizerKernelMetadata.cpp
+
   ADDITIONAL_HEADER_DIRS
   ${LLVM_MAIN_INCLUDE_DIR}/llvm/SYCLLowerIR
   ${LLVM_MAIN_SRC_DIR}/projects/vc-intrinsics/GenXIntrinsics/include
diff --git a/llvm/lib/SYCLLowerIR/ComputeModuleRuntimeInfo.cpp b/llvm/lib/SYCLLowerIR/ComputeModuleRuntimeInfo.cpp
index c61728bc511fb..9526f7ec61009 100644
--- a/llvm/lib/SYCLLowerIR/ComputeModuleRuntimeInfo.cpp
+++ b/llvm/lib/SYCLLowerIR/ComputeModuleRuntimeInfo.cpp
@@ -50,6 +50,10 @@ bool isModuleUsingAsan(const Module &M) {
   return M.getNamedGlobal("__AsanKernelMetadata");
 }
 
+bool isModuleUsingMsan(const Module &M) {
+  return M.getNamedGlobal("__MsanKernelMetadata");
+}
+
 // This function traverses over reversed call graph by BFS algorithm.
 // It means that an edge links some function @func with functions
 // which contain call of function @func. It starts from
@@ -399,7 +403,9 @@ PropSetRegTy computeModuleProperties(const Module &M,
 
   {
     if (isModuleUsingAsan(M))
-      PropSet.add(PropSetRegTy::SYCL_MISC_PROP, "asanUsed", true);
+      PropSet.add(PropSetRegTy::SYCL_MISC_PROP, "sanUsed", "asan");
+    else if (isModuleUsingMsan(M))
+      PropSet.add(PropSetRegTy::SYCL_MISC_PROP, "sanUsed", "msan");
   }
 
   if (GlobProps.EmitDeviceGlobalPropSet) {
diff --git a/llvm/lib/SYCLLowerIR/AsanKernelMetadata.cpp b/llvm/lib/SYCLLowerIR/SanitizerKernelMetadata.cpp
similarity index 70%
rename from llvm/lib/SYCLLowerIR/AsanKernelMetadata.cpp
rename to llvm/lib/SYCLLowerIR/SanitizerKernelMetadata.cpp
index 2edaff49c471d..c99fc0ed83509 100644
--- a/llvm/lib/SYCLLowerIR/AsanKernelMetadata.cpp
+++ b/llvm/lib/SYCLLowerIR/SanitizerKernelMetadata.cpp
@@ -1,4 +1,4 @@
-//===-- AsanKernelMetadata.cpp - fix kernel medatadata for sanitizer -===//
+//===-- SanitizerKernelMetadata.cpp - fix kernel medatadata for sanitizer -===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
@@ -6,17 +6,17 @@
 //
 //===----------------------------------------------------------------------===//
 // This pass fixes attributes and metadata of global variable
-// "__AsanKernelMetadata".
-// We treat "__AsanKernelMetadata" as a device global variable, so that it can
-// be read by runtime.
+// "__AsanKernelMetadata" or "__MsanKernelMetadata".
+// We treat "KernelMetadata" as a device global variable, so that it
+// can be read by runtime.
 // "spirv.Decorations" is removed by llvm-link, so we add it here again.
 //===----------------------------------------------------------------------===//
 
-#include "llvm/SYCLLowerIR/AsanKernelMetadata.h"
+#include "llvm/SYCLLowerIR/SanitizerKernelMetadata.h"
 
 #include "llvm/IR/IRBuilder.h"
 
-#define DEBUG_TYPE "AsanKernelMetadata"
+#define DEBUG_TYPE "SanitizerKernelMetadata"
 
 using namespace llvm;
 
@@ -25,12 +25,15 @@ namespace llvm {
 constexpr StringRef SPIRV_DECOR_MD_KIND = "spirv.Decorations";
 constexpr uint32_t SPIRV_HOST_ACCESS_DECOR = 6147;
 
-PreservedAnalyses AsanKernelMetadataPass::run(Module &M,
-                                              ModuleAnalysisManager &MAM) {
+PreservedAnalyses SanitizerKernelMetadataPass::run(Module &M,
+                                                   ModuleAnalysisManager &MAM) {
   auto *KernelMetadata = M.getNamedGlobal("__AsanKernelMetadata");
-  if (!KernelMetadata) {
+
+  if (!KernelMetadata)
+    KernelMetadata = M.getNamedGlobal("__MsanKernelMetadata");
+
+  if (!KernelMetadata)
     return PreservedAnalyses::all();
-  }
 
   auto &DL = M.getDataLayout();
   auto &Ctx = M.getContext();
@@ -51,7 +54,7 @@ PreservedAnalyses AsanKernelMetadataPass::run(Module &M,
       Constant::getIntegerValue(Ty, APInt(32, SPIRV_HOST_ACCESS_DECOR))));
   MD.push_back(
       ConstantAsMetadata::get(Constant::getIntegerValue(Ty, APInt(32, 0))));
-  MD.push_back(MDString::get(Ctx, "_Z20__AsanKernelMetadata"));
+  MD.push_back(MDString::get(Ctx, "_Z20__SanitizerKernelMetadata"));
 
   MDOps.push_back(MDNode::get(Ctx, MD));
 
diff --git a/llvm/lib/Transforms/Instrumentation/MemorySanitizer.cpp b/llvm/lib/Transforms/Instrumentation/MemorySanitizer.cpp
index 2e6568225b908..27fcf66186453 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"
@@ -483,6 +485,14 @@ static const MemoryMapParams NetBSD_X86_64_MemoryMapParams = {
     0x100000000000, // OriginBase
 };
 
+// SPIR64 Intel
+static const MemoryMapParams Intel_SPIR64_MemoryMapParams = {
+    0, // AndMask
+    0, // XorMask
+    0, // ShadowBase
+    0, // OriginBase
+};
+
 static const PlatformMemoryMapParams Linux_X86_MemoryMapParams = {
     &Linux_I386_MemoryMapParams,
     &Linux_X86_64_MemoryMapParams,
@@ -528,6 +538,18 @@ static const PlatformMemoryMapParams NetBSD_X86_MemoryMapParams = {
     &NetBSD_X86_64_MemoryMapParams,
 };
 
+static const PlatformMemoryMapParams Intel_SPIR_MemoryMapParams = {
+    nullptr,
+    &Intel_SPIR64_MemoryMapParams,
+};
+
+// Spir memory address space
+static constexpr unsigned kSpirOffloadPrivateAS = 0;
+static constexpr unsigned kSpirOffloadGlobalAS = 1;
+static constexpr unsigned kSpirOffloadConstantAS = 2;
+static constexpr unsigned kSpirOffloadLocalAS = 3;
+static constexpr unsigned kSpirOffloadGenericAS = 4;
+
 namespace {
 
 /// Instrument functions of a module to detect uninitialized reads.
@@ -655,6 +677,9 @@ class MemorySanitizer {
   FunctionCallee MsanMetadataPtrForStore_1_8[4];
   FunctionCallee MsanInstrumentAsmStoreFn;
 
+  /// Get shadow memory address
+  FunctionCallee MsanGetShadowFn;
+
   /// Storage for return values of the MsanMetadataPtrXxx functions.
   Value *MsanMetadataAlloca;
 
@@ -705,17 +730,83 @@ MemorySanitizerOptions::MemorySanitizerOptions(int TO, bool R, bool K,
       Recover(getOptOrDefault(ClKeepGoing, Kernel || R)),
       EagerChecks(getOptOrDefault(ClEagerChecks, EagerChecks)) {}
 
+Constant *getOrCreateGlobalString(Module &M, StringRef Name, StringRef Value,
+                                  unsigned AddressSpace) {
+  auto StringName = (Twine(Name) + "_" + Value).str();
+  auto *Ty = ArrayType::get(Type::getInt8Ty(M.getContext()), Value.size() + 1);
+
+  return M.getOrInsertGlobal(StringName, Ty, [&] {
+    return new GlobalVariable(
+        M, Ty, true, GlobalValue::InternalLinkage,
+        ConstantDataArray::getString(M.getContext(), Value), StringName,
+        nullptr, GlobalValue::NotThreadLocal, AddressSpace);
+  });
+}
+
+static void extendSpirKernelArgs(Module &M) {
+  SmallVector<Constant *, 8> SpirKernelsMetadata;
+
+  auto DL = M.getDataLayout();
+  Type *IntptrTy = DL.getIntPtrType(M.getContext());
+
+  // SpirKernelsMetadata only saves fixed kernels, and is described by
+  // following structure:
+  //  uptr unmangled_kernel_name
+  //  uptr unmangled_kernel_name_size
+  StructType *StructTy = StructType::get(IntptrTy, IntptrTy);
+  for (Function &F : M) {
+    if (F.getCallingConv() != CallingConv::SPIR_KERNEL)
+      continue;
+
+    if (!F.hasFnAttribute(Attribute::SanitizeMemory) ||
+        F.hasFnAttribute(Attribute::DisableSanitizerInstrumentation))
+      continue;
+
+    auto KernelName = F.getName();
+    auto *KernelNameGV = getOrCreateGlobalString(M, "__msan_kernel", KernelName,
+                                                 kSpirOffloadGlobalAS);
+    SpirKernelsMetadata.emplace_back(ConstantStruct::get(
+        StructTy, ConstantExpr::getPointerCast(KernelNameGV, IntptrTy),
+        ConstantInt::get(IntptrTy, KernelName.size())));
+  }
+
+  // Create global variable to record spirv kernels' information
+  ArrayType *ArrayTy = ArrayType::get(StructTy, SpirKernelsMetadata.size());
+  Constant *MetadataInitializer =
+      ConstantArray::get(ArrayTy, SpirKernelsMetadata);
+  GlobalVariable *MsanSpirKernelMetadata = new GlobalVariable(
+      M, MetadataInitializer->getType(), false, GlobalValue::AppendingLinkage,
+      MetadataInitializer, "__MsanKernelMetadata", nullptr,
+      GlobalValue::NotThreadLocal, 1);
+  MsanSpirKernelMetadata->setUnnamedAddr(GlobalValue::UnnamedAddr::Local);
+  // Add device global attributes
+  MsanSpirKernelMetadata->addAttribute(
+      "sycl-device-global-size", std::to_string(DL.getTypeAllocSize(ArrayTy)));
+  MsanSpirKernelMetadata->addAttribute("sycl-device-image-scope");
+  MsanSpirKernelMetadata->addAttribute("sycl-host-access", "0"); // read only
+  MsanSpirKernelMetadata->addAttribute("sycl-unique-id",
+                                       "_Z20__MsanKernelMetadata");
+  MsanSpirKernelMetadata->setDSOLocal(true);
+}
+
 PreservedAnalyses MemorySanitizerPass::run(Module &M,
                                            ModuleAnalysisManager &AM) {
   // Return early if nosanitize_memory module flag is present for the module.
   if (checkIfAlreadyInstrumented(M, "nosanitize_memory"))
     return PreservedAnalyses::all();
   bool Modified = false;
-  if (!Options.Kernel) {
+  auto TargetTriple = Triple(M.getTargetTriple());
+
+  if (!Options.Kernel && !TargetTriple.isSPIROrSPIRV()) {
     insertModuleCtor(M);
     Modified = true;
   }
 
+  if (TargetTriple.isSPIROrSPIRV()) {
+    extendSpirKernelArgs(M);
+    Modified = true;
+  }
+
   auto &FAM = AM.getResult<FunctionAnalysisManagerModuleProxy>(M).getManager();
   for (Function &F : M) {
     if (F.empty())
@@ -888,9 +979,24 @@ void MemorySanitizer::createUserspaceApi(Module &M, const TargetLibraryInfo &TLI
        AccessSizeIndex++) {
     unsigned AccessSize = 1 << AccessSizeIndex;
     std::string FunctionName = "__msan_maybe_warning_" + itostr(AccessSize);
-    MaybeWarningFn[AccessSizeIndex] = M.getOrInsertFunction(
-        FunctionName, TLI.getAttrList(C, {0, 1}, /*Signed=*/false),
-        IRB.getVoidTy(), IRB.getIntNTy(AccessSize * 8), IRB.getInt32Ty());
+
+    if (!TargetTriple.isSPIROrSPIRV()) {
+      MaybeWarningFn[AccessSizeIndex] = M.getOrInsertFunction(
+          FunctionName, TLI.getAttrList(C, {0, 1}, /*Signed=*/false),
+          IRB.getVoidTy(), IRB.getIntNTy(AccessSize * 8), IRB.getInt32Ty());
+    } else { // SPIR or SPIR-V
+      // __msan_maybe_warning_N(
+      //   ...
+      //   char* file,
+      //   unsigned int line,
+      //   char* func
+      // )
+      MaybeWarningFn[AccessSizeIndex] = M.getOrInsertFunction(
+          FunctionName, TLI.getAttrList(C, {0, 1}, /*Signed=*/false),
+          IRB.getVoidTy(), IRB.getIntNTy(AccessSize * 8), IRB.getInt32Ty(),
+          IRB.getInt8PtrTy(kSpirOffloadConstantAS), IRB.getInt32Ty(),
+          IRB.getInt8PtrTy(kSpirOffloadConstantAS));
+    }
 
     FunctionName = "__msan_maybe_store_origin_" + itostr(AccessSize);
     MaybeStoreOriginFn[AccessSizeIndex] = M.getOrInsertFunction(
@@ -936,6 +1042,9 @@ void MemorySanitizer::initializeCallbacks(Module &M, const TargetLibraryInfo &TL
   MsanInstrumentAsmStoreFn = M.getOrInsertFunction(
       "__msan_instrument_asm_store", IRB.getVoidTy(), PtrTy, IntptrTy);
 
+  MsanGetShadowFn = M.getOrInsertFunction("__msan_get_shadow", IntptrTy,
+                                          IntptrTy, IRB.getInt32Ty());
+
   if (CompileKernel) {
     createKernelApi(M, TLI);
   } else {
@@ -1035,6 +1144,12 @@ void MemorySanitizer::initializeModule(Module &M) {
         report_fatal_error("unsupported architecture");
       }
       break;
+    case Triple::UnknownOS:
+      // NOTE: Support SPIR or SPIRV only, without MapParams
+      if (!TargetTriple.isSPIROrSPIRV())
+        report_fatal_error("unsupported architecture");
+      MapParams = Intel_SPIR_MemoryMapParams.bits64;
+      break;
     default:
       report_fatal_error("unsupported operating system");
     }
@@ -1112,6 +1227,72 @@ static unsigned TypeSizeToSizeIndex(TypeSize TS) {
   return Log2_32_Ceil((TypeSizeFixed + 7) / 8);
 }
 
+static bool isUnsupportedSPIRAccess(const Value *Addr, Instruction *I) {
+  // Skip SPIR-V built-in varibles
+  auto *OrigValue = Addr->stripInBoundsOffsets();
+  assert(OrigValue != nullptr);
+  if (OrigValue->getName().starts_with("__spirv_BuiltIn"))
+    return true;
+
+  Type *PtrTy = cast<PointerType>(Addr->getType()->getScalarType());
+  switch (PtrTy->getPointerAddressSpace()) {
+  case kSpirOffloadPrivateAS:
+  case kSpirOffloadLocalAS:
+  case kSpirOffloadGenericAS:
+    return true;
+  }
+
+  return false;
+}
+
+static void setNoSanitizedMetadataSPIR(Instruction &I) {
+  const Value *Addr = nullptr;
+  if (const auto *LI = dyn_cast<LoadInst>(&I))
+    Addr = LI->getPointerOperand();
+  else if (const auto *SI = dyn_cast<StoreInst>(&I))
+    Addr = SI->getPointerOperand();
+  else if (const auto *RMW = dyn_cast<AtomicRMWInst>(&I))
+    Addr = RMW->getPointerOperand();
+  else if (const auto *XCHG = dyn_cast<AtomicCmpXchgInst>(&I))
+    Addr = XCHG->getPointerOperand();
+  else if (isa<MemCpyInst>(&I))
+    I.setNoSanitizeMetadata();
+  else if (const auto *CI = dyn_cast<CallInst>(&I)) {
+    auto *Func = CI->getCalledFunction();
+    if (Func) {
+      if (Func->isIntrinsic()) {
+        switch (CI->getIntrinsicID()) {
+        case Intrinsic::masked_load:
+        case Intrinsic::masked_store:
+        case Intrinsic::masked_gather:
+        case Intrinsic::masked_scatter: {
+          bool IsWrite = CI->getType()->isVoidTy();
+          // Masked store has an initial operand for the value.
+          unsigned OpOffset = IsWrite ? 1 : 0;
+          Addr = CI->getOperand(OpOffset);
+          break;
+        }
+        case Intrinsic::masked_expandload:
+        case Intrinsic::masked_compressstore: {
+          bool IsWrite =
+              CI->getIntrinsicID() == Intrinsic::masked_compressstore;
+          unsigned OpOffset = IsWrite ? 1 : 0;
+          Addr = CI->getOperand(OpOffset);
+          break;
+        }
+        }
+      } else {
+        auto FuncName = Func->getName();
+        if (FuncName.contains("__spirv_"))
+          I.setNoSanitizeMetadata();
+      }
+    }
+  }
+
+  if (Addr && isUnsupportedSPIRAccess(Addr, &I))
+    I.setNoSanitizeMetadata();
+}
+
 namespace {
 
 /// Helper class to attach debug information of the given instruction onto new
@@ -1146,6 +1327,8 @@ struct MemorySanitizerVisitor : public InstVisitor<MemorySanitizerVisitor> {
   bool PoisonStack;
   bool PoisonUndef;
 
+  bool SpirOrSpirv;
+
   struct ShadowOriginAndInsertPoint {
     Value *Shadow;
     Value *Origin;
@@ -1171,6 +1354,7 @@ struct MemorySanitizerVisitor : public InstVisitor<MemorySanitizerVisitor> {
     PropagateShadow = SanitizeFunction;
     PoisonStack = SanitizeFunction && ClPoisonStack;
     PoisonUndef = SanitizeFunction && ClPoisonUndef;
+    SpirOrSpirv = Triple(F.getParent()->getTargetTriple()).isSPIROrSPIRV();
 
     // In the presence of unreachable blocks, we may see Phi nodes with
     // incoming nodes from such blocks. Since InstVisitor skips unreachable
@@ -1402,11 +1586,60 @@ struct MemorySanitizerVisitor : public InstVisitor<MemorySanitizerVisitor> {
       ConvertedShadow = convertShadowToScalar(ConvertedShadow, IRB);
       Value *ConvertedShadow2 =
           IRB.CreateZExt(ConvertedShadow, IRB.getIntNTy(8 * (1 << SizeIndex)));
-      CallBase *CB = IRB.CreateCall(
-          Fn, {ConvertedShadow2,
-               MS.TrackOrigins && Origin ? Origin : (Value *)IRB.getInt32(0)});
-      CB->addParamAttr(0, Attribute::ZExt);
-      CB->addParamAttr(1, Attribute::ZExt);
+      if (!SpirOrSpirv) {
+        CallBase *CB = IRB.CreateCall(
+            Fn,
+            {ConvertedShadow2,
+             MS.TrackOrigins && Origin ? Origin : (Value *)IRB.getInt32(0)});
+        CB->addParamAttr(0, Attribute::ZExt);
+        CB->addParamAttr(1, Attribute::ZExt);
+      } else { // SPIR or SPIR-V
+        SmallVector<Value *, 5> Args = {
+            ConvertedShadow2,
+            MS.TrackOrigins && Origin ? Origin : (Value *)IRB.getInt32(0)};
+
+        {
+          auto *M = F.getParent();
+          auto *ConstASPtrTy = IRB.getInt8PtrTy(kSpirOffloadConstantAS);
+
+          // file name and line number
+          {
+            bool HasDebugLoc = false;
+            auto *ConvertedShadowInst = dyn_cast<Instruction>(ConvertedShadow);
+
+            if (ConvertedShadowInst) {
+              if (auto &Loc = ConvertedShadowInst->getDebugLoc()) {
+                llvm::SmallString<128> Source = Loc->getDirectory();
+                sys::path::append(Source, Loc->getFilename());
+                auto *FileNameGV = getOrCreateGlobalString(
+                    *M, "__asan_file", Source, kSpirOffloadConstantAS);
+                Args.push_back(
+                    ConstantExpr::getPointerCast(FileNameGV, ConstASPtrTy));
+                Args.push_back(
+                    ConstantInt::get(IRB.getInt32Ty(), Loc.getLine()));
+
+                HasDebugLoc = true;
+              }
+            }
+
+            if (!HasDebugLoc) {
+              Args.push_back(ConstantPointerNull::get(ConstASPtrTy));
+              Args.push_back(ConstantInt::get(IRB.getInt32Ty(), 0));
+            }
+          }
+
+          // function name
+          auto FuncName = F.getName();
+          auto *FuncNameGV = getOrCreateGlobalString(
+              *M, "__asan_func", demangle(FuncName), kSpirOffloadConstantAS);
+          Args.push_back(
+              ConstantExpr::getPointerCast(FuncNameGV, ConstASPtrTy));
+        }
+
+        CallBase *CB = IRB.CreateCall(Fn, Args);
+        CB->addParamAttr(0, Attribute::ZExt);
+        CB->addParamAttr(1, Attribute::ZExt);
+      }
     } else {
       Value *Cmp = convertToBool(ConvertedShadow, IRB, "_mscmp");
       Instruction *CheckTerm = SplitBlockAndInsertIfThen(
@@ -1566,6 +1799,10 @@ struct MemorySanitizerVisitor : public InstVisitor<MemorySanitizerVisitor> {
     // This may not add new address checks.
     materializeStores();
 
+    // Fix unexpected llvm intrinsic
+    if (SpirOrSpirv)
+      FnPrologueEnd->eraseFromParent();
+
     return true;
   }
 
@@ -1678,14 +1915,15 @@ struct MemorySanitizerVisitor : public InstVisitor<MemorySanitizerVisitor> {
     return MS.IntptrTy;
   }
 
-  Type *getPtrToShadowPtrType(Type *IntPtrTy, Type *ShadowTy) const {
+  Type *getPtrToShadowPtrType(Type *IntPtrTy, Type *ShadowTy,
+                              unsigned int AddressSapce = 0) const {
     if (VectorType *VectTy = dyn_cast<VectorType>(IntPtrTy)) {
       return VectorType::get(
           getPtrToShadowPtrType(VectTy->getElementType(), ShadowTy),
           VectTy->getElementCount());
     }
     assert(IntPtrTy == MS.IntptrTy);
-    return MS.PtrTy;
+    return PointerType::get(*MS.C, AddressSapce);
   }
 
   Constant *constToIntPtr(Type *IntPtrTy, uint64_t C) const {
@@ -1708,11 +1946,21 @@ struct MemorySanitizerVisitor : public InstVisitor<MemorySanitizerVisitor> {
     Type *IntptrTy = ptrToIntPtrType(Addr->getType());
     Value *OffsetLong = IRB.CreatePointerCast(Addr, IntptrTy);
 
-    if (uint64_t AndMask = MS.MapParams->AndMask)
-      OffsetLong = IRB.CreateAnd(OffsetLong, constToIntPtr(IntptrTy, ~AndMask));
+    if (!SpirOrSpirv) {
+      if (uint64_t AndMask = MS.MapParams->AndMask)
+        OffsetLong =
+            IRB.CreateAnd(OffsetLong, constToIntPtr(IntptrTy, ~AndMask));
+
+      if (uint64_t XorMask = MS.MapParams->XorMask)
+        OffsetLong =
+            IRB.CreateXor(OffsetLong, constToIntPtr(IntptrTy, XorMask));
+    } else { // SPIR or SPIR-V
+      OffsetLong = IRB.CreateCall(
+          MS.MsanGetShadowFn,
+          {OffsetLong,
+           IRB.getInt32(Addr->getType()->getPointerAddressSpace())});
+    }
 
-    if (uint64_t XorMask = MS.MapParams->XorMask)
-      OffsetLong = IRB.CreateXor(OffsetLong, constToIntPtr(IntptrTy, XorMask));
     return OffsetLong;
   }
 
@@ -1741,7 +1989,9 @@ struct MemorySanitizerVisitor : public InstVisitor<MemorySanitizerVisitor> {
           IRB.CreateAdd(ShadowLong, constToIntPtr(IntptrTy, ShadowBase));
     }
     Value *ShadowPtr = IRB.CreateIntToPtr(
-        ShadowLong, getPtrToShadowPtrType(IntptrTy, ShadowTy));
+        ShadowLong,
+        getPtrToShadowPtrType(IntptrTy, ShadowTy,
+                              SpirOrSpirv ? kSpirOffloadGlobalAS : 0));
 
     Value *OriginPtr = nullptr;
     if (MS.TrackOrigins) {
@@ -2195,6 +2445,9 @@ struct MemorySanitizerVisitor : public InstVisitor<MemorySanitizerVisitor> {
   // ------------------- Visitors.
   using InstVisitor<MemorySanitizerVisitor>::visit;
   void visit(Instruction &I) {
+    if (SpirOrSpirv)
+      setNoSanitizedMetadataSPIR(I);
+
     if (I.getMetadata(LLVMContext::MD_nosanitize))
       return;
     // Don't want to visit if we're in the prologue
diff --git a/llvm/test/Instrumentation/MemorySanitizer/SPIRV/instrument_global_address_space.ll b/llvm/test/Instrumentation/MemorySanitizer/SPIRV/instrument_global_address_space.ll
new file mode 100644
index 0000000000000..51c20c938c72d
--- /dev/null
+++ b/llvm/test/Instrumentation/MemorySanitizer/SPIRV/instrument_global_address_space.ll
@@ -0,0 +1,143 @@
+; RUN: opt < %s -passes=msan -msan-instrumentation-with-call-threshold=0 -msan-eager-checks=1 -S | FileCheck %s
+
+; ModuleID = 'check_call.cpp'
+source_filename = "check_call.cpp"
+target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64-G1"
+target triple = "spir64-unknown-unknown"
+
+$_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_E8MyKernel = comdat any
+
+; CHECK: @__MsanKernelMetadata = appending dso_local local_unnamed_addr addrspace(1) global
+; CHECK-SAME: [[ATTR0:#[0-9]+]]
+
+; Function Attrs: mustprogress norecurse nounwind sanitize_memory uwtable
+define weak_odr dso_local spir_kernel void @_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_E8MyKernel(ptr addrspace(1) noundef align 4 %_arg_array) local_unnamed_addr #0 comdat !srcloc !85 !kernel_arg_buffer_location !86 !sycl_fixed_targets !87 {
+; CHECK-LABEL: @_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_E8MyKernel
+entry:
+; CHECK-NOT: @__msan_param_tls
+  %0 = load i32, ptr addrspace(1) %_arg_array, align 4, !tbaa !88
+  %arrayidx3.i = getelementptr inbounds i8, ptr addrspace(1) %_arg_array, i64 4
+; CHECK: @__msan_get_shadow
+  %1 = load i32, ptr addrspace(1) %arrayidx3.i, align 4, !tbaa !88
+  %conv.i = sext i32 %1 to i64
+  %call.i = tail call spir_func noundef i64 @_Z3fooix(i32 noundef %0, i64 noundef %conv.i) #2
+  %conv4.i = trunc i64 %call.i to i32
+  store i32 %conv4.i, ptr addrspace(1) %_arg_array, align 4, !tbaa !88
+  ret void
+}
+
+; Function Attrs: mustprogress noinline norecurse nounwind sanitize_memory uwtable
+define linkonce_odr dso_local spir_func noundef i64 @_Z3fooix(i32 noundef %data1, i64 noundef %data2) local_unnamed_addr #1 !srcloc !92 {
+; CHECK-LABEL: @_Z3fooix
+entry:
+  %conv = sext i32 %data1 to i64
+  %add = add nsw i64 %data2, %conv
+  ret i64 %add
+}
+
+; CHECK: attributes [[ATTR0]]
+; CHECK-SAME: "sycl-device-global-size"="16" "sycl-device-image-scope" "sycl-host-access"="0" "sycl-unique-id"="_Z20__MsanKernelMetadata"
+
+attributes #0 = { mustprogress norecurse nounwind sanitize_memory uwtable "frame-pointer"="all" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="check_call.cpp" "sycl-single-task" "uniform-work-group-size"="true" }
+attributes #1 = { mustprogress noinline norecurse nounwind sanitize_memory uwtable "frame-pointer"="all" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
+attributes #2 = { nounwind }
+
+!llvm.module.flags = !{!0, !1, !2, !3}
+!opencl.spir.version = !{!4}
+!spirv.Source = !{!5}
+!sycl_aspects = !{!6, !7, !8, !9, !10, !11, !12, !13, !14, !15, !16, !17, !18, !19, !20, !21, !22, !23, !24, !25, !26, !27, !28, !29, !30, !31, !32, !33, !34, !35, !36, !37, !38, !39, !40, !41, !42, !43, !44, !45, !46, !47, !48, !49, !50, !51, !52, !53, !54, !55, !56, !57, !58, !59, !60, !61, !62, !63, !64, !65, !66, !67, !68, !69, !70, !71, !72, !73, !74, !75, !76, !77, !78, !79, !80, !81, !82, !83}
+!llvm.ident = !{!84}
+
+!0 = !{i32 1, !"wchar_size", i32 4}
+!1 = !{i32 1, !"sycl-device", i32 1}
+!2 = !{i32 7, !"uwtable", i32 2}
+!3 = !{i32 7, !"frame-pointer", i32 2}
+!4 = !{i32 1, i32 2}
+!5 = !{i32 4, i32 100000}
+!6 = !{!"cpu", i32 1}
+!7 = !{!"gpu", i32 2}
+!8 = !{!"accelerator", i32 3}
+!9 = !{!"custom", i32 4}
+!10 = !{!"fp16", i32 5}
+!11 = !{!"fp64", i32 6}
+!12 = !{!"image", i32 9}
+!13 = !{!"online_compiler", i32 10}
+!14 = !{!"online_linker", i32 11}
+!15 = !{!"queue_profiling", i32 12}
+!16 = !{!"usm_device_allocations", i32 13}
+!17 = !{!"usm_host_allocations", i32 14}
+!18 = !{!"usm_shared_allocations", i32 15}
+!19 = !{!"usm_system_allocations", i32 17}
+!20 = !{!"ext_intel_pci_address", i32 18}
+!21 = !{!"ext_intel_gpu_eu_count", i32 19}
+!22 = !{!"ext_intel_gpu_eu_simd_width", i32 20}
+!23 = !{!"ext_intel_gpu_slices", i32 21}
+!24 = !{!"ext_intel_gpu_subslices_per_slice", i32 22}
+!25 = !{!"ext_intel_gpu_eu_count_per_subslice", i32 23}
+!26 = !{!"ext_intel_max_mem_bandwidth", i32 24}
+!27 = !{!"ext_intel_mem_channel", i32 25}
+!28 = !{!"usm_atomic_host_allocations", i32 26}
+!29 = !{!"usm_atomic_shared_allocations", i32 27}
+!30 = !{!"atomic64", i32 28}
+!31 = !{!"ext_intel_device_info_uuid", i32 29}
+!32 = !{!"ext_oneapi_srgb", i32 30}
+!33 = !{!"ext_oneapi_native_assert", i32 31}
+!34 = !{!"host_debuggable", i32 32}
+!35 = !{!"ext_intel_gpu_hw_threads_per_eu", i32 33}
+!36 = !{!"ext_oneapi_cuda_async_barrier", i32 34}
+!37 = !{!"ext_intel_free_memory", i32 36}
+!38 = !{!"ext_intel_device_id", i32 37}
+!39 = !{!"ext_intel_memory_clock_rate", i32 38}
+!40 = !{!"ext_intel_memory_bus_width", i32 39}
+!41 = !{!"emulated", i32 40}
+!42 = !{!"ext_intel_legacy_image", i32 41}
+!43 = !{!"ext_oneapi_bindless_images", i32 42}
+!44 = !{!"ext_oneapi_bindless_images_shared_usm", i32 43}
+!45 = !{!"ext_oneapi_bindless_images_1d_usm", i32 44}
+!46 = !{!"ext_oneapi_bindless_images_2d_usm", i32 45}
+!47 = !{!"ext_oneapi_external_memory_import", i32 46}
+!48 = !{!"ext_oneapi_external_semaphore_import", i32 48}
+!49 = !{!"ext_oneapi_mipmap", i32 50}
+!50 = !{!"ext_oneapi_mipmap_anisotropy", i32 51}
+!51 = !{!"ext_oneapi_mipmap_level_reference", i32 52}
+!52 = !{!"ext_intel_esimd", i32 53}
+!53 = !{!"ext_oneapi_ballot_group", i32 54}
+!54 = !{!"ext_oneapi_fixed_size_group", i32 55}
+!55 = !{!"ext_oneapi_opportunistic_group", i32 56}
+!56 = !{!"ext_oneapi_tangle_group", i32 57}
+!57 = !{!"ext_intel_matrix", i32 58}
+!58 = !{!"ext_oneapi_is_composite", i32 59}
+!59 = !{!"ext_oneapi_is_component", i32 60}
+!60 = !{!"ext_oneapi_graph", i32 61}
+!61 = !{!"ext_intel_fpga_task_sequence", i32 62}
+!62 = !{!"ext_oneapi_limited_graph", i32 63}
+!63 = !{!"ext_oneapi_private_alloca", i32 64}
+!64 = !{!"ext_oneapi_cubemap", i32 65}
+!65 = !{!"ext_oneapi_cubemap_seamless_filtering", i32 66}
+!66 = !{!"ext_oneapi_bindless_sampled_image_fetch_1d_usm", i32 67}
+!67 = !{!"ext_oneapi_bindless_sampled_image_fetch_1d", i32 68}
+!68 = !{!"ext_oneapi_bindless_sampled_image_fetch_2d_usm", i32 69}
+!69 = !{!"ext_oneapi_bindless_sampled_image_fetch_2d", i32 70}
+!70 = !{!"ext_oneapi_bindless_sampled_image_fetch_3d", i32 72}
+!71 = !{!"ext_oneapi_queue_profiling_tag", i32 73}
+!72 = !{!"ext_oneapi_virtual_mem", i32 74}
+!73 = !{!"ext_oneapi_cuda_cluster_group", i32 75}
+!74 = !{!"ext_oneapi_image_array", i32 76}
+!75 = !{!"ext_oneapi_unique_addressing_per_dim", i32 77}
+!76 = !{!"ext_oneapi_bindless_images_sample_1d_usm", i32 78}
+!77 = !{!"ext_oneapi_bindless_images_sample_2d_usm", i32 79}
+!78 = !{!"ext_oneapi_atomic16", i32 80}
+!79 = !{!"ext_oneapi_virtual_functions", i32 81}
+!80 = !{!"host", i32 0}
+!81 = !{!"int64_base_atomics", i32 7}
+!82 = !{!"int64_extended_atomics", i32 8}
+!83 = !{!"usm_restricted_shared_allocations", i32 16}
+!84 = !{!"clang version 20.0.0git (https://github.com/intel/llvm.git 7384106e6410c6f038b2a9d6367a32b55278c638)"}
+!85 = !{i32 563}
+!86 = !{i32 -1}
+!87 = !{}
+!88 = !{!89, !89, i64 0}
+!89 = !{!"int", !90, i64 0}
+!90 = !{!"omnipotent char", !91, i64 0}
+!91 = !{!"Simple C++ TBAA"}
+!92 = !{i32 345}
diff --git a/llvm/test/tools/sycl-post-link/device-sanitizer/asan.ll b/llvm/test/tools/sycl-post-link/device-sanitizer/asan.ll
index 7fa89b236fa03..e956876032084 100644
--- a/llvm/test/tools/sycl-post-link/device-sanitizer/asan.ll
+++ b/llvm/test/tools/sycl-post-link/device-sanitizer/asan.ll
@@ -1,11 +1,11 @@
-; This test checks that the post-link tool properly generates "asanUsed=1"
+; This test checks that the post-link tool properly generates "sanUsed=asan"
 ; in prop file, and fixes the attributes and metadata of @__AsanKernelMetadata
 
 ; RUN: sycl-post-link -properties -split=kernel -symbols -S < %s -o %t.table
 
 ; RUN: FileCheck %s -input-file=%t_0.prop --check-prefix CHECK-PROP
 ; CHECK-PROP: [SYCL/misc properties]
-; CHECK-PROP: asanUsed=1
+; CHECK-PROP: sanUsed=2|gAAAAAAAAAQYzFmb
 
 ; RUN: FileCheck %s -input-file=%t_0.ll --check-prefix CHECK-IR
 
diff --git a/llvm/test/tools/sycl-post-link/device-sanitizer/msan.ll b/llvm/test/tools/sycl-post-link/device-sanitizer/msan.ll
new file mode 100644
index 0000000000000..41110bb30af2f
--- /dev/null
+++ b/llvm/test/tools/sycl-post-link/device-sanitizer/msan.ll
@@ -0,0 +1,84 @@
+; This test checks that the post-link tool properly generates "sanUsed=msan"
+; in [SYCL/misc properties], and fixes the attributes and metadata of @__MsanKernelMetadata
+
+; RUN: sycl-post-link -properties -split=kernel -symbols -S < %s -o %t.table
+
+; RUN: FileCheck %s -input-file=%t_0.prop --check-prefix CHECK-PROP
+; CHECK-PROP: [SYCL/misc properties]
+; CHECK-PROP: sanUsed=2|gAAAAAAAAAQbzFmb
+
+; RUN: FileCheck %s -input-file=%t_0.ll --check-prefix CHECK-IR
+
+; ModuleID = 'check_call.cpp'
+source_filename = "check_call.cpp"
+target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64"
+target triple = "spir64-unknown-unknown"
+
+$_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_E8MyKernel = comdat any
+
+@__msan_kernel = internal addrspace(1) constant [55 x i8] c"_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_E8MyKernel\00"
+@__MsanKernelMetadata = appending dso_local local_unnamed_addr addrspace(1) global [1 x { i64, i64 }] [{ i64, i64 } { i64 ptrtoint (ptr addrspace(1) @__msan_kernel to i64), i64 54 }] #0
+; CHECK-IR: @__MsanKernelMetadata {{.*}} !spirv.Decorations
+@__spirv_BuiltInGlobalInvocationId = external dso_local local_unnamed_addr addrspace(1) constant <3 x i64>, align 32
+@__asan_func = internal addrspace(2) constant [106 x i8] c"typeinfo name for main::'lambda'(sycl::_V1::handler&)::operator()(sycl::_V1::handler&) const::MyKernelR_4\00"
+
+; Function Attrs: mustprogress norecurse nounwind sanitize_memory uwtable
+define weak_odr dso_local spir_kernel void @_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_E8MyKernel(ptr addrspace(1) noundef align 4 %_arg_array) local_unnamed_addr #1 comdat !srcloc !6 !kernel_arg_buffer_location !7 !sycl_fixed_targets !8 {
+entry:
+  %0 = load i32, ptr addrspace(1) %_arg_array, align 4
+  %1 = ptrtoint ptr addrspace(1) %_arg_array to i64
+  %2 = call i64 @__msan_get_shadow(i64 %1, i32 1)
+  %3 = inttoptr i64 %2 to ptr addrspace(1)
+  %_msld = load i32, ptr addrspace(1) %3, align 4
+  %arrayidx3.i = getelementptr inbounds i8, ptr addrspace(1) %_arg_array, i64 4
+  %4 = load i32, ptr addrspace(1) %arrayidx3.i, align 4
+  %5 = ptrtoint ptr addrspace(1) %arrayidx3.i to i64
+  %6 = call i64 @__msan_get_shadow(i64 %5, i32 1)
+  %7 = inttoptr i64 %6 to ptr addrspace(1)
+  %_msld2 = load i32, ptr addrspace(1) %7, align 4
+  %_msprop = sext i32 %_msld2 to i64
+  %conv.i = sext i32 %4 to i64
+  %_mscmp = icmp ne i32 %_msld, 0
+  %_mscmp3 = icmp ne i64 %_msprop, 0
+  %_msor = or i1 %_mscmp, %_mscmp3
+  %8 = zext i1 %_msor to i8
+  call void @__msan_maybe_warning_1(i8 zeroext %8, i32 zeroext 0, ptr addrspace(2) null, i32 0, ptr addrspace(2) @__asan_func)
+  %call.i = tail call spir_func noundef i64 @_Z3fooix(i32 noundef %0, i64 noundef %conv.i) #4
+  %conv4.i = trunc i64 %call.i to i32
+  %9 = ptrtoint ptr addrspace(1) %_arg_array to i64
+  %10 = call i64 @__msan_get_shadow(i64 %9, i32 1)
+  %11 = inttoptr i64 %10 to ptr addrspace(1)
+  store i32 0, ptr addrspace(1) %11, align 4
+  store i32 %conv4.i, ptr addrspace(1) %_arg_array, align 4
+  ret void
+}
+
+; Function Attrs: mustprogress noinline norecurse nounwind sanitize_memory uwtable
+define linkonce_odr dso_local spir_func noundef i64 @_Z3fooix(i32 noundef %data1, i64 noundef %data2) local_unnamed_addr #2  {
+entry:
+  %conv = sext i32 %data1 to i64
+  %add = add nsw i64 %data2, %conv
+  ret i64 %add
+}
+
+declare i64 @__msan_get_shadow(i64, i32)
+declare void @__msan_maybe_warning_1(i8, i32, ptr addrspace(2), i32, ptr addrspace(2))
+
+attributes #0 = { "sycl-device-global-size"="16" "sycl-device-image-scope" "sycl-host-access"="0" "sycl-unique-id"="_Z20__MsanKernelMetadata" }
+attributes #1 = { mustprogress norecurse nounwind sanitize_memory uwtable "frame-pointer"="all" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="check_call.cpp" "sycl-single-task" "uniform-work-group-size"="true" }
+attributes #2 = { mustprogress noinline norecurse nounwind sanitize_memory uwtable "frame-pointer"="all" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
+
+!llvm.module.flags = !{!0, !1, !2}
+!opencl.spir.version = !{!3}
+!spirv.Source = !{!4}
+!llvm.ident = !{!5}
+
+!0 = !{i32 1, !"wchar_size", i32 4}
+!1 = !{i32 7, !"uwtable", i32 2}
+!2 = !{i32 7, !"frame-pointer", i32 2}
+!3 = !{i32 1, i32 2}
+!4 = !{i32 4, i32 100000}
+!5 = !{!"clang version 19.0.0git (https://github.com/intel/llvm f8eada76c08c6a5e6c5842842ac5b98fa72669be)"}
+!6 = !{i32 563}
+!7 = !{i32 -1}
+!8 = !{}
diff --git a/llvm/tools/sycl-post-link/sycl-post-link.cpp b/llvm/tools/sycl-post-link/sycl-post-link.cpp
index 91083cea614c7..6948277eabd73 100644
--- a/llvm/tools/sycl-post-link/sycl-post-link.cpp
+++ b/llvm/tools/sycl-post-link/sycl-post-link.cpp
@@ -29,7 +29,6 @@
 #include "llvm/IRReader/IRReader.h"
 #include "llvm/Linker/Linker.h"
 #include "llvm/Passes/PassBuilder.h"
-#include "llvm/SYCLLowerIR/AsanKernelMetadata.h"
 #include "llvm/SYCLLowerIR/CompileTimePropertiesPass.h"
 #include "llvm/SYCLLowerIR/ComputeModuleRuntimeInfo.h"
 #include "llvm/SYCLLowerIR/DeviceConfigFile.hpp"
@@ -41,6 +40,7 @@
 #include "llvm/SYCLLowerIR/ModuleSplitter.h"
 #include "llvm/SYCLLowerIR/SYCLJointMatrixTransform.h"
 #include "llvm/SYCLLowerIR/SYCLUtils.h"
+#include "llvm/SYCLLowerIR/SanitizerKernelMetadata.h"
 #include "llvm/SYCLLowerIR/SpecConstants.h"
 #include "llvm/SYCLLowerIR/Support.h"
 #include "llvm/Support/CommandLine.h"
@@ -800,11 +800,10 @@ processInputModule(std::unique_ptr<Module> M) {
   if (M->getTargetTriple().find("spir") != std::string::npos)
     Modified |= removeDeviceGlobalFromCompilerUsed(*M.get());
 
-  // AddressSanitizer specific passes
-  if (isModuleUsingAsan(*M)) {
-    // Fix attributes and metadata of the global variable
-    // "__AsanKernelMetadata"
-    Modified |= runModulePass<AsanKernelMetadataPass>(*M);
+  // MemorySanitizer specific passes
+  if (isModuleUsingAsan(*M) || isModuleUsingMsan(*M)) {
+    // Fix attributes and metadata of KernelMetadata
+    Modified |= runModulePass<SanitizerKernelMetadataPass>(*M);
   }
 
   // Transform Joint Matrix builtin calls to align them with SPIR-V friendly
diff --git a/sycl/cmake/modules/UnifiedRuntimeTag.cmake b/sycl/cmake/modules/UnifiedRuntimeTag.cmake
index 809852d25b41a..80bf04a00c321 100644
--- a/sycl/cmake/modules/UnifiedRuntimeTag.cmake
+++ b/sycl/cmake/modules/UnifiedRuntimeTag.cmake
@@ -1,7 +1,7 @@
-# commit 45f3d8ad95d7acada8e3bc4aa49bfae07dabcba5
-# Merge: 8f33719fa4fc 08acf05f2d66
-# Author:  Martin Grant <martin.morrisongrant@codeplay.com>
-# Date:   Wed Dec 11 09:12:33 2024 +0000
-#     Merge pull request #2430 from nrspruit/fix_inorder_initList
-#     [UR] Fix correct usage of In Order sync list given counting events
-set(UNIFIED_RUNTIME_TAG 45f3d8ad95d7acada8e3bc4aa49bfae07dabcba5)
+# commit 8818ab5bf11e94ffd967bed697b13e6f82ea87c9
+# Merge: 45f3d8ad 064da157
+# Author: Kenneth Benzie (Benie) <k.benzie@codeplay.com>
+# Date:   Wed Dec 11 11:04:29 2024 +0000
+#     Merge pull request #2325 from AllanZyne/review/yang/restructure_asan_msan
+#     [DeviceMSAN] Support MemorySanitizer for device offloading
+set(UNIFIED_RUNTIME_TAG 8818ab5bf11e94ffd967bed697b13e6f82ea87c9)
diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp
index c7a3a2491f72c..841ef9f562db5 100644
--- a/sycl/source/detail/program_manager/program_manager.cpp
+++ b/sycl/source/detail/program_manager/program_manager.cpp
@@ -1350,7 +1350,8 @@ loadDeviceLibFallback(const ContextImplPtr Context, DeviceLibExt Extension,
   return URProgram;
 }
 
-ProgramManager::ProgramManager() : m_AsanFoundInImage(false) {
+ProgramManager::ProgramManager()
+    : m_SanitizerFoundInImage(SanitizerType::None) {
   const char *SpvFile = std::getenv(UseSpvEnv);
   // If a SPIR-V file is specified with an environment variable,
   // register the corresponding image
@@ -1905,11 +1906,21 @@ void ProgramManager::addImages(sycl_device_binaries DeviceBinary) {
 
     cacheKernelUsesAssertInfo(*Img);
 
-    // check if kernel uses asan
+    // check if kernel uses sanitizer
     {
-      sycl_device_binary_property Prop = Img->getProperty("asanUsed");
-      m_AsanFoundInImage |=
-          Prop && (detail::DeviceBinaryProperty(Prop).asUint32() != 0);
+      sycl_device_binary_property SanProp = Img->getProperty("sanUsed");
+      if (SanProp) {
+        std::string SanValue =
+            detail::DeviceBinaryProperty(SanProp).asCString();
+
+        if (SanValue.rfind("asan", 0) == 0) { // starts_with
+          m_SanitizerFoundInImage = SanitizerType::AddressSanitizer;
+        } else if (SanValue.rfind("msan", 0) == 0) {
+          m_SanitizerFoundInImage = SanitizerType::MemorySanitizer;
+        } else if (SanValue.rfind("tsan", 0) == 0) {
+          m_SanitizerFoundInImage = SanitizerType::ThreadSanitizer;
+        }
+      }
     }
 
     cacheKernelImplicitLocalArg(*Img);
diff --git a/sycl/source/detail/program_manager/program_manager.hpp b/sycl/source/detail/program_manager/program_manager.hpp
index abfdb1144105b..14467a1dd26b8 100644
--- a/sycl/source/detail/program_manager/program_manager.hpp
+++ b/sycl/source/detail/program_manager/program_manager.hpp
@@ -87,6 +87,13 @@ enum class DeviceLibExt : std::uint32_t {
   cl_intel_devicelib_bfloat16,
 };
 
+enum class SanitizerType {
+  None,
+  AddressSanitizer,
+  MemorySanitizer,
+  ThreadSanitizer
+};
+
 // A helper class for storing image/program objects and their dependencies
 // and making their handling a bit more readable.
 template <typename T> class ObjectWithDeps {
@@ -330,11 +337,11 @@ class ProgramManager {
 
   bool kernelUsesAssert(const std::string &KernelName) const;
 
+  SanitizerType kernelUsesSanitizer() const { return m_SanitizerFoundInImage; }
+
   std::optional<int>
   kernelImplicitLocalArgPos(const std::string &KernelName) const;
 
-  bool kernelUsesAsan() const { return m_AsanFoundInImage; }
-
   std::set<RTDeviceBinaryImage *>
   getRawDeviceImages(const std::vector<kernel_id> &KernelIDs);
 
@@ -465,8 +472,8 @@ class ProgramManager {
   std::set<std::string> m_KernelUsesAssert;
   std::unordered_map<std::string, int> m_KernelImplicitLocalArgPos;
 
-  // True iff there is a device image compiled with AddressSanitizer
-  bool m_AsanFoundInImage;
+  // Sanitizer type used in device image
+  SanitizerType m_SanitizerFoundInImage;
 
   // Maps between device_global identifiers and associated information.
   std::unordered_map<std::string, std::unique_ptr<DeviceGlobalMapEntry>>
diff --git a/sycl/source/detail/ur.cpp b/sycl/source/detail/ur.cpp
index 74906e6588a6c..33249576265bb 100644
--- a/sycl/source/detail/ur.cpp
+++ b/sycl/source/detail/ur.cpp
@@ -70,9 +70,8 @@ void *getAdapterOpaqueData([[maybe_unused]] void *OpaqueDataParam) {
   // entry point introduced for the now deleted ESIMD adapter. All calls to this
   // entry point returned a similar error code to INVALID_OPERATION and would
   // have resulted in a similar throw to this one
-  throw exception(
-      make_error_code(errc::feature_not_supported),
-      "This operation is not supported by any existing backends.");
+  throw exception(make_error_code(errc::feature_not_supported),
+                  "This operation is not supported by any existing backends.");
   return nullptr;
 }
 
@@ -142,7 +141,7 @@ static void initializeAdapters(std::vector<AdapterPtr> &Adapters,
 
   bool OwnLoaderConfig = false;
   // If we weren't provided with a custom config handle create our own.
-  if(!LoaderConfig) {
+  if (!LoaderConfig) {
     CHECK_UR_SUCCESS(loaderConfigCreate(&LoaderConfig))
     OwnLoaderConfig = true;
   }
@@ -168,12 +167,18 @@ static void initializeAdapters(std::vector<AdapterPtr> &Adapters,
   CHECK_UR_SUCCESS(loaderConfigSetCodeLocationCallback(
       LoaderConfig, codeLocationCallback, nullptr));
 
-  if (ProgramManager::getInstance().kernelUsesAsan()) {
-    if (loaderConfigEnableLayer(LoaderConfig, "UR_LAYER_ASAN")) {
-      loaderConfigRelease(LoaderConfig);
-      std::cerr << "Failed to enable ASAN layer\n";
-      return;
-    }
+  switch (ProgramManager::getInstance().kernelUsesSanitizer()) {
+  case SanitizerType::AddressSanitizer:
+    CHECK_UR_SUCCESS(loaderConfigEnableLayer(LoaderConfig, "UR_LAYER_ASAN"));
+    break;
+  case SanitizerType::MemorySanitizer:
+    CHECK_UR_SUCCESS(loaderConfigEnableLayer(LoaderConfig, "UR_LAYER_MSAN"));
+    break;
+  case SanitizerType::ThreadSanitizer:
+    CHECK_UR_SUCCESS(loaderConfigEnableLayer(LoaderConfig, "UR_LAYER_TSAN"));
+    break;
+  default:
+    break;
   }
 
   ur_device_init_flags_t device_flags = 0;
diff --git a/sycl/test-e2e/MemorySanitizer/check_buffer.cpp b/sycl/test-e2e/MemorySanitizer/check_buffer.cpp
new file mode 100644
index 0000000000000..dfce88803a96e
--- /dev/null
+++ b/sycl/test-e2e/MemorySanitizer/check_buffer.cpp
@@ -0,0 +1,32 @@
+// REQUIRES: linux, cpu || (gpu && level_zero)
+// RUN: %{build} %device_msan_flags -O1 -g -o %t2.out
+// RUN: %{run} not %t2.out 2>&1 | FileCheck %s
+// RUN: %{build} %device_msan_flags -O2 -g -o %t3.out
+// RUN: %{run} not %t3.out 2>&1 | FileCheck %s
+
+// XFAIL: gpu-intel-gen12 || gpu-intel-dg2
+// XFAIL-TRACKER: https://github.com/intel/llvm/issues/16184
+
+#include <sycl/detail/core.hpp>
+
+__attribute__((noinline)) long long foo(int data1, long long data2) {
+  return data1 + data2;
+}
+
+int main() {
+  sycl::queue q;
+
+  sycl::buffer<int, 1> buf1(sycl::range<1>(1));
+  sycl::buffer<long long, 1> buf2(sycl::range<1>(1));
+  q.submit([&](sycl::handler &h) {
+     auto array1 = buf1.get_access<sycl::access::mode::read_write>(h);
+     auto array2 = buf2.get_access<sycl::access::mode::read_write>(h);
+     h.single_task<class MyKernel>(
+         [=]() { array1[0] = foo(array1[0], array2[0]); });
+   }).wait();
+  // CHECK: use-of-uninitialized-value
+  // CHECK: kernel <{{.*MyKernel}}>
+  // CHECK: #0 {{.*}} {{.*check_buffer.cpp}}:[[@LINE-4]]
+
+  return 0;
+}
diff --git a/sycl/test-e2e/MemorySanitizer/check_call.cpp b/sycl/test-e2e/MemorySanitizer/check_call.cpp
new file mode 100644
index 0000000000000..fc1d41ac4c1e8
--- /dev/null
+++ b/sycl/test-e2e/MemorySanitizer/check_call.cpp
@@ -0,0 +1,32 @@
+// REQUIRES: linux, cpu || (gpu && level_zero)
+// RUN: %{build} %device_msan_flags -O1 -g -o %t2.out
+// RUN: %{run} not %t2.out 2>&1 | FileCheck %s
+// RUN: %{build} %device_msan_flags -O2 -g -o %t3.out
+// RUN: %{run} not %t3.out 2>&1 | FileCheck %s
+
+// XFAIL: gpu-intel-gen12 || gpu-intel-dg2
+// XFAIL-TRACKER: https://github.com/intel/llvm/issues/16184
+
+#include <sycl/detail/core.hpp>
+#include <sycl/usm.hpp>
+
+__attribute__((noinline)) long long foo(int data1, long long data2) {
+  return data1 + data2;
+}
+
+int main() {
+  sycl::queue Q;
+  auto *array = sycl::malloc_device<int>(2, Q);
+
+  Q.submit([&](sycl::handler &h) {
+    h.single_task<class MyKernel>(
+        [=]() { 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(array, Q);
+  return 0;
+}
diff --git a/sycl/test-e2e/MemorySanitizer/check_divide.cpp b/sycl/test-e2e/MemorySanitizer/check_divide.cpp
new file mode 100644
index 0000000000000..4769e1a3b9d44
--- /dev/null
+++ b/sycl/test-e2e/MemorySanitizer/check_divide.cpp
@@ -0,0 +1,37 @@
+// REQUIRES: linux, cpu || (gpu && level_zero)
+// RUN: %{build} %device_msan_flags -O1 -g -o %t2.out
+// RUN: %{run} not %t2.out 2>&1 | FileCheck %s
+// RUN: %{build} %device_msan_flags -O2 -g -o %t3.out
+// RUN: %{run} not %t3.out 2>&1 | FileCheck %s
+
+// XFAIL: gpu-intel-gen12 || gpu-intel-dg2
+// XFAIL-TRACKER: https://github.com/intel/llvm/issues/16184
+
+#include <sycl/detail/core.hpp>
+#include <sycl/usm.hpp>
+
+int main() {
+  sycl::queue Q;
+
+  auto *array = sycl::malloc_device<int>(3, Q);
+
+  Q.submit([&](sycl::handler &h) { h.single_task([=]() { array[1] = 1; }); });
+  Q.wait();
+
+  Q.submit([&](sycl::handler &h) {
+    h.single_task<class MyKernel1>([=]() { array[0] = array[0] / array[1]; });
+  });
+  Q.wait();
+  // CHECK-NOT: kernel <{{.*MyKernel1}}>
+
+  Q.submit([&](sycl::handler &h) {
+    h.single_task<class MyKernel2>([=]() { array[0] = array[0] / array[2]; });
+  });
+  Q.wait();
+  // CHECK: use-of-uninitialized-value
+  // CHECK: kernel <{{.*MyKernel2}}>
+  // CHECK: #0 {{.*}} {{.*check_divide.cpp}}:[[@LINE-5]]
+  sycl::free(array, Q);
+
+  return 0;
+}
diff --git a/sycl/test-e2e/MemorySanitizer/lit.local.cfg b/sycl/test-e2e/MemorySanitizer/lit.local.cfg
new file mode 100644
index 0000000000000..8817570e9ae06
--- /dev/null
+++ b/sycl/test-e2e/MemorySanitizer/lit.local.cfg
@@ -0,0 +1,6 @@
+config.substitutions.append(
+    ("%device_msan_flags", "-Xarch_device -fsanitize=memory")
+)
+config.substitutions.append(
+    ("%force_device_msan_rt", "env UR_ENABLE_LAYERS=UR_LAYER_MSAN")
+)