diff --git a/sycl/include/CL/sycl/accessor.hpp b/sycl/include/CL/sycl/accessor.hpp
index 6ab337be1d075..aad08a26a2bf0 100644
--- a/sycl/include/CL/sycl/accessor.hpp
+++ b/sycl/include/CL/sycl/accessor.hpp
@@ -810,11 +810,11 @@ class accessor
            buffer<DataT, Dimensions>>::type &bufferRef)
 #ifdef __SYCL_DEVICE_ONLY__
       : impl((dataT *)detail::getSyclObjImpl(bufferRef)->BufPtr,
-             bufferRef.get_range(), bufferRef.get_range()) {
+             bufferRef.MemRange, bufferRef.MemRange) {
 #else
       : impl(std::make_shared<_ImplT>(
                  (dataT *)detail::getSyclObjImpl(bufferRef)->BufPtr,
-                 bufferRef.get_range(), bufferRef.get_range())) {
+                 bufferRef.MemRange, bufferRef.MemRange)) {
 #endif
     auto BufImpl = detail::getSyclObjImpl(bufferRef);
     if (AccessTarget == access::target::host_buffer) {
@@ -858,12 +858,12 @@ class accessor
       // Pass nullptr as a pointer to mem and use buffers from the ctor
       // arguments to avoid the need in adding utility functions for
       // dummy/default initialization of range fields.
-      : impl(nullptr, bufferRef.get_range(), bufferRef.get_range(),
+      : impl(nullptr, bufferRef.MemRange, bufferRef.MemRange,
              &commandGroupHandlerRef) {}
 #else
       : impl(std::make_shared<_ImplT>(
                  (dataT *)detail::getSyclObjImpl(bufferRef)->BufPtr,
-                 bufferRef.get_range(), bufferRef.get_range(),
+                 bufferRef.MemRange, bufferRef.MemRange,
                  &commandGroupHandlerRef)) {
     auto BufImpl = detail::getSyclObjImpl(bufferRef);
     if (BufImpl->OpenCLInterop && !BufImpl->isValidAccessToMem(accessMode)) {
@@ -906,11 +906,11 @@ class accessor
       // arguments to avoid the need in adding utility functions for
       // dummy/default initialization of range<Dimensions> and
       // id<Dimension> fields.
-      : impl(nullptr, Range, bufferRef.get_range(), Offset) {}
+      : impl(nullptr, Range, bufferRef.MemRange, Offset) {}
 #else   // !__SYCL_DEVICE_ONLY__
       : impl(std::make_shared<_ImplT>(
                  (dataT *)detail::getSyclObjImpl(bufferRef)->BufPtr, Range,
-                 bufferRef.get_range(), Offset)) {
+                 bufferRef.MemRange, Offset)) {
     auto BufImpl = detail::getSyclObjImpl(bufferRef);
     if (AccessTarget == access::target::host_buffer) {
       if (BufImpl->OpenCLInterop) {
@@ -956,12 +956,12 @@ class accessor
       // arguments to avoid the need in adding utility functions for
       // dummy/default initialization of range<Dimensions> and
       // id<Dimension> fields.
-      : impl(nullptr, Range, bufferRef.get_range(),
+      : impl(nullptr, Range, bufferRef.MemRange,
              &commandGroupHandlerRef, Offset) {}
 #else   // !__SYCL_DEVICE_ONLY__
       : impl(std::make_shared<_ImplT>(
                  (dataT *)detail::getSyclObjImpl(bufferRef)->BufPtr, Range,
-                 bufferRef.get_range(), &commandGroupHandlerRef, Offset)) {
+                 bufferRef.MemRange, &commandGroupHandlerRef, Offset)) {
     auto BufImpl = detail::getSyclObjImpl(bufferRef);
     if (BufImpl->OpenCLInterop && !BufImpl->isValidAccessToMem(accessMode)) {
       throw cl::sycl::runtime_error(
diff --git a/sycl/include/CL/sycl/buffer.hpp b/sycl/include/CL/sycl/buffer.hpp
index d3a43837a666f..1d9b2f19c2b07 100644
--- a/sycl/include/CL/sycl/buffer.hpp
+++ b/sycl/include/CL/sycl/buffer.hpp
@@ -33,46 +33,50 @@ class buffer {
 
   buffer(const range<dimensions> &bufferRange,
          const property_list &propList = {})
-      : Range(bufferRange) {
+      : Range(bufferRange), MemRange(bufferRange) {
     impl = std::make_shared<detail::buffer_impl<AllocatorT>>(
         get_count() * sizeof(T), propList);
   }
 
   buffer(const range<dimensions> &bufferRange, AllocatorT allocator,
-         const property_list &propList = {}) {
+         const property_list &propList = {})
+      : Range(bufferRange), MemRange(bufferRange) {
     impl = std::make_shared<detail::buffer_impl<AllocatorT>>(
         get_count() * sizeof(T), propList, allocator);
   }
 
   buffer(T *hostData, const range<dimensions> &bufferRange,
          const property_list &propList = {})
-      : Range(bufferRange) {
+      : Range(bufferRange), MemRange(bufferRange) {
     impl = std::make_shared<detail::buffer_impl<AllocatorT>>(
         hostData, get_count() * sizeof(T), propList);
   }
 
   buffer(T *hostData, const range<dimensions> &bufferRange,
-         AllocatorT allocator, const property_list &propList = {}) {
+         AllocatorT allocator, const property_list &propList = {})
+      : Range(bufferRange), MemRange(bufferRange) {
     impl = std::make_shared<detail::buffer_impl<AllocatorT>>(
         hostData, get_count() * sizeof(T), propList, allocator);
   }
 
   buffer(const T *hostData, const range<dimensions> &bufferRange,
          const property_list &propList = {})
-      : Range(bufferRange) {
+      : Range(bufferRange), MemRange(bufferRange) {
     impl = std::make_shared<detail::buffer_impl<AllocatorT>>(
         hostData, get_count() * sizeof(T), propList);
   }
 
   buffer(const T *hostData, const range<dimensions> &bufferRange,
-         AllocatorT allocator, const property_list &propList = {}) {
+         AllocatorT allocator, const property_list &propList = {})
+      : Range(bufferRange), MemRange(bufferRange) {
     impl = std::make_shared<detail::buffer_impl<AllocatorT>>(
         hostData, get_count() * sizeof(T), propList, allocator);
   }
 
   buffer(const shared_ptr_class<T> &hostData,
          const range<dimensions> &bufferRange, AllocatorT allocator,
-         const property_list &propList = {}) {
+         const property_list &propList = {})
+      : Range(bufferRange), MemRange(bufferRange) {
     impl = std::make_shared<detail::buffer_impl<AllocatorT>>(
         hostData, get_count() * sizeof(T), propList, allocator);
   }
@@ -80,7 +84,7 @@ class buffer {
   buffer(const shared_ptr_class<T> &hostData,
          const range<dimensions> &bufferRange,
          const property_list &propList = {})
-      : Range(bufferRange) {
+      : Range(bufferRange), MemRange(bufferRange) {
     impl = std::make_shared<detail::buffer_impl<AllocatorT>>(
         hostData, get_count() * sizeof(T), propList);
   }
@@ -89,7 +93,8 @@ class buffer {
             typename = EnableIfOneDimension<N>>
   buffer(InputIterator first, InputIterator last, AllocatorT allocator,
          const property_list &propList = {})
-      : Range(range<1>(std::distance(first, last))) {
+      : Range(range<1>(std::distance(first, last))),
+        MemRange(range<1>(std::distance(first, last))) {
     impl = std::make_shared<detail::buffer_impl<AllocatorT>>(
         first, last, get_count() * sizeof(T), propList, allocator);
   }
@@ -98,15 +103,16 @@ class buffer {
             typename = EnableIfOneDimension<N>>
   buffer(InputIterator first, InputIterator last,
          const property_list &propList = {})
-      : Range(range<1>(std::distance(first, last))) {
+      : Range(range<1>(std::distance(first, last))),
+        MemRange(range<1>(std::distance(first, last))) {
     impl = std::make_shared<detail::buffer_impl<AllocatorT>>(
         first, last, get_count() * sizeof(T), propList);
   }
 
-  // buffer(buffer<T, dimensions, AllocatorT> b, const id<dimensions>
-  // &baseIndex, const range<dimensions> &subRange) {
-  //     impl = std::make_shared<detail::buffer_impl>(b, baseIndex, subRange);
-  // }
+  buffer(buffer<T, dimensions, AllocatorT> &b, const id<dimensions> &baseIndex,
+         const range<dimensions> &subRange)
+      : impl(b.impl), Offset(baseIndex + b.Offset), Range(subRange), MemRange(b.MemRange),
+        IsSubBuffer(true) {}
 
   template <int N = dimensions, typename = EnableIfOneDimension<N>>
   buffer(cl_mem MemObject, const context &SyclContext,
@@ -116,6 +122,7 @@ class buffer {
     CHECK_OCL_CODE(clGetMemObjectInfo(MemObject, CL_MEM_SIZE, sizeof(size_t),
                                       &BufSize, nullptr));
     Range[0] = BufSize / sizeof(T);
+    MemRange[0] = BufSize / sizeof(T);
     impl = std::make_shared<detail::buffer_impl<AllocatorT>>(
         MemObject, SyclContext, BufSize, AvailableEvent);
   }
@@ -150,6 +157,9 @@ class buffer {
             access::target target = access::target::global_buffer>
   accessor<T, dimensions, mode, target, access::placeholder::false_t>
   get_access(handler &commandGroupHandler) {
+    if (IsSubBuffer)
+      return impl->template get_access<T, dimensions, mode, target>(
+          *this, commandGroupHandler, Range, Offset);
     return impl->template get_access<T, dimensions, mode, target>(
         *this, commandGroupHandler);
   }
@@ -158,6 +168,9 @@ class buffer {
   accessor<T, dimensions, mode, access::target::host_buffer,
            access::placeholder::false_t>
   get_access() {
+    if (IsSubBuffer)
+      return impl->template get_access<T, dimensions, mode>(*this, Range,
+                                                            Offset);
     return impl->template get_access<T, dimensions, mode>(*this);
   }
 
@@ -185,7 +198,7 @@ class buffer {
 
   void set_write_back(bool flag = true) { return impl->set_write_back(flag); }
 
-  // bool is_sub_buffer() const { return impl->is_sub_buffer(); }
+  bool is_sub_buffer() const { return IsSubBuffer; }
 
   template <typename ReinterpretT, int ReinterpretDim>
   buffer<ReinterpretT, ReinterpretDim, AllocatorT>
@@ -212,12 +225,22 @@ class buffer {
   template <class Obj>
   friend decltype(Obj::impl) detail::getSyclObjImpl(const Obj &SyclObject);
   template <typename A, int dims, typename C> friend class buffer;
+  template <typename DataT, int dims, access::mode mode,
+            access::target target, access::placeholder isPlaceholder>
+  friend class accessor;
+  // If this buffer is subbuffer - this range represents range of the parent
+  // buffer
+  range<dimensions> MemRange;
+  bool IsSubBuffer = false;
   range<dimensions> Range;
+  // If this buffer is sub-buffer - offset field specifies the origin of the
+  // sub-buffer inside the parent buffer
+  id<dimensions> Offset;
 
   // Reinterpret contructor
   buffer(shared_ptr_class<detail::buffer_impl<AllocatorT>> Impl,
          range<dimensions> reinterpretRange)
-      : impl(Impl), Range(reinterpretRange){};
+      : impl(Impl), Range(reinterpretRange), MemRange(reinterpretRange) {};
 };
 } // namespace sycl
 } // namespace cl
diff --git a/sycl/test/basic_tests/buffer/subbuffer.cpp b/sycl/test/basic_tests/buffer/subbuffer.cpp
new file mode 100644
index 0000000000000..92b091313b13d
--- /dev/null
+++ b/sycl/test/basic_tests/buffer/subbuffer.cpp
@@ -0,0 +1,89 @@
+// RUN: %clang -std=c++11 -fsycl %s -o %t.out -lstdc++ -lOpenCL
+// RUN: env SYCL_DEVICE_TYPE=HOST %t.out
+// RUN: %CPU_RUN_PLACEHOLDER %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// RUN: %ACC_RUN_PLACEHOLDER %t.out
+//==---------- subbuffer.cpp --- sub-buffer basic test ---------------------==//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+#include <CL/sycl.hpp>
+
+using namespace cl::sycl;
+
+int main() {
+
+  bool Failed = false;
+  // Basic test case
+  {
+    const int M = 6;
+    const int N = 7;
+    int Result[M][N] = {0};
+    {
+      auto OrigRange = range<2>(M, N);
+      buffer<int, 2> Buffer(OrigRange);
+      Buffer.set_final_data((int *)Result);
+      auto Offset = id<2>(1, 1);
+      auto SubRange = range<2>(M - 2, N - 2);
+      queue MyQueue;
+      buffer<int, 2> SubBuffer(Buffer, Offset, SubRange);
+      MyQueue.submit([&](handler &cgh) {
+        auto B = SubBuffer.get_access<access::mode::read_write>(cgh);
+        cgh.parallel_for<class Subbuf_test>(SubRange,
+                                            [=](id<2> Index) { B[Index] = 1; });
+      });
+    }
+
+    // Check that we filled correct subset of buffer:
+    // 0000000     0000000
+    // 0000000     0111110
+    // 0000000 --> 0111110
+    // 0000000     0111110
+    // 0000000     0111110
+    // 0000000     0000000
+
+    for (size_t i = 0; i < M; ++i) {
+      for (size_t j = 0; j < N; ++j) {
+        size_t Expected =
+            ((i == 0) || (i == M - 1) || (j == 0) || (j == N - 1)) ? 0 : 1;
+        if (Result[i][j] != Expected) {
+          std::cout << "line: " << __LINE__ << " Result[" << i << "][" << j
+                    << "] is " << Result[i][j] << " expected " << Expected
+                    << std::endl;
+          Failed = true;
+        }
+      }
+    }
+  }
+  // Try to create subbuffer from subbuffer
+  {
+    const int M = 10;
+    int Data[M] = {0};
+    auto OrigRange = range<1>(M);
+    buffer<int, 1> Buffer(Data, OrigRange);
+    auto Offset = id<1>(1);
+    auto SubRange = range<1>(M - 2);
+    auto SubSubRange = range<1>(M - 4);
+    queue MyQueue;
+    buffer<int, 1> SubBuffer(Buffer, Offset, SubRange);
+    buffer<int, 1> SubSubBuffer(SubBuffer, Offset, SubSubRange);
+    MyQueue.submit([&](handler &cgh) {
+      auto B = SubSubBuffer.get_access<access::mode::read_write>(cgh);
+      cgh.parallel_for<class Subsubbuf_test>(SubSubRange,
+                                          [=](id<1> Index) { B[Index] = 1; });
+    });
+    auto Acc = Buffer.get_access<cl::sycl::access::mode::read>();
+    for (size_t i = 0; i < M; ++i) {
+      size_t Expected = (i > 1 && i < M - 2) ? 1 : 0;
+      if (Acc[i] != Expected) {
+        std::cout << "line: " << __LINE__ << " Data[" << i << "] is " << Acc[i]
+                  << " expected " << Expected << std::endl;
+        Failed = true;
+      }
+    }
+  }
+  return Failed;
+}