[SYCL][ESIMD] Reenable local_accessor atomic_update() tests on DG2 (#12932)

Author: fineg74

sycl/test-e2e/ESIMD/dword_local_accessor_atomic_smoke.cpp

@@ -8,8 +8,6 @@
 // This test checks DWORD local accessor atomic operations.
 //===----------------------------------------------------------------------===//
 // REQUIRES-INTEL-DRIVER: lin: 26690, win: 101.4576
-// TODO: disabled temporarily because of flaky issue.
-// UNSUPPORTED: windows
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out
 //

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

@@ -180,8 +180,8 @@ bool test(queue q) {
       auto accessor = local_accessor<T, 1>(size, cgh);
 
       cgh.parallel_for<TestID<T, N, ImplF>>(
-          rng, [=](id<1> ii) SYCL_ESIMD_KERNEL {
-            int i = ii;
+          rng, [=](sycl::nd_item<1> ndi) SYCL_ESIMD_KERNEL {
+            int i = ndi.get_global_id(0);
 #ifndef USE_SCALAR_OFFSET
             simd<uint32_t, N> offsets(start_ind * sizeof(T),
                                       stride * sizeof(T));
@@ -192,7 +192,8 @@ bool test(queue q) {
             data.copy_from(arr);
 
             simd<uint32_t, size> LocalOffsets(0, sizeof(T));
-            scatter<T, size>(accessor, LocalOffsets, data, 0, 1);
+            if (ndi.get_local_id(0) == 0)
+              scatter<T, size>(accessor, LocalOffsets, data, 0, 1);
             simd_mask<N> m = 1;
             if (masked_lane < N)
               m[masked_lane] = 0;
@@ -221,8 +222,11 @@ bool test(queue q) {
                   ;
               }
             }
-            auto data0 = gather<T, size>(accessor, LocalOffsets, 0);
-            data0.copy_to(arr);
+            barrier();
+            if (ndi.get_local_id(0) == 0) {
+              auto data0 = gather<T, size>(accessor, LocalOffsets, 0);
+              data0.copy_to(arr);
+            }
           });
     });
     e.wait();

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

@@ -7,7 +7,7 @@
 //===----------------------------------------------------------------------===//
 // This test checks local accessor cmpxchg atomic operations.
 //===----------------------------------------------------------------------===//
-// REQUIRES: gpu-intel-pvc
+// REQUIRES: gpu-intel-pvc || gpu-intel-dg2
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out
 //

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

@@ -7,7 +7,7 @@
 //===----------------------------------------------------------------------===//
 // This test checks LSC SLM atomic operations.
 //===----------------------------------------------------------------------===//
-// REQUIRES: gpu-intel-pvc
+// REQUIRES: gpu-intel-pvc || gpu-intel-dg2
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out
 
@@ -123,16 +123,17 @@ bool test(queue q) {
   try {
     auto e = q.submit([&](handler &cgh) {
       cgh.parallel_for<TestID<T, N, ImplF>>(
-          rng, [=](id<1> ii) SYCL_ESIMD_KERNEL {
-            int i = ii;
+          rng, [=](sycl::nd_item<1> ndi) SYCL_ESIMD_KERNEL {
+            int i = ndi.get_global_id(0);
             slm_init<32768>();
             simd<uint32_t, N> offsets(start_ind * sizeof(T),
                                       stride * sizeof(T));
             simd<T, size> data;
             data.copy_from(arr);
 
             simd<uint32_t, size> slm_offsets(0, sizeof(T));
-            lsc_slm_scatter(slm_offsets, data);
+            if (ndi.get_local_id(0) == 0)
+              lsc_slm_scatter(slm_offsets, data);
 
             simd_mask<N> m = 1;
             if (masked_lane < N)
@@ -161,8 +162,11 @@ bool test(queue q) {
                   ;
               }
             }
-            auto data0 = lsc_slm_gather<T>(slm_offsets);
-            data0.copy_to(arr);
+            barrier();
+            if (ndi.get_local_id(0) == 0) {
+              auto data0 = lsc_slm_gather<T>(slm_offsets);
+              data0.copy_to(arr);
+            }
           });
     });
     e.wait();