From 309ca0d7d1ded83c82b59c480c966c486d5c46d1 Mon Sep 17 00:00:00 2001
From: KornevNikita <nikita.kornev@intel.com>
Date: Thu, 1 Sep 2022 04:54:36 -0700
Subject: [PATCH 1/6] [SYCL] Add accessor iterator tests

---
 SYCL/Basic/accessor/accessor.cpp | 287 +++++++++++++++++++++++++++++++
 1 file changed, 287 insertions(+)

diff --git a/SYCL/Basic/accessor/accessor.cpp b/SYCL/Basic/accessor/accessor.cpp
index 5ca49f0067..d18e768b5a 100644
--- a/SYCL/Basic/accessor/accessor.cpp
+++ b/SYCL/Basic/accessor/accessor.cpp
@@ -787,5 +787,292 @@ int main() {
     }
   }
 
+  // Accessor begin/end member functions
+  {
+    // 0-dim accessor iteration
+    {
+      int data = 8;
+      try {
+        sycl::buffer<int> buf(&data, 1);
+
+        sycl::queue q;
+        q.submit([&](sycl::handler &cgh) {
+          auto acc = buf.get_access<sycl::access::mode::read_write>(cgh);
+          cgh.single_task([=]() {
+            for (auto It = acc.begin(); It != acc.end(); It++)
+              *It *= 2;
+            for (auto &It : acc)
+              It *= 2;
+          });
+        });
+        q.wait();
+      } catch (sycl::exception &e) {
+        std::cout << e.what() << std::endl;
+      }
+
+      assert(data == 32);
+    }
+
+    // Simple iteration through the accessor
+    {
+      constexpr int N = 32;
+      std::vector<int> vec(N, 0);
+      try {
+        sycl::buffer<int> buf(vec.data(), vec.size());
+        sycl::queue q;
+
+        q.submit([&](sycl::handler &cgh) {
+          auto acc = buf.get_access<sycl::access::mode::read_write>(cgh);
+
+          cgh.single_task([=]() {
+            for (auto It = acc.begin(); It != acc.end(); It++) {
+              (*It)++;
+            }
+            for (auto &It : acc)
+              It++;
+          });
+        });
+        q.wait();
+      } catch (sycl::exception &e) {
+        std::cout << e.what() << std::endl;
+      }
+
+      for (const auto &i : vec)
+        assert(i == 2);
+    }
+
+    // Const iterator
+    {
+      constexpr int N = 32;
+      std::vector<int> vecProd(N, 1), vecCons(N, 0);
+      try {
+        sycl::buffer<int> bufProd(vecProd.data(), vecProd.size());
+        sycl::buffer<int> bufCons(vecCons.data(), vecCons.size());
+        sycl::queue q;
+
+        q.submit([&](sycl::handler &cgh) {
+          auto accProd = bufProd.get_access<sycl::access::mode::read>(cgh);
+          auto accCons = bufCons.get_access<sycl::access::mode::write>(cgh);
+
+          cgh.single_task([=]() {
+            auto ItCons = accCons.begin();
+            for (auto ItProd = accProd.cbegin(); ItProd != accProd.cend();
+                 ItProd++)
+              *(ItCons++) += *ItProd;
+            ItCons = accCons.begin();
+            for (auto &It : accProd)
+              *(ItCons++) += It;
+          });
+        });
+        q.wait();
+      } catch (sycl::exception &e) {
+        std::cout << e.what() << std::endl;
+      }
+
+      for (const auto &i : vecCons)
+        assert(i == 2);
+    }
+
+    // Reverse iterator
+    {
+      constexpr int N = 32;
+      std::vector<int> vec(N);
+      try {
+        sycl::buffer<int> buf(vec.data(), vec.size());
+        sycl::queue q;
+
+        q.submit([&](sycl::handler &cgh) {
+          auto acc1 = buf.get_access<sycl::access::mode::write>(cgh);
+
+          cgh.single_task([=]() {
+            size_t count = N;
+            for (auto It = acc1.rbegin(); It != acc1.rend(); It++) {
+              *It = N - count;
+              count--;
+            }
+          });
+        });
+        q.wait();
+      } catch (sycl::exception &e) {
+        std::cout << e.what() << std::endl;
+      }
+
+      int count = N - 1;
+      for (const auto &i : vec) {
+        assert(i == count);
+        count--;
+      }
+    }
+
+    // Const reverse iterator
+    {
+      constexpr int N = 32;
+      std::vector<int> vecProd(N), vecCons(N);
+
+      int count = N;
+      for (auto &i : vecProd) {
+        i = N - count;
+        count--;
+      }
+
+      try {
+        sycl::buffer<int> bufProd(vecProd.data(), vecProd.size());
+        sycl::buffer<int> bufCons(vecCons.data(), vecCons.size());
+        sycl::queue q;
+
+        q.submit([&](sycl::handler &cgh) {
+          auto accProd = bufProd.get_access<sycl::access::mode::read>(cgh);
+          auto accCons = bufCons.get_access<sycl::access::mode::write>(cgh);
+
+          cgh.single_task([=]() {
+            auto ItCons = accCons.begin();
+            for (auto ItProd = accProd.crbegin(); ItProd != accProd.crend();
+                 ItProd++) {
+              *(ItCons++) = N - *ItProd - 1;
+            }
+          });
+        });
+
+        q.wait();
+      } catch (sycl::exception &e) {
+        std::cout << e.what() << std::endl;
+      }
+
+      count = 0;
+      for (const auto &i : vecCons)
+        assert(i == count++);
+    }
+
+    // 3-dim accessor simple iteration
+    {
+      constexpr int N = 24;
+      std::vector<int> vec(N, 0);
+      try {
+        sycl::buffer<int, 3> buf(vec.data(), sycl::range<3>(2, 3, 4));
+        sycl::queue q;
+
+        q.submit([&](sycl::handler &cgh) {
+          auto acc = buf.get_access<sycl::access::mode::read_write>(cgh);
+
+          cgh.single_task([=]() {
+            for (auto It = acc.begin(); It != acc.end(); It++) {
+              (*It)++;
+            }
+            for (auto &It : acc)
+              It++;
+          });
+        });
+        q.wait();
+      } catch (sycl::exception &e) {
+        std::cout << e.what() << std::endl;
+      }
+
+      for (const auto &i : vec)
+        assert(i == 2);
+    }
+
+    // 3-dim accessor const iterator
+    {
+      constexpr int N = 24;
+      std::vector<int> vecProd(N, 1), vecCons(N, 0);
+      try {
+        sycl::buffer<int, 3> bufProd(vecProd.data(), sycl::range<3>(2, 3, 4));
+        sycl::buffer<int, 3> bufCons(vecCons.data(), sycl::range<3>(2, 3, 4));
+        sycl::queue q;
+
+        q.submit([&](sycl::handler &cgh) {
+          auto accProd = bufProd.get_access<sycl::access::mode::read>(cgh);
+          auto accCons = bufCons.get_access<sycl::access::mode::write>(cgh);
+
+          cgh.single_task([=]() {
+            auto ItCons = accCons.begin();
+            for (auto ItProd = accProd.cbegin(); ItProd != accProd.cend();
+                 ItProd++)
+              *(ItCons++) += *ItProd;
+            ItCons = accCons.begin();
+            for (auto &It : accProd)
+              *(ItCons++) += It;
+          });
+        });
+        q.wait();
+      } catch (sycl::exception &e) {
+        std::cout << e.what() << std::endl;
+      }
+
+      for (const auto &i : vecCons)
+        assert(i == 2);
+    }
+
+    // 3-dim accessor reverse iterator
+    {
+      constexpr int N = 24;
+      std::vector<int> vec(N);
+      try {
+        sycl::buffer<int, 3> buf(vec.data(), sycl::range<3>(2, 3, 4));
+        sycl::queue q;
+
+        q.submit([&](sycl::handler &cgh) {
+          auto acc1 = buf.get_access<sycl::access::mode::write>(cgh);
+
+          cgh.single_task([=]() {
+            size_t count = N;
+            for (auto It = acc1.rbegin(); It != acc1.rend(); It++) {
+              *It = N - count;
+              count--;
+            }
+          });
+        });
+        q.wait();
+      } catch (sycl::exception &e) {
+        std::cout << e.what() << std::endl;
+      }
+
+      int count = N - 1;
+      for (const auto &i : vec) {
+        assert(i == count);
+        count--;
+      }
+    }
+
+    // 3-dim accessor const reverse iterator
+    {
+      constexpr int N = 24;
+      std::vector<int> vecProd(N), vecCons(N);
+
+      int count = N;
+      for (auto &i : vecProd) {
+        i = N - count;
+        count--;
+      }
+
+      try {
+        sycl::buffer<int, 3> bufProd(vecProd.data(), sycl::range<3>(2, 3, 4));
+        sycl::buffer<int, 3> bufCons(vecCons.data(), sycl::range<3>(2, 3, 4));
+        sycl::queue q;
+
+        q.submit([&](sycl::handler &cgh) {
+          auto accProd = bufProd.get_access<sycl::access::mode::read>(cgh);
+          auto accCons = bufCons.get_access<sycl::access::mode::write>(cgh);
+
+          cgh.single_task([=]() {
+            auto ItCons = accCons.begin();
+            for (auto ItProd = accProd.crbegin(); ItProd != accProd.crend();
+                 ItProd++) {
+              *(ItCons++) = N - *ItProd - 1;
+            }
+          });
+        });
+
+        q.wait();
+      } catch (sycl::exception &e) {
+        std::cout << e.what() << std::endl;
+      }
+
+      count = 0;
+      for (const auto &i : vecCons)
+        assert(i == count++);
+    }
+  }
+
   std::cout << "Test passed" << std::endl;
 }

