diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index c7301d5970b6e..14907bb03e325 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -636,15 +636,15 @@ static void buildArgTys(ASTContext &Context, CXXRecordDecl *KernelObj,
 
       CreateAndAddPrmDsc(Fld, PointerType);
 
-      FieldDecl *RangeFld = getFieldDeclByName(RecordDecl, {"__impl", "Range"});
+      FieldDecl *RangeFld = getFieldDeclByName(RecordDecl, {"__implx", "Range"});
       assert(RangeFld &&
              "The accessor must contain the Range from the __impl field");
       CreateAndAddPrmDsc(RangeFld, RangeFld->getType());
 
       FieldDecl *OffsetFld =
-          getFieldDeclByName(RecordDecl, {"__impl", "Offset"});
+          getFieldDeclByName(RecordDecl, {"__implx", "Offset"});
       assert(OffsetFld &&
-             "The accessor must contain the Offset from the __impl field");
+             "The accessor must contain the Offset from the __implx field");
       CreateAndAddPrmDsc(OffsetFld, OffsetFld->getType());
     } else if (Util::isSyclStreamType(ArgTy)) {
       // the parameter is a SYCL stream object
@@ -700,7 +700,7 @@ static void populateIntHeader(SYCLIntegrationHeader &H, const StringRef Name,
                      getAccessTarget(AccTmplTy), Offset);
       // ... second descriptor (translated to range kernel parameter):
       FieldDecl *RngFld =
-          getFieldDeclByName(AccTy, {"__impl", "Range"}, &Offset);
+          getFieldDeclByName(AccTy, {"__implx", "Range"}, &Offset);
       uint64_t Sz = Ctx.getTypeSizeInChars(RngFld->getType()).getQuantity();
       H.addParamDesc(SYCLIntegrationHeader::kind_std_layout,
                      static_cast<unsigned>(Sz), static_cast<unsigned>(Offset));
@@ -708,7 +708,7 @@ static void populateIntHeader(SYCLIntegrationHeader &H, const StringRef Name,
       // Get offset in bytes
       Offset = Layout.getFieldOffset(Fld->getFieldIndex()) / 8;
       FieldDecl *OffstFld =
-          getFieldDeclByName(AccTy, {"__impl", "Offset"}, &Offset);
+          getFieldDeclByName(AccTy, {"__implx", "Offset"}, &Offset);
       Sz = Ctx.getTypeSizeInChars(OffstFld->getType()).getQuantity();
       H.addParamDesc(SYCLIntegrationHeader::kind_std_layout,
                      static_cast<unsigned>(Sz), static_cast<unsigned>(Offset));
diff --git a/sycl/include/CL/sycl/accessor.hpp b/sycl/include/CL/sycl/accessor.hpp
index 413196a9cf076..e27d91cd58cd2 100644
--- a/sycl/include/CL/sycl/accessor.hpp
+++ b/sycl/include/CL/sycl/accessor.hpp
@@ -355,6 +355,9 @@ SYCL_ACCESSOR_SUBCLASS(accessor_common, accessor_base, true /* always */) {
   template <int Dimensions = dimensions>
   typename std::enable_if<(Dimensions > 0), id<Dimensions>>::type
   get_offset() const { return this->__impl()->Offset; }
+
+  friend class ::cl::sycl::simple_scheduler::Node;
+  friend class ::cl::sycl::simple_scheduler::Scheduler;
 };
 
 SYCL_ACCESSOR_SUBCLASS(accessor_opdata_w, accessor_common,
@@ -366,6 +369,8 @@ SYCL_ACCESSOR_SUBCLASS(accessor_opdata_w, accessor_common,
   operator dataT &() const {
     return this->__impl()->Data[0];
   }
+  friend class ::cl::sycl::simple_scheduler::Node;
+  friend class ::cl::sycl::simple_scheduler::Scheduler;
 };
 
 SYCL_ACCESSOR_SUBCLASS(accessor_subscript_wn, accessor_opdata_w,
@@ -387,6 +392,8 @@ SYCL_ACCESSOR_SUBCLASS(accessor_subscript_wn, accessor_opdata_w,
     return subscript_obj<dimensions, dataT, dimensions - 1, accessMode,
                          accessTarget, isPlaceholder>(*this, ids);
   }
+  friend class ::cl::sycl::simple_scheduler::Node;
+  friend class ::cl::sycl::simple_scheduler::Scheduler;
 };
 
 SYCL_ACCESSOR_SUBCLASS(accessor_subscript_w, accessor_subscript_wn,
@@ -408,6 +415,8 @@ SYCL_ACCESSOR_SUBCLASS(accessor_subscript_w, accessor_subscript_wn,
   dataT &operator[](size_t index) const {
     return this->__impl()->Data[index];
   }
+  friend class ::cl::sycl::simple_scheduler::Node;
+  friend class ::cl::sycl::simple_scheduler::Scheduler;
 };
 
 SYCL_ACCESSOR_SUBCLASS(accessor_opdata_r, accessor_subscript_w,
@@ -416,6 +425,8 @@ SYCL_ACCESSOR_SUBCLASS(accessor_opdata_r, accessor_subscript_w,
   operator PureType() const {
     return this->__impl()->Data[0];
   }
+  friend class ::cl::sycl::simple_scheduler::Node;
+  friend class ::cl::sycl::simple_scheduler::Scheduler;
 };
 
 SYCL_ACCESSOR_SUBCLASS(accessor_subscript_rn, accessor_opdata_r,
@@ -434,6 +445,8 @@ SYCL_ACCESSOR_SUBCLASS(accessor_subscript_rn, accessor_opdata_r,
     return subscript_obj<dimensions, dataT, dimensions - 1, accessMode,
                          accessTarget, isPlaceholder>(*this, ids);
   }
+  friend class ::cl::sycl::simple_scheduler::Node;
+  friend class ::cl::sycl::simple_scheduler::Scheduler;
 };
 
 SYCL_ACCESSOR_SUBCLASS(accessor_subscript_r, accessor_subscript_rn,
@@ -447,6 +460,8 @@ SYCL_ACCESSOR_SUBCLASS(accessor_subscript_r, accessor_subscript_rn,
   operator[](size_t index) const {
     return this->__impl()->Data[index];
   }
+  friend class ::cl::sycl::simple_scheduler::Node;
+  friend class ::cl::sycl::simple_scheduler::Scheduler;
 };
 
 template <access::target accessTarget> struct getAddressSpace {
@@ -469,6 +484,8 @@ SYCL_ACCESSOR_SUBCLASS(accessor_subscript_atomic_eq0, accessor_subscript_r,
     return atomic<PureType, addressSpace>(
         multi_ptr<PureType, addressSpace>(&(this->__impl()->Data[0])));
   }
+  friend class ::cl::sycl::simple_scheduler::Node;
+  friend class ::cl::sycl::simple_scheduler::Scheduler;
 };
 
 // Available when: accessMode == access::mode::atomic && dimensions > 0
@@ -483,6 +500,8 @@ SYCL_ACCESSOR_SUBCLASS(accessor_subscript_atomic_gt0,
         multi_ptr<PureType, addressSpace>(&(this->__impl()->Data[getOffsetForId(
             this->__impl()->Range, index, this->__impl()->Offset)])));
   }
+  friend class ::cl::sycl::simple_scheduler::Node;
+  friend class ::cl::sycl::simple_scheduler::Scheduler;
 };
 
 // Available only when: accessMode == access::mode::atomic && dimensions == 1
@@ -496,6 +515,8 @@ SYCL_ACCESSOR_SUBCLASS(accessor_subscript_atomic_eq1,
     return atomic<PureType, addressSpace>(
         multi_ptr<PureType, addressSpace>(&(this->__impl()->Data[index])));
   }
+  friend class ::cl::sycl::simple_scheduler::Node;
+  friend class ::cl::sycl::simple_scheduler::Scheduler;
 };
 
 // TODO:
@@ -535,6 +556,8 @@ SYCL_ACCESSOR_SUBCLASS(accessor_pointer, accessor_subscript_atomic_eq1, true) {
   get_pointer() const {
     return local_ptr<DataT>(this->__impl()->Data);
   }
+  friend class ::cl::sycl::simple_scheduler::Node;
+  friend class ::cl::sycl::simple_scheduler::Scheduler;
 };
 
 } // namespace detail
@@ -557,13 +580,13 @@ class accessor
   // Make sure Impl field is the first in the class, so that it is
   // safe to reinterpret a pointer to accessor as a pointer to the
   // implementation.
-  _ImplT __impl;
+  _ImplT __implx;
 
   void __init(_ValueType *Ptr, range<dimensions> Range,
       id<dimensions> Offset) {
-    __impl.Data = Ptr;
-    __impl.Range = Range;
-    __impl.Offset = Offset;
+    __implx.Data = Ptr;
+    __implx.Range = Range;
+    __implx.Offset = Offset;
   }
 
 public:
@@ -593,7 +616,7 @@ class accessor
                 AccessTarget == access::target::constant_buffer))) &&
             Dimensions == 0),
            buffer<DataT, 1>>::type &bufferRef)
