diff --git a/sycl/include/CL/sycl.hpp b/sycl/include/CL/sycl.hpp
index 8ca7d28223cad..f5d9aa76ff093 100644
--- a/sycl/include/CL/sycl.hpp
+++ b/sycl/include/CL/sycl.hpp
@@ -71,3 +71,5 @@
 #include <sycl/ext/oneapi/reduction.hpp>
 #include <sycl/ext/oneapi/sub_group.hpp>
 #include <sycl/ext/oneapi/sub_group_mask.hpp>
+
+#include <sycl/ext/intel/experimental/pipe_properties.hpp>
diff --git a/sycl/include/sycl/ext/intel/experimental/pipe_properties.hpp b/sycl/include/sycl/ext/intel/experimental/pipe_properties.hpp
new file mode 100644
index 0000000000000..e4e012ccd972d
--- /dev/null
+++ b/sycl/include/sycl/ext/intel/experimental/pipe_properties.hpp
@@ -0,0 +1,197 @@
+//==----- pipe_properties.hpp - SYCL properties associated with data flow pipe
+//---==//
+//
+// 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>
+
+__SYCL_INLINE_NAMESPACE(cl) {
+namespace sycl {
+namespace ext {
+namespace intel {
+namespace experimental {
+
+struct min_capacity_key {
+  template <int Capacity>
+  using value_t = oneapi::experimental::property_value<
+      min_capacity_key, std::integral_constant<int, Capacity>>;
+};
+
+struct ready_latency_key {
+  template <int Latency>
+  using value_t = oneapi::experimental::property_value<
+      ready_latency_key, std::integral_constant<int, Latency>>;
+};
+
+struct bits_per_symbol_key {
+  template <int Bits>
+  using value_t =
+      oneapi::experimental::property_value<bits_per_symbol_key,
+                                           std::integral_constant<int, Bits>>;
+};
+
+struct uses_valid_key {
+  template <bool Valid>
+  using value_t =
+      oneapi::experimental::property_value<uses_valid_key,
+                                           sycl::detail::bool_constant<Valid>>;
+};
+
+struct uses_ready_key {
+  template <bool Ready>
+  using value_t =
+      oneapi::experimental::property_value<uses_ready_key,
+                                           sycl::detail::bool_constant<Ready>>;
+};
+
+struct in_csr_key {
+  template <bool Enable>
+  using value_t =
+      oneapi::experimental::property_value<in_csr_key,
+                                           sycl::detail::bool_constant<Enable>>;
+};
+
+struct first_symbol_in_high_order_bits_key {
+  template <bool HighOrder>
+  using value_t = oneapi::experimental::property_value<
+      first_symbol_in_high_order_bits_key,
+      sycl::detail::bool_constant<HighOrder>>;
+};
+
+enum class protocol_name : std::uint16_t { AVALON, AXI };
+struct protocol_key {
+  template <protocol_name Protocol>
+  using value_t = oneapi::experimental::property_value<
+      protocol_key, std::integral_constant<protocol_name, Protocol>>;
+};
+
+template <int Capacity>
+inline constexpr min_capacity_key::value_t<Capacity> min_capacity;
+
+template <int Latency>
+inline constexpr ready_latency_key::value_t<Latency> ready_latency;
+
+template <int Bits>
+inline constexpr bits_per_symbol_key::value_t<Bits> bits_per_symbol;
+
+template <bool Valid>
+inline constexpr uses_valid_key::value_t<Valid> uses_valid;
+inline constexpr uses_valid_key::value_t<true> uses_valid_on;
+inline constexpr uses_valid_key::value_t<false> uses_valid_off;
+
+template <bool Ready>
+inline constexpr uses_ready_key::value_t<Ready> uses_ready;
+inline constexpr uses_ready_key::value_t<true> uses_ready_on;
+inline constexpr uses_ready_key::value_t<false> uses_ready_off;
+
+template <bool Enable> inline constexpr in_csr_key::value_t<Enable> in_csr;
+inline constexpr in_csr_key::value_t<true> in_csr_on;
+inline constexpr in_csr_key::value_t<false> in_csr_off;
+
+template <bool HighOrder>
+inline constexpr first_symbol_in_high_order_bits_key::value_t<HighOrder>
+    first_symbol_in_high_order_bits;
+inline constexpr first_symbol_in_high_order_bits_key::value_t<true>
+    first_symbol_in_high_order_bits_on;
+inline constexpr first_symbol_in_high_order_bits_key::value_t<false>
+    first_symbol_in_high_order_bits_off;
+
+template <protocol_name Protocol>
+inline constexpr protocol_key::value_t<Protocol> protocol;
+inline constexpr protocol_key::value_t<protocol_name::AVALON> protocol_avalon;
+inline constexpr protocol_key::value_t<protocol_name::AXI> protocol_axi;
+
+} // namespace experimental
+} // namespace intel
+
+namespace oneapi {
+namespace experimental {
+
+template <>
+struct is_property_key<intel::experimental::min_capacity_key> : std::true_type {
+};
+template <>
+struct is_property_key<intel::experimental::ready_latency_key>
+    : std::true_type {};
+template <>
+struct is_property_key<intel::experimental::bits_per_symbol_key>
+    : std::true_type {};
+template <>
+struct is_property_key<intel::experimental::uses_valid_key> : std::true_type {};
+template <>
+struct is_property_key<intel::experimental::uses_ready_key> : std::true_type {};
+template <>
+struct is_property_key<intel::experimental::in_csr_key> : std::true_type {};
+template <>
+struct is_property_key<intel::experimental::first_symbol_in_high_order_bits_key>
+    : std::true_type {};
+template <>
+struct is_property_key<intel::experimental::protocol_key> : std::true_type {};
+
+namespace detail {
+template <> struct PropertyToKind<intel::experimental::min_capacity_key> {
+  static constexpr PropKind Kind = PropKind::MinCapacity;
+};
+template <> struct PropertyToKind<intel::experimental::ready_latency_key> {
+  static constexpr PropKind Kind = PropKind::ReadyLatency;
+};
+template <> struct PropertyToKind<intel::experimental::bits_per_symbol_key> {
+  static constexpr PropKind Kind = PropKind::BitsPerSymbol;
+};
+template <> struct PropertyToKind<intel::experimental::uses_valid_key> {
+  static constexpr PropKind Kind = PropKind::UsesValid;
+};
+template <> struct PropertyToKind<intel::experimental::uses_ready_key> {
+  static constexpr PropKind Kind = PropKind::UsesReady;
+};
+template <> struct PropertyToKind<intel::experimental::in_csr_key> {
+  static constexpr PropKind Kind = PropKind::ImplementInCSR;
+};
+template <>
+struct PropertyToKind<
+    intel::experimental::first_symbol_in_high_order_bits_key> {
+  static constexpr PropKind Kind = PropKind::FirstSymbolInHigherOrderBit;
+};
+template <> struct PropertyToKind<intel::experimental::protocol_key> {
+  static constexpr PropKind Kind = PropKind::PipeProtocol;
+};
+
+template <>
+struct IsCompileTimeProperty<intel::experimental::min_capacity_key>
+    : std::true_type {};
+template <>
+struct IsCompileTimeProperty<intel::experimental::ready_latency_key>
+    : std::true_type {};
+template <>
+struct IsCompileTimeProperty<intel::experimental::bits_per_symbol_key>
+    : std::true_type {};
+template <>
+struct IsCompileTimeProperty<intel::experimental::uses_valid_key>
+    : std::true_type {};
+template <>
+struct IsCompileTimeProperty<intel::experimental::uses_ready_key>
+    : std::true_type {};
+template <>
+struct IsCompileTimeProperty<intel::experimental::in_csr_key> : std::true_type {
+};
+template <>
+struct IsCompileTimeProperty<
+    intel::experimental::first_symbol_in_high_order_bits_key> : std::true_type {
+};
+template <>
+struct IsCompileTimeProperty<intel::experimental::protocol_key>
+    : std::true_type {};
+
+} // namespace detail
+} // namespace experimental
+} // namespace oneapi
+} // namespace ext
+} // namespace sycl
+} // __SYCL_INLINE_NAMESPACE(cl)
diff --git a/sycl/include/sycl/ext/oneapi/properties/property.hpp b/sycl/include/sycl/ext/oneapi/properties/property.hpp
index 658fa6721ca20..bde7ccd85c6b1 100644
--- a/sycl/include/sycl/ext/oneapi/properties/property.hpp
+++ b/sycl/include/sycl/ext/oneapi/properties/property.hpp
@@ -148,7 +148,14 @@ enum PropKind : uint32_t {
   HostAccess = 1,
   InitMode = 2,
   ImplementInCSR = 3,
-  PropKindSize = 4,
+  BitsPerSymbol = 4,
+  FirstSymbolInHigherOrderBit = 5,
+  MinCapacity = 6,
+  PipeProtocol = 7,
+  ReadyLatency = 8,
+  UsesReady = 9,
+  UsesValid = 10,
+  PropKindSize = 11,
 };
 
 // This trait must be specialized for all properties and must have a unique
