Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom

[SYCL][DOC] Update pipe and pipe properties specifications #14290

Closed
wants to merge 6 commits into from
Closed

[SYCL][DOC] Update pipe and pipe properties specifications #14290

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
6 commits
Select commit Hold shift + click to select a range
a5ba5ef
Update property list -> properties object
justin-rosner Jun 24, 2024
536b06b
Update pipe and pipe properties spec
justin-rosner Jun 24, 2024
7a768fb
Rename avalon_mm to implement_in_csr
justin-rosner Jun 25, 2024
dc3ca96
Revert name of implement_in_csr back to avalon_mm
justin-rosner Jul 5, 2024
d093c5d
Fix list formatting
justin-rosner Jul 10, 2024
3383a56
Break up properties table
justin-rosner Jul 26, 2024
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
166 changes: 162 additions & 4 deletions .../doc/extensions/experimental/sycl_ext_intel_data_flow_pipes_properties.asciidoc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -32,7 +32,7 @@ to pipes.

== Notice

Copyright (c) 2022-2023 Intel Corporation. All rights reserved.
Copyright (c) 2022-2024 Intel Corporation. All rights reserved.

== Status

Expand Down Expand Up @@ -61,7 +61,8 @@ Aditi Kumaraswamy, Intel +
Robert Ho, Intel +
Sherry Yuan, Intel +
Peter Colberg, Intel +
Zibai Wang, Intel
Zibai Wang, Intel +
Justin Rosner, Intel

== Dependencies

Expand Down Expand Up @@ -100,7 +101,7 @@ value to determine which of the extension's APIs the implementation supports.

=== Pipe properties

Below is a list of compile-time-constant properties which `pipe` supports.
Below is a list of compile-time-constant properties which `pipe` supports. This includes properties that can be placed on the read/write calls, or on the pipe type itself.

```c++
namespace sycl {
Expand Down Expand Up @@ -154,6 +155,28 @@ struct protocol_key {
protocol_key, std::integral_constant<protocol_name, Protocol>>;
};

enum class latency_control_type {
none, //default
exact,
max,
min
};

struct latency_anchor_id_key {
template <int Anchor>
using value_t =
oneapi::experimental::property_value<latency_anchor_id_key,
std::integral_constant<int, Anchor>>;
};

struct latency_constraint_key {
template <int Target, latency_control_type Type, int Cycle>
using value_t = oneapi::experimental::property_value<
latency_constraint_key, std::integral_constant<int, Target>,
std::integral_constant<latency_control_type, Type>,
std::integral_constant<int, Cycle>>;
};

template <uint32_t Latency>
inline constexpr ready_latency_key::value_t<Latency> ready_latency;

Expand All @@ -173,6 +196,16 @@ inline constexpr first_symbol_in_high_order_bits_key::value_t<HighOrder>
template <protocol_name Protocol>
inline constexpr protocol_key::value_t<Protocol> protocol;

template <int Anchor>
inline constexpr latency_anchor_id_key::value_t<Anchor> latency_anchor_id;

template <int Target, latency_control_type Type, int Cycle>
inline constexpr latency_constraint_key::value_t<Target, Type, Cycle>
latency_constraint;

template <class Name, class DataT, int32_t MinCapacity = 0,
class PropertiesT = decltype(oneapi::experimental::properties{})>

namespace avalon-st {
using sycl::ext::intel::experimental::bits_per_symbol;
using sycl::ext::intel::experimental::bits_per_symbol_key;
Expand All @@ -194,6 +227,9 @@ namespace avalon-st {
```

--

The following is a table of properties that can be applied to the pipe type.

[options="header"]
|====
| Property | Description
Expand Down Expand Up @@ -273,7 +309,6 @@ or *avalon_mm*.
|`protocol`
| Specifies the protocol for the pipe interface. Currently, the protocols supported
are: *avalon_streaming*, *avalon_mm*, and *axi_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].
Expand All @@ -287,9 +322,131 @@ Provide an Avalon memory mapped interface as described in https://www.intel.com/
Provide an AXI4-Stream interface as described in https://documentation-service.arm.com/static/642583d7314e245d086bc8c9[AMBA 4 AXI4-Stream Protocol Specification].

The default protocol is *avalon_streaming*

|====
--

The following is a table of properties that can be applied on the read/write calls to the pipe.

[options="header"]
|====
| Property | Description

|`latency_anchor_id<N>`
| Valid values: Non-negative integer values of N

This property associated an ID with the current read/write function call, which can then be referenced by other `latency_constraint` properties elsewhere in the program to define relative latency constraints. The ID must be unique within the application, and a diagnostic is required if that condition is not met.

This property is only valid when passed as a function argument of a device side read/write method as a properties object.

|`latency_constraint<A, B, C>`
| Valid values: `A` and `C` are non-negative integer values, while `B` is a enum value from latency_control_type.