-      : __impl(detail::getSyclObjImpl(bufferRef)->BufPtr) {
+      : __implx(detail::getSyclObjImpl(bufferRef)->BufPtr) {
     auto BufImpl = detail::getSyclObjImpl(bufferRef);
     if (AccessTarget == access::target::host_buffer) {
       if (BufImpl->OpenCLInterop) {
@@ -633,7 +656,7 @@ class accessor
 #ifdef __SYCL_DEVICE_ONLY__
       ; // This ctor can't be used in device code, so no need to define it.
 #else // !__SYCL_DEVICE_ONLY__
-      : __impl(detail::getSyclObjImpl(bufferRef)->BufPtr,
+      : __implx(detail::getSyclObjImpl(bufferRef)->BufPtr,
                detail::getSyclObjImpl(bufferRef)->Range,
                &commandGroupHandlerRef) {
     auto BufImpl = detail::getSyclObjImpl(bufferRef);
@@ -643,7 +666,7 @@ class accessor
           "interoperability buffer");
     }
     commandGroupHandlerRef.AddBufDep<AccessMode, AccessTarget>(*BufImpl);
-    __impl.m_Buf = BufImpl.get();
+    __implx.m_Buf = BufImpl.get();
   }
 #endif // !__SYCL_DEVICE_ONLY__
 
@@ -669,7 +692,7 @@ class accessor
                AccessTarget == access::target::constant_buffer))) &&
             Dimensions > 0),
            buffer<DataT, Dimensions>>::type &bufferRef)
