diff --git a/include/simsycl/sycl.hh b/include/simsycl/sycl.hh
index 9c3a70c..3c38de5 100644
--- a/include/simsycl/sycl.hh
+++ b/include/simsycl/sycl.hh
@@ -27,3 +27,4 @@
 #include "sycl/reduction.hh"
 #include "sycl/sub_group.hh"
 #include "sycl/type_traits.hh"
+#include "sycl/usm.hh"
diff --git a/include/simsycl/sycl/enums.hh b/include/simsycl/sycl/enums.hh
index 1e9d174..dcb42a2 100644
--- a/include/simsycl/sycl/enums.hh
+++ b/include/simsycl/sycl/enums.hh
@@ -1,5 +1,8 @@
 #pragma once
 
+#include <cstdint>
+
+
 namespace simsycl::sycl {
 
 enum class addressing_mode { mirrored_repeat, repeat, clamp_to_edge, clamp, none };
@@ -307,9 +310,17 @@ struct command_execution_status;
 
 namespace simsycl::sycl::info::event_profiling {
 
-struct command_submit;
-struct command_start;
-struct command_end;
+struct command_submit {
+    using return_type = uint64_t;
+};
+
+struct command_start {
+    using return_type = uint64_t;
+};
+
+struct command_end {
+    using return_type = uint64_t;
+};
 
 } // namespace simsycl::sycl::info::event_profiling
 
diff --git a/include/simsycl/sycl/event.hh b/include/simsycl/sycl/event.hh
index 932ed36..386d73a 100644
--- a/include/simsycl/sycl/event.hh
+++ b/include/simsycl/sycl/event.hh
@@ -1,36 +1,110 @@
 #pragma once
 
 #include "forward.hh"
+#include "type_traits.hh"
 
+#include "../detail/reference_type.hh"
+
+#include <chrono>
 #include <vector>
 
+
+namespace simsycl::detail {
+
+struct execution_status {
+    std::chrono::steady_clock::time_point t_submit{};
+    std::chrono::steady_clock::time_point t_start{};
+    std::chrono::steady_clock::time_point t_end{};
+
+    void start() { t_start = std::chrono::steady_clock::now(); }
+
+    [[nodiscard]] static execution_status submit() {
+        execution_status status;
+        status.t_submit = std::chrono::steady_clock::now();
+        return status;
+    }
+
+    [[nodiscard]] static execution_status submit_and_start() {
+        auto status = submit();
+        status.start();
+        return status;
+    }
+
+    [[nodiscard]] sycl::event end();
+
+    [[nodiscard]] static sycl::event instant();
+};
+
+template <typename Clock, typename Dur>
+uint64_t nanoseconds_since_epoch(std::chrono::time_point<Clock, Dur> time_point) {
+    return std::chrono::duration_cast<std::chrono::nanoseconds>(time_point.time_since_epoch()).count();
+}
+
+} // namespace simsycl::detail
+
 namespace simsycl::sycl {
 
-class event {
- public:
-  event() = default;
+class event : detail::reference_type<event, detail::execution_status> {
+  public:
+    event() = default;
 
-  /* -- common interface members -- */
+    backend get_backend() const noexcept;
 
-  backend get_backend() const noexcept;
+    std::vector<event> get_wait_list();
 
-  std::vector<event> get_wait_list();
+    void wait() {}
 
-  void wait() {}
+    static void wait(const std::vector<event> & /* event_list */) {}
 
-  static void wait(const std::vector<event>& /* event_list */) {}
+    void wait_and_throw() {}
 
-  void wait_and_throw() {}
+    static void wait_and_throw(const std::vector<event> & /* event_list */) {}
 
-  static void wait_and_throw(const std::vector<event>& /* event_list */) {}
+    template <typename Param>
+    typename Param::return_type get_info() const {
+        if constexpr(std::is_same_v<Param, info::event::command_execution_status>) {
+            return info::event_command_status::complete;
+        } else {
+            static_assert(detail::always_false<Param>, "Unknown event::get_info() parameter");
+        }
+    }
 
-  template <typename Param> typename Param::return_type get_info() const;
+    template <typename Param>
+    typename Param::return_type get_backend_info() const {
+        static_assert(detail::always_false<Param>, "Unknown event::get_backend_info() parameter");
+    }
 
-  template <typename Param>
-  typename Param::return_type get_backend_info() const;
+    template <typename Param>
+    typename Param::return_type get_profiling_info() const {
+        if constexpr(std::is_same_v<Param, info::event_profiling::command_submit>) {
+            return detail::nanoseconds_since_epoch(state().t_submit);
+        } else if constexpr(std::is_same_v<Param, info::event_profiling::command_start>) {
+            return detail::nanoseconds_since_epoch(state().t_start);
+        } else if constexpr(std::is_same_v<Param, info::event_profiling::command_end>) {
+            return detail::nanoseconds_since_epoch(state().t_end);
+        } else {
+            static_assert(detail::always_false<Param>, "Unknown event::get_profiling_info() parameter");
+        }
+    }
 
-  template <typename Param>
-  typename Param::return_type get_profiling_info() const;
+  private:
+    friend event detail::make_event(const detail::execution_status &status);
+
+    explicit event(const detail::execution_status &status)
+        : detail::reference_type<event, detail::execution_status>(std::in_place, status) {}
 };
 
-} // namespace sycl
+} // namespace simsycl::sycl
+
+namespace simsycl::detail {
+
+inline sycl::event make_event(const execution_status &status) { return sycl::event(status); }
+
+inline sycl::event execution_status::end() {
+    t_end = std::chrono::steady_clock::now();
+    return make_event(*this);
+}
+
+inline sycl::event execution_status::instant() { return submit_and_start().end(); }
+
+} // namespace simsycl::detail
diff --git a/include/simsycl/sycl/forward.hh b/include/simsycl/sycl/forward.hh
index 30fa04d..3a10707 100644
--- a/include/simsycl/sycl/forward.hh
+++ b/include/simsycl/sycl/forward.hh
@@ -110,4 +110,8 @@ sycl::handler make_handler();
 
 void **require_local_memory(sycl::handler &cgh, size_t size, size_t align);
 
+struct execution_status;
+
+sycl::event make_event(const execution_status &status);
+
 } // namespace simsycl::detail
