[DeviceMSAN] Add e2e tests for memcpy2d (#19335)

Author: AllanZyne

sycl/test-e2e/MemorySanitizer/check_usm.cpp

@@ -13,15 +13,16 @@
 #include <sycl/detail/core.hpp>
 #include <sycl/usm.hpp>
 
-__attribute__((noinline)) long long foo(int data1, long long data2) {
+__attribute__((noinline)) int check(int data1, int data2) {
   return data1 + data2;
 }
 
 void check_memset(sycl::queue &Q) {
   std::cout << "check_memset" << std::endl;
   auto *array = sycl::malloc_device<int>(2, Q);
   auto ev1 = Q.memset(array, 0, 2 * sizeof(int));
-  auto ev2 = Q.single_task(ev1, [=]() { array[0] = foo(array[0], array[1]); });
+  auto ev2 =
+      Q.single_task(ev1, [=]() { array[0] = check(array[0], array[1]); });
   Q.wait();
   sycl::free(array, Q);
   std::cout << "PASS" << std::endl;
@@ -30,25 +31,41 @@ void check_memset(sycl::queue &Q) {
 // CHECK-NOT: use-of-uninitialized-value
 // CHECK: PASS
 
+void check_fill(sycl::queue &Q) {
+  std::cout << "check_fill" << std::endl;
+  int *array = sycl::malloc_device<int>(2, Q);
+  uint32_t pattern = 0;
+  auto ev1 = Q.fill(array, pattern, 2);
+  auto ev2 =
+      Q.single_task(ev1, [=]() { array[0] = check(array[0], array[1]); });
+  Q.wait();
+  sycl::free(array, Q);
+  std::cout << "PASS" << std::endl;
+}
+// CHECK-LABEL: check_fill
+// CHECK-NOT: use-of-uninitialized-value
+// CHECK: PASS
+
 void check_memcpy(sycl::queue &Q) {
-  std::cout << "check_memcpy2" << std::endl;
+  std::cout << "check_memcpy" << std::endl;
   auto *source = sycl::malloc_device<int>(2, Q);
   auto *array = sycl::malloc_device<int>(2, Q);
-  // FIXME: We don't support shadow propagation on host/shared usm
   auto ev1 = Q.memcpy(array, source, 2 * sizeof(int));
-  auto ev2 = Q.single_task(ev1, [=]() { array[0] = foo(array[0], array[1]); });
+  auto ev2 =
+      Q.single_task(ev1, [=]() { array[0] = check(array[0], array[1]); });
   Q.wait();
   sycl::free(array, Q);
   sycl::free(source, Q);
   std::cout << "PASS" << std::endl;
 }
-// CHECK-LABEL: check_memcpy2
+// CHECK-LABEL: check_memcpy
 // CHECK: use-of-uninitialized-value
 // CHECK-NOT: PASS
 
 int main() {
   sycl::queue Q;
   check_memset(Q);
+  check_fill(Q);
   check_memcpy(Q);
   return 0;
 }

sycl/test-e2e/MemorySanitizer/check_usm2d.cpp

@@ -0,0 +1,57 @@
+// REQUIRES: linux, cpu || (gpu && level_zero)
+// RUN: %{build} %device_msan_flags -O0 -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
+
+// UNSUPPORTED: cpu
+// UNSUPPORTED-TRACKER: CMPLRLLVM-64618
+
+// XFAIL: spirv-backend && gpu
+// XFAIL-TRACKER: CMPLRLLVM-64705
+
+#include <sycl/detail/core.hpp>
+#include <sycl/ext/oneapi/memcpy2d.hpp>
+#include <sycl/usm.hpp>
+
+constexpr size_t Pitch = 4;
+constexpr size_t Width = 2;
+constexpr size_t Height = 2;
+constexpr size_t Size = Pitch * Height;
+
+__attribute__((noinline)) int check(int data1, int data2) {
+  return data1 + data2;
+}
+
+void check_memcpy2d(sycl::queue &Q) {
+  std::cout << "check_memcpy2d" << std::endl;
+
+  auto *source = sycl::malloc_device<int>(Size, Q);
+  Q.memset(source, 0, Size * sizeof(int)).wait();
+
+  auto *dest = sycl::malloc_device<int>(Size, Q);
+  Q.ext_oneapi_memcpy2d(dest, Pitch * sizeof(int), source, Pitch * sizeof(int),
+                        Width * sizeof(int), Height)
+      .wait();
+
+  Q.single_task<class Test1>([=]() {
+     dest[0] = check(dest[0], dest[1]);
+     dest[0] = check(dest[4], dest[5]);
+   }).wait();
+  // CHECK-NOT: check_usm2d.cpp:[[@LINE-3]]
+
+  Q.single_task<class Test2>([=]() {
+     dest[0] = check(dest[2], dest[3]);
+   }).wait();
+  // CHECK: use-of-uninitialized-value
+  // CHECK: check_usm2d.cpp:[[@LINE-3]]
+
+  sycl::free(dest, Q);
+  sycl::free(source, Q);
+}
+
+int main() {
+  sycl::queue Q;
+  check_memcpy2d(Q);
+  return 0;
+}

unified-runtime/source/loader/layers/sanitizer/msan/msan_ddi.cpp

@@ -61,8 +61,7 @@ ur_result_t urEnqueueUSMFill2DFallback(ur_queue_handle_t hQueue, void *pMem,
   ur_result_t Result = getContext()->urDdiTable.Enqueue.pfnUSMFill2D(
       hQueue, pMem, pitch, patternSize, pPattern, width, height,
       numEventsInWaitList, phEventWaitList, phEvent);
-  if (Result == UR_RESULT_SUCCESS ||
-      Result != UR_RESULT_ERROR_UNSUPPORTED_FEATURE) {
+  if (Result != UR_RESULT_ERROR_UNSUPPORTED_FEATURE) {
     return Result;
   }