-      : __impl(detail::getSyclObjImpl(bufferRef)->BufPtr,
+      : __implx(detail::getSyclObjImpl(bufferRef)->BufPtr,
                detail::getSyclObjImpl(bufferRef)->Range) {
     auto BufImpl = detail::getSyclObjImpl(bufferRef);
     if (AccessTarget == access::target::host_buffer) {
@@ -710,7 +733,7 @@ class accessor
 #ifdef __SYCL_DEVICE_ONLY__
       ; // This ctor can't be used in device code, so no need to define it.
 #else
-      : __impl(detail::getSyclObjImpl(bufferRef)->BufPtr,
+      : __implx(detail::getSyclObjImpl(bufferRef)->BufPtr,
                detail::getSyclObjImpl(bufferRef)->Range,
                &commandGroupHandlerRef) {
     auto BufImpl = detail::getSyclObjImpl(bufferRef);
@@ -720,7 +743,7 @@ class accessor
           "interoperability buffer");
     }
     commandGroupHandlerRef.AddBufDep<AccessMode, AccessTarget>(*BufImpl);
-    __impl.m_Buf = BufImpl.get();
+    __implx.m_Buf = BufImpl.get();
   }
 #endif
 
@@ -752,7 +775,7 @@ class accessor
 #ifdef __SYCL_DEVICE_ONLY__
       ; // This ctor can't be used in device code, so no need to define it.
 #else // !__SYCL_DEVICE_ONLY__
-      : __impl(detail::getSyclObjImpl(bufferRef)->BufPtr, Range, Offset) {
+      : __implx(detail::getSyclObjImpl(bufferRef)->BufPtr, Range, Offset) {
     auto BufImpl = detail::getSyclObjImpl(bufferRef);
     if (AccessTarget == access::target::host_buffer) {
       if (BufImpl->OpenCLInterop) {
@@ -796,7 +819,7 @@ class accessor
 #ifdef __SYCL_DEVICE_ONLY__
       ; // This ctor can't be used in device code, so no need to define it.
 #else // !__SYCL_DEVICE_ONLY__
-      : __impl(detail::getSyclObjImpl(bufferRef)->BufPtr, Range,
+      : __implx(detail::getSyclObjImpl(bufferRef)->BufPtr, Range,
                &commandGroupHandlerRef, Offset) {
     auto BufImpl = detail::getSyclObjImpl(bufferRef);
     if (BufImpl->OpenCLInterop && !BufImpl->isValidAccessToMem(accessMode)) {
@@ -805,7 +828,7 @@ class accessor
           "interoperability buffer");
     }
     commandGroupHandlerRef.AddBufDep<AccessMode, AccessTarget>(*BufImpl);
-    __impl.m_Buf = BufImpl.get();
+    __implx.m_Buf = BufImpl.get();
   }
 #endif // !__SYCL_DEVICE_ONLY__
 
@@ -835,7 +858,7 @@ class accessor
                                     Dimensions > 0),
                                    range<Dimensions>>::type allocationSize,
            handler &commandGroupHandlerRef)
-      : __impl(allocationSize, &commandGroupHandlerRef) {}
+      : __implx(allocationSize, &commandGroupHandlerRef) {}
 };
 
 } // namespace sycl
diff --git a/sycl/include/CL/sycl/detail/scheduler/scheduler.cpp b/sycl/include/CL/sycl/detail/scheduler/scheduler.cpp
index 0509196da4bd0..dc8161ebe35e2 100644
--- a/sycl/include/CL/sycl/detail/scheduler/scheduler.cpp
+++ b/sycl/include/CL/sycl/detail/scheduler/scheduler.cpp
@@ -61,9 +61,7 @@ void Node::addAccRequirement(
     accessor<dataT, dimensions, accessMode, accessTarget, isPlaceholder> &&Acc,
     int argIndex) {
   detail::buffer_impl<dataT, dimensions> *buf =
-      Acc.template accessor_base<dataT, dimensions, accessMode, accessTarget,
-                                 isPlaceholder>::__impl()
-          ->m_Buf;
+      Acc.__impl()->m_Buf;
   addBufRequirement<accessMode, accessTarget, dataT, dimensions>(*buf);
   addInteropArg(nullptr, buf->get_size(), argIndex,
                 getReqForBuffer(m_Bufs, *buf));
@@ -128,8 +126,7 @@ template <typename T, int Dimensions, access::mode mode, access::target tgt,
           access::placeholder isPlaceholder>
 void Node::addExplicitMemOp(
     accessor<T, Dimensions, mode, tgt, isPlaceholder> &Dest, T Src) {
-  auto *DestBase = Dest.template accessor_base<T, Dimensions, mode, tgt,
-                                               isPlaceholder>::__impl();
+  auto *DestBase = Dest.__impl();
   assert(DestBase != nullptr &&
          "Accessor should have an initialized accessor_base");
   detail::buffer_impl<T, Dimensions> *Buf = DestBase->m_Buf;
@@ -153,13 +150,10 @@ template <typename T_src, int dim_src, access::mode mode_src,
 void Node::addExplicitMemOp(
     accessor<T_src, dim_src, mode_src, tgt_src, isPlaceholder_src> Src,
     accessor<T_dest, dim_dest, mode_dest, tgt_dest, isPlaceholder_dest> Dest) {
-  auto *SrcBase = Src.template accessor_base<T_src, dim_src, mode_src, tgt_src,
-                                             isPlaceholder_src>::__impl();
+  auto *SrcBase = Src.__impl();
   assert(SrcBase != nullptr &&
          "Accessor should have an initialized accessor_base");
-  auto *DestBase =
-      Dest.template accessor_base<T_dest, dim_dest, mode_dest, tgt_dest,
-                                  isPlaceholder_dest>::__impl();
+  auto *DestBase = Dest.__impl();
   assert(DestBase != nullptr &&
          "Accessor should have an initialized accessor_base");
 
@@ -191,8 +185,8 @@ template <typename T, int Dimensions, access::mode mode, access::target tgt,
 void Scheduler::updateHost(
     accessor<T, Dimensions, mode, tgt, isPlaceholder> &Acc,
     cl::sycl::event &Event) {
-  auto *AccBase = Acc.template accessor_base<T, Dimensions, mode, tgt,
-                                             isPlaceholder>::__impl();
+  auto *AccBase = Acc.impl();
+
   assert(AccBase != nullptr &&
          "Accessor should have an initialized accessor_base");
   detail::buffer_impl<T, Dimensions> *Buf = AccBase->m_Buf;