From 534f667e9925715cfdf47965b176f892940bf198 Mon Sep 17 00:00:00 2001
From: "aidan.belton" <aidan.belton@codeplay.com>
Date: Tue, 21 Jun 2022 09:51:58 +0100
Subject: [PATCH 01/13] Add local_accessor and deprecate target::local accessor

---
 sycl/include/CL/sycl/accessor.hpp                    | 12 ++++++++++--
 .../test/basic_tests/accessor/addrspace_exposure.cpp |  5 ++++-
 sycl/test/basic_tests/set_arg_error.cpp              |  5 +++++
 sycl/test/check_device_code/kernel_arguments_as.cpp  | 11 +++++++++++
 sycl/test/esimd/simd_copy_to_copy_from.cpp           |  4 +---
 sycl/test/extensions/sub_group_as.cpp                |  4 +---
 sycl/test/multi_ptr/ctad.cpp                         |  5 ++++-
 7 files changed, 36 insertions(+), 10 deletions(-)

diff --git a/sycl/include/CL/sycl/accessor.hpp b/sycl/include/CL/sycl/accessor.hpp
index e39268fcad55f..1c76ecce07749 100644
--- a/sycl/include/CL/sycl/accessor.hpp
+++ b/sycl/include/CL/sycl/accessor.hpp
@@ -2015,8 +2015,10 @@ accessor(buffer<DataT, Dimensions, AllocatorT>, handler, Type1, Type2, Type3,
 /// \ingroup sycl_api_acc
 template <typename DataT, int Dimensions, access::mode AccessMode,
           access::placeholder IsPlaceholder>
-class __SYCL_SPECIAL_CLASS accessor<DataT, Dimensions, AccessMode,
-                                    access::target::local, IsPlaceholder> :
+class __SYCL_SPECIAL_CLASS __SYCL2020_DEPRECATED(
+    "'accessor' with 'target::local' is deprecated: use 'local_accessor' "
+    "instead.") accessor<DataT, Dimensions, AccessMode, access::target::local,
+                         IsPlaceholder> :
 #ifndef __SYCL_DEVICE_ONLY__
     public detail::LocalAccessorBaseHost,
 #endif
@@ -2218,6 +2220,12 @@ class __SYCL_SPECIAL_CLASS accessor<DataT, Dimensions, AccessMode,
   bool operator!=(const accessor &Rhs) const { return !(*this == Rhs); }
 };
 
+// local_accessor
+// Same semantics and restrictions as accessor with target::local
+template <typename DataT, int Dimensions = 1>
+using local_accessor =
+    accessor<DataT, Dimensions, access_mode::read_write, access::target::local>;
+
 /// Image accessors.
 ///
 /// Available only when accessTarget == access::target::image.
diff --git a/sycl/test/basic_tests/accessor/addrspace_exposure.cpp b/sycl/test/basic_tests/accessor/addrspace_exposure.cpp
index 0e793032f60e5..eba2462cb152a 100644
--- a/sycl/test/basic_tests/accessor/addrspace_exposure.cpp
+++ b/sycl/test/basic_tests/accessor/addrspace_exposure.cpp
@@ -31,7 +31,8 @@ int main() {
         GlobalBuf.get_access<mode::read, target::global_buffer>(Cgh);
     auto ConstantAcc =
         ConstantBuf.get_access<mode::read, target::constant_buffer>(Cgh);
-    accessor<int, 1, mode::read_write, target::local> LocalAcc(Range, Cgh);
+    local_accessor<int, 1> LocalAcc(Range, Cgh);
+    accessor<int, 1, mode::read_write, target::local> LocalAccDep(Range, Cgh);
 
     Cgh.single_task<class test>([=]() {
       static_assert(std::is_same<decltype(DeviceRWAcc[0]), int &>::value,
@@ -46,6 +47,8 @@ int main() {
                     "Incorrect type from constant accessor");
       static_assert(std::is_same<decltype(LocalAcc[0]), int &>::value,
                     "Incorrect type from local accessor");
+      static_assert(std::is_same<decltype(LocalAccDep[0]), int &>::value,
+                    "Incorrect type from access target::local");
     });
   });
 }
diff --git a/sycl/test/basic_tests/set_arg_error.cpp b/sycl/test/basic_tests/set_arg_error.cpp
index c7ea5be42808e..ccab1a8fce1ad 100644
--- a/sycl/test/basic_tests/set_arg_error.cpp
+++ b/sycl/test/basic_tests/set_arg_error.cpp
@@ -1,3 +1,4 @@
+// RUN: %clangxx %fsycl-host-only -DUSE_DEPRECATED_LOCAL_ACC -fsyntax-only -Xclang -verify -Xclang -verify-ignore-unexpected=note,warning %s
 // RUN: %clangxx %fsycl-host-only -fsyntax-only -Xclang -verify -Xclang -verify-ignore-unexpected=note,warning %s
 
 #include <CL/sycl.hpp>
