From cf19cb98bf2fc3a105f6063c9cd289dd733e7e02 Mon Sep 17 00:00:00 2001
From: mdimakov <maxim.dimakov@intel.com>
Date: Mon, 5 Apr 2021 16:47:28 +0300
Subject: [PATCH 1/4] [SYCL] Test for multi-dim subscript operator for atomic
 accessor

Signed-off-by: mdimakov <maxim.dimakov@intel.com>
---
 SYCL/Basic/accessor/accessor.cpp | 22 ++++++++++++++++++++++
 1 file changed, 22 insertions(+)

diff --git a/SYCL/Basic/accessor/accessor.cpp b/SYCL/Basic/accessor/accessor.cpp
index e03186279a..43ea804393 100644
--- a/SYCL/Basic/accessor/accessor.cpp
+++ b/SYCL/Basic/accessor/accessor.cpp
@@ -166,6 +166,28 @@ int main() {
     }
   }
 
+  // Device accessor with 2-dimensional subscript operators for atomic accessor
+  // check compile error
+  {
+    cl::sycl::queue queue;
+    if (!queue.is_host()) {
+      cl::sycl::range<2> range(1, 1);
+      int Arr[] = {2};
+      {
+        cl::sycl::buffer<int, 1> Buf(Arr, 1);
+        queue.submit([&](cl::sycl::handler &cgh) {
+          auto acc =
+              cl::sycl::accessor<int, 2, cl::sycl::access::mode::atomic,
+              cl::sycl::access::target::local>(range, cgh);
+          cgh.parallel_for<class dim2_subscr_atomic>(range, [=] (cl::sycl::id<2>) {
+            sycl::atomic<int, sycl::access::address_space::local_space> value =
+                acc[0][0];
+          });
+        });
+      }
+    }
+  }
+
   // Device accessor with 3-dimensional subscript operators.
   {
     sycl::queue Queue;

From ce3266f80d76eed28d40267cc8d1dce2ac71c5cb Mon Sep 17 00:00:00 2001
From: mdimakov <maxim.dimakov@intel.com>
Date: Mon, 5 Apr 2021 16:56:41 +0300
Subject: [PATCH 2/4] Clang-format fix

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

diff --git a/SYCL/Basic/accessor/accessor.cpp b/SYCL/Basic/accessor/accessor.cpp
index 43ea804393..b912b2cd83 100644
--- a/SYCL/Basic/accessor/accessor.cpp
+++ b/SYCL/Basic/accessor/accessor.cpp
@@ -178,11 +178,12 @@ int main() {
         queue.submit([&](cl::sycl::handler &cgh) {
           auto acc =
               cl::sycl::accessor<int, 2, cl::sycl::access::mode::atomic,
-              cl::sycl::access::target::local>(range, cgh);
-          cgh.parallel_for<class dim2_subscr_atomic>(range, [=] (cl::sycl::id<2>) {
-            sycl::atomic<int, sycl::access::address_space::local_space> value =
-                acc[0][0];
-          });
+                                cl::sycl::access::target::local>(range, cgh);
+          cgh.parallel_for<class dim2_subscr_atomic>(
+              range, [=] (cl::sycl::id<2>) {
+                sycl::atomic<int, sycl::access::address_space::local_space>
+                    value = acc[0][0];
+              });
         });
       }
     }

From 7fdfafb113b75dbbc10af14af3610589b1396b99 Mon Sep 17 00:00:00 2001
From: mdimakov <maxim.dimakov@intel.com>
Date: Mon, 5 Apr 2021 17:31:23 +0300
Subject: [PATCH 3/4] Clang-format fix

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

diff --git a/SYCL/Basic/accessor/accessor.cpp b/SYCL/Basic/accessor/accessor.cpp
index b912b2cd83..74b4c202c8 100644
--- a/SYCL/Basic/accessor/accessor.cpp
+++ b/SYCL/Basic/accessor/accessor.cpp
@@ -178,9 +178,9 @@ int main() {
         queue.submit([&](cl::sycl::handler &cgh) {
           auto acc =
               cl::sycl::accessor<int, 2, cl::sycl::access::mode::atomic,
-                                cl::sycl::access::target::local>(range, cgh);
+                                 cl::sycl::access::target::local>(range, cgh);
           cgh.parallel_for<class dim2_subscr_atomic>(
-              range, [=] (cl::sycl::id<2>) {
+              range, [=](cl::sycl::id<2>) {
                 sycl::atomic<int, sycl::access::address_space::local_space>
                     value = acc[0][0];
               });

From 2fdc2143e9a7201dfa75b02557df2f7f5df35e77 Mon Sep 17 00:00:00 2001
From: mdimakov <maxim.dimakov@intel.com>
Date: Tue, 6 Apr 2021 15:25:17 +0300
Subject: [PATCH 4/4] Address to review comment

---
 SYCL/Basic/accessor/accessor.cpp | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/SYCL/Basic/accessor/accessor.cpp b/SYCL/Basic/accessor/accessor.cpp
index 74b4c202c8..b30cbcbe2e 100644
--- a/SYCL/Basic/accessor/accessor.cpp
+++ b/SYCL/Basic/accessor/accessor.cpp
@@ -180,7 +180,7 @@ int main() {
               cl::sycl::accessor<int, 2, cl::sycl::access::mode::atomic,
                                  cl::sycl::access::target::local>(range, cgh);
           cgh.parallel_for<class dim2_subscr_atomic>(
-              range, [=](cl::sycl::id<2>) {
+              cl::sycl::nd_range<2>{range, range}, [=](cl::sycl::nd_item<2>) {
                 sycl::atomic<int, sycl::access::address_space::local_space>
                     value = acc[0][0];
               });