diff --git a/sycl/test/extensions/properties/properties_pipe.cpp b/sycl/test/extensions/properties/properties_pipe.cpp
new file mode 100644
index 0000000000000..9179d9f66c003
--- /dev/null
+++ b/sycl/test/extensions/properties/properties_pipe.cpp
@@ -0,0 +1,168 @@
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -fsyntax-only -Xclang -verify -Xclang -verify-ignore-unexpected=note,warning %s
+// expected-no-diagnostics
+
+#include <CL/sycl.hpp>
+
+#include <sycl/ext/intel/experimental/pipe_properties.hpp>
+
+using namespace sycl::ext;
+
+constexpr sycl::ext::intel::experimental::protocol_name TestProtocol =
+    sycl::ext::intel::experimental::protocol_name::AVALON;
+
+int main() {
+  // Check that is_property_key is correctly specialized.
+  static_assert(sycl::ext::oneapi::experimental::is_property_key<
+                sycl::ext::intel::experimental::min_capacity_key>::value);
+  static_assert(sycl::ext::oneapi::experimental::is_property_key<
+                sycl::ext::intel::experimental::ready_latency_key>::value);
+  static_assert(sycl::ext::oneapi::experimental::is_property_key<
+                sycl::ext::intel::experimental::bits_per_symbol_key>::value);
+  static_assert(sycl::ext::oneapi::experimental::is_property_key<
+                sycl::ext::intel::experimental::uses_valid_key>::value);
+  static_assert(sycl::ext::oneapi::experimental::is_property_key<
+                sycl::ext::intel::experimental::uses_ready_key>::value);
+  static_assert(sycl::ext::oneapi::experimental::is_property_key<
+                sycl::ext::intel::experimental::in_csr_key>::value);
+  static_assert(
+      sycl::ext::oneapi::experimental::is_property_key<
+          sycl::ext::intel::experimental::first_symbol_in_high_order_bits_key>::
+          value);
+  static_assert(sycl::ext::oneapi::experimental::is_property_key<
+                sycl::ext::intel::experimental::protocol_key>::value);
+
+  // Check that is_property_value is correctly specialized.
+  static_assert(
+      sycl::ext::oneapi::experimental::is_property_value<
+          decltype(sycl::ext::intel::experimental::min_capacity<3>)>::value);
+  static_assert(
+      sycl::ext::oneapi::experimental::is_property_value<
+          decltype(sycl::ext::intel::experimental::ready_latency<3>)>::value);
+  static_assert(
+      sycl::ext::oneapi::experimental::is_property_value<
+          decltype(sycl::ext::intel::experimental::bits_per_symbol<3>)>::value);
+
+  static_assert(
+      sycl::ext::oneapi::experimental::is_property_value<
+          decltype(sycl::ext::intel::experimental::uses_valid<true>)>::value);
+  static_assert(
+      sycl::ext::oneapi::experimental::is_property_value<
+          decltype(sycl::ext::intel::experimental::uses_valid_on)>::value);
+  static_assert(
+      sycl::ext::oneapi::experimental::is_property_value<
+          decltype(sycl::ext::intel::experimental::uses_valid_off)>::value);
+
+  static_assert(
+      sycl::ext::oneapi::experimental::is_property_value<
+          decltype(sycl::ext::intel::experimental::uses_ready<true>)>::value);
+  static_assert(
+      sycl::ext::oneapi::experimental::is_property_value<
+          decltype(sycl::ext::intel::experimental::uses_ready_on)>::value);
+  static_assert(
+      sycl::ext::oneapi::experimental::is_property_value<
+          decltype(sycl::ext::intel::experimental::uses_ready_off)>::value);
+
+  static_assert(sycl::ext::oneapi::experimental::is_property_value<
+                decltype(sycl::ext::intel::experimental::in_csr<true>)>::value);
+  static_assert(sycl::ext::oneapi::experimental::is_property_value<
+                decltype(sycl::ext::intel::experimental::in_csr_on)>::value);
+  static_assert(sycl::ext::oneapi::experimental::is_property_value<
+                decltype(sycl::ext::intel::experimental::in_csr_off)>::value);
+
+  static_assert(sycl::ext::oneapi::experimental::is_property_value<
+                decltype(sycl::ext::intel::experimental::
+                             first_symbol_in_high_order_bits<true>)>::value);
+  static_assert(sycl::ext::oneapi::experimental::is_property_value<
+                decltype(sycl::ext::intel::experimental::
+                             first_symbol_in_high_order_bits_on)>::value);
+  static_assert(sycl::ext::oneapi::experimental::is_property_value<
+                decltype(sycl::ext::intel::experimental::
+                             first_symbol_in_high_order_bits_off)>::value);
+
+  static_assert(
+      sycl::ext::oneapi::experimental::is_property_value<
+          decltype(sycl::ext::intel::experimental::protocol<TestProtocol>)>::
+          value);
+  static_assert(
+      sycl::ext::oneapi::experimental::is_property_value<
+          decltype(sycl::ext::intel::experimental::protocol_avalon)>::value);
+  static_assert(sycl::ext::oneapi::experimental::is_property_value<
+                decltype(sycl::ext::intel::experimental::protocol_axi)>::value);
+
+  // Checks that fully specialized properties are the same as the templated
+  // variants.
+  static_assert(std::is_same_v<
+                decltype(sycl::ext::intel::experimental::uses_valid_on),
+                decltype(sycl::ext::intel::experimental::uses_valid<true>)>);
+  static_assert(std::is_same_v<
+                decltype(sycl::ext::intel::experimental::uses_ready_off),
+                decltype(sycl::ext::intel::experimental::uses_ready<false>)>);
+  static_assert(
+      std::is_same_v<decltype(sycl::ext::intel::experimental::in_csr_on),
+                     decltype(sycl::ext::intel::experimental::in_csr<true>)>);
+  static_assert(
+      std::is_same_v<decltype(sycl::ext::intel::experimental::
+                                  first_symbol_in_high_order_bits_on),
+                     decltype(sycl::ext::intel::experimental::
+                                  first_symbol_in_high_order_bits<true>)>);
+  static_assert(
+      std::is_same_v<
+          decltype(sycl::ext::intel::experimental::protocol_avalon),
+          decltype(sycl::ext::intel::experimental::protocol<TestProtocol>)>);
+  static_assert(std::is_same_v<
+                decltype(sycl::ext::intel::experimental::protocol_axi),
+                decltype(sycl::ext::intel::experimental::protocol<
+                         sycl::ext::intel::experimental::protocol_name::AXI>)>);
+
+  // Check that property lists will accept the new properties.
+  using P = decltype(sycl::ext::oneapi::experimental::properties(
+      sycl::ext::intel::experimental::min_capacity<0>,
+      sycl::ext::intel::experimental::ready_latency<1>,
+      sycl::ext::intel::experimental::bits_per_symbol<2>,
+      sycl::ext::intel::experimental::uses_valid<true>,
+      sycl::ext::intel::experimental::uses_ready<false>,
+      sycl::ext::intel::experimental::in_csr<true>,
+      sycl::ext::intel::experimental::first_symbol_in_high_order_bits_off,
+      sycl::ext::intel::experimental::protocol_avalon));
+  static_assert(sycl::ext::oneapi::experimental::is_property_list_v<P>);
+  static_assert(
+      P::has_property<sycl::ext::intel::experimental::min_capacity_key>());
+  static_assert(
+      P::has_property<sycl::ext::intel::experimental::ready_latency_key>());
+  static_assert(
+      P::has_property<sycl::ext::intel::experimental::bits_per_symbol_key>());
+  static_assert(
+      P::has_property<sycl::ext::intel::experimental::uses_valid_key>());
+  static_assert(
+      P::has_property<sycl::ext::intel::experimental::uses_ready_key>());
+  static_assert(P::has_property<sycl::ext::intel::experimental::in_csr_key>());
+  static_assert(P::has_property<sycl::ext::intel::experimental::
+                                    first_symbol_in_high_order_bits_key>());
+  static_assert(
+      P::has_property<sycl::ext::intel::experimental::protocol_key>());
+
+  static_assert(
+      P::get_property<sycl::ext::intel::experimental::min_capacity_key>() ==
+      sycl::ext::intel::experimental::min_capacity<0>);
+  static_assert(
+      P::get_property<sycl::ext::intel::experimental::ready_latency_key>() ==
+      sycl::ext::intel::experimental::ready_latency<1>);
+  static_assert(
+      P::get_property<sycl::ext::intel::experimental::bits_per_symbol_key>() ==
+      sycl::ext::intel::experimental::bits_per_symbol<2>);
+  static_assert(
+      P::get_property<sycl::ext::intel::experimental::uses_valid_key>() ==
+      sycl::ext::intel::experimental::uses_valid<true>);
+  static_assert(
+      P::get_property<sycl::ext::intel::experimental::uses_ready_key>() ==
+      sycl::ext::intel::experimental::uses_ready<false>);
+  static_assert(P::get_property<sycl::ext::intel::experimental::in_csr_key>() ==
+                sycl::ext::intel::experimental::in_csr<true>);
+  static_assert(
+      P::get_property<sycl::ext::intel::experimental::
+                          first_symbol_in_high_order_bits_key>() ==
+      sycl::ext::intel::experimental::first_symbol_in_high_order_bits_off);
+  static_assert(
+      P::get_property<sycl::ext::intel::experimental::protocol_key>() ==
+      sycl::ext::intel::experimental::protocol_avalon);
+}