From 2cfe8e6b15b03322704b6473bd44089f1a3905c3 Mon Sep 17 00:00:00 2001
From: Nick Sarnie <nick.sarnie@intel.com>
Date: Sat, 11 Jan 2025 07:16:04 +0900
Subject: [PATCH] [SYCL][ESIMD][E2E] Split some atomic update tests (#16580)

These are really slow so split them into two parts, we get like a 40
second speedup.

Signed-off-by: Sarnie, Nick <nick.sarnie@intel.com>
---
 .../Inputs/atomic_update.hpp                  | 43 ++++++++++++-------
 .../Inputs/atomic_update_slm.hpp              | 40 +++++++++++------
 .../atomic_update_acc_dg2_pvc.cpp             |  2 +-
 .../atomic_update_acc_dg2_pvc_2.cpp           | 26 +++++++++++
 .../atomic_update_acc_dg2_pvc_64_2.cpp        | 17 ++++++++
 .../atomic_update_acc_dg2_pvc_cmpxchg_2.cpp   | 17 ++++++++
 .../atomic_update_acc_dg2_pvc_stateless_2.cpp | 14 ++++++
 .../atomic_update_slm_acc_pvc.cpp             |  2 +-
 .../atomic_update_slm_acc_pvc_2.cpp           | 26 +++++++++++
 .../atomic_update_slm_acc_pvc_cmpxchg_2.cpp   | 16 +++++++
 .../atomic_update_slm_pvc.cpp                 |  2 +-
 .../atomic_update_slm_pvc_2.cpp               | 26 +++++++++++
 .../atomic_update_slm_pvc_cmpxchg_2.cpp       | 16 +++++++
 .../atomic_update_usm_dg2_pvc.cpp             |  2 +-
 .../atomic_update_usm_dg2_pvc_2.cpp           | 30 +++++++++++++
 .../atomic_update_usm_dg2_pvc_64_2.cpp        | 16 +++++++
 .../atomic_update_usm_dg2_pvc_cmpxchg_2.cpp   | 17 ++++++++
 17 files changed, 278 insertions(+), 34 deletions(-)
 create mode 100644 sycl/test-e2e/ESIMD/unified_memory_api/atomic_update_acc_dg2_pvc_2.cpp
 create mode 100644 sycl/test-e2e/ESIMD/unified_memory_api/atomic_update_acc_dg2_pvc_64_2.cpp
 create mode 100644 sycl/test-e2e/ESIMD/unified_memory_api/atomic_update_acc_dg2_pvc_cmpxchg_2.cpp
 create mode 100644 sycl/test-e2e/ESIMD/unified_memory_api/atomic_update_acc_dg2_pvc_stateless_2.cpp
 create mode 100644 sycl/test-e2e/ESIMD/unified_memory_api/atomic_update_slm_acc_pvc_2.cpp
 create mode 100644 sycl/test-e2e/ESIMD/unified_memory_api/atomic_update_slm_acc_pvc_cmpxchg_2.cpp
 create mode 100644 sycl/test-e2e/ESIMD/unified_memory_api/atomic_update_slm_pvc_2.cpp
 create mode 100644 sycl/test-e2e/ESIMD/unified_memory_api/atomic_update_slm_pvc_cmpxchg_2.cpp
 create mode 100644 sycl/test-e2e/ESIMD/unified_memory_api/atomic_update_usm_dg2_pvc_2.cpp
 create mode 100644 sycl/test-e2e/ESIMD/unified_memory_api/atomic_update_usm_dg2_pvc_64_2.cpp
 create mode 100644 sycl/test-e2e/ESIMD/unified_memory_api/atomic_update_usm_dg2_pvc_cmpxchg_2.cpp

diff --git a/sycl/test-e2e/ESIMD/unified_memory_api/Inputs/atomic_update.hpp b/sycl/test-e2e/ESIMD/unified_memory_api/Inputs/atomic_update.hpp
index cf2bccc166ac6..e31f5e923dcf6 100644
--- a/sycl/test-e2e/ESIMD/unified_memory_api/Inputs/atomic_update.hpp
+++ b/sycl/test-e2e/ESIMD/unified_memory_api/Inputs/atomic_update.hpp
@@ -118,7 +118,7 @@ bool verify(T *arr, const Config &cfg, size_t size) {
 
 template <class T, int N, template <class, int> class ImplF, bool UseMask,
           bool UseProperties>
-bool test_usm(queue q, const Config &cfg) {
+bool test_usm(queue &q, const Config &cfg) {
   constexpr auto op = ImplF<T, N>::atomic_op;
   using CurAtomicOpT = decltype(op);
   constexpr int n_args = ImplF<T, N>::n_args;
@@ -245,7 +245,7 @@ bool test_usm(queue q, const Config &cfg) {
 
 template <class T, int N, template <class, int> class ImplF, bool UseMask,
           bool UseProperties>
-bool test_acc(queue q, const Config &cfg) {
+bool test_acc(queue &q, const Config &cfg) {
   constexpr auto op = ImplF<T, N>::atomic_op;
   using CurAtomicOpT = decltype(op);
   constexpr int n_args = ImplF<T, N>::n_args;
@@ -615,7 +615,7 @@ struct ImplFcmpwr : ImplCmpxchgBase<T, N, atomic_op, atomic_op::fcmpxchg> {};
 
 template <bool UseAcc, class T, int N, template <class, int> class ImplF,
           bool UseMask, bool UseLSCFeatures>
-auto run_test(queue q, const Config &cfg) {
+auto run_test(queue &q, const Config &cfg) {
   if constexpr (UseAcc)
     return test_acc<T, N, ImplF, UseMask, UseLSCFeatures>(q, cfg);
   else
@@ -624,7 +624,7 @@ auto run_test(queue q, const Config &cfg) {
 
 template <int N, template <class, int> class Op, bool UseMask,
           bool UseLSCFeatures, bool UseAcc, int SignMask = (Signed | Unsigned)>
-bool test_int_types(queue q, const Config &cfg) {
+bool test_int_types(queue &q, const Config &cfg) {
   bool passed = true;
   if constexpr (SignMask & Signed) {
     // Supported by LSC atomic:
@@ -662,7 +662,7 @@ bool test_int_types(queue q, const Config &cfg) {
 
 template <int N, template <class, int> class Op, bool UseMask,
           bool UseLSCFeatures, bool UseAcc>
-bool test_fp_types(queue q, const Config &cfg) {
+bool test_fp_types(queue &q, const Config &cfg) {
   bool passed = true;
   // TODO: Enable FADD/FSUB on DG2/PVC when the error in GPU driver is resolved.
   if constexpr (UseLSCFeatures &&
@@ -685,7 +685,7 @@ bool test_fp_types(queue q, const Config &cfg) {
 
 template <template <class, int> class Op, bool UseMask, bool UseLSCFeatures,
           bool UseAcc, int SignMask = (Signed | Unsigned)>
-bool test_int_types_and_sizes(queue q, const Config &cfg) {
+bool test_int_types_and_sizes(queue &q, const Config &cfg) {
   bool passed = true;
   passed &=
       test_int_types<1, Op, UseMask, UseLSCFeatures, UseAcc, SignMask>(q, cfg);
@@ -715,7 +715,7 @@ bool test_int_types_and_sizes(queue q, const Config &cfg) {
 
 template <template <class, int> class Op, bool UseMask, bool UseLSCFeatures,
           bool UseAcc>
-bool test_fp_types_and_sizes(queue q, const Config &cfg) {
+bool test_fp_types_and_sizes(queue &q, const Config &cfg) {
   bool passed = true;
   passed &= test_fp_types<1, Op, UseMask, UseLSCFeatures, UseAcc>(q, cfg);
   passed &= test_fp_types<2, Op, UseMask, UseLSCFeatures, UseAcc>(q, cfg);
@@ -733,7 +733,7 @@ bool test_fp_types_and_sizes(queue q, const Config &cfg) {
 }
 
 template <bool UseMask, bool UseLSCFeatures, bool UseAcc>
-bool test_with_mask(queue q) {
+bool test_with_mask(queue &q) {
   bool passed = true;
 
   Config cfg{
@@ -818,26 +818,37 @@ bool test_with_mask(queue q) {
   return passed;
 }
 
-template <bool UseLSCFeatures> bool test_main(queue q) {
+template <bool UseLSCFeatures, bool PartOne> bool test_main(queue &q) {
   bool passed = true;
 
   constexpr const bool UseMask = true;
   constexpr const bool UseAcc = true;
-
-  passed &= test_with_mask<UseMask, UseLSCFeatures, !UseAcc>(q);
-  passed &= test_with_mask<!UseMask, UseLSCFeatures, !UseAcc>(q);
+  if constexpr (PartOne)
+    passed &= test_with_mask<UseMask, UseLSCFeatures, !UseAcc>(q);
+  else
+    passed &= test_with_mask<!UseMask, UseLSCFeatures, !UseAcc>(q);
 
   return passed;
 }
 
-template <bool UseLSCFeatures> bool test_main_acc(queue q) {
+template <bool UseLSCFeature> bool test_main(queue &q) {
+  return test_main<UseLSCFeature, true>(q) &&
+         test_main<UseLSCFeature, false>(q);
+}
+
+template <bool UseLSCFeatures, bool PartOne> bool test_main_acc(queue &q) {
   bool passed = true;
 
   constexpr const bool UseMask = true;
   constexpr const bool UseAcc = true;
-
-  passed &= test_with_mask<UseMask, UseLSCFeatures, UseAcc>(q);
-  passed &= test_with_mask<!UseMask, UseLSCFeatures, UseAcc>(q);
+  if constexpr (PartOne)
+    passed &= test_with_mask<UseMask, UseLSCFeatures, UseAcc>(q);
+  else
+    passed &= test_with_mask<!UseMask, UseLSCFeatures, UseAcc>(q);
 
   return passed;
 }
+template <bool UseLSCFeatures> bool test_main_acc(queue &q) {
+  return test_main_acc<UseLSCFeatures, true>(q) &&
+         test_main_acc<UseLSCFeatures, false>(q);
+}
diff --git a/sycl/test-e2e/ESIMD/unified_memory_api/Inputs/atomic_update_slm.hpp b/sycl/test-e2e/ESIMD/unified_memory_api/Inputs/atomic_update_slm.hpp
index 16b95fc555c7a..ee094381a8d6c 100644
--- a/sycl/test-e2e/ESIMD/unified_memory_api/Inputs/atomic_update_slm.hpp
+++ b/sycl/test-e2e/ESIMD/unified_memory_api/Inputs/atomic_update_slm.hpp
@@ -78,7 +78,7 @@ template <int N> inline bool any(simd_mask<N> m, simd_mask<N> ignore_mask) {
 // The main test function
 
 template <class T, int N, template <class, int> class ImplF, bool UseMask>
-bool test_slm(queue q) {
+bool test_slm(queue &q) {
   constexpr auto op = ImplF<T, N>::atomic_op;
   using CurAtomicOpT = decltype(op);
   constexpr int n_args = ImplF<T, N>::n_args;
@@ -202,7 +202,7 @@ bool test_slm(queue q) {
 }
 
 template <class T, int N, template <class, int> class ImplF, bool UseMask>
-bool test_slm_acc(queue q) {
+bool test_slm_acc(queue &q) {
   constexpr auto op = ImplF<T, N>::atomic_op;
   using CurAtomicOpT = decltype(op);
   constexpr int n_args = ImplF<T, N>::n_args;
@@ -542,7 +542,7 @@ struct ImplLSCFcmpwr : ImplCmpxchgBase<T, N, atomic_op, atomic_op::fcmpxchg> {};
 
 template <bool UseAcc, class T, int N, template <class, int> class ImplF,
           bool UseMask>
-auto run_test(queue q) {
+auto run_test(queue &q) {
   if constexpr (UseAcc) {
     return test_slm_acc<T, N, ImplF, UseMask>(q);
   } else {
@@ -553,7 +553,7 @@ auto run_test(queue q) {
 template <int N, template <class, int> class Op, bool UseMask,
           TestFeatures Features, bool UseAcc,
           int SignMask = (Signed | Unsigned)>
-bool test_int_types(queue q) {
+bool test_int_types(queue &q) {
   bool passed = true;
   if constexpr (SignMask & Signed) {
     if constexpr (Features == TestFeatures::DG2 ||
@@ -585,7 +585,7 @@ bool test_int_types(queue q) {
 
 template <int N, template <class, int> class Op, bool UseMask,
           TestFeatures Features, bool UseAcc>
-bool test_fp_types(queue q) {
+bool test_fp_types(queue &q) {
   bool passed = true;
 
   // TODO: Enable 'half' FADD/FSUB on DG2 when the error in GPU driver is fixed.
@@ -612,7 +612,7 @@ bool test_fp_types(queue q) {
 
 template <template <class, int> class Op, bool UseMask, TestFeatures Features,
           bool UseAcc, int SignMask = (Signed | Unsigned)>
-bool test_int_types_and_sizes(queue q) {
+bool test_int_types_and_sizes(queue &q) {
   bool passed = true;
   passed &= test_int_types<1, Op, UseMask, Features, UseAcc, SignMask>(q);
   passed &= test_int_types<2, Op, UseMask, Features, UseAcc, SignMask>(q);
@@ -636,7 +636,7 @@ bool test_int_types_and_sizes(queue q) {
 
 template <template <class, int> class Op, bool UseMask, TestFeatures Features,
           bool UseAcc>
-bool test_fp_types_and_sizes(queue q) {
+bool test_fp_types_and_sizes(queue &q) {
   bool passed = true;
   passed &= test_fp_types<1, Op, UseMask, Features, UseAcc>(q);
   passed &= test_fp_types<2, Op, UseMask, Features, UseAcc>(q);
@@ -656,7 +656,7 @@ bool test_fp_types_and_sizes(queue q) {
 }
 
 template <bool UseMask, TestFeatures Features, bool UseAcc>
-int test_with_mask(queue q) {
+int test_with_mask(queue &q) {
   bool passed = true;
 #ifndef CMPXCHG_TEST
   passed &= test_int_types_and_sizes<ImplInc, UseMask, Features, UseAcc>(q);
@@ -708,26 +708,38 @@ int test_with_mask(queue q) {
   return passed;
 }
 
-template <TestFeatures Features> bool test_main(queue q) {
+template <TestFeatures Features, bool PartOne> bool test_main(queue &q) {
   bool passed = true;
 
   constexpr const bool UseMask = true;
   constexpr const bool UseAcc = true;
 
-  passed &= test_with_mask<UseMask, Features, !UseAcc>(q);
-  passed &= test_with_mask<!UseMask, Features, !UseAcc>(q);
+  if constexpr (PartOne)
+    passed &= test_with_mask<UseMask, Features, !UseAcc>(q);
+  else
+    passed &= test_with_mask<!UseMask, Features, !UseAcc>(q);
 
   return passed;
 }
 
-template <TestFeatures Features> bool test_main_acc(queue q) {
+template <TestFeatures Features> bool test_main(queue &q) {
+  return test_main<Features, true>(q) && test_main<Features, false>(q);
+}
+
+template <TestFeatures Features, bool PartOne> bool test_main_acc(queue &q) {
   bool passed = true;
 
   constexpr const bool UseMask = true;
   constexpr const bool UseAcc = true;
 
-  passed &= test_with_mask<UseMask, Features, UseAcc>(q);
-  passed &= test_with_mask<!UseMask, Features, UseAcc>(q);
+  if constexpr (PartOne)
+    passed &= test_with_mask<UseMask, Features, UseAcc>(q);
+  else
+    passed &= test_with_mask<!UseMask, Features, UseAcc>(q);
 
   return passed;
 }
+
+template <TestFeatures Features> bool test_main_acc(queue &q) {
+  return test_main_acc<Features, true>(q) && test_main_acc<Features, false>(q);
+}
diff --git a/sycl/test-e2e/ESIMD/unified_memory_api/atomic_update_acc_dg2_pvc.cpp b/sycl/test-e2e/ESIMD/unified_memory_api/atomic_update_acc_dg2_pvc.cpp
index ff99c8f687bef..a115b6f75d752 100644
--- a/sycl/test-e2e/ESIMD/unified_memory_api/atomic_update_acc_dg2_pvc.cpp
+++ b/sycl/test-e2e/ESIMD/unified_memory_api/atomic_update_acc_dg2_pvc.cpp
@@ -19,7 +19,7 @@ int main(void) {
   esimd_test::printTestLabel(q);
 
   constexpr bool TestCacheHintProperties = true;
-  bool passed = test_main_acc<TestCacheHintProperties>(q);
+  bool passed = test_main_acc<TestCacheHintProperties, true>(q);
 
   std::cout << (passed ? "Passed\n" : "FAILED\n");
   return passed ? 0 : 1;
diff --git a/sycl/test-e2e/ESIMD/unified_memory_api/atomic_update_acc_dg2_pvc_2.cpp b/sycl/test-e2e/ESIMD/unified_memory_api/atomic_update_acc_dg2_pvc_2.cpp
new file mode 100644
index 0000000000000..4fa38dd16b8ae
--- /dev/null
+++ b/sycl/test-e2e/ESIMD/unified_memory_api/atomic_update_acc_dg2_pvc_2.cpp
@@ -0,0 +1,26 @@
+//==---- atomic_update_acc_dg2_pvc_2.cpp - DPC++ ESIMD on-device test ---==//
+//
+// 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
+//
+//===-------------------------------------------------------------------===//
+
+// REQUIRES: arch-intel_gpu_pvc || gpu-intel-dg2
+
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+
+#include "Inputs/atomic_update.hpp"
+
+int main(void) {
+  queue q(esimd_test::ESIMDSelector, esimd_test::createExceptionHandler());
+
+  esimd_test::printTestLabel(q);
+
+  constexpr bool TestCacheHintProperties = true;
+  bool passed = test_main_acc<TestCacheHintProperties, false>(q);
+
+  std::cout << (passed ? "Passed\n" : "FAILED\n");
+  return passed ? 0 : 1;
+}
diff --git a/sycl/test-e2e/ESIMD/unified_memory_api/atomic_update_acc_dg2_pvc_64_2.cpp b/sycl/test-e2e/ESIMD/unified_memory_api/atomic_update_acc_dg2_pvc_64_2.cpp
new file mode 100644
index 0000000000000..85e54d6c6b541
--- /dev/null
+++ b/sycl/test-e2e/ESIMD/unified_memory_api/atomic_update_acc_dg2_pvc_64_2.cpp
@@ -0,0 +1,17 @@
+//==-- atomic_update_acc_dg2_pvc_64_2.cpp - DPC++ ESIMD on-device test----==//
+//
+// 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
+//
+//===--------------------------------------------------------------===//
+
+// REQUIRES: arch-intel_gpu_pvc || gpu-intel-dg2
+
+// RUN: %{build} -fsycl-esimd-force-stateless-mem -o %t.out
+// RUN: %{run} %t.out
+
+// 64-bit offset is supported for accessors only in stateless mode
+#define USE_64_BIT_OFFSET
+
+#include "atomic_update_acc_dg2_pvc_2.cpp"
diff --git a/sycl/test-e2e/ESIMD/unified_memory_api/atomic_update_acc_dg2_pvc_cmpxchg_2.cpp b/sycl/test-e2e/ESIMD/unified_memory_api/atomic_update_acc_dg2_pvc_cmpxchg_2.cpp
new file mode 100644
index 0000000000000..59ee6720cd27b
--- /dev/null
+++ b/sycl/test-e2e/ESIMD/unified_memory_api/atomic_update_acc_dg2_pvc_cmpxchg_2.cpp
@@ -0,0 +1,17 @@
+//=- atomic_update_acc_dg2_pvc_cmpxchg_2.cpp- DPC++ ESIMD on-device test -=//
+//
+// 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
+//
+//===---------------------------------------------------------------------===//
+
+// REQUIRES: arch-intel_gpu_pvc || gpu-intel-dg2
+// REQUIRES-INTEL-DRIVER: lin: 29803
+
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+
+#define CMPXCHG_TEST
+
+#include "atomic_update_acc_dg2_pvc_2.cpp"
diff --git a/sycl/test-e2e/ESIMD/unified_memory_api/atomic_update_acc_dg2_pvc_stateless_2.cpp b/sycl/test-e2e/ESIMD/unified_memory_api/atomic_update_acc_dg2_pvc_stateless_2.cpp
new file mode 100644
index 0000000000000..8402273b328e5
--- /dev/null
+++ b/sycl/test-e2e/ESIMD/unified_memory_api/atomic_update_acc_dg2_pvc_stateless_2.cpp
@@ -0,0 +1,14 @@
+//=- atomic_update_acc_dg2_pvc_stateless_2.cpp - DPC++ ESIMD on-device test -=//
+//
+// 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
+//
+//=-----------------------------------------------------------------------=//
+
+// REQUIRES: arch-intel_gpu_pvc || gpu-intel-dg2
+
+// RUN: %{build} -fsycl-esimd-force-stateless-mem -o %t.out
+// RUN: %{run} %t.out
+
+#include "atomic_update_acc_dg2_pvc_2.cpp"
diff --git a/sycl/test-e2e/ESIMD/unified_memory_api/atomic_update_slm_acc_pvc.cpp b/sycl/test-e2e/ESIMD/unified_memory_api/atomic_update_slm_acc_pvc.cpp
index ab6fa4ca97373..b70531a36b0ca 100644
--- a/sycl/test-e2e/ESIMD/unified_memory_api/atomic_update_slm_acc_pvc.cpp
+++ b/sycl/test-e2e/ESIMD/unified_memory_api/atomic_update_slm_acc_pvc.cpp
@@ -19,7 +19,7 @@ int main(void) {
   esimd_test::printTestLabel(q);
 
   constexpr auto Features = TestFeatures::PVC;
-  bool passed = test_main_acc<Features>(q);
+  bool passed = test_main_acc<Features, true>(q);
 
   std::cout << (passed ? "Passed\n" : "FAILED\n");
   return passed ? 0 : 1;
diff --git a/sycl/test-e2e/ESIMD/unified_memory_api/atomic_update_slm_acc_pvc_2.cpp b/sycl/test-e2e/ESIMD/unified_memory_api/atomic_update_slm_acc_pvc_2.cpp
new file mode 100644
index 0000000000000..9e6517964adeb
--- /dev/null
+++ b/sycl/test-e2e/ESIMD/unified_memory_api/atomic_update_slm_acc_pvc_2.cpp
@@ -0,0 +1,26 @@
+//==- atomic_update_slm_acc_pvc_2.cpp - DPC++ ESIMD on-device test -==//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+
+// REQUIRES: arch-intel_gpu_pvc
+
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+
+#include "Inputs/atomic_update_slm.hpp"
+
+int main(void) {
+  queue q(esimd_test::ESIMDSelector, esimd_test::createExceptionHandler());
+
+  esimd_test::printTestLabel(q);
+
+  constexpr auto Features = TestFeatures::PVC;
+  bool passed = test_main_acc<Features, false>(q);
+
+  std::cout << (passed ? "Passed\n" : "FAILED\n");
+  return passed ? 0 : 1;
+}
diff --git a/sycl/test-e2e/ESIMD/unified_memory_api/atomic_update_slm_acc_pvc_cmpxchg_2.cpp b/sycl/test-e2e/ESIMD/unified_memory_api/atomic_update_slm_acc_pvc_cmpxchg_2.cpp
new file mode 100644
index 0000000000000..387839d8acdc2
--- /dev/null
+++ b/sycl/test-e2e/ESIMD/unified_memory_api/atomic_update_slm_acc_pvc_cmpxchg_2.cpp
@@ -0,0 +1,16 @@
+//=- atomic_update_slm_acc_pvc_cmpxchg_2.cpp -- DPC++ ESIMD on-device test -=//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+
+// REQUIRES: arch-intel_gpu_pvc
+
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+
+#define CMPXCHG_TEST
+
+#include "atomic_update_slm_acc_pvc_2.cpp"
diff --git a/sycl/test-e2e/ESIMD/unified_memory_api/atomic_update_slm_pvc.cpp b/sycl/test-e2e/ESIMD/unified_memory_api/atomic_update_slm_pvc.cpp
index eea3a923a90a8..562e01868a6b0 100644
--- a/sycl/test-e2e/ESIMD/unified_memory_api/atomic_update_slm_pvc.cpp
+++ b/sycl/test-e2e/ESIMD/unified_memory_api/atomic_update_slm_pvc.cpp
@@ -19,7 +19,7 @@ int main(void) {
   esimd_test::printTestLabel(q);
 
   constexpr auto Features = TestFeatures::PVC;
-  bool passed = test_main<Features>(q);
+  bool passed = test_main<Features, true>(q);
 
   std::cout << (passed ? "Passed\n" : "FAILED\n");
   return passed ? 0 : 1;
diff --git a/sycl/test-e2e/ESIMD/unified_memory_api/atomic_update_slm_pvc_2.cpp b/sycl/test-e2e/ESIMD/unified_memory_api/atomic_update_slm_pvc_2.cpp
new file mode 100644
index 0000000000000..5efe6966fe1b2
--- /dev/null
+++ b/sycl/test-e2e/ESIMD/unified_memory_api/atomic_update_slm_pvc_2.cpp
@@ -0,0 +1,26 @@
+//==------- atomic_update_slm_pvc_2.cpp - DPC++ ESIMD on-device test -------==//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+
+// REQUIRES: arch-intel_gpu_pvc
+
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+
+#include "Inputs/atomic_update_slm.hpp"
+
+int main(void) {
+  queue q(esimd_test::ESIMDSelector, esimd_test::createExceptionHandler());
+
+  esimd_test::printTestLabel(q);
+
+  constexpr auto Features = TestFeatures::PVC;
+  bool passed = test_main<Features, false>(q);
+
+  std::cout << (passed ? "Passed\n" : "FAILED\n");
+  return passed ? 0 : 1;
+}
diff --git a/sycl/test-e2e/ESIMD/unified_memory_api/atomic_update_slm_pvc_cmpxchg_2.cpp b/sycl/test-e2e/ESIMD/unified_memory_api/atomic_update_slm_pvc_cmpxchg_2.cpp
new file mode 100644
index 0000000000000..088ab474c2e04
--- /dev/null
+++ b/sycl/test-e2e/ESIMD/unified_memory_api/atomic_update_slm_pvc_cmpxchg_2.cpp
@@ -0,0 +1,16 @@
+//==---- atomic_update_slm_pvc_cmpxchg_2.cpp -- DPC++ ESIMD on-device test -==//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+
+// REQUIRES: arch-intel_gpu_pvc
+
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+
+#define CMPXCHG_TEST
+
+#include "atomic_update_slm_pvc_2.cpp"
diff --git a/sycl/test-e2e/ESIMD/unified_memory_api/atomic_update_usm_dg2_pvc.cpp b/sycl/test-e2e/ESIMD/unified_memory_api/atomic_update_usm_dg2_pvc.cpp
index 281727c610184..85ce505045a27 100644
--- a/sycl/test-e2e/ESIMD/unified_memory_api/atomic_update_usm_dg2_pvc.cpp
+++ b/sycl/test-e2e/ESIMD/unified_memory_api/atomic_update_usm_dg2_pvc.cpp
@@ -23,7 +23,7 @@ int main(void) {
   bool passed = true;
 
   constexpr bool TestCacheHintProperties = true;
-  passed &= test_main<TestCacheHintProperties>(q);
+  passed &= test_main<TestCacheHintProperties, true>(q);
 
   std::cout << (passed ? "Passed\n" : "FAILED\n");
   return passed ? 0 : 1;
diff --git a/sycl/test-e2e/ESIMD/unified_memory_api/atomic_update_usm_dg2_pvc_2.cpp b/sycl/test-e2e/ESIMD/unified_memory_api/atomic_update_usm_dg2_pvc_2.cpp
new file mode 100644
index 0000000000000..f6527550268f2
--- /dev/null
+++ b/sycl/test-e2e/ESIMD/unified_memory_api/atomic_update_usm_dg2_pvc_2.cpp
@@ -0,0 +1,30 @@
+//==-- atomic_update_usm_dg2_pvc_2.cpp - DPC++ ESIMD on-device test----==//
+//
+// 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
+//
+//===---------------------------------------------------------------===//
+
+// REQUIRES: arch-intel_gpu_pvc || gpu-intel-dg2
+
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+
+#include "Inputs/atomic_update.hpp"
+
+int main(void) {
+  queue q(esimd_test::ESIMDSelector, esimd_test::createExceptionHandler());
+
+  auto dev = q.get_device();
+  std::cout << "Running on " << dev.get_info<sycl::info::device::name>()
+            << "\n";
+
+  bool passed = true;
+
+  constexpr bool TestCacheHintProperties = true;
+  passed &= test_main<TestCacheHintProperties, false>(q);
+
+  std::cout << (passed ? "Passed\n" : "FAILED\n");
+  return passed ? 0 : 1;
+}
diff --git a/sycl/test-e2e/ESIMD/unified_memory_api/atomic_update_usm_dg2_pvc_64_2.cpp b/sycl/test-e2e/ESIMD/unified_memory_api/atomic_update_usm_dg2_pvc_64_2.cpp
new file mode 100644
index 0000000000000..d0fb7d274e9e4
--- /dev/null
+++ b/sycl/test-e2e/ESIMD/unified_memory_api/atomic_update_usm_dg2_pvc_64_2.cpp
@@ -0,0 +1,16 @@
+//==-- atomic_update_usm_dg2_pvc_64_2.cpp - DPC++ ESIMD on-device test -==//
+//
+// 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
+//
+//===----------------------------------------------------------------===//
+
+// REQUIRES: arch-intel_gpu_pvc || gpu-intel-dg2
+
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+
+#define USE_64_BIT_OFFSET
+
+#include "atomic_update_usm_dg2_pvc_2.cpp"
diff --git a/sycl/test-e2e/ESIMD/unified_memory_api/atomic_update_usm_dg2_pvc_cmpxchg_2.cpp b/sycl/test-e2e/ESIMD/unified_memory_api/atomic_update_usm_dg2_pvc_cmpxchg_2.cpp
new file mode 100644
index 0000000000000..7d7b29a31d061
--- /dev/null
+++ b/sycl/test-e2e/ESIMD/unified_memory_api/atomic_update_usm_dg2_pvc_cmpxchg_2.cpp
@@ -0,0 +1,17 @@
+//=- atomic_update_usm_dg2_pvc_cmpxchg_2.cpp- DPC++ ESIMD on-device test -=//
+//
+// 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
+//
+//===---------------------------------------------------------------------===//
+
+// REQUIRES: arch-intel_gpu_pvc || gpu-intel-dg2
+// REQUIRES-INTEL-DRIVER: lin: 29803
+
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+
+#define CMPXCHG_TEST
+
+#include "atomic_update_usm_dg2_pvc_2.cpp"