diff --git a/sycl/include/sycl/accessor.hpp b/sycl/include/sycl/accessor.hpp
index c060625d3329a..a63b6da725646 100644
--- a/sycl/include/sycl/accessor.hpp
+++ b/sycl/include/sycl/accessor.hpp
@@ -11,6 +11,7 @@
 #include <CL/__spirv/spirv_types.hpp>
 #include <sycl/atomic.hpp>
 #include <sycl/buffer.hpp>
+#include <sycl/detail/accessor_iterator.hpp>
 #include <sycl/detail/cl.h>
 #include <sycl/detail/common.hpp>
 #include <sycl/detail/export.hpp>
@@ -30,6 +31,7 @@
 #include <sycl/property_list_conversion.hpp>
 #include <sycl/sampler.hpp>
 
+#include <iterator>
 #include <type_traits>
 
 #include <utility>
@@ -334,7 +336,7 @@ class accessor_common {
 
   public:
     AccessorSubscript(AccType Accessor, id<Dims> IDs)
-        : MAccessor(Accessor), MIDs(IDs) {}
+        : MIDs(IDs), MAccessor(Accessor) {}
 
     // Only accessor class is supposed to use this c'tor for the first
     // operator[].
@@ -1192,7 +1194,12 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor :
   using value_type = DataT;
   using reference = DataT &;
   using const_reference = const DataT &;
-  using difference_type = size_t;
+
+  using iterator = typename detail::accessor_iterator<DataT, Dimensions>;
+  using const_iterator =
+      typename detail::accessor_iterator<const DataT, Dimensions>;
+  using difference_type =
+      typename std::iterator_traits<iterator>::difference_type;
 
   // The list of accessor constructors with their arguments
   // -------+---------+-------+----+-----+--------------
@@ -2091,6 +2098,34 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor :
   bool operator==(const accessor &Rhs) const { return impl == Rhs.impl; }
   bool operator!=(const accessor &Rhs) const { return !(*this == Rhs); }
 
+  iterator begin() const noexcept {
+    return iterator::getBegin(
+        get_pointer(),
+        detail::convertToArrayOfN<Dimensions, 1>(getMemoryRange()), get_range(),
+        get_offset());
+  }
+
+  iterator end() const noexcept {
+    return iterator::getEnd(
+        get_pointer(),
+        detail::convertToArrayOfN<Dimensions, 1>(getMemoryRange()), get_range(),
+        get_offset());
+  }
+
+  const_iterator cbegin() const noexcept {
+    return const_iterator::getBegin(
+        get_pointer(),
+        detail::convertToArrayOfN<Dimensions, 1>(getMemoryRange()), get_range(),
+        get_offset());
+  }
+
+  const_iterator cend() const noexcept {
+    return const_iterator::getEnd(
+        get_pointer(),
+        detail::convertToArrayOfN<Dimensions, 1>(getMemoryRange()), get_range(),
+        get_offset());
+  }
+
 private:
 #ifdef __SYCL_DEVICE_ONLY__
   size_t getTotalOffset() const {
diff --git a/sycl/include/sycl/detail/accessor_iterator.hpp b/sycl/include/sycl/detail/accessor_iterator.hpp
new file mode 100644
index 0000000000000..244350b5cb7fb
--- /dev/null
+++ b/sycl/include/sycl/detail/accessor_iterator.hpp
@@ -0,0 +1,351 @@
+//==------------ accessor_iterator.hpp - SYCL standard header file ---------==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+#include <sycl/id.hpp>
+
+#include <cstddef>
+#include <iterator>
+#include <ostream>
+#include <type_traits>
+
+/// \file accessor_iterator.hpp
+/// The file contains implementation of accessor iterator class.
+///
+/// The reason why we can't use a plain pointer as an interator and have to
+/// implement a custom class here is explained in section 4.7.6.8. Ranged
+/// accessors of SYCL 2020 specification. A couple of quotes from there:
+///
+/// > Accessors of type accessor and host_accessor can be constructed from a
+/// > sub-range of a buffer by providing a range and offset to the constructor.
+/// >
+/// > If the ranged accessor is multi-dimensional, the sub-range is allowed to
+/// > describe a region of memory in the underlying buffer that is not
+/// > contiguous in the linear address space.
+/// >
+/// > Most of the accessor member functions which provide a reference to the
+/// > underlying buffer elements are affected by a ranged accessor’s offset and
+/// > range. ... In addition, the accessor’s iterator functions iterate only
+/// > over the elements that are within the sub-range.
+
+namespace sycl {
+__SYCL_INLINE_VER_NAMESPACE(_V1) {
+
+template <typename AccessorDataT, int AccessorDimensions,
+          access::mode AccessMode, access::target AccessTarget,
+          access::placeholder IsPlaceholder, typename PropertyListT>
+class accessor;
+
+namespace detail {
+
+template <typename DataT, int Dimensions> class accessor_iterator {
+public:
+  using difference_type = std::ptrdiff_t;
+  using value_type = DataT;
+  // FIXME: this should likely include address space
+  using pointer = DataT *;
+  using reference = DataT &;
+  using iterator_category = std::random_access_iterator_tag;
+
+  accessor_iterator() = default;
+
+  reference operator*() { return *(MDataPtr + getAbsoluteOffsetToBuffer()); }
+
+  accessor_iterator &operator++() {
+    ++MLinearId;
+    return *this;
+  }
+
+  accessor_iterator operator++(int) {
+    auto Old = *this;
+    ++(*this);
+    return Old;
+  }
+
+  accessor_iterator &operator--() {
+    --MLinearId;
+    return *this;
+  }
+
+  accessor_iterator operator--(int) {
+    auto Old = *this;
+    --(*this);
+    return Old;
+  }
+
+  accessor_iterator &operator+=(difference_type N) {
+    MLinearId += N;
+
+    return *this;
+  }
+
+  friend accessor_iterator operator+(const accessor_iterator &Lhs,
+                                     difference_type N) {
+    auto Ret = Lhs;
+    Ret += N;
+    return Ret;
+  }
+
+  friend accessor_iterator operator+(difference_type N,
+                                     const accessor_iterator &Rhs) {
+    auto Ret = Rhs;
+    Ret += N;
+    return Ret;
+  }
+
+  accessor_iterator &operator-=(difference_type N) {
+    MLinearId -= N;
+
+    return *this;
+  }
+
+  friend accessor_iterator operator-(accessor_iterator &Lhs,
+                                     difference_type N) {
+    Lhs -= N;
+    return Lhs;
+  }
+
+  reference &operator[](difference_type N) {
+    auto Copy = *this;
+    Copy += N;
+    return *Copy;
+  }
+
+  bool operator<(const accessor_iterator &Other) const {
+    return MLinearId < Other.MLinearId;
+  }
+
+  bool operator>(const accessor_iterator &Other) const { return Other < *this; }
+
+  bool operator<=(const accessor_iterator &Other) const {
+    return !(*this > Other);
+  }
+
+  bool operator>=(const accessor_iterator &Other) const {
+    return !(*this < Other);
+  }
+
+  bool operator==(const accessor_iterator &Other) const {
+    return MLinearId == Other.MLinearId;
+  }
+
+  bool operator!=(const accessor_iterator &Other) const {
+    return !(*this == Other);
+  }
+
+  difference_type operator-(const accessor_iterator &Rhs) {
+    return MLinearId - Rhs.MLinearId;
+  }
+
+private:
+  template <typename AccessorDataT, int AccessorDimensions,
+            access::mode AccessMode, access::target AccessTarget,
+            access::placeholder IsPlaceholder, typename PropertyListT>
+  friend class sycl::accessor;
+
+  DataT *MDataPtr = nullptr;
+
+  // Stores a linear id of an accessor's buffer element the iterator points to.
+  // This id is relative to a range accessible through an accessor, i.e. it is
+  // limited by a space with top left corner defiend as accessor::get_offset()
+  // and bottom right corner defined as accesor::get_range().
+  size_t MLinearId = 0;
+
+  // Describes range of linear IDs accessible by the iterator. MEnd corresponds
+  // to ID of en element past the last accessible element of accessors's
+  // buffer.
+  size_t MBegin = 0;
+  size_t MEnd = 0;
+
+  // If set to true, then it indicates that accessor has its offset and/or range
+  // set to non-zero, i.e. it is a ranged accessor.
+  bool MAccessorIsRanged = false;
+
+  // Fields below are used (and changed to be non-zero) only if we deal with
+  // a ranged accessor.
+  //
+  // TODO: consider making their existance dependable on Dimensions template
+  // parameter, because not all of them are needed for all possible dimensions.
+
+  // Three field below allow us to calculate an absolute offset to an accessor's
+  // buffer to correctly identify a memory region which this iterator should
+  // point to. Comments below describe them using an iterator to the following
+  // accessor as an example:
+  //
+  //   buffer<int, 2> buf(input.data(), range<2>{5, 5});
+  //   auto acc = buf.get_access(range<2>{3, 3}, id<2>{1, 1});
+  //
+  // Such combination of buffer size, access range and offset is visualized
+  // below. Dot (.) symbols represent buffer elements NOT reacheable by the
+  // accessor; X symbols represent buffer elements which ARE reachable by the
+  // the accessor.
+  //
+  //   . . . . .
+  //   . X X X .
+  //   . X X X .
+  //   . X X X .
+  //   . . . . .
+  //
+  // MStaticOffset stores a number of elements which precede the first
+  // accessible element, calculated as if the buffer was linearized.
+  // For the example above, MStaticOffset would be equal to 6, because
+  // there is one full row before the first accessible element and a one more on
+  // the second line. "Static" in the name highlights that this is a constant
+  // element in an equation which calculates an absoulte offset to an accessor's
+  // buffer, it doesn't depend on the current state of the iterator.
+  //
+  // NOTE: MStaticOffset is set to 0 in 1D case even if the accessor was
+  // created with offset: it is done to further optimize 1D case by
+  // incorporating that offset into MLinearId right away.
+  //
+  // MPerRowOffset stores a number of _inaccessible_ elements in each
+  // _accessible_ row. For the example above it would be equal to 2 (leftmost
+  // and the rightmost elements of a row).
+  //
+  // MPerSliceOffset stores a number of _inaccessible_ elements in each
+  // _accessible_ slice. Slice here means a single 2D layer in a 3D buffer. For
+  // the example above it would be equal to 0, because we are not looking at a
+  // 3D buffer. However, if we had two slices like visualized above,
+  // MPerSliceOffset would be equal to 16 (elements on the "perimeter" of the
+  // slice, i.e. ones represented as dots (.)).
+
+  size_t MStaticOffset = 0;
+  size_t MPerRowOffset = 0;
+  size_t MPerSliceOffset = 0;
+
+  // Contains a number of _accessible_ elements in a row
+  size_t MRowSize = 0;
+  // Contains a number of _accessible_ elements in a slice
+  size_t MSliceSize = 0;
+
+  // MLinearId stores an offset which is relative to the accessible range of
+  // the accessor, which means that it could be the case that MlinearId equal
+  // to 0 should not correspond to the beginning of the underlying buffer, but
+  // instead should be re-adjusted to account for an offset passed to the
+  // accessor constructor.
+  //
+  // This function performs necessary calculations to make sure that all
+  // access ranges and offsets are taken into account.
+  size_t getAbsoluteOffsetToBuffer() {
+    // For 1D case, any possible offsets are already incorporated into
+    // MLinearId, so 1D is always treated as a non-ranged accessor
+    if (!MAccessorIsRanged || Dimensions == 1)
+      return MLinearId;
+
+    // Here we need to deal with 2D or 3D ranged accessor.
+    // MLinearId points to an element relative to the accessible range. It
+    // should be adjusted to account for elements which are outside of the
+    // accessible range of the accessor.
+
+    // We start with static offset: that is a number of elements in full rows
+    // and full slices before the first accessible element.
+    size_t AbsoluteId = MLinearId + MStaticOffset;
+
+    // Then we account for inaccessible elements in each full slice
+    size_t Remaining = MLinearId;
+    if constexpr (Dimensions == 3) {
+      AbsoluteId += MPerSliceOffset * (Remaining / MSliceSize);
+      Remaining %= MSliceSize;
+    }
+
+    // Then we account for inaccessible elements in each full row
+    AbsoluteId += MPerRowOffset * (Remaining / MRowSize);
+    Remaining %= MRowSize;
+
+    return AbsoluteId;
+  }
+
+  accessor_iterator(DataT *DataPtr, const range<Dimensions> &MemoryRange,
+                    const range<Dimensions> &AccessRange,
+                    const id<Dimensions> &Offset)
+      : MDataPtr(DataPtr) {
+    constexpr int XIndex = Dimensions - 1;
+    constexpr int YIndex = Dimensions - 2;
+    (void)YIndex;
+    constexpr int ZIndex = Dimensions - 3;
+    (void)ZIndex;
+
+    if constexpr (Dimensions > 1)
+      MRowSize = AccessRange[XIndex];
+    if constexpr (Dimensions > 2)
+      MSliceSize = AccessRange[YIndex] * MRowSize;
+
+    if (id<Dimensions>{} != Offset)
+      MAccessorIsRanged = true;
+    else {
+      for (size_t I = 0; I < Dimensions; ++I)
+        if (AccessRange[I] != MemoryRange[I])
+          MAccessorIsRanged = true;
+    }
+
+    if (MAccessorIsRanged) {
+      if constexpr (Dimensions > 2) {
+        MStaticOffset +=
+            MemoryRange[XIndex] * MemoryRange[YIndex] * Offset[ZIndex];
+        MPerSliceOffset =
+            MemoryRange[XIndex] * MemoryRange[YIndex] - MSliceSize;
+      }
+      if constexpr (Dimensions > 1) {
+        // Elements in fully inaccessible rows
+        MStaticOffset += MemoryRange[XIndex] * Offset[YIndex];
+        MPerRowOffset = MemoryRange[XIndex] - MRowSize;
+      }
+
+      // Elements from the first accessible row
+      if constexpr (Dimensions == 1)
+        // To further optimize 1D case, offset is already included into Begin
+        MBegin = Offset[XIndex];
+      else
+        MStaticOffset += Offset[XIndex];
+    }
+
+    MEnd = MBegin + AccessRange.size();
+  }
+
+  static accessor_iterator getBegin(DataT *DataPtr,
+                                    const range<Dimensions> &MemoryRange,
+                                    const range<Dimensions> &AccessRange,
+                                    const id<Dimensions> &Offset) {
+    auto It = accessor_iterator(DataPtr, MemoryRange, AccessRange, Offset);
+    It.MLinearId = It.MBegin;
+    return It;
+  }
+
+  static accessor_iterator getEnd(DataT *DataPtr,
+                                  const range<Dimensions> &MemoryRange,
+                                  const range<Dimensions> &AccessRange,
+                                  const id<Dimensions> &Offset) {
+    auto It = accessor_iterator(DataPtr, MemoryRange, AccessRange, Offset);
+    It.MLinearId = It.MEnd;
+    return It;
+  }
+
+public:
+#ifndef NDEBUG
+  // Could be useful for debugging, but not a part of the official API,
+  // therefore only available in builds with assertions enabled.
+  friend std::ostream &operator<<(std::ostream &os,
+                                  const accessor_iterator &it) {
+    os << "accessor_iterator {\n";
+    os << "\tMLinearId: " << it.MLinearId << "\n";
+    os << "\tMEnd: " << it.MEnd << "\n";
+    os << "\tMStaticOffset: " << it.MStaticOffset << "\n";
+    os << "\tMPerRowOffset: " << it.MPerRowOffset << "\n";
+    os << "\tMPerSliceOffset: " << it.MPerSliceOffset << "\n";
+    os << "\tMRowSize: " << it.MRowSize << "\n";
+    os << "\tMSliceSize: " << it.MSliceSize << "\n";
+    os << "\tMAccessorIsRanged: " << it.MAccessorIsRanged << "\n";
+    os << "}";
+    return os;
+  }
+#endif // NDEBUG
+};
+} // namespace detail
+} // __SYCL_INLINE_VER_NAMESPACE(_V1)
+} // namespace sycl
diff --git a/sycl/unittests/CMakeLists.txt b/sycl/unittests/CMakeLists.txt
index 812e06cde7356..b30a6680365e3 100644
--- a/sycl/unittests/CMakeLists.txt
+++ b/sycl/unittests/CMakeLists.txt
@@ -47,3 +47,4 @@ add_subdirectory(windows)
 add_subdirectory(event)
 add_subdirectory(buffer)
 add_subdirectory(context)
+add_subdirectory(accessor)
diff --git a/sycl/unittests/accessor/AccessorIterator.cpp b/sycl/unittests/accessor/AccessorIterator.cpp
new file mode 100644
index 0000000000000..c67dbc7af51cd
--- /dev/null
+++ b/sycl/unittests/accessor/AccessorIterator.cpp
@@ -0,0 +1,513 @@
+#include <gtest/gtest.h>
+
+#include <sycl/sycl.hpp>
+
+#include <iterator>
+#include <numeric>
+#include <type_traits>
+#include <vector>
+
+class AccessorIteratorTest : public ::testing::Test {
+public:
+  template <int Dimensions, typename T = int>
+  void checkWriteThroughIterator(const sycl::range<Dimensions> &fullShape,
+                                 const sycl::range<Dimensions> &fillShape,
+                                 const sycl::id<Dimensions> &offset) {
+    std::vector<T> data(fullShape.size(), T{});
+    sycl::buffer buffer(data.data(), fullShape);
+    {
+      auto accessor = buffer.template get_access<sycl::access_mode::write>(
+          fillShape, offset);
+      T linear_id = 1;
+      for (auto it = accessor.begin(), e = accessor.end(); it != e; ++it) {
+        *it = linear_id;
+        linear_id += 1;
+      }
+    }
+
+    sycl::id<3> offsetToUse(Dimensions > 2 ? offset[Dimensions - 3] : 0,
+                            Dimensions > 1 ? offset[Dimensions - 2] : 0,
+                            offset[Dimensions - 1]);
+
+    sycl::id<3> shapeToCheck(
+        (Dimensions > 2 ? fillShape[Dimensions - 3] : 1) + offsetToUse[0],
+        (Dimensions > 1 ? fillShape[Dimensions - 2] : 1) + offsetToUse[1],
+        fillShape[Dimensions - 1] + offsetToUse[2]);
+
+    auto fullAccessor = buffer.template get_access<sycl::access_mode::read>();
+    T linear_id = 1;
+    for (size_t z = offsetToUse[0]; z < shapeToCheck[0]; ++z) {
+      for (size_t y = offsetToUse[1]; y < shapeToCheck[1]; ++y) {
+        for (size_t x = offsetToUse[2]; x < shapeToCheck[2]; ++x) {
+          auto value = accessHelper<Dimensions>(fullAccessor, z, y, x);
+          ASSERT_EQ(linear_id, value);
+          linear_id += 1;
+        }
+      }
+    }
+
+    sycl::id<3> adjustedFullShape(
+        Dimensions > 2 ? fullShape[Dimensions - 3] : 1,
+        Dimensions > 1 ? fullShape[Dimensions - 2] : 1,
+        fullShape[Dimensions - 1]);
+
+    for (size_t z = 0; z < adjustedFullShape[0]; ++z) {
+      for (size_t y = 0; y < adjustedFullShape[1]; ++y) {
+        for (size_t x = 0; x < adjustedFullShape[2]; ++x) {
+          // Skip elements which we previously checked
+          if (z >= offsetToUse[0] && z < shapeToCheck[0] &&
+              y >= offsetToUse[1] && y < shapeToCheck[1] &&
+              x >= offsetToUse[2] && x < shapeToCheck[2])
+            continue;
+          auto value = accessHelper<Dimensions>(fullAccessor, z, y, x);
+          ASSERT_EQ(T{}, value) << "at (" << z << "; " << y << "; " << x << ")";
+        }
+      }
+    }
+  }
+
+  template <int Dimensions, typename T = int>
+  void checkFullCopyThroughIterator(const sycl::range<Dimensions> &shape) {
+    std::vector<T> reference(shape.size());
+    std::iota(reference.begin(), reference.end(), 0);
+    sycl::buffer<T, Dimensions> buffer(reference.data(), shape);
+    auto accessor = buffer.template get_access<sycl::access_mode::read_write>();
+
+    ASSERT_NO_FATAL_FAILURE(checkFullCopyThroughIteratorImpl(
+        reference, accessor.begin(), accessor.end()));
+    ASSERT_NO_FATAL_FAILURE(checkFullCopyThroughIteratorImpl(
+        reference, accessor.cbegin(), accessor.cend()));
+  }
+
+  template <int Dimensions, typename T = int>
+  void
+  checkPartialCopyThroughIterator(const sycl::range<Dimensions> &fullShape,
+                                  const sycl::range<Dimensions> &copyShape,
+                                  const sycl::id<Dimensions> &offset = {}) {
+    std::vector<T> reference(fullShape.size());
+    std::iota(reference.begin(), reference.end(), 0);
+    sycl::buffer<T, Dimensions> buffer(reference.data(), fullShape);
+    std::vector<T> copied;
+
+    {
+      auto accessor = buffer.template get_access<sycl::access_mode::read_write>(
+          copyShape, offset);
+      copied = copyThroughIterators<T>(accessor.begin(), accessor.end());
+    }
+    ASSERT_NO_FATAL_FAILURE(
+        validatePartialCopyThroughIterator(copied, buffer, copyShape, offset));
+
+    {
+      auto accessor = buffer.template get_access<sycl::access_mode::read_write>(
+          copyShape, offset);
+      copied = copyThroughIterators<T>(accessor.cbegin(), accessor.cend());
+    }
+    ASSERT_NO_FATAL_FAILURE(
+        validatePartialCopyThroughIterator(copied, buffer, copyShape, offset));
+  }
+
+private:
+  template <typename IteratorT, typename T = int>
+  void checkFullCopyThroughIteratorImpl(const std::vector<T> &reference,
+                                        IteratorT begin, IteratorT end) {
+    std::vector<T> copied = copyThroughIterators<T>(begin, end);
+
+    ASSERT_EQ(copied.size(), reference.size());
+    for (size_t i = 0, e = reference.size(); i < e; ++i) {
+      ASSERT_EQ(copied[i], reference[i]);
+    }
+  }
+
+  template <typename T, typename IteratorT>
+  std::vector<T> copyThroughIterators(IteratorT begin, IteratorT end) {
+    std::vector<T> copied;
+    for (auto it = begin; it != end; ++it)
+      copied.push_back(*it);
+
+    return copied;
+  }
+
+  template <int Dimensions, typename T = int>
+  void
+  validatePartialCopyThroughIterator(const std::vector<T> &copied,
+                                     sycl::buffer<T, Dimensions> &buffer,
+                                     const sycl::range<Dimensions> &copyShape,
+                                     const sycl::id<Dimensions> &offset = {}) {
+    auto fullAccessor = buffer.template get_access<sycl::access_mode::read>();
+    size_t linearId = 0;
+
+    sycl::id<3> offsetToUse(Dimensions > 2 ? offset[Dimensions - 3] : 1,
+                            Dimensions > 1 ? offset[Dimensions - 2] : 1,
+                            offset[Dimensions - 1]);
+
+    sycl::id<3> shapeToCheck(
+        (Dimensions > 2 ? copyShape[Dimensions - 3] : 1) + offsetToUse[0],
+        (Dimensions > 1 ? copyShape[Dimensions - 2] : 1) + offsetToUse[1],
+        copyShape[Dimensions - 1] + offsetToUse[2]);
+
+    for (size_t z = offsetToUse[0]; z < shapeToCheck[0]; ++z) {
+      for (size_t y = offsetToUse[1]; y < shapeToCheck[1]; ++y) {
+        for (size_t x = offsetToUse[2]; x < shapeToCheck[2]; ++x) {
+          auto value = accessHelper<Dimensions>(fullAccessor, z, y, x);
+          ASSERT_EQ(copied[linearId], value);
+          ++linearId;
+        }
+      }
+    }
+  }
+
+  template <int TotalDimensions, int CurrentDimension = 3, typename Container,
+            typename... Indices>
+  auto &&accessHelper(Container &&C, int Idx, Indices... Ids) {
+    if constexpr (CurrentDimension > TotalDimensions) {
+      (void)Idx;
+      return accessHelper<TotalDimensions, CurrentDimension - 1>(C, Ids...);
+    } else
+      return accessHelper<TotalDimensions, CurrentDimension - 1>(C[Idx],
+                                                                 Ids...);
+  }
+
+  template <int TotalDimensions, int CurrentDimension = 3, typename Container>
+  auto &&accessHelper(Container &&C, int Idx) {
+    return C[Idx];
+  }
+};
+
+// FIXME: consider turning this into parameterized test to check various
+// accessor types
+TEST_F(AccessorIteratorTest, IteratorTraits) {
+  using IteratorT = sycl::accessor<int>::iterator;
+  ASSERT_TRUE(
+      (std::is_same_v<sycl::accessor<int>::difference_type,
+                      std::iterator_traits<IteratorT>::difference_type>));
+  ASSERT_TRUE((std::is_same_v<sycl::accessor<int>::value_type,
+                              std::iterator_traits<IteratorT>::value_type>));
+  ASSERT_TRUE((std::is_same_v<sycl::accessor<int>::value_type *,
+                              std::iterator_traits<IteratorT>::pointer>));
+  ASSERT_TRUE((std::is_same_v<sycl::accessor<int>::reference,
+                              std::iterator_traits<IteratorT>::reference>));
+  ASSERT_TRUE(
+      (std::is_same_v<std::random_access_iterator_tag,
+                      std::iterator_traits<IteratorT>::iterator_category>));
+}
+
+// Based on requirements listed at
+// https://en.cppreference.com/w/cpp/named_req/RandomAccessIterator
+TEST_F(AccessorIteratorTest, LegacyRandomAccessIteratorRequirements) {
+  using IteratorT = sycl::accessor<int>::iterator;
+  IteratorT It;
+  auto &RefToIt = It;
+  ASSERT_TRUE((std::is_same_v<IteratorT &, decltype(RefToIt += 3)>));
+  ASSERT_TRUE((std::is_same_v<IteratorT, decltype(It + 3)>));
+  ASSERT_TRUE((std::is_same_v<IteratorT, decltype(3 + It)>));
+  ASSERT_TRUE((std::is_same_v<IteratorT &, decltype(RefToIt -= 3)>));
+  ASSERT_TRUE((std::is_same_v<IteratorT, decltype(It - 3)>));
+  IteratorT It2;
+  ASSERT_TRUE((std::is_same_v<std::iterator_traits<IteratorT>::difference_type,
+                              decltype(It - It2)>));
+  ASSERT_TRUE(
+      (std::is_convertible_v<decltype(It[3]),
+                             std::iterator_traits<IteratorT>::reference>));
+  ASSERT_TRUE((std::is_convertible_v<decltype(It < It2), bool>));
+  ASSERT_TRUE((std::is_convertible_v<decltype(It <= It2), bool>));
+  ASSERT_TRUE((std::is_convertible_v<decltype(It > It2), bool>));
+  ASSERT_TRUE((std::is_convertible_v<decltype(It >= It2), bool>));
+}
+
+// Based on notes listed at
+// https://en.cppreference.com/w/cpp/named_req/RandomAccessIterator
+TEST_F(AccessorIteratorTest, LegacyRandomAccessIteratorRequirementsExtra) {
+  std::vector<int> reference(6);
+  std::iota(reference.begin(), reference.end(), 0);
+  sycl::buffer<int> buffer(reference.data(), sycl::range<1>{reference.size()});
+  auto accessor = buffer.template get_access<sycl::access_mode::read_write>();
+  auto It = accessor.begin();
+  It += 3;
+
+  { // It += n should be equivalent to incrementint/decrementing It n times
+    // The test also checks the same for operator +, i.e. It + n
+    for (int n = -3; n <= 3; ++n) {
+      auto It1 = It;
+      auto It2 = It;
+      It1 += n;
+
+      if (n < 0) {
+        int i = n;
+        while (i++)
+          --It2;
+      } else {
+        int i = n;
+        while (i--)
+          ++It2;
+      }
+
+      ASSERT_EQ(It1, It2) << " with n = " << n;
+      ASSERT_EQ(It + n, It2) << " with n = " << n;
+    }
+  }
+
+  { // It + n == n + It
+    for (int n = -3; n <= 3; ++n) {
+      ASSERT_EQ(It + n, n + It);
+    }
+  }
+
+  {
+    auto It1 = accessor.begin();
+    auto It2 = accessor.end();
+    ASSERT_EQ(std::abs(It - It1), std::abs(It1 - It));
+    ASSERT_EQ(std::abs(It - It2), std::abs(It2 - It));
+    ASSERT_EQ(It1 - It, -3);
+    ASSERT_EQ(It - It1, 3);
+    ASSERT_EQ(It2, It + (It2 - It));
+    ASSERT_EQ(It, It1 + (It - It1));
+  }
+
+  {
+    auto It1 = accessor.begin();
+    auto It2 = accessor.begin();
+    auto It3 = accessor.end();
+
+    ASSERT_TRUE(!(It1 < It2));
+    ASSERT_TRUE(It1 < It); // precondition for the next check
+    ASSERT_TRUE(!(It < It1));
+    ASSERT_TRUE(It < It3); // precondition for the next check
+    ASSERT_TRUE(It1 < It3);
+
+    ASSERT_FALSE(It3 < It);
+    ASSERT_FALSE(It == It3);
+  }
+}
+
+// Based on requirements listed at
+// https://en.cppreference.com/w/cpp/named_req/BidirectionalIterator
+TEST_F(AccessorIteratorTest, LegacyBidirectionalIteratorRequirements) {
+  using IteratorT = sycl::accessor<int>::iterator;
+  IteratorT It;
+  ASSERT_TRUE((std::is_same_v<IteratorT &, decltype(--It)>));
+  ASSERT_TRUE((std::is_convertible_v<decltype(It--), const IteratorT &>));
+  ASSERT_TRUE((std::is_same_v<std::iterator_traits<IteratorT>::reference,
+                              decltype(*It--)>));
+}
+
+// Based on requirements listed at
+// https://en.cppreference.com/w/cpp/named_req/ForwardIterator
+TEST_F(AccessorIteratorTest, LegacyForwardIteratorRequirements) {
+  using IteratorT = sycl::accessor<int>::iterator;
+  ASSERT_TRUE(std::is_default_constructible_v<IteratorT>);
+  IteratorT It;
+  ASSERT_TRUE((std::is_same_v<IteratorT, decltype(It++)>));
+  ASSERT_TRUE((std::is_same_v<std::iterator_traits<IteratorT>::reference,
+                              decltype(*It++)>));
+  IteratorT It2;
+  ASSERT_TRUE((std::is_convertible_v<decltype(It == It2), bool>));
+  ASSERT_TRUE((std::is_convertible_v<decltype(It != It2), bool>));
+}
+
+TEST_F(AccessorIteratorTest, MultipassGuarantee) {
+  std::vector<int> reference(5);
+  std::iota(reference.begin(), reference.end(), 0);
+  sycl::buffer<int> buffer(reference.data(), sycl::range<1>{reference.size()});
+  auto accessor = buffer.template get_access<sycl::access_mode::read_write>();
+  auto It1 = accessor.begin();
+  auto It2 = accessor.begin();
+
+  while (It1 != accessor.end()) {
+    ASSERT_EQ(It1, It2);
+    ASSERT_EQ(*It1, *It2);
+    ASSERT_EQ(++It1, ++It2);
+  }
+
+  It1 = accessor.begin();
+  It2 = It1;
+  ASSERT_EQ(((void)++It2, *It1), *It1);
+}
+
+// Based on requirements listead at
+// https://en.cppreference.com/w/cpp/named_req/Iterator
+TEST_F(AccessorIteratorTest, LegacyIteratorRequirements) {
+  using IteratorT = sycl::accessor<int>::iterator;
+  ASSERT_TRUE(std::is_copy_constructible_v<IteratorT>);
+  ASSERT_TRUE(std::is_copy_assignable_v<IteratorT>);
+  ASSERT_TRUE(std::is_destructible_v<IteratorT>);
+  ASSERT_TRUE(std::is_swappable_v<IteratorT>);
+  IteratorT It;
+  ASSERT_TRUE((std::is_same_v<IteratorT &, decltype(++It)>));
+  ASSERT_TRUE((std::is_same_v<std::iterator_traits<IteratorT>::reference,
+                              decltype(*It)>));
+}
+
+TEST_F(AccessorIteratorTest, FullCopy1D) {
+  ASSERT_NO_FATAL_FAILURE(checkFullCopyThroughIterator(sycl::range<1>{10}));
+}
+
+TEST_F(AccessorIteratorTest, FullCopy2D) {
+  ASSERT_NO_FATAL_FAILURE(checkFullCopyThroughIterator(sycl::range<2>{2, 5}));
+  ASSERT_NO_FATAL_FAILURE(checkFullCopyThroughIterator(sycl::range<2>{5, 2}));
+  ASSERT_NO_FATAL_FAILURE(checkFullCopyThroughIterator(sycl::range<2>{1, 10}));
+  ASSERT_NO_FATAL_FAILURE(checkFullCopyThroughIterator(sycl::range<2>{10, 1}));
+}
+
+TEST_F(AccessorIteratorTest, FullCopy3D) {
+  ASSERT_NO_FATAL_FAILURE(
+      checkFullCopyThroughIterator(sycl::range<3>{3, 3, 3}));
+  ASSERT_NO_FATAL_FAILURE(
+      checkFullCopyThroughIterator(sycl::range<3>{1, 3, 3}));
+  ASSERT_NO_FATAL_FAILURE(
+      checkFullCopyThroughIterator(sycl::range<3>{3, 1, 3}));
+  ASSERT_NO_FATAL_FAILURE(
+      checkFullCopyThroughIterator(sycl::range<3>{3, 3, 1}));
+}
+
+TEST_F(AccessorIteratorTest, PartialCopyWithoutOffset1D) {
+  ASSERT_NO_FATAL_FAILURE(
+      checkPartialCopyThroughIterator(sycl::range<1>{10}, sycl::range<1>{5}));
+  ASSERT_NO_FATAL_FAILURE(
+      checkPartialCopyThroughIterator(sycl::range<1>{10}, sycl::range<1>{10}));
+}
+
+TEST_F(AccessorIteratorTest, PartialCopyWithoutOffset2D) {
+  ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIterator(
+      sycl::range<2>{5, 5}, sycl::range<2>{3, 3}));
+  ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIterator(
+      sycl::range<2>{5, 5}, sycl::range<2>{5, 5}));
+  ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIterator(
+      sycl::range<2>{5, 5}, sycl::range<2>{2, 5}));
+  ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIterator(
+      sycl::range<2>{5, 5}, sycl::range<2>{5, 2}));
+  ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIterator(
+      sycl::range<2>{5, 5}, sycl::range<2>{3, 2}));
+}
+
+TEST_F(AccessorIteratorTest, PartialCopyWithoutOffset3D) {
+  ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIterator(
+      sycl::range<3>{5, 5, 5}, sycl::range<3>{3, 3, 3}));
+  ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIterator(
+      sycl::range<3>{5, 5, 5}, sycl::range<3>{5, 3, 3}));
+  ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIterator(
+      sycl::range<3>{5, 5, 5}, sycl::range<3>{3, 5, 3}));
+  ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIterator(
+      sycl::range<3>{5, 5, 5}, sycl::range<3>{3, 3, 5}));
+  ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIterator(
+      sycl::range<3>{5, 5, 5}, sycl::range<3>{5, 5, 5}));
+  ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIterator(
+      sycl::range<3>{5, 5, 5}, sycl::range<3>{1, 2, 3}));
+}
+
+TEST_F(AccessorIteratorTest, PartialCopyWithOffset1D) {
+  ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIterator(
+      sycl::range<1>{10}, sycl::range<1>{5}, sycl::id<1>{3}));
+  ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIterator(
+      sycl::range<1>{10}, sycl::range<1>{5}, sycl::id<1>{5}));
+}
+
+TEST_F(AccessorIteratorTest, PartialCopyWithOffset2D) {
+  ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIterator(
+      sycl::range<2>{10, 10}, sycl::range<2>{5, 5}, sycl::id<2>{3, 3}));
+  ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIterator(
+      sycl::range<2>{10, 10}, sycl::range<2>{5, 5}, sycl::id<2>{3, 0}));
+  ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIterator(
+      sycl::range<2>{10, 10}, sycl::range<2>{5, 5}, sycl::id<2>{0, 3}));
+  ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIterator(
+      sycl::range<2>{10, 10}, sycl::range<2>{5, 5}, sycl::id<2>{5, 5}));
+  ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIterator(
+      sycl::range<2>{5, 10}, sycl::range<2>{3, 5}, sycl::id<2>{1, 4}));
+  ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIterator(
+      sycl::range<2>{10, 5}, sycl::range<2>{5, 3}, sycl::id<2>{5, 1}));
+}
+
+TEST_F(AccessorIteratorTest, PartialCopyWithOffset3D) {
+  ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIterator(
+      sycl::range<3>{7, 7, 7}, sycl::range<3>{3, 3, 3}, sycl::id<3>{2, 2, 2}));
+  ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIterator(
+      sycl::range<3>{8, 8, 8}, sycl::range<3>{4, 4, 4}, sycl::id<3>{4, 4, 4}));
+  // FIXME: figure out why the test below fails
+  // ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIterator(
+  //     sycl::range<3>{7, 7, 7}, sycl::range<3>{3, 3, 3}, sycl::id<3>{4, 4,
+  //     4}));
+  ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIterator(
+      sycl::range<3>{7, 7, 7}, sycl::range<3>{3, 4, 5}, sycl::id<3>{3, 2, 1}));
+  ASSERT_NO_FATAL_FAILURE(checkPartialCopyThroughIterator(
+      sycl::range<3>{9, 8, 7}, sycl::range<3>{3, 4, 5}, sycl::id<3>{3, 2, 1}));
+}
+
+TEST_F(AccessorIteratorTest, FullWrite1D) {
+  ASSERT_NO_FATAL_FAILURE(checkWriteThroughIterator(
+      sycl::range<1>{10}, sycl::range<1>{10}, sycl::id<1>{0}));
+}
+
+TEST_F(AccessorIteratorTest, FullWrite2D) {
+  ASSERT_NO_FATAL_FAILURE(checkWriteThroughIterator(
+      sycl::range<2>{5, 5}, sycl::range<2>{5, 5}, sycl::id<2>{0, 0}));
+  ASSERT_NO_FATAL_FAILURE(checkWriteThroughIterator(
+      sycl::range<2>{2, 5}, sycl::range<2>{2, 5}, sycl::id<2>{0, 0}));
+  ASSERT_NO_FATAL_FAILURE(checkWriteThroughIterator(
+      sycl::range<2>{5, 2}, sycl::range<2>{5, 2}, sycl::id<2>{0, 0}));
+}
+
+TEST_F(AccessorIteratorTest, FullWrite3D) {
+  ASSERT_NO_FATAL_FAILURE(checkWriteThroughIterator(
+      sycl::range<3>{5, 5, 5}, sycl::range<3>{5, 5, 5}, sycl::id<3>{0, 0, 0}));
+  ASSERT_NO_FATAL_FAILURE(checkWriteThroughIterator(
+      sycl::range<3>{1, 5, 5}, sycl::range<3>{1, 5, 5}, sycl::id<3>{0, 0, 0}));
+  ASSERT_NO_FATAL_FAILURE(checkWriteThroughIterator(
+      sycl::range<3>{5, 1, 5}, sycl::range<3>{5, 1, 5}, sycl::id<3>{0, 0, 0}));
+  ASSERT_NO_FATAL_FAILURE(checkWriteThroughIterator(
+      sycl::range<3>{5, 5, 1}, sycl::range<3>{5, 5, 1}, sycl::id<3>{0, 0, 0}));
+  ASSERT_NO_FATAL_FAILURE(checkWriteThroughIterator(
+      sycl::range<3>{3, 6, 4}, sycl::range<3>{3, 6, 4}, sycl::id<3>{0, 0, 0}));
+}
+
+TEST_F(AccessorIteratorTest, PartialWriteWithoutOffset1D) {
+  ASSERT_NO_FATAL_FAILURE(checkWriteThroughIterator(
+      sycl::range<1>{10}, sycl::range<1>{5}, sycl::id<1>{0}));
+}
+
+TEST_F(AccessorIteratorTest, PartialWriteWithoutOffset2D) {
+  ASSERT_NO_FATAL_FAILURE(checkWriteThroughIterator(
+      sycl::range<2>{5, 5}, sycl::range<2>{3, 3}, sycl::id<2>{0, 0}));
+  ASSERT_NO_FATAL_FAILURE(checkWriteThroughIterator(
+      sycl::range<2>{2, 5}, sycl::range<2>{1, 3}, sycl::id<2>{0, 0}));
+  ASSERT_NO_FATAL_FAILURE(checkWriteThroughIterator(
+      sycl::range<2>{5, 2}, sycl::range<2>{3, 1}, sycl::id<2>{0, 0}));
+}
+
+TEST_F(AccessorIteratorTest, PartialWriteWithoutOffset3D) {
+  ASSERT_NO_FATAL_FAILURE(checkWriteThroughIterator(
+      sycl::range<3>{5, 5, 5}, sycl::range<3>{3, 3, 3}, sycl::id<3>{0, 0, 0}));
+  ASSERT_NO_FATAL_FAILURE(checkWriteThroughIterator(
+      sycl::range<3>{1, 5, 5}, sycl::range<3>{0, 3, 3}, sycl::id<3>{0, 0, 0}));
+  ASSERT_NO_FATAL_FAILURE(checkWriteThroughIterator(
+      sycl::range<3>{5, 1, 5}, sycl::range<3>{3, 1, 3}, sycl::id<3>{0, 0, 0}));
+  ASSERT_NO_FATAL_FAILURE(checkWriteThroughIterator(
+      sycl::range<3>{5, 5, 1}, sycl::range<3>{3, 3, 1}, sycl::id<3>{0, 0, 0}));
+  ASSERT_NO_FATAL_FAILURE(checkWriteThroughIterator(
+      sycl::range<3>{3, 6, 4}, sycl::range<3>{1, 3, 2}, sycl::id<3>{0, 0, 0}));
+}
+
+TEST_F(AccessorIteratorTest, PartialWriteWithOffset1D) {
+  ASSERT_NO_FATAL_FAILURE(checkWriteThroughIterator(
+      sycl::range<1>{10}, sycl::range<1>{5}, sycl::id<1>{3}));
+}
+
+TEST_F(AccessorIteratorTest, PartialWriteWithOffset2D) {
+  ASSERT_NO_FATAL_FAILURE(checkWriteThroughIterator(
+      sycl::range<2>{5, 5}, sycl::range<2>{3, 3}, sycl::id<2>{1, 1}));
+  ASSERT_NO_FATAL_FAILURE(checkWriteThroughIterator(
+      sycl::range<2>{3, 5}, sycl::range<2>{1, 3}, sycl::id<2>{1, 2}));
+  ASSERT_NO_FATAL_FAILURE(checkWriteThroughIterator(
+      sycl::range<2>{5, 3}, sycl::range<2>{3, 1}, sycl::id<2>{1, 1}));
+}
+
+TEST_F(AccessorIteratorTest, PartialWriteWithOffset3D) {
+  ASSERT_NO_FATAL_FAILURE(checkWriteThroughIterator(
+      sycl::range<3>{5, 5, 5}, sycl::range<3>{3, 3, 3}, sycl::id<3>{1, 1, 1}));
+  ASSERT_NO_FATAL_FAILURE(checkWriteThroughIterator(
+      sycl::range<3>{3, 5, 5}, sycl::range<3>{0, 3, 3}, sycl::id<3>{1, 2, 2}));
+  ASSERT_NO_FATAL_FAILURE(checkWriteThroughIterator(
+      sycl::range<3>{5, 2, 5}, sycl::range<3>{3, 1, 3}, sycl::id<3>{1, 1, 2}));
+  ASSERT_NO_FATAL_FAILURE(checkWriteThroughIterator(
+      sycl::range<3>{5, 5, 3}, sycl::range<3>{3, 3, 1}, sycl::id<3>{1, 1, 1}));
+  ASSERT_NO_FATAL_FAILURE(checkWriteThroughIterator(
+      sycl::range<3>{3, 6, 4}, sycl::range<3>{1, 3, 2}, sycl::id<3>{1, 3, 2}));
+}
diff --git a/sycl/unittests/accessor/CMakeLists.txt b/sycl/unittests/accessor/CMakeLists.txt
new file mode 100644
index 0000000000000..8080fd43cfe81
--- /dev/null
+++ b/sycl/unittests/accessor/CMakeLists.txt
@@ -0,0 +1,3 @@
+add_sycl_unittest(AccessorTests OBJECT
+  AccessorIterator.cpp
+)