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..74d36204e92c0 --- /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 : /* unspecified */ { + 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 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` +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. +//************************************************************************ 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() ==