Remove accessor & add local_accessor

Author: KornevNikita

SYCL/Basic/accessor/accessor.cpp

@@ -61,103 +61,44 @@ template <typename Acc> struct Wrapper2 {
 
 template <typename Acc> struct Wrapper3 { Wrapper2<Acc> w2; };
 
-void fillVec(std::vector<int> &vec) {
-  for (size_t i = 0; i < vec.size(); ++i) {
-    vec[i] = i;
-  }
-}
-
-template <typename T> void modifyAccBuf(const T &acc, bool useReverse) {
-  if (!useReverse) {
-    for (auto It = acc.begin(); It != acc.end(); It++)
-      *It *= 2;
-    for (auto &It : acc)
-      It += 1;
-  } else {
-    for (auto It = acc.rbegin(); It != acc.rend(); It++) {
-      *It *= 2;
-      (*It)++;
-    }
-  }
-}
-
-template <typename T1, typename T2>
-void copyAccBuf(const T1 &accProd, T2 &accCons, bool useReverse) {
-  if (!useReverse) {
-    auto ItCons = accCons.begin();
-    for (auto ItProd = accProd.cbegin(); ItProd != accProd.cend(); ++ItProd)
-      *(ItCons++) = *ItProd;
-  } else {
-    auto ItCons = accCons.rbegin();
-    for (auto ItProd = accProd.crbegin(); ItProd != accProd.crend(); ++ItProd)
-      *(ItCons++) = *ItProd;
-  }
-}
-
-template <typename T> void testAccModImpl(T &buf, bool useReverse) {
-  sycl::queue q;
-  q.submit([&](sycl::handler &cgh) {
-    auto devAcc = buf.template get_access<sycl::access::mode::read_write>(cgh);
-    cgh.single_task([=]() { modifyAccBuf(devAcc, useReverse); });
-  });
-  q.wait();
-
-  auto hostAcc = buf.template get_access<sycl::access::mode::read_write>();
-  modifyAccBuf(hostAcc, useReverse);
-}
-
-void testAccMod(std::vector<int> &vec, bool useReverse = false, int Dim = 1) {
+void testLocalAccIters(std::vector<int> &vec, bool testConstIter = false) {
   try {
-    if (Dim == 1) {
-      sycl::buffer<int> buf(vec.data(), vec.size());
-      testAccModImpl(buf, useReverse);
-    } else if (Dim == 3) {
-      sycl::buffer<int, 3> buf(vec.data(), sycl::range<3>(2, 3, 4));
-      testAccModImpl(buf, useReverse);
-    }
-  } catch (sycl::exception &e) {
-    std::cout << e.what() << std::endl;
-  }
-
-  for (size_t i = 0; i < vec.size(); ++i)
-    assert(vec[i] == (i * 2 + 1) * 2 + 1);
-}
-
-template <typename T>
-void testAccCopyImpl(T &buf1, T &buf2, T &buf3, bool useReverse) {
-  sycl::queue q;
-  q.submit([&](sycl::handler &cgh) {
-    auto accProd = buf1.template get_access<sycl::access::mode::read>(cgh);
-    auto accCons = buf2.template get_access<sycl::access::mode::write>(cgh);
-    cgh.single_task([=]() { copyAccBuf(accProd, accCons, useReverse); });
-  });
-  q.wait();
-
-  auto accProd = buf2.template get_access<sycl::access::mode::read>();
-  auto accCons = buf3.template get_access<sycl::access::mode::write>();
-  copyAccBuf(accProd, accCons, useReverse);
-}
-
-void testAccCopy(std::vector<int> &vec1, std::vector<int> &vec2,
-                 std::vector<int> &vec3, bool useReverse = false, int Dim = 1) {
-  try {
-    if (Dim == 1) {
-      sycl::buffer<int> bufProd(vec1.data(), vec1.size());
-      sycl::buffer<int> bufCons1(vec2.data(), vec2.size());
-      sycl::buffer<int> bufCons2(vec3.data(), vec3.size());
-      testAccCopyImpl(bufProd, bufCons1, bufCons2, useReverse);
-    } else if (Dim == 3) {
-      sycl::buffer<int, 3> bufProd(vec1.data(), sycl::range<3>(2, 3, 4));
-      sycl::buffer<int, 3> bufCons1(vec2.data(), sycl::range<3>(2, 3, 4));
-      sycl::buffer<int, 3> bufCons2(vec3.data(), sycl::range<3>(2, 3, 4));
-      testAccCopyImpl(bufProd, bufCons1, bufCons2, useReverse);
-    }
+    sycl::queue queue;
+    sycl::buffer<int, 1> buf(vec.data(), vec.size());
+    queue.submit([&](sycl::handler &cgh) {
+      sycl::local_accessor<int, 1> locAcc(32, cgh);
+      auto globAcc = buf.get_access<sycl::access::mode::read_write>(cgh);
+      if (testConstIter) {
+        cgh.single_task([=]() {
+          for (int i = 0; i < locAcc.size(); ++i)
+            locAcc[i] = globAcc[i];
+          size_t Idx = 0;
+          for (auto ItLoc = locAcc.cbegin(); ItLoc != locAcc.cend(); ItLoc++)
+            globAcc[Idx++] = *ItLoc * 2 + 1;
+          Idx = locAcc.size() - 1;
+          for (auto ItLoc = locAcc.crbegin(); ItLoc != locAcc.crend(); ItLoc++)
+            globAcc[Idx--] += *ItLoc;
+        });
+      } else {
+        cgh.single_task([=]() {
+        size_t Idx = 0;
+        for (auto ItLoc = locAcc.begin(); ItLoc != locAcc.end(); ItLoc++)
+          *ItLoc = globAcc[Idx++] * 2;
+        for (auto &ItLoc : locAcc)
+          ItLoc++;
+        for (auto ItLoc = locAcc.rbegin(); ItLoc != locAcc.rend(); ItLoc++) {
+          *ItLoc *= 2;
+          *ItLoc += 1;
+        }
+        Idx = 0;
+        for (auto &ItLoc : locAcc)
+          globAcc[Idx++] = ItLoc;
+        });
+      }
+    });
   } catch (sycl::exception &e) {
     std::cout << e.what() << std::endl;
   }
-
-  for (const auto &i : vec3)
-    assert(i == 1);
 }
 
 int main() {
@@ -886,71 +827,19 @@ int main() {
     }
   }
 
-  // Accessor begin/end member function
+  // Test local_accessor begin(), end(), range_based loop, rbegin() & rend()
   {
-    // 0-dim accessor iteration
-    {
-      std::vector<int> vec(1, 0);
-      testAccMod(vec);
-    }
-
-    // Simple iteration through the accessor
-    {
-      std::vector<int> vec(32);
-      fillVec(vec);
-      testAccMod(vec);
-    }
-
-    // Const iterator
-    {
-      constexpr int N = 32;
-      std::vector<int> vecProd(N, 1), vecCons1(N, 0), vecCons2(N, 0);
-      testAccCopy(vecProd, vecCons1, vecCons2);
-    }
-
-    // Reverse iterator
-    {
-      std::vector<int> vec(32);
-      fillVec(vec);
-      testAccMod(vec, true);
-    }
-
-    // Const reverse iterator
-    {
-      constexpr int N = 32;
-      std::vector<int> vecProd(N, 1), vecCons1(N, 0), vecCons2(N, 0);
-      testAccCopy(vecProd, vecCons1, vecCons2, true);
-    }
-
-    // 3-dim accessor simple iteration
-    {
-      constexpr int N = 24;
-      std::vector<int> vec(N);
-      fillVec(vec);
-      testAccMod(vec, false, 3);
-    }
-
-    // 3-dim accessor const iterator
-    {
-      constexpr int N = 24;
-      std::vector<int> vecProd(N, 1), vecCons1(N, 0), vecCons2(N, 0);
-      testAccCopy(vecProd, vecCons1, vecCons2, false, 3);
-    }
-
-    // 3-dim accessor reverse iterator
-    {
-      constexpr int N = 24;
-      std::vector<int> vec(N);
-      fillVec(vec);
-      testAccMod(vec, true, 3);
-    }
-
-    // 3-dim accessor const reverse iterator
-    {
-      constexpr int N = 24;
-      std::vector<int> vecProd(N, 1), vecCons1(N, 0), vecCons2(N, 0);
-      testAccCopy(vecProd, vecCons1, vecCons2, true, 3);
-    }
+    std::vector<int> v(32);
+    for (int i = 0; i < v.size(); ++i) { v[i] = i; }
+    testLocalAccIters(v);
+    for (int i = 0; i < v.size(); ++i) { assert(v[i] == (i * 2 + 1) * 2 + 1); }
+  }
+  // Test local_accessor cbegin(), cend(), crbegin(), crend()
+  {
+    std::vector<int> v(32);
+    for (int i = 0; i < v.size(); ++i) { v[i] = i; }
+    testLocalAccIters(v, true);
+    for (int i = 0; i < v.size(); ++i) { assert(v[i] == (i * 2 + 1) + i); }
   }
 
   std::cout << "Test passed" << std::endl;