diff --git a/sycl/include/sycl/ext/oneapi/annotated_arg/annotated_arg.hpp b/sycl/include/sycl/ext/oneapi/annotated_arg/annotated_arg.hpp
new file mode 100644
index 0000000000000..f58d1596248c8
--- /dev/null
+++ b/sycl/include/sycl/ext/oneapi/annotated_arg/annotated_arg.hpp
@@ -0,0 +1,277 @@
+//==----------- annotated_arg.hpp - SYCL annotated_arg extension -----------==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+#include <cstddef>
+#include <type_traits>
+
+#include <sycl/detail/stl_type_traits.hpp>
+#include <sycl/exception.hpp>
+#include <sycl/ext/oneapi/annotated_arg/properties.hpp>
+#include <sycl/ext/oneapi/properties/properties.hpp>
+
+namespace sycl {
+__SYCL_INLINE_VER_NAMESPACE(_V1) {
+namespace ext {
+namespace oneapi {
+namespace experimental {
+
+namespace detail {
+
+// Type-trait for checking if a type defines `operator[]`.
+template <typename T>
+struct HasSubscriptOperator
+    : std::bool_constant<
+          !std::is_void<decltype(std::declval<T>().operator[](0))>::value>{};
+
+// Deduce a `properties<>` type from given variadic properties
+template <typename... Args> struct DeducedProperties {
+  using type = decltype(properties{std::declval<Args>()...});
+};
+
+// Partial specialization for deducing a `properties<>` type by forwarding the
+// given `properties<>` type
+template <typename... Args>
+struct DeducedProperties<detail::properties_t<Args...>> {
+  using type = detail::properties_t<Args...>;
+};
+
+} // namespace detail
+
+// Deduction guide
+template <typename T, typename... Args>
+annotated_arg(T, Args...)
+    -> annotated_arg<T, typename detail::DeducedProperties<Args...>::type>;
+
+template <typename T, typename old, typename... ArgT>
+annotated_arg(annotated_arg<T, old>, properties<std::tuple<ArgT...>>)
+    -> annotated_arg<
+        T, detail::merged_properties_t<old, detail::properties_t<ArgT...>>>;
+
+template <typename T, typename PropertyListT = detail::empty_properties_t>
+class annotated_arg {
+  // This should always fail when instantiating the unspecialized version.
+  static_assert(is_property_list<PropertyListT>::value,
+                "Property list is invalid.");
+};
+
+// Partial specialization for pointer type
+template <typename T, typename... Props>
+class __SYCL_SPECIAL_CLASS
+__SYCL_TYPE(annotated_arg) annotated_arg<T *, detail::properties_t<Props...>> {
+  using property_list_t = detail::properties_t<Props...>;
+
+#ifdef __SYCL_DEVICE_ONLY__
+  using global_pointer_t = typename decorated_global_ptr<T>::pointer;
+#else
+  using global_pointer_t = T *;
+#endif
+
+  global_pointer_t obj;
+
+  template <typename T2, typename PropertyListT> friend class annotated_arg;
+
+#ifdef __SYCL_DEVICE_ONLY__
+  void __init([[__sycl_detail__::add_ir_attributes_kernel_parameter(
+      detail::PropertyMetaInfo<Props>::name...,
+      detail::PropertyMetaInfo<Props>::value...)]] global_pointer_t _obj) {
+    obj = _obj;
+  }
+#endif
+
+public:
+  static_assert(is_property_list<property_list_t>::value,
+                "Property list is invalid.");
+
+  annotated_arg() noexcept = default;
+  annotated_arg(const annotated_arg &) = default;
+  annotated_arg &operator=(annotated_arg &) = default;
+
+  annotated_arg(T *_ptr,
+                const property_list_t &PropList = properties{}) noexcept
+      : obj(global_pointer_t(_ptr)) {}
+
+  // Constructs an annotated_arg object from a raw pointer and variadic
+  // properties. The new property set contains all properties of the input
+  // variadic properties. The same property in `Props...` and
+  // `PropertyValueTs...` must have the same property value.
+  template <typename... PropertyValueTs>
+  annotated_arg(T *_ptr, const PropertyValueTs &...props) noexcept
+      : obj(global_pointer_t(_ptr)) {
+    static_assert(
+        std::is_same<
+            property_list_t,
+            detail::merged_properties_t<property_list_t,
+                                        decltype(properties{props...})>>::value,
+        "The property list must contain all properties of the input of the "
+        "constructor");
+  }
+
+  // Constructs an annotated_arg object from another annotated_arg object.
+  // The new property set contains all properties of the input
+  // annotated_arg object. The same property in `Props...` and `PropertyList2`
+  // must have the same property value.
+  template <typename T2, typename PropertyList2>
+  explicit annotated_arg(const annotated_arg<T2, PropertyList2> &other) noexcept
+      : obj(other.obj) {
+    static_assert(std::is_convertible<T2, T *>::value,
+                  "The underlying data type of the input annotated_arg is not "
+                  "compatible");
+
+    static_assert(
+        std::is_same<
+            property_list_t,
+            detail::merged_properties_t<property_list_t, PropertyList2>>::value,
+        "The constructed annotated_arg type must contain all the properties of "
+        "the input annotated_arg");
+  }
+
+  // Constructs an annotated_arg object from another annotated_arg object and a
+  // property list. The new property set is the union of property lists
+  // `PropertyListU` and `PropertyListV`. The same property in `PropertyListU`
+  // and `PropertyListV` must have the same property value.
+  template <typename T2, typename PropertyListU, typename PropertyListV>
+  explicit annotated_arg(const annotated_arg<T2, PropertyListU> &other,
+                         const PropertyListV &proplist) noexcept
+      : obj(other.obj) {
+    static_assert(std::is_convertible<T2, T *>::value,
+                  "The underlying data type of the input annotated_arg is not "
+                  "compatible");
+
+    static_assert(
+        std::is_same<property_list_t, detail::merged_properties_t<
+                                          PropertyListU, PropertyListV>>::value,
+        "The property list of constructed annotated_arg type must be the union "
+        "of the input property lists");
+  }
+
+  operator T *() noexcept { return obj; }
+  operator T *() const noexcept { return obj; }
+
+  T &operator[](std::ptrdiff_t idx) const noexcept { return obj[idx]; }
+
+  template <typename PropertyT> static constexpr bool has_property() {
+    return property_list_t::template has_property<PropertyT>();
+  }
+
+  template <typename PropertyT> static constexpr auto get_property() {
+    return property_list_t::template get_property<PropertyT>();
+  }
+};
+
+// Partial specialization for non-pointer type
+template <typename T, typename... Props>
+class __SYCL_SPECIAL_CLASS
+__SYCL_TYPE(annotated_arg) annotated_arg<T, detail::properties_t<Props...>> {
+  using property_list_t = detail::properties_t<Props...>;
+
+  template <typename T2, typename PropertyListT> friend class annotated_arg;
+
+  T obj;
+
+#ifdef __SYCL_DEVICE_ONLY__
+  void __init([[__sycl_detail__::add_ir_attributes_kernel_parameter(
+      detail::PropertyMetaInfo<Props>::name...,
+      detail::PropertyMetaInfo<Props>::value...)]] T _obj) {
+    obj = _obj;
+  }
+#endif
+
+public:
+  static_assert(is_device_copyable_v<T>, "Type T must be device copyable.");
+  static_assert(is_property_list<property_list_t>::value,
+                "Property list is invalid.");
+  static_assert(check_property_list<T, Props...>::value,
+                "The property list contains invalid property.");
+
+  annotated_arg() noexcept = default;
+  annotated_arg(const annotated_arg &) = default;
+  annotated_arg &operator=(annotated_arg &) = default;
+
+  annotated_arg(const T &_obj,
+                const property_list_t &PropList = properties{}) noexcept
+      : obj(_obj) {}
+
+  // Constructs an annotated_arg object from a raw pointer and variadic
+  // properties. The new property set contains all properties of the input
+  // variadic properties. The same property in `Props...` and
+  // `PropertyValueTs...` must have the same property value.
+  template <typename... PropertyValueTs>
+  annotated_arg(const T &_obj, PropertyValueTs... props) noexcept : obj(_obj) {
+    static_assert(
+        std::is_same<
+            property_list_t,
+            detail::merged_properties_t<property_list_t,
+                                        decltype(properties{props...})>>::value,
+        "The property list must contain all properties of the input of the "
+        "constructor");
+  }
+
+  // Constructs an annotated_arg object from another annotated_arg object.
+  // The new property set contains all properties of the input
+  // annotated_arg object. The same property in `Props...` and `PropertyList2`
+  // must have the same property value.
+  template <typename T2, typename PropertyList2>
+  explicit annotated_arg(const annotated_arg<T2, PropertyList2> &other) noexcept
+      : obj(other.obj) {
+    static_assert(std::is_convertible<T2, T>::value,
+                  "The underlying data type of the input annotated_arg is not "
+                  "compatible");
+
+    static_assert(
+        std::is_same<
+            property_list_t,
+            detail::merged_properties_t<property_list_t, PropertyList2>>::value,
+        "The constructed annotated_arg type must contain all the properties of "
+        "the input annotated_arg");
+  }
+
+  // Constructs an annotated_arg object from another annotated_arg object and a
+  // property list. The new property set is the union of property lists
+  // `PropertyListU` and `PropertyListV`. The same property in `PropertyListU`
+  // and `PropertyListV` must have the same property value.
+  template <typename T2, typename PropertyListU, typename PropertyListV>
+  explicit annotated_arg(const annotated_arg<T2, PropertyListU> &other,
+                         const PropertyListV &proplist) noexcept
+      : obj(other.obj) {
+    static_assert(std::is_convertible<T2, T>::value,
+                  "The underlying data type of the input annotated_arg is not "
+                  "compatible");
+
+    static_assert(
+        std::is_same<property_list_t, detail::merged_properties_t<
+                                          PropertyListU, PropertyListV>>::value,
+        "The property list of constructed annotated_arg type must be the union "
+        "of the input property lists");
+  }
+
+  operator T() noexcept { return obj; }
+  operator T() const noexcept { return obj; }
+
+  template <class RelayT = T>
+  std::enable_if_t<detail::HasSubscriptOperator<RelayT>::value,
+                   decltype(std::declval<RelayT>().operator[](0))> &
+  operator[](std::ptrdiff_t idx) const noexcept {
+    return obj.operator[](idx);
+  }
+
+  template <typename PropertyT> static constexpr bool has_property() {
+    return property_list_t::template has_property<PropertyT>();
+  }
+
+  template <typename PropertyT> static constexpr auto get_property() {
+    return property_list_t::template get_property<PropertyT>();
+  }
+};
+
+} // namespace experimental
+} // namespace oneapi
+} // namespace ext
+} // __SYCL_INLINE_VER_NAMESPACE(_V1)
+} // namespace sycl
diff --git a/sycl/include/sycl/ext/oneapi/annotated_arg/properties.hpp b/sycl/include/sycl/ext/oneapi/annotated_arg/properties.hpp
new file mode 100644
index 0000000000000..eb21a01e6d229
--- /dev/null
+++ b/sycl/include/sycl/ext/oneapi/annotated_arg/properties.hpp
@@ -0,0 +1,304 @@
+//==-- properties.hpp - SYCL properties associated with annotated_arg/ptr --==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+#include <sycl/ext/oneapi/properties/property.hpp>
+#include <sycl/ext/oneapi/properties/property_value.hpp>
+
+namespace sycl {
+__SYCL_INLINE_VER_NAMESPACE(_V1) {
+namespace ext {
+namespace oneapi {
+namespace experimental {
+
+template <typename T, typename PropertyListT> class annotated_arg;
+
+//===----------------------------------------------------------------------===//
+//        Common properties of annotated_arg/annotated_ptr
+//===----------------------------------------------------------------------===//
+struct register_map_key {
+  using value_t = property_value<register_map_key>;
+};
+
+struct conduit_key {
+  using value_t = property_value<conduit_key>;
+};
+
+struct stable_key {
+  using value_t = property_value<stable_key>;
+};
+
+struct buffer_location_key {
+  template <int K>
+  using value_t =
+      property_value<buffer_location_key, std::integral_constant<int, K>>;
+};
+
+struct awidth_key {
+  template <int K>
+  using value_t = property_value<awidth_key, std::integral_constant<int, K>>;
+};
+
+struct dwidth_key {
+  template <int K>
+  using value_t = property_value<dwidth_key, std::integral_constant<int, K>>;
+};
+
+struct latency_key {
+  template <int K>
+  using value_t = property_value<latency_key, std::integral_constant<int, K>>;
+};
+
+enum class read_write_mode_enum : std::uint16_t { read_write, read, write };
+
+struct read_write_mode_key {
+  template <read_write_mode_enum Mode>
+  using value_t =
+      property_value<read_write_mode_key,
+                     std::integral_constant<read_write_mode_enum, Mode>>;
+};
+
+struct maxburst_key {
+  template <int K>
+  using value_t = property_value<maxburst_key, std::integral_constant<int, K>>;
+};
+
+struct wait_request_key {
+  template <int K>
+  using value_t =
+      property_value<wait_request_key, std::integral_constant<int, K>>;
+};
+
+// non-mmhost properties
+inline constexpr register_map_key::value_t register_map;
+inline constexpr conduit_key::value_t conduit;
+inline constexpr stable_key::value_t stable;
+
+// mmhost properties
+template <int N>
+inline constexpr buffer_location_key::value_t<N> buffer_location;
+template <int W> inline constexpr awidth_key::value_t<W> awidth;
+template <int W> inline constexpr dwidth_key::value_t<W> dwidth;
+template <int N> inline constexpr latency_key::value_t<N> latency;
+template <int N> inline constexpr maxburst_key::value_t<N> maxburst;
+template <int Enable>
+inline constexpr wait_request_key::value_t<Enable> wait_request;
+inline constexpr wait_request_key::value_t<1> wait_request_requested;
+inline constexpr wait_request_key::value_t<0> wait_request_not_requested;
+
+template <read_write_mode_enum Mode>
+inline constexpr read_write_mode_key::value_t<Mode> read_write_mode;
+inline constexpr read_write_mode_key::value_t<read_write_mode_enum::read>
+    read_write_mode_read;
+inline constexpr read_write_mode_key::value_t<read_write_mode_enum::write>
+    read_write_mode_write;
+inline constexpr read_write_mode_key::value_t<read_write_mode_enum::read_write>
+    read_write_mode_readwrite;
+
+template <> struct is_property_key<register_map_key> : std::true_type {};
+template <> struct is_property_key<conduit_key> : std::true_type {};
+template <> struct is_property_key<stable_key> : std::true_type {};
+
+template <> struct is_property_key<buffer_location_key> : std::true_type {};
+template <> struct is_property_key<awidth_key> : std::true_type {};
+template <> struct is_property_key<dwidth_key> : std::true_type {};
+template <> struct is_property_key<latency_key> : std::true_type {};
+template <> struct is_property_key<read_write_mode_key> : std::true_type {};
+template <> struct is_property_key<maxburst_key> : std::true_type {};
+template <> struct is_property_key<wait_request_key> : std::true_type {};
+
+template <typename T, typename PropertyListT>
+struct is_property_key_of<register_map_key, annotated_arg<T, PropertyListT>>
+    : std::true_type {};
+
+template <typename T, typename PropertyListT>
+struct is_property_key_of<conduit_key, annotated_arg<T, PropertyListT>>
+    : std::true_type {};
+
+template <typename T, typename PropertyListT>
+struct is_property_key_of<stable_key, annotated_arg<T, PropertyListT>>
+    : std::true_type {};
+
+template <typename T, typename PropertyListT>
+struct is_property_key_of<buffer_location_key, annotated_arg<T, PropertyListT>>
+    : std::true_type {};
+
+template <typename T, typename PropertyListT>
+struct is_property_key_of<awidth_key, annotated_arg<T, PropertyListT>>
+    : std::true_type {};
+
+template <typename T, typename PropertyListT>
+struct is_property_key_of<dwidth_key, annotated_arg<T, PropertyListT>>
+    : std::true_type {};
+
+template <typename T, typename PropertyListT>
+struct is_property_key_of<latency_key, annotated_arg<T, PropertyListT>>
+    : std::true_type {};
+
+template <typename T, typename PropertyListT>
+struct is_property_key_of<read_write_mode_key, annotated_arg<T, PropertyListT>>
+    : std::true_type {};
+
+template <typename T, typename PropertyListT>
+struct is_property_key_of<maxburst_key, annotated_arg<T, PropertyListT>>
+    : std::true_type {};
+
+template <typename T, typename PropertyListT>
+struct is_property_key_of<wait_request_key, annotated_arg<T, PropertyListT>>
+    : std::true_type {};
+
+namespace detail {
+template <> struct PropertyToKind<register_map_key> {
+  static constexpr PropKind Kind = PropKind::RegisterMap;
+};
+template <> struct PropertyToKind<conduit_key> {
+  static constexpr PropKind Kind = PropKind::Conduit;
+};
+template <> struct PropertyToKind<stable_key> {
+  static constexpr PropKind Kind = PropKind::Stable;
+};
+template <> struct PropertyToKind<buffer_location_key> {
+  static constexpr PropKind Kind = PropKind::BufferLocation;
+};
+template <> struct PropertyToKind<awidth_key> {
+  static constexpr PropKind Kind = PropKind::AddrWidth;
+};
+template <> struct PropertyToKind<dwidth_key> {
+  static constexpr PropKind Kind = PropKind::DataWidth;
+};
+template <> struct PropertyToKind<latency_key> {
+  static constexpr PropKind Kind = PropKind::Latency;
+};
+template <> struct PropertyToKind<read_write_mode_key> {
+  static constexpr PropKind Kind = PropKind::RWMode;
+};
+template <> struct PropertyToKind<maxburst_key> {
+  static constexpr PropKind Kind = PropKind::MaxBurst;
+};
+template <> struct PropertyToKind<wait_request_key> {
+  static constexpr PropKind Kind = PropKind::WaitRequest;
+};
+
+template <> struct IsCompileTimeProperty<register_map_key> : std::true_type {};
+template <> struct IsCompileTimeProperty<conduit_key> : std::true_type {};
+template <> struct IsCompileTimeProperty<stable_key> : std::true_type {};
+
+template <>
+struct IsCompileTimeProperty<buffer_location_key> : std::true_type {};
+template <> struct IsCompileTimeProperty<awidth_key> : std::true_type {};
+template <> struct IsCompileTimeProperty<dwidth_key> : std::true_type {};
+template <>
+struct IsCompileTimeProperty<read_write_mode_key> : std::true_type {};
+template <> struct IsCompileTimeProperty<latency_key> : std::true_type {};
+template <> struct IsCompileTimeProperty<maxburst_key> : std::true_type {};
+template <> struct IsCompileTimeProperty<wait_request_key> : std::true_type {};
+
+template <> struct PropertyMetaInfo<register_map_key::value_t> {
+  static constexpr const char *name = "sycl-register-map";
+  static constexpr std::nullptr_t value = nullptr;
+};
+template <> struct PropertyMetaInfo<conduit_key::value_t> {
+  static constexpr const char *name = "sycl-conduit";
+  static constexpr std::nullptr_t value = nullptr;
+};
+template <> struct PropertyMetaInfo<stable_key::value_t> {
+  static constexpr const char *name = "sycl-stable";
+  static constexpr std::nullptr_t value = nullptr;
+};
+
+template <int N> struct PropertyMetaInfo<buffer_location_key::value_t<N>> {
+  static constexpr const char *name = "sycl-buffer-location";
+  static constexpr int value = N;
+};
+template <int W> struct PropertyMetaInfo<awidth_key::value_t<W>> {
+  static constexpr const char *name = "sycl-awidth";
+  static constexpr int value = W;
+};
+template <int W> struct PropertyMetaInfo<dwidth_key::value_t<W>> {
+  static constexpr const char *name = "sycl-dwidth";
+  static constexpr int value = W;
+};
+template <int N> struct PropertyMetaInfo<latency_key::value_t<N>> {
+  static constexpr const char *name = "sycl-latency";
+  static constexpr int value = N;
+};
+template <int N> struct PropertyMetaInfo<maxburst_key::value_t<N>> {
+  static constexpr const char *name = "sycl-maxburst";
+  static constexpr int value = N;
+};
+template <int Enable>
+struct PropertyMetaInfo<wait_request_key::value_t<Enable>> {
+  static constexpr const char *name = "sycl-wait-request";
+  static constexpr int value = Enable;
+};
+template <read_write_mode_enum Mode>
+struct PropertyMetaInfo<read_write_mode_key::value_t<Mode>> {
+  static constexpr const char *name = "sycl-read-write-mode";
+  static constexpr read_write_mode_enum value = Mode;
+};
+
+} // namespace detail
+
+// 'buffer_location' and mmhost properties are pointers-only
+template <typename T, typename PropertyValueT>
+struct is_valid_property : std::false_type {};
+
+template <typename T, int N>
+struct is_valid_property<T, buffer_location_key::value_t<N>>
+    : std::bool_constant<std::is_pointer<T>::value> {};
+
+template <typename T, int W>
+struct is_valid_property<T, awidth_key::value_t<W>>
+    : std::bool_constant<std::is_pointer<T>::value> {};
+
+template <typename T, int W>
+struct is_valid_property<T, dwidth_key::value_t<W>>
+    : std::bool_constant<std::is_pointer<T>::value> {};
+
+template <typename T, int N>
+struct is_valid_property<T, latency_key::value_t<N>>
+    : std::bool_constant<std::is_pointer<T>::value> {};
+
+template <typename T, read_write_mode_enum Mode>
+struct is_valid_property<T, read_write_mode_key::value_t<Mode>>
+    : std::bool_constant<std::is_pointer<T>::value> {};
+
+template <typename T, int N>
+struct is_valid_property<T, maxburst_key::value_t<N>>
+    : std::bool_constant<std::is_pointer<T>::value> {};
+
+template <typename T, int Enable>
+struct is_valid_property<T, wait_request_key::value_t<Enable>>
+    : std::bool_constant<std::is_pointer<T>::value> {};
+
+// 'register_map',  'conduit',  'stable' are common properties for pointers
+// and non pointers;
+template <typename T>
+struct is_valid_property<T, register_map_key::value_t> : std::true_type {};
+template <typename T>
+struct is_valid_property<T, conduit_key::value_t> : std::true_type {};
+template <typename T>
+struct is_valid_property<T, stable_key::value_t> : std::true_type {};
+
+template <typename T, typename... Props>
+struct check_property_list : std::true_type {};
+
+template <typename T, typename Prop, typename... Props>
+struct check_property_list<T, Prop, Props...>
+    : std::conditional_t<is_valid_property<T, Prop>::value,
+                         check_property_list<T, Props...>, std::false_type> {
+  static_assert(is_valid_property<T, Prop>::value,
+                "Property is invalid for the given type.");
+};
+
+} // namespace experimental
+} // namespace oneapi
+} // namespace ext
+} // __SYCL_INLINE_VER_NAMESPACE(_V1)
+} // namespace sycl
diff --git a/sycl/include/sycl/ext/oneapi/properties/property.hpp b/sycl/include/sycl/ext/oneapi/properties/property.hpp
index 01df3abf6c296..08f6d41ecffbf 100644
--- a/sycl/include/sycl/ext/oneapi/properties/property.hpp
+++ b/sycl/include/sycl/ext/oneapi/properties/property.hpp
@@ -173,11 +173,21 @@ enum PropKind : uint32_t {
   WorkGroupSizeHint = 7,
   SubGroupSize = 8,
   DeviceHas = 9,
-  StreamingInterface = 10,
+  StreamingInterface = 10, // kernel attribute
   RegisterMapInterface = 11,
   Pipelined = 12,
+  RegisterMap = 13, // kernel argument attribute
+  Conduit = 14,
+  Stable = 15,
+  BufferLocation = 16,
+  AddrWidth = 17,
+  DataWidth = 18,
+  Latency = 19,
+  RWMode = 20,
+  MaxBurst = 21,
+  WaitRequest = 22,
   // PropKindSize must always be the last value.
-  PropKindSize = 13,
+  PropKindSize = 23,
 };
 
 // This trait must be specialized for all properties and must have a unique
