[DevMSAN] Propagate shadow memory in buffer related APIs

sycl/cmake/modules/UnifiedRuntimeTag.cmake

@@ -1,9 +1,7 @@
-# commit ad88f0a1e3a22807866af4f6dad53e8986733abb
-# Merge: 68aed2d5 599a28e1
-# Author: Ross Brunton <[email protected]>
-# Date:   Tue Jan 14 10:28:07 2025 +0000
-
-#    Merge pull request #2527 from RossBrunton/ross/wrapper
-#    
-#    Wrap urEventSetCallback when ran through loader
-set(UNIFIED_RUNTIME_TAG ad88f0a1e3a22807866af4f6dad53e8986733abb)
+# commit afbb289aa8d4f3b27b1536ba33ca618b0aba65c7
+# Merge: ef70004f d7c33f88
+# Author: Kenneth Benzie (Benie) <[email protected]>
+# Date:   Wed Jan 15 11:54:25 2025 +0000
+#     Merge pull request #2520 from zhaomaosu/fix-buffer-shadow
+#     [DevMSAN] Propagate shadow memory in buffer related APIs
+set(UNIFIED_RUNTIME_TAG afbb289aa8d4f3b27b1536ba33ca618b0aba65c7)

sycl/test-e2e/AddressSanitizer/lit.local.cfg

@@ -8,5 +8,8 @@ config.substitutions.append(
     ("%force_device_asan_rt", "env UR_ENABLE_LAYERS=UR_LAYER_ASAN")
 )
 
+if "-fsanitize=memory" in config.cxx_flags:
+    config.unsupported=True
+
 # https://github.com/intel/llvm/issues/15953
 config.unsupported_features += ['gpu-intel-gen12']

sycl/test-e2e/MemorySanitizer/check_buffer_memset_memcpy.cpp

@@ -0,0 +1,64 @@
+// REQUIRES: linux, cpu || (gpu && level_zero)
+// RUN: %{build} %device_msan_flags -O0 -g -o %t1.out
+// RUN: %{run} %t1.out 2>&1 | FileCheck %s
+// RUN: %{build} %device_msan_flags -O2 -g -o %t2.out
+// RUN: %{run} %t2.out 2>&1 | FileCheck %s
+
+#include <sycl/detail/core.hpp>
+
+__attribute__((noinline)) int foo(int data1, int data2) {
+  return data1 + data2;
+}
+
+void check_memset(sycl::queue &q) {
+  std::cout << "check_memset" << std::endl;
+  sycl::buffer<int, 1> buf(sycl::range<1>(2));
+  const int Pattern = 0;
+
+  q.submit([&](sycl::handler &h) {
+     auto array = buf.get_access<sycl::access::mode::read_write>(h);
+     h.fill(array, Pattern);
+   }).wait();
+
+  q.submit([&](sycl::handler &h) {
+     auto array = buf.get_access<sycl::access::mode::read_write>(h);
+     h.single_task<class MyKernel1>(
+         [=]() { array[0] = foo(array[0], array[1]); });
+   }).wait();
+  std::cout << "PASS" << std::endl;
+  // CHECK-LABEL: check_memset
+  // CHECK-NOT: use-of-uninitialized-value
+  // CHECK: PASS
+}
+
+void check_memcpy(sycl::queue &q) {
+  std::cout << "check_memcpy" << std::endl;
+  int host[2] = {1, 2};
+  sycl::buffer<int, 1> buf1(sycl::range<1>(2));
+  sycl::buffer<int, 1> buf2(host, sycl::range<1>(2));
+
+  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.copy(array2, array1);
+   }).wait();
+
+  q.submit([&](sycl::handler &h) {
+     auto array = buf1.get_access<sycl::access::mode::read_write>(h);
+     h.single_task<class MyKernel2>(
+         [=]() { array[0] = foo(array[0], array[1]); });
+   }).wait();
+  std::cout << "PASS" << std::endl;
+  // CHECK-LABEL: check_memcpy
+  // CHECK-NOT: use-of-uninitialized-value
+  // CHECK: PASS
+}
+
+int main() {
+  sycl::queue q;
+
+  check_memset(q);
+  check_memcpy(q);
+
+  return 0;
+}

sycl/test-e2e/MemorySanitizer/lit.local.cfg

@@ -12,3 +12,6 @@ config.substitutions.append(
 config.substitutions.append(
     ("%force_device_msan_rt", "env UR_ENABLE_LAYERS=UR_LAYER_MSAN")
 )
+
+if "-fsanitize=address" in config.cxx_flags:
+    config.unsupported=True