diff --git a/clang/test/dpct/helper_files_ref/include/dpl_extras/memory.h b/clang/test/dpct/helper_files_ref/include/dpl_extras/memory.h
index 945795b9c659..b57a225dd268 100644
--- a/clang/test/dpct/helper_files_ref/include/dpl_extras/memory.h
+++ b/clang/test/dpct/helper_files_ref/include/dpl_extras/memory.h
@@ -60,30 +60,65 @@ template <typename T> struct device_reference {
   template <typename OtherT>
   device_reference(const device_reference<OtherT> &input)
       : value(input.value) {}
-  device_reference(const pointer &input) : value((*input).value) {}
-  device_reference(value_type &input) : value(input) {}
+  device_reference(const pointer &input) : value(input.get()) {}
+  device_reference(value_type *input) : value(input) {}
+  device_reference(value_type &input) : value(&input) {}
   template <typename OtherT>
   device_reference &operator=(const device_reference<OtherT> &input) {
-    value = input;
-    return *this;
-  };
-  device_reference &operator=(const device_reference &input) {
-    T val = input.value;
-    value = val;
+    __assign_from(input.__get_value());
     return *this;
   };
   device_reference &operator=(const value_type &x) {
-    value = x;
+    __assign_from(x);
     return *this;
   };
-  pointer operator&() const { return pointer(&value); };
-  operator value_type() const { return T(value); }
+  pointer operator&() const { return pointer(value); };
+
+  operator value_type() const { return __get_value(); }
+
+  __attribute__((__used__))
+  void PlusPlus_helper()
+  {
+    value_type *tmp = value;
+    sycl::queue default_queue = dpct::get_default_queue();
+    default_queue.submit([&](sycl::handler& h) {
+        h.single_task<dpct_kernel_name<class dv_PlusPlus, value_type>>([=] () {
+          (*tmp)++;
+        });
+      })
+      .wait();
+  }
+
   device_reference &operator++() {
-    ++value;
+#if __SYCL_DEVICE_ONLY__
+    (*value)++;
+#else
+    PlusPlus_helper();
+#endif
     return *this;
   };
+
+
+  __attribute__((__used__))
+  void MinusMinus_helper()
+  {
+    value_type *tmp = value;
+    sycl::queue default_queue = dpct::get_default_queue();
+    default_queue.submit([&](sycl::handler& h) {
+        h.single_task<dpct_kernel_name<class dv_MinusMinus, value_type>>([=] () {
+          (*tmp)--;
+        });
+      })
+      .wait();
+  }
+
+
   device_reference &operator--() {
-    --value;
+#if __SYCL_DEVICE_ONLY__
+    (*value)--;
+#else
+    MinusMinus_helper();
+#endif
     return *this;
   };
   device_reference operator++(int) {
@@ -96,52 +131,314 @@ template <typename T> struct device_reference {
     --(*this);
     return ref;
   };
-  device_reference &operator+=(const T &input) {
-    value += input;
-    return *this;
+
+  __attribute__((__used__))
+  void PlusEqual_helper(const value_type &input)
+  {
+    value_type *tmp = value;
+    sycl::queue default_queue = dpct::get_default_queue();
+    default_queue.submit([&](sycl::handler& h) {
+        h.single_task<dpct_kernel_name<class dv_PlusEqual, value_type>>([=] () {
+          (*tmp)+=input;
+        });
+      })
+      .wait();
+  }
+
+  device_reference &operator+=(const value_type &input) {
+#if __SYCL_DEVICE_ONLY__
+    (*value)+=input;
+#else
+    PlusEqual_helper(input);
+#endif
+   return *this;
   };
-  device_reference &operator-=(const T &input) {
-    value -= input;
+
+  __attribute__((__used__))
+  void MinusEqual_helper(const value_type &input)
+  {
+    value_type *tmp = value;
+    sycl::queue default_queue = dpct::get_default_queue();
+    default_queue.submit([&](sycl::handler& h) {
+        h.single_task<dpct_kernel_name<class dv_MinusEqual, value_type>>([=] () {
+          (*tmp)-=input;
+        });
+      })
+      .wait();
+  }
+
+  device_reference &operator-=(const value_type &input) {
+#if __SYCL_DEVICE_ONLY__
+    (*value)-=input;
+#else
+    MinusEqual_helper(input);
+#endif
     return *this;
   };
-  device_reference &operator*=(const T &input) {
-    value *= input;
+
+  __attribute__((__used__))
+  void TimesEqual_helper(const value_type &input)
+  {
+    value_type *tmp = value;
+    sycl::queue default_queue = dpct::get_default_queue();
+    default_queue.submit([&](sycl::handler& h) {
+        h.single_task<dpct_kernel_name<class dv_TimesEqual, value_type>>([=] () {
+          (*tmp) *= input;
+        });
+      })
+      .wait();
+  }
+
+  device_reference &operator*=(const value_type &input) {
+#if __SYCL_DEVICE_ONLY__
+    (*value) *= input;
+#else
+    TimesEqual_helper(input);
+#endif
     return *this;
-  };
-  device_reference &operator/=(const T &input) {
-    value /= input;
+  }
+
+  __attribute__((__used__))
+  void DivideEqual_helper(const value_type &input)
+  {
+    value_type *tmp = value;
+    sycl::queue default_queue = dpct::get_default_queue();
+    default_queue.submit([&](sycl::handler& h) {
+        h.single_task<dpct_kernel_name<class dv_DivideEqual, value_type>>([=] () {
+          (*tmp) /= input;
+        });
+      })
+      .wait();
+  }
+
+  device_reference &operator/=(const value_type &input) {
+#if __SYCL_DEVICE_ONLY__
+    (*value)/=input;
+#else
+    DivideEqual_helper(input);
+#endif
     return *this;
-  };
-  device_reference &operator%=(const T &input) {
-    value %= input;
+  }
+
+  __attribute__((__used__))
+  void ModEqual_helper(const value_type &input)
+  {
+    value_type *tmp = value;
+    sycl::queue default_queue = dpct::get_default_queue();
+    default_queue.submit([&](sycl::handler& h) {
+        h.single_task<dpct_kernel_name<class dv_ModEqual, value_type>>([=] () {
+          (*tmp) %= input;
+        });
+      })
+      .wait();
+  }
+
+  device_reference &operator%=(const value_type &input) {
+#if __SYCL_DEVICE_ONLY__
+    (*value) %= input;
+#else
+    ModEqual_helper(input);
+#endif
     return *this;
-  };
-  device_reference &operator&=(const T &input) {
-    value &= input;
+  }
+  
+  __attribute__((__used__))
+  void AndEqual_helper(const value_type &input)
+  {
+    value_type *tmp = value;
+
+    sycl::queue default_queue = dpct::get_default_queue();
+    default_queue.submit([&](sycl::handler& h) {
+        h.single_task<dpct_kernel_name<class dv_AndEqual, value_type>>([=] () {
+          (*tmp) &= input;
+        });
+      })
+      .wait();
+  }
+
+  device_reference &operator&=(const value_type &input) {
+#if __SYCL_DEVICE_ONLY__
+    (*value) &= input;
+#else
+    AndEqual_helper(input);
+#endif
     return *this;
-  };
-  device_reference &operator|=(const T &input) {
-    value |= input;
+  }
+
+  __attribute__((__used__))
+  void OrEqual_helper(const value_type &input)
+  {
+    value_type *tmp = value;
+
+    sycl::queue default_queue = dpct::get_default_queue();
+    default_queue.submit([&](sycl::handler& h) {
+        h.single_task<dpct_kernel_name<class dv_OrEqual, value_type>>([=] () {
+          (*tmp)|=input;
+        });
+      })
+      .wait();
+  }
+
+  device_reference &operator|=(const value_type &input) {
+#if __SYCL_DEVICE_ONLY__
+    (*value)|=input;
+#else
+    OrEqual_helper(input);
+#endif
     return *this;
   };
-  device_reference &operator^=(const T &input) {
-    value ^= input;
+
+  __attribute__((__used__))
+  void CrossEqual_helper(const value_type &input)
+  {
+    value_type *tmp = value;
+    sycl::queue default_queue = dpct::get_default_queue();
+    default_queue.submit([&](sycl::handler& h) {
+        h.single_task<dpct_kernel_name<class dv_CrossEqual, value_type>>([=] () {
+          (*tmp)^=input;
+        });
+      })
+      .wait();
+  }
+  device_reference &operator^=(const value_type &input) {
+#if __SYCL_DEVICE_ONLY__
+    (*value)^=input;
+#else
+    CrossEqual_helper(input);
+#endif
     return *this;
   };
-  device_reference &operator<<=(const T &input) {
-    value <<= input;
+
+  __attribute__((__used__))
+  void ShiftLeftEqual_helper(const int &input)
+  {
+    value_type *tmp = value;
+
+    sycl::queue default_queue = dpct::get_default_queue();
+    default_queue.submit([&](sycl::handler& h) {
+        h.single_task<dpct_kernel_name<class dv_ShiftLeftEqual, value_type>>([=] () {
+          (*tmp)<<=input;
+        });
+      })
+      .wait();
+  }
+  device_reference &operator<<=(const value_type &input) {
+#if __SYCL_DEVICE_ONLY__
+    (*value)<<=input;
+#else
+    ShiftLeftEqual_helper(input);
+#endif
     return *this;
   };
-  device_reference &operator>>=(const T &input) {
-    value >>= input;
+
+  __attribute__((__used__))
+  void ShiftRightEqual_helper(const int &input)
+  {
+    value_type *tmp = value;
+
+    sycl::queue default_queue = dpct::get_default_queue();
+    default_queue.submit([&](sycl::handler& h) {
+        h.single_task<dpct_kernel_name<class dv_ShiftRightEqual, value_type>>([=] () {
+          (*tmp)>>=input;
+        });
+      })
+      .wait();
+  }
+
+  device_reference &operator>>=(const value_type &input) {
+#if __SYCL_DEVICE_ONLY__
+    (*value)>>=input;
+#else
+    ShiftRightEqual_helper(input);
+#endif
     return *this;
   };
+
+  __attribute__((__used__))
+  void
+  swap_helper(device_reference &input)
+  {
+    value_type *my_val = value;
+    value_type *input_val = input.value;
+
+    sycl::queue default_queue = dpct::get_default_queue();
+    default_queue.submit([&](sycl::handler& h) {
+        h.single_task<dpct_kernel_name<class dv_swap, value_type>>([=]() {
+          T tmp = *my_val;
+          *my_val = *(input_val);
+          *(input_val) = tmp;
+        });
+    }).wait();
+  }
+
   void swap(device_reference &input) {
-    T tmp = (*this);
-    *this = (input);
-    input = (tmp);
+#ifdef __SYCL_DEVICE_ONLY__
+    T tmp = *value;
+    *value = *(input.value);
+    *(input.value) = tmp;
+#else
+    swap_helper(input);
+#endif
   }
-  T &value;
+
+  __attribute__((__used__))
+  void operator_equal_helper(const device_reference &input)
+  {
+    value_type *tmp = value;
+    value_type *input_val = input.value;
+    sycl::queue default_queue = dpct::get_default_queue();
+    default_queue.submit([&](sycl::handler& h) {
+        h.single_task<dpct_kernel_name<class dv_equal_op, value_type>>([=]() {
+          *tmp = *(input_val);
+        });
+    }).wait();
+  }
+
+  device_reference &operator=(const device_reference &input) {
+#ifdef __SYCL_DEVICE_ONLY__
+    *value = *(input.value);
+#else
+    operator_equal_helper(input);
+#endif
+    return *this;
+  };
+
+  value_type *value;
+private: 
+
+  __attribute__((__used__))
+  void assign_from_helper(const value_type& from)
+  {
+    dpct::get_default_queue().fill(value, from, 1);
+  }
+
+  device_reference &__assign_from(const value_type& from)
+  {
+#ifdef __SYCL_DEVICE_ONLY__ //call from the device
+    *value = from;
+#else
+    assign_from_helper(from);
+#endif
+    return *this;
+  }
+
+  __attribute__((__used__))
+  value_type get_value_helper() const
+  {
+      value_type tmp;
+      sycl::queue default_queue = dpct::get_default_queue();
+      default_queue.memcpy(&tmp, value, sizeof(value_type)*1).wait();
+      return value_type(tmp);
+  }
+  value_type __get_value() const
+  {
+#ifdef __SYCL_DEVICE_ONLY__ //call from the device
+    return *value;
+#else // call from the host
+    return get_value_helper();
+#endif
+  }
+
 };
 
 template <typename T>
@@ -365,9 +662,6 @@ template <typename ValueType, typename Derived> class device_pointer_base {
   operator ValueType *() { return ptr; }
   operator ValueType *() const { return ptr; }
 
-  ValueType &operator[](difference_type idx) { return ptr[idx]; }
-  ValueType &operator[](difference_type idx) const { return ptr[idx]; }
-
   Derived operator+(difference_type forward) const {
     return Derived{ptr + forward};
   }
@@ -396,7 +690,7 @@ class device_pointer : public device_pointer_base<T, device_pointer<T>> {
   using value_type = T;
   using difference_type = std::make_signed<std::size_t>::type;
   using pointer = T *;
-  using reference = T &;
+  using reference = device_reference<value_type>;
   using const_reference = const T &;
   using iterator_category = std::random_access_iterator_tag;
   using is_hetero = std::false_type;         // required
@@ -411,6 +705,9 @@ class device_pointer : public device_pointer_base<T, device_pointer<T>> {
     return *this;
   }
 
+  reference operator[](difference_type idx) { return reference(this->ptr + idx); }
+  reference operator[](difference_type idx) const { return reference(this->ptr + idx); }
+
   // include operators from base class
   using base_type::operator++;
   using base_type::operator--;
@@ -442,7 +739,7 @@ class device_pointer<void>
   using value_type = dpct::byte_t;
   using difference_type = std::make_signed<std::size_t>::type;
   using pointer = void *;
-  using reference = value_type &;
+  using reference = device_reference<value_type>;
   using const_reference = const value_type &;
   using iterator_category = std::random_access_iterator_tag;
   using is_hetero = std::false_type;         // required
@@ -456,6 +753,9 @@ class device_pointer<void>
   operator void *() { return this->ptr; }
   operator void *() const { return this->ptr; }
 
+  reference operator[](difference_type idx) { return reference(this->ptr + idx); }
+  reference operator[](difference_type idx) const { return reference(this->ptr + idx); }
+
   // include operators from base class
   using base_type::operator++;
   using base_type::operator--;
@@ -476,6 +776,8 @@ class device_pointer<void>
     return *this;
   }
 };
+
+
 #endif
 
 #ifdef DPCT_USM_LEVEL_NONE
@@ -600,10 +902,10 @@ template <typename T> class device_iterator : public device_pointer<T> {
     return *this;
   }
 
-  reference operator*() const { return *(Base::ptr + idx); }
+  reference operator*() const { return reference(Base::ptr + idx); }
 
-  reference operator[](difference_type i) { return Base::ptr[idx + i]; }
-  reference operator[](difference_type i) const { return Base::ptr[idx + i]; }
+  reference operator[](difference_type i) { return reference(Base::ptr + idx + i); }
+  reference operator[](difference_type i) const { return reference(Base::ptr + idx + i); }
   device_iterator &operator++() {
     ++idx;
     return *this;
@@ -738,4 +1040,10 @@ template <typename T> T &get_raw_reference(T &ref) {
 
 } // namespace dpct
 
+template<typename T>
+struct ::sycl::is_device_copyable<dpct::device_pointer<T>> : std::true_type {};
+template<typename T>
+struct ::sycl::is_device_copyable<dpct::device_iterator<T>> : std::true_type {};
+
+
 #endif
diff --git a/clang/test/dpct/helper_files_ref/include/dpl_extras/vector.h b/clang/test/dpct/helper_files_ref/include/dpl_extras/vector.h
index 4480a060774e..53aa0ce1c123 100644
--- a/clang/test/dpct/helper_files_ref/include/dpl_extras/vector.h
+++ b/clang/test/dpct/helper_files_ref/include/dpl_extras/vector.h
@@ -25,6 +25,191 @@
 namespace dpct {
 
 namespace internal {
+
+// usm_device_allocator is provided here specifically for dpct::device_vector.
+//  Warning: It may be dangerous to use usm_device_allocator in other settings,
+//  because containers may use the supplied allocator to allocate side
+//  information which needs to be available on the host.  Data allocated with
+//  this allocator is by definition not available on the host, and would result
+//  in an error if accessed from the host without proper handling.
+template <typename T, size_t Alignment = 0> class usm_device_allocator {
+public:
+  using value_type = T;
+  using propagate_on_container_copy_assignment = std::true_type;
+  using propagate_on_container_move_assignment = std::true_type;
+  using propagate_on_container_swap = std::true_type;
+
+public:
+  template <typename U> struct rebind {
+    typedef usm_device_allocator<U, Alignment> other;
+  };
+
+  usm_device_allocator() = delete;
+  usm_device_allocator(const sycl::context &Ctxt, const sycl::device &Dev,
+                       const sycl::property_list &PropList = {})
+      : MContext(Ctxt), MDevice(Dev), MPropList(PropList) {}
+  usm_device_allocator(const sycl::queue &Q,
+                       const sycl::property_list &PropList = {})
+      : MContext(Q.get_context()), MDevice(Q.get_device()),
+        MPropList(PropList) {}
+  usm_device_allocator(const usm_device_allocator &) = default;
+  usm_device_allocator(usm_device_allocator &&) noexcept = default;
+  usm_device_allocator &operator=(const usm_device_allocator &Other) {
+    MContext = Other.MContext;
+    MDevice = Other.MDevice;
+    MPropList = Other.MPropList;
+    return *this;
+  }
+  usm_device_allocator &operator=(usm_device_allocator &&Other) {
+    MContext = std::move(Other.MContext);
+    MDevice = std::move(Other.MDevice);
+    MPropList = std::move(Other.MPropList);
+    return *this;
+  }
+
+  template <class U>
+  usm_device_allocator(const usm_device_allocator<U, Alignment> &Other) noexcept
+      : MContext(Other.MContext), MDevice(Other.MDevice),
+        MPropList(Other.MPropList) {}
+
+  /// Allocates memory.
+  ///
+  /// \param NumberOfElements is a count of elements to allocate memory for.
+  T *allocate(size_t NumberOfElements,
+              const sycl::detail::code_location CodeLoc =
+                  sycl::detail::code_location::current()) {
+
+    auto Result = reinterpret_cast<T *>(aligned_alloc(
+        getAlignment(), NumberOfElements * sizeof(value_type), MDevice,
+        MContext, sycl::usm::alloc::device, MPropList, CodeLoc));
+    if (!Result) {
+      throw sycl::exception(sycl::errc::memory_allocation);
+    }
+    return Result;
+  }
+
+  /// Deallocates memory.
+  ///
+  /// \param Ptr is a pointer to memory being deallocated.
+  /// \param Size is a number of elements previously passed to allocate.
+  void deallocate(T *Ptr, size_t,
+                  const sycl::detail::code_location CodeLoc =
+                      sycl::detail::code_location::current()) {
+    if (Ptr) {
+      free(Ptr, MContext, CodeLoc);
+    }
+  }
+
+  template <class U, size_t AlignmentU>
+  friend bool operator==(const usm_device_allocator<T, Alignment> &One,
+                         const usm_device_allocator<U, AlignmentU> &Two) {
+    return ((One.MContext == Two.MContext) && (One.MDevice == Two.MDevice));
+  }
+
+  template <class U, size_t AlignmentU>
+  friend bool operator!=(const usm_device_allocator<T, Alignment> &One,
+                         const usm_device_allocator<U, AlignmentU> &Two) {
+    return !((One.MContext == Two.MContext) && (One.MDevice == Two.MDevice));
+  }
+
+  template <typename Property> bool has_property() const noexcept {
+    return MPropList.has_property<Property>();
+  }
+
+  template <typename Property> Property get_property() const {
+    return MPropList.get_property<Property>();
+  }
+
+private:
+  constexpr size_t getAlignment() const {
+    return sycl::max(alignof(T), Alignment);
+  }
+
+  template <class U, size_t AlignmentU> friend class usm_device_allocator;
+
+  sycl::context MContext;
+  sycl::device MDevice;
+  sycl::property_list MPropList;
+};
+
+template <class, class _Alloc, class... _Args>
+struct __has_construct_impl : ::std::false_type {};
+
+template <class _Alloc, class... _Args>
+struct __has_construct_impl<decltype((void)std::declval<_Alloc>().construct(
+                                std::declval<_Args>()...)),
+                            _Alloc, _Args...> : ::std::true_type {};
+
+// check if the provided allocator has a construct() member function
+template <class _Alloc, class... _Args>
+struct __has_construct : __has_construct_impl<void, _Alloc, _Args...> {};
+
+// check if the provided allocator has a destroy() member function
+template <class _Alloc, class _Pointer, class = void>
+struct __has_destroy : ::std::false_type {};
+
+template <class _Alloc, class _Pointer>
+struct __has_destroy<_Alloc, _Pointer,
+                     decltype((void)std::declval<_Alloc>().destroy(
+                         std::declval<_Pointer>()))> : ::std::true_type {};
+
+// device_allocator_traits is a device-friendly subset of the functionality of
+//  std::allocator_traits which uses static construct and destroy functions
+//  and is usable inside of sycl kernels without passing the allocator to the
+//  kernel.
+template <typename _Allocator> struct device_allocator_traits {
+
+  // apply default constructor if no override is provided
+  template <typename DataT>
+  static
+      typename ::std::enable_if_t<!__has_construct<_Allocator, DataT *>::value,
+                                  void>
+      construct(DataT *p) {
+    ::new ((void *)p) DataT();
+  }
+
+  // use provided default construct call if it exists
+  template <typename DataT>
+  static
+      typename ::std::enable_if_t<__has_construct<_Allocator, DataT *>::value,
+                                  void>
+      construct(DataT *p) {
+    _Allocator::construct(p);
+  }
+
+  // apply constructor if no override is provided
+  template <typename DataT, typename T_in>
+  static typename ::std::enable_if_t<
+      !__has_construct<_Allocator, DataT *, T_in>::value, void>
+  construct(DataT *p, T_in arg) {
+    ::new ((void *)p) DataT(arg);
+  }
+
+  // use provided construct call if it exists
+  template <typename DataT, typename T_in>
+  static typename ::std::enable_if_t<
+      __has_construct<_Allocator, DataT *, T_in>::value, void>
+  construct(DataT *p, T_in arg) {
+    _Allocator::construct(p, arg);
+  }
+
+  // apply default destructor if no destroy override is provided
+  template <typename DataT>
+  static typename ::std::enable_if_t<!__has_destroy<_Allocator, DataT *>::value,
+                                     void>
+  destroy(DataT *p) {
+    p->~DataT();
+  }
+
+  // use provided destroy call if it exists
+  template <typename DataT>
+  static typename ::std::enable_if_t<__has_destroy<_Allocator, DataT *>::value,
+                                     void>
+  destroy(DataT *p) {
+    _Allocator::destroy(p);
+  }
+};
+
 template <typename Iter, typename Void = void> // for non-iterators
 struct is_iterator : std::false_type {};
 
@@ -42,7 +227,7 @@ struct is_iterator<T *> : std::true_type {};
 #ifndef DPCT_USM_LEVEL_NONE
 
 template <typename T,
-          typename Allocator = sycl::usm_allocator<T, sycl::usm::alloc::shared>>
+          typename Allocator = dpct::internal::usm_device_allocator<T>>
 class device_vector {
 public:
   using iterator = device_iterator<T>;
@@ -55,6 +240,7 @@ class device_vector {
   using difference_type =
       typename ::std::iterator_traits<iterator>::difference_type;
   using size_type = ::std::size_t;
+  using allocator_type = Allocator;
 
 private:
   Allocator _alloc;
@@ -66,143 +252,164 @@ class device_vector {
 
   void _set_capacity_and_alloc() {
     _capacity = ::std::max(_size * 2, _min_capacity());
-    _storage = _alloc.allocate(_capacity);
+    _storage = ::std::allocator_traits<Allocator>::allocate(_alloc, _capacity);
+  }
+
+  void _construct(size_type n, size_type start_idx = 0) {
+    if (n > 0) {
+      pointer p = _storage;
+      get_default_queue()
+          .submit([&](sycl::handler &cgh) {
+            cgh.parallel_for(n, [=](sycl::id<1> i) {
+              ::dpct::internal::device_allocator_traits<Allocator>::construct(
+                  p + start_idx + i);
+            });
+          })
+          .wait();
+    }
+  }
+
+  void _construct(size_type n, const T &value, size_type start_idx = 0) {
+    if (n > 0) {
+      pointer p = _storage;
+      get_default_queue()
+          .submit([&](sycl::handler &cgh) {
+            cgh.parallel_for(n, [=](sycl::id<1> i) {
+              ::dpct::internal::device_allocator_traits<Allocator>::construct(
+                  p + start_idx + i, value);
+            });
+          })
+          .wait();
+    }
+  }
+
+  template <typename Iter>
+  void _construct(Iter first, Iter last, size_type start_idx = 0) {
+    int num_eles = ::std::distance(first, last);
+    if (num_eles > 0) {
+      // this should properly handle host or device input data
+      auto read_input =
+          oneapi::dpl::__ranges::__get_sycl_range<sycl::access_mode::read,
+                                                  Iter>();
+      auto input_rng = read_input(first, last).all_view();
+      pointer p = _storage;
+      get_default_queue()
+          .submit([&](sycl::handler &cgh) {
+            oneapi::dpl::__ranges::__require_access(cgh, input_rng);
+            cgh.parallel_for(num_eles, [=](sycl::id<1> i) {
+              ::dpct::internal::device_allocator_traits<Allocator>::construct(
+                  p + start_idx + i, input_rng[i]);
+            });
+          })
+          .wait();
+    }
+  }
+
+  void _destroy(size_type n, size_type start_idx = 0) {
+    // only call destroy kernel *only* if custom destroy function is provided to
+    //  prevent extra unnecessary kernel call
+    if constexpr (::dpct::internal::__has_destroy<Allocator, pointer>::value) {
+      if (n > 0) {
+        pointer p = _storage;
+        get_default_queue()
+            .submit([&](sycl::handler &cgh) {
+              cgh.parallel_for(n, [=](sycl::id<1> i) {
+                ::dpct::internal::device_allocator_traits<Allocator>::destroy(
+                    p + start_idx + i);
+              });
+            })
+            .wait();
+      }
+    }
   }
 
 public:
   template <typename OtherA> operator ::std::vector<T, OtherA>() const {
     auto __tmp = ::std::vector<T, OtherA>(this->size());
     ::std::copy(oneapi::dpl::execution::make_device_policy(get_default_queue()),
-              this->begin(), this->end(), __tmp.begin());
+                this->begin(), this->end(), __tmp.begin());
     return __tmp;
   }
   device_vector()
       : _alloc(get_default_queue()), _size(0), _capacity(_min_capacity()) {
     _set_capacity_and_alloc();
   }
-  ~device_vector() /*= default*/ { _alloc.deallocate(_storage, _capacity); };
-  explicit device_vector(size_type n) : device_vector(n, T()) {}
+  ~device_vector() /*= default*/ {
+    _destroy(size());
+    ::std::allocator_traits<Allocator>::deallocate(_alloc, _storage, _capacity);
+  }
+  explicit device_vector(size_type n) : _alloc(get_default_queue()), _size(n) {
+    _set_capacity_and_alloc();
+    _construct(n);
+  }
   explicit device_vector(size_type n, const T &value)
       : _alloc(get_default_queue()), _size(n) {
     _set_capacity_and_alloc();
-    if (_size > 0) {
-      ::std::fill(oneapi::dpl::execution::make_device_policy(get_default_queue()),
-                begin(), end(), T(value));
-    }
-  }
-  device_vector(const device_vector &other) : _alloc(get_default_queue()) {
-    _size = other.size();
-    _capacity = other.capacity();
-    _storage = _alloc.allocate(_capacity);
-    ::std::copy(oneapi::dpl::execution::make_device_policy(get_default_queue()),
-              other.begin(), other.end(), begin());
+    _construct(n, value);
   }
+
   device_vector(device_vector &&other)
-      : _alloc(get_default_queue()), _size(other.size()),
+      : _alloc(std::move(other._alloc)), _size(other.size()),
         _capacity(other.capacity()), _storage(other._storage) {
     other._size = 0;
-    other._capacity = 0; 
+    other._capacity = 0;
     other._storage = nullptr;
   }
 
   template <typename InputIterator>
-  device_vector(InputIterator first,
-                typename ::std::enable_if<
-                    internal::is_iterator<InputIterator>::value &&
-                        !::std::is_pointer<InputIterator>::value &&
-                        ::std::is_same<typename ::std::iterator_traits<
-                                         InputIterator>::iterator_category,
-                                     ::std::random_access_iterator_tag>::value,
-                    InputIterator>::type last)
-      : _alloc(get_default_queue()) {
-    _size = ::std::distance(first, last);
-    _set_capacity_and_alloc();
-    if (_size > 0) {
-      ::std::copy(oneapi::dpl::execution::make_device_policy(get_default_queue()),
-                first, last, begin());
-    }
-  }
-
-  template <typename InputIterator>
-  device_vector(InputIterator first,
-                typename ::std::enable_if<::std::is_pointer<InputIterator>::value,
-                                        InputIterator>::type last)
+  device_vector(InputIterator first, 
+  typename ::std::enable_if_t<dpct::internal::is_iterator<InputIterator>::value, InputIterator> last)
       : _alloc(get_default_queue()) {
     _size = ::std::distance(first, last);
     _set_capacity_and_alloc();
-    if (_size > 0) {
-      auto ptr_type = sycl::get_pointer_type(first, get_default_context());
-      if (ptr_type != sycl::usm::alloc::host &&
-          ptr_type != sycl::usm::alloc::unknown) {
-        ::std::copy(
-            oneapi::dpl::execution::make_device_policy(get_default_queue()),
-            first, last, begin());
-      } else {
-        sycl::buffer<T, 1> buf(first, last);
-        auto buf_first = oneapi::dpl::begin(buf);
-        auto buf_last = oneapi::dpl::end(buf);
-        ::std::copy(
-            oneapi::dpl::execution::make_device_policy(get_default_queue()),
-            buf_first, buf_last, begin());
-      }
-    }
+    _construct(first, last);
   }
 
-  template <typename InputIterator>
-  device_vector(InputIterator first,
-                typename ::std::enable_if<
-                    internal::is_iterator<InputIterator>::value &&
-                        !::std::is_pointer<InputIterator>::value &&
-                        !::std::is_same<typename ::std::iterator_traits<
-                                          InputIterator>::iterator_category,
-                                      ::std::random_access_iterator_tag>::value,
-                    InputIterator>::type last)
-      : _alloc(get_default_queue()), _size(::std::distance(first, last)) {
-    _set_capacity_and_alloc();
-    ::std::vector<T> _tmp(first, last);
-    if (_size > 0) {
-      ::std::copy(oneapi::dpl::execution::make_device_policy(get_default_queue()),
-                _tmp.begin(), _tmp.end(), this->begin());
-    }
+  device_vector(const device_vector &other)
+      : _alloc(std::allocator_traits<Allocator>::select_on_container_copy_construction(other._alloc)) {
+    _size = other.size();
+    _capacity = other.capacity();
+    _storage = ::std::allocator_traits<Allocator>::allocate(_alloc, _capacity);
+    _construct(other.begin(), other.end());
   }
 
   template <typename OtherAllocator>
-  device_vector(const device_vector<T, OtherAllocator> &v)
-      : _alloc(get_default_queue()), _storage(v.real_begin()), _size(v.size()),
-        _capacity(v.capacity()) {}
+  device_vector(const device_vector<T, OtherAllocator> &other)
+      : _alloc(get_default_queue()) {
+    _size = other.size();
+    _capacity = other.capacity();
+    _storage = ::std::allocator_traits<Allocator>::allocate(_alloc, _capacity);
+    _construct(other.begin(), other.end());
+  }
 
   template <typename OtherAllocator>
   device_vector(::std::vector<T, OtherAllocator> &v)
-      : _alloc(get_default_queue()), _size(v.size()) {
-    _set_capacity_and_alloc();
-    if (_size > 0) {
-      ::std::copy(oneapi::dpl::execution::make_device_policy(get_default_queue()),
-                v.begin(), v.end(), this->begin());
-    }
-  }
-
+      : device_vector(v.begin(), v.end()) {}
   template <typename OtherAllocator>
   device_vector &operator=(const ::std::vector<T, OtherAllocator> &v) {
     resize(v.size());
-    if (_size > 0) {
-      ::std::copy(oneapi::dpl::execution::make_device_policy(get_default_queue()),
-                v.begin(), v.end(), begin());
-    }
+    _construct(v.begin(), v.end());
     return *this;
   }
   device_vector &operator=(const device_vector &other) {
     // Copy assignment operator:
-    resize(other.size());
-    if (_size > 0) {
-      ::std::copy(oneapi::dpl::execution::make_device_policy(get_default_queue()),
-                other.begin(), other.end(), begin());
+    if constexpr(::std::allocator_traits<Allocator>::propagate_on_container_copy_assignment::value)
+    {
+      _alloc = other._alloc;
     }
+    resize(other.size());
+    _construct(other.begin(), other.end());
     return *this;
   }
   device_vector &operator=(device_vector &&other) {
     // Move assignment operator:
-    device_vector dummy(::std::move(other));
-    this->swap(dummy);
+    if constexpr(::std::allocator_traits<Allocator>::propagate_on_container_move_assignment::value)
+    {
+      _alloc = std::move(other._alloc);
+    }
+    _size = std::move(other._size);
+    _capacity = std::move(other._capacity);
+    _storage = std::move(other._storage);
     return *this;
   }
   size_type size() const { return _size; }
@@ -220,28 +427,32 @@ class device_vector {
     ::std::swap(_size, v._size);
     ::std::swap(_capacity, v._capacity);
     ::std::swap(_storage, v._storage);
-    ::std::swap(_alloc, v._alloc);
+    if constexpr (::std::allocator_traits<Allocator>::propagate_on_container_swap::value)
+    {
+      ::std::swap(_alloc, v._alloc);
+    }
   }
-  reference operator[](size_type n) { return _storage[n]; }
-  const_reference operator[](size_type n) const { return _storage[n]; }
+  reference operator[](size_type n) { return reference(_storage + n); }
+  const_reference operator[](size_type n) const { return reference(_storage + n); }
   void reserve(size_type n) {
     if (n > capacity()) {
       // allocate buffer for new size
-      auto tmp = _alloc.allocate(2 * n);
+      auto tmp = ::std::allocator_traits<Allocator>::allocate(_alloc, 2 * n);
       // copy content (old buffer to new buffer)
-      ::std::copy(oneapi::dpl::execution::make_device_policy(get_default_queue()),
-                begin(), end(), tmp);
+      ::std::copy(
+          oneapi::dpl::execution::make_device_policy(get_default_queue()),
+          begin(), end(), tmp);
       // deallocate old memory
-      _alloc.deallocate(_storage, _capacity);
+      ::std::allocator_traits<Allocator>::deallocate(_alloc, _storage,
+                                                     _capacity);
       _storage = tmp;
       _capacity = 2 * n;
     }
   }
   void resize(size_type new_size, const T &x = T()) {
     reserve(new_size);
-    if (_size < new_size) {
-      ::std::fill(oneapi::dpl::execution::make_device_policy(get_default_queue()),
-                begin() + _size, begin() + new_size, x);
+    if (new_size > size()) {
+      _construct(new_size - size(), x, size());
     }
     _size = new_size;
   }
@@ -249,22 +460,24 @@ class device_vector {
     return ::std::numeric_limits<size_type>::max() / sizeof(T);
   }
   size_type capacity() const { return _capacity; }
-  const_reference front() const { return *begin(); }
-  reference front() { return *begin(); }
-  const_reference back(void) const { return *(end() - 1); }
-  reference back(void) { return *(end() - 1); }
+  const_reference front() const { return reference(begin()); }
+  reference front() { return reference(begin()); }
+  const_reference back(void) const { return reference(end() - 1); }
+  reference back(void) { return reference(end() - 1); }
   pointer data(void) { return _storage; }
   const_pointer data(void) const { return _storage; }
   void shrink_to_fit(void) {
     if (_size != capacity()) {
       size_type tmp_capacity = ::std::max(_size, _min_capacity());
-      auto tmp = _alloc.allocate(tmp_capacity);
+      auto tmp =
+          ::std::allocator_traits<Allocator>::allocate(_alloc, tmp_capacity);
       if (_size > 0) {
         ::std::copy(
             oneapi::dpl::execution::make_device_policy(get_default_queue()),
             begin(), end(), tmp);
       }
-      _alloc.deallocate(_storage, _capacity);
+      ::std::allocator_traits<Allocator>::deallocate(_alloc, _storage,
+                                                     _capacity);
       _storage = tmp;
       _capacity = tmp_capacity;
     }
@@ -272,20 +485,22 @@ class device_vector {
   void assign(size_type n, const T &x) {
     resize(n);
     if (_size > 0) {
-      ::std::fill(oneapi::dpl::execution::make_device_policy(get_default_queue()),
-                begin(), begin() + n, x);
+      ::std::fill(
+          oneapi::dpl::execution::make_device_policy(get_default_queue()),
+          begin(), begin() + n, x);
     }
   }
   template <typename InputIterator>
   void
   assign(InputIterator first,
          typename ::std::enable_if<internal::is_iterator<InputIterator>::value,
-                                 InputIterator>::type last) {
+                                   InputIterator>::type last) {
     auto n = ::std::distance(first, last);
     resize(n);
     if (_size > 0) {
-      ::std::copy(oneapi::dpl::execution::make_device_policy(get_default_queue()),
-                first, last, begin());
+      ::std::copy(
+          oneapi::dpl::execution::make_device_policy(get_default_queue()),
+          first, last, begin());
     }
   }
   void clear(void) { _size = 0; }
@@ -305,14 +520,18 @@ class device_vector {
     if (m <= 0) {
       return end();
     }
-    auto tmp = _alloc.allocate(m);
+    auto tmp = ::std::allocator_traits<Allocator>::allocate(_alloc, m);
     // copy remainder to temporary buffer.
     ::std::copy(oneapi::dpl::execution::make_device_policy(get_default_queue()),
-              last, end(), tmp);
+                last, end(), tmp);
+
+    auto position = ::std::distance(begin(), first);
+    _destroy(n, position);
+
     // override (erase) subsequence in storage.
     ::std::copy(oneapi::dpl::execution::make_device_policy(get_default_queue()),
-              tmp, tmp + m, first);
-    _alloc.deallocate(tmp, m);
+                tmp, tmp + m, first);
+    ::std::allocator_traits<Allocator>::deallocate(_alloc, tmp, m);
     _size -= n;
     return begin() + first.get_idx() + n;
   }
@@ -325,57 +544,61 @@ class device_vector {
   void insert(iterator position, size_type n, const T &x) {
     if (position == end()) {
       resize(size() + n);
-      ::std::fill(oneapi::dpl::execution::make_device_policy(get_default_queue()),
-                end() - n, end(), x);
+      _construct(n, x, size() - n);
     } else {
       auto i_n = ::std::distance(begin(), position);
       // allocate temporary storage
       auto m = ::std::distance(position, end());
       // will throw if position is not inside active vector
-      auto tmp = _alloc.allocate(m);
+      auto tmp = ::std::allocator_traits<Allocator>::allocate(_alloc, m);
       // copy remainder
-      ::std::copy(oneapi::dpl::execution::make_device_policy(get_default_queue()),
-                position, end(), tmp);
+      ::std::copy(
+          oneapi::dpl::execution::make_device_policy(get_default_queue()),
+          position, end(), tmp);
 
       resize(size() + n);
       // resizing might invalidate position
       position = begin() + position.get_idx();
 
-      ::std::fill(oneapi::dpl::execution::make_device_policy(get_default_queue()),
-                position, position + n, x);
+      _construct(n, x, position.get_idx());
 
-      ::std::copy(oneapi::dpl::execution::make_device_policy(get_default_queue()),
-                tmp, tmp + m, position + n);
-      _alloc.deallocate(tmp, m);
+      ::std::copy(
+          oneapi::dpl::execution::make_device_policy(get_default_queue()), tmp,
+          tmp + m, position + n);
+      ::std::allocator_traits<Allocator>::deallocate(_alloc, tmp, m);
     }
   }
   template <typename InputIterator>
   void
   insert(iterator position, InputIterator first,
          typename ::std::enable_if<internal::is_iterator<InputIterator>::value,
-                                 InputIterator>::type last) {
+                                   InputIterator>::type last) {
     auto n = ::std::distance(first, last);
     if (position == end()) {
       resize(size() + n);
-      ::std::copy(oneapi::dpl::execution::make_device_policy(get_default_queue()),
-                first, last, end());
+      _construct(first, last, size() - n);
+      ::std::copy(
+          oneapi::dpl::execution::make_device_policy(get_default_queue()),
+          first, last, end());
     } else {
       auto m = ::std::distance(position, end());
       // will throw if position is not inside active vector
-      auto tmp = _alloc.allocate(m);
+      auto tmp = ::std::allocator_traits<Allocator>::allocate(_alloc, m);
 
-      ::std::copy(oneapi::dpl::execution::make_device_policy(get_default_queue()),
-                position, end(), tmp);
+      ::std::copy(
+          oneapi::dpl::execution::make_device_policy(get_default_queue()),
+          position, end(), tmp);
 
       resize(size() + n);
       // resizing might invalidate position
       position = begin() + position.get_idx();
 
-      ::std::copy(oneapi::dpl::execution::make_device_policy(get_default_queue()),
-                first, last, position);
-      ::std::copy(oneapi::dpl::execution::make_device_policy(get_default_queue()),
-                tmp, tmp + m, position + n);
-      _alloc.deallocate(tmp, m);
+      _construct(first, last, position.get_idx());
+
+      ::std::copy(
+          oneapi::dpl::execution::make_device_policy(get_default_queue()), tmp,
+          tmp + m, position + n);
+      ::std::allocator_traits<Allocator>::deallocate(_alloc, tmp, m);
     }
   }
   Allocator get_allocator() const { return _alloc; }