diff --git a/sycl/include/sycl/sycl.hpp b/sycl/include/sycl/sycl.hpp
index 8b1be0aabc3b1..f9396b7a05b67 100644
--- a/sycl/include/sycl/sycl.hpp
+++ b/sycl/include/sycl/sycl.hpp
@@ -62,6 +62,8 @@
 #endif
 #include <sycl/ext/codeplay/experimental/fusion_wrapper.hpp>
 #include <sycl/ext/intel/usm_pointers.hpp>
+#include <sycl/ext/oneapi/annotated_arg/annotated_arg.hpp>
+#include <sycl/ext/oneapi/annotated_arg/properties.hpp>
 #include <sycl/ext/oneapi/device_global/device_global.hpp>
 #include <sycl/ext/oneapi/device_global/properties.hpp>
 #include <sycl/ext/oneapi/experimental/bfloat16_math.hpp>
diff --git a/sycl/test/extensions/annotated_arg/annotated_arg_for_non_ptr.cpp b/sycl/test/extensions/annotated_arg/annotated_arg_for_non_ptr.cpp
new file mode 100644
index 0000000000000..9d3d535bc408e
--- /dev/null
+++ b/sycl/test/extensions/annotated_arg/annotated_arg_for_non_ptr.cpp
@@ -0,0 +1,161 @@
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -fsyntax-only -Xclang -verify -Xclang -verify-ignore-unexpected=note %s
+// expected-no-diagnostics
+
+#include "sycl/sycl.hpp"
+#include <sycl/ext/intel/fpga_extensions.hpp>
+
+#include <iostream>
+
+// clang-format on
+
+using namespace sycl;
+using namespace ext::oneapi::experimental;
+
+struct B {};
+
+struct A : public B {
+  int x;
+  A() {}
+  A(int x_) : x(x_) {}
+
+  int &operator[](std::ptrdiff_t idx) { return x; }
+  int &operator[](std::ptrdiff_t idx) const { return const_cast<int &>(x); }
+};
+
+using annotated_arg_t1 =
+    annotated_arg<A, decltype(properties(conduit, stable))>;
+
+using annotated_arg_t3 = annotated_arg<A, decltype(properties(conduit))>;
+
+struct MyIP {
+  int *a;
+
+  annotated_arg_t1 b;
+
+  MyIP(int *a_, const A &b_) : a(a_), b(b_) {}
+
+  void operator()() const {
+    b[0] = 10;
+
+    const A &tmp = b;
+    A tmp2 = b;
+
+    for (int i = 0; i < tmp.x; i++) {
+      *a += 1;
+    }
+    *a += b[0];
+  }
+};
+
+template <typename T> T foo() {
+  A obj(5);
+  return annotated_arg(obj, conduit);
+}
+
+void TestVectorAddWithAnnotatedMMHosts() {
+  // Create the SYCL device queue
+  queue q(sycl::ext::intel::fpga_selector_v);
+  auto raw = malloc_shared<int>(1, q);
+
+  A obj(0);
+  // default ctor
+  annotated_arg_t3 a1(obj);
+
+  // copy ctor
+  auto a2(a1);
+  auto a3(foo<annotated_arg_t3>());
+  // // assign ctor
+  auto a4 = a3;
+
+  // Construct from A instance
+  auto tmp11 = annotated_arg(obj); // empty property list
+
+  // Construct from A instance and a property list
+  // auto tmp12 = annotated_arg(obj, properties{conduit});
+  auto tmp12 = annotated_arg(obj, conduit);
+  static_assert(std::is_same<decltype(tmp12), annotated_arg_t3>::value,
+                "deduction guide failed 1");
+
+  // Construct from A instance and variadic properties
+  auto tmp13 = annotated_arg(obj, stable, conduit); // deduction guide
+  static_assert(std::is_same<decltype(tmp13), annotated_arg_t1>::value,
+                "deduction guide failed 2");
+
+  // property list contains invalid property
+  // auto tmp14 = annotated_arg(obj, awidth<32>);  // ERR
+
+  // Construct from another annotated_arg
+  // templated copy constructor
+  annotated_arg<A, decltype(properties{conduit, stable})> arg11(tmp11);
+  annotated_arg<B, decltype(properties{stable, conduit})> arg14(
+      tmp11); // convertible type
+  auto arg12 = annotated_arg<A, decltype(properties{stable, conduit})>(tmp11);
+
+  // default copy constructor
+  auto arg13 = annotated_arg(tmp12);
+  static_assert(std::is_same<decltype(arg13), annotated_arg_t3>::value,
+                "deduction guide failed 3");
+
+  // Construct from another annotated_arg and a property list
+  // annotated_arg<A, decltype(properties{conduit, stable})> arg21(tmp11,
+  // properties{stable});   // ERR:  the type properties should be the union of
+  // the inputs
+  annotated_arg<A, decltype(properties{conduit, stable})> arg22(
+      tmp12, properties{stable});
+  auto arg23 = annotated_arg(tmp12, properties{stable}); // deduction guide
+  static_assert(std::is_same<decltype(arg22), annotated_arg_t1>::value,
+                "deduction guide failed 4");
+  static_assert(std::is_same<decltype(arg23), decltype(arg22)>::value,
+                "deduction guide failed 5");
+  annotated_arg<B, decltype(properties{stable, conduit})> arg24(
+      tmp12, properties{stable}); // convertible type
+
+  // Property merge
+  auto arg31 = annotated_arg_t3(obj, conduit);                    // OK
+  auto arg32 = annotated_arg(arg31, properties{stable});          // OK
+  auto arg33 = annotated_arg(arg32, properties{stable, conduit}); // OK
+  // auto arg34 = annotated_arg(arg32, properties{conduit, latency<22>});  //
+  // ERR: invalid property
+  static_assert(std::is_same<decltype(arg32), annotated_arg_t1>::value,
+                "deduction guide failed 6");
+  static_assert(std::is_same<decltype(arg33), annotated_arg_t1>::value,
+                "deduction guide failed 7");
+  // auto arg35 = annotated_arg(arg32, properties{conduit, dwidth<22>});  //
+  // ERR: two input property lists are conflict
+  // annotated_arg<A, decltype(properties{conduit, stable})>
+  //    arg36(arg31, properties{latency<32>, stable}); // ERR: input
+  // property list is conflict with the declared type
+
+  // Implicit Conversion
+  const A &x13 = arg32; // OK
+  A x14 = arg32;        // OK
+  // A& x11 = arg32;   // ERR: non-const lvalue reference to type 'A' cannot
+  // bind to a value of unrelated type
+
+  // operator[]
+  a1[0] = 5;
+
+  // has/get property
+  static_assert(annotated_arg_t1::has_property<conduit_key>(),
+                "has property 1");
+  static_assert(annotated_arg_t1::get_property<conduit_key>() == conduit,
+                "get property 1");
+  static_assert(annotated_arg_t1::has_property<latency_key>() == false,
+                "has property 2");
+
+  static_assert(annotated_arg_t3::has_property<stable_key>() == false,
+                "has property 3");
+  // auto stable_prop = annotated_arg_t3::get_property<stable_key>();   // ERR:
+  // can't get non-existing property
+
+  *raw = 0;
+  q.submit([&](handler &h) { h.single_task(MyIP{raw, a1}); }).wait();
+
+  std::cout << raw[0] << std::endl;
+  free(raw, q);
+}
+
+int main() {
+  TestVectorAddWithAnnotatedMMHosts();
+  return 0;
+}
diff --git a/sycl/test/extensions/annotated_arg/annotated_arg_for_ptr.cpp b/sycl/test/extensions/annotated_arg/annotated_arg_for_ptr.cpp
new file mode 100644
index 0000000000000..12549bbbe4ebf
--- /dev/null
+++ b/sycl/test/extensions/annotated_arg/annotated_arg_for_ptr.cpp
@@ -0,0 +1,159 @@
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -fsyntax-only -Xclang -verify -Xclang -verify-ignore-unexpected=note %s
+// expected-no-diagnostics
+
+#include "sycl/sycl.hpp"
+#include <sycl/ext/intel/fpga_extensions.hpp>
+
+#include <iostream>
+
+// clang-format on
+
+using namespace sycl;
+using namespace ext::oneapi::experimental;
+
+using annotated_arg_t1 =
+    annotated_arg<int *, decltype(properties(awidth<32>, dwidth<32>))>;
+
+using annotated_arg_t2 =
+    annotated_arg<int, decltype(properties(conduit, register_map))>;
+
+using annotated_arg_t3 = annotated_arg<int *, decltype(properties(awidth<32>))>;
+
+struct MyIP {
+  annotated_arg<int *, decltype(properties(awidth<32>, dwidth<32>))> a;
+
+  int b;
+
+  MyIP(int *a_, int b_) : a(a_), b(b_) {}
+
+  void operator()() const {
+    int *p = a;
+    const int *p2 = a;
+
+    for (int i = 0; i < b; i++) {
+      p[i] = i;
+      a[i] += 1;
+    }
+  }
+};
+
+template <typename T> T foo() {
+  auto raw = new int;
+  return annotated_arg(raw, awidth<32>);
+}
+
+void TestVectorAddWithAnnotatedMMHosts() {
+  // Create the SYCL device queue
+  queue q(sycl::ext::intel::fpga_selector_v);
+  auto raw = malloc_shared<int>(5, q);
+
+  // default ctor
+  annotated_arg_t3 a1;
+  // copy ctor
+  auto a2(a1);
+  auto a3(foo<annotated_arg_t3>());
+  // // assign ctor
+  auto a4 = a1;
+
+  // Construct from raw pointers
+  auto tmp11 = annotated_arg(raw); // empty property list
+  // Construct from raw pointers and a property list
+  auto tmp12 = annotated_arg<int *, decltype(properties{awidth<32>})>(
+      raw, properties{awidth<32>});
+  auto tmp14 = annotated_arg(raw, properties{awidth<32>}); // deduction guide
+  static_assert(std::is_same<decltype(tmp14), annotated_arg_t3>::value,
+                "deduction guide failed 1");
+  // Construct from raw pointers and variadic properties
+  auto tmp13 = annotated_arg(raw, dwidth<32>, awidth<32>); // deduction guide
+  static_assert(std::is_same<decltype(tmp13), annotated_arg_t1>::value,
+                "deduction guide failed 2");
+  auto tmp15 = annotated_arg(raw, awidth<32>);
+  static_assert(std::is_same<decltype(tmp15), annotated_arg_t3>::value,
+                "deduction guide failed 1");
+
+  // Property list can't have duplicated properties
+  // auto tmp16 = annotated_arg(raw, awidth<32>, awidth<32>);   // ERR
+  // auto tmp17 = annotated_arg(raw, awidth<32>, awidth<22>);  // ERR
+
+  // auto tmp18 = annotated_arg(raw, properties{awidth<32>, dwidth<32>,
+  // awidth<32>}); // ERR: Duplicate properties in property list auto tmp19 =
+  // annotated_arg(raw, properties{awidth<32>, awidth<22>});  // ERR
+
+  // Construct from another annotated_arg
+  // templated copy constructor
+  annotated_arg<int *, decltype(properties{awidth<32>, dwidth<32>})> arg11(
+      tmp11);
+  auto arg12 =
+      annotated_arg<int *, decltype(properties{dwidth<32>, awidth<32>})>(tmp11);
+
+  // default copy constructor
+  auto arg13 = annotated_arg(tmp12);
+  static_assert(std::is_same<decltype(arg13), annotated_arg_t3>::value,
+                "deduction guide failed 3");
+
+  // Construct from another annotated_arg and a property list
+  // annotated_arg<int*, decltype(properties{awidth<32>, dwidth<32>})>
+  // arg21(tmp11, properties{dwidth<32>});   // ERR:  the type properties should
+  // be the union of the inputs
+  annotated_arg<int *, decltype(properties{awidth<32>, dwidth<32>})> arg22(
+      tmp12, properties{dwidth<32>});
+  auto arg23 = annotated_arg(tmp12, properties{dwidth<32>}); // deduction guide
+  static_assert(std::is_same<decltype(arg22), annotated_arg_t1>::value,
+                "deduction guide failed 4");
+  static_assert(std::is_same<decltype(arg23), decltype(arg22)>::value,
+                "deduction guide failed 5");
+
+  // Construct from inconvertible type
+  // annotated_arg<int> tmp21;
+  // annotated_arg<int*, decltype(properties{dwidth<32>})> arg24(tmp21,
+  // properties{dwidth<32>});   // ERR
+
+  // Property merge
+  auto arg31 = annotated_arg_t3(raw, awidth<32>);                         // OK
+  auto arg32 = annotated_arg(arg31, properties{dwidth<32>});              // OK
+  auto arg33 = annotated_arg(arg32, properties{dwidth<32>, awidth<32>});  // OK
+  auto arg34 = annotated_arg(arg32, properties{awidth<32>, latency<22>}); // OK
+  static_assert(std::is_same<decltype(arg32), annotated_arg_t1>::value,
+                "deduction guide failed 6");
+  static_assert(std::is_same<decltype(arg33), annotated_arg_t1>::value,
+                "deduction guide failed 7");
+  // auto arg34 = annotated_arg(arg32, properties{awidth<32>, dwidth<22>});  //
+  // ERR: two input property lists are conflict
+  // annotated_arg<int*, decltype(properties{awidth<32>, dwidth<32>})>
+  //    arg35(arg31, properties{latency<32>, dwidth<32>}); // ERR: input
+  // property list is conflict with the declared type
+
+  // Implicit Conversion
+  int *x11 = arg13;
+  const int *x13 = arg32;
+
+  // operator[]
+  arg31[0] = 1;
+  for (int i = 1; i < 5; i++) {
+    arg31[i] = arg31[i - 1];
+  }
+
+  // has/get property
+  static_assert(annotated_arg_t1::has_property<awidth_key>(), "has property 1");
+  static_assert(annotated_arg_t1::get_property<awidth_key>() == awidth<32>,
+                "get property 1");
+  static_assert(annotated_arg_t1::has_property<latency_key>() == false,
+                "has property 2");
+
+  static_assert(annotated_arg_t3::has_property<dwidth_key>() == false,
+                "has property 3");
+  // auto dwidth_prop = annotated_arg_t3::get_property<dwidth_key>();   // ERR
+
+  q.submit([&](handler &h) { h.single_task(MyIP{raw, 5}); }).wait();
+
+  for (int i = 0; i < 5; i++) {
+    std::cout << raw[i] << std::endl;
+  }
+
+  free(raw, q);
+}
+
+int main() {
+  TestVectorAddWithAnnotatedMMHosts();
+  return 0;
+}
diff --git a/sycl/test/extensions/annotated_arg/annotated_arg_properties.cpp b/sycl/test/extensions/annotated_arg/annotated_arg_properties.cpp
new file mode 100644
index 0000000000000..e622e67652808
--- /dev/null
+++ b/sycl/test/extensions/annotated_arg/annotated_arg_properties.cpp
@@ -0,0 +1,94 @@
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -fsyntax-only -Xclang -verify -Xclang -verify-ignore-unexpected=note %s
+// expected-no-diagnostics
+
+#include <sycl/sycl.hpp>
+
+using namespace sycl::ext::oneapi::experimental;
+
+static annotated_arg<int, decltype(properties())> AnnotatedArg1;
+static annotated_arg<int, decltype(properties(register_map))> AnnotatedArg2;
+static annotated_arg<int *, decltype(properties(conduit, stable))>
+    AnnotatedArg3;
+static annotated_arg<int *, decltype(properties(buffer_location<1>,
+                                                read_write_mode_read, stable,
+                                                conduit))>
+    AnnotatedArg4;
+
+struct A {};
+
+// Checks is_property_key_of and is_property_value_of for T.
+template <typename T> void checkIsPropertyOf() {
+  static_assert(is_property_key_of<register_map_key, T>::value);
+  static_assert(is_property_key_of<conduit_key, T>::value);
+  static_assert(is_property_key_of<stable_key, T>::value);
+
+  static_assert(is_property_key_of<buffer_location_key, T>::value);
+  static_assert(is_property_key_of<awidth_key, T>::value);
+  static_assert(is_property_key_of<dwidth_key, T>::value);
+  static_assert(is_property_key_of<latency_key, T>::value);
+  static_assert(is_property_key_of<read_write_mode_key, T>::value);
+  static_assert(is_property_key_of<maxburst_key, T>::value);
+  static_assert(is_property_key_of<wait_request_key, T>::value);
+
+  static_assert(is_property_value_of<decltype(register_map), T>::value);
+  static_assert(is_property_value_of<decltype(conduit), T>::value);
+  static_assert(is_property_value_of<decltype(stable), T>::value);
+
+  static_assert(is_property_value_of<decltype(buffer_location<1>), T>::value);
+  static_assert(is_property_value_of<decltype(awidth<2>), T>::value);
+  static_assert(is_property_value_of<decltype(dwidth<8>), T>::value);
+  static_assert(is_property_value_of<decltype(latency<0>), T>::value);
+  static_assert(is_property_value_of<decltype(read_write_mode_read), T>::value);
+  static_assert(is_property_value_of<decltype(maxburst<1>), T>::value);
+  static_assert(
+      is_property_value_of<decltype(wait_request_requested), T>::value);
+}
+
+// Checks is_property_key_of and is_property_value_of are false for non-pointer
+// type T.
+template <typename T> void checkIsValidPropertyOfNonPtr() {
+  static_assert(
+      is_valid_property<T, decltype(wait_request_not_requested)>::value ==
+      false);
+  static_assert(is_valid_property<T, decltype(latency<8>)>::value == false);
+}
+
+int main() {
+  static_assert(is_property_key<register_map_key>::value);
+  static_assert(is_property_key<buffer_location_key>::value);
+
+  checkIsPropertyOf<decltype(AnnotatedArg1)>();
+  static_assert(!AnnotatedArg1.has_property<register_map_key>());
+  static_assert(!AnnotatedArg1.has_property<buffer_location_key>());
+
+  checkIsPropertyOf<decltype(AnnotatedArg2)>();
+  static_assert(AnnotatedArg2.has_property<register_map_key>());
+  static_assert(!AnnotatedArg2.has_property<conduit_key>());
+  static_assert(!AnnotatedArg2.has_property<buffer_location_key>());
+  static_assert(AnnotatedArg2.get_property<register_map_key>() == register_map);
+
+  checkIsPropertyOf<decltype(AnnotatedArg3)>();
+  static_assert(!AnnotatedArg3.has_property<register_map_key>());
+  static_assert(AnnotatedArg3.has_property<conduit_key>());
+  static_assert(AnnotatedArg3.has_property<stable_key>());
+  static_assert(!AnnotatedArg3.has_property<buffer_location_key>());
+  static_assert(AnnotatedArg3.get_property<stable_key>() == stable);
+  static_assert(AnnotatedArg3.get_property<conduit_key>() == conduit);
+
+  checkIsPropertyOf<decltype(AnnotatedArg4)>();
+  static_assert(!AnnotatedArg4.has_property<register_map_key>());
+  static_assert(AnnotatedArg4.has_property<conduit_key>());
+  static_assert(AnnotatedArg4.has_property<stable_key>());
+  static_assert(AnnotatedArg4.has_property<buffer_location_key>());
+  static_assert(AnnotatedArg4.has_property<read_write_mode_key>());
+  static_assert(AnnotatedArg4.get_property<conduit_key>() == conduit);
+  static_assert(AnnotatedArg4.get_property<stable_key>() == stable);
+  static_assert(AnnotatedArg4.get_property<buffer_location_key>() ==
+                buffer_location<1>);
+  static_assert(AnnotatedArg4.get_property<read_write_mode_key>() ==
+                read_write_mode_read);
+
+  // Check if a property is valid for a given type
+  checkIsValidPropertyOfNonPtr<A>();
+  return 0;
+}