From cbecda28702d39360138f90304cf03a908819985 Mon Sep 17 00:00:00 2001
From: Andrei Elovikov <andrei.elovikov@intel.com>
Date: Fri, 30 Aug 2024 09:24:24 -0700
Subject: [PATCH] [SYCL][E2E] P2P must use Device USM, not Shared

---
 sycl/test-e2e/USM/P2P/p2p_atomics.cpp | 9 ++++++---
 1 file changed, 6 insertions(+), 3 deletions(-)

diff --git a/sycl/test-e2e/USM/P2P/p2p_atomics.cpp b/sycl/test-e2e/USM/P2P/p2p_atomics.cpp
index 3975573394c0d..b81405f316e7c 100644
--- a/sycl/test-e2e/USM/P2P/p2p_atomics.cpp
+++ b/sycl/test-e2e/USM/P2P/p2p_atomics.cpp
@@ -50,15 +50,15 @@ int main() {
     h_sum += value;
   }
 
-  int *d_sum = malloc_shared<int>(1, Queues[0]);
+  int *d_sum = malloc_device<int>(1, Queues[0]);
   int *d_in = malloc_device<int>(N, Queues[0]);
 
+  Queues[0].single_task([=]() { *d_sum = 0; });
   Queues[0].memcpy(d_in, &input[0], N * sizeof(int));
   Queues[0].wait();
 
   range global_range{N};
 
-  *d_sum = 0.;
   Queues[1].submit([&](handler &h) {
     h.parallel_for<class peer_atomic>(global_range, [=](id<1> i) {
       sycl::atomic_ref<int, sycl::memory_order::relaxed,
@@ -68,7 +68,10 @@ int main() {
   });
   Queues[1].wait();
 
-  assert(*d_sum == h_sum);
+  int result = 0;
+  Queues[0].memcpy(&result, d_sum, sizeof(int)).wait();
+
+  assert(result == h_sum);
 
   free(d_sum, Queues[0]);
   free(d_in, Queues[0]);