diff --git a/include/simsycl/sycl/group_algorithms.hh b/include/simsycl/sycl/group_algorithms.hh
index cc2fe29..5e7f49c 100644
--- a/include/simsycl/sycl/group_algorithms.hh
+++ b/include/simsycl/sycl/group_algorithms.hh
@@ -216,6 +216,10 @@ T permute_group(G g, T x, typename G::linear_id_type mask) {
                 }});
 }
 
+template <typename Group, typename T>
+T permute_group_by_xor(Group g, T x, typename Group::linear_id_type mask); // TODO
+
+
 // select
 
 template <SubGroup G, TriviallyCopyable T>
diff --git a/include/simsycl/sycl/multi_ptr.hh b/include/simsycl/sycl/multi_ptr.hh
index 3f2841c..1d7e5f8 100644
--- a/include/simsycl/sycl/multi_ptr.hh
+++ b/include/simsycl/sycl/multi_ptr.hh
@@ -1,6 +1,5 @@
 #pragma once
 
-#include "accessor.hh"
 #include "enums.hh"
 #include "forward.hh"
 
@@ -306,4 +305,36 @@ using global_ptr = multi_ptr<ElementType, access::address_space::global_space, I
 template <typename ElementType, access::decorated IsDecorated = access::decorated::legacy>
 using local_ptr = multi_ptr<ElementType, access::address_space::local_space, IsDecorated>;
 
+// Deprecated in SYCL 2020
+template <typename ElementType>
+
+using constant_ptr [[deprecated]]
+= multi_ptr<ElementType, access::address_space::constant_space, access::decorated::legacy>;
+
+template <typename ElementType, access::decorated IsDecorated = access::decorated::legacy>
+using private_ptr = multi_ptr<ElementType, access::address_space::private_space, IsDecorated>;
+
+// Template specialization aliases for different pointer address spaces.
+// The interface exposes non-decorated pointer while keeping the address space information internally.
+
+template <typename ElementType>
+using raw_global_ptr = multi_ptr<ElementType, access::address_space::global_space, access::decorated::no>;
+
+template <typename ElementType>
+using raw_local_ptr = multi_ptr<ElementType, access::address_space::local_space, access::decorated::no>;
+
+template <typename ElementType>
+using raw_private_ptr = multi_ptr<ElementType, access::address_space::private_space, access::decorated::no>;
+
+// Template specialization aliases for different pointer address spaces.
+// The interface exposes decorated pointer.
+template <typename ElementType>
+using decorated_global_ptr = multi_ptr<ElementType, access::address_space::global_space, access::decorated::yes>;
+
+template <typename ElementType>
+using decorated_local_ptr = multi_ptr<ElementType, access::address_space::local_space, access::decorated::yes>;
+
+template <typename ElementType>
+using decorated_private_ptr = multi_ptr<ElementType, access::address_space::private_space, access::decorated::yes>;
+
 } // namespace simsycl::sycl
diff --git a/include/simsycl/sycl/nd_item.hh b/include/simsycl/sycl/nd_item.hh
index eadcdf5..5c746a4 100644
--- a/include/simsycl/sycl/nd_item.hh
+++ b/include/simsycl/sycl/nd_item.hh
@@ -6,6 +6,7 @@
 
 #include "group.hh"
 #include "id.hh"