From 328e794ae00ffe9b94e197c691e9eef1a1f03699 Mon Sep 17 00:00:00 2001
From: KornevNikita <nikita.kornev@intel.com>
Date: Mon, 5 Sep 2022 08:32:04 -0700
Subject: [PATCH 2/6] Apply suggestions

---
 SYCL/Basic/accessor/accessor.cpp | 254 +++++++++++++++----------------
 1 file changed, 122 insertions(+), 132 deletions(-)

diff --git a/SYCL/Basic/accessor/accessor.cpp b/SYCL/Basic/accessor/accessor.cpp
index d18e768b5a..4ae429a977 100644
--- a/SYCL/Basic/accessor/accessor.cpp
+++ b/SYCL/Basic/accessor/accessor.cpp
@@ -61,6 +61,39 @@ 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 = false) {
+  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 = false) {
+  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;
+  }
+}
+
 int main() {
   // Host accessor.
   {
@@ -794,283 +827,240 @@ int main() {
       int data = 8;
       try {
         sycl::buffer<int> buf(&data, 1);
-
         sycl::queue q;
+
         q.submit([&](sycl::handler &cgh) {
-          auto acc = buf.get_access<sycl::access::mode::read_write>(cgh);
-          cgh.single_task([=]() {
-            for (auto It = acc.begin(); It != acc.end(); It++)
-              *It *= 2;
-            for (auto &It : acc)
-              It *= 2;
-          });
+          auto devAcc = buf.get_access<sycl::access::mode::read_write>(cgh);
+          cgh.single_task([=]() { modifyAccBuf(devAcc); });
         });
         q.wait();
+
+        auto hostAcc = buf.get_access<sycl::access::mode::read_write>();
+        modifyAccBuf(hostAcc);
       } catch (sycl::exception &e) {
         std::cout << e.what() << std::endl;
       }
 
-      assert(data == 32);
+      assert(data == (8 * 2 + 1) * 2 + 1);
     }
 
     // Simple iteration through the accessor
     {
       constexpr int N = 32;
-      std::vector<int> vec(N, 0);
+      std::vector<int> vec(N);
+      fillVec(vec);
+
       try {
         sycl::buffer<int> buf(vec.data(), vec.size());
         sycl::queue q;
 
         q.submit([&](sycl::handler &cgh) {
-          auto acc = buf.get_access<sycl::access::mode::read_write>(cgh);
-
-          cgh.single_task([=]() {
-            for (auto It = acc.begin(); It != acc.end(); It++) {
-              (*It)++;
-            }
-            for (auto &It : acc)
-              It++;
-          });
+          auto devAcc = buf.get_access<sycl::access::mode::read_write>(cgh);
+          cgh.single_task([=]() { modifyAccBuf(devAcc); });
         });
         q.wait();
+
+        auto hostAcc = buf.get_access<sycl::access::mode::read_write>();
+        modifyAccBuf(hostAcc);
       } catch (sycl::exception &e) {
         std::cout << e.what() << std::endl;
       }
 
-      for (const auto &i : vec)
-        assert(i == 2);
+      for (size_t i = 0; i < vec.size(); ++i)
+        assert(vec[i] == (i * 2 + 1) * 2 + 1);
     }
 
     // Const iterator
     {
       constexpr int N = 32;
-      std::vector<int> vecProd(N, 1), vecCons(N, 0);
+      std::vector<int> vecProd(N, 1), vecCons1(N, 0), vecCons2(N, 0);
       try {
         sycl::buffer<int> bufProd(vecProd.data(), vecProd.size());
-        sycl::buffer<int> bufCons(vecCons.data(), vecCons.size());
+        sycl::buffer<int> bufCons1(vecCons1.data(), vecCons1.size());
+        sycl::buffer<int> bufCons2(vecCons2.data(), vecCons2.size());
         sycl::queue q;
 
         q.submit([&](sycl::handler &cgh) {
           auto accProd = bufProd.get_access<sycl::access::mode::read>(cgh);
-          auto accCons = bufCons.get_access<sycl::access::mode::write>(cgh);
-
-          cgh.single_task([=]() {
-            auto ItCons = accCons.begin();
-            for (auto ItProd = accProd.cbegin(); ItProd != accProd.cend();
-                 ItProd++)
-              *(ItCons++) += *ItProd;
-            ItCons = accCons.begin();
-            for (auto &It : accProd)
-              *(ItCons++) += It;
-          });
+          auto accCons = bufCons1.get_access<sycl::access::mode::write>(cgh);
+          cgh.single_task([=]() { copyAccBuf(accProd, accCons); });
         });
         q.wait();
+
+        auto accProd = bufCons1.get_access<sycl::access::mode::read>();
+        auto accCons = bufCons2.get_access<sycl::access::mode::write>();
+        copyAccBuf(accProd, accCons);
       } catch (sycl::exception &e) {
         std::cout << e.what() << std::endl;
       }
 
-      for (const auto &i : vecCons)
-        assert(i == 2);
+      for (const auto &i : vecCons2)
+        assert(i == 1);
     }
 
     // Reverse iterator
     {
       constexpr int N = 32;
       std::vector<int> vec(N);
+      fillVec(vec);
+
       try {
         sycl::buffer<int> buf(vec.data(), vec.size());
         sycl::queue q;
 
         q.submit([&](sycl::handler &cgh) {
-          auto acc1 = buf.get_access<sycl::access::mode::write>(cgh);
-
-          cgh.single_task([=]() {
-            size_t count = N;
-            for (auto It = acc1.rbegin(); It != acc1.rend(); It++) {
-              *It = N - count;
-              count--;
-            }
-          });
+          auto acc = buf.get_access<sycl::access::mode::read_write>(cgh);
+          cgh.single_task([=]() { modifyAccBuf(acc, true); });
         });
         q.wait();
+
+        auto acc = buf.get_access<sycl::access::mode::read_write>();
+        modifyAccBuf(acc, true);
       } catch (sycl::exception &e) {
         std::cout << e.what() << std::endl;
       }
 
-      int count = N - 1;
-      for (const auto &i : vec) {
-        assert(i == count);
-        count--;
-      }
+      for (size_t i = 0; i < vec.size(); ++i)
+        assert(vec[i] == (i * 2 + 1) * 2 + 1);
     }
 
     // Const reverse iterator
     {
       constexpr int N = 32;
-      std::vector<int> vecProd(N), vecCons(N);
-
-      int count = N;
-      for (auto &i : vecProd) {
-        i = N - count;
-        count--;
-      }
+      std::vector<int> vecProd(N, 1), vecCons1(N, 0), vecCons2(N, 0);
 
       try {
         sycl::buffer<int> bufProd(vecProd.data(), vecProd.size());
-        sycl::buffer<int> bufCons(vecCons.data(), vecCons.size());
+        sycl::buffer<int> bufCons1(vecCons1.data(), vecCons1.size());
+        sycl::buffer<int> bufCons2(vecCons2.data(), vecCons2.size());
         sycl::queue q;
 
         q.submit([&](sycl::handler &cgh) {
           auto accProd = bufProd.get_access<sycl::access::mode::read>(cgh);
-          auto accCons = bufCons.get_access<sycl::access::mode::write>(cgh);
-
-          cgh.single_task([=]() {
-            auto ItCons = accCons.begin();
-            for (auto ItProd = accProd.crbegin(); ItProd != accProd.crend();
-                 ItProd++) {
-              *(ItCons++) = N - *ItProd - 1;
-            }
-          });
+          auto accCons1 = bufCons1.get_access<sycl::access::mode::write>(cgh);
+          cgh.single_task([=]() { copyAccBuf(accProd, accCons1, true); });
         });
-
         q.wait();
+
+        auto accCons1 = bufCons1.get_access<sycl::access::mode::read>();
+        auto accCons2 = bufCons2.get_access<sycl::access::mode::write>();
+        copyAccBuf(accCons1, accCons2, true);
       } catch (sycl::exception &e) {
         std::cout << e.what() << std::endl;
       }
 
-      count = 0;
-      for (const auto &i : vecCons)
-        assert(i == count++);
+      for (const auto &i : vecCons2)
+        assert(i == 1);
     }
 
     // 3-dim accessor simple iteration
     {
       constexpr int N = 24;
-      std::vector<int> vec(N, 0);
+      std::vector<int> vec(N);
+      fillVec(vec);
+
       try {
         sycl::buffer<int, 3> buf(vec.data(), sycl::range<3>(2, 3, 4));
         sycl::queue q;
 
         q.submit([&](sycl::handler &cgh) {
           auto acc = buf.get_access<sycl::access::mode::read_write>(cgh);
-
-          cgh.single_task([=]() {
-            for (auto It = acc.begin(); It != acc.end(); It++) {
-              (*It)++;
-            }
-            for (auto &It : acc)
-              It++;
-          });
+          cgh.single_task([=]() { modifyAccBuf(acc); });
         });
         q.wait();
+
+        auto acc = buf.get_access<sycl::access::mode::read_write>();
+        modifyAccBuf(acc);
       } catch (sycl::exception &e) {
         std::cout << e.what() << std::endl;
       }
 
-      for (const auto &i : vec)
-        assert(i == 2);
+      for (size_t i = 0; i < vec.size(); ++i)
+        assert(vec[i] == (i * 2 + 1) * 2 + 1);
     }
 
     // 3-dim accessor const iterator
     {
       constexpr int N = 24;
-      std::vector<int> vecProd(N, 1), vecCons(N, 0);
+      std::vector<int> vecProd(N, 1), vecCons1(N, 0), vecCons2(N, 0);
+
       try {
         sycl::buffer<int, 3> bufProd(vecProd.data(), sycl::range<3>(2, 3, 4));
-        sycl::buffer<int, 3> bufCons(vecCons.data(), sycl::range<3>(2, 3, 4));
+        sycl::buffer<int, 3> bufCons1(vecCons1.data(), sycl::range<3>(2, 3, 4));
+        sycl::buffer<int, 3> bufCons2(vecCons2.data(), sycl::range<3>(2, 3, 4));
         sycl::queue q;
 
         q.submit([&](sycl::handler &cgh) {
           auto accProd = bufProd.get_access<sycl::access::mode::read>(cgh);
-          auto accCons = bufCons.get_access<sycl::access::mode::write>(cgh);
-
-          cgh.single_task([=]() {
-            auto ItCons = accCons.begin();
-            for (auto ItProd = accProd.cbegin(); ItProd != accProd.cend();
-                 ItProd++)
-              *(ItCons++) += *ItProd;
-            ItCons = accCons.begin();
-            for (auto &It : accProd)
-              *(ItCons++) += It;
-          });
+          auto accCons = bufCons1.get_access<sycl::access::mode::write>(cgh);
+          cgh.single_task([=]() { copyAccBuf(accProd, accCons); });
         });
         q.wait();
+
+        auto accProd = bufCons1.get_access<sycl::access::mode::read>();
+        auto accCons = bufCons2.get_access<sycl::access::mode::write>();
+        copyAccBuf(accProd, accCons);
       } catch (sycl::exception &e) {
         std::cout << e.what() << std::endl;
       }
 
-      for (const auto &i : vecCons)
-        assert(i == 2);
+      for (const auto &i : vecCons2)
+        assert(i == 1);
     }
 
     // 3-dim accessor reverse iterator
     {
       constexpr int N = 24;
       std::vector<int> vec(N);
+      fillVec(vec);
+
       try {
         sycl::buffer<int, 3> buf(vec.data(), sycl::range<3>(2, 3, 4));
         sycl::queue q;
 
         q.submit([&](sycl::handler &cgh) {
-          auto acc1 = buf.get_access<sycl::access::mode::write>(cgh);
-
-          cgh.single_task([=]() {
-            size_t count = N;
-            for (auto It = acc1.rbegin(); It != acc1.rend(); It++) {
-              *It = N - count;
-              count--;
-            }
-          });
+          auto acc = buf.get_access<sycl::access::mode::write>(cgh);
+          cgh.single_task([=]() { modifyAccBuf(acc, true); });
         });
         q.wait();
+
+        auto acc = buf.get_access<sycl::access::mode::write>();
+        modifyAccBuf(acc, true);
       } catch (sycl::exception &e) {
         std::cout << e.what() << std::endl;
       }
 
-      int count = N - 1;
-      for (const auto &i : vec) {
-        assert(i == count);
-        count--;
-      }
+      for (size_t i = 0; i < vec.size(); ++i)
+        assert(vec[i] == (i * 2 + 1) * 2 + 1);
     }
 
     // 3-dim accessor const reverse iterator
     {
       constexpr int N = 24;
-      std::vector<int> vecProd(N), vecCons(N);
-
-      int count = N;
-      for (auto &i : vecProd) {
-        i = N - count;
-        count--;
-      }
+      std::vector<int> vecProd(N, 1), vecCons1(N, 0), vecCons2(N, 0);
 
       try {
         sycl::buffer<int, 3> bufProd(vecProd.data(), sycl::range<3>(2, 3, 4));
-        sycl::buffer<int, 3> bufCons(vecCons.data(), sycl::range<3>(2, 3, 4));
+        sycl::buffer<int, 3> bufCons1(vecCons1.data(), sycl::range<3>(2, 3, 4));
+        sycl::buffer<int, 3> bufCons2(vecCons2.data(), sycl::range<3>(2, 3, 4));
         sycl::queue q;
 
         q.submit([&](sycl::handler &cgh) {
           auto accProd = bufProd.get_access<sycl::access::mode::read>(cgh);
-          auto accCons = bufCons.get_access<sycl::access::mode::write>(cgh);
-
-          cgh.single_task([=]() {
-            auto ItCons = accCons.begin();
-            for (auto ItProd = accProd.crbegin(); ItProd != accProd.crend();
-                 ItProd++) {
-              *(ItCons++) = N - *ItProd - 1;
-            }
-          });
+          auto accCons = bufCons1.get_access<sycl::access::mode::write>(cgh);
+          cgh.single_task([=]() { copyAccBuf(accProd, accCons, true); });
         });
-
         q.wait();
+
+        auto accProd = bufCons1.get_access<sycl::access::mode::read>();
+        auto accCons = bufCons2.get_access<sycl::access::mode::write>();
+        copyAccBuf(accProd, accCons, true);
       } catch (sycl::exception &e) {
         std::cout << e.what() << std::endl;
       }
 
-      count = 0;
-      for (const auto &i : vecCons)
-        assert(i == count++);
+      for (const auto &i : vecCons2)
+        assert(i == 1);
     }
   }
 

From 2a0f4b6d00ce873c67a1b0df001927f1fba4666a Mon Sep 17 00:00:00 2001
From: KornevNikita <nikita.kornev@intel.com>
Date: Tue, 6 Sep 2022 04:54:21 -0700
Subject: [PATCH 3/6] Apply suggestions

---
 SYCL/Basic/accessor/accessor.cpp | 273 +++++++++----------------------
 1 file changed, 81 insertions(+), 192 deletions(-)

diff --git a/SYCL/Basic/accessor/accessor.cpp b/SYCL/Basic/accessor/accessor.cpp
index 4ae429a977..28eeb3730e 100644
--- a/SYCL/Basic/accessor/accessor.cpp
+++ b/SYCL/Basic/accessor/accessor.cpp
@@ -67,7 +67,7 @@ void fillVec(std::vector<int> &vec) {
   }
 }
 
