[SYCL][ESIMD][E2E] Add DG2 atomic_update tests (#12131)

Add atomic_update tests for DG2. I manually tested these on PVC Linux
and DG2 Linux/Windows.

For the normal accessor and USM tests, all tests work on DG2 with no
changes, so I just moved the PVC tests to be DG2/PVC tests instead of
making new files since there's no code change required besides the
`REQUIRED` line.

For the SLM and SLM accessor tests, we have to skip 64-bit types as
those are not supported for LSC SLM atomics on DG2, if we try to use
them we get an error explicitly saying that. So for these I did
something similar to what I did for block_load/block_store and just have
an if check to skip that case.

The diff is a little bit weird, some of the file matching is a bit
wrong, so you might have to look for the correct files to compare in the
file list.

---------

Signed-off-by: Sarnie, Nick <[email protected]>

Author: sarnex

sycl/test-e2e/ESIMD/unified_memory_api/Inputs/atomic_update.hpp

@@ -6,7 +6,7 @@
 //
 //===----------------------------------------------------------------------===//
 
-#include "../../esimd_test_utils.hpp"
+#include "common.hpp"
 
 #include <iostream>
 #include <sycl/ext/intel/esimd.hpp>
@@ -553,38 +553,44 @@ auto run_test(queue q) {
 }
 
 template <int N, template <class, int> class Op, bool UseMask,
-          bool UsePVCFeatures, bool UseAcc, int SignMask = (Signed | Unsigned)>
+          TestFeatures Features, bool UseAcc,
+          int SignMask = (Signed | Unsigned)>
 bool test_int_types(queue q) {
   bool passed = true;
   if constexpr (SignMask & Signed) {
-    if constexpr (UsePVCFeatures)
+    if constexpr (Features == TestFeatures::DG2 ||
+                  Features == TestFeatures::PVC)
       passed &= run_test<UseAcc, int16_t, N, Op, UseMask>(q);
 
     passed &= run_test<UseAcc, int32_t, N, Op, UseMask>(q);
 
-    if constexpr (UsePVCFeatures) {
+    // int64_t not supported on DG2
+    if constexpr (Features == TestFeatures::PVC) {
       passed &= run_test<UseAcc, int64_t, N, Op, UseMask>(q);
     }
   }
 
   if constexpr (SignMask & Unsigned) {
-    if constexpr (UsePVCFeatures)
+    if constexpr (Features == TestFeatures::DG2 ||
+                  Features == TestFeatures::PVC)
       passed &= run_test<UseAcc, uint16_t, N, Op, UseMask>(q);
 
     passed &= run_test<UseAcc, uint32_t, N, Op, UseMask>(q);
 
-    if constexpr (UsePVCFeatures) {
+    // uint64_t not supported on DG2
+    if constexpr (Features == TestFeatures::PVC) {
       passed &= run_test<UseAcc, uint64_t, N, Op, UseMask>(q);
     }
   }
   return passed;
 }
 
 template <int N, template <class, int> class Op, bool UseMask,
-          bool UsePVCFeatures, bool UseAcc>
+          TestFeatures Features, bool UseAcc>
 bool test_fp_types(queue q) {
   bool passed = true;
-  if constexpr (UsePVCFeatures) {
+  if constexpr (Features == TestFeatures::DG2 ||
+                Features == TestFeatures::PVC) {
     if constexpr (std::is_same_v<Op<sycl::half, N>,
                                  ImplLSCFmax<sycl::half, N>> ||
                   std::is_same_v<Op<sycl::half, N>,
@@ -600,7 +606,8 @@ bool test_fp_types(queue q) {
 
   passed &= run_test<UseAcc, float, N, Op, UseMask>(q);
 
-  if constexpr (UsePVCFeatures) {
+  if constexpr (Features == TestFeatures::DG2 ||
+                Features == TestFeatures::PVC) {
     // TODO: fmin/max for double does not pass validation likely due to
     // a driver bug. fcmpwr is hanging.
     if constexpr (!std::is_same_v<Op<double, N>, ImplLSCFmax<double, N>> &&
@@ -615,142 +622,133 @@ bool test_fp_types(queue q) {
   return passed;
 }
 
-template <template <class, int> class Op, bool UseMask, bool UsePVCFeatures,
+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 passed = true;
-  passed &= test_int_types<1, Op, UseMask, UsePVCFeatures, UseAcc, SignMask>(q);
-  passed &= test_int_types<2, Op, UseMask, UsePVCFeatures, UseAcc, SignMask>(q);
-  passed &= test_int_types<4, Op, UseMask, UsePVCFeatures, UseAcc, SignMask>(q);
-  passed &= test_int_types<8, Op, UseMask, UsePVCFeatures, UseAcc, SignMask>(q);
+  passed &= test_int_types<1, Op, UseMask, Features, UseAcc, SignMask>(q);
+  passed &= test_int_types<2, Op, UseMask, Features, UseAcc, SignMask>(q);
+  passed &= test_int_types<4, Op, UseMask, Features, UseAcc, SignMask>(q);
+  passed &= test_int_types<8, Op, UseMask, Features, UseAcc, SignMask>(q);
   // TODO: N=16 and N=32 does not pass on Gen12 with mask due to older driver.
-  if (UseMask && !UsePVCFeatures &&
+  if (UseMask && Features == TestFeatures::Generic &&
       esimd_test::isGPUDriverGE(q, esimd_test::GPUDriverOS::LinuxAndWindows,
                                 "26918", "101.4953", false)) {
-    passed &=
-        test_int_types<16, Op, UseMask, UsePVCFeatures, UseAcc, SignMask>(q);
-    passed &=
-        test_int_types<32, Op, UseMask, UsePVCFeatures, UseAcc, SignMask>(q);
+    passed &= test_int_types<16, Op, UseMask, Features, UseAcc, SignMask>(q);
+    passed &= test_int_types<32, Op, UseMask, Features, UseAcc, SignMask>(q);
   }
 
   // Supported by LSC atomic:
-  if constexpr (UsePVCFeatures) {
-    passed &=
-        test_int_types<64, Op, UseMask, UsePVCFeatures, UseAcc, SignMask>(q);
+  if constexpr (Features == TestFeatures::DG2 ||
+                Features == TestFeatures::PVC) {
+    passed &= test_int_types<64, Op, UseMask, Features, UseAcc, SignMask>(q);
     // non power of two values are supported only in newer driver.
     // TODO: Enable this when the new driver reaches test infrastructure
     // (v27556).
 #if 0
-    passed &= test_int_types<12, Op, UseMask, UsePVCFeatures, UseAcc, SignMask>(q);
-    passed &= test_int_types<33, Op, UseMask, UsePVCFeatures, UseAcc, SignMask>(q);
+    passed &= test_int_types<12, Op, UseMask, Features, UseAcc, SignMask>(q);
+    passed &= test_int_types<33, Op, UseMask, Features, UseAcc, SignMask>(q);
 #endif
   }
 
   return passed;
 }
 
-template <template <class, int> class Op, bool UseMask, bool UsePVCFeatures,
+template <template <class, int> class Op, bool UseMask, TestFeatures Features,
           bool UseAcc>
 bool test_fp_types_and_sizes(queue q) {
   bool passed = true;
-  passed &= test_fp_types<1, Op, UseMask, UsePVCFeatures, UseAcc>(q);
-  passed &= test_fp_types<2, Op, UseMask, UsePVCFeatures, UseAcc>(q);
-  passed &= test_fp_types<4, Op, UseMask, UsePVCFeatures, UseAcc>(q);
-  passed &= test_fp_types<8, Op, UseMask, UsePVCFeatures, UseAcc>(q);
-  passed &= test_fp_types<16, Op, UseMask, UsePVCFeatures, UseAcc>(q);
-  passed &= test_fp_types<32, Op, UseMask, UsePVCFeatures, UseAcc>(q);
+  passed &= test_fp_types<1, Op, UseMask, Features, UseAcc>(q);
+  passed &= test_fp_types<2, Op, UseMask, Features, UseAcc>(q);
+  passed &= test_fp_types<4, Op, UseMask, Features, UseAcc>(q);
+  passed &= test_fp_types<8, Op, UseMask, Features, UseAcc>(q);
+  passed &= test_fp_types<16, Op, UseMask, Features, UseAcc>(q);
+  passed &= test_fp_types<32, Op, UseMask, Features, UseAcc>(q);
 
   // Supported by LSC atomic:
-  if constexpr (UsePVCFeatures) {
-    passed &= test_fp_types<64, Op, UseMask, UsePVCFeatures, UseAcc>(q);
+  if constexpr (Features == TestFeatures::DG2 ||
+                Features == TestFeatures::PVC) {
+    passed &= test_fp_types<64, Op, UseMask, Features, UseAcc>(q);
     // non power of two values are supported only in newer driver.
     // TODO: Enable this when the new driver reaches test infrastructure
     // (v27556).
 #if 0
-    passed &= test_fp_types<33, Op, UseMask, UsePVCFeatures, UseAcc>(q);
-    passed &= test_fp_types<65, Op, UseMask, UsePVCFeatures, UseAcc>(q);
+    passed &= test_fp_types<33, Op, UseMask, Features, UseAcc>(q);
+    passed &= test_fp_types<65, Op, UseMask, Features, UseAcc>(q);
 #endif
   }
   return passed;
 }
 
-template <bool UseMask, bool UsePVCFeatures, bool UseAcc>
+template <bool UseMask, TestFeatures Features, bool UseAcc>
 int test_with_mask(queue q) {
   bool passed = true;
 #ifndef CMPXCHG_TEST
+  passed &= test_int_types_and_sizes<ImplInc, UseMask, Features, UseAcc>(q);
+  passed &= test_int_types_and_sizes<ImplDec, UseMask, Features, UseAcc>(q);
+
+  passed &= test_int_types_and_sizes<ImplIntAdd, UseMask, Features, UseAcc>(q);
+  passed &= test_int_types_and_sizes<ImplIntSub, UseMask, Features, UseAcc>(q);
+
   passed &=
-      test_int_types_and_sizes<ImplInc, UseMask, UsePVCFeatures, UseAcc>(q);
+      test_int_types_and_sizes<ImplSMax, UseMask, Features, UseAcc, Signed>(q);
   passed &=
-      test_int_types_and_sizes<ImplDec, UseMask, UsePVCFeatures, UseAcc>(q);
+      test_int_types_and_sizes<ImplSMin, UseMask, Features, UseAcc, Signed>(q);
 
   passed &=
-      test_int_types_and_sizes<ImplIntAdd, UseMask, UsePVCFeatures, UseAcc>(q);
+      test_int_types_and_sizes<ImplUMax, UseMask, Features, UseAcc, Unsigned>(
+          q);
   passed &=
-      test_int_types_and_sizes<ImplIntSub, UseMask, UsePVCFeatures, UseAcc>(q);
-
-  passed &= test_int_types_and_sizes<ImplSMax, UseMask, UsePVCFeatures, UseAcc,
-                                     Signed>(q);
-  passed &= test_int_types_and_sizes<ImplSMin, UseMask, UsePVCFeatures, UseAcc,
-                                     Signed>(q);
-
-  passed &= test_int_types_and_sizes<ImplUMax, UseMask, UsePVCFeatures, UseAcc,
-                                     Unsigned>(q);
-  passed &= test_int_types_and_sizes<ImplUMin, UseMask, UsePVCFeatures, UseAcc,
-                                     Unsigned>(q);
+      test_int_types_and_sizes<ImplUMin, UseMask, Features, UseAcc, Unsigned>(
+          q);
 
-  if constexpr (UsePVCFeatures) {
+  if constexpr (Features == TestFeatures::DG2 ||
+                Features == TestFeatures::PVC) {
     passed &=
-        test_fp_types_and_sizes<ImplLSCFmax, UseMask, UsePVCFeatures, UseAcc>(
-            q);
+        test_fp_types_and_sizes<ImplLSCFmax, UseMask, Features, UseAcc>(q);
     passed &=
-        test_fp_types_and_sizes<ImplLSCFmin, UseMask, UsePVCFeatures, UseAcc>(
-            q);
+        test_fp_types_and_sizes<ImplLSCFmin, UseMask, Features, UseAcc>(q);
 
     // TODO: fadd/fsub are emulated in the newer driver, but do not pass
     // validation.
 #if 0
-    passed &= test_fp_types_and_sizes<ImplFadd, UseMask, UsePVCFeatures, UseAcc>(q);
-    passed &= test_fp_types_and_sizes<ImplFsub, UseMask, UsePVCFeatures, UseAcc>(q);
+    passed &= test_fp_types_and_sizes<ImplFadd, UseMask, Features, UseAcc>(q);
+    passed &= test_fp_types_and_sizes<ImplFsub, UseMask, Features, UseAcc>(q);
 #endif
 
     // Check load/store operations.
-    passed &=
-        test_int_types_and_sizes<ImplLoad, UseMask, UsePVCFeatures, UseAcc>(q);
-    passed &=
-        test_int_types_and_sizes<ImplStore, UseMask, UsePVCFeatures, UseAcc>(q);
-    passed &=
-        test_fp_types_and_sizes<ImplStore, UseMask, UsePVCFeatures, UseAcc>(q);
+    passed &= test_int_types_and_sizes<ImplLoad, UseMask, Features, UseAcc>(q);
+    passed &= test_int_types_and_sizes<ImplStore, UseMask, Features, UseAcc>(q);
+    passed &= test_fp_types_and_sizes<ImplStore, UseMask, Features, UseAcc>(q);
   }
 #else
+  passed &= test_int_types_and_sizes<ImplCmpxchg, UseMask, Features, UseAcc>(q);
   passed &=
-      test_int_types_and_sizes<ImplCmpxchg, UseMask, UsePVCFeatures, UseAcc>(q);
-  passed &=
-      test_fp_types_and_sizes<ImplLSCFcmpwr, UseMask, UsePVCFeatures, UseAcc>(
-          q);
+      test_fp_types_and_sizes<ImplLSCFcmpwr, UseMask, Features, UseAcc>(q);
 #endif
   return passed;
 }
 
-template <bool UsePVCFeatures> bool test_main(queue q) {
+template <TestFeatures Features> bool test_main(queue q) {
   bool passed = true;
 
   constexpr const bool UseMask = true;
   constexpr const bool UseAcc = true;
 
-  passed &= test_with_mask<UseMask, UsePVCFeatures, !UseAcc>(q);
-  passed &= test_with_mask<!UseMask, UsePVCFeatures, !UseAcc>(q);
+  passed &= test_with_mask<UseMask, Features, !UseAcc>(q);
+  passed &= test_with_mask<!UseMask, Features, !UseAcc>(q);
 
   return passed;
 }
 
-template <bool UsePVCFeatures> bool test_main_acc(queue q) {
+template <TestFeatures Features> bool test_main_acc(queue q) {
   bool passed = true;
 
   constexpr const bool UseMask = true;
   constexpr const bool UseAcc = true;
 
-  passed &= test_with_mask<UseMask, UsePVCFeatures, UseAcc>(q);
-  passed &= test_with_mask<!UseMask, UsePVCFeatures, UseAcc>(q);
+  passed &= test_with_mask<UseMask, Features, UseAcc>(q);
+  passed &= test_with_mask<!UseMask, Features, UseAcc>(q);
 
   return passed;
 }

sycl/test-e2e/ESIMD/unified_memory_api/Inputs/atomic_update_slm.hpp

@@ -1,12 +1,12 @@
-//==------- atomic_update_acc_pvc.cpp - DPC++ ESIMD on-device test ---------==//
+//==----- atomic_update_acc_dg2_pvc.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: gpu-intel-pvc
+// REQUIRES: gpu-intel-pvc || gpu-intel-dg2
 
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out

sycl/test-e2e/ESIMD/unified_memory_api/atomic_update_acc_pvc.cpp renamed to sycl/test-e2e/ESIMD/unified_memory_api/atomic_update_acc_dg2_pvc.cpp

@@ -1,17 +1,17 @@
-//==-- atomic_update_acc_pvc_64.cpp - DPC++ ESIMD on-device test----==//
+//==-- atomic_update_acc_dg2_pvc_64.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: gpu-intel-pvc
+// REQUIRES: gpu-intel-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_pvc.cpp"
+#include "atomic_update_acc_dg2_pvc.cpp"

sycl/test-e2e/ESIMD/unified_memory_api/atomic_update_acc_pvc_64.cpp renamed to sycl/test-e2e/ESIMD/unified_memory_api/atomic_update_acc_dg2_pvc_64.cpp

@@ -1,14 +1,16 @@
-//==-- atomic_update_acc_pvc_stateless.cpp - DPC++ ESIMD on-device test --==//
+//==-- atomic_update_acc_dg2_pvc_cmpxchg.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: gpu-intel-pvc
+// REQUIRES: gpu-intel-pvc || gpu-intel-dg2
 
-// RUN: %{build} -fsycl-esimd-force-stateless-mem -o %t.out
+// RUN: %{build} -o %t.out
 // RUN: %{run} %t.out
 
-#include "atomic_update_acc_pvc.cpp"
+#define CMPXCHG_TEST
+
+#include "atomic_update_acc_dg2_pvc.cpp"

sycl/test-e2e/ESIMD/unified_memory_api/atomic_update_acc_pvc_stateless.cpp renamed to sycl/test-e2e/ESIMD/unified_memory_api/atomic_update_acc_dg2_pvc_cmpxchg.cpp

@@ -0,0 +1,14 @@
+//=- atomic_update_acc_dg2_pvc_stateless.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: gpu-intel-pvc || gpu-intel-dg2
+
+// RUN: %{build} -fsycl-esimd-force-stateless-mem -o %t.out
+// RUN: %{run} %t.out
+
+#include "atomic_update_acc_dg2_pvc.cpp"

sycl/test-e2e/ESIMD/unified_memory_api/atomic_update_acc_dg2_pvc_stateless.cpp

@@ -15,8 +15,8 @@ int main(void) {
 
   esimd_test::printTestLabel(q);
 
-  constexpr bool TestCacheHintProperties = true;
-  bool passed = test_main<!TestCacheHintProperties>(q);
+  constexpr auto Features = TestFeatures::Generic;
+  bool passed = test_main<Features>(q);
 
   std::cout << (passed ? "Passed\n" : "FAILED\n");
   return passed ? 0 : 1;

sycl/test-e2e/ESIMD/unified_memory_api/atomic_update_slm.cpp

@@ -16,8 +16,8 @@ int main(void) {
 
   esimd_test::printTestLabel(q);
 
-  constexpr bool TestCacheHintProperties = true;
-  bool passed = test_main_acc<!TestCacheHintProperties>(q);
+  constexpr auto Features = TestFeatures::Generic;
+  bool passed = test_main_acc<Features>(q);
 
   std::cout << (passed ? "Passed\n" : "FAILED\n");
   return passed ? 0 : 1;

sycl/test-e2e/ESIMD/unified_memory_api/atomic_update_slm_acc.cpp

@@ -0,0 +1,26 @@
+//==------- atomic_update_slm_acc_dg2.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: gpu-intel-dg2
+
+// 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::DG2;
+  bool passed = test_main_acc<Features>(q);
+
+  std::cout << (passed ? "Passed\n" : "FAILED\n");
+  return passed ? 0 : 1;
+}