From 2503a625ef8ef7386579c446fed0da774c233102 Mon Sep 17 00:00:00 2001
From: M Moadeli <mahmoud.moadeli@codeplay.com>
Date: Tue, 31 Jan 2023 09:49:09 +0000
Subject: [PATCH 01/15] Diagnostic for  const qualified DataT with non
 read-only accessor

---
 sycl/include/sycl/accessor.hpp | 4 ++++
 1 file changed, 4 insertions(+)

diff --git a/sycl/include/sycl/accessor.hpp b/sycl/include/sycl/accessor.hpp
index 897216306fd67..aad162f5da9b1 100644
--- a/sycl/include/sycl/accessor.hpp
+++ b/sycl/include/sycl/accessor.hpp
@@ -304,6 +304,10 @@ class accessor_common {
       AccessMode == access::mode::discard_read_write;
 
   constexpr static bool IsAccessReadOnly = AccessMode == access::mode::read;
+  static constexpr bool isConst = std::is_const<DataT>::value;
+  static_assert(
+      !isConst || IsAccessReadOnly,
+      "A const qualified DataT is only allowed for a read-only accessor");
 
   constexpr static bool IsAccessReadWrite =
       AccessMode == access::mode::read_write;

From 1edb94f82cc1d369e12262fb43b2a103d12b276f Mon Sep 17 00:00:00 2001
From: M Moadeli <mahmoud.moadeli@codeplay.com>
Date: Tue, 31 Jan 2023 22:52:19 +0000
Subject: [PATCH 02/15] Add test for const qualified DataT with non readonly
 accessor. Update failing tests due to adding diagnostic for const qualified
 DataT only allowed for non readonly accessor.

---
 .../accessor/accessor_ptr_alias.cpp           |  6 +++--
 .../const-type-non-readonly-accessor.cpp      | 27 +++++++++++++++++++
 .../host_accessor_get_pointer_type.cpp        |  6 +++--
 .../accessor_subscript_and_ref_type.cpp       | 11 ++++----
 4 files changed, 41 insertions(+), 9 deletions(-)
 create mode 100644 sycl/test/basic_tests/accessor/const-type-non-readonly-accessor.cpp

diff --git a/sycl/test/basic_tests/accessor/accessor_ptr_alias.cpp b/sycl/test/basic_tests/accessor/accessor_ptr_alias.cpp
index b55ababb8612a..fa15ebb4307aa 100644
--- a/sycl/test/basic_tests/accessor/accessor_ptr_alias.cpp
+++ b/sycl/test/basic_tests/accessor/accessor_ptr_alias.cpp
@@ -40,9 +40,11 @@ template <typename DataT, int Dims> void CheckLocalAccessor() {
 
 template <typename DataT, int Dims> void CheckAccessorForModes() {
   CheckDeviceAccessor<DataT, Dims, sycl::access::mode::read>();
-  CheckDeviceAccessor<DataT, Dims, sycl::access::mode::read_write>();
-  CheckDeviceAccessor<DataT, Dims, sycl::access::mode::write>();
   CheckLocalAccessor<DataT, Dims>();
+  if constexpr (!std::is_const<DataT>::value) {
+    CheckDeviceAccessor<DataT, Dims, sycl::access::mode::read_write>();
+    CheckDeviceAccessor<DataT, Dims, sycl::access::mode::write>();
+  }
 }
 
 template <typename DataT> void CheckAccessorForAllDimsAndModes() {
diff --git a/sycl/test/basic_tests/accessor/const-type-non-readonly-accessor.cpp b/sycl/test/basic_tests/accessor/const-type-non-readonly-accessor.cpp
new file mode 100644
index 0000000000000..31fb92642c8ac
--- /dev/null
+++ b/sycl/test/basic_tests/accessor/const-type-non-readonly-accessor.cpp
@@ -0,0 +1,27 @@
+// RUN: not %clangxx -fsycl -fsycl-device-only -fsyntax-only \
+// RUN: %s -I %sycl_include 2>&1 | FileCheck %s
+
+#include <sycl/sycl.hpp>
+
+using namespace sycl;
+
+constexpr size_t dataSize = 1;
+
+int main() {
+  const int a[dataSize] = {1};
+
+  try {
+    auto defaultQueue = queue{};
+    auto bufA = buffer<const int, 1>{a, range{dataSize}};
+    defaultQueue.submit([&](handler &cgh) {
+      sycl::accessor accA{bufA, cgh, read_write};
+    });
+
+    defaultQueue.throw_asynchronous();
+  } catch (const exception &e) {
+    std::cout << "Exception caught: " << e.what() << std::endl;
+  }
+  return 0;
+}
+
+// CHECK: static assertion failed due to requirement '!isConst || IsAccessReadOnly': A const qualified DataT is only allowed for a read-only accessor
\ No newline at end of file
diff --git a/sycl/test/basic_tests/accessor/host_accessor_get_pointer_type.cpp b/sycl/test/basic_tests/accessor/host_accessor_get_pointer_type.cpp
index 50c830f6240d1..7a4bdfcfd423c 100644
--- a/sycl/test/basic_tests/accessor/host_accessor_get_pointer_type.cpp
+++ b/sycl/test/basic_tests/accessor/host_accessor_get_pointer_type.cpp
@@ -18,8 +18,10 @@ void CheckHostAccessor() {
 
 template <typename DataT, int Dims> void CheckHostAccessorForModes() {
   CheckHostAccessor<DataT, Dims, sycl::access::mode::read>();
-  CheckHostAccessor<DataT, Dims, sycl::access::mode::read_write>();
-  CheckHostAccessor<DataT, Dims, sycl::access::mode::write>();
+  if constexpr (!std::is_const<DataT>::value) {
+    CheckHostAccessor<DataT, Dims, sycl::access::mode::read_write>();
+    CheckHostAccessor<DataT, Dims, sycl::access::mode::write>();
+  }
 }
 
 template <typename DataT> void CheckHostAccessorForAllDimsAndModes() {
diff --git a/sycl/test/regression/accessor_subscript_and_ref_type.cpp b/sycl/test/regression/accessor_subscript_and_ref_type.cpp
index d2b818d55a1cb..e2c6f74515e68 100644
--- a/sycl/test/regression/accessor_subscript_and_ref_type.cpp
+++ b/sycl/test/regression/accessor_subscript_and_ref_type.cpp
@@ -81,13 +81,14 @@ template <typename DataT, access::mode AccessMode> void CheckAccAllDims() {
 }
 
 template <typename DataT> void CheckAccAllAccessModesAndDims() {
-  CheckAccAllDims<DataT, access::mode::write>();
-  CheckAccAllDims<DataT, access::mode::read_write>();
-  CheckAccAllDims<DataT, access::mode::discard_write>();
-  CheckAccAllDims<DataT, access::mode::discard_read_write>();
   CheckAccAllDims<DataT, access::mode::read>();
-  if constexpr (!std::is_const_v<DataT>)
+  if constexpr (!std::is_const_v<DataT>) {
+    CheckAccAllDims<DataT, access::mode::write>();
+    CheckAccAllDims<DataT, access::mode::read_write>();
+    CheckAccAllDims<DataT, access::mode::discard_write>();
+    CheckAccAllDims<DataT, access::mode::discard_read_write>();
     CheckAccAllDims<DataT, access::mode::atomic>();
+  }
 }
 
 int main() {

From 7ba64c6a688f50fd04d9a32ad6529a5cb66c360f Mon Sep 17 00:00:00 2001
From: M Moadeli <mahmoud.moadeli@codeplay.com>
Date: Wed, 1 Feb 2023 00:47:49 +0000
Subject: [PATCH 03/15] Assign valid access mode for local accessor, depending
 on the type of data.

---
 sycl/include/sycl/accessor.hpp | 14 ++++++++++----
 1 file changed, 10 insertions(+), 4 deletions(-)

diff --git a/sycl/include/sycl/accessor.hpp b/sycl/include/sycl/accessor.hpp
index 17b32d03e3383..31199fc666c0f 100644
--- a/sycl/include/sycl/accessor.hpp
+++ b/sycl/include/sycl/accessor.hpp
@@ -2645,15 +2645,21 @@ class __SYCL_EBO __SYCL_SPECIAL_CLASS accessor<
 #endif
 };
 
+template <typename DataT> constexpr access::mode accessMode() {
+  if constexpr (std::is_const<DataT>::value)
+    return access::mode::read;
+  else
+    return access::mode::read_write;
+}
+
 template <typename DataT, int Dimensions = 1>
 class __SYCL_EBO __SYCL_SPECIAL_CLASS __SYCL_TYPE(local_accessor) local_accessor
-    : public local_accessor_base<DataT, Dimensions, access::mode::read_write,
+    : public local_accessor_base<DataT, Dimensions, accessMode<DataT>(),
                                  access::placeholder::false_t>,
       public detail::OwnerLessBase<local_accessor<DataT, Dimensions>> {
 
-  using local_acc =
-      local_accessor_base<DataT, Dimensions, access::mode::read_write,
-                          access::placeholder::false_t>;
+  using local_acc = local_accessor_base<DataT, Dimensions, accessMode<DataT>(),
+                                        access::placeholder::false_t>;
 
   // Use base classes constructors
   using local_acc::local_acc;

From ef0038e9a6a1e7d8d79fd8af1e506da296945092 Mon Sep 17 00:00:00 2001
From: M Moadeli <mahmoud.moadeli@codeplay.com>
Date: Wed, 1 Feb 2023 01:00:08 +0000
Subject: [PATCH 04/15] Minor test update.

---
 .../basic_tests/accessor/const-type-non-readonly-accessor.cpp | 4 +---
 1 file changed, 1 insertion(+), 3 deletions(-)

diff --git a/sycl/test/basic_tests/accessor/const-type-non-readonly-accessor.cpp b/sycl/test/basic_tests/accessor/const-type-non-readonly-accessor.cpp
index 31fb92642c8ac..962dc9a59008c 100644
--- a/sycl/test/basic_tests/accessor/const-type-non-readonly-accessor.cpp
+++ b/sycl/test/basic_tests/accessor/const-type-non-readonly-accessor.cpp
@@ -16,12 +16,10 @@ int main() {
     defaultQueue.submit([&](handler &cgh) {
       sycl::accessor accA{bufA, cgh, read_write};
     });
-
-    defaultQueue.throw_asynchronous();
   } catch (const exception &e) {
     std::cout << "Exception caught: " << e.what() << std::endl;
   }
   return 0;
 }
 
-// CHECK: static assertion failed due to requirement '!isConst || IsAccessReadOnly': A const qualified DataT is only allowed for a read-only accessor
\ No newline at end of file
+// CHECK: static assertion failed due to requirement '!isConst || IsAccessReadOnly': A const qualified DataT is only allowed for a read-only accessor

From 36971f8b80fd8021aec60fac640c37d008a97dab Mon Sep 17 00:00:00 2001
From: M Moadeli <mahmoud.moadeli@codeplay.com>
Date: Thu, 2 Feb 2023 13:54:15 +0000
Subject: [PATCH 05/15] - Rename AccessMode to AccessModeFromConstness - Use
 -Xclang -verify for testing

---
 sycl/include/sycl/accessor.hpp                 | 18 +++++++++---------
 .../const-type-non-readonly-accessor.cpp       |  9 ++++-----
 2 files changed, 13 insertions(+), 14 deletions(-)

diff --git a/sycl/include/sycl/accessor.hpp b/sycl/include/sycl/accessor.hpp
index 31199fc666c0f..1ea4b76e8798d 100644
--- a/sycl/include/sycl/accessor.hpp
+++ b/sycl/include/sycl/accessor.hpp
@@ -365,6 +365,13 @@ class accessor_common {
   };
 };
 
+template <typename DataT> constexpr access::mode accessModeFromConstness() {
+  if constexpr (std::is_const<DataT>::value)
+    return access::mode::read;
+  else
+    return access::mode::read_write;
+}
+
 template <typename MayBeTag1, typename MayBeTag2>
 constexpr access::mode deduceAccessMode() {
   // property_list = {} is not properly detected by deduction guide,
@@ -2645,20 +2652,13 @@ class __SYCL_EBO __SYCL_SPECIAL_CLASS accessor<
 #endif
 };
 
-template <typename DataT> constexpr access::mode accessMode() {
-  if constexpr (std::is_const<DataT>::value)
-    return access::mode::read;
-  else
-    return access::mode::read_write;
-}
-
 template <typename DataT, int Dimensions = 1>
 class __SYCL_EBO __SYCL_SPECIAL_CLASS __SYCL_TYPE(local_accessor) local_accessor
-    : public local_accessor_base<DataT, Dimensions, accessMode<DataT>(),
+    : public local_accessor_base<DataT, Dimensions, detail::accessModeFromConstness<DataT>(),
                                  access::placeholder::false_t>,
       public detail::OwnerLessBase<local_accessor<DataT, Dimensions>> {
 
-  using local_acc = local_accessor_base<DataT, Dimensions, accessMode<DataT>(),
+  using local_acc = local_accessor_base<DataT, Dimensions, detail::accessModeFromConstness<DataT>(),
                                         access::placeholder::false_t>;
 
   // Use base classes constructors
diff --git a/sycl/test/basic_tests/accessor/const-type-non-readonly-accessor.cpp b/sycl/test/basic_tests/accessor/const-type-non-readonly-accessor.cpp
index 962dc9a59008c..a72be5626daeb 100644
--- a/sycl/test/basic_tests/accessor/const-type-non-readonly-accessor.cpp
+++ b/sycl/test/basic_tests/accessor/const-type-non-readonly-accessor.cpp
@@ -1,5 +1,5 @@
-// RUN: not %clangxx -fsycl -fsycl-device-only -fsyntax-only \
-// RUN: %s -I %sycl_include 2>&1 | FileCheck %s
+// RUN: %clangxx -fsyntax-only %fsycl-host-only -Xclang -verify -Xclang -verify-ignore-unexpected=error,note %s
+// RUN: %clangxx -fsyntax-only -fsycl -fsycl-device-only -Xclang -verify -Xclang -verify-ignore-unexpected=error,note %s
 
 #include <sycl/sycl.hpp>
 
@@ -14,12 +14,11 @@ int main() {
     auto defaultQueue = queue{};
     auto bufA = buffer<const int, 1>{a, range{dataSize}};
     defaultQueue.submit([&](handler &cgh) {
-      sycl::accessor accA{bufA, cgh, read_write};
+      accessor accA{bufA, cgh, read_write};
+    // expected-error@sycl/accessor.hpp:* {{A const qualified DataT is only allowed for a read-only accessor}}
     });
   } catch (const exception &e) {
     std::cout << "Exception caught: " << e.what() << std::endl;
   }
   return 0;
 }
-
-// CHECK: static assertion failed due to requirement '!isConst || IsAccessReadOnly': A const qualified DataT is only allowed for a read-only accessor

From 29e3036b3363d24484eb6f3bff7291e6169a09df Mon Sep 17 00:00:00 2001
From: M Moadeli <mahmoud.moadeli@codeplay.com>
Date: Thu, 2 Feb 2023 15:59:42 +0000
Subject: [PATCH 06/15] - Revert test to be without -Xclang -verify. Having it
 requires handling over 20 expected-erros.

---
 .../accessor/const-type-non-readonly-accessor.cpp        | 9 ++++++---
 1 file changed, 6 insertions(+), 3 deletions(-)

diff --git a/sycl/test/basic_tests/accessor/const-type-non-readonly-accessor.cpp b/sycl/test/basic_tests/accessor/const-type-non-readonly-accessor.cpp
index a72be5626daeb..a60cdee1469c7 100644
--- a/sycl/test/basic_tests/accessor/const-type-non-readonly-accessor.cpp
+++ b/sycl/test/basic_tests/accessor/const-type-non-readonly-accessor.cpp
@@ -1,5 +1,5 @@
-// RUN: %clangxx -fsyntax-only %fsycl-host-only -Xclang -verify -Xclang -verify-ignore-unexpected=error,note %s
-// RUN: %clangxx -fsyntax-only -fsycl -fsycl-device-only -Xclang -verify -Xclang -verify-ignore-unexpected=error,note %s
+// RUN: not %clangxx -fsycl -fsycl-device-only -fsyntax-only \
+// RUN: %s -I %sycl_include 2>&1 | FileCheck %s
 
 #include <sycl/sycl.hpp>
 
@@ -15,10 +15,13 @@ int main() {
     auto bufA = buffer<const int, 1>{a, range{dataSize}};
     defaultQueue.submit([&](handler &cgh) {
       accessor accA{bufA, cgh, read_write};
-    // expected-error@sycl/accessor.hpp:* {{A const qualified DataT is only allowed for a read-only accessor}}
     });
+
+    defaultQueue.throw_asynchronous();
   } catch (const exception &e) {
     std::cout << "Exception caught: " << e.what() << std::endl;
   }
   return 0;
 }
+
+// CHECK: static assertion failed due to requirement '!isConst || IsAccessReadOnly': A const qualified DataT is only allowed for a read-only accessor
\ No newline at end of file

From 92debe8e9afd6a49a213c032944cee6c76a82707 Mon Sep 17 00:00:00 2001
From: M Moadeli <mahmoud.moadeli@codeplay.com>
Date: Fri, 3 Feb 2023 14:48:25 +0000
Subject: [PATCH 07/15] Refactors implementation of diagnostic to avoid
 compiler errors due to inheritance from base class hitting assert.

---
 sycl/include/sycl/accessor.hpp | 26 +++++++++++++++++++++-----
 1 file changed, 21 insertions(+), 5 deletions(-)

diff --git a/sycl/include/sycl/accessor.hpp b/sycl/include/sycl/accessor.hpp
index 1ea4b76e8798d..9dbdfd6cd9fff 100644
--- a/sycl/include/sycl/accessor.hpp
+++ b/sycl/include/sycl/accessor.hpp
@@ -304,10 +304,7 @@ class accessor_common {
       AccessMode == access::mode::discard_read_write;
 
   constexpr static bool IsAccessReadOnly = AccessMode == access::mode::read;
-  static constexpr bool isConst = std::is_const<DataT>::value;
-  static_assert(
-      !isConst || IsAccessReadOnly,
-      "A const qualified DataT is only allowed for a read-only accessor");
+  static constexpr bool IsConst = std::is_const<DataT>::value;
 
   constexpr static bool IsAccessReadWrite =
       AccessMode == access::mode::read_write;
@@ -1014,10 +1011,15 @@ class __SYCL_EBO __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor :
   static constexpr bool IsGlobalBuf = AccessorCommonT::IsGlobalBuf;
   static constexpr bool IsHostBuf = AccessorCommonT::IsHostBuf;
   static constexpr bool IsPlaceH = AccessorCommonT::IsPlaceH;
+  static constexpr bool IsConst = AccessorCommonT::IsConst;
   template <int Dims>
   using AccessorSubscript =
       typename AccessorCommonT::template AccessorSubscript<Dims>;
 
+  static_assert(
+      !IsConst || IsAccessReadOnly,
+      "A const qualified DataT is only allowed for a read-only accessor");
+
   using ConcreteASPtrType = typename detail::DecoratedType<DataT, AS>::type *;
 
   using RefType = detail::const_if_const_AS<AS, DataT> &;
@@ -2383,7 +2385,13 @@ class __SYCL_SPECIAL_CLASS local_accessor_base :
                               access::target::local, IsPlaceholder>;
 
   using AccessorCommonT::AS;
-  using AccessorCommonT::IsAccessAnyWrite;
+
+  // Cannot do "using AccessorCommonT::Flag" as it doesn't work with g++ as host
+  // compiler, for some reason.
+  static constexpr bool IsAccessAnyWrite = AccessorCommonT::IsAccessAnyWrite;
+  static constexpr bool IsAccessReadOnly = AccessorCommonT::IsAccessReadOnly;
+  static constexpr bool IsConst = AccessorCommonT::IsConst;
+
   template <int Dims>
   using AccessorSubscript =
       typename AccessorCommonT::template AccessorSubscript<
@@ -2625,6 +2633,10 @@ class __SYCL_EBO __SYCL_SPECIAL_CLASS accessor<
   using local_acc =
       local_accessor_base<DataT, Dimensions, AccessMode, IsPlaceholder>;
 
+  static_assert(
+      !local_acc::IsConst || local_acc::IsAccessReadOnly,
+      "A const qualified DataT is only allowed for a read-only accessor");
+
   // Use base classes constructors
   using local_acc::local_acc;
 
@@ -2661,6 +2673,10 @@ class __SYCL_EBO __SYCL_SPECIAL_CLASS __SYCL_TYPE(local_accessor) local_accessor
   using local_acc = local_accessor_base<DataT, Dimensions, detail::accessModeFromConstness<DataT>(),
                                         access::placeholder::false_t>;
 
+  static_assert(
+      !local_acc::IsConst || local_acc::IsAccessReadOnly,
+      "A const qualified DataT is only allowed for a read-only accessor");
+
   // Use base classes constructors
   using local_acc::local_acc;
 

From 7eac372090e737d18b2e66e1d6929ad57504ccf6 Mon Sep 17 00:00:00 2001
From: M Moadeli <mahmoud.moadeli@codeplay.com>
Date: Fri, 3 Feb 2023 14:54:39 +0000
Subject: [PATCH 08/15] Updates the test to use -Xclang -verify

---
 .../accessor/const-type-non-readonly-accessor.cpp           | 6 ++----
 1 file changed, 2 insertions(+), 4 deletions(-)

diff --git a/sycl/test/basic_tests/accessor/const-type-non-readonly-accessor.cpp b/sycl/test/basic_tests/accessor/const-type-non-readonly-accessor.cpp
index a60cdee1469c7..32c20a0bff257 100644
--- a/sycl/test/basic_tests/accessor/const-type-non-readonly-accessor.cpp
+++ b/sycl/test/basic_tests/accessor/const-type-non-readonly-accessor.cpp
@@ -1,5 +1,4 @@
-// RUN: not %clangxx -fsycl -fsycl-device-only -fsyntax-only \
-// RUN: %s -I %sycl_include 2>&1 | FileCheck %s
+// RUN: %clangxx -fsycl -fsyntax-only -Xclang -verify -Xclang -verify-ignore-unexpected=note %s
 
 #include <sycl/sycl.hpp>
 
@@ -14,6 +13,7 @@ int main() {
     auto defaultQueue = queue{};
     auto bufA = buffer<const int, 1>{a, range{dataSize}};
     defaultQueue.submit([&](handler &cgh) {
+      // expected-error@sycl/accessor.hpp:* {{A const qualified DataT is only allowed for a read-only accessor}}
       accessor accA{bufA, cgh, read_write};
     });
 
@@ -23,5 +23,3 @@ int main() {
   }
   return 0;
 }
-
-// CHECK: static assertion failed due to requirement '!isConst || IsAccessReadOnly': A const qualified DataT is only allowed for a read-only accessor
\ No newline at end of file

From dcac70615d553e69057a85edf7151292c3b59952 Mon Sep 17 00:00:00 2001
From: M Moadeli <mahmoud.moadeli@codeplay.com>
Date: Fri, 3 Feb 2023 16:07:27 +0000
Subject: [PATCH 09/15] Fix style-checl

---
 sycl/include/sycl/accessor.hpp | 35 ++++++++++++++++++----------------
 1 file changed, 19 insertions(+), 16 deletions(-)

diff --git a/sycl/include/sycl/accessor.hpp b/sycl/include/sycl/accessor.hpp
index 9dbdfd6cd9fff..6be701f36890f 100644
--- a/sycl/include/sycl/accessor.hpp
+++ b/sycl/include/sycl/accessor.hpp
@@ -2484,7 +2484,8 @@ class __SYCL_SPECIAL_CLASS local_accessor_base :
   local_accessor_base(handler &, const detail::code_location CodeLoc =
                                      detail::code_location::current())
 #ifdef __SYCL_DEVICE_ONLY__
-      : impl(range<AdjustedDim>{1}){}
+      : impl(range<AdjustedDim>{1}) {
+  }
 #else
       : LocalAccessorBaseHost(range<3>{1, 1, 1}, AdjustedDim, sizeof(DataT)) {
     detail::constructorNotification(nullptr, LocalAccessorBaseHost::impl.get(),
@@ -2493,11 +2494,10 @@ class __SYCL_SPECIAL_CLASS local_accessor_base :
   }
 #endif
 
-        template <int Dims = Dimensions,
-                  typename = detail::enable_if_t<Dims == 0>>
-        local_accessor_base(handler &, const property_list &propList,
-                            const detail::code_location CodeLoc =
-                                detail::code_location::current())
+  template <int Dims = Dimensions, typename = detail::enable_if_t<Dims == 0>>
+  local_accessor_base(
+      handler &, const property_list &propList,
+      const detail::code_location CodeLoc = detail::code_location::current())
 #ifdef __SYCL_DEVICE_ONLY__
       : impl(range<AdjustedDim>{1}) {
     (void)propList;
@@ -2516,7 +2516,8 @@ class __SYCL_SPECIAL_CLASS local_accessor_base :
       range<Dimensions> AllocationSize, handler &,
       const detail::code_location CodeLoc = detail::code_location::current())
 #ifdef __SYCL_DEVICE_ONLY__
-      : impl(AllocationSize){}
+      : impl(AllocationSize) {
+  }
 #else
       : LocalAccessorBaseHost(detail::convertToArrayOfN<3, 1>(AllocationSize),
                               AdjustedDim, sizeof(DataT)) {
@@ -2526,12 +2527,11 @@ class __SYCL_SPECIAL_CLASS local_accessor_base :
   }
 #endif
 
-        template <int Dims = Dimensions,
-                  typename = detail::enable_if_t<(Dims > 0)>>
-        local_accessor_base(range<Dimensions> AllocationSize, handler &,
-                            const property_list &propList,
-                            const detail::code_location CodeLoc =
-                                detail::code_location::current())
+  template <int Dims = Dimensions, typename = detail::enable_if_t<(Dims > 0)>>
+  local_accessor_base(
+      range<Dimensions> AllocationSize, handler &,
+      const property_list &propList,
+      const detail::code_location CodeLoc = detail::code_location::current())
 #ifdef __SYCL_DEVICE_ONLY__
       : impl(AllocationSize) {
     (void)propList;
@@ -2666,12 +2666,15 @@ class __SYCL_EBO __SYCL_SPECIAL_CLASS accessor<
 
 template <typename DataT, int Dimensions = 1>
 class __SYCL_EBO __SYCL_SPECIAL_CLASS __SYCL_TYPE(local_accessor) local_accessor
-    : public local_accessor_base<DataT, Dimensions, detail::accessModeFromConstness<DataT>(),
+    : public local_accessor_base<DataT, Dimensions,
+                                 detail::accessModeFromConstness<DataT>(),
                                  access::placeholder::false_t>,
       public detail::OwnerLessBase<local_accessor<DataT, Dimensions>> {
 
-  using local_acc = local_accessor_base<DataT, Dimensions, detail::accessModeFromConstness<DataT>(),
-                                        access::placeholder::false_t>;
+  using local_acc =
+      local_accessor_base<DataT, Dimensions,
+                          detail::accessModeFromConstness<DataT>(),
+                          access::placeholder::false_t>;
 
   static_assert(
       !local_acc::IsConst || local_acc::IsAccessReadOnly,

From b142c5df7b6a06e6a723d0a1f3bb4f0f39f8f394 Mon Sep 17 00:00:00 2001
From: m moadeli <mahmoud.moadeli@codeplay.com>
Date: Wed, 8 Mar 2023 23:47:40 +0000
Subject: [PATCH 10/15] Throw exception when using local_accessor in a SYCL
 kernel function that is invoked via single_task or via the simple form of
 parallel_for that takes a range parameter.

---
 sycl/include/sycl/handler.hpp                 | 27 +++++++++++++++
 .../accessor/local_accessor_misuse.cpp        | 34 +++++++++++++++++++
 2 files changed, 61 insertions(+)
 create mode 100644 sycl/test/basic_tests/accessor/local_accessor_misuse.cpp

diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp
index 6ffd3922cc205..868c65de61c7c 100644
--- a/sycl/include/sycl/handler.hpp
+++ b/sycl/include/sycl/handler.hpp
@@ -354,6 +354,31 @@ class __SYCL_EXPORT handler {
                                 PI_ERROR_INVALID_OPERATION);
   }
 
+  /// According to section 4.7.6.11. of the SYCL specification, a local accessor
+  /// must not be used in a SYCL kernel function that is invoked via single_task
+  /// or via the simple form of parallel_for that takes a range parameter.
+  template <typename KernelName, typename KernelType>
+  void throwOnLocalAccessorMisuse() {
+    using NameT =
+        typename detail::get_kernel_name_t<KernelName, KernelType>::name;
+    using KI = sycl::detail::KernelInfo<NameT>;
+
+    auto *KernelArgs = &KI::getParamDesc(0);
+
+    for (unsigned I = 0; I < KI::getNumParams(); ++I) {
+      const detail::kernel_param_kind_t &Kind = KernelArgs[I].kind;
+      const access::target AccTarget =
+          static_cast<access::target>(KernelArgs[I].info & 0x7ff);
+      if ((Kind == detail::kernel_param_kind_t::kind_accessor) &&
+          (AccTarget == target::local))
+        throw sycl::exception(
+            make_error_code(errc::kernel_argument),
+            "A local accessor must not be used in a SYCL kernel function "
+            "that is invoked via single_task or via the simple form of "
+            "parallel_for that takes a range parameter.");
+    }
+  }
+
   /// Extracts and prepares kernel arguments from the lambda using integration
   /// header.
   void
@@ -908,6 +933,7 @@ class __SYCL_EXPORT handler {
   void parallel_for_lambda_impl(range<Dims> NumWorkItems,
                                 KernelType KernelFunc) {
     throwIfActionIsCreated();
+    throwOnLocalAccessorMisuse<KernelName, KernelType>();
     using LambdaArgType = sycl::detail::lambda_arg_type<KernelType, item<Dims>>;
 
     // If 1D kernel argument is an integral type, convert it to sycl::item<1>
@@ -1385,6 +1411,7 @@ class __SYCL_EXPORT handler {
                 ext::oneapi::experimental::detail::empty_properties_t>
   void single_task_lambda_impl(_KERNELFUNCPARAM(KernelFunc)) {
     throwIfActionIsCreated();
+    throwOnLocalAccessorMisuse<KernelName, KernelType>();
     // TODO: Properties may change the kernel function, so in order to avoid
     //       conflicts they should be included in the name.
     using NameT =
diff --git a/sycl/test/basic_tests/accessor/local_accessor_misuse.cpp b/sycl/test/basic_tests/accessor/local_accessor_misuse.cpp
new file mode 100644
index 0000000000000..4c143811c9437
--- /dev/null
+++ b/sycl/test/basic_tests/accessor/local_accessor_misuse.cpp
@@ -0,0 +1,34 @@
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
+// RUN: %t.out
+
+#include <cassert>
+#include <sycl/sycl.hpp>
+
+using namespace sycl;
+
+int main() {
+  constexpr static int size = 1;
+  queue testQueue;
+
+  try {
+
+    testQueue.submit([&](handler &cgh) {
+      auto local_acc = local_accessor<int, 1>({size}, cgh);
+      cgh.single_task<class kernel>([=]() { (void)local_acc; });
+    });
+    assert(0);
+  } catch (sycl::exception) {
+  }
+
+  try {
+    testQueue.submit([&](sycl::handler &cgh) {
+      auto local_acc = local_accessor<int, 1>({42}, cgh);
+      cgh.parallel_for<class parallel_kernel>(
+          sycl::range<1>{size}, [=](sycl::id<1> ID) { (void)local_acc; });
+    });
+    assert(0);
+  } catch (sycl::exception) {
+  }
+
+  return 0;
+}

From 98ea1ee9e4db9438d8ba6830269f66a4996b02ab Mon Sep 17 00:00:00 2001
From: m moadeli <mahmoud.moadeli@codeplay.com>
Date: Wed, 8 Mar 2023 23:59:54 +0000
Subject: [PATCH 11/15] Minor size update.

---
 sycl/test/basic_tests/accessor/local_accessor_misuse.cpp | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/sycl/test/basic_tests/accessor/local_accessor_misuse.cpp b/sycl/test/basic_tests/accessor/local_accessor_misuse.cpp
index 4c143811c9437..e3eb72d472911 100644
--- a/sycl/test/basic_tests/accessor/local_accessor_misuse.cpp
+++ b/sycl/test/basic_tests/accessor/local_accessor_misuse.cpp
@@ -22,7 +22,7 @@ int main() {
 
   try {
     testQueue.submit([&](sycl::handler &cgh) {
-      auto local_acc = local_accessor<int, 1>({42}, cgh);
+      auto local_acc = local_accessor<int, 1>({size}, cgh);
       cgh.parallel_for<class parallel_kernel>(
           sycl::range<1>{size}, [=](sycl::id<1> ID) { (void)local_acc; });
     });

From 2d0f1d608690dc0952d84188ec3bfabb34615a6f Mon Sep 17 00:00:00 2001
From: m moadeli <mahmoud.moadeli@codeplay.com>
Date: Thu, 9 Mar 2023 09:38:52 +0000
Subject: [PATCH 12/15] Limit target to spir64

---
 sycl/test/basic_tests/accessor/local_accessor_misuse.cpp | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/sycl/test/basic_tests/accessor/local_accessor_misuse.cpp b/sycl/test/basic_tests/accessor/local_accessor_misuse.cpp
index e3eb72d472911..0ee2893734916 100644
--- a/sycl/test/basic_tests/accessor/local_accessor_misuse.cpp
+++ b/sycl/test/basic_tests/accessor/local_accessor_misuse.cpp
@@ -1,4 +1,4 @@
-// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
+// RUN: %clangxx -fsycl -fsycl-targets=spir64 %s -o %t.out
 // RUN: %t.out
 
 #include <cassert>

From 03107237c91ec49f2cbe9198d07ff69338c05b16 Mon Sep 17 00:00:00 2001
From: m moadeli <mahmoud.moadeli@codeplay.com>
Date: Thu, 9 Mar 2023 10:37:22 +0000
Subject: [PATCH 13/15] Moves test on local accessor diagnostics to
 llvm-test-suite.

---
 .../accessor/local_accessor_misuse.cpp        | 34 -------------------
 1 file changed, 34 deletions(-)
 delete mode 100644 sycl/test/basic_tests/accessor/local_accessor_misuse.cpp

diff --git a/sycl/test/basic_tests/accessor/local_accessor_misuse.cpp b/sycl/test/basic_tests/accessor/local_accessor_misuse.cpp
deleted file mode 100644
index 0ee2893734916..0000000000000
--- a/sycl/test/basic_tests/accessor/local_accessor_misuse.cpp
+++ /dev/null
@@ -1,34 +0,0 @@
-// RUN: %clangxx -fsycl -fsycl-targets=spir64 %s -o %t.out
-// RUN: %t.out
-
-#include <cassert>
-#include <sycl/sycl.hpp>
-
-using namespace sycl;
-
-int main() {
-  constexpr static int size = 1;
-  queue testQueue;
-
-  try {
-
-    testQueue.submit([&](handler &cgh) {
-      auto local_acc = local_accessor<int, 1>({size}, cgh);
-      cgh.single_task<class kernel>([=]() { (void)local_acc; });
-    });
-    assert(0);
-  } catch (sycl::exception) {
-  }
-
-  try {
-    testQueue.submit([&](sycl::handler &cgh) {
-      auto local_acc = local_accessor<int, 1>({size}, cgh);
-      cgh.parallel_for<class parallel_kernel>(
-          sycl::range<1>{size}, [=](sycl::id<1> ID) { (void)local_acc; });
-    });
-    assert(0);
-  } catch (sycl::exception) {
-  }
-
-  return 0;
-}

From 854423b57586e1916ebbe2677ebf51c311aa0fe9 Mon Sep 17 00:00:00 2001
From: M Moadeli <mahmoud.moadeli@codeplay.com>
Date: Thu, 16 Mar 2023 13:10:16 +0000
Subject: [PATCH 14/15] Replace multiple instances of literal value with
 consexpr static definition.

---
 sycl/include/sycl/handler.hpp | 3 ++-
 sycl/source/handler.cpp       | 5 +++--
 2 files changed, 5 insertions(+), 3 deletions(-)

diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp
index c652b53942974..4a48255ebe39a 100644
--- a/sycl/include/sycl/handler.hpp
+++ b/sycl/include/sycl/handler.hpp
@@ -354,6 +354,7 @@ class __SYCL_EXPORT handler {
                                 PI_ERROR_INVALID_OPERATION);
   }
 
+  constexpr static int AccessTargetMask = 0x7ff;
   /// According to section 4.7.6.11. of the SYCL specification, a local accessor
   /// must not be used in a SYCL kernel function that is invoked via single_task
   /// or via the simple form of parallel_for that takes a range parameter.
@@ -368,7 +369,7 @@ class __SYCL_EXPORT handler {
     for (unsigned I = 0; I < KI::getNumParams(); ++I) {
       const detail::kernel_param_kind_t &Kind = KernelArgs[I].kind;
       const access::target AccTarget =
-          static_cast<access::target>(KernelArgs[I].info & 0x7ff);
+          static_cast<access::target>(KernelArgs[I].info & AccessTargetMask);
       if ((Kind == detail::kernel_param_kind_t::kind_accessor) &&
           (AccTarget == target::local))
         throw sycl::exception(
diff --git a/sycl/source/handler.cpp b/sycl/source/handler.cpp
index 7fbb962df6b5a..4cdae7f896264 100644
--- a/sycl/source/handler.cpp
+++ b/sycl/source/handler.cpp
@@ -491,7 +491,8 @@ void handler::processArg(void *Ptr, const detail::kernel_param_kind_t &Kind,
   case kernel_param_kind_t::kind_accessor: {
     // For args kind of accessor Size is information about accessor.
     // The first 11 bits of Size encodes the accessor target.
-    const access::target AccTarget = static_cast<access::target>(Size & 0x7ff);
+    const access::target AccTarget =
+        static_cast<access::target>(Size & AccessTargetMask);
     switch (AccTarget) {
     case access::target::device:
     case access::target::constant_buffer: {
@@ -616,7 +617,7 @@ void handler::extractArgsAndReqsFromLambda(
       // For args kind of accessor Size is information about accessor.
       // The first 11 bits of Size encodes the accessor target.
       const access::target AccTarget =
-          static_cast<access::target>(Size & 0x7ff);
+          static_cast<access::target>(Size & AccessTargetMask);
       if ((AccTarget == access::target::device ||
            AccTarget == access::target::constant_buffer) ||
           (AccTarget == access::target::image ||

From 46f380ce4acb56270f412069424aca1f0cdacad5 Mon Sep 17 00:00:00 2001
From: M Moadeli <mahmoud.moadeli@codeplay.com>
Date: Thu, 16 Mar 2023 13:12:40 +0000
Subject: [PATCH 15/15] Declare throwOnLocalAccessorMisuse as const

---
 sycl/include/sycl/handler.hpp | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp
index 4a48255ebe39a..0f0b1b192dc57 100644
--- a/sycl/include/sycl/handler.hpp
+++ b/sycl/include/sycl/handler.hpp
@@ -359,7 +359,7 @@ class __SYCL_EXPORT handler {
   /// must not be used in a SYCL kernel function that is invoked via single_task
   /// or via the simple form of parallel_for that takes a range parameter.
   template <typename KernelName, typename KernelType>
-  void throwOnLocalAccessorMisuse() {
+  void throwOnLocalAccessorMisuse() const {
     using NameT =
         typename detail::get_kernel_name_t<KernelName, KernelType>::name;
     using KI = sycl::detail::KernelInfo<NameT>;