-template <typename T> void modifyAccBuf(const T &acc, bool useReverse = false) {
+template <typename T> void modifyAccBuf(const T &acc, bool useReverse) {
   if (!useReverse) {
     for (auto It = acc.begin(); It != acc.end(); It++)
       *It *= 2;
@@ -82,7 +82,7 @@ template <typename T> void modifyAccBuf(const T &acc, bool useReverse = false) {
 }
 
 template <typename T1, typename T2>
-void copyAccBuf(const T1 &accProd, T2 &accCons, bool useReverse = false) {
+void copyAccBuf(const T1 &accProd, T2 &accCons, bool useReverse) {
   if (!useReverse) {
     auto ItCons = accCons.begin();
     for (auto ItProd = accProd.cbegin(); ItProd != accProd.cend(); ++ItProd)
@@ -94,6 +94,72 @@ void copyAccBuf(const T1 &accProd, T2 &accCons, bool useReverse = false) {
   }
 }
 
+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) {
+  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);
+    }
+  } catch (sycl::exception &e) {
+    std::cout << e.what() << std::endl;
+  }
+
+  for (const auto &i : vec3)
+    assert(i == 1);
+}
+
 int main() {
   // Host accessor.
   {
@@ -820,137 +886,40 @@ int main() {
     }
   }
 
-  // Accessor begin/end member functions
+  // Accessor begin/end member function
   {
     // 0-dim accessor iteration
     {
-      int data = 8;
-      try {
-        sycl::buffer<int> buf(&data, 1);
-        sycl::queue q;
-
-        q.submit([&](sycl::handler &cgh) {
-          auto devAcc = buf.get_access<sycl::access::mode::read_write>(cgh);
-          cgh.single_task([=]() { modifyAccBuf(devAcc); });
-        });
-        q.wait();
-
-        auto hostAcc = buf.get_access<sycl::access::mode::read_write>();
-        modifyAccBuf(hostAcc);
-      } catch (sycl::exception &e) {
-        std::cout << e.what() << std::endl;
-      }
-
-      assert(data == (8 * 2 + 1) * 2 + 1);
+      std::vector<int> vec(1, 0);
+      testAccMod(vec);
     }
 
     // Simple iteration through the accessor
     {
-      constexpr int N = 32;
-      std::vector<int> vec(N);
+      std::vector<int> vec(32);
       fillVec(vec);
-
-      try {
-        sycl::buffer<int> buf(vec.data(), vec.size());
-        sycl::queue q;
-
-        q.submit([&](sycl::handler &cgh) {
-          auto devAcc = buf.get_access<sycl::access::mode::read_write>(cgh);
-          cgh.single_task([=]() { modifyAccBuf(devAcc); });
-        });
-        q.wait();
-
-        auto hostAcc = buf.get_access<sycl::access::mode::read_write>();
-        modifyAccBuf(hostAcc);
-      } 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);
+      testAccMod(vec);
     }
 
     // Const iterator
     {
       constexpr int N = 32;
       std::vector<int> vecProd(N, 1), vecCons1(N, 0), vecCons2(N, 0);
-      try {
-        sycl::buffer<int> bufProd(vecProd.data(), vecProd.size());
-        sycl::buffer<int> bufCons1(vecCons1.data(), vecCons1.size());
-        sycl::buffer<int> bufCons2(vecCons2.data(), vecCons2.size());
-        sycl::queue q;
-
-        q.submit([&](sycl::handler &cgh) {
-          auto accProd = bufProd.get_access<sycl::access::mode::read>(cgh);
-          auto accCons = bufCons1.get_access<sycl::access::mode::write>(cgh);
-          cgh.single_task([=]() { copyAccBuf(accProd, accCons); });
-        });
-        q.wait();
-
-        auto accProd = bufCons1.get_access<sycl::access::mode::read>();
-        auto accCons = bufCons2.get_access<sycl::access::mode::write>();
-        copyAccBuf(accProd, accCons);
-      } catch (sycl::exception &e) {
-        std::cout << e.what() << std::endl;
-      }
-
-      for (const auto &i : vecCons2)
-        assert(i == 1);
+      testAccCopy(vecProd, vecCons1, vecCons2);
     }
 
     // Reverse iterator
     {
-      constexpr int N = 32;
-      std::vector<int> vec(N);
+      std::vector<int> vec(32);
       fillVec(vec);
-
-      try {
-        sycl::buffer<int> buf(vec.data(), vec.size());
-        sycl::queue q;
-
-        q.submit([&](sycl::handler &cgh) {
-          auto acc = buf.get_access<sycl::access::mode::read_write>(cgh);
-          cgh.single_task([=]() { modifyAccBuf(acc, true); });
-        });
-        q.wait();
-
-        auto acc = buf.get_access<sycl::access::mode::read_write>();
-        modifyAccBuf(acc, true);
-      } 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);
+      testAccMod(vec, true);
     }
 
     // Const reverse iterator
     {
       constexpr int N = 32;
       std::vector<int> vecProd(N, 1), vecCons1(N, 0), vecCons2(N, 0);
-
-      try {
-        sycl::buffer<int> bufProd(vecProd.data(), vecProd.size());
-        sycl::buffer<int> bufCons1(vecCons1.data(), vecCons1.size());
-        sycl::buffer<int> bufCons2(vecCons2.data(), vecCons2.size());
-        sycl::queue q;
-
-        q.submit([&](sycl::handler &cgh) {
-          auto accProd = bufProd.get_access<sycl::access::mode::read>(cgh);
-          auto accCons1 = bufCons1.get_access<sycl::access::mode::write>(cgh);
-          cgh.single_task([=]() { copyAccBuf(accProd, accCons1, true); });
-        });
-        q.wait();
-
-        auto accCons1 = bufCons1.get_access<sycl::access::mode::read>();
-        auto accCons2 = bufCons2.get_access<sycl::access::mode::write>();
-        copyAccBuf(accCons1, accCons2, true);
-      } catch (sycl::exception &e) {
-        std::cout << e.what() << std::endl;
-      }
-
-      for (const auto &i : vecCons2)
-        assert(i == 1);
+      testAccCopy(vecProd, vecCons1, vecCons2, true);
     }
 
     // 3-dim accessor simple iteration
@@ -958,54 +927,14 @@ int main() {
       constexpr int N = 24;
       std::vector<int> vec(N);
       fillVec(vec);
-
-      try {
-        sycl::buffer<int, 3> buf(vec.data(), sycl::range<3>(2, 3, 4));
-        sycl::queue q;
-
-        q.submit([&](sycl::handler &cgh) {
-          auto acc = buf.get_access<sycl::access::mode::read_write>(cgh);
-          cgh.single_task([=]() { modifyAccBuf(acc); });
-        });
-        q.wait();
-
-        auto acc = buf.get_access<sycl::access::mode::read_write>();
-        modifyAccBuf(acc);
-      } 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);
+      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);
-
-      try {
-        sycl::buffer<int, 3> bufProd(vecProd.data(), sycl::range<3>(2, 3, 4));
-        sycl::buffer<int, 3> bufCons1(vecCons1.data(), sycl::range<3>(2, 3, 4));
-        sycl::buffer<int, 3> bufCons2(vecCons2.data(), sycl::range<3>(2, 3, 4));
-        sycl::queue q;
-
-        q.submit([&](sycl::handler &cgh) {
-          auto accProd = bufProd.get_access<sycl::access::mode::read>(cgh);
-          auto accCons = bufCons1.get_access<sycl::access::mode::write>(cgh);
-          cgh.single_task([=]() { copyAccBuf(accProd, accCons); });
-        });
-        q.wait();
-
-        auto accProd = bufCons1.get_access<sycl::access::mode::read>();
-        auto accCons = bufCons2.get_access<sycl::access::mode::write>();
-        copyAccBuf(accProd, accCons);
-      } catch (sycl::exception &e) {
-        std::cout << e.what() << std::endl;
-      }
-
-      for (const auto &i : vecCons2)
-        assert(i == 1);
+      testAccCopy(vecProd, vecCons1, vecCons2, false, 3);
     }
 
     // 3-dim accessor reverse iterator