@@ -30,9 +31,13 @@ int main() {
     cl::sycl::sampler samp(cl::sycl::coordinate_normalization_mode::normalized,
                            cl::sycl::addressing_mode::clamp,
                            cl::sycl::filtering_mode::nearest);
+#ifdef USE_DEPRECATED_LOCAL_ACC
     cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write,
                        cl::sycl::access::target::local>
         local_acc({size}, h);
+#else
+    cl::sycl::local_accessor<int, 1> local_acc({size}, h);
+#endif
     TriviallyCopyable tc{1, 2};
     NonTriviallyCopyable ntc;
     h.set_arg(0, local_acc);
diff --git a/sycl/test/check_device_code/kernel_arguments_as.cpp b/sycl/test/check_device_code/kernel_arguments_as.cpp
index fbbecaf891d8d..9fce24301ca5a 100644
--- a/sycl/test/check_device_code/kernel_arguments_as.cpp
+++ b/sycl/test/check_device_code/kernel_arguments_as.cpp
@@ -1,5 +1,12 @@
+// RUN: %clangxx -DUSE_DEPRECATED_LOCAL_ACC -fsycl-device-only -Xclang -fsycl-is-device -emit-llvm %s -S -o %t.ll -I %sycl_include -Wno-sycl-strict -Xclang -verify-ignore-unexpected=note,warning -Xclang -disable-llvm-passes
+// RUN: FileCheck %s --input-file %t.ll --check-prefixes=CHECK,CHECK-DISABLE
+//
 // RUN: %clangxx -fsycl-device-only -Xclang -fsycl-is-device -emit-llvm %s -S -o %t.ll -I %sycl_include -Wno-sycl-strict -Xclang -verify-ignore-unexpected=note,warning -Xclang -disable-llvm-passes
 // RUN: FileCheck %s --input-file %t.ll --check-prefixes=CHECK,CHECK-DISABLE
+//
+// RUN: %clangxx -DUSE_DPRECATED_LOCAL_ACC -fsycl-device-only -Xclang -fsycl-is-device -emit-llvm %s -S -o %t.ll -I %sycl_include -Wno-sycl-strict -Xclang -verify-ignore-unexpected=note,warning -Xclang -disable-llvm-passes -D__ENABLE_USM_ADDR_SPACE__
+// RUN: FileCheck %s --input-file %t.ll --check-prefixes=CHECK,CHECK-ENABLE
+//
 // RUN: %clangxx -fsycl-device-only -Xclang -fsycl-is-device -emit-llvm %s -S -o %t.ll -I %sycl_include -Wno-sycl-strict -Xclang -verify-ignore-unexpected=note,warning -Xclang -disable-llvm-passes -D__ENABLE_USM_ADDR_SPACE__
 // RUN: FileCheck %s --input-file %t.ll --check-prefixes=CHECK,CHECK-ENABLE
 //
@@ -31,9 +38,13 @@ int main() {
                                  {cl::sycl::property::buffer::use_host_ptr()});
     queue.submit([&](cl::sycl::handler &cgh) {
       auto acc = buf.get_access<cl::sycl::access::mode::read_write>(cgh);
+#ifdef USE_DEPRECATED_LOCAL_ACC
       cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write,
                          cl::sycl::access::target::local>
           local_acc(cl::sycl::range<1>(10), cgh);
+#else
+      cl::sycl::local_accessor<int, 1> local_acc(cl::sycl::range<1>(10), cgh);
+#endif // USE_DEPRECATED_LOCAL_ACC
       auto acc_wrapped = AccWrapper<decltype(acc)>{acc};
       auto local_acc_wrapped = AccWrapper<decltype(local_acc)>{local_acc};
       cgh.parallel_for<class check_adress_space>(
diff --git a/sycl/test/esimd/simd_copy_to_copy_from.cpp b/sycl/test/esimd/simd_copy_to_copy_from.cpp
index 8d636b951e3f0..489bd32c81757 100644
--- a/sycl/test/esimd/simd_copy_to_copy_from.cpp
+++ b/sycl/test/esimd/simd_copy_to_copy_from.cpp
@@ -37,9 +37,7 @@ SYCL_EXTERNAL void kernel2(int *ptr) SYCL_ESIMD_FUNCTION {
 // --- Negative tests.
 
 // Incompatible target.
-SYCL_EXTERNAL void
-kernel3(accessor<int, 1, access::mode::read_write, access::target::local> &buf)
-    SYCL_ESIMD_FUNCTION {
+SYCL_EXTERNAL void kernel3(local_accessor<int, 1> &buf) SYCL_ESIMD_FUNCTION {
   simd<int, 32> v1(0, 1);
   simd<int, 32> v0;
   // CHECK: simd_copy_to_copy_from.cpp:46{{.*}}error: no matching member function for call to 'copy_from'
diff --git a/sycl/test/extensions/sub_group_as.cpp b/sycl/test/extensions/sub_group_as.cpp
index 43e79ac5aeb9a..566f02e63c7a6 100644
--- a/sycl/test/extensions/sub_group_as.cpp
+++ b/sycl/test/extensions/sub_group_as.cpp
@@ -26,9 +26,7 @@ int main(int argc, char *argv[]) {
     queue.submit([&](cl::sycl::handler &cgh) {
       auto global = buf.get_access<cl::sycl::access::mode::read_write,
                                    cl::sycl::access::target::device>(cgh);
-      sycl::accessor<int, 1, sycl::access::mode::read_write,
-                     sycl::access::target::local>
-          local(N, cgh);
+      sycl::local_accessor<int, 1> local(N, cgh);
 
       cgh.parallel_for<class test>(
           cl::sycl::nd_range<1>(N, 32), [=](cl::sycl::nd_item<1> it) {
diff --git a/sycl/test/multi_ptr/ctad.cpp b/sycl/test/multi_ptr/ctad.cpp
index a37c94fda737a..b37a4702c1f01 100644
--- a/sycl/test/multi_ptr/ctad.cpp
+++ b/sycl/test/multi_ptr/ctad.cpp
@@ -18,11 +18,13 @@ int main() {
   using deviceAcc = sycl::accessor<int, 1, mode::read, target::device>;
   using globlAcc = sycl::accessor<int, 1, mode::read, target::global_buffer>;
   using constAcc = sycl::accessor<int, 1, mode::read, target::constant_buffer>;
-  using localAcc = sycl::accessor<int, 1, mode::read, target::local>;
+  using localAcc = sycl::local_accessor<int, 1>;
+  using localAccDep = sycl::accessor<int, 1, mode::read, target::local>;
   using deviceCTAD = decltype(sycl::multi_ptr(std::declval<deviceAcc>()));
   using globlCTAD = decltype(sycl::multi_ptr(std::declval<globlAcc>()));
   using constCTAD = decltype(sycl::multi_ptr(std::declval<constAcc>()));
   using localCTAD = decltype(sycl::multi_ptr(std::declval<localAcc>()));
+  using localCTADDep = decltype(sycl::multi_ptr(std::declval<localAccDep>()));
   using deviceMPtr = sycl::multi_ptr<int, address_space::global_space>;
   using globlMPtr = sycl::multi_ptr<int, address_space::global_space>;
   using constMPtr = sycl::multi_ptr<int, address_space::constant_space>;
@@ -32,4 +34,5 @@ int main() {
   static_assert(std::is_same<globlCTAD, globlMPtr>::value);
   static_assert(std::is_same<constCTAD, constMPtr>::value);
   static_assert(std::is_same<localCTAD, localMPtr>::value);
+  static_assert(std::is_same<localCTADDep, localMPtr>::value);
 }

From aec26e8ab66af4a422184e2f61ac0ef1cfc4212b Mon Sep 17 00:00:00 2001
From: "aidan.belton" <aidan.belton@codeplay.com>
Date: Wed, 22 Jun 2022 11:34:07 +0100
Subject: [PATCH 02/13] Update simd_copy_to_copy_from.cpp line numbers

---
 sycl/test/esimd/simd_copy_to_copy_from.cpp | 8 ++++----
 1 file changed, 4 insertions(+), 4 deletions(-)

diff --git a/sycl/test/esimd/simd_copy_to_copy_from.cpp b/sycl/test/esimd/simd_copy_to_copy_from.cpp
index 489bd32c81757..279ca9290d324 100644
--- a/sycl/test/esimd/simd_copy_to_copy_from.cpp
+++ b/sycl/test/esimd/simd_copy_to_copy_from.cpp
@@ -40,10 +40,10 @@ SYCL_EXTERNAL void kernel2(int *ptr) SYCL_ESIMD_FUNCTION {
 SYCL_EXTERNAL void kernel3(local_accessor<int, 1> &buf) SYCL_ESIMD_FUNCTION {
   simd<int, 32> v1(0, 1);
   simd<int, 32> v0;
-  // CHECK: simd_copy_to_copy_from.cpp:46{{.*}}error: no matching member function for call to 'copy_from'
+  // CHECK: simd_copy_to_copy_from.cpp:44{{.*}}error: no matching member function for call to 'copy_from'
   v0.copy_from(buf, 0);
   v0 = v0 + v1;
-  // CHECK: simd_copy_to_copy_from.cpp:49{{.*}}error: no matching member function for call to 'copy_to'
+  // CHECK: simd_copy_to_copy_from.cpp:47{{.*}}error: no matching member function for call to 'copy_to'
   v0.copy_to(buf, 0);
 }
 
@@ -52,7 +52,7 @@ SYCL_EXTERNAL void
 kernel4(accessor<int, 1, access::mode::write, access::target::device> &buf)
     SYCL_ESIMD_FUNCTION {
   simd<int, 32> v;
-  // CHECK: simd_copy_to_copy_from.cpp:58{{.*}}error: no matching member function for call to 'copy_from'
+  // CHECK: simd_copy_to_copy_from.cpp:56{{.*}}error: no matching member function for call to 'copy_from'
   v.copy_from(buf, 0);
 }
 
@@ -61,6 +61,6 @@ SYCL_EXTERNAL void
 kernel5(accessor<int, 1, access::mode::read, access::target::device> &buf)
     SYCL_ESIMD_FUNCTION {
   simd<int, 32> v(0, 1);
-  // CHECK: simd_copy_to_copy_from.cpp:67{{.*}}error: no matching member function for call to 'copy_to'
+  // CHECK: simd_copy_to_copy_from.cpp:65{{.*}}error: no matching member function for call to 'copy_to'
   v.copy_to(buf, 0);
 }

From 712664ea44e6c272ed10fe74c4eb824bd412a089 Mon Sep 17 00:00:00 2001
From: "aidan.belton" <aidan.belton@codeplay.com>
Date: Tue, 9 Aug 2022 13:29:33 +0100
Subject: [PATCH 03/13] Create local_accessor_base, make access target::local
 and local_accessor children

---
 sycl/include/sycl/access/access.hpp           |   2 +-
 sycl/include/sycl/accessor.hpp                | 116 ++++++++++++++----
 sycl/include/sycl/handler.hpp                 |   5 +
 sycl/include/sycl/multi_ptr.hpp               |  44 ++++++-
 sycl/test/abi/layout_accessors_device.cpp     |  47 +++++--
 sycl/test/abi/layout_accessors_host.cpp       |  39 ++++--
 sycl/test/abi/user_mangling.cpp               |   6 +
 sycl/test/basic_tests/set_arg_error.cpp       |   1 +
 sycl/test/warnings/sycl_2020_deprecations.cpp |   6 +
 9 files changed, 221 insertions(+), 45 deletions(-)

diff --git a/sycl/include/sycl/access/access.hpp b/sycl/include/sycl/access/access.hpp
index 8efc8b55567dd..8830f4c65d409 100644
--- a/sycl/include/sycl/access/access.hpp
+++ b/sycl/include/sycl/access/access.hpp
@@ -17,7 +17,7 @@ namespace access {
 enum class target {
   global_buffer __SYCL2020_DEPRECATED("use 'target::device' instead") = 2014,
   constant_buffer = 2015,
-  local = 2016,
+  local __SYCL2020_DEPRECATED("use `local_accessor` instead") = 2016,
   image = 2017,
   host_buffer = 2018,
   host_image = 2019,
diff --git a/sycl/include/sycl/accessor.hpp b/sycl/include/sycl/accessor.hpp
index 07cda2d1eace7..36f48b3812f2d 100644
--- a/sycl/include/sycl/accessor.hpp
+++ b/sycl/include/sycl/accessor.hpp
@@ -300,12 +300,13 @@ class accessor_common {
   using ConstRefType = const DataT &;
   using PtrType = detail::const_if_const_AS<AS, DataT> *;
 
-  using AccType = accessor<DataT, Dimensions, AccessMode, AccessTarget,
-                           IsPlaceholder, PropertyListT>;
-
   // The class which allows to access value of N dimensional accessor using N
   // subscript operators, e.g. accessor[2][2][3]
-  template <int SubDims> class AccessorSubscript {
+  template <int SubDims,
+            typename AccType =
+                accessor<DataT, Dimensions, AccessMode, AccessTarget,
+                         IsPlaceholder, PropertyListT>>
+  class AccessorSubscript {
     static constexpr int Dims = Dimensions;
 
     mutable id<Dims> MIDs;
@@ -2016,8 +2017,7 @@ accessor(buffer<DataT, Dimensions, AllocatorT>, handler, Type1, Type2, Type3,
 /// \ingroup sycl_api_acc
 template <typename DataT, int Dimensions, access::mode AccessMode,
           access::placeholder IsPlaceholder>
-class __SYCL_SPECIAL_CLASS accessor<DataT, Dimensions, AccessMode,
-                                    access::target::local, IsPlaceholder> :
+class __SYCL_SPECIAL_CLASS local_accessor_base :
 #ifndef __SYCL_DEVICE_ONLY__
     public detail::LocalAccessorBaseHost,
 #endif
@@ -2034,7 +2034,9 @@ class __SYCL_SPECIAL_CLASS accessor<DataT, Dimensions, AccessMode,
   using AccessorCommonT::IsAccessAnyWrite;
   template <int Dims>
   using AccessorSubscript =
-      typename AccessorCommonT::template AccessorSubscript<Dims>;
+      typename AccessorCommonT::template AccessorSubscript<
+          Dims,
+          local_accessor_base<DataT, Dimensions, AccessMode, IsPlaceholder>>;
 
   using ConcreteASPtrType = typename detail::DecoratedType<DataT, AS>::type *;
 
@@ -2057,7 +2059,7 @@ class __SYCL_SPECIAL_CLASS accessor<DataT, Dimensions, AccessMode,
 
 public:
   // Default constructor for objects later initialized with __init member.
-  accessor()
+  local_accessor_base()
       : impl(detail::InitializedVal<AdjustedDim, range>::template get<0>()) {}
 
 protected:
@@ -2091,8 +2093,8 @@ class __SYCL_SPECIAL_CLASS accessor<DataT, Dimensions, AccessMode,
   using const_reference = const DataT &;
 
   template <int Dims = Dimensions, typename = detail::enable_if_t<Dims == 0>>
-  accessor(handler &, const detail::code_location CodeLoc =
-                          detail::code_location::current())
+  local_accessor_base(handler &, const detail::code_location CodeLoc =
+                                     detail::code_location::current())
 #ifdef __SYCL_DEVICE_ONLY__
       : impl(range<AdjustedDim>{1}){}
 #else
@@ -2104,9 +2106,9 @@ class __SYCL_SPECIAL_CLASS accessor<DataT, Dimensions, AccessMode,
 
         template <int Dims = Dimensions,
                   typename = detail::enable_if_t<Dims == 0>>
-        accessor(handler &, const property_list &propList,
-                 const detail::code_location CodeLoc =
-                     detail::code_location::current())
+        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;
@@ -2120,7 +2122,7 @@ class __SYCL_SPECIAL_CLASS accessor<DataT, Dimensions, AccessMode,
 #endif
 
   template <int Dims = Dimensions, typename = detail::enable_if_t<(Dims > 0)>>
-  accessor(
+  local_accessor_base(
       range<Dimensions> AllocationSize, handler &,
       const detail::code_location CodeLoc = detail::code_location::current())
 #ifdef __SYCL_DEVICE_ONLY__
@@ -2135,10 +2137,10 @@ class __SYCL_SPECIAL_CLASS accessor<DataT, Dimensions, AccessMode,
 
         template <int Dims = Dimensions,
                   typename = detail::enable_if_t<(Dims > 0)>>
-        accessor(range<Dimensions> AllocationSize, handler &,
-                 const property_list &propList,
-                 const detail::code_location CodeLoc =
-                     detail::code_location::current())
+        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;
@@ -2206,7 +2208,9 @@ class __SYCL_SPECIAL_CLASS accessor<DataT, Dimensions, AccessMode,
   }
 
   template <int Dims = Dimensions, typename = detail::enable_if_t<(Dims > 1)>>
-  typename AccessorCommonT::template AccessorSubscript<Dims - 1>
+  typename AccessorCommonT::template AccessorSubscript<
+      Dims - 1,
+      local_accessor_base<DataT, Dimensions, AccessMode, IsPlaceholder>>
   operator[](size_t Index) const {
     return AccessorSubscript<Dims - 1>(*this, Index);
   }
@@ -2215,8 +2219,78 @@ class __SYCL_SPECIAL_CLASS accessor<DataT, Dimensions, AccessMode,
     return local_ptr<DataT>(getQualifiedPtr());
   }
 
-  bool operator==(const accessor &Rhs) const { return impl == Rhs.impl; }
-  bool operator!=(const accessor &Rhs) const { return !(*this == Rhs); }
+  bool operator==(const local_accessor_base &Rhs) const {
+    return impl == Rhs.impl;
+  }
+  bool operator!=(const local_accessor_base &Rhs) const {
+    return !(*this == Rhs);
+  }
+};
+
+template <typename DataT, int Dimensions, access::mode AccessMode,
+          access::placeholder IsPlaceholder>
+class __SYCL_SPECIAL_CLASS accessor<DataT, Dimensions, AccessMode,
+                                    access::target::local, IsPlaceholder>
+    : public local_accessor_base<DataT, Dimensions, AccessMode, IsPlaceholder> {
+
+  using local_acc =
+      local_accessor_base<DataT, Dimensions, AccessMode, IsPlaceholder>;
+
+  // Use base classes constructors
+  using local_acc::local_acc;
+
+#ifdef __SYCL_DEVICE_ONLY__
+
+  // __init needs to be defined within the class not through inheritance.
+  // Map this function to inherited func.
+  void __init(typename local_acc::ConcreteASPtrType Ptr,
+              range<local_acc::AdjustedDim> AccessRange,
+              range<local_acc::AdjustedDim> range,
+              id<local_acc::AdjustedDim> id) {
+    local_acc::__init(Ptr, AccessRange, range, id);
+  }
+
+public:
+  // Default constructor for objects later initialized with __init member.
+  accessor() {
+    local_acc::impl = detail::InitializedVal<local_acc::AdjustedDim,
+                                             range>::template get<0>();
+  }
+
+#endif
+};
+
+template <typename DataT, int Dimensions = 1>
+class __SYCL_SPECIAL_CLASS local_accessor
+    : public local_accessor_base<DataT, Dimensions, access::mode::read_write,
+                                 access::placeholder::false_t> {
+
+  using local_acc =
+      local_accessor_base<DataT, Dimensions, access::mode::read_write,
+                          access::placeholder::false_t>;
+
+  // Use base classes constructors
+  using local_acc::local_acc;
+
+#ifdef __SYCL_DEVICE_ONLY__
+
+  // __init needs to be defined within the class not through inheritance.
+  // Map this function to inherited func.
+  void __init(typename local_acc::ConcreteASPtrType Ptr,
+              range<local_acc::AdjustedDim> AccessRange,
+              range<local_acc::AdjustedDim> range,
+              id<local_acc::AdjustedDim> id) {
+    local_acc::__init(Ptr, AccessRange, range, id);
+  }
+
+public:
+  // Default constructor for objects later initialized with __init member.
+  local_accessor() {
+    local_acc::impl = detail::InitializedVal<local_acc::AdjustedDim,
+                                             range>::template get<0>();
+  }
+
+#endif
 };
 
 /// Image accessors.
diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp
index 6425d59a9b5fe..1dde2144172bb 100644
--- a/sycl/include/sycl/handler.hpp
+++ b/sycl/include/sycl/handler.hpp
@@ -1445,6 +1445,11 @@ class __SYCL_EXPORT handler {
     setArgHelper(ArgIndex, std::move(Arg));
   }
 
+  template <typename DataT, int Dims>
+  void set_arg(int ArgIndex, local_accessor<DataT, Dims> Arg) {
+    setArgHelper(ArgIndex, std::move(Arg));
+  }
+
   /// Sets arguments for OpenCL interoperability kernels.
   ///
   /// Registers pack of arguments(Args) with indexes starting from 0.
diff --git a/sycl/include/sycl/multi_ptr.hpp b/sycl/include/sycl/multi_ptr.hpp
index 325db5b7dacfd..940ccaa56d93d 100644
--- a/sycl/include/sycl/multi_ptr.hpp
+++ b/sycl/include/sycl/multi_ptr.hpp
@@ -16,11 +16,12 @@
 
 __SYCL_INLINE_NAMESPACE(cl) {
 namespace sycl {
-// Forward declaration
+// Forward declarations
 template <typename dataT, int dimensions, access::mode accessMode,
           access::target accessTarget, access::placeholder isPlaceholder,
           typename PropertyListT>
 class accessor;
+template <typename dataT, int dimensions> class local_accessor;
 
 /// Provides constructors for address space qualified and non address space
 /// qualified pointers to allow interoperability between plain C++ and OpenCL C.
@@ -151,6 +152,11 @@ template <typename ElementType, access::address_space Space> class multi_ptr {
                 Accessor)
       : multi_ptr(Accessor.get_pointer()) {}
 
+  // Only if Space == local_space || generic_space
+  template <int dimensions>
+  multi_ptr(local_accessor<ElementType, dimensions> Accessor)
+      : multi_ptr(Accessor.get_pointer()) {}
+
   // Only if Space == constant_space
   template <
       int dimensions, access::mode Mode, access::placeholder isPlaceholder,
@@ -204,6 +210,19 @@ template <typename ElementType, access::address_space Space> class multi_ptr {
                 Accessor)
       : multi_ptr(Accessor.get_pointer()) {}
 
+  // Only if Space == local_space || generic_space and element type is const
+  template <
+      int dimensions, access::address_space _Space = Space,
+      typename ET = ElementType,
+      typename = typename detail::enable_if_t<
+          _Space == Space &&
+          (Space == access::address_space::generic_space ||
+           Space == access::address_space::local_space) &&
+          std::is_const<ET>::value && std::is_same<ET, ElementType>::value>>
+  multi_ptr(
+      local_accessor<typename detail::remove_const_t<ET>, dimensions> Accessor)
+      : multi_ptr(Accessor.get_pointer()) {}
+
   // Only if Space == constant_space and element type is const
   template <
       int dimensions, access::mode Mode, access::placeholder isPlaceholder,
@@ -422,6 +441,16 @@ template <access::address_space Space> class multi_ptr<void, Space> {
                 Accessor)
       : multi_ptr(Accessor.get_pointer()) {}
 
+  // Only if Space == local_space || generic_space
+  template <
+      typename ElementType, int dimensions,
+      access::address_space _Space = Space,
+      typename = typename detail::enable_if_t<
+          _Space == Space && (Space == access::address_space::generic_space ||
+                              Space == access::address_space::local_space)>>
+  multi_ptr(local_accessor<ElementType, dimensions> Accessor)
+      : multi_ptr(Accessor.get_pointer()) {}
+
   // Only if Space == constant_space
   template <
       typename ElementType, int dimensions, access::mode Mode,
@@ -546,6 +575,16 @@ template <access::address_space Space> class multi_ptr<const void, Space> {
                 Accessor)
       : multi_ptr(Accessor.get_pointer()) {}
 
+  // Only if Space == local_space || generic_space
+  template <
+      typename ElementType, int dimensions,
+      access::address_space _Space = Space,
+      typename = typename detail::enable_if_t<
+          _Space == Space && (Space == access::address_space::generic_space ||
+                              Space == access::address_space::local_space)>>
+  multi_ptr(local_accessor<ElementType, dimensions> Accessor)
+      : multi_ptr(Accessor.get_pointer()) {}
+
   // Only if Space == constant_space
   template <
       typename ElementType, int dimensions, access::mode Mode,
@@ -597,6 +636,9 @@ template <int dimensions, access::mode Mode, access::placeholder isPlaceholder,
 multi_ptr(accessor<T, dimensions, Mode, access::target::local, isPlaceholder,
                    PropertyListT>)
     -> multi_ptr<T, access::address_space::local_space>;
+template <int dimensions, class T>
+multi_ptr(local_accessor<T, dimensions>)
+    -> multi_ptr<T, access::address_space::local_space>;
 #endif
 
 template <typename ElementType, access::address_space Space>
diff --git a/sycl/test/abi/layout_accessors_device.cpp b/sycl/test/abi/layout_accessors_device.cpp
index 0519a7bedb040..fa5c5346efd0c 100644
--- a/sycl/test/abi/layout_accessors_device.cpp
+++ b/sycl/test/abi/layout_accessors_device.cpp
@@ -39,19 +39,42 @@ SYCL_EXTERNAL void hostAcc(accessor<int, 1, access::mode::read, access::target::
 SYCL_EXTERNAL void hostAcc(accessor<int, 1, access::mode::read_write, access::target::local> Acc) {
   (void)Acc.get_size();
 }
+
 // CHECK:      0 | class sycl::accessor<int, 1, sycl::access::mode::read_write, sycl::access::target::local, sycl::access::placeholder::false_t>
-// CHECK-NEXT: 0 |   class sycl::detail::accessor_common<int, 1, sycl::access::mode::read_write, sycl::access::target::local, sycl::access::placeholder::false_t> (base) (empty)
-// CHECK-NEXT: 0 |   class sycl::detail::LocalAccessorBaseDevice<1> impl
-// CHECK-NEXT: 0 |     class sycl::range<1> AccessRange
-// CHECK-NEXT: 0 |       class sycl::detail::array<1> (base)
-// CHECK-NEXT: 0 |         size_t[1] common_array
-// CHECK-NEXT: 8 |     class sycl::range<1> MemRange
-// CHECK-NEXT: 8 |       class sycl::detail::array<1> (base)
-// CHECK-NEXT: 8 |         size_t[1] common_array
-// CHECK-NEXT: 16 |     class sycl::id<1> Offset
-// CHECK-NEXT: 16 |       class sycl::detail::array<1> (base)
-// CHECK-NEXT: 16 |         size_t[1] common_array
-// CHECK-NEXT: 24 |   sycl::accessor<int, 1, sycl::access::mode::read_write, sycl::access::target::local, sycl::access::placeholder::false_t>::ConcreteASPtrType MData
+// CHECK-NEXT: 0 |   class sycl::local_accessor_base<int, 1, sycl::access::mode::read_write, sycl::access::placeholder::false_t> (base)
+// CHECK-NEXT: 0 |     class sycl::detail::accessor_common<int, 1, sycl::access::mode::read_write, sycl::access::target::local, sycl::access::placeholder::false_t> (base) (empty)
+// CHECK-NEXT: 0 |     class sycl::detail::LocalAccessorBaseDevice<1> impl
+// CHECK-NEXT: 0 |       class sycl::range<1> AccessRange
+// CHECK-NEXT: 0 |         class sycl::detail::array<1> (base)
+// CHECK-NEXT: 0 |           size_t[1] common_array
+// CHECK-NEXT: 8 |       class sycl::range<1> MemRange
+// CHECK-NEXT: 8 |         class sycl::detail::array<1> (base)
+// CHECK-NEXT: 8 |           size_t[1] common_array
+// CHECK-NEXT: 16 |       class sycl::id<1> Offset
+// CHECK-NEXT: 16 |         class sycl::detail::array<1> (base)
+// CHECK-NEXT: 16 |           size_t[1] common_array
+// CHECK-NEXT: 24 |     sycl::local_accessor_base<int, 1, sycl::access::mode::read_write, sycl::access::placeholder::false_t>::ConcreteASPtrType MData
+// CHECK-NEXT: | [sizeof=32, dsize=32, align=8,
+// CHECK-NEXT: |  nvsize=32, nvalign=8]
+
+SYCL_EXTERNAL void hostAcc(local_accessor<int, 1> Acc) {
+  (void)Acc.get_size();
+}
+
+// CHECK:      0 | class sycl::local_accessor<int, 1>
+// CHECK-NEXT: 0 |   class sycl::local_accessor_base<int, 1, sycl::access::mode::read_write, sycl::access::placeholder::false_t> (base)
+// CHECK-NEXT: 0 |     class sycl::detail::accessor_common<int, 1, sycl::access::mode::read_write, sycl::access::target::local, sycl::access::placeholder::false_t> (base) (empty)
+// CHECK-NEXT: 0 |     class sycl::detail::LocalAccessorBaseDevice<1> impl
+// CHECK-NEXT: 0 |       class sycl::range<1> AccessRange
+// CHECK-NEXT: 0 |         class sycl::detail::array<1> (base)
+// CHECK-NEXT: 0 |           size_t[1] common_array
+// CHECK-NEXT: 8 |       class sycl::range<1> MemRange
+// CHECK-NEXT: 8 |         class sycl::detail::array<1> (base)
+// CHECK-NEXT: 8 |           size_t[1] common_array
+// CHECK-NEXT: 16 |       class sycl::id<1> Offset
+// CHECK-NEXT: 16 |         class sycl::detail::array<1> (base)
+// CHECK-NEXT: 16 |           size_t[1] common_array
+// CHECK-NEXT: 24 |     sycl::local_accessor_base<int, 1, sycl::access::mode::read_write, sycl::access::placeholder::false_t>::ConcreteASPtrType MData
 // CHECK-NEXT: | [sizeof=32, dsize=32, align=8,
 // CHECK-NEXT: |  nvsize=32, nvalign=8]
 
diff --git a/sycl/test/abi/layout_accessors_host.cpp b/sycl/test/abi/layout_accessors_host.cpp
index 2110cd565cccf..7a854c05a9185 100644
--- a/sycl/test/abi/layout_accessors_host.cpp
+++ b/sycl/test/abi/layout_accessors_host.cpp
@@ -98,16 +98,35 @@ void hostAcc(accessor<int, 1, access::mode::read_write, access::target::local> A
   (void)Acc.get_size();
 }
 
-// CHECK: 0      | class sycl::accessor<int, 1, sycl::access::mode::read_write, sycl::access::target::local, sycl::access::placeholder::false_t>
-// CHECK-NEXT: 0 |   class sycl::detail::LocalAccessorBaseHost (base)
-// CHECK-NEXT: 0 |     class std::shared_ptr<class sycl::detail::LocalAccessorImplHost> impl
-// CHECK-NEXT: 0 |       class std::__shared_ptr<class sycl::detail::LocalAccessorImplHost, __gnu_cxx::_S_atomic> (base)
-// CHECK-NEXT: 0 |         class std::__shared_ptr_access<class sycl::detail::LocalAccessorImplHost, __gnu_cxx::_S_atomic, false, false> (base) (empty)
-// CHECK-NEXT: 0 |         std::__shared_ptr<class sycl::detail::LocalAccessorImplHost, __gnu_cxx::_S_atomic>::element_type * _M_ptr
-// CHECK-NEXT: 8 |         class std::__shared_count<__gnu_cxx::_S_atomic> _M_refcount
-// CHECK-NEXT: 8 |           _Sp_counted_base<(enum __gnu_cxx::_Lock_policy)2U> * _M_pi
-// CHECK-NEXT: 0 |   class sycl::detail::accessor_common<int, 1, sycl::access::mode::read_write, sycl::access::target::local, sycl::access::placeholder::false_t> (base) (empty)
-// CHECK-NEXT: 16 |   char[16] padding
+// CHECK: 0 | class sycl::accessor<int, 1, sycl::access::mode::read_write, sycl::access::target::local, sycl::access::placeholder::false_t>
+// CHECK-NEXT: 0 |   class sycl::local_accessor_base<int, 1, sycl::access::mode::read_write, sycl::access::placeholder::false_t> (base)
+// CHECK-NEXT: 0 |     class sycl::detail::LocalAccessorBaseHost (base)
+// CHECK-NEXT: 0 |       class std::shared_ptr<class sycl::detail::LocalAccessorImplHost> impl
+// CHECK-NEXT: 0 |         class std::__shared_ptr<class sycl::detail::LocalAccessorImplHost, __gnu_cxx::_S_atomic> (base)
+// CHECK-NEXT: 0 |           class std::__shared_ptr_access<class sycl::detail::LocalAccessorImplHost, __gnu_cxx::_S_atomic, false, false> (base) (empty)
+// CHECK-NEXT: 0 |           std::__shared_ptr<class sycl::detail::LocalAccessorImplHost, __gnu_cxx::_S_atomic>::element_type * _M_ptr
+// CHECK-NEXT: 8 |           class std::__shared_count<__gnu_cxx::_S_atomic> _M_refcount
+// CHECK-NEXT: 8 |             _Sp_counted_base<(enum __gnu_cxx::_Lock_policy)2U> * _M_pi
+// CHECK-NEXT: 0 |     class sycl::detail::accessor_common<int, 1, sycl::access::mode::read_write, sycl::access::target::local, sycl::access::placeholder::false_t> (base) (empty)
+// CHECK-NEXT: 16 |     char[16] padding
+// CHECK-NEXT: [sizeof=32, dsize=32, align=8,
+// CHECK-NEXT: nvsize=32, nvalign=8]
+
+void hostAcc(local_accessor<int, 1> Acc) {
+  (void)Acc.get_size();
+}
+
+// CHECK: 0 | class sycl::local_accessor<int, 1>
+// CHECK-NEXT: 0 |   class sycl::local_accessor_base<int, 1, sycl::access::mode::read_write, sycl::access::placeholder::false_t> (base)
+// CHECK-NEXT: 0 |     class sycl::detail::LocalAccessorBaseHost (base)
+// CHECK-NEXT: 0 |       class std::shared_ptr<class sycl::detail::LocalAccessorImplHost> impl
+// CHECK-NEXT: 0 |         class std::__shared_ptr<class sycl::detail::LocalAccessorImplHost, __gnu_cxx::_S_atomic> (base)
+// CHECK-NEXT: 0 |           class std::__shared_ptr_access<class sycl::detail::LocalAccessorImplHost, __gnu_cxx::_S_atomic, false, false> (base) (empty)
+// CHECK-NEXT: 0 |           std::__shared_ptr<class sycl::detail::LocalAccessorImplHost, __gnu_cxx::_S_atomic>::element_type * _M_ptr
+// CHECK-NEXT: 8 |           class std::__shared_count<__gnu_cxx::_S_atomic> _M_refcount
+// CHECK-NEXT: 8 |             _Sp_counted_base<(enum __gnu_cxx::_Lock_policy)2U> * _M_pi
+// CHECK-NEXT: 0 |     class sycl::detail::accessor_common<int, 1, sycl::access::mode::read_write, sycl::access::target::local, sycl::access::placeholder::false_t> (base) (empty)
+// CHECK-NEXT: 16 |     char[16] padding
 // CHECK-NEXT: [sizeof=32, dsize=32, align=8,
 // CHECK-NEXT: nvsize=32, nvalign=8]
 
diff --git a/sycl/test/abi/user_mangling.cpp b/sycl/test/abi/user_mangling.cpp
index af66df4569f40..abb194f42a719 100644
--- a/sycl/test/abi/user_mangling.cpp
+++ b/sycl/test/abi/user_mangling.cpp
@@ -16,6 +16,9 @@ SYCL_EXTERNAL void acc(sycl::accessor<int, 1, sycl::access::mode::read, sycl::ac
 // CHK-DEVICE: define dso_local spir_func void @_Z3accN2cl4sycl8accessorIiLi1ELNS0_6access4modeE1024ELNS2_6targetE2016ELNS2_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEE({{.*}})
 SYCL_EXTERNAL void acc(sycl::accessor<int, 1, sycl::access::mode::read, sycl::access::target::local>) {} 
 
+// CHK_DEVICE: define dso_local void @_Z3accN2cl4sycl14local_accessorIiLi1EEE({{.*}})
+SYCL_EXTERNAL void acc(sycl::local_accessor<int, 1>) {}
+
 // CHK-DEVICE: define dso_local spir_func void @_Z3accN2cl4sycl8accessorINS0_3vecIiLi4EEELi1ELNS0_6access4modeE1024ELNS4_6targetE2017ELNS4_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEE({{.*}})
 SYCL_EXTERNAL void acc(sycl::accessor<sycl::cl_int4, 1, sycl::access::mode::read, sycl::access::target::image>) {} 
 
@@ -46,6 +49,9 @@ void acc(sycl::accessor<int, 1, sycl::access::mode::read, sycl::access::target::
 // CHK-HOST: define dso_local void @_Z3accN2cl4sycl8accessorIiLi1ELNS0_6access4modeE1024ELNS2_6targetE2016ELNS2_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEE({{.*}})
 void acc(sycl::accessor<int, 1, sycl::access::mode::read, sycl::access::target::local>) {} 
 
+// CHK-HOST: define dso_local void @_Z3accN2cl4sycl14local_accessorIiLi1EEE({{.*}})
+void acc(sycl::local_accessor<int, 1>) {}
+
 // CHK-HOST: define dso_local void @_Z3accN2cl4sycl8accessorINS0_3vecIiLi4EEELi1ELNS0_6access4modeE1024ELNS4_6targetE2019ELNS4_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEE({{.*}})
 void acc(sycl::accessor<sycl::cl_int4, 1, sycl::access::mode::read, sycl::access::target::host_image>) {} 
 
diff --git a/sycl/test/basic_tests/set_arg_error.cpp b/sycl/test/basic_tests/set_arg_error.cpp
index 9b4f7341a9046..6800714b914bf 100644
--- a/sycl/test/basic_tests/set_arg_error.cpp
+++ b/sycl/test/basic_tests/set_arg_error.cpp
@@ -38,6 +38,7 @@ int main() {
 #else
     cl::sycl::local_accessor<int, 1> local_acc({size}, h);
 #endif
+
     TriviallyCopyable tc{1, 2};
     NonTriviallyCopyable ntc;
     h.set_arg(0, local_acc);
diff --git a/sycl/test/warnings/sycl_2020_deprecations.cpp b/sycl/test/warnings/sycl_2020_deprecations.cpp
index c39db53cf9358..cb07977f19723 100644
--- a/sycl/test/warnings/sycl_2020_deprecations.cpp
+++ b/sycl/test/warnings/sycl_2020_deprecations.cpp
@@ -184,5 +184,11 @@ int main() {
   // expected-warning@+1{{'get_linear_id' is deprecated: use sycl::group::get_group_linear_id() instead}}
   group.get_linear_id();
 
+  // expected-warning@+1{{'local' is deprecated: use `local_accessor` instead}}
+  Queue.submit([&](sycl::handler &CGH) {
+    sycl::accessor<int, 1, sycl::access::mode::read_write, sycl::target::local>
+        LocalAcc(sycl::range<1>(1), CGH);
+  });
+
   return 0;
 }

From 3dbfd0ebb5ad44c5c66434a5538f054a3fb7c833 Mon Sep 17 00:00:00 2001
From: "aidan.belton" <aidan.belton@codeplay.com>
Date: Tue, 9 Aug 2022 17:33:22 +0100
Subject: [PATCH 04/13] Update SemaSYCL to recognize and handle local_accessor

---
 clang/lib/Sema/SemaSYCL.cpp                   | 60 +++++++++++++------
 clang/test/CodeGenSYCL/Inputs/sycl.hpp        | 20 +++++++
 .../kernel-arg-accessor-pointer.cpp           | 22 ++++++-
 .../no_opaque_kernel-arg-accessor-pointer.cpp | 21 ++++++-
 clang/test/SemaSYCL/Inputs/sycl.hpp           | 20 +++++++
 clang/test/SemaSYCL/accessors-targets.cpp     | 12 +++-
 6 files changed, 131 insertions(+), 24 deletions(-)

diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 5dbaed297dd69..c2870844398f5 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -111,6 +111,11 @@ class Util {
   /// \param Tmpl  whether the class is template instantiation or simple record
   static bool isSyclType(QualType Ty, StringRef Name, bool Tmpl = false);
 
+  /// Checks whether given clang type is a standard SYCL API accessor class.
+  /// \param Ty    the clang type being checked
+  /// \param Tmpl  whether the class is template instantiation or simple record
+  static bool isSyclAccessorType(QualType Ty, bool Tmpl = false);
+
   /// Checks whether given clang type is a full specialization of the SYCL
   /// specialization constant class.
   static bool isSyclSpecConstantType(QualType Ty);
@@ -1021,7 +1026,11 @@ static ParamDesc makeParamDesc(ASTContext &Ctx, StringRef Name, QualType Ty) {
 }
 
 /// \return the target of given SYCL accessor type
-static target getAccessTarget(const ClassTemplateSpecializationDecl *AccTy) {
+static target getAccessTarget(QualType FieldTy,
+                              const ClassTemplateSpecializationDecl *AccTy) {
+  if (Util::isSyclType(FieldTy, "local_accessor", true /*Tmpl*/))
+    return local;
+
   return static_cast<target>(
       AccTy->getTemplateArgs()[3].getAsIntegral().getExtValue());
 }
@@ -1615,7 +1624,7 @@ class SyclKernelFieldChecker : public SyclKernelFieldHandler {
     assert(Util::isSyclSpecialType(Ty) &&
            "Should only be called on sycl special class types.");
     const RecordDecl *RecD = Ty->getAsRecordDecl();
-    if (IsSIMD && !Util::isSyclType(Ty, "accessor", true /*Tmp*/))
+    if (IsSIMD && !Util::isSyclAccessorType(Ty, true /*Tmp*/))
       return SemaRef.Diag(Loc.getBegin(),
                           diag::err_sycl_esimd_not_supported_for_type)
              << RecD;
@@ -1927,17 +1936,25 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler {
   }
 
   // Additional processing is required for accessor type.
-  void handleAccessorType(const CXXRecordDecl *RecordDecl, SourceLocation Loc) {
+  void handleAccessorType(QualType FieldTy, const CXXRecordDecl *RecordDecl,
+                          SourceLocation Loc) {
     handleAccessorPropertyList(Params.back(), RecordDecl, Loc);
-    // Get access mode of accessor.
-    const auto *AccessorSpecializationDecl =
-        cast<ClassTemplateSpecializationDecl>(RecordDecl);
-    const TemplateArgument &AccessModeArg =
-        AccessorSpecializationDecl->getTemplateArgs().get(2);
+    bool isReadOnly = false;
+
+    // If "accessor" type check if read only
+    if (Util::isSyclType(FieldTy, "accessor", true /*Tmpl*/)) {
+      // Get access mode of accessor.
+      const auto *AccessorSpecializationDecl =
+          cast<ClassTemplateSpecializationDecl>(RecordDecl);
+      const TemplateArgument &AccessModeArg =
+          AccessorSpecializationDecl->getTemplateArgs().get(2);
+      if (isReadOnlyAccessor(AccessModeArg))
+        isReadOnly = true;
+    }
 
     // Add implicit attribute to parameter decl when it is a read only
     // SYCL accessor.
-    if (isReadOnlyAccessor(AccessModeArg))
+    if (isReadOnly)
       Params.back()->addAttr(
           SYCLAccessorReadonlyAttr::CreateImplicit(SemaRef.getASTContext()));
     Params.back()->addAttr(
@@ -1953,7 +1970,7 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler {
     assert(RecordDecl && "The type must be a RecordDecl");
     llvm::StringLiteral MethodName =
         KernelDecl->hasAttr<SYCLSimdAttr>() &&
-                Util::isSyclType(FieldTy, "accessor", true /*Tmp*/)
+                Util::isSyclAccessorType(FieldTy, true /*Tmp*/)
             ? InitESIMDMethodName
             : InitMethodName;
     CXXMethodDecl *InitMethod = getMethodByName(RecordDecl, MethodName);
@@ -1978,8 +1995,8 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler {
       // added, this code needs to be refactored to call
       // handleAccessorPropertyList for each class which requires it.
       if (ParamTy.getTypePtr()->isPointerType() &&
-          Util::isSyclType(FieldTy, "accessor", true /*Tmp*/))
-        handleAccessorType(RecordDecl, FD->getBeginLoc());
+          Util::isSyclAccessorType(FieldTy, true /*Tmp*/))
+        handleAccessorType(FieldTy, RecordDecl, FD->getBeginLoc());
     }
     LastParamIndex = ParamIndex;
     return true;
@@ -2074,7 +2091,7 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler {
     assert(RecordDecl && "The type must be a RecordDecl");
     llvm::StringLiteral MethodName =
         KernelDecl->hasAttr<SYCLSimdAttr>() &&
-                Util::isSyclType(FieldTy, "accessor", true /*Tmp*/)
+                Util::isSyclAccessorType(FieldTy, true /*Tmp*/)
             ? InitESIMDMethodName
             : InitMethodName;
     CXXMethodDecl *InitMethod = getMethodByName(RecordDecl, MethodName);
@@ -2093,8 +2110,8 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler {
       // added, this code needs to be refactored to call
       // handleAccessorPropertyList for each class which requires it.
       if (ParamTy.getTypePtr()->isPointerType() &&
-          Util::isSyclType(FieldTy, "accessor", true /*Tmp*/))
-        handleAccessorType(RecordDecl, BS.getBeginLoc());
+          Util::isSyclAccessorType(FieldTy, true /*Tmp*/))
+        handleAccessorType(FieldTy, RecordDecl, BS.getBeginLoc());
     }
     LastParamIndex = ParamIndex;
     return true;
@@ -2215,7 +2232,7 @@ class SyclKernelArgsSizeChecker : public SyclKernelFieldHandler {
     const CXXRecordDecl *RecordDecl = FieldTy->getAsCXXRecordDecl();
     assert(RecordDecl && "The type must be a RecordDecl");
     llvm::StringLiteral MethodName =
-        (IsSIMD && Util::isSyclType(FieldTy, "accessor", true /*Tmp*/))
+        (IsSIMD && Util::isSyclAccessorType(FieldTy, true /*Tmp*/))
             ? InitESIMDMethodName
             : InitMethodName;
     CXXMethodDecl *InitMethod = getMethodByName(RecordDecl, MethodName);
@@ -3124,7 +3141,7 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler {
            "Incorrect template args for Accessor Type");
     int Dims = static_cast<int>(
         AccTy->getTemplateArgs()[1].getAsIntegral().getExtValue());
-    int Info = getAccessTarget(AccTy) | (Dims << 11);
+    int Info = getAccessTarget(FieldTy, AccTy) | (Dims << 11);
     Header.addParamDesc(SYCLIntegrationHeader::kind_accessor, Info,
                         CurOffset +
                             offsetOf(RD, BC.getType()->getAsCXXRecordDecl()));
@@ -3134,14 +3151,14 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler {
   bool handleSyclSpecialType(FieldDecl *FD, QualType FieldTy) final {
     const auto *ClassTy = FieldTy->getAsCXXRecordDecl();
     assert(ClassTy && "Type must be a C++ record type");
-    if (Util::isSyclType(FieldTy, "accessor", true /*Tmp*/)) {
+    if (Util::isSyclAccessorType(FieldTy, true /*Tmp*/)) {
       const auto *AccTy =
           cast<ClassTemplateSpecializationDecl>(FieldTy->getAsRecordDecl());
       assert(AccTy->getTemplateArgs().size() >= 2 &&
              "Incorrect template args for Accessor Type");
       int Dims = static_cast<int>(
           AccTy->getTemplateArgs()[1].getAsIntegral().getExtValue());
-      int Info = getAccessTarget(AccTy) | (Dims << 11);
+      int Info = getAccessTarget(FieldTy, AccTy) | (Dims << 11);
 
       Header.addParamDesc(SYCLIntegrationHeader::kind_accessor, Info,
                           CurOffset + offsetOf(FD, FieldTy));
@@ -5195,6 +5212,11 @@ bool Util::isSyclType(QualType Ty, StringRef Name, bool Tmpl) {
   return matchQualifiedTypeName(Ty, Scopes);
 }
 
+bool Util::isSyclAccessorType(QualType Ty, bool Tmpl) {
+  return isSyclType(Ty, "accessor", Tmpl) ||
+         isSyclType(Ty, "local_accessor", Tmpl);
+}
+
 bool Util::isAccessorPropertyListType(QualType Ty) {
   std::array<DeclContextDesc, 5> Scopes = {
       Util::MakeDeclContextDesc(Decl::Kind::Namespace, "cl"),
diff --git a/clang/test/CodeGenSYCL/Inputs/sycl.hpp b/clang/test/CodeGenSYCL/Inputs/sycl.hpp
index 6ecbd802387cf..cef420afc4eb7 100644
--- a/clang/test/CodeGenSYCL/Inputs/sycl.hpp
+++ b/clang/test/CodeGenSYCL/Inputs/sycl.hpp
@@ -291,6 +291,26 @@ class accessor<dataT, dimensions, accessmode, access::target::host_image, access
   _ImageImplT<dimensions, accessmode, access::target::host_image> impl;
 };
 
+template <typename dataT, int dimensions>
+class __attribute__((sycl_special_class))
+local_accessor: public accessor<dataT,
+        dimensions, access::mode::read_write,
+        access::target::local> {
+public:
+  void use(void) const {}
+  template <typename... T>
+  void use(T... args) {}
+  template <typename... T>
+  void use(T... args) const {}
+  _ImplT<dimensions> impl;
+
+private:
+#ifdef __SYCL_DEVICE_ONLY__
+  void __init(__attribute__((opencl_local)) dataT *Ptr, range<dimensions> AccessRange,
+              range<dimensions> MemRange, id<dimensions> Offset) {}
+#endif
+};
+
 // TODO: Add support for image_array accessor.
 // template <typename dataT, int dimensions, access::mode accessmode>
 //class accessor<dataT, dimensions, accessmode, access::target::image_array, access::placeholder::false_t>
diff --git a/clang/test/CodeGenSYCL/kernel-arg-accessor-pointer.cpp b/clang/test/CodeGenSYCL/kernel-arg-accessor-pointer.cpp
index 1b7aeb55f3feb..63116886e153f 100644
--- a/clang/test/CodeGenSYCL/kernel-arg-accessor-pointer.cpp
+++ b/clang/test/CodeGenSYCL/kernel-arg-accessor-pointer.cpp
@@ -22,6 +22,8 @@ int main() {
            access::placeholder::true_t>
       acc3;
 
+  local_accessor<float, 2> acc4;
+
   // kernel_A parameters : int*, sycl::range<1>, sycl::range<1>, sycl::id<1>,
   // int*, sycl::range<1>, sycl::range<1>,sycl::id<1>.
   q.submit([&](handler &h) {
@@ -67,11 +69,19 @@ int main() {
   // Using local accessor as a kernel parameter.
   // kernel_arg_runtime_aligned is generated for pointers from local accessors.
   q.submit([&](handler &h) {
-    h.single_task<class localAccessor>([=]() {
+    h.single_task<class localAccessorDep>([=]() {
       acc3.use();
     });
   });
 
+  // Using local_accessor as a kernel parameter.
+  // kernel_arg_runtime_aligned is generated for pointers from local accessors.
+  q.submit([&](handler &h) {
+    h.single_task<class localAccessor>([=]() {
+      acc4.use();
+    });
+  });
+
   // kernel_acc_raw_ptr parameters : int*, sycl::range<1>, sycl::range<1>, sycl::id<1>, int*.
   int *rawPtr;
   q.submit([&](handler &h) {
@@ -130,7 +140,7 @@ int main() {
 // CHECK-NOT: kernel_arg_runtime_aligned
 // CHECK-NOT: kernel_arg_exclusive_ptr
 
-// CHECK: define {{.*}}spir_kernel void @{{.*}}localAccessor
+// CHECK: define {{.*}}spir_kernel void @{{.*}}localAccessorDep
 // CHECK-SAME: ptr addrspace(1) noundef align 4 [[MEM_ARG1:%[a-zA-Z0-9_]+]],
 // CHECK-SAME: ptr noundef byval{{.*}}align 4 [[ACC_RANGE1:%[a-zA-Z0-9_]+1]],
 // CHECK-SAME: ptr noundef byval{{.*}}align 4 [[MEM_RANGE1:%[a-zA-Z0-9_]+2]],
@@ -138,6 +148,14 @@ int main() {
 // CHECK-SAME: !kernel_arg_runtime_aligned ![[#ACCESSORMD2]]
 // CHECK-SAME: !kernel_arg_exclusive_ptr ![[#ACCESSORMD2]]
 
+// CHECK: define {{.*}}spir_kernel void @{{.*}}localAccessor
+// CHECK-SAME: ptr addrspace(3) noundef align 4 [[MEM_ARG1:%[a-zA-Z0-9_]+]],
+// CHECK-SAME: ptr noundef byval{{.*}}align 4 [[ACC_RANGE1:%[a-zA-Z0-9_]+1]],
+// CHECK-SAME: ptr noundef byval{{.*}}align 4 [[MEM_RANGE1:%[a-zA-Z0-9_]+2]],
+// CHECK-SAME: ptr noundef byval{{.*}}align 4 [[OFFSET1:%[a-zA-Z0-9_]+3]]
+// CHECK-SAME: !kernel_arg_runtime_aligned ![[#ACCESSORMD2]]
+// CHECK-SAME: !kernel_arg_exclusive_ptr ![[#ACCESSORMD2]]
+
 // Check kernel_acc_raw_ptr parameters
 // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_acc_raw_ptr
 // CHECK-SAME: ptr addrspace(1) noundef readonly align 4 [[MEM_ARG1:%[a-zA-Z0-9_]+]],
diff --git a/clang/test/CodeGenSYCL/no_opaque_kernel-arg-accessor-pointer.cpp b/clang/test/CodeGenSYCL/no_opaque_kernel-arg-accessor-pointer.cpp
index 434a6499ba586..7564ff285fb4d 100644
--- a/clang/test/CodeGenSYCL/no_opaque_kernel-arg-accessor-pointer.cpp
+++ b/clang/test/CodeGenSYCL/no_opaque_kernel-arg-accessor-pointer.cpp
@@ -22,6 +22,8 @@ int main() {
            access::placeholder::true_t>
       acc3;
 
+  local_accessor<float, 2> acc4;
+
   // kernel_A parameters : int*, sycl::range<1>, sycl::range<1>, sycl::id<1>,
   // int*, sycl::range<1>, sycl::range<1>,sycl::id<1>.
   q.submit([&](handler &h) {
@@ -67,11 +69,19 @@ int main() {
   // Using local accessor as a kernel parameter.
   // kernel_arg_runtime_aligned is generated for pointers from local accessors.
   q.submit([&](handler &h) {
-    h.single_task<class localAccessor>([=]() {
+    h.single_task<class localAccessorDep>([=]() {
       acc3.use();
     });
   });
 
+  // Using local accessor as a kernel parameter.
+  // kernel_arg_runtime_aligned is generated for pointers from local accessors.
+  q.submit([&](handler &h) {
+    h.single_task<class localAccessor>([=]() {
+      acc4.use();
+    });
+  });
+
   // kernel_acc_raw_ptr parameters : int*, sycl::range<1>, sycl::range<1>, sycl::id<1>, int*.
   int *rawPtr;
   q.submit([&](handler &h) {
@@ -125,13 +135,20 @@ int main() {
 // CHECK-SAME: float addrspace(1)* noundef align 4 [[MEM_ARG1:%[a-zA-Z0-9_]+]]
 // CHECK-NOT: kernel_arg_runtime_aligned
 
-// CHECK: define {{.*}}spir_kernel void @{{.*}}localAccessor
+// CHECK: define {{.*}}spir_kernel void @{{.*}}localAccessorDep
 // CHECK-SAME: float addrspace(1)* noundef align 4 [[MEM_ARG1:%[a-zA-Z0-9_]+]],
 // CHECK-SAME: %"struct.cl::sycl::range.5"* noundef byval{{.*}}align 4 [[ACC_RANGE1:%[a-zA-Z0-9_]+1]],
 // CHECK-SAME: %"struct.cl::sycl::range.5"* noundef byval{{.*}}align 4 [[MEM_RANGE1:%[a-zA-Z0-9_]+2]],
 // CHECK-SAME: %"struct.cl::sycl::id.6"* noundef byval{{.*}}align 4 [[OFFSET1:%[a-zA-Z0-9_]+3]]
 // CHECK-SAME: !kernel_arg_runtime_aligned ![[#RTALIGNED2]]
 
+// CHECK: define {{.*}}spir_kernel void @{{.*}}localAccessor
+// CHECK-SAME: float addrspace(3)* noundef align 4 [[MEM_ARG1:%[a-zA-Z0-9_]+]],
+// CHECK-SAME: %"struct.cl::sycl::range.5"* noundef byval{{.*}}align 4 [[ACC_RANGE1:%[a-zA-Z0-9_]+1]],
+// CHECK-SAME: %"struct.cl::sycl::range.5"* noundef byval{{.*}}align 4 [[MEM_RANGE1:%[a-zA-Z0-9_]+2]],
+// CHECK-SAME: %"struct.cl::sycl::id.6"* noundef byval{{.*}}align 4 [[OFFSET1:%[a-zA-Z0-9_]+3]]
+// CHECK-SAME: !kernel_arg_runtime_aligned ![[#RTALIGNED2]]
+
 // Check kernel_acc_raw_ptr parameters
 // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_acc_raw_ptr
 // CHECK-SAME: i32 addrspace(1)* noundef readonly align 4 [[MEM_ARG1:%[a-zA-Z0-9_]+]],
diff --git a/clang/test/SemaSYCL/Inputs/sycl.hpp b/clang/test/SemaSYCL/Inputs/sycl.hpp
index abc4358d739ff..0a54af0626b7f 100644
--- a/clang/test/SemaSYCL/Inputs/sycl.hpp
+++ b/clang/test/SemaSYCL/Inputs/sycl.hpp
@@ -207,6 +207,26 @@ class __attribute__((sycl_special_class)) accessor<dataT, dimensions, accessmode
 #endif
 };
 
+template <typename dataT, int dimensions>
+class __attribute__((sycl_special_class))
+local_accessor: public accessor<dataT,
+        dimensions, access::mode::read_write,
+        access::target::local> {
+public:
+  void use(void) const {}
+  template <typename... T>
+  void use(T... args) {}
+  template <typename... T>
+  void use(T... args) const {}
+  _ImplT<dimensions> impl;
+
+private:
+#ifdef __SYCL_DEVICE_ONLY__
+  void __init(__attribute__((opencl_local)) dataT *Ptr, range<dimensions> AccessRange,
+              range<dimensions> MemRange, id<dimensions> Offset) {}
+#endif
+};
+
 struct sampler_impl {
 #ifdef __SYCL_DEVICE_ONLY__
   __ocl_sampler_t m_Sampler;
diff --git a/clang/test/SemaSYCL/accessors-targets.cpp b/clang/test/SemaSYCL/accessors-targets.cpp
index 901e148b0785a..6a06fc86687cf 100644
--- a/clang/test/SemaSYCL/accessors-targets.cpp
+++ b/clang/test/SemaSYCL/accessors-targets.cpp
@@ -11,7 +11,9 @@ int main() {
   // Access work-group local memory with read and write access.
   sycl::accessor<int, 1, sycl::access::mode::read_write,
                  sycl::access::target::local>
-      local_acc;
+      local_acc_dep;
+  // Access work-group local memory with read and write access.
+  sycl::local_accessor<int, 1> local_acc;
   // Access buffer via global memory with read and write access.
   sycl::accessor<int, 1, sycl::access::mode::read_write,
                  sycl::access::target::global_buffer>
@@ -21,6 +23,13 @@ int main() {
                  sycl::access::target::constant_buffer>
       constant_acc;
 
+  q.submit([&](sycl::handler &h) {
+    h.single_task<class use_local_dep>(
+        [=] {
+          local_acc_dep.use();
+        });
+  });
+
   q.submit([&](sycl::handler &h) {
     h.single_task<class use_local>(
         [=] {
@@ -42,6 +51,7 @@ int main() {
         });
   });
 }
+// CHECK: {{.*}}use_local_dep{{.*}} 'void (__local int *, sycl::range<1>, sycl::range<1>, sycl::id<1>)'
 // CHECK: {{.*}}use_local{{.*}} 'void (__local int *, sycl::range<1>, sycl::range<1>, sycl::id<1>)'
 // CHECK: {{.*}}use_global{{.*}} 'void (__global int *, sycl::range<1>, sycl::range<1>, sycl::id<1>)'
 // CHECK: {{.*}}use_constant{{.*}} 'void (__constant int *, sycl::range<1>, sycl::range<1>, sycl::id<1>)'

From 8d0f764c4dc148dfb1a33dcd5bf1ae6ec5a577bf Mon Sep 17 00:00:00 2001
From: "aidan.belton" <aidan.belton@codeplay.com>
Date: Wed, 10 Aug 2022 14:55:59 +0100
Subject: [PATCH 05/13] Update kernel_arguments_as.cpp

---
 .../check_device_code/kernel_arguments_as.cpp | 24 ++++++++++++-------
 1 file changed, 15 insertions(+), 9 deletions(-)

diff --git a/sycl/test/check_device_code/kernel_arguments_as.cpp b/sycl/test/check_device_code/kernel_arguments_as.cpp
index fc351d59ff081..ae876af527c0d 100644
--- a/sycl/test/check_device_code/kernel_arguments_as.cpp
+++ b/sycl/test/check_device_code/kernel_arguments_as.cpp
@@ -1,14 +1,14 @@
-// RUN: %clangxx -DUSE_DEPRECATED_LOCAL_ACC -opaque-pointers -Xclang -fsycl-is-device -emit-llvm %s -S -o %t.ll -I %sycl_include -Wno-sycl-strict -Xclang -verify-ignore-unexpected=note,warning -Xclang -disable-llvm-passes
-// RUN: FileCheck %s --input-file %t.ll --check-prefixes=CHECK,CHECK-DISABLE
+// RUN: %clangxx -DUSE_DEPRECATED_LOCAL_ACC -fsycl-device-only -Xclang -opaque-pointers -emit-llvm %s -S -o %t.ll -I %sycl_include -Wno-sycl-strict -Xclang -verify-ignore-unexpected=note,warning -Xclang -disable-llvm-passes
+// RUN: FileCheck %s --input-file %t.ll --check-prefixes=CHECK,CHECK-DISABLE,CHECK-DEP
 //
 // RUN: %clangxx -fsycl-device-only -Xclang -opaque-pointers -emit-llvm %s -S -o %t.ll -I %sycl_include -Wno-sycl-strict -Xclang -verify-ignore-unexpected=note,warning -Xclang -disable-llvm-passes
-// RUN: FileCheck %s --input-file %t.ll --check-prefixes=CHECK,CHECK-DISABLE
+// RUN: FileCheck %s --input-file %t.ll --check-prefixes=CHECK,CHECK-DISABLE,CHECK-SYCL2020
 //
-// RUN: %clangxx -DUSE_DPRECATED_LOCAL_ACC -opaque-pointers -Xclang -fsycl-is-device -emit-llvm %s -S -o %t.ll -I %sycl_include -Wno-sycl-strict -Xclang -verify-ignore-unexpected=note,warning -Xclang -disable-llvm-passes -D__ENABLE_USM_ADDR_SPACE__
-// RUN: FileCheck %s --input-file %t.ll --check-prefixes=CHECK,CHECK-ENABLE
+// RUN: %clangxx -DUSE_DEPRECATED_LOCAL_ACC -fsycl-device-only -Xclang -opaque-pointers -emit-llvm %s -S -o %t.ll -I %sycl_include -Wno-sycl-strict -Xclang -verify-ignore-unexpected=note,warning -Xclang -disable-llvm-passes -D__ENABLE_USM_ADDR_SPACE__
+// RUN: FileCheck %s --input-file %t.ll --check-prefixes=CHECK,CHECK-ENABLE,CHECK-DEP
 //
 // RUN: %clangxx -fsycl-device-only -Xclang -opaque-pointers -emit-llvm %s -S -o %t.ll -I %sycl_include -Wno-sycl-strict -Xclang -verify-ignore-unexpected=note,warning -Xclang -disable-llvm-passes -D__ENABLE_USM_ADDR_SPACE__
-// RUN: FileCheck %s --input-file %t.ll --check-prefixes=CHECK,CHECK-ENABLE
+// RUN: FileCheck %s --input-file %t.ll --check-prefixes=CHECK,CHECK-ENABLE,CHECK-SYCL2020
 //
 // Check the address space of the pointer in accessor class.
 //
@@ -16,8 +16,12 @@
 // CHECK: %"class.cl::sycl::accessor[[NUMBER_SUFFIX]]" = type { %"class{{.*}}AccessorImplDevice", %[[UNION:.*]] }
 // CHECK-DISABLE: %[[UNION]] = type { ptr addrspace(1) }
 // CHECK-ENABLE: %[[UNION]] = type { ptr addrspace(5) }
-// CHECK: %struct.AccWrapper.{{[0-9]+}} = type { %"class.cl::sycl::accessor.[[NUM:[0-9]+]]" }
-// CHECK-NEXT: %"class.cl::sycl::accessor.[[NUM]]" = type { %"class{{.*}}LocalAccessorBaseDevice", ptr addrspace(3) }
+// CHECK-DEP: %struct.AccWrapper.{{[0-9]+}} = type { %"class.cl::sycl::accessor.[[NUM:[0-9]+]]" }
+// CHECK-DEP-NEXT: %"class.cl::sycl::accessor.[[NUM]]" = type { %"class{{.*}}local_accessor_base" }
+// CHECK-DEP-NEXT: %"class.cl::sycl::local_accessor_base" = type { %"class{{.*}}LocalAccessorBaseDevice", ptr addrspace(3) }
+// CHECK-SYCL2020: %struct.AccWrapper.{{[0-9]+}} = type { %"class.cl::sycl::local_accessor" }
+// CHECK-SYCL2020-NEXT: %"class.cl::sycl::local_accessor" = type { %"class{{.*}}local_accessor_base" }
+// CHECK-SYCL2020-NEXT: %"class.cl::sycl::local_accessor_base" = type { %"class{{.*}}LocalAccessorBaseDevice", ptr addrspace(3) }
 //
 // Check that kernel arguments doesn't have generic address space.
 //
@@ -27,7 +31,9 @@
 
 using namespace sycl;
 
-template <typename Acc> struct AccWrapper { Acc accessor; };
+template <typename Acc> struct AccWrapper {
+  Acc accessor;
+};
 
 int main() {
 

From dec6b1fa4fa85cd765f43fd7339e30a98f9c9188 Mon Sep 17 00:00:00 2001
From: "aidan.belton" <aidan.belton@codeplay.com>
Date: Thu, 11 Aug 2022 09:55:29 +0100
Subject: [PATCH 06/13] Remove isReadOnly

---
 clang/lib/Sema/SemaSYCL.cpp | 7 ++-----
 1 file changed, 2 insertions(+), 5 deletions(-)

diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index c2870844398f5..2e501261a08df 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -1939,7 +1939,6 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler {
   void handleAccessorType(QualType FieldTy, const CXXRecordDecl *RecordDecl,
                           SourceLocation Loc) {
     handleAccessorPropertyList(Params.back(), RecordDecl, Loc);
-    bool isReadOnly = false;
 
     // If "accessor" type check if read only
     if (Util::isSyclType(FieldTy, "accessor", true /*Tmpl*/)) {
@@ -1949,14 +1948,12 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler {
       const TemplateArgument &AccessModeArg =
           AccessorSpecializationDecl->getTemplateArgs().get(2);
       if (isReadOnlyAccessor(AccessModeArg))
-        isReadOnly = true;
+        Params.back()->addAttr(
+            SYCLAccessorReadonlyAttr::CreateImplicit(SemaRef.getASTContext()));
     }
 
     // Add implicit attribute to parameter decl when it is a read only
     // SYCL accessor.
-    if (isReadOnly)
-      Params.back()->addAttr(
-          SYCLAccessorReadonlyAttr::CreateImplicit(SemaRef.getASTContext()));
     Params.back()->addAttr(
         SYCLAccessorPtrAttr::CreateImplicit(SemaRef.getASTContext()));
   }

From 70090768e69a8edcaf94ecd0289fb2e39f49d037 Mon Sep 17 00:00:00 2001
From: "aidan.belton" <aidan.belton@codeplay.com>
Date: Thu, 11 Aug 2022 09:55:40 +0100
Subject: [PATCH 07/13] Clang format

---
 sycl/include/sycl/accessor.hpp                | 24 +++++++++----------
 .../check_device_code/kernel_arguments_as.cpp |  4 +---
 2 files changed, 12 insertions(+), 16 deletions(-)

diff --git a/sycl/include/sycl/accessor.hpp b/sycl/include/sycl/accessor.hpp
index 110e6a4eaa3a4..82fddb7798a7b 100644
--- a/sycl/include/sycl/accessor.hpp
+++ b/sycl/include/sycl/accessor.hpp
@@ -2096,7 +2096,7 @@ 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(),
@@ -2104,11 +2104,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;
@@ -2126,7 +2125,7 @@ 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)) {
@@ -2135,12 +2134,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;
diff --git a/sycl/test/check_device_code/kernel_arguments_as.cpp b/sycl/test/check_device_code/kernel_arguments_as.cpp
index ae876af527c0d..fe9671917d286 100644
--- a/sycl/test/check_device_code/kernel_arguments_as.cpp
+++ b/sycl/test/check_device_code/kernel_arguments_as.cpp
@@ -31,9 +31,7 @@
 
 using namespace sycl;
 
-template <typename Acc> struct AccWrapper {
-  Acc accessor;
-};
+template <typename Acc> struct AccWrapper { Acc accessor; };
 
 int main() {
 

From 3648d5caa02bc13c01a5800331524b39ded0e197 Mon Sep 17 00:00:00 2001
From: "aidan.belton" <aidan.belton@codeplay.com>
Date: Thu, 11 Aug 2022 10:01:31 +0100
Subject: [PATCH 08/13] Clang-format accessor

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

diff --git a/sycl/include/sycl/accessor.hpp b/sycl/include/sycl/accessor.hpp
index 82fddb7798a7b..b03fccd7fd35d 100644
--- a/sycl/include/sycl/accessor.hpp
+++ b/sycl/include/sycl/accessor.hpp
@@ -2096,7 +2096,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(),
@@ -2125,7 +2126,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)) {

From 626a101036e7d70239e4b45c39455a2a4029909e Mon Sep 17 00:00:00 2001
From: "aidan.belton" <aidan.belton@codeplay.com>
Date: Thu, 11 Aug 2022 10:14:54 +0100
Subject: [PATCH 09/13] Format with more upto date version

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

diff --git a/sycl/include/sycl/accessor.hpp b/sycl/include/sycl/accessor.hpp
index b03fccd7fd35d..110e6a4eaa3a4 100644
--- a/sycl/include/sycl/accessor.hpp
+++ b/sycl/include/sycl/accessor.hpp
@@ -2096,8 +2096,7 @@ 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(),
@@ -2105,10 +2104,11 @@ 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;
@@ -2126,8 +2126,7 @@ 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)) {
@@ -2136,11 +2135,12 @@ 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;

From e6622623698a81e588d806528a3449e054104966 Mon Sep 17 00:00:00 2001
From: "aidan.belton" <aidan.belton@codeplay.com>
Date: Fri, 12 Aug 2022 15:16:58 +0100
Subject: [PATCH 10/13] Remove unneccessary parameter in isSyclAccessorType

---
 clang/lib/Sema/SemaSYCL.cpp | 31 ++++++++++++++-----------------
 1 file changed, 14 insertions(+), 17 deletions(-)

diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 2e501261a08df..5f9b7872dd456 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -111,10 +111,10 @@ class Util {
   /// \param Tmpl  whether the class is template instantiation or simple record
   static bool isSyclType(QualType Ty, StringRef Name, bool Tmpl = false);
 
-  /// Checks whether given clang type is a standard SYCL API accessor class.
+  /// Checks whether given clang type is a standard SYCL API accessor class,
+  /// the check assumes the type is templated.
   /// \param Ty    the clang type being checked
-  /// \param Tmpl  whether the class is template instantiation or simple record
-  static bool isSyclAccessorType(QualType Ty, bool Tmpl = false);
+  static bool isSyclAccessorType(QualType Ty);
 
   /// Checks whether given clang type is a full specialization of the SYCL
   /// specialization constant class.
@@ -1624,7 +1624,7 @@ class SyclKernelFieldChecker : public SyclKernelFieldHandler {
     assert(Util::isSyclSpecialType(Ty) &&
            "Should only be called on sycl special class types.");
     const RecordDecl *RecD = Ty->getAsRecordDecl();
-    if (IsSIMD && !Util::isSyclAccessorType(Ty, true /*Tmp*/))
+    if (IsSIMD && !Util::isSyclAccessorType(Ty))
       return SemaRef.Diag(Loc.getBegin(),
                           diag::err_sycl_esimd_not_supported_for_type)
              << RecD;
@@ -1966,8 +1966,7 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler {
     const auto *RecordDecl = FieldTy->getAsCXXRecordDecl();
     assert(RecordDecl && "The type must be a RecordDecl");
     llvm::StringLiteral MethodName =
-        KernelDecl->hasAttr<SYCLSimdAttr>() &&
-                Util::isSyclAccessorType(FieldTy, true /*Tmp*/)
+        KernelDecl->hasAttr<SYCLSimdAttr>() && Util::isSyclAccessorType(FieldTy)
             ? InitESIMDMethodName
             : InitMethodName;
     CXXMethodDecl *InitMethod = getMethodByName(RecordDecl, MethodName);
@@ -1992,7 +1991,7 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler {
       // added, this code needs to be refactored to call
       // handleAccessorPropertyList for each class which requires it.
       if (ParamTy.getTypePtr()->isPointerType() &&
-          Util::isSyclAccessorType(FieldTy, true /*Tmp*/))
+          Util::isSyclAccessorType(FieldTy))
         handleAccessorType(FieldTy, RecordDecl, FD->getBeginLoc());
     }
     LastParamIndex = ParamIndex;
@@ -2087,8 +2086,7 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler {
     const auto *RecordDecl = FieldTy->getAsCXXRecordDecl();
     assert(RecordDecl && "The type must be a RecordDecl");
     llvm::StringLiteral MethodName =
-        KernelDecl->hasAttr<SYCLSimdAttr>() &&
-                Util::isSyclAccessorType(FieldTy, true /*Tmp*/)
+        KernelDecl->hasAttr<SYCLSimdAttr>() && Util::isSyclAccessorType(FieldTy)
             ? InitESIMDMethodName
             : InitMethodName;
     CXXMethodDecl *InitMethod = getMethodByName(RecordDecl, MethodName);
@@ -2107,7 +2105,7 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler {
       // added, this code needs to be refactored to call
       // handleAccessorPropertyList for each class which requires it.
       if (ParamTy.getTypePtr()->isPointerType() &&
-          Util::isSyclAccessorType(FieldTy, true /*Tmp*/))
+          Util::isSyclAccessorType(FieldTy))
         handleAccessorType(FieldTy, RecordDecl, BS.getBeginLoc());
     }
     LastParamIndex = ParamIndex;
@@ -2229,9 +2227,8 @@ class SyclKernelArgsSizeChecker : public SyclKernelFieldHandler {
     const CXXRecordDecl *RecordDecl = FieldTy->getAsCXXRecordDecl();
     assert(RecordDecl && "The type must be a RecordDecl");
     llvm::StringLiteral MethodName =
-        (IsSIMD && Util::isSyclAccessorType(FieldTy, true /*Tmp*/))
-            ? InitESIMDMethodName
-            : InitMethodName;
+        (IsSIMD && Util::isSyclAccessorType(FieldTy)) ? InitESIMDMethodName
+                                                      : InitMethodName;
     CXXMethodDecl *InitMethod = getMethodByName(RecordDecl, MethodName);
     assert(InitMethod && "The type must have the __init method");
     for (const ParmVarDecl *Param : InitMethod->parameters())
@@ -3148,7 +3145,7 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler {
   bool handleSyclSpecialType(FieldDecl *FD, QualType FieldTy) final {
     const auto *ClassTy = FieldTy->getAsCXXRecordDecl();
     assert(ClassTy && "Type must be a C++ record type");
-    if (Util::isSyclAccessorType(FieldTy, true /*Tmp*/)) {
+    if (Util::isSyclAccessorType(FieldTy)) {
       const auto *AccTy =
           cast<ClassTemplateSpecializationDecl>(FieldTy->getAsRecordDecl());
       assert(AccTy->getTemplateArgs().size() >= 2 &&
@@ -5209,9 +5206,9 @@ bool Util::isSyclType(QualType Ty, StringRef Name, bool Tmpl) {
   return matchQualifiedTypeName(Ty, Scopes);
 }
 
-bool Util::isSyclAccessorType(QualType Ty, bool Tmpl) {
-  return isSyclType(Ty, "accessor", Tmpl) ||
-         isSyclType(Ty, "local_accessor", Tmpl);
+bool Util::isSyclAccessorType(QualType Ty) {
+  return isSyclType(Ty, "accessor", true /* Tmpl */) ||
+         isSyclType(Ty, "local_accessor", true /* Tmpl */);
 }
 
 bool Util::isAccessorPropertyListType(QualType Ty) {

From 69820f4be5e4a9af96c732a2f15491fde02cfde4 Mon Sep 17 00:00:00 2001
From: AidanBeltonS <87009434+AidanBeltonS@users.noreply.github.com>
Date: Fri, 12 Aug 2022 15:25:00 +0100
Subject: [PATCH 11/13] Apply suggestions from code review

remove unnecessary cl namespace

Co-authored-by: Steffen Larsen <steffen.larsen@intel.com>
---
 sycl/test/basic_tests/set_arg_error.cpp             | 2 +-
 sycl/test/check_device_code/kernel_arguments_as.cpp | 2 +-
 2 files changed, 2 insertions(+), 2 deletions(-)

diff --git a/sycl/test/basic_tests/set_arg_error.cpp b/sycl/test/basic_tests/set_arg_error.cpp
index c88feaa4b942e..5febdede1533b 100644
--- a/sycl/test/basic_tests/set_arg_error.cpp
+++ b/sycl/test/basic_tests/set_arg_error.cpp
@@ -36,7 +36,7 @@ int main() {
                    sycl::access::target::local>
         local_acc({size}, h);
 #else
-    cl::sycl::local_accessor<int, 1> local_acc({size}, h);
+    sycl::local_accessor<int, 1> local_acc({size}, h);
 #endif
 
     TriviallyCopyable tc{1, 2};
diff --git a/sycl/test/check_device_code/kernel_arguments_as.cpp b/sycl/test/check_device_code/kernel_arguments_as.cpp
index fe9671917d286..0a5f1eafbc31a 100644
--- a/sycl/test/check_device_code/kernel_arguments_as.cpp
+++ b/sycl/test/check_device_code/kernel_arguments_as.cpp
@@ -47,7 +47,7 @@ int main() {
                      sycl::access::target::local>
           local_acc(sycl::range<1>(10), cgh);
 #else
-      cl::sycl::local_accessor<int, 1> local_acc(cl::sycl::range<1>(10), cgh);
+      sycl::local_accessor<int, 1> local_acc(sycl::range<1>(10), cgh);
 #endif // USE_DEPRECATED_LOCAL_ACC
       auto acc_wrapped = AccWrapper<decltype(acc)>{acc};
       auto local_acc_wrapped = AccWrapper<decltype(local_acc)>{local_acc};

From 627e2f781fb12bbd9ee0811b62fbf99167c9f3dd Mon Sep 17 00:00:00 2001
From: "aidan.belton" <aidan.belton@codeplay.com>
Date: Mon, 15 Aug 2022 09:50:09 +0100
Subject: [PATCH 12/13] Add TODO

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

diff --git a/sycl/include/sycl/accessor.hpp b/sycl/include/sycl/accessor.hpp
index b14b4045028cd..3031cf6a395ac 100644
--- a/sycl/include/sycl/accessor.hpp
+++ b/sycl/include/sycl/accessor.hpp
@@ -2227,6 +2227,7 @@ class __SYCL_SPECIAL_CLASS local_accessor_base :
   }
 };
 
+// TODO: Remove deprecated specialization once no longer needed
 template <typename DataT, int Dimensions, access::mode AccessMode,
           access::placeholder IsPlaceholder>
 class __SYCL_SPECIAL_CLASS accessor<DataT, Dimensions, AccessMode,

From f1f82cde30306378789a97c43d9cdc6edf547458 Mon Sep 17 00:00:00 2001
From: "aidan.belton" <aidan.belton@codeplay.com>
Date: Mon, 15 Aug 2022 12:07:27 +0100
Subject: [PATCH 13/13] Update test namespace

---
 .../CodeGenSYCL/no_opaque_kernel-arg-accessor-pointer.cpp   | 6 +++---
 1 file changed, 3 insertions(+), 3 deletions(-)

diff --git a/clang/test/CodeGenSYCL/no_opaque_kernel-arg-accessor-pointer.cpp b/clang/test/CodeGenSYCL/no_opaque_kernel-arg-accessor-pointer.cpp
index 30bd93858f28e..95d6b687a0d58 100644
--- a/clang/test/CodeGenSYCL/no_opaque_kernel-arg-accessor-pointer.cpp
+++ b/clang/test/CodeGenSYCL/no_opaque_kernel-arg-accessor-pointer.cpp
@@ -144,9 +144,9 @@ int main() {
 
 // CHECK: define {{.*}}spir_kernel void @{{.*}}localAccessor
 // CHECK-SAME: float addrspace(3)* noundef align 4 [[MEM_ARG1:%[a-zA-Z0-9_]+]],
-// CHECK-SAME: %"struct.cl::sycl::range.5"* noundef byval{{.*}}align 4 [[ACC_RANGE1:%[a-zA-Z0-9_]+1]],
-// CHECK-SAME: %"struct.cl::sycl::range.5"* noundef byval{{.*}}align 4 [[MEM_RANGE1:%[a-zA-Z0-9_]+2]],
-// CHECK-SAME: %"struct.cl::sycl::id.6"* noundef byval{{.*}}align 4 [[OFFSET1:%[a-zA-Z0-9_]+3]]
+// CHECK-SAME: %"struct.sycl::_V1::range.5"* noundef byval{{.*}}align 4 [[ACC_RANGE1:%[a-zA-Z0-9_]+1]],
+// CHECK-SAME: %"struct.sycl::_V1::range.5"* noundef byval{{.*}}align 4 [[MEM_RANGE1:%[a-zA-Z0-9_]+2]],
+// CHECK-SAME: %"struct.sycl::_V1::id.6"* noundef byval{{.*}}align 4 [[OFFSET1:%[a-zA-Z0-9_]+3]]
 // CHECK-SAME: !kernel_arg_runtime_aligned ![[#RTALIGNED2]]
 
 // Check kernel_acc_raw_ptr parameters