+#include "multi_ptr.hh"
 #include "range.hh"
 #include "sub_group.hh"
 
@@ -41,6 +42,11 @@ sycl::nd_item<Dimensions> make_nd_item(const sycl::item<Dimensions, false> &glob
 
 namespace simsycl::sycl {
 
+class device_event {
+  public:
+    void wait() noexcept {}
+};
+
 template <int Dimensions>
 class nd_item {
   public:
@@ -104,6 +110,46 @@ class nd_item {
         SIMSYCL_NOT_IMPLEMENTED(access_space);
     }
 
+    // Deprecated in SYCL 2020.
+    template <typename DataT>
+    [[deprecated]] device_event async_work_group_copy(
+        local_ptr<DataT> dest, global_ptr<DataT> src, size_t num_elements) const;
+
+    // Deprecated in SYCL 2020.
+    template <typename DataT>
+    [[deprecated]] device_event async_work_group_copy(
+        global_ptr<DataT> dest, local_ptr<DataT> src, size_t num_elements) const;
+
+    // Deprecated in SYCL 2020.
+    template <typename DataT>
+    [[deprecated]] device_event async_work_group_copy(
+        local_ptr<DataT> dest, global_ptr<DataT> src, size_t num_elements, size_t src_stride) const;
+
+    // Deprecated in SYCL 2020.
+    template <typename DataT>
+    [[deprecated]] device_event async_work_group_copy(
+        global_ptr<DataT> dest, local_ptr<DataT> src, size_t num_elements, size_t dest_stride) const;
+
+    template <typename DestDataT, typename SrcDataT>
+        requires(std::is_same_v<DestDataT, std::remove_const_t<SrcDataT>>)
+    device_event async_work_group_copy(
+        decorated_local_ptr<DestDataT> dest, decorated_global_ptr<SrcDataT> src, size_t num_elements) const;
+
+    template <typename DestDataT, typename SrcDataT>
+        requires(std::is_same_v<DestDataT, std::remove_const_t<SrcDataT>>)
+    device_event async_work_group_copy(
+        decorated_global_ptr<DestDataT> dest, decorated_local_ptr<SrcDataT> src, size_t num_elements) const;
+
+    template <typename DestDataT, typename SrcDataT>
+        requires(std::is_same_v<DestDataT, std::remove_const_t<SrcDataT>>)
+    device_event async_work_group_copy(decorated_local_ptr<DestDataT> dest, decorated_global_ptr<SrcDataT> src,
+        size_t num_elements, size_t src_stride) const;
+
+    template <typename DestDataT, typename SrcDataT>
+        requires(std::is_same_v<DestDataT, std::remove_const_t<SrcDataT>>)
+    device_event async_work_group_copy(decorated_global_ptr<DestDataT> dest, decorated_local_ptr<SrcDataT> src,
+        size_t num_elements, size_t dest_stride) const;
+
     template <typename... Events>
     void wait_for(Events... events) const {
         m_group.wait_for(events...);
diff --git a/include/simsycl/sycl/property.hh b/include/simsycl/sycl/property.hh
index a695135..5628b39 100644
--- a/include/simsycl/sycl/property.hh
+++ b/include/simsycl/sycl/property.hh
@@ -34,6 +34,22 @@ class property_list {
         requires(is_property_v<Properties> && ...)
     property_list(Properties... props) : m_properties{props...} {}
 
+    // Implemented by hipSYCL and DPC++ although the spec does not mention any members beside the constructor
+    template <typename Property>
+    bool has_property() const noexcept {
+        return std::any_of(m_properties.begin(), m_properties.end(),
+            [](const std::any &prop) { return prop.type() == typeid(Property); });
+    }
+
+    // Implemented by hipSYCL and DPC++ although the spec does not mention any members beside the constructor
+    template <typename Property>
+    Property get_property() const {
+        const auto iter = std::find_if(m_properties.begin(), m_properties.end(),
+            [](const std::any &prop) { return prop.type() == typeid(Property); });
+        SIMSYCL_CHECK(iter != m_properties.end());
+        return std::any_cast<Property>(*iter);
+    }
+
   private:
     friend class detail::property_interface;
 
diff --git a/include/simsycl/sycl/queue.hh b/include/simsycl/sycl/queue.hh
index 40fe0f9..61341a2 100644
--- a/include/simsycl/sycl/queue.hh
+++ b/include/simsycl/sycl/queue.hh
@@ -5,6 +5,9 @@
 #include "handler.hh"
 #include "property.hh"
 
+#include "../detail/reference_type.hh"
+
+
 namespace simsycl::sycl::property::queue {
 
 class enable_profiling {};
@@ -26,18 +29,27 @@ struct is_property_of<property::queue::in_order, queue> : std::true_type {};
 
 } // namespace simsycl::sycl
 
+namespace simsycl::detail {
+
+struct queue_state {};
+
+} // namespace simsycl::detail
+
 namespace simsycl::sycl {
 
-class queue : public simsycl::detail::property_interface {
+class queue : public detail::reference_type<queue, detail::queue_state>, public simsycl::detail::property_interface {
   private:
-    using property_compatibility
-        = simsycl::detail::property_compatibility_with<queue, property::queue::enable_profiling, property::queue::in_order>;
+    using reference_type = detail::reference_type<queue, detail::queue_state>;
+    using property_compatibility = simsycl::detail::property_compatibility_with<queue,
+        property::queue::enable_profiling, property::queue::in_order>;
 
   public:
-    explicit queue(const property_list &prop_list = {}) : property_interface(prop_list, property_compatibility()) {}
+    explicit queue(const property_list &prop_list = {})
+        : reference_type(std::in_place), property_interface(prop_list, property_compatibility()) {}
 
     explicit queue(const async_handler &async_handler, const property_list &prop_list = {})
-        : property_interface(prop_list, property_compatibility()), m_async_handler(async_handler) {}
+        : reference_type(std::in_place), property_interface(prop_list, property_compatibility()),
+          m_async_handler(async_handler) {}
 
     template <typename DeviceSelector>
     explicit queue(const DeviceSelector &device_selector, const property_list &prop_list = {});
@@ -63,8 +75,6 @@ class queue : public simsycl::detail::property_interface {
     explicit queue(const context &sycl_context, const device &sycl_device, const async_handler &async_handler,
         const property_list &prop_list = {});
 
-    /* -- common interface members -- */
-
     backend get_backend() const noexcept;
 
     context get_context() const;
@@ -81,9 +91,11 @@ class queue : public simsycl::detail::property_interface {
 
     template <typename T>
     event submit(T cgf) {
+        auto status = detail::execution_status::submit();
         auto cgh = simsycl::detail::make_handler();
+        status.start();
         cgf(cgh);
-        return event();
+        return status.end();
     }
 
     template <typename T>
@@ -100,145 +112,170 @@ class queue : public simsycl::detail::property_interface {
 
     template <typename KernelName, typename KernelType>
     event single_task(const KernelType &kernel_func) {
+        auto status = detail::execution_status::submit_and_start();
         kernel_func();
-        return event();
+        return status.end();
     }
 
     template <typename KernelName, typename KernelType>
     event single_task(event /* dep_event */, const KernelType &kernel_func) {
+        auto status = detail::execution_status::submit_and_start();
         kernel_func();
-        return event();
+        return status.end();
     }
 
     template <typename KernelName, typename KernelType>
     event single_task(const std::vector<event> & /* dep_events */, const KernelType &kernel_func) {
+        auto status = detail::execution_status::submit_and_start();
         kernel_func();
-        return event();
+        return status.end();
     }
 
     template <typename KernelName, int Dims, typename... Rest, std::enable_if_t<(sizeof...(Rest) > 0), int> = 0>
     event parallel_for(range<Dims> num_work_items, Rest &&...rest) {
+        auto status = detail::execution_status::submit_and_start();
         simsycl::detail::parallel_for(num_work_items, std::forward<Rest>(rest)...);
-        return event();
+        return status.end();
     }
 
     template <typename KernelName, int Dims, typename... Rest, std::enable_if_t<(sizeof...(Rest) > 0), int> = 0>
     event parallel_for(range<Dims> num_work_items, event /* dep_event */, Rest &&...rest) {
+        auto status = detail::execution_status::submit_and_start();
         simsycl::detail::parallel_for(num_work_items, std::forward<Rest>(rest)...);
-        return event();
+        return status.end();
     }
 
     template <typename KernelName, int Dims, typename... Rest, std::enable_if_t<(sizeof...(Rest) > 0), int> = 0>
     event parallel_for(range<Dims> num_work_items, const std::vector<event> & /* dep_events */, Rest &&...rest) {
+        auto status = detail::execution_status::submit_and_start();
         simsycl::detail::parallel_for(num_work_items, std::forward<Rest>(rest)...);
-        return event();
+        return status.end();
     }
 
     template <typename KernelName, int Dims, typename... Rest, std::enable_if_t<(sizeof...(Rest) > 0), int> = 0>
     event parallel_for(nd_range<Dims> execution_range, Rest &&...rest) {
+        auto status = detail::execution_status::submit_and_start();
         simsycl::detail::parallel_for(execution_range, std::forward<Rest>(rest)...);
-        return event();
+        return status.end();
     }
 
     template <typename KernelName, int Dims, typename... Rest, std::enable_if_t<(sizeof...(Rest) > 0), int> = 0>
     event parallel_for(nd_range<Dims> execution_range, event /* dep_event */, Rest &&...rest) {
+        auto status = detail::execution_status::submit_and_start();
         simsycl::detail::parallel_for(execution_range, std::forward<Rest>(rest)...);
-        return event();
+        return status.end();
     }
 
     template <typename KernelName, int Dims, typename... Rest, std::enable_if_t<(sizeof...(Rest) > 0), int> = 0>
     event parallel_for(nd_range<Dims> execution_range, const std::vector<event> & /* dep_events */, Rest &&...rest) {
+        auto status = detail::execution_status::submit_and_start();
         simsycl::detail::parallel_for(execution_range, std::forward<Rest>(rest)...);
-        return event();
+        return status.end();
     }
 
     /* -- USM functions -- */
 
     event memcpy(void *dest, const void *src, size_t num_bytes) {
+        auto status = detail::execution_status::submit_and_start();
         ::memcpy(dest, src, num_bytes);
-        return event();
+        return status.end();
     }
 
     event memcpy(void *dest, const void *src, size_t num_bytes, event /* dep_event */) {
+        auto status = detail::execution_status::submit_and_start();
         ::memcpy(dest, src, num_bytes);
-        return event();
+        return status.end();
     }
 
     event memcpy(void *dest, const void *src, size_t num_bytes, const std::vector<event> & /* dep_events */) {
+        auto status = detail::execution_status::submit_and_start();
         ::memcpy(dest, src, num_bytes);
-        return event();
+        return status.end();
     }
 
     template <typename T>
     event copy(const T *src, T *dest, size_t count) {
+        auto status = detail::execution_status::submit_and_start();
         std::copy_n(src, count, dest);
-        return event();
+        return status.end();
     }
 
     template <typename T>
     event copy(const T *src, T *dest, size_t count, event dep_event) {
         (void)(dep_event);
+        auto status = detail::execution_status::submit_and_start();
         std::copy_n(src, count, dest);
-        return event();
+        return status.end();
     }
 
     template <typename T>
     event copy(const T *src, T *dest, size_t count, const std::vector<event> &dep_events) {
         (void)(dep_events);
+        auto status = detail::execution_status::submit_and_start();
         std::copy_n(src, count, dest);
-        return event();
+        return status.end();
     }
 
     event memset(void *ptr, int value, size_t num_bytes) {
+        auto status = detail::execution_status::submit_and_start();
         ::memset(ptr, value, num_bytes);
-        return event();
+        return status.end();
     }
 
     event memset(void *ptr, int value, size_t num_bytes, event /* dep_event */) {
+        auto status = detail::execution_status::submit_and_start();
         ::memset(ptr, value, num_bytes);
-        return event();
+        return status.end();
     }
 
     event memset(void *ptr, int value, size_t num_bytes, const std::vector<event> & /* dep_events */) {
+        auto status = detail::execution_status::submit_and_start();
         ::memset(ptr, value, num_bytes);
-        return event();
+        return status.end();
     }
 
     template <typename T>
     event fill(void *ptr, const T &pattern, size_t count) {
+        auto status = detail::execution_status::submit_and_start();
         std::fill_n(ptr, count, pattern);
-        return event();
+        return status.end();
     }
 
     template <typename T>
     event fill(void *ptr, const T &pattern, size_t count, event /* dep_event */) {
+        auto status = detail::execution_status::submit_and_start();
         std::fill_n(ptr, count, pattern);
-        return event();
+        return status.end();
     }
 
     template <typename T>
     event fill(void *ptr, const T &pattern, size_t count, const std::vector<event> & /* dep_events */) {
+        auto status = detail::execution_status::submit_and_start();
         std::fill_n(ptr, count, pattern);
-        return event();
+        return status.end();
     }
 
-    event prefetch(void * /* ptr */, size_t /* num_bytes */) { return event(); }
+    event prefetch(void * /* ptr */, size_t /* num_bytes */) { return detail::execution_status::instant(); }
 
-    event prefetch(void * /* ptr */, size_t /* num_bytes */, event /* dep_event */) { return event(); }
+    event prefetch(void * /* ptr */, size_t /* num_bytes */, event /* dep_event */) {
+        return detail::execution_status::instant();
+    }
 
     event prefetch(void * /* ptr */, size_t /* num_bytes */, const std::vector<event> & /* dep_events */) {
-        return event();
+        return detail::execution_status::instant();
     }
 
-    event mem_advise(void * /* ptr */, size_t /* num_bytes */, int /* advice */) { return event(); }
+    event mem_advise(void * /* ptr */, size_t /* num_bytes */, int /* advice */) {
+        return detail::execution_status::instant();
+    }
 
     event mem_advise(void * /* ptr */, size_t /* num_bytes */, int /* advice */, event /* dep_event */) {
-        return event();
+        return detail::execution_status::instant();
     }
 
     event mem_advise(
         void * /* ptr */, size_t /* num_bytes */, int /* advice */, const std::vector<event> & /* dep_events */) {
-        return event();
+        return detail::execution_status::instant();
     }
 
     /// Placeholder accessor shortcuts
diff --git a/include/simsycl/sycl/type_traits.hh b/include/simsycl/sycl/type_traits.hh
index 85b37b1..56317f5 100644
--- a/include/simsycl/sycl/type_traits.hh
+++ b/include/simsycl/sycl/type_traits.hh
@@ -29,5 +29,7 @@ struct is_arithmetic : std::bool_constant<std::is_arithmetic_v<T>> {};
 template <class T>
 inline constexpr bool is_arithmetic_v = is_arithmetic<T>::value;
 
+template<typename...>
+constexpr bool always_false = false;
 
 } // namespace simsycl::sycl
diff --git a/include/simsycl/sycl/usm.hh b/include/simsycl/sycl/usm.hh
new file mode 100644
index 0000000..a2fc115
--- /dev/null
+++ b/include/simsycl/sycl/usm.hh
@@ -0,0 +1,331 @@
+#pragma once
+
+#include "enums.hh"
+#include "forward.hh"
+#include "property.hh"
+
+
+namespace simsycl::sycl {
+
+template <typename T, usm::alloc AllocKind, size_t Alignment = 0>
+class usm_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;
+
+    template <typename U>
+    struct rebind {
+        typedef usm_allocator<U, AllocKind, Alignment> other;
+    };
+
+    usm_allocator() = delete;
+
+    usm_allocator(const context &sycl_context, const device &sycl_device, const property_list &prop_list = {}) {
+        (void)sycl_context;
+        (void)sycl_device;
+        (void)prop_list;
+    }
+
+    usm_allocator(const queue &sycl_queue, const property_list &prop_list = {}) {
+        (void)sycl_queue;
+        (void)prop_list;
+    }
+
+    template <class U>
+    usm_allocator(const usm_allocator<U, AllocKind, Alignment> &other) noexcept {
+        (void)other;
+    }
+
+    usm_allocator(const usm_allocator &other) = default;
+    usm_allocator(usm_allocator &&) noexcept;
+    usm_allocator &operator=(const usm_allocator &) = default;
+    usm_allocator &operator=(usm_allocator &&) = default;
+
+    /// Allocate memory
+    T *allocate(size_t count) { return static_cast<T *>(std::aligned_alloc(Alignment, count * sizeof(T))); }
+
+    /// Deallocate memory
+    void deallocate(T *ptr, size_t count) {
+        (void)count;
+        std::free(ptr);
+    }
+
+    /// Equality Comparison
+    ///
+    /// Allocators only compare equal if they are of the same USM kind, alignment, context, and device
+    template <class U, usm::alloc AllocKindU, size_t AlignmentU>
+    friend bool operator==(
+        const usm_allocator<T, AllocKind, Alignment> &lhs, const usm_allocator<U, AllocKindU, AlignmentU> &rhs);
+
+    /// Inequality Comparison
+    /// Allocators only compare unequal if they are not of the same USM kind, alignment, context, or device
+    template <class U, usm::alloc AllocKindU, size_t AlignmentU>
+    friend bool operator!=(
+        const usm_allocator<T, AllocKind, Alignment> &lhs, const usm_allocator<U, AllocKindU, AlignmentU> &rhs);
+};
+
+
+inline void *malloc_device(
+    size_t num_bytes, const device &sycl_device, const context &sycl_context, const property_list &prop_list = {}) {
+    (void)sycl_device;
+    (void)sycl_context;
+    (void)prop_list;
+    return std::malloc(num_bytes);
+}
+
+template <typename T>
+T *malloc_device(
+    size_t count, const device &sycl_device, const context &sycl_context, const property_list &prop_list = {}) {
+    (void)sycl_device;
+    (void)sycl_context;
+    (void)prop_list;
+    return static_cast<T *>(std::aligned_alloc(alignof(T), count * sizeof(T)));
+}
+
+inline void *malloc_device(size_t num_bytes, const queue &sycl_queue, const property_list &prop_list = {}) {
+    (void)sycl_queue;
+    (void)prop_list;
+    return std::malloc(num_bytes);
+}
+
+template <typename T>
+T *malloc_device(size_t count, const queue &sycl_queue, const property_list &prop_list = {}) {
+    (void)sycl_queue;
+    (void)prop_list;
+    return static_cast<T *>(std::aligned_alloc(alignof(T), count * sizeof(T)));
+}
+
+inline void *aligned_alloc_device(size_t alignment, size_t num_bytes, const device &sycl_device,
+    const context &sycl_context, const property_list &prop_list = {}) {
+    (void)sycl_device;
+    (void)sycl_context;
+    (void)prop_list;
+    return std::aligned_alloc(alignment, num_bytes);
+}
+
+template <typename T>
+T *aligned_alloc_device(size_t alignment, size_t count, const device &sycl_device, const context &sycl_context,
+    const property_list &prop_list = {}) {
+    (void)sycl_device;
+    (void)sycl_context;
+    (void)prop_list;
+    return static_cast<T *>(std::aligned_alloc(alignment, count * sizeof(T)));
+}
+
+inline void *aligned_alloc_device(
+    size_t alignment, size_t num_bytes, const queue &sycl_queue, const property_list &prop_list = {}) {
+    (void)sycl_queue;
+    (void)prop_list;
+    return std::aligned_alloc(alignment, num_bytes);
+}
+
+template <typename T>
+T *aligned_alloc_device(size_t alignment, size_t count, const queue &sycl_queue, const property_list &prop_list = {}) {
+    (void)sycl_queue;
+    (void)prop_list;
+    return static_cast<T *>(std::aligned_alloc(alignment, count * sizeof(T)));
+};
+
+inline void *malloc_host(size_t num_bytes, const context &sycl_context, const property_list &prop_list = {}) {
+    (void)sycl_context;
+    (void)prop_list;
+    return std::malloc(num_bytes);
+}
+
+template <typename T>
+T *malloc_host(size_t count, const context &sycl_context, const property_list &prop_list = {}) {
+    (void)sycl_context;
+    (void)prop_list;
+    return static_cast<T *>(std::aligned_alloc(alignof(T), count * sizeof(T)));
+}
+
+inline void *malloc_host(size_t num_bytes, const queue &sycl_queue, const property_list &prop_list = {}) {
+    (void)sycl_queue;
+    (void)prop_list;
+    return std::malloc(num_bytes);
+}
+
+template <typename T>
+T *malloc_host(size_t count, const queue &sycl_queue, const property_list &prop_list = {}) {
+    (void)sycl_queue;
+    (void)prop_list;
+    return static_cast<T *>(std::aligned_alloc(alignof(T), count * sizeof(T)));
+}
+
+inline void *aligned_alloc_host(
+    size_t alignment, size_t num_bytes, const context &sycl_context, const property_list &prop_list = {}) {
+    (void)sycl_context;
+    (void)prop_list;
+    return std::aligned_alloc(alignment, num_bytes);
+}
+
+template <typename T>
+T *aligned_alloc_host(
+    size_t alignment, size_t count, const context &sycl_context, const property_list &prop_list = {}) {
+    (void)sycl_context;
+    (void)prop_list;
+    return static_cast<T *>(std::aligned_alloc(alignment, count * sizeof(T)));
+}
+
+inline void *aligned_alloc_host(
+    size_t alignment, size_t num_bytes, const queue &sycl_queue, const property_list &prop_list = {}) {
+    (void)sycl_queue;
+    (void)prop_list;
+    return std::aligned_alloc(alignment, num_bytes);
+}
+
+template <typename T>
+void *aligned_alloc_host(size_t alignment, size_t count, const queue &sycl_queue, const property_list &prop_list = {}) {
+    (void)sycl_queue;
+    (void)prop_list;
+    return static_cast<T *>(std::aligned_alloc(alignment, count * sizeof(T)));
+}
+
+inline void *malloc_shared(
+    size_t num_bytes, const device &sycl_device, const context &sycl_context, const property_list &prop_list = {}) {
+    (void)sycl_device;
+    (void)sycl_context;
+    (void)prop_list;
+    return std::malloc(num_bytes);
+}
+
+template <typename T>
+T *malloc_shared(
+    size_t count, const device &sycl_device, const context &sycl_context, const property_list &prop_list = {}) {
+    (void)sycl_device;
+    (void)sycl_context;
+    (void)prop_list;
+    return static_cast<T *>(std::aligned_alloc(alignof(T), count * sizeof(T)));
+}
+
+inline void *malloc_shared(size_t num_bytes, const queue &sycl_queue, const property_list &prop_list = {}) {
+    (void)sycl_queue;
+    (void)prop_list;
+    return std::malloc(num_bytes);
+}
+
+template <typename T>
+T *malloc_shared(size_t count, const queue &sycl_queue, const property_list &prop_list = {}) {
+    (void)sycl_queue;
+    (void)prop_list;
+    return static_cast<T *>(std::aligned_alloc(alignof(T), count * sizeof(T)));
+}
+
+inline void *aligned_alloc_shared(size_t alignment, size_t num_bytes, const device &sycl_device,
+    const context &sycl_context, const property_list &prop_list = {}) {
+    (void)sycl_device;
+    (void)sycl_context;
+    (void)prop_list;
+    return std::aligned_alloc(alignment, num_bytes);
+}
+
+template <typename T>
+T *aligned_alloc_shared(size_t alignment, size_t count, const device &sycl_device, const context &sycl_context,
+    const property_list &prop_list = {}) {
+    (void)sycl_device;
+    (void)sycl_context;
+    (void)prop_list;
+    return static_cast<T *>(std::aligned_alloc(alignment, count * sizeof(T)));
+}
+
+inline void *aligned_alloc_shared(
+    size_t alignment, size_t num_bytes, const queue &sycl_queue, const property_list &prop_list = {}) {
+    (void)sycl_queue;
+    (void)prop_list;
+    return std::aligned_alloc(alignment, num_bytes);
+}
+
+template <typename T>
+T *aligned_alloc_shared(size_t alignment, size_t count, const queue &sycl_queue, const property_list &prop_list = {}) {
+    (void)sycl_queue;
+    (void)prop_list;
+    return static_cast<T *>(std::aligned_alloc(alignment, count * sizeof(T)));
+}
+
+inline void *malloc(size_t num_bytes, const device &sycl_device, const context &sycl_context, usm::alloc kind,
+    const property_list &prop_list = {}) {
+    (void)sycl_device;
+    (void)sycl_context;
+    (void)kind;
+    (void)prop_list;
+    return std::malloc(num_bytes);
+}
+
+template <typename T>
+T *malloc(size_t count, const device &sycl_device, const context &sycl_context, usm::alloc kind,
+    const property_list &prop_list = {}) {
+    (void)sycl_device;
+    (void)sycl_context;
+    (void)kind;
+    (void)prop_list;
+    return static_cast<T *>(std::aligned_alloc(alignof(T), count * sizeof(T)));
+}
+
+inline void *malloc(size_t num_bytes, const queue &sycl_queue, usm::alloc kind, const property_list &prop_list = {}) {
+    (void)sycl_queue;
+    (void)kind;
+    (void)prop_list;
+    return std::malloc(num_bytes);
+}
+
+template <typename T>
+T *malloc(size_t count, const queue &sycl_queue, usm::alloc kind, const property_list &prop_list = {}) {
+    (void)sycl_queue;
+    (void)kind;
+    (void)prop_list;
+    return static_cast<T *>(std::aligned_alloc(alignof(T), count * sizeof(T)));
+}
+
+inline void *aligned_alloc(size_t alignment, size_t num_bytes, const device &sycl_device, const context &sycl_context,
+    usm::alloc kind, const property_list &prop_list = {}) {
+    (void)sycl_device;
+    (void)sycl_context;
+    (void)kind;
+    (void)prop_list;
+    return std::aligned_alloc(alignment, num_bytes);
+}
+
+template <typename T>
+T *aligned_alloc(size_t alignment, size_t count, const device &sycl_device, const context &sycl_context,
+    usm::alloc kind, const property_list &prop_list = {}) {
+    (void)sycl_device;
+    (void)sycl_context;
+    (void)kind;
+    (void)prop_list;
+    return static_cast<T *>(std::aligned_alloc(alignment, count * sizeof(T)));
+}
+
+inline void *aligned_alloc(
+    size_t alignment, size_t num_bytes, const queue &sycl_queue, usm::alloc kind, const property_list &prop_list = {}) {
+    (void)sycl_queue;
+    (void)kind;
+    (void)prop_list;
+    return std::aligned_alloc(alignment, num_bytes);
+}
+
+template <typename T>
+T *aligned_alloc(
+    size_t alignment, size_t count, const queue &sycl_queue, usm::alloc kind, const property_list &prop_list = {}) {
+    (void)sycl_queue;
+    (void)kind;
+    (void)prop_list;
+    return static_cast<T *>(std::aligned_alloc(alignment, count * sizeof(T)));
+}
+
+inline void free(void *ptr, const context &sycl_context) {
+    (void)sycl_context;
+    std::free(ptr);
+}
+
+inline void free(void *ptr, const queue &sycl_queue) {
+    (void)sycl_queue;
+    std::free(ptr);
+}
+
+usm::alloc get_pointer_type(const void *ptr, const context &sycl_context);
+
+device get_pointer_device(const void *ptr, const context &sycl_context);
+
+} // namespace simsycl::sycl