From 258abfd4d091f7415336e371b35c79045fca626c Mon Sep 17 00:00:00 2001 From: Robert Ho <84344325+rho180@users.noreply.github.com> Date: Tue, 11 Apr 2023 13:06:36 -0400 Subject: [PATCH 1/4] Create sycl_ext_intel_data_flow_pipes_properties.asciidoc --- ..._intel_data_flow_pipes_properties.asciidoc | 274 ++++++++++++++++++ 1 file changed, 274 insertions(+) create mode 100644 sycl/doc/extensions/experimental/sycl_ext_intel_data_flow_pipes_properties.asciidoc diff --git a/sycl/doc/extensions/experimental/sycl_ext_intel_data_flow_pipes_properties.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_intel_data_flow_pipes_properties.asciidoc new file mode 100644 index 0000000000000..fb293b4d9f2f0 --- /dev/null +++ b/sycl/doc/extensions/experimental/sycl_ext_intel_data_flow_pipes_properties.asciidoc @@ -0,0 +1,274 @@ += sycl_ext_intel_data_flow_pipes_properties + +:source-highlighter: coderay +:coderay-linenums-mode: table + +// This section needs to be after the document title. +:doctype: book +:toc2: +:toc: left +:encoding: utf-8 +:lang: en +:dpcpp: pass:[DPC++] +:blank: pass:[ +] + +// Set the default source code type in this document to C++, +// for syntax highlighting purposes. This is needed because +// docbook uses c++ and html5 uses cpp. +:language: {basebackend@docbook:c++:cpp} + +// This is necessary for asciidoc, but not for asciidoctor +:cpp: C++ + +NOTE: Khronos(R) is a registered trademark and SYCL(TM) and SPIR(TM) are +trademarks of The Khronos Group Inc. OpenCL(TM) is a trademark of Apple Inc. +used by permission by Khronos. + +NOTE: This document is better viewed when rendered as html with asciidoctor. +GitHub does not render image icons. + +This document describes an extension that adds compile-time constant properties +to pipes. + +== Notice + +Copyright (c) 2022-2023 Intel Corporation. All rights reserved. + +== Status + +This is an experimental extension specification, intended to provide early +access to features and gather community feedback. Interfaces defined in this +specification are implemented in {dpcpp}, but they are not finalized and may +change incompatibly in future versions of {dpcpp} without prior notice. +*Shipping software products should not rely on APIs defined in this +specification.* + +== Version + +Built On: {docdate} + +Revision: A + +== Contact + +Robert Ho, Intel (robert 'dot' ho 'at' intel 'dot' com) + +== Contributors + +Bo Lei, Intel + +Marco Jacques, Intel + +Joe Garvey, Intel + +Aditi Kumaraswamy, Intel + +Robert Ho, Intel + +Sherry Yuan, Intel + +Peter Colberg, Intel + +Zibai Wang, Intel + +== Dependencies + +This extension is written against the SYCL 2020 specification, Revision 6 and +the following extensions: + +- link:../supported/sycl_ext_intel_dataflow_pipes.asciidoc[SYCL_INTEL_data_flow_pipes] +- link:../experimental/sycl_ext_oneapi_properties.asciidoc[sycl_ext_oneapi_properties] + +== Overview + +This extension introduces properties that establish differences in the +implementation of `sycl::ext::intel::experimental::pipe`. These properties are FPGA specific. An example +of the syntax can be seen below. + +[source,c++] +---- +using pipe = pipe})>; +---- + +== Feature test macro + +This extension provides a feature-test macro as described in the core SYCL +specification section 6.3.3 "Feature test macros". Therefore, an implementation +supporting this extension must predefine the macro +`SYCL_EXT_INTEL_FPGA_PIPE_PROPERTIES` to one of the values defined in the table +below. Applications can test for the existence of this macro to determine if +the implementation supports this feature, or applications can test the macro's +value to determine which of the extension's APIs the implementation supports. + +[%header,cols="1,5"] +|=== +|Value |Description +|1 |Initial extension version. Base features are supported. +|=== + +=== Pipe properties + +Below is a list of compile-time-constant properties which `pipe` supports. + +```c++ +namespace sycl::ext::intel::experimental { + +struct ready_latency_key { + template + using value_t = oneapi::experimental::property_value< + ready_latency_key, std::integral_constant>; +}; + +struct bits_per_symbol_key { + template + using value_t = + oneapi::experimental::property_value>; +}; + +struct uses_valid_key { + template + using value_t = + oneapi::experimental::property_value>; +}; + +struct first_symbol_in_high_order_bits_key { + template + using value_t = + oneapi::experimental::property_value>; +}; + +enum class protocol_name : std::uint16_t { + AVALON_STREAMING = 0, + AVALON_STREAMING_USES_READY = 1, + AVALON_MM = 2, + AVALON_MM_USES_READY = 3 +}; + +struct protocol_key { + template + using value_t = oneapi::experimental::property_value< + protocol_key, std::integral_constant>; +}; + +template +inline constexpr ready_latency_key::value_t ready_latency; + +template +inline constexpr bits_per_symbol_key::value_t bits_per_symbol; + +template +inline constexpr uses_valid_key::value_t uses_valid; + +template +inline constexpr first_symbol_in_high_order_bits_key::value_t + first_symbol_in_high_order_bits; + +template +inline constexpr protocol_key::value_t protocol; + +} // namespace sycl::ext::intel::experimental +``` + +-- +[options="header"] +|==== +| Property | Description + +|`ready_latency` +| Valid values: Non-negative integer value. + +Default value: 0 + +The number of cycles between when the ready signal is deasserted and when the +pipe can no longer accept new inputs. + +This property is not guaranteed to be respected if the pipe is an inter-kernel +pipe. The compiler is allowed to optimize the pipe if both sides are visible. + +|`bits_per_symbol` +| Valid values: A positive integer value that evenly divides by the data type size. + +Default value: 8 + +Describes how the data is broken into symbols on the data bus. + +Data is broken down according to how you set the first_symbol_in_high_order_bits +property. By default, data is broken down in little endian order. + +This property is not guaranteed to be respected if the pipe is an inter-kernel +pipe. The compiler is allowed to optimize the pipe if both sides are visible. + +|`uses_valid` +| Valid values: true or false + +Default value: true + +Controls whether a valid signal is present on the pipe interface. If false, the +upstream source must provide valid data on every cycle that ready is asserted. + +This is equivalent to changing the pipe read calls to tryRead and assuming that +success is always true. + +If set to false, the min_capacity pipe class template parameter and ready_latency +property must be 0. + +This property is not guaranteed to be respected if the pipe is an inter-kernel +pipe. The compiler is allowed to optimize the pipe if both sides are visible. + +|`first_symbol_in_high_order_bits` +| Valid values: true or false + +Default value: false + +Specifies whether the data symbols in the pipe are in big-endian +order. + +This property is not guaranteed to be respected if the pipe is an inter-kernel +pipe. The compiler is allowed to optimize the pipe if both sides are visible. + +|`protocol` +| Specifies the protocol for the pipe interface. Currently, the protocols supported +are: *AVALON_STREAMING*, *AVALON_STREAMING_USES_READY*, *AVALON_MM*, and *AVALON_MM_USES_READY*. + +*AVALON_STREAMING* + +Provide an Avalon streaming interface as described in https://www.intel.com/content/www/us/en/docs/programmable/683091/22-3/introduction-to-the-interface-specifications.html[Intel® Avalon Interface Specifications]. + +With this choice of protocol, no ready signal is exposed by the host pipe, and the sink cannot backpressure. + +*AVALON_STREAMING_USES_READY* + +Provide an Avalon streaming interface as described in https://www.intel.com/content/www/us/en/docs/programmable/683091/22-3/introduction-to-the-interface-specifications.html[Intel® Avalon Interface Specifications]. + +This protocol allows the sink to backpressure by deasserting the ready signal asserted. The sink signifies that it is ready to consume data by asserting the ready signal. + +*AVALON_MM* + +Provide an Avalon memory mapped interface as described in https://www.intel.com/content/www/us/en/docs/programmable/683091/22-3/introduction-to-the-interface-specifications.html[Intel® Avalon Interface Specifications]. + +With this protocol, an implicit ready signal is held high, and the sink cannot backpressure. + +*AVALON_MM_USES_READY* + +Provide an Avalon memory mapped interface as described in https://www.intel.com/content/www/us/en/docs/programmable/683091/22-3/introduction-to-the-interface-specifications.html[Intel® Avalon Interface Specifications]. + +With this protocol, an additional memory mapped location is created to hold the ready signal. You must set the uses_valid property to true. + +The default protocol is *AVALON_STREAMING_USES_READY* +|==== +-- + +== Revision History + +[cols="5,15,15,70"] +[grid="rows"] +[options="header"] +|======================================== +|Rev|Date|Author|Changes +|1|2022-03-18|Peter Colberg|*Initial public working draft* +|2|2023-04-06|Robert Ho|Removal of unused properties, update protocols +|======================================== + +//************************************************************************ +//Other formatting suggestions: +// +//* Use *bold* text for host APIs, or [source] syntax highlighting. +//* Use +mono+ text for device APIs, or [source] syntax highlighting. +//* Use +mono+ text for extension names, types, or enum values. +//* Use _italics_ for parameters. +//************************************************************************ From f9e876912bf41595341bf3d774b7826be8ea4207 Mon Sep 17 00:00:00 2001 From: Robert Ho <84344325+rho180@users.noreply.github.com> Date: Wed, 12 Apr 2023 09:41:20 -0400 Subject: [PATCH 2/4] Apply suggestions from code review Co-authored-by: Steffen Larsen --- ...l_ext_intel_data_flow_pipes_properties.asciidoc | 14 +++++++------- 1 file changed, 7 insertions(+), 7 deletions(-) diff --git a/sycl/doc/extensions/experimental/sycl_ext_intel_data_flow_pipes_properties.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_intel_data_flow_pipes_properties.asciidoc index fb293b4d9f2f0..ac837c37f17a5 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_intel_data_flow_pipes_properties.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_intel_data_flow_pipes_properties.asciidoc @@ -132,7 +132,7 @@ struct first_symbol_in_high_order_bits_key { std::bool_constant>; }; -enum class protocol_name : std::uint16_t { +enum class protocol_name : /* unspecified */ { AVALON_STREAMING = 0, AVALON_STREAMING_USES_READY = 1, AVALON_MM = 2, @@ -187,24 +187,24 @@ Default value: 8 Describes how the data is broken into symbols on the data bus. -Data is broken down according to how you set the first_symbol_in_high_order_bits +Data is broken down according to how you set the `first_symbol_in_high_order_bits` property. By default, data is broken down in little endian order. This property is not guaranteed to be respected if the pipe is an inter-kernel pipe. The compiler is allowed to optimize the pipe if both sides are visible. |`uses_valid` -| Valid values: true or false +| Valid values: `true` or `false` -Default value: true +Default value: `true` -Controls whether a valid signal is present on the pipe interface. If false, the +Controls whether a valid signal is present on the pipe interface. If `false`, the upstream source must provide valid data on every cycle that ready is asserted. This is equivalent to changing the pipe read calls to tryRead and assuming that success is always true. -If set to false, the min_capacity pipe class template parameter and ready_latency +If set to `false`, the `min_capacity` pipe class template parameter and `ready_latency` property must be 0. This property is not guaranteed to be respected if the pipe is an inter-kernel @@ -247,7 +247,7 @@ With this protocol, an implicit ready signal is held high, and the sink cannot b Provide an Avalon memory mapped interface as described in https://www.intel.com/content/www/us/en/docs/programmable/683091/22-3/introduction-to-the-interface-specifications.html[Intel® Avalon Interface Specifications]. -With this protocol, an additional memory mapped location is created to hold the ready signal. You must set the uses_valid property to true. +With this protocol, an additional memory mapped location is created to hold the ready signal. You must set the `uses_valid` property to `true`. The default protocol is *AVALON_STREAMING_USES_READY* |==== From e0e3816f101c41b874bd09e714d21fadfda84284 Mon Sep 17 00:00:00 2001 From: Robert Ho <84344325+rho180@users.noreply.github.com> Date: Wed, 12 Apr 2023 09:46:43 -0400 Subject: [PATCH 3/4] Apply more review changes --- ..._intel_data_flow_pipes_properties.asciidoc | 22 +++++++++---------- 1 file changed, 11 insertions(+), 11 deletions(-) diff --git a/sycl/doc/extensions/experimental/sycl_ext_intel_data_flow_pipes_properties.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_intel_data_flow_pipes_properties.asciidoc index ac837c37f17a5..74d36204e92c0 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_intel_data_flow_pipes_properties.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_intel_data_flow_pipes_properties.asciidoc @@ -133,10 +133,10 @@ struct first_symbol_in_high_order_bits_key { }; enum class protocol_name : /* unspecified */ { - AVALON_STREAMING = 0, - AVALON_STREAMING_USES_READY = 1, - AVALON_MM = 2, - AVALON_MM_USES_READY = 3 + avalon_streaming = 0, + avalon_streaming_uses_ready = 1, + avalon_mm = 2, + avalon_mm_uses_ready = 3 }; struct protocol_key { @@ -201,7 +201,7 @@ Default value: `true` Controls whether a valid signal is present on the pipe interface. If `false`, the upstream source must provide valid data on every cycle that ready is asserted. -This is equivalent to changing the pipe read calls to tryRead and assuming that +This is equivalent to changing the pipe read calls to a non-blocking call and assuming that success is always true. If set to `false`, the `min_capacity` pipe class template parameter and `ready_latency` @@ -223,33 +223,33 @@ pipe. The compiler is allowed to optimize the pipe if both sides are visible. |`protocol` | Specifies the protocol for the pipe interface. Currently, the protocols supported -are: *AVALON_STREAMING*, *AVALON_STREAMING_USES_READY*, *AVALON_MM*, and *AVALON_MM_USES_READY*. +are: *avalon_streaming*, *avalon_streaming_uses_ready*, *avalon_mm*, and *avalon_mm_uses_ready*. -*AVALON_STREAMING* +*avalon_streaming* Provide an Avalon streaming interface as described in https://www.intel.com/content/www/us/en/docs/programmable/683091/22-3/introduction-to-the-interface-specifications.html[Intel® Avalon Interface Specifications]. With this choice of protocol, no ready signal is exposed by the host pipe, and the sink cannot backpressure. -*AVALON_STREAMING_USES_READY* +*avalon_streaming_uses_ready* Provide an Avalon streaming interface as described in https://www.intel.com/content/www/us/en/docs/programmable/683091/22-3/introduction-to-the-interface-specifications.html[Intel® Avalon Interface Specifications]. This protocol allows the sink to backpressure by deasserting the ready signal asserted. The sink signifies that it is ready to consume data by asserting the ready signal. -*AVALON_MM* +*avalon_mm* Provide an Avalon memory mapped interface as described in https://www.intel.com/content/www/us/en/docs/programmable/683091/22-3/introduction-to-the-interface-specifications.html[Intel® Avalon Interface Specifications]. With this protocol, an implicit ready signal is held high, and the sink cannot backpressure. -*AVALON_MM_USES_READY* +*avalon_mm_uses_ready* Provide an Avalon memory mapped interface as described in https://www.intel.com/content/www/us/en/docs/programmable/683091/22-3/introduction-to-the-interface-specifications.html[Intel® Avalon Interface Specifications]. With this protocol, an additional memory mapped location is created to hold the ready signal. You must set the `uses_valid` property to `true`. -The default protocol is *AVALON_STREAMING_USES_READY* +The default protocol is *avalon_streaming_uses_ready* |==== -- From c4871482f4d4da0bd2a282536a25ad947e574925 Mon Sep 17 00:00:00 2001 From: "Ho, Robert" Date: Mon, 24 Apr 2023 05:25:53 -0700 Subject: [PATCH 4/4] Align pipe properties implementation with spec --- .../intel/experimental/pipe_properties.hpp | 53 +++---------------- .../sycl/ext/intel/experimental/pipes.hpp | 2 +- .../sycl/ext/oneapi/properties/property.hpp | 10 ++-- .../extensions/properties/properties_pipe.cpp | 34 ++---------- 4 files changed, 17 insertions(+), 82 deletions(-) diff --git a/sycl/include/sycl/ext/intel/experimental/pipe_properties.hpp b/sycl/include/sycl/ext/intel/experimental/pipe_properties.hpp index 747c6359c7012..5d90b88202750 100644 --- a/sycl/include/sycl/ext/intel/experimental/pipe_properties.hpp +++ b/sycl/include/sycl/ext/intel/experimental/pipe_properties.hpp @@ -18,12 +18,6 @@ namespace ext { namespace intel { namespace experimental { -struct min_capacity_key { - template - using value_t = oneapi::experimental::property_value< - min_capacity_key, std::integral_constant>; -}; - struct ready_latency_key { template using value_t = oneapi::experimental::property_value< @@ -44,13 +38,6 @@ struct uses_valid_key { std::bool_constant>; }; -struct in_csr_key { - template - using value_t = - oneapi::experimental::property_value>; -}; - struct first_symbol_in_high_order_bits_key { template using value_t = @@ -59,10 +46,10 @@ struct first_symbol_in_high_order_bits_key { }; enum class protocol_name : std::uint16_t { - AVALON_STREAMING = 0, - AVALON_STREAMING_USES_READY = 1, - AVALON_MM = 2, - AVALON_MM_USES_READY = 3 + avalon_streaming = 0, + avalon_streaming_uses_ready = 1, + avalon_mm = 2, + avalon_mm_uses_ready = 3 }; struct protocol_key { @@ -71,9 +58,6 @@ struct protocol_key { protocol_key, std::integral_constant>; }; -template -inline constexpr min_capacity_key::value_t min_capacity; - template inline constexpr ready_latency_key::value_t ready_latency; @@ -85,10 +69,6 @@ inline constexpr uses_valid_key::value_t uses_valid; inline constexpr uses_valid_key::value_t uses_valid_on; inline constexpr uses_valid_key::value_t uses_valid_off; -template inline constexpr in_csr_key::value_t in_csr; -inline constexpr in_csr_key::value_t in_csr_on; -inline constexpr in_csr_key::value_t in_csr_off; - template inline constexpr first_symbol_in_high_order_bits_key::value_t first_symbol_in_high_order_bits; @@ -99,14 +79,14 @@ inline constexpr first_symbol_in_high_order_bits_key::value_t template inline constexpr protocol_key::value_t protocol; -inline constexpr protocol_key::value_t +inline constexpr protocol_key::value_t protocol_avalon_streaming; inline constexpr protocol_key::value_t< - protocol_name::AVALON_STREAMING_USES_READY> + protocol_name::avalon_streaming_uses_ready> protocol_avalon_streaming_uses_ready; -inline constexpr protocol_key::value_t +inline constexpr protocol_key::value_t protocol_avalon_mm; -inline constexpr protocol_key::value_t +inline constexpr protocol_key::value_t protocol_avalon_mm_uses_ready; } // namespace experimental @@ -115,9 +95,6 @@ inline constexpr protocol_key::value_t namespace oneapi { namespace experimental { -template <> -struct is_property_key : std::true_type { -}; template <> struct is_property_key : std::true_type {}; @@ -127,17 +104,12 @@ struct is_property_key template <> struct is_property_key : std::true_type {}; template <> -struct is_property_key : std::true_type {}; -template <> struct is_property_key : std::true_type {}; template <> struct is_property_key : std::true_type {}; namespace detail { -template <> struct PropertyToKind { - static constexpr PropKind Kind = PropKind::MinCapacity; -}; template <> struct PropertyToKind { static constexpr PropKind Kind = PropKind::ReadyLatency; }; @@ -147,9 +119,6 @@ template <> struct PropertyToKind { template <> struct PropertyToKind { static constexpr PropKind Kind = PropKind::UsesValid; }; -template <> struct PropertyToKind { - static constexpr PropKind Kind = PropKind::ImplementInCSR; -}; template <> struct PropertyToKind< intel::experimental::first_symbol_in_high_order_bits_key> { @@ -159,9 +128,6 @@ template <> struct PropertyToKind { static constexpr PropKind Kind = PropKind::PipeProtocol; }; -template <> -struct IsCompileTimeProperty - : std::true_type {}; template <> struct IsCompileTimeProperty : std::true_type {}; @@ -172,9 +138,6 @@ template <> struct IsCompileTimeProperty : std::true_type {}; template <> -struct IsCompileTimeProperty : std::true_type { -}; -template <> struct IsCompileTimeProperty< intel::experimental::first_symbol_in_high_order_bits_key> : std::true_type { }; diff --git a/sycl/include/sycl/ext/intel/experimental/pipes.hpp b/sycl/include/sycl/ext/intel/experimental/pipes.hpp index 37bbef0c66bea..15478aa993674 100644 --- a/sycl/include/sycl/ext/intel/experimental/pipes.hpp +++ b/sycl/include/sycl/ext/intel/experimental/pipes.hpp @@ -385,7 +385,7 @@ class pipe : public pipe_base { first_symbol_in_high_order_bits_key>::template get(0); static constexpr protocol_name m_protocol = oneapi::experimental::detail:: ValueOrDefault<_propertiesT, protocol_key>::template get( - protocol_name::AVALON_STREAMING_USES_READY); + protocol_name::avalon_streaming_uses_ready); public: static constexpr struct ConstantPipeStorageExp m_Storage = { diff --git a/sycl/include/sycl/ext/oneapi/properties/property.hpp b/sycl/include/sycl/ext/oneapi/properties/property.hpp index 7c2e3063ace13..1c63a3966f6fc 100644 --- a/sycl/include/sycl/ext/oneapi/properties/property.hpp +++ b/sycl/include/sycl/ext/oneapi/properties/property.hpp @@ -190,13 +190,11 @@ enum PropKind : uint32_t { CacheConfig = 24, BitsPerSymbol = 25, FirstSymbolInHigherOrderBit = 26, - MinCapacity = 27, - PipeProtocol = 28, - ReadyLatency = 29, - UsesReady = 30, - UsesValid = 31, + PipeProtocol = 27, + ReadyLatency = 28, + UsesValid = 29, // PropKindSize must always be the last value. - PropKindSize = 32, + PropKindSize = 30, }; // 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 index 70e2c30078db3..892a311cab13d 100644 --- a/sycl/test/extensions/properties/properties_pipe.cpp +++ b/sycl/test/extensions/properties/properties_pipe.cpp @@ -8,20 +8,16 @@ using namespace sycl::ext; constexpr sycl::ext::intel::experimental::protocol_name TestProtocol = - sycl::ext::intel::experimental::protocol_name::AVALON_STREAMING; + sycl::ext::intel::experimental::protocol_name::avalon_streaming; 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::in_csr_key>::value); static_assert( sycl::ext::oneapi::experimental::is_property_key< sycl::ext::intel::experimental::first_symbol_in_high_order_bits_key>:: @@ -30,9 +26,6 @@ int main() { 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); @@ -49,12 +42,6 @@ int main() { 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::in_csr)>::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:: @@ -89,9 +76,6 @@ int main() { static_assert(std::is_same_v< decltype(sycl::ext::intel::experimental::uses_valid_on), decltype(sycl::ext::intel::experimental::uses_valid)>); - static_assert( - std::is_same_v)>); static_assert( std::is_same_v)>); + avalon_streaming_uses_ready>)>); static_assert( std::is_same_v< decltype(sycl::ext::intel::experimental::protocol_avalon_mm), decltype(sycl::ext::intel::experimental::protocol< - sycl::ext::intel::experimental::protocol_name::AVALON_MM>)>); + sycl::ext::intel::experimental::protocol_name::avalon_mm>)>); static_assert( std::is_same_v)>); + avalon_mm_uses_ready>)>); // 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, - sycl::ext::intel::experimental::in_csr, sycl::ext::intel::experimental::first_symbol_in_high_order_bits_off, sycl::ext::intel::experimental::protocol_avalon_streaming)); static_assert(sycl::ext::oneapi::experimental::is_property_list_v

); - static_assert( - P::has_property()); static_assert( P::has_property()); static_assert( P::has_property()); static_assert( P::has_property()); - static_assert(P::has_property()); static_assert(P::has_property()); static_assert( P::has_property()); - static_assert( - P::get_property() == - sycl::ext::intel::experimental::min_capacity<0>); static_assert( P::get_property() == sycl::ext::intel::experimental::ready_latency<1>); @@ -155,8 +131,6 @@ int main() { static_assert( P::get_property() == sycl::ext::intel::experimental::uses_valid); - static_assert(P::get_property() == - sycl::ext::intel::experimental::in_csr); static_assert( P::get_property() ==