diff --git a/sycl/include/CL/sycl/accessor.hpp b/sycl/include/CL/sycl/accessor.hpp
index 4584a8c3c955c..9bdc641371b14 100644
--- a/sycl/include/CL/sycl/accessor.hpp
+++ b/sycl/include/CL/sycl/accessor.hpp
@@ -150,9 +150,9 @@ SYCL_ACCESSOR_IMPL(!isTargetHostAccess(accessTarget) &&
   // reinterpret casting while setting kernel arguments in order to get cl_mem
   // value from the buffer regardless of the accessor's dimensionality.
 #ifndef __SYCL_DEVICE_ONLY__
-  detail::buffer_impl<buffer_allocator<char>> *m_Buf = nullptr;
+  detail::buffer_impl<buffer_allocator> *m_Buf = nullptr;
 #else
-  char padding[sizeof(detail::buffer_impl<buffer_allocator<char>> *)];
+  char padding[sizeof(detail::buffer_impl<buffer_allocator> *)];
 #endif // __SYCL_DEVICE_ONLY__
 
   dataT *Data;
@@ -185,9 +185,9 @@ SYCL_ACCESSOR_IMPL(!isTargetHostAccess(accessTarget) &&
   // reinterpret casting while setting kernel arguments in order to get cl_mem
   // value from the buffer regardless of the accessor's dimensionality.
 #ifndef __SYCL_DEVICE_ONLY__
-  detail::buffer_impl<buffer_allocator<char>> *m_Buf = nullptr;
+  detail::buffer_impl<buffer_allocator> *m_Buf = nullptr;
 #else
-  char padding[sizeof(detail::buffer_impl<buffer_allocator<char>> *)];
+  char padding[sizeof(detail::buffer_impl<buffer_allocator> *)];
 #endif // __SYCL_DEVICE_ONLY__
 
   dataT *Data;
@@ -763,7 +763,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,
+      : __impl((dataT *)detail::getSyclObjImpl(bufferRef)->BufPtr, Range,
                bufferRef.get_range(), Offset) {
     auto BufImpl = detail::getSyclObjImpl(bufferRef);
     if (AccessTarget == access::target::host_buffer) {
diff --git a/sycl/include/CL/sycl/buffer.hpp b/sycl/include/CL/sycl/buffer.hpp
index eac3d1fee317e..9a43504d5de72 100644
--- a/sycl/include/CL/sycl/buffer.hpp
+++ b/sycl/include/CL/sycl/buffer.hpp
@@ -21,7 +21,7 @@ class queue;
 template <int dimentions> class range;
 
 template <typename T, int dimensions = 1,
-          typename AllocatorT = cl::sycl::buffer_allocator<char>>
+          typename AllocatorT = cl::sycl::buffer_allocator>
 class buffer {
 public:
   using value_type = T;
@@ -36,11 +36,11 @@ class buffer {
         get_count() * sizeof(T), propList);
   }
 
-  // buffer(const range<dimensions> &bufferRange, AllocatorT allocator,
-  // const property_list &propList = {}) {
-  //     impl = std::make_shared<detail::buffer_impl>(bufferRange, allocator,
-  //     propList);
-  // }
+  buffer(const range<dimensions> &bufferRange, AllocatorT allocator,
+         const property_list &propList = {}) {
+    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 = {})
@@ -49,11 +49,11 @@ class buffer {
         hostData, get_count() * sizeof(T), propList);
   }
 
-  // buffer(T *hostData, const range<dimensions> &bufferRange,
-  // AllocatorT allocator, const property_list &propList = {}) {
-  //     impl = std::make_shared<detail::buffer_impl>(hostData, bufferRange,
-  //     allocator, propList);
-  // }
+  buffer(T *hostData, const range<dimensions> &bufferRange,
+         AllocatorT allocator, const property_list &propList = {}) {
+    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 = {})
@@ -62,18 +62,18 @@ class buffer {
         hostData, get_count() * sizeof(T), propList);
   }
 
-  // buffer(const T *hostData, const range<dimensions> &bufferRange,
-  // AllocatorT allocator, const property_list &propList = {}) {
-  //     impl = std::make_shared<detail::buffer_impl>(hostData, bufferRange,
-  //     allocator, propList);
-  // }
+  buffer(const T *hostData, const range<dimensions> &bufferRange,
+         AllocatorT allocator, const property_list &propList = {}) {
+    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 = {}) {
-  //     impl = std::make_shared<detail::buffer_impl>(hostData, bufferRange,
-  //     allocator, propList);
-  // }
+  buffer(const shared_ptr_class<T> &hostData,
+         const range<dimensions> &bufferRange, AllocatorT allocator,
+         const property_list &propList = {}) {
+    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,
@@ -83,12 +83,13 @@ class buffer {
         hostData, get_count() * sizeof(T), propList);
   }
 
-  // template <class InputIterator>
-  // buffer<T, 1>(InputIterator first, InputIterator last, AllocatorT allocator,
-  // const property_list &propList = {}) {
-  //     impl = std::make_shared<detail::buffer_impl>(first, last, allocator,
-  //     propList);
-  // }
+  template <class InputIterator>
+  buffer(InputIterator first, InputIterator last, AllocatorT allocator,
+         const property_list &propList = {})
+      : Range(range<1>(std::distance(first, last))) {
+    impl = std::make_shared<detail::buffer_impl<AllocatorT>>(
+        first, last, get_count() * sizeof(T), propList, allocator);
+  }
 
   template <class InputIterator, int N = dimensions,
             typename = std::enable_if<N == 1>>
@@ -135,7 +136,7 @@ class buffer {
 
   size_t get_size() const { return impl->get_size(); }
 
-  // AllocatorT get_allocator() const { return impl->get_allocator(); }
+  AllocatorT get_allocator() const { return impl->get_allocator(); }
 
   template <access::mode mode,
             access::target target = access::target::global_buffer>
@@ -152,28 +153,29 @@ class buffer {
     return impl->template get_access<T, dimensions, mode>(*this);
   }
 
-  // template <access::mode mode, access::target target =
-  // access::target::global_buffer> accessor<T, dimensions, mode, target,
-  // access::placeholder::false_t> get_access( handler &commandGroupHandler,
-  // range<dimensions> accessRange, id<dimensions> accessOffset = {}) {
-  //     return impl->get_access(commandGroupHandler, accessRange,
-  //     accessOffset);
-  // }
+  template <access::mode mode,
+            access::target target = access::target::global_buffer>
+  accessor<T, dimensions, mode, target, access::placeholder::false_t>
+  get_access(handler &commandGroupHandler, range<dimensions> accessRange,
+             id<dimensions> accessOffset = {}) {
+    return impl->template get_access<T, dimensions, mode, target>(
+        *this, commandGroupHandler, accessRange, accessOffset);
+  }
 
-  // template <access::mode mode>
-  // accessor<T, dimensions, mode, access::target::host_buffer,
-  // access::placeholder::false_t> get_access( range<dimensions> accessRange,
-  // id<dimensions> accessOffset = {}) {
-  //     return impl->get_access(accessRange, accessOffset);
-  // }
+  template <access::mode mode>
+  accessor<T, dimensions, mode, access::target::host_buffer,
+           access::placeholder::false_t>
+  get_access(range<dimensions> accessRange, id<dimensions> accessOffset = {}) {
+    return impl->template get_access<T, dimensions, mode>(*this, accessRange,
+                                                          accessOffset);
+  }
 
   template <typename Destination = std::nullptr_t>
   void set_final_data(Destination finalData = nullptr) {
     impl->set_final_data(finalData);
   }
 
-  // void set_write_back(bool flag = true) { return impl->set_write_back(flag);
-  // }
+  void set_write_back(bool flag = true) { return impl->set_write_back(flag); }
 
   // bool is_sub_buffer() const { return impl->is_sub_buffer(); }
 
@@ -189,6 +191,14 @@ class buffer {
                                                             reinterpretRange);
   }
 
+  template <typename propertyT> bool has_property() const {
+    return impl->template has_property<propertyT>();
+  }
+
+  template <typename propertyT> propertyT get_property() const {
+    return impl->template get_property<propertyT>();
+  }
+
 private:
   shared_ptr_class<detail::buffer_impl<AllocatorT>> impl;
   template <class Obj>
diff --git a/sycl/include/CL/sycl/detail/buffer_impl.hpp b/sycl/include/CL/sycl/detail/buffer_impl.hpp
index 4921c203f533f..321bacb2f5aae 100644
--- a/sycl/include/CL/sycl/detail/buffer_impl.hpp
+++ b/sycl/include/CL/sycl/detail/buffer_impl.hpp
@@ -37,16 +37,18 @@ class handler;
 class queue;
 template <int dimentions> class id;
 template <int dimentions> class range;
-template <class T> using buffer_allocator = std::allocator<T>;
+using buffer_allocator = std::allocator<char>;
 namespace detail {
 template <typename AllocatorT> class buffer_impl {
 public:
-  buffer_impl(const size_t sizeInBytes, const property_list &propList)
-      : buffer_impl((void *)nullptr, sizeInBytes, propList) {}
+  buffer_impl(const size_t sizeInBytes, const property_list &propList,
+              AllocatorT allocator = AllocatorT())
+      : buffer_impl((void *)nullptr, sizeInBytes, propList, allocator) {}
 
   buffer_impl(void *hostData, const size_t sizeInBytes,
-              const property_list &propList)
-      : SizeInBytes(sizeInBytes), Props(propList) {
+              const property_list &propList,
+              AllocatorT allocator = AllocatorT())
+      : SizeInBytes(sizeInBytes), Props(propList), MAllocator(allocator) {
     if (Props.has_property<property::buffer::use_host_ptr>()) {
       BufPtr = hostData;
     } else {
@@ -62,8 +64,9 @@ template <typename AllocatorT> class buffer_impl {
 
   // TODO temporary solution for allowing initialisation with const data
   buffer_impl(const void *hostData, const size_t sizeInBytes,
-              const property_list &propList)
-      : SizeInBytes(sizeInBytes), Props(propList) {
+              const property_list &propList,
+              AllocatorT allocator = AllocatorT())
+      : SizeInBytes(sizeInBytes), Props(propList), MAllocator(allocator) {
     if (Props.has_property<property::buffer::use_host_ptr>()) {
       // TODO make this buffer read only
       BufPtr = const_cast<void *>(hostData);
@@ -79,8 +82,9 @@ template <typename AllocatorT> class buffer_impl {
 
   template <typename T>
   buffer_impl(const shared_ptr_class<T> &hostData, const size_t sizeInBytes,
-              const property_list &propList)
-      : SizeInBytes(sizeInBytes), Props(propList) {
+              const property_list &propList,
+              AllocatorT allocator = AllocatorT())
+      : SizeInBytes(sizeInBytes), Props(propList), MAllocator(allocator) {
     if (Props.has_property<property::buffer::use_host_ptr>()) {
       BufPtr = hostData.get();
     } else {
@@ -97,8 +101,9 @@ template <typename AllocatorT> class buffer_impl {
 
   template <class InputIterator>
   buffer_impl(InputIterator first, InputIterator last, const size_t sizeInBytes,
-              const property_list &propList)
-      : SizeInBytes(sizeInBytes), Props(propList) {
+              const property_list &propList,
+              AllocatorT allocator = AllocatorT())
+      : SizeInBytes(sizeInBytes), Props(propList), MAllocator(allocator) {
     if (Props.has_property<property::buffer::use_host_ptr>()) {
       // TODO next line looks unsafe
       BufPtr = &*first;
@@ -140,7 +145,7 @@ template <typename AllocatorT> class buffer_impl {
           .copyBack<access::mode::read_write, access::target::host_buffer>(
               *this);
 
-    if (uploadData != nullptr) {
+    if (uploadData != nullptr && NeedWriteBack) {
       uploadData();
     }
 
@@ -170,7 +175,7 @@ template <typename AllocatorT> class buffer_impl {
       throw cl::sycl::runtime_error(
           "set_final_data could not be used with interoperability buffer");
     static_assert(!std::is_const<Destination>::value,
-                  "Сan not write in a constant Destination. Destination should "
+                  "Can not write in a constant Destination. Destination should "
                   "not be const.");
     uploadData = [this, final_data]() mutable {
       auto *Ptr =
@@ -182,6 +187,10 @@ template <typename AllocatorT> class buffer_impl {
     };
   }
 
+  void set_write_back(bool flag) { NeedWriteBack = flag; }
+
+  AllocatorT get_allocator() const { return MAllocator; }
+
   template <typename T, int dimensions, access::mode mode,
             access::target target = access::target::global_buffer>
   accessor<T, dimensions, mode, target, access::placeholder::false_t>
@@ -199,6 +208,34 @@ template <typename AllocatorT> class buffer_impl {
                     access::placeholder::false_t>(Buffer);
   }
 
+  template <typename T, int dimensions, access::mode mode,
+            access::target target = access::target::global_buffer>
+  accessor<T, dimensions, mode, target, access::placeholder::false_t>
+  get_access(buffer<T, dimensions, AllocatorT> &Buffer,
+             handler &commandGroupHandler, range<dimensions> accessRange,
+             id<dimensions> accessOffset) {
+    return accessor<T, dimensions, mode, target, access::placeholder::false_t>(
+        Buffer, commandGroupHandler, accessRange, accessOffset);
+  }
+
+  template <typename T, int dimensions, access::mode mode>
+  accessor<T, dimensions, mode, access::target::host_buffer,
+           access::placeholder::false_t>
+  get_access(buffer<T, dimensions, AllocatorT> &Buffer,
+             range<dimensions> accessRange, id<dimensions> accessOffset) {
+    return accessor<T, dimensions, mode, access::target::host_buffer,
+                    access::placeholder::false_t>(Buffer, accessRange,
+                                                  accessOffset);
+  }
+
+  template <typename propertyT> bool has_property() const {
+    return Props.has_property<propertyT>();
+  }
+
+  template <typename propertyT> propertyT get_property() const {
+    return Props.get_property<propertyT>();
+  }
+
 public:
   void moveMemoryTo(QueueImplPtr Queue, std::vector<cl::sycl::event> DepEvents,
                     EventImplPtr Event);
@@ -243,8 +280,10 @@ template <typename AllocatorT> class buffer_impl {
   // This field must be the first to guarantee that it's safe to use
   // reinterpret casting while setting kernel arguments in order to get cl_mem
   // value from the buffer regardless of its dimensionality.
+  AllocatorT MAllocator;
   OpenCLMemState OCLState;
   bool OpenCLInterop = false;
+  bool NeedWriteBack = true;
   event AvailableEvent;
   cl_context OpenCLContext = nullptr;
   void *BufPtr = nullptr;
diff --git a/sycl/include/CL/sycl/detail/scheduler/scheduler.cpp b/sycl/include/CL/sycl/detail/scheduler/scheduler.cpp
index 19c49a1c1e6e4..b724fac653cd8 100644
--- a/sycl/include/CL/sycl/detail/scheduler/scheduler.cpp
+++ b/sycl/include/CL/sycl/detail/scheduler/scheduler.cpp
@@ -61,8 +61,7 @@ template <typename dataT, int dimensions, access::mode accessMode,
 void Node::addAccRequirement(
     accessor<dataT, dimensions, accessMode, accessTarget, isPlaceholder> &&Acc,
     int argIndex) {
-  detail::buffer_impl<buffer_allocator<char>> *buf =
-      Acc.__get_impl()->m_Buf;
+  detail::buffer_impl<buffer_allocator> *buf = Acc.__get_impl()->m_Buf;
   addBufRequirement<accessMode, accessTarget>(*buf);
   addInteropArg(nullptr, buf->get_size(), argIndex,
                 getReqForBuffer(m_Bufs, *buf));
@@ -134,7 +133,7 @@ void Node::addExplicitMemOp(
   auto *DestBase = Dest.__get_impl();
   assert(DestBase != nullptr &&
          "Accessor should have an initialized accessor_base");
-  detail::buffer_impl<buffer_allocator<char>> *Buf = DestBase->m_Buf;
+  detail::buffer_impl<buffer_allocator> *Buf = DestBase->m_Buf;
 
   range<Dimensions> Range = DestBase->AccessRange;
   id<Dimensions> Offset = DestBase->Offset;
@@ -162,10 +161,10 @@ void Node::addExplicitMemOp(
   assert(DestBase != nullptr &&
          "Accessor should have an initialized accessor_base");
 
-  detail::buffer_impl<buffer_allocator<char>> *SrcBuf = SrcBase->m_Buf;
+  detail::buffer_impl<buffer_allocator> *SrcBuf = SrcBase->m_Buf;
   assert(SrcBuf != nullptr &&
          "Accessor should have an initialized buffer_impl");
-  detail::buffer_impl<buffer_allocator<char>> *DestBuf = DestBase->m_Buf;
+  detail::buffer_impl<buffer_allocator> *DestBuf = DestBase->m_Buf;
   assert(DestBuf != nullptr &&
          "Accessor should have an initialized buffer_impl");
 
@@ -195,7 +194,7 @@ void Scheduler::updateHost(
   auto *AccBase = Acc.__get_impl();
   assert(AccBase != nullptr &&
          "Accessor should have an initialized accessor_base");
-  detail::buffer_impl<buffer_allocator<char>> *Buf = AccBase->m_Buf;
+  detail::buffer_impl<buffer_allocator> *Buf = AccBase->m_Buf;
 
   updateHost<mode, tgt>(*Buf, Event);
 }
diff --git a/sycl/test/basic_tests/buffer/buffer.cpp b/sycl/test/basic_tests/buffer/buffer.cpp
index 02450b91abbfc..574d938fee1d5 100644
--- a/sycl/test/basic_tests/buffer/buffer.cpp
+++ b/sycl/test/basic_tests/buffer/buffer.cpp
@@ -19,6 +19,7 @@ using namespace cl::sycl;
 
 int main() {
   int data = 5;
+  bool failed = false;
   buffer<int, 1> buf(&data, range<1>(1));
   {
     int data1[10] = {-1, -1, -1, -1, -1, -1, -1, -1, -1, -1};
@@ -453,5 +454,55 @@ int main() {
     for (int i = 5; i < 10; i++)
       assert(data1[i] == -1);
   }
+
+  // Check that data is copied back after forcing write-back using
+  // set_write_back
+  {
+    std::vector<int> data1(10, -1);
+    {
+      buffer<int, 1> b(range<1>(10));
+      b.set_final_data(data1.data());
+      b.set_write_back(true);
+      queue myQueue;
+      myQueue.submit([&](handler &cgh) {
+        auto B = b.get_access<access::mode::read_write>(cgh);
+        cgh.parallel_for<class wb>(range<1>{10},
+                                       [=](id<1> index) { B[index] = 0; });
+      });
+
+    }
+    // Data is copied back because there is a user side ptr and write-back is
+    // enabled
+    for (int i = 0; i < 10; i++)
+      if (data1[i] != 0) {
+        assert(false);
+        failed = true;
+      }
+  }
+
+  // Check that data is not copied back after canceling write-back using
+  // set_write_back
+  {
+    std::vector<int> data1(10, -1);
+    {
+      buffer<int, 1> b(range<1>(10));
+      b.set_final_data(data1.data());
+      b.set_write_back(false);
+      queue myQueue;
+      myQueue.submit([&](handler &cgh) {
+        auto B = b.get_access<access::mode::read_write>(cgh);
+        cgh.parallel_for<class notwb>(range<1>{10},
+                                       [=](id<1> index) { B[index] = 0; });
+      });
+
+    }
+    // Data is not copied back because write-back is canceled
+    for (int i = 0; i < 10; i++)
+      if (data1[i] != -1) {
+        assert(false);
+        failed = true;
+      }
+  }
   // TODO tests with mutex property
+  return failed;
 }