[ESIMD]Enable tests for DG2 (intel#14099)

sycl/test-e2e/ESIMD/lsc/Inputs/lsc_usm_block_load_prefetch.hpp

@@ -145,7 +145,14 @@ template <typename T> bool test_lsc_block_load() {
             << Q.get_device().get_info<sycl::info::device::name>() << std::endl;
 
   bool Passed = true;
-  Passed &= test<T, 64, DS, L1H, L2H, NoPrefetch, NoCheckMerge>(Q, 1, 4);
+  if constexpr (sizeof(T) * 64 < 256)
+    Passed &= test<T, 64, DS, L1H, L2H, NoPrefetch, NoCheckMerge>(Q, 1, 4);
+  else {
+#ifdef USE_PVC
+    Passed &= test<T, 64, DS, L1H, L2H, NoPrefetch, NoCheckMerge>(Q, 1, 4);
+#endif
+  }
+
   Passed &= test<T, 32, DS, L1H, L2H, NoPrefetch, NoCheckMerge>(Q, 1, 4);
   Passed &= test<T, 16, DS, L1H, L2H, NoPrefetch, NoCheckMerge>(Q, 2, 2);
   Passed &= test<T, 8, DS, L1H, L2H, NoPrefetch, NoCheckMerge>(Q, 2, 8);
@@ -154,6 +161,7 @@ template <typename T> bool test_lsc_block_load() {
     Passed &= test<T, 2, DS, L1H, L2H, NoPrefetch, NoCheckMerge>(Q, 5, 5);
   if constexpr (sizeof(T) >= sizeof(int))
     Passed &= test<T, 1, DS, L1H, L2H, NoPrefetch, CheckMerge>(Q, 3, 5);
+#ifdef USE_PVC
   if constexpr (sizeof(T) <= 4) {
     Passed &= test<T, 128, DS, L1H, L2H, NoPrefetch, CheckMerge,
                    __ESIMD_NS::overaligned_tag<8>>(Q, 1, 4);
@@ -172,8 +180,15 @@ template <typename T> bool test_lsc_block_load() {
                      __ESIMD_NS::overaligned_tag<8>>(Q, 1, 4);
     }
   }
-
-  Passed &= test<T, 64, DS, L1H, L2H, NoPrefetch, CheckMerge>(Q, 1, 4);
+#endif
+
+  if constexpr (sizeof(T) * 64 < 256)
+    Passed &= test<T, 64, DS, L1H, L2H, NoPrefetch, CheckMerge>(Q, 1, 4);
+  else {
+#ifdef USE_PVC
+    Passed &= test<T, 64, DS, L1H, L2H, NoPrefetch, CheckMerge>(Q, 1, 4);
+#endif
+  }
   Passed &= test<T, 32, DS, L1H, L2H, NoPrefetch, CheckMerge>(Q, 2, 2);
   Passed &= test<T, 16, DS, L1H, L2H, NoPrefetch, CheckMerge>(Q, 4, 4);
   Passed &= test<T, 8, DS, L1H, L2H, NoPrefetch, CheckMerge>(Q, 2, 8);
@@ -182,6 +197,7 @@ template <typename T> bool test_lsc_block_load() {
     Passed &= test<T, 2, DS, L1H, L2H, NoPrefetch, CheckMerge>(Q, 5, 5);
   if constexpr (sizeof(T) >= sizeof(int))
     Passed &= test<T, 1, DS, L1H, L2H, NoPrefetch, CheckMerge>(Q, 3, 5);
+#ifdef USE_PVC
   // Only 512-bits maximum can be loaded at once (i.e. 4*128 bytes).
   if constexpr (sizeof(T) <= 4)
     Passed &= test<T, 128, DS, L1H, L2H, NoPrefetch, CheckMerge,
@@ -192,7 +208,7 @@ template <typename T> bool test_lsc_block_load() {
   if constexpr (sizeof(T) == 1)
     Passed &= test<T, 512, DS, L1H, L2H, NoPrefetch, CheckMerge,
                    __ESIMD_NS::overaligned_tag<8>>(Q, 1, 4);
-
+#endif
   return Passed;
 }
 
@@ -209,7 +225,13 @@ std::enable_if_t<!IsGatherLikePrefetch, bool> test_lsc_prefetch() {
             << Q.get_device().get_info<sycl::info::device::name>() << std::endl;
 
   bool Passed = true;
-  Passed &= test<T, 64, DS, L1H, L2H, DoPrefetch>(Q, 1, 4);
+  if constexpr (sizeof(T) * 64 < 256)
+    Passed &= test<T, 64, DS, L1H, L2H, DoPrefetch>(Q, 1, 4);
+  else {
+#ifdef USE_PVC
+    Passed &= test<T, 64, DS, L1H, L2H, DoPrefetch>(Q, 1, 4);
+#endif
+  }
   Passed &= test<T, 32, DS, L1H, L2H, DoPrefetch>(Q, 1, 4);
   Passed &= test<T, 16, DS, L1H, L2H, DoPrefetch>(Q, 2, 2);
   Passed &= test<T, 8, DS, L1H, L2H, DoPrefetch>(Q, 2, 8);

sycl/test-e2e/ESIMD/lsc/lsc_usm_block_load_u32.cpp

@@ -5,7 +5,7 @@
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
-// REQUIRES: gpu-intel-pvc
+// REQUIRES: gpu-intel-dg2
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out
 
@@ -19,8 +19,10 @@ int main(void) {
 
   Passed &= test_lsc_block_load<uint32_t>();
   Passed &= test_lsc_block_load<float>();
+#ifdef USE_PVC
   Passed &=
       test_lsc_block_load<sycl::ext::intel::experimental::esimd::tfloat32>();
+#endif
 
   std::cout << (Passed ? "Passed\n" : "FAILED\n");
   return Passed ? 0 : 1;

sycl/test-e2e/ESIMD/lsc/lsc_usm_block_load_u32_64.cpp

@@ -5,7 +5,7 @@
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
-// REQUIRES: gpu-intel-pvc
+// REQUIRES: gpu-intel-dg2
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out
 

sycl/test-e2e/ESIMD/lsc/lsc_usm_block_load_u32_64_pvc.cpp

@@ -0,0 +1,17 @@
+//==--- lsc_usm_block_load_u32_64_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
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+
+// PVC variant of the test
+
+#define USE_64_BIT_OFFSET
+#define USE_PVC
+
+#include "lsc_usm_block_load_u32.cpp"

sycl/test-e2e/ESIMD/lsc/lsc_usm_block_load_u32_pvc.cpp

@@ -0,0 +1,17 @@
+//==--- lsc_usm_block_load_u32_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
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+
+// PVC variant of the test
+
+#define USE_64_BIT_OFFSET
+#define USE_PVC
+
+#include "lsc_usm_block_load_u32.cpp"

sycl/test-e2e/ESIMD/lsc/lsc_usm_block_load_u32_scalar_off.cpp

@@ -5,7 +5,7 @@
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
-// REQUIRES: gpu-intel-pvc
+// REQUIRES: gpu-intel-dg2
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out
 

sycl/test-e2e/ESIMD/lsc/lsc_usm_block_load_u32_scalar_off_pvc.cpp

@@ -0,0 +1,17 @@
+//==-lsc_usm_block_load_u32_scalar_off_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
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+
+// PVC variant of the test
+
+#define USE_SCALAR_OFFSET
+#define USE_PVC
+
+#include "lsc_usm_block_load_u32.cpp"

sycl/test-e2e/ESIMD/lsc/lsc_usm_block_load_u64.cpp

@@ -5,7 +5,7 @@
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
-// REQUIRES: gpu-intel-pvc
+// REQUIRES: gpu-intel-dg2
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out
 
@@ -15,10 +15,11 @@ constexpr uint32_t Seed = 187;
 
 int main(void) {
   srand(Seed);
-
   bool Passed = true;
   Passed &= test_lsc_block_load<uint64_t>();
+#ifdef USE_PVC
   Passed &= test_lsc_block_load<double>();
+#endif
 
   std::cout << (Passed ? "Passed\n" : "FAILED\n");
   return Passed ? 0 : 1;

sycl/test-e2e/ESIMD/lsc/lsc_usm_block_load_u64_64.cpp

@@ -5,7 +5,7 @@
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
-// REQUIRES: gpu-intel-pvc
+// REQUIRES: gpu-intel-dg2
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out
 

sycl/test-e2e/ESIMD/lsc/lsc_usm_block_load_u64_64_pvc.cpp

@@ -0,0 +1,17 @@
+//==------- lsc_usm_block_load_u64_64_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
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+
+// PVC variant of the test
+
+#define USE_64_BIT_OFFSET
+#define USE_PVC
+
+#include "lsc_usm_block_load_u64.cpp"

sycl/test-e2e/ESIMD/lsc/lsc_usm_block_load_u64_pvc.cpp

@@ -0,0 +1,16 @@
+//==------- lsc_usm_block_load_u64_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
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+
+// PVC variant of the test
+
+#define USE_PVC
+
+#include "lsc_usm_block_load_u64.cpp"

sycl/test-e2e/ESIMD/lsc/lsc_usm_block_load_u64_scalar_off.cpp

@@ -5,7 +5,7 @@
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
-// REQUIRES: gpu-intel-pvc
+// REQUIRES:  gpu-intel-dg2
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out
 

sycl/test-e2e/ESIMD/lsc/lsc_usm_block_load_u64_scalar_off_pvc.cpp

@@ -0,0 +1,17 @@
+//==---- lsc_usm_load_u64_scalar_off_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
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+
+// PVC variant of the test
+
+#define USE_SCALAR_OFFSET
+#define USE_PVC
+
+#include "lsc_usm_block_load_u64.cpp"

sycl/test-e2e/ESIMD/lsc/lsc_usm_block_load_u8_u16.cpp

@@ -5,7 +5,7 @@
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
-// REQUIRES: gpu-intel-pvc
+// REQUIRES: gpu-intel-dg2
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out
 

sycl/test-e2e/ESIMD/lsc/lsc_usm_block_load_u8_u16_64.cpp

@@ -5,7 +5,7 @@
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
-// REQUIRES: gpu-intel-pvc
+// REQUIRES: gpu-intel-dg2
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out
 

sycl/test-e2e/ESIMD/lsc/lsc_usm_block_load_u8_u16_64_pvc.cpp

@@ -0,0 +1,17 @@
+//==-- lsc_usm_block_load_u8_u16_64_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
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+
+// PVC variant of the test
+
+#define USE_64_BIT_OFFSET
+#define USE_PVC
+
+#include "lsc_usm_block_load_u8_u16.cpp"

sycl/test-e2e/ESIMD/lsc/lsc_usm_block_load_u8_u16_pvc.cpp

@@ -0,0 +1,17 @@
+//==----- lsc_usm_block_load_u8_u16_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
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+
+// PVC variant of the test
+
+#define USE_64_BIT_OFFSET
+#define USE_PVC
+
+#include "lsc_usm_block_load_u8_u16.cpp"

sycl/test-e2e/ESIMD/lsc/lsc_usm_prefetch_u32.cpp

@@ -5,7 +5,7 @@
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
-// REQUIRES: gpu-intel-pvc
+// REQUIRES: gpu-intel-dg2
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out
 
@@ -30,7 +30,9 @@ int main(void) {
 
   Passed &= tests<uint32_t>();
   Passed &= tests<float>();
+#ifdef USE_PVC
   Passed &= tests<sycl::ext::intel::experimental::esimd::tfloat32>();
+#endif
 
   std::cout << (Passed ? "Passed\n" : "FAILED\n");
   return Passed ? 0 : 1;

sycl/test-e2e/ESIMD/lsc/lsc_usm_prefetch_u32_64.cpp

@@ -5,7 +5,7 @@
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
-// REQUIRES: gpu-intel-pvc
+// REQUIRES: gpu-intel-dg2
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out
 

sycl/test-e2e/ESIMD/lsc/lsc_usm_prefetch_u32_64_pvc.cpp

@@ -0,0 +1,17 @@
+//==----- lsc_usm_prefetch_u32_64_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
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+
+// PVC variant of the test
+
+#define USE_64_BIT_OFFSET
+#define USE_PVC
+
+#include "lsc_usm_prefetch_u32.cpp"

sycl/test-e2e/ESIMD/lsc/lsc_usm_prefetch_u32_pvc.cpp

@@ -0,0 +1,16 @@
+//==------- lsc_usm_prefetch_u32_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
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+
+// PVC variant of the test
+
+#define USE_PVC
+
+#include "lsc_usm_prefetch_u32.cpp"

sycl/test-e2e/ESIMD/lsc/lsc_usm_prefetch_u32_scalar_off.cpp

@@ -5,7 +5,7 @@
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
-// REQUIRES: gpu-intel-pvc
+// REQUIRES: gpu-intel-dg2
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out