This property is a tuple of three values which cause the current read/write function call to act as an endpoint of a latency constraint relative to a specified `latency_anchor_id` defined by a different instruction.

`A` is the ID of the target anchor defined on a different instruction through a `latency_anchor_id` property.

`B` is an enum value: The type of control from the set {`latency_control_type::exact`, `latency_control_type::max`, `latency_control_type::min`}.

`C` is the relative clock cycle difference between the target anchor and the current function call that the constraint should infer subject to the type of the control (exact, min, max).

This property is only valud when passed as a function argument of a device side read/write method as a properties object.

GarveyJoe marked this conversation as resolved.
Show resolved Hide resolved
|====

== Latency Control Example

[source,c++]
----
#include <sycl/ext/intel/fpga_extensions.hpp>
...
using Pipe1 = ext::intel::experimental::pipe<class PipeClass1, int, 8>;
using Pipe2 = ext::intel::experimental::pipe<class PipeClass2, int, 8>;
using Pipe3 = ext::intel::experimental::pipe<class PipeClass2, int, 8>;

myQueue.submit([&](handler &cgh) {
cgh.single_task<class foo>([=] {
// The following Pipe1::read is anchor 0
int value = Pipe1::read(
ext::oneapi::experimental::properties(latency_anchor_id<0>));

// The following Pipe2::write is anchor 1
// The following Pipe2::write occurs exactly 2 cycles after anchor 0
Pipe2::write(value,
ext::oneapi::experimental::properties(
latency_anchor_id<1>,
latency_constraint<0, latency_control_type::exact, 2>));

// The following Pipe3::write occurs at least 2 cycles after anchor 1
Pipe3::write(value,
ext::oneapi::experimental::properties(
latency_constraint<1, latency_control_type::min, 2>));
});
});
----

=== Device side pipe read/write

[source,c++]
----
namespace sycl::ext::intel::experimental {
enum class latency_control_type {
none, // default
exact,
max,
min
};

struct latency_anchor_id_key {
template <int Anchor>
using value_t =
oneapi::experimental::property_value<latency_anchor_id_key,
std::integral_constant<int, Anchor>>;
};

struct latency_constraint_key {
template <int Target, latency_control_type Type, int Cycle>
using value_t = oneapi::experimental::property_value<
latency_constraint_key, std::integral_constant<int, Target>,
std::integral_constant<latency_control_type, Type>,
std::integral_constant<int, Cycle>>;
};

template <int Anchor>
inline constexpr latency_anchor_id_key::value_t<Anchor> latency_anchor_id;

template <int Target, latency_control_type Type, int Cycle>
inline constexpr latency_constraint_key::value_t<Target, Type, Cycle>
latency_constraint;

template <class Name, class DataT, int32_t MinCapacity = 0,
class PropertiesT = decltype(oneapi::experimental::properties{})>
class pipe {
// Blocking
static DataT read();

template <typename PropertiesT>
static DataT read( PropertiesT Properties );

static void write( const DataT &Data);

template <typename PropertiesT>
static void write( const DataT &Data, PropertiesT Properties );

// Non-blocking
static DataT read( bool &Success );

template <typename PropertiesT>
static DataT read( bool &Success, PropertiesT Properties );

static void write( const DataT &Data, bool &Success );

template <typename PropertiesT>
static void write( const DataT &Data, bool &Success, PropertiesT Properties );
}
} // namespace sycl::ext::intel::experimental
----

== Revision History

[cols="5,15,15,70"]
Expand All @@ -300,6 +457,7 @@ The default protocol is *avalon_streaming*
|1|2022-03-18|Peter Colberg|*Initial public working draft*
|2|2023-04-06|Robert Ho|Removal of unused properties, update protocols
|3|2023-08-30|Robert Ho|Add axi_streaming protocol
|4|2024-06-24|Justin Rosner|Add latency controls and update avalon_mm description
|========================================

//************************************************************************
Expand Down
115 changes: 4 additions & 111 deletions sycl/doc/extensions/supported/sycl_ext_intel_dataflow_pipes.asciidoc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -31,7 +31,7 @@ This document describes an extension that adds pipes to SYCL. Pipes are first i

== Notice

Copyright (c) 2019-2023 Intel Corporation. All rights reserved.
Copyright (c) 2019-2024 Intel Corporation. All rights reserved.

== Status

Expand Down Expand Up @@ -60,7 +60,8 @@ Aditi Kumaraswamy, Intel +
Robert Ho, Intel +
Sherry Yuan, Intel +
Peter Colberg, Intel +
Zibai Wang, Intel
Zibai Wang, Intel +
Justin Rosner, Intel

== Dependencies

Expand Down Expand Up @@ -653,115 +654,6 @@ Automated mechanisms are possible to provide uniquification across calls, and co
*RESOLUTION*: Resolved. Abstraction/libraries on top enable functionality like this. We will make public a library that enables arrays of pipes.
--

== Experimental APIs

*NOTE*: The APIs described in this section are experimental. Future versions of this extension may change these APIs in ways that are incompatible with the versions described here.

The Intel FPGA experimental `pipe` class is implemented in `sycl/ext/intel/experimental/pipes.hpp` which is included in `sycl/ext/intel/fpga_extensions.hpp`.

In the experimental API version, the device side read/write methods take in a property list as function argument, which can contain the latency control properties `latency_anchor_id` and/or `latency_constraint`.

* `sycl::ext::intel::experimental::latency_anchor_id<N>`, where `N` is an integer: An ID to associate with the current read/write function call, which can then be referenced by other `latency_constraint` properties elsewhere in the program to define relative latency constaints. ID must be unique within the application, and a diagnostic is required if that condition is not met.
* `sycl::ext::intel::experimental::latency_constraint<A, B, C>`: A tuple of three values which cause the current read/write function call to act as an endpoint of a latency constraint relative to a specified `latency_anchor_id` defined by a different instruction.
** `A` is an integer: The ID of the target anchor defined on a different instruction through a `latency_anchor_id` property.
** `B` is an enum value: The type of control from the set {`latency_control_type::exact`, `latency_control_type::max`, `latency_control_type::min`}.
** `C` is an integer: The relative clock cycle difference between the target anchor and the current function call, that the constraint should infer subject to the type of the control (exact, max, min).

=== Device side pipe read/write

[source,c++]
----
// Added in version 2 of this extension.
namespace sycl::ext::intel::experimental {
enum class latency_control_type {
none, // default
exact,
max,
min
};

struct latency_anchor_id_key {
template <int Anchor>
using value_t =
oneapi::experimental::property_value<latency_anchor_id_key,
std::integral_constant<int, Anchor>>;
};

struct latency_constraint_key {
template <int Target, latency_control_type Type, int Cycle>
using value_t = oneapi::experimental::property_value<
latency_constraint_key, std::integral_constant<int, Target>,
std::integral_constant<latency_control_type, Type>,
std::integral_constant<int, Cycle>>;
};

template <int Anchor>
inline constexpr latency_anchor_id_key::value_t<Anchor> latency_anchor_id;

template <int Target, latency_control_type Type, int Cycle>
inline constexpr latency_constraint_key::value_t<Target, Type, Cycle>
latency_constraint;

template <class Name, class DataT, int32_t MinCapacity = 0,
class PropertiesT = decltype(oneapi::experimental::properties{})>
class pipe {
// Blocking
static DataT read();

template <typename PropertiesT>
static DataT read( PropertiesT Properties );

static void write( const DataT &Data);

template <typename PropertiesT>
static void write( const DataT &Data, PropertiesT Properties );

// Non-blocking
static DataT read( bool &Success );

template <typename PropertiesT>
static DataT read( bool &Success, PropertiesT Properties );

static void write( const DataT &Data, bool &Success );

template <typename PropertiesT>
static void write( const DataT &Data, bool &Success, PropertiesT Properties );
}
} // namespace sycl::ext::intel::experimental
----

=== Latency Control example

[source,c++]
----
// Added in version 2 of this extension.
#include <sycl/ext/intel/fpga_extensions.hpp>
...
using Pipe1 = ext::intel::experimental::pipe<class PipeClass1, int, 8>;
using Pipe2 = ext::intel::experimental::pipe<class PipeClass2, int, 8>;
using Pipe3 = ext::intel::experimental::pipe<class PipeClass2, int, 8>;

myQueue.submit([&](handler &cgh) {
cgh.single_task<class foo>([=] {
// The following Pipe1::read is anchor 0
int value = Pipe1::read(
ext::oneapi::experimental::properties(latency_anchor_id<0>));

// The following Pipe2::write is anchor 1
// The following Pipe2::write occurs exactly 2 cycles after anchor 0
Pipe2::write(value,
ext::oneapi::experimental::properties(
latency_anchor_id<1>,
latency_constraint<0, latency_control_type::exact, 2>));

// The following Pipe3::write occurs at least 2 cycles after anchor 1
Pipe3::write(value,
ext::oneapi::experimental::properties(
latency_constraint<1, latency_control_type::min, 2>));
});
});
----

== Host Side pipe read/write

If the read/write member functions of a pipe are called from the host side, a `sycl::queue` is added to the parameters. The `memory_order` parameter is also added to the parameters for future work.
Expand Down Expand Up @@ -854,6 +746,7 @@ extension's APIs the implementation supports.
|3|2020-04-27|Michael Kinsner|Clarify that pipe operations behave as-if they are relaxed atomic operations. Make SYCL2020 the baseline
|4|2021-12-02|Shuo Niu|Add experimental latency control API
|5|2023-03-27|Zibai Wang|Experimental API change only. Add memory order parameter and compile-time properties. Add host pipe read/write functions.
|6|2024-06-24|Justin Rosner|Move latency controls to the host pipe interface controls extension
|========================================

//************************************************************************
Expand Down