@@ -1013,54 +942,14 @@ int main() {
       constexpr int N = 24;
       std::vector<int> vec(N);
       fillVec(vec);
-
-      try {
-        sycl::buffer<int, 3> buf(vec.data(), sycl::range<3>(2, 3, 4));
-        sycl::queue q;
-
-        q.submit([&](sycl::handler &cgh) {
-          auto acc = buf.get_access<sycl::access::mode::write>(cgh);
-          cgh.single_task([=]() { modifyAccBuf(acc, true); });
-        });
-        q.wait();
-
-        auto acc = buf.get_access<sycl::access::mode::write>();
-        modifyAccBuf(acc, true);
-      } 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);
+      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);
-
-      try {
-        sycl::buffer<int, 3> bufProd(vecProd.data(), sycl::range<3>(2, 3, 4));
-        sycl::buffer<int, 3> bufCons1(vecCons1.data(), sycl::range<3>(2, 3, 4));
-        sycl::buffer<int, 3> bufCons2(vecCons2.data(), sycl::range<3>(2, 3, 4));
-        sycl::queue q;
-
-        q.submit([&](sycl::handler &cgh) {
-          auto accProd = bufProd.get_access<sycl::access::mode::read>(cgh);
-          auto accCons = bufCons1.get_access<sycl::access::mode::write>(cgh);
-          cgh.single_task([=]() { copyAccBuf(accProd, accCons, true); });
-        });
-        q.wait();
-
-        auto accProd = bufCons1.get_access<sycl::access::mode::read>();
-        auto accCons = bufCons2.get_access<sycl::access::mode::write>();
-        copyAccBuf(accProd, accCons, true);
-      } catch (sycl::exception &e) {
-        std::cout << e.what() << std::endl;
-      }
-
-      for (const auto &i : vecCons2)
-        assert(i == 1);
+      testAccCopy(vecProd, vecCons1, vecCons2, true, 3);
     }
   }
 

From 8a7bbc2bef263a2c308b5efcb5c58056a7218e9c Mon Sep 17 00:00:00 2001
From: KornevNikita <nikita.kornev@intel.com>
Date: Mon, 19 Sep 2022 09:34:38 -0700
Subject: [PATCH 4/6] Remove accessor & add local_accessor

---
 SYCL/Basic/accessor/accessor.cpp | 203 +++++++------------------------
 1 file changed, 46 insertions(+), 157 deletions(-)

diff --git a/SYCL/Basic/accessor/accessor.cpp b/SYCL/Basic/accessor/accessor.cpp
index 28eeb3730e..86f54f00b7 100644
--- a/SYCL/Basic/accessor/accessor.cpp
+++ b/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;

From b3fde259c14afe93766d7ea9e7e13867b67c2869 Mon Sep 17 00:00:00 2001
From: KornevNikita <nikita.kornev@intel.com>
Date: Wed, 21 Sep 2022 05:18:46 -0700
Subject: [PATCH 5/6] Update to test 2D local_accessor

---
 SYCL/Basic/accessor/accessor.cpp | 96 +++++++++++++++++++++-----------
 1 file changed, 64 insertions(+), 32 deletions(-)

diff --git a/SYCL/Basic/accessor/accessor.cpp b/SYCL/Basic/accessor/accessor.cpp
index 86f54f00b7..888357a70c 100644
--- a/SYCL/Basic/accessor/accessor.cpp
+++ b/SYCL/Basic/accessor/accessor.cpp
@@ -61,39 +61,55 @@ template <typename Acc> struct Wrapper2 {
 
 template <typename Acc> struct Wrapper3 { Wrapper2<Acc> w2; };
 
-void testLocalAccIters(std::vector<int> &vec, bool testConstIter = false) {
+template <typename GlobAcc, typename LocAcc>
+void testLocalAccItersImpl(sycl::handler &cgh, GlobAcc &globAcc, LocAcc &locAcc,
+                           bool testConstIter) {
+  if (testConstIter) {
+    cgh.single_task([=]() {
+      size_t Idx = 0;
+      for (auto &It : locAcc) {
+        It = globAcc[Idx++];
+      }
+      Idx = 0;
+      for (auto It = locAcc.cbegin(); It != locAcc.cend(); It++)
+        globAcc[Idx++] = *It * 2 + 1;
+      Idx = locAcc.size() - 1;
+      for (auto It = locAcc.crbegin(); It != locAcc.crend(); It++)
+        globAcc[Idx--] += *It;
+    });
+  } else {
+    cgh.single_task([=]() {
+      size_t Idx = 0;
+      for (auto It = locAcc.begin(); It != locAcc.end(); It++)
+        *It = globAcc[Idx++] * 2;
+      for (auto &It : locAcc) {
+        It++;
+      }
+      for (auto It = locAcc.rbegin(); It != locAcc.rend(); It++) {
+        *It *= 2;
+        *It += 1;
+      }
+      Idx = 0;
+      for (auto &It : locAcc) {
+        globAcc[Idx++] = It;
+      }
+    });
+  }
+}
+
+void testLocalAccIters(std::vector<int> &vec, bool testConstIter = false,
+                       bool test2D = false) {
   try {
     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;
-        });
+      if (test2D) {
+        sycl::local_accessor<int, 2> locAcc(sycl::range<2>{2, 16}, cgh);
+        testLocalAccItersImpl(cgh, globAcc, locAcc, testConstIter);
       } 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;
-        });
+        sycl::local_accessor<int, 1> locAcc(32, cgh);
+        testLocalAccItersImpl(cgh, globAcc, locAcc, testConstIter);
       }
     });
   } catch (sycl::exception &e) {
@@ -827,19 +843,35 @@ int main() {
     }
   }
 
-  // Test local_accessor begin(), end(), range_based loop, rbegin() & rend()
+  // Test iterator methods with 1D local_accessor
   {
     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); }
+    for (int i = 0; i < v.size(); ++i)
+      assert(v[i] == ((i * 2 + 1) * 2 + 1));
+
+    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));
   }
-  // Test local_accessor cbegin(), cend(), crbegin(), crend()
+  // Test iterator methods with 2D local_accessor
   {
     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); }
+    testLocalAccIters(v, false, true);
+    for (int i = 0; i < v.size(); ++i)
+      assert(v[i] == ((i * 2 + 1) * 2 + 1));
+
+    for (int i = 0; i < v.size(); ++i) {
+      v[i] = i;
+    }
+    testLocalAccIters(v, true, true);
+    for (int i = 0; i < v.size(); ++i)
+      assert(v[i] == ((i * 2 + 1) + i));
   }
 
   std::cout << "Test passed" << std::endl;

From 909f4980bfb41681b5a2ed5f00c529cbdd25f03b Mon Sep 17 00:00:00 2001
From: KornevNikita <nikita.kornev@intel.com>
Date: Wed, 21 Sep 2022 05:33:55 -0700
Subject: [PATCH 6/6] Fix clang-format

---
 SYCL/Basic/accessor/accessor.cpp | 8 ++++++--
 1 file changed, 6 insertions(+), 2 deletions(-)

diff --git a/SYCL/Basic/accessor/accessor.cpp b/SYCL/Basic/accessor/accessor.cpp
index 888357a70c..b9231519cb 100644
--- a/SYCL/Basic/accessor/accessor.cpp
+++ b/SYCL/Basic/accessor/accessor.cpp
@@ -846,7 +846,9 @@ int main() {
   // Test iterator methods with 1D local_accessor
   {
     std::vector<int> v(32);
-    for (int i = 0; i < v.size(); ++i) { v[i] = i; }
+    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));
@@ -861,7 +863,9 @@ int main() {
   // Test iterator methods with 2D local_accessor
   {
     std::vector<int> v(32);
-    for (int i = 0; i < v.size(); ++i) { v[i] = i; }
+    for (int i = 0; i < v.size(); ++i) {
+      v[i] = i;
+    }
     testLocalAccIters(v, false, true);
     for (int i = 0; i < v.size(); ++i)
       assert(v[i] == ((i * 2 + 1) * 2 + 1));