[ESIMD] Add test for accessor based gather/scatter_rgba

SYCL/ESIMD/PrefixSum.cpp

@@ -160,8 +160,8 @@ void cmk_acum_iterative(unsigned *buf, unsigned h_pos,
   cnt_table.select<1, 1, TUPLE_SZ, 1>(0, 0) +=
       cnt_table.select<1, 1, TUPLE_SZ, 1>(1, 0);
 
-  simd<unsigned, 8> voff(0, 1);        // 0, 1, 2, 3
-  simd_mask<8> p = voff < TUPLE_SZ;    // predicate
+  simd<unsigned, 8> voff(0, 1);     // 0, 1, 2, 3
+  simd_mask<8> p = voff < TUPLE_SZ; // predicate
   voff = (voff + (global_offset + stride_threads * TUPLE_SZ - TUPLE_SZ)) *
          sizeof(unsigned);
   scatter<unsigned, 8>(buf, voff, S.select<8, 1>(0), p);
@@ -183,8 +183,7 @@ void cmk_acum_final(unsigned *buf, unsigned h_pos, unsigned int stride_elems,
 
     simd_mask<32> p = elm32 < remaining;
 
-    S = gather_rgba<unsigned int, 32, GATHER_SCATTER_MASK>(buf, element_offset,
-                                                           p);
+    S = gather_rgba<GATHER_SCATTER_MASK>(buf, element_offset, p);
 
     auto cnt_table = S.bit_cast_view<unsigned int, TUPLE_SZ, 32>();
     cnt_table.column(0) += prev;
@@ -214,8 +213,7 @@ void cmk_acum_final(unsigned *buf, unsigned h_pos, unsigned int stride_elems,
       cnt_table.select<1, 1, 16, 1>(j, 16) +=
           cnt_table.replicate_vs_w_hs<1, 0, 16, 0>(j, 15);
     }
-    scatter_rgba<unsigned int, 32, GATHER_SCATTER_MASK>(buf, element_offset, S,
-                                                        p);
+    scatter_rgba<GATHER_SCATTER_MASK>(buf, element_offset, S, p);
     elm32 += 32;
     element_offset += stride_elems * TUPLE_SZ * sizeof(unsigned) * 32;
     prev = cnt_table.column(31);
@@ -253,7 +251,7 @@ void cmk_prefix_iterative(unsigned *buf, unsigned h_pos,
   unsigned n_iter = n_entries / 32;
   for (unsigned i = 0; i < n_iter; i++) {
 
-    S = gather_rgba<unsigned int, 32, GATHER_SCATTER_MASK>(buf, element_offset);
+    S = gather_rgba<GATHER_SCATTER_MASK>(buf, element_offset);
 
     auto cnt_table = S.bit_cast_view<unsigned int, TUPLE_SZ, 32>();
     cnt_table.column(0) += prev;
@@ -289,7 +287,7 @@ void cmk_prefix_iterative(unsigned *buf, unsigned h_pos,
     if (i == n_iter - 1)
       cnt_table.column(31) -= cnt_table.column(30);
 
-    scatter_rgba<unsigned int, 32, GATHER_SCATTER_MASK>(buf, element_offset, S);
+    scatter_rgba<GATHER_SCATTER_MASK>(buf, element_offset, S);
 
     element_offset += stride_elems * TUPLE_SZ * sizeof(unsigned) * 32;
     prev = cnt_table.column(31);

SYCL/ESIMD/Prefix_Local_sum2.cpp

@@ -72,13 +72,13 @@ void cmk_acum_iterative(unsigned *buf, unsigned h_pos,
 
   simd<unsigned int, 32 * TUPLE_SZ> S, T;
 
-  S = gather_rgba<unsigned int, 32, GATHER_SCATTER_MASK>(buf, element_offset);
+  S = gather_rgba<GATHER_SCATTER_MASK>(buf, element_offset);
 
 #pragma unroll
   for (int i = 1; i < PREFIX_ENTRIES / 32; i++) {
     element_offset += (stride_elems * 32 * TUPLE_SZ) * sizeof(unsigned);
     // scattered read, each inst reads 16 entries
-    T = gather_rgba<unsigned int, 32, GATHER_SCATTER_MASK>(buf, element_offset);
+    T = gather_rgba<GATHER_SCATTER_MASK>(buf, element_offset);
     S += T;
   }
 
@@ -93,8 +93,8 @@ void cmk_acum_iterative(unsigned *buf, unsigned h_pos,
 
   simd<unsigned, 8> result = 0;
   result.select<TUPLE_SZ, 1>(0) = sum;
-  simd<unsigned, 8> voff(0, 1);        // 0, 1, 2, 3
-  simd_mask<8> p = voff < TUPLE_SZ;    // predicate
+  simd<unsigned, 8> voff(0, 1);     // 0, 1, 2, 3
+  simd_mask<8> p = voff < TUPLE_SZ; // predicate
   voff = (voff + (global_offset + stride_threads * TUPLE_SZ - TUPLE_SZ)) *
          sizeof(unsigned);
   scatter<unsigned, 8>(buf, voff, result, p);

SYCL/ESIMD/Prefix_Local_sum3.cpp

@@ -122,8 +122,8 @@ void cmk_acum_iterative(unsigned *buf, unsigned h_pos,
   cnt_table.select<1, 1, TUPLE_SZ, 1>(0, 0) +=
       cnt_table.select<1, 1, TUPLE_SZ, 1>(1, 0);
 
-  simd<unsigned, 8> voff(0, 1);        // 0, 1, 2, 3
-  simd_mask<8> p = voff < TUPLE_SZ;    // predicate
+  simd<unsigned, 8> voff(0, 1);     // 0, 1, 2, 3
+  simd_mask<8> p = voff < TUPLE_SZ; // predicate
   voff = (voff + (global_offset + stride_threads * TUPLE_SZ - TUPLE_SZ)) *
          sizeof(unsigned);
   scatter<unsigned, 8>(buf, voff, S.select<8, 1>(0), p);
@@ -173,8 +173,8 @@ void cmk_acum_iterative_low(unsigned *buf, unsigned h_pos,
   cnt_table.select<1, 1, TUPLE_SZ, 1>(0, 0) +=
       cnt_table.select<1, 1, TUPLE_SZ, 1>(1, 0);
 
-  simd<unsigned, 8> voff(0, 1);        // 0, 1, 2, 3
-  simd_mask<8> p = voff < TUPLE_SZ;    // predicate
+  simd<unsigned, 8> voff(0, 1);     // 0, 1, 2, 3
+  simd_mask<8> p = voff < TUPLE_SZ; // predicate
   voff = (voff + (global_offset + stride_threads * TUPLE_SZ - TUPLE_SZ)) *
          sizeof(unsigned);
   scatter<unsigned, 8>(buf, voff, S.select<8, 1>(0), p);
@@ -196,8 +196,7 @@ void cmk_acum_final(unsigned *buf, unsigned h_pos, unsigned int stride_elems,
 
     simd_mask<32> p = elm32 < remaining;
 
-    S = gather_rgba<unsigned int, 32, GATHER_SCATTER_MASK>(buf, element_offset,
-                                                           p);
+    S = gather_rgba<GATHER_SCATTER_MASK>(buf, element_offset, p);
 
     auto cnt_table = S.bit_cast_view<unsigned int, TUPLE_SZ, 32>();
     cnt_table.column(0) += prev;
@@ -226,8 +225,7 @@ void cmk_acum_final(unsigned *buf, unsigned h_pos, unsigned int stride_elems,
       cnt_table.select<1, 1, 16, 1>(j, 16) +=
           cnt_table.replicate_vs_w_hs<1, 0, 16, 0>(j, 15);
     }
-    scatter_rgba<unsigned int, 32, GATHER_SCATTER_MASK>(buf, element_offset, S,
-                                                        p);
+    scatter_rgba<GATHER_SCATTER_MASK>(buf, element_offset, S, p);
     elm32 += 32;
     element_offset += stride_elems * TUPLE_SZ * sizeof(unsigned) * 32;
     prev = cnt_table.column(31);

SYCL/ESIMD/acc_gather_scatter_rgba.cpp

@@ -0,0 +1,196 @@
+//==-------- acc_gather_scatter_rgba.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
+// UNSUPPORTED: cuda || hip
+// RUN: %clangxx -fsycl %s -o %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+//
+// The test checks functionality of the gather_rgba/scatter_rgba accessor-based
+// ESIMD intrinsics.
+
+#include "esimd_test_utils.hpp"
+
+#include <CL/sycl.hpp>
+#include <iostream>
+#include <sycl/ext/intel/esimd.hpp>
+
+using namespace cl::sycl;
+
+template <typename T>
+using AccT = accessor<T, 1, access_mode::read_write, access::target::device>;
+
+constexpr int MASKED_LANE_NUM_REV = 1;
+constexpr int NUM_RGBA_CHANNELS =
+    get_num_channels_enabled(sycl::ext::intel::esimd::rgba_channel_mask::ABGR);
+
+template <typename T, unsigned VL, unsigned STRIDE, auto CH_MASK>
+struct Kernel {
+  AccT<T> InAcc;
+  AccT<T> OutAcc;
+  Kernel(AccT<T> InAcc, AccT<T> OutAcc) : InAcc(InAcc), OutAcc(OutAcc) {}
+
+  void operator()(id<1> i) const SYCL_ESIMD_KERNEL {
+    using namespace sycl::ext::intel::esimd;
+    constexpr int numChannels = get_num_channels_enabled(CH_MASK);
+
+    // Every workitem accesses contiguous block of VL * STRIDE elements,
+    // where each element consists of RGBA channels.
+    uint32_t global_offset = i * VL * STRIDE * NUM_RGBA_CHANNELS * sizeof(T);
+
+    simd<uint32_t, VL> byteOffsets(0, STRIDE * sizeof(T) * NUM_RGBA_CHANNELS);
+    simd<T, VL * numChannels> v;
+    if constexpr (CH_MASK == rgba_channel_mask::ABGR)
+      // Check that the default mask value is ABGR.
+      v = gather_rgba(InAcc, byteOffsets, global_offset);
+    else
+      v = gather_rgba<CH_MASK>(InAcc, byteOffsets, global_offset);
+    v += (int)i;
+
+    simd_mask<VL> pred = 1;
+    pred[VL - MASKED_LANE_NUM_REV] = 0; // mask out the last lane
+    scatter_rgba<CH_MASK>(OutAcc, byteOffsets, v, global_offset, pred);
+  }
+};
+
+std::string convertMaskToStr(sycl::ext::intel::esimd::rgba_channel_mask mask) {
+  using namespace sycl::ext::intel::esimd;
+  switch (mask) {
+  case rgba_channel_mask::R:
+    return "R";
+  case rgba_channel_mask::GR:
+    return "GR";
+  case rgba_channel_mask::ABGR:
+    return "ABGR";
+  default:
+    return "";
+  }
+  return "";
+}
+
+template <typename T, unsigned VL, unsigned STRIDE, auto CH_MASK>
+bool test(queue q) {
+  size_t numWorkItems = 2;
+  size_t size = VL * STRIDE * NUM_RGBA_CHANNELS * numWorkItems;
+  using namespace sycl::ext::intel::esimd;
+  constexpr int numChannels = get_num_channels_enabled(CH_MASK);
+
+  std::cout << "Testing T=" << typeid(T).name() << " VL=" << VL
+            << " STRIDE=" << STRIDE << " MASK=" << convertMaskToStr(CH_MASK)
+            << "...\t";
+
+  T *A = new T[size];
+  T *B = new T[size];
+  T *gold = new T[size];
+
+  for (int i = 0; i < size; ++i) {
+    A[i] = (T)i;
+    B[i] = (T)-i;
+    gold[i] = (T)-i;
+  }
+
+  // Fill out the array with gold values. The kernel only writes the elements
+  // that are not masked. For example,
+  // for STRIDE=1 and MASK=R, we have the following indices written:
+  //   0, 4, 8, 12 ...
+  // for STRIDE=2 and MASK=RG, we have the following indices written:
+  //   0, 1, 8, 9, 16, 17 ...
+  // All the other elements will be equal to '-A[i]'.
+  auto blockSize = VL * STRIDE * NUM_RGBA_CHANNELS;
+  for (unsigned i = 0; i < size; i += NUM_RGBA_CHANNELS * STRIDE)
+    for (unsigned j = 0; j < numChannels; j++)
+      gold[i + j] = A[i + j] + (i / (blockSize));
+
+  // Account for masked out last lanes (with pred argument to scatter_rgba).
+  auto maskedElementOffset = (VL - 1) * STRIDE * NUM_RGBA_CHANNELS;
+  for (unsigned i = maskedElementOffset; i < size; i += blockSize)
+    for (unsigned j = 0; j < numChannels; j++)
+      gold[i + j] = -A[i + j];
+
+  try {
+    buffer<T, 1> InBuf(A, range<1>(size));
+    buffer<T, 1> OutBuf(B, range<1>(size));
+    range<1> glob_range{numWorkItems};
+    auto e = q.submit([&](handler &cgh) {
+      auto InAcc = InBuf.template get_access<access::mode::read_write>(cgh);
+      auto OutAcc = OutBuf.template get_access<access::mode::read_write>(cgh);
+      Kernel<T, VL, STRIDE, CH_MASK> kernel(InAcc, OutAcc);
+      cgh.parallel_for(glob_range, kernel);
+    });
+    e.wait();
+  } catch (sycl::exception const &e) {
+    std::cerr << "SYCL exception caught: " << e.what() << '\n';
+    delete[] A;
+    delete[] B;
+    delete[] gold;
+    return false; // not success
+  }
+
+  int err_cnt = 0;
+  for (unsigned i = 0; i < size; ++i) {
+    if (B[i] != gold[i]) {
+      if (++err_cnt < 35) {
+        std::cout << "\nFAILED at index " << i << ": " << B[i]
+                  << " != " << gold[i] << " (gold)";
+      }
+    }
+  }
+
+  if (err_cnt > 0) {
+    std::cout << "\n  pass rate: "
+              << ((float)(size - err_cnt) / (float)size) * 100.0f << "% ("
+              << (size - err_cnt) << "/" << size << ")\n";
+  }
+
+  delete[] A;
+  delete[] B;
+  delete[] gold;
+
+  if (err_cnt == 0)
+    std::cout << "Passed\n";
+  return err_cnt == 0;
+}
+
+template <typename T, unsigned VL, unsigned STRIDE> bool test(queue q) {
+  using namespace sycl::ext::intel::esimd;
+  bool passed = true;
+  passed &= test<T, VL, STRIDE, rgba_channel_mask::R>(q);
+  passed &= test<T, VL, STRIDE, rgba_channel_mask::GR>(q);
+  passed &= test<T, VL, STRIDE, rgba_channel_mask::ABGR>(q);
+  return passed;
+}
+
+int main(void) {
+  queue q(esimd_test::ESIMDSelector{}, esimd_test::createExceptionHandler());
+
+  auto dev = q.get_device();
+  std::cout << "Running on " << dev.get_info<info::device::name>() << "\n";
+
+  bool passed = true;
+  passed &= test<int, 16, 1>(q);
+  passed &= test<int, 16, 2>(q);
+  passed &= test<int, 16, 4>(q);
+  passed &= test<int, 32, 1>(q);
+  passed &= test<int, 32, 3>(q);
+  passed &= test<int, 32, 8>(q);
+  passed &= test<float, 16, 1>(q);
+  passed &= test<float, 16, 2>(q);
+  passed &= test<float, 16, 4>(q);
+  passed &= test<float, 32, 1>(q);
+  passed &= test<float, 32, 3>(q);
+  passed &= test<float, 32, 8>(q);
+
+  passed &= test<int, 8, 1>(q);
+  passed &= test<int, 8, 3>(q);
+  passed &= test<int, 8, 8>(q);
+  passed &= test<float, 8, 1>(q);
+  passed &= test<float, 8, 2>(q);
+  passed &= test<float, 8, 4>(q);
+
+  std::cout << (passed ? "All tests passed.\n" : "Some tests failed!\n");
+  return passed ? 0 : 1;
+}

SYCL/ESIMD/api/esimd_rgba_smoke.cpp

@@ -92,11 +92,11 @@ bool test_impl(queue q) {
       cgh.single_task<TestID<T, NPixels, static_cast<int>(ChMask)>>(
           [=]() SYCL_ESIMD_KERNEL {
             constexpr unsigned NElems = NPixels * NOnChs;
-            simd<T, NPixels> offsets(0, sizeof(T) * NAllChs);
-            simd<T, NElems> p = gather_rgba<T, NPixels, ChMask>(A, offsets);
+            simd<unsigned int, NPixels> offsets(0, sizeof(T) * NAllChs);
+            simd<T, NElems> p = gather_rgba<ChMask>(A, offsets);
             // simply scatter back to B - should give same results as A in
             // enabled channels, the rest should remain zero:
-            scatter_rgba<T, NPixels, ChMask>(B, offsets, p);
+            scatter_rgba<ChMask>(B, offsets, p);
             // copy instead of scattering to C - thus getting AOS to SOA layout
             // layout conversion:
             //   R0 R1 ... G0 G1 ... B0 B1 ... A0 A1 ...