[SYCL] Fix image_array.cpp on GPU (#17166)

Fix issue #15398 by adding group
barriers between write and read statements. Also revised the parameters
of the second write-read test case (from `42, 42` to `53, 55`) to
prevent failures from it been covered by success of the first write-read
test case.

---------

Signed-off-by: peisenhu <peisen.hu@intel.com>
Signed-off-by: Hu, Peisen <peisen.hu@intel.com>

Author: HPS-1

sycl/test-e2e/Basic/image/image_array.cpp

@@ -4,9 +4,6 @@
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out
 
-// UNSUPPORTED: gpu
-// UNSUPPORTED-TRACKER: https://github.com/intel/llvm/issues/15398
-
 //==------------------- image.cpp - SYCL image basic test -----------------==//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
@@ -18,6 +15,7 @@
 #include <sycl/accessor_image.hpp>
 #include <sycl/builtins.hpp>
 #include <sycl/detail/core.hpp>
+#include <sycl/group_barrier.hpp>
 
 #include <iostream>
 #include <vector>
@@ -32,6 +30,7 @@ int main() {
   const sycl::image_channel_type ChanType = sycl::image_channel_type::fp32;
 
   const sycl::range<2> ImgSize(4, 4);
+  const sycl::nd_range<2> NDImgSize{ImgSize, ImgSize};
 
   std::vector<sycl::float4> Img1HostData(ImgSize.size(), {1, 2, 3, 4});
 
@@ -119,12 +118,13 @@ int main() {
 
       auto ResAcc = ResBuf.get_access<SYCLReadWrite>(CGH);
 
-      CGH.parallel_for<class Check2>(ImgSize, [=](sycl::item<2> Item) {
-        sycl::int2 CoordI{Item[0], Item[1]};
+      CGH.parallel_for<class Check2>(NDImgSize, [=](sycl::nd_item<2> Item) {
+        sycl::int2 CoordI{Item.get_global_id(0), Item.get_global_id(1)};
 
         // CHeck that data written using image array
         const sycl::float4 ValRef{CoordI.x(), 42, 42, CoordI.y()};
         ImgArrayAcc[CoordI.y()].write((int)CoordI.x(), ValRef);
+        sycl::group_barrier(Item.get_group());
         auto Val = ImgAcc.read(CoordI);
 
         ResAcc[WRITE1] |= sycl::any(sycl::isnotequal(Val, ValRef));
@@ -139,12 +139,13 @@ int main() {
 
       auto ResAcc = ResBuf.get_access<SYCLReadWrite>(CGH);
 
-      CGH.parallel_for<class Check3>(ImgSize, [=](sycl::item<2> Item) {
-        sycl::int2 CoordI{Item[0], Item[1]};
+      CGH.parallel_for<class Check3>(NDImgSize, [=](sycl::nd_item<2> Item) {
+        sycl::int2 CoordI{Item.get_global_id(0), Item.get_global_id(1)};
 
         // CHeck that data read using image array
-        const sycl::float4 ValRef{CoordI.x(), 42, 42, CoordI.y()};
+        const sycl::float4 ValRef{CoordI.x(), 53, 55, CoordI.y()};
         ImgAcc.write(CoordI, ValRef);
+        sycl::group_barrier(Item.get_group());
         auto Val = ImgArrayAcc[CoordI.y()].read((int)CoordI.x());
 
         ResAcc[WRITE2] |= sycl::any(sycl::isnotequal(Val, ValRef));