[SYCL][ESIMD] Fix local accessor scatter test failure on PVC (#12745)

intel · Feb 22, 2024 · a261ac1 · a261ac1
1 parent 0cfe7e3
commit a261ac1
Showing 1 changed file with 17 additions and 6 deletions.
diff --git a/sycl/test-e2e/ESIMD/unified_memory_api/Inputs/scatter.hpp b/sycl/test-e2e/ESIMD/unified_memory_api/Inputs/scatter.hpp
@@ -672,7 +672,7 @@ bool testLACC(queue Q, uint32_t MaskStride,
 
   try {
     Q.submit([&](handler &cgh) {
-       constexpr uint32_t SLMSize = N;
+       constexpr uint32_t SLMSize = (Threads * N + 8);
        auto LocalAcc = local_accessor<T, 1>(SLMSize, cgh);
 
        cgh.parallel_for(Range, [=](sycl::nd_item<1> ndi) SYCL_ESIMD_KERNEL {
@@ -682,10 +682,17 @@ bool testLACC(queue Q, uint32_t MaskStride,
          uint32_t GlobalElemOffset = GlobalID * N;
          uint32_t LocalElemOffset = LocalID * N;
 
-         simd<T, N> InVec(GlobalElemOffset, 1);
+         if (LocalID == 0) {
+           for (int I = 0; I < Threads * N; I += 8) {
+             simd<T, 8> InVec(Out + GlobalElemOffset + I);
+             simd<uint32_t, 8> Offsets(I * sizeof(T), sizeof(T));
+             scatter<T>(LocalAcc, Offsets, InVec);
+           }
+         }
+         barrier();
 
-         simd<uint32_t, NOffsets> ByteOffsets(0, VS * sizeof(T));
-         scatter<T, N, VS>(LocalAcc, ByteOffsets, InVec);
+         simd<uint32_t, NOffsets> ByteOffsets(LocalElemOffset * sizeof(T),
+                                              VS * sizeof(T));
          auto ByteOffsetsView = ByteOffsets.template select<NOffsets, 1>();
          simd<T, N> Vals = gather<T, N, VS>(LocalAcc, ByteOffsets, Props);
 
@@ -786,8 +793,12 @@ bool testLACC(queue Q, uint32_t MaskStride,
            }
          }
 
-         simd<T, N> OutVec = gather<T, N, VS>(LocalAcc, ByteOffsets, Props);
-         OutVec.copy_to(Out + GlobalElemOffset);
+         barrier();
+         if (LocalID == 0) {
+           for (int I = 0; I < Threads * N; I++) {
+             Out[GlobalElemOffset + I] = LocalAcc[I];
+           }
+         }
        });
      }).wait();
   } catch (sycl::exception const &e) {