Skip to content

Commit 8712084

Browse files
sherry-yuansherry-yuan
committed
[SYCL] Add data_flow_pipe properties
Defines new properties for data flow pipes Properties defined to align with spec in [1] Existing data flow pipe defined in [2] data flow pipe will be accepting new property list, spec in [3] [1] #5839 [2] https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/proposed/sycl_ext_intel_dataflow_pipes.asciidoc [3] #5838
1 parent 86cf56a commit 8712084

File tree

4 files changed

+375
-1
lines changed

4 files changed

+375
-1
lines changed

‎sycl/include/CL/sycl.hpp

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -71,3 +71,5 @@
7171
#include <sycl/ext/oneapi/reduction.hpp>
7272
#include <sycl/ext/oneapi/sub_group.hpp>
7373
#include <sycl/ext/oneapi/sub_group_mask.hpp>
74+
75+
#include <sycl/ext/intel/experimental/pipe_properties.hpp>

‎sycl/include/sycl/ext/intel/experimental/pipe_properties.hpp

Lines changed: 197 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,197 @@
1+
//==----- pipe_properties.hpp - SYCL properties associated with data flow pipe
2+
//---==//
3+
//
4+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
5+
// See https://llvm.org/LICENSE.txt for license information.
6+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
7+
//
8+
//===----------------------------------------------------------------------===//
9+
10+
#pragma once
11+
12+
#include <sycl/ext/oneapi/properties/property.hpp>
13+
#include <sycl/ext/oneapi/properties/property_value.hpp>
14+
15+
__SYCL_INLINE_NAMESPACE(cl) {
16+
namespace sycl {
17+
namespace ext {
18+
namespace intel {
19+
namespace experimental {
20+
21+
struct min_capacity_key {
22+
template <int Capacity>
23+
using value_t = oneapi::experimental::property_value<
24+
min_capacity_key, std::integral_constant<int, Capacity>>;
25+
};
26+
27+
struct ready_latency_key {
28+
template <int Latency>
29+
using value_t = oneapi::experimental::property_value<
30+
ready_latency_key, std::integral_constant<int, Latency>>;
31+
};
32+
33+
struct bits_per_symbol_key {
34+
template <int Bits>
35+
using value_t =
36+
oneapi::experimental::property_value<bits_per_symbol_key,
37+
std::integral_constant<int, Bits>>;
38+
};
39+
40+
struct uses_valid_key {
41+
template <bool Valid>
42+
using value_t =
43+
oneapi::experimental::property_value<uses_valid_key,
44+
sycl::detail::bool_constant<Valid>>;
45+
};
46+
47+
struct uses_ready_key {
48+
template <bool Ready>
49+
using value_t =
50+
oneapi::experimental::property_value<uses_ready_key,
51+
sycl::detail::bool_constant<Ready>>;
52+
};
53+
54+
struct in_csr_key {
55+
template <bool Enable>
56+
using value_t =
57+
oneapi::experimental::property_value<in_csr_key,
58+
sycl::detail::bool_constant<Enable>>;
59+
};
60+
61+
struct first_symbol_in_high_order_bits_key {
62+
template <bool HighOrder>
63+
using value_t = oneapi::experimental::property_value<
64+
first_symbol_in_high_order_bits_key,
65+
sycl::detail::bool_constant<HighOrder>>;
66+
};
67+
68+
enum class protocol_name : std::uint16_t { AVALON, AXI };
69+
struct protocol_key {
70+
template <protocol_name Protocol>
71+
using value_t = oneapi::experimental::property_value<
72+
protocol_key, std::integral_constant<protocol_name, Protocol>>;
73+
};
74+
75+
template <int Capacity>
76+
inline constexpr min_capacity_key::value_t<Capacity> min_capacity;
77+
78+
template <int Latency>
79+
inline constexpr ready_latency_key::value_t<Latency> ready_latency;
80+
81+
template <int Bits>
82+
inline constexpr bits_per_symbol_key::value_t<Bits> bits_per_symbol;
83+
84+
template <bool Valid>
85+
inline constexpr uses_valid_key::value_t<Valid> uses_valid;
86+
inline constexpr uses_valid_key::value_t<true> uses_valid_on;
87+
inline constexpr uses_valid_key::value_t<false> uses_valid_off;
88+
89+
template <bool Ready>
90+
inline constexpr uses_ready_key::value_t<Ready> uses_ready;
91+
inline constexpr uses_ready_key::value_t<true> uses_ready_on;
92+
inline constexpr uses_ready_key::value_t<false> uses_ready_off;
93+
94+
template <bool Enable> inline constexpr in_csr_key::value_t<Enable> in_csr;
95+
inline constexpr in_csr_key::value_t<true> in_csr_on;
96+
inline constexpr in_csr_key::value_t<false> in_csr_off;
97+
98+
template <bool HighOrder>
99+
inline constexpr first_symbol_in_high_order_bits_key::value_t<HighOrder>
100+
first_symbol_in_high_order_bits;
101+
inline constexpr first_symbol_in_high_order_bits_key::value_t<true>
102+
first_symbol_in_high_order_bits_on;
103+
inline constexpr first_symbol_in_high_order_bits_key::value_t<false>
104+
first_symbol_in_high_order_bits_off;
105+
106+
template <protocol_name Protocol>
107+
inline constexpr protocol_key::value_t<Protocol> protocol;
108+
inline constexpr protocol_key::value_t<protocol_name::AVALON> protocol_avalon;
109+
inline constexpr protocol_key::value_t<protocol_name::AXI> protocol_axi;
110+
111+
} // namespace experimental
112+
} // namespace intel
113+
114+
namespace oneapi {
115+
namespace experimental {
116+
117+
template <>
118+
struct is_property_key<intel::experimental::min_capacity_key> : std::true_type {
119+
};
120+
template <>
121+
struct is_property_key<intel::experimental::ready_latency_key>
122+
: std::true_type {};
123+
template <>
124+
struct is_property_key<intel::experimental::bits_per_symbol_key>
125+
: std::true_type {};
126+
template <>
127+
struct is_property_key<intel::experimental::uses_valid_key> : std::true_type {};
128+
template <>
129+
struct is_property_key<intel::experimental::uses_ready_key> : std::true_type {};
130+
template <>
131+
struct is_property_key<intel::experimental::in_csr_key> : std::true_type {};
132+
template <>
133+
struct is_property_key<intel::experimental::first_symbol_in_high_order_bits_key>
134+
: std::true_type {};
135+
template <>
136+
struct is_property_key<intel::experimental::protocol_key> : std::true_type {};
137+
138+
namespace detail {
139+
template <> struct PropertyToKind<intel::experimental::min_capacity_key> {
140+
static constexpr PropKind Kind = PropKind::MinCapacity;
141+
};
142+
template <> struct PropertyToKind<intel::experimental::ready_latency_key> {
143+
static constexpr PropKind Kind = PropKind::ReadyLatency;
144+
};
145+
template <> struct PropertyToKind<intel::experimental::bits_per_symbol_key> {
146+
static constexpr PropKind Kind = PropKind::BitsPerSymbol;
147+
};
148+
template <> struct PropertyToKind<intel::experimental::uses_valid_key> {
149+
static constexpr PropKind Kind = PropKind::UsesValid;
150+
};
151+
template <> struct PropertyToKind<intel::experimental::uses_ready_key> {
152+
static constexpr PropKind Kind = PropKind::UsesReady;
153+
};
154+
template <> struct PropertyToKind<intel::experimental::in_csr_key> {
155+
static constexpr PropKind Kind = PropKind::ImplementInCSR;
156+
};
157+
template <>
158+
struct PropertyToKind<
159+
intel::experimental::first_symbol_in_high_order_bits_key> {
160+
static constexpr PropKind Kind = PropKind::FirstSymbolInHigherOrderBit;
161+
};
162+
template <> struct PropertyToKind<intel::experimental::protocol_key> {
163+
static constexpr PropKind Kind = PropKind::PipeProtocol;
164+
};
165+
166+
template <>
167+
struct IsCompileTimeProperty<intel::experimental::min_capacity_key>
168+
: std::true_type {};
169+
template <>
170+
struct IsCompileTimeProperty<intel::experimental::ready_latency_key>
171+
: std::true_type {};
172+
template <>
173+
struct IsCompileTimeProperty<intel::experimental::bits_per_symbol_key>
174+
: std::true_type {};
175+
template <>
176+
struct IsCompileTimeProperty<intel::experimental::uses_valid_key>
177+
: std::true_type {};
178+
template <>
179+
struct IsCompileTimeProperty<intel::experimental::uses_ready_key>
180+
: std::true_type {};
181+
template <>
182+
struct IsCompileTimeProperty<intel::experimental::in_csr_key> : std::true_type {
183+
};
184+
template <>
185+
struct IsCompileTimeProperty<
186+
intel::experimental::first_symbol_in_high_order_bits_key> : std::true_type {
187+
};
188+
template <>
189+
struct IsCompileTimeProperty<intel::experimental::protocol_key>
190+
: std::true_type {};
191+
192+
} // namespace detail
193+
} // namespace experimental
194+
} // namespace oneapi
195+
} // namespace ext
196+
} // namespace sycl
197+
} // __SYCL_INLINE_NAMESPACE(cl)

‎sycl/include/sycl/ext/oneapi/properties/property.hpp

Lines changed: 8 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -148,7 +148,14 @@ enum PropKind : uint32_t {
148148
HostAccess = 1,
149149
InitMode = 2,
150150
ImplementInCSR = 3,
151-
PropKindSize = 4,
151+
BitsPerSymbol = 4,
152+
FirstSymbolInHigherOrderBit = 5,
153+
MinCapacity = 6,
154+
PipeProtocol = 7,
155+
ReadyLatency = 8,
156+
UsesReady = 9,
157+
UsesValid = 10,
158+
PropKindSize = 11,
152159
};
153160

154161
// This trait must be specialized for all properties and must have a unique

‎sycl/test/extensions/properties/properties_pipe.cpp

Lines changed: 168 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,168 @@
1+
// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -fsyntax-only -Xclang -verify -Xclang -verify-ignore-unexpected=note,warning %s
2+
// expected-no-diagnostics
3+
4+
#include <CL/sycl.hpp>
5+
6+
#include <sycl/ext/intel/experimental/pipe_properties.hpp>
7+
8+
using namespace sycl::ext;
9+
10+
constexpr sycl::ext::intel::experimental::protocol_name TestProtocol =
11+
sycl::ext::intel::experimental::protocol_name::AVALON;
12+
13+
int main() {
14+
// Check that is_property_key is correctly specialized.
15+
static_assert(sycl::ext::oneapi::experimental::is_property_key<
16+
sycl::ext::intel::experimental::min_capacity_key>::value);
17+
static_assert(sycl::ext::oneapi::experimental::is_property_key<
18+
sycl::ext::intel::experimental::ready_latency_key>::value);
19+
static_assert(sycl::ext::oneapi::experimental::is_property_key<
20+
sycl::ext::intel::experimental::bits_per_symbol_key>::value);
21+
static_assert(sycl::ext::oneapi::experimental::is_property_key<
22+
sycl::ext::intel::experimental::uses_valid_key>::value);
23+
static_assert(sycl::ext::oneapi::experimental::is_property_key<
24+
sycl::ext::intel::experimental::uses_ready_key>::value);
25+
static_assert(sycl::ext::oneapi::experimental::is_property_key<
26+
sycl::ext::intel::experimental::in_csr_key>::value);
27+
static_assert(
28+
sycl::ext::oneapi::experimental::is_property_key<
29+
sycl::ext::intel::experimental::first_symbol_in_high_order_bits_key>::
30+
value);
31+
static_assert(sycl::ext::oneapi::experimental::is_property_key<
32+
sycl::ext::intel::experimental::protocol_key>::value);
33+
34+
// Check that is_property_value is correctly specialized.
35+
static_assert(
36+
sycl::ext::oneapi::experimental::is_property_value<
37+
decltype(sycl::ext::intel::experimental::min_capacity<3>)>::value);
38+
static_assert(
39+
sycl::ext::oneapi::experimental::is_property_value<
40+
decltype(sycl::ext::intel::experimental::ready_latency<3>)>::value);
41+
static_assert(
42+
sycl::ext::oneapi::experimental::is_property_value<
43+
decltype(sycl::ext::intel::experimental::bits_per_symbol<3>)>::value);
44+
45+
static_assert(
46+
sycl::ext::oneapi::experimental::is_property_value<
47+
decltype(sycl::ext::intel::experimental::uses_valid<true>)>::value);
48+
static_assert(
49+
sycl::ext::oneapi::experimental::is_property_value<
50+
decltype(sycl::ext::intel::experimental::uses_valid_on)>::value);
51+
static_assert(
52+
sycl::ext::oneapi::experimental::is_property_value<
53+
decltype(sycl::ext::intel::experimental::uses_valid_off)>::value);
54+
55+
static_assert(
56+
sycl::ext::oneapi::experimental::is_property_value<
57+
decltype(sycl::ext::intel::experimental::uses_ready<true>)>::value);
58+
static_assert(
59+
sycl::ext::oneapi::experimental::is_property_value<
60+
decltype(sycl::ext::intel::experimental::uses_ready_on)>::value);
61+
static_assert(
62+
sycl::ext::oneapi::experimental::is_property_value<
63+
decltype(sycl::ext::intel::experimental::uses_ready_off)>::value);
64+
65+
static_assert(sycl::ext::oneapi::experimental::is_property_value<
66+
decltype(sycl::ext::intel::experimental::in_csr<true>)>::value);
67+
static_assert(sycl::ext::oneapi::experimental::is_property_value<
68+
decltype(sycl::ext::intel::experimental::in_csr_on)>::value);
69+
static_assert(sycl::ext::oneapi::experimental::is_property_value<
70+
decltype(sycl::ext::intel::experimental::in_csr_off)>::value);
71+
72+
static_assert(sycl::ext::oneapi::experimental::is_property_value<
73+
decltype(sycl::ext::intel::experimental::
74+
first_symbol_in_high_order_bits<true>)>::value);
75+
static_assert(sycl::ext::oneapi::experimental::is_property_value<
76+
decltype(sycl::ext::intel::experimental::
77+
first_symbol_in_high_order_bits_on)>::value);
78+
static_assert(sycl::ext::oneapi::experimental::is_property_value<
79+
decltype(sycl::ext::intel::experimental::
80+
first_symbol_in_high_order_bits_off)>::value);
81+
82+
static_assert(
83+
sycl::ext::oneapi::experimental::is_property_value<
84+
decltype(sycl::ext::intel::experimental::protocol<TestProtocol>)>::
85+
value);
86+
static_assert(
87+
sycl::ext::oneapi::experimental::is_property_value<
88+
decltype(sycl::ext::intel::experimental::protocol_avalon)>::value);
89+
static_assert(sycl::ext::oneapi::experimental::is_property_value<
90+
decltype(sycl::ext::intel::experimental::protocol_axi)>::value);
91+
92+
// Checks that fully specialized properties are the same as the templated
93+
// variants.
94+
static_assert(std::is_same_v<
95+
decltype(sycl::ext::intel::experimental::uses_valid_on),
96+
decltype(sycl::ext::intel::experimental::uses_valid<true>)>);
97+
static_assert(std::is_same_v<
98+
decltype(sycl::ext::intel::experimental::uses_ready_off),
99+
decltype(sycl::ext::intel::experimental::uses_ready<false>)>);
100+
static_assert(
101+
std::is_same_v<decltype(sycl::ext::intel::experimental::in_csr_on),
102+
decltype(sycl::ext::intel::experimental::in_csr<true>)>);
103+
static_assert(
104+
std::is_same_v<decltype(sycl::ext::intel::experimental::
105+
first_symbol_in_high_order_bits_on),
106+
decltype(sycl::ext::intel::experimental::
107+
first_symbol_in_high_order_bits<true>)>);
108+
static_assert(
109+
std::is_same_v<
110+
decltype(sycl::ext::intel::experimental::protocol_avalon),
111+
decltype(sycl::ext::intel::experimental::protocol<TestProtocol>)>);
112+
static_assert(std::is_same_v<
113+
decltype(sycl::ext::intel::experimental::protocol_axi),
114+
decltype(sycl::ext::intel::experimental::protocol<
115+
sycl::ext::intel::experimental::protocol_name::AXI>)>);
116+
117+
// Check that property lists will accept the new properties.
118+
using P = decltype(sycl::ext::oneapi::experimental::properties(
119+
sycl::ext::intel::experimental::min_capacity<0>,
120+
sycl::ext::intel::experimental::ready_latency<1>,
121+
sycl::ext::intel::experimental::bits_per_symbol<2>,
122+
sycl::ext::intel::experimental::uses_valid<true>,
123+
sycl::ext::intel::experimental::uses_ready<false>,
124+
sycl::ext::intel::experimental::in_csr<true>,
125+
sycl::ext::intel::experimental::first_symbol_in_high_order_bits_off,
126+
sycl::ext::intel::experimental::protocol_avalon));
127+
static_assert(sycl::ext::oneapi::experimental::is_property_list_v<P>);
128+
static_assert(
129+
P::has_property<sycl::ext::intel::experimental::min_capacity_key>());
130+
static_assert(
131+
P::has_property<sycl::ext::intel::experimental::ready_latency_key>());
132+
static_assert(
133+
P::has_property<sycl::ext::intel::experimental::bits_per_symbol_key>());
134+
static_assert(
135+
P::has_property<sycl::ext::intel::experimental::uses_valid_key>());
136+
static_assert(
137+
P::has_property<sycl::ext::intel::experimental::uses_ready_key>());
138+
static_assert(P::has_property<sycl::ext::intel::experimental::in_csr_key>());
139+
static_assert(P::has_property<sycl::ext::intel::experimental::
140+
first_symbol_in_high_order_bits_key>());
141+
static_assert(
142+
P::has_property<sycl::ext::intel::experimental::protocol_key>());
143+
144+
static_assert(
145+
P::get_property<sycl::ext::intel::experimental::min_capacity_key>() ==
146+
sycl::ext::intel::experimental::min_capacity<0>);
147+
static_assert(
148+
P::get_property<sycl::ext::intel::experimental::ready_latency_key>() ==
149+
sycl::ext::intel::experimental::ready_latency<1>);
150+
static_assert(
151+
P::get_property<sycl::ext::intel::experimental::bits_per_symbol_key>() ==
152+
sycl::ext::intel::experimental::bits_per_symbol<2>);
153+
static_assert(
154+
P::get_property<sycl::ext::intel::experimental::uses_valid_key>() ==
155+
sycl::ext::intel::experimental::uses_valid<true>);
156+
static_assert(
157+
P::get_property<sycl::ext::intel::experimental::uses_ready_key>() ==
158+
sycl::ext::intel::experimental::uses_ready<false>);
159+
static_assert(P::get_property<sycl::ext::intel::experimental::in_csr_key>() ==
160+
sycl::ext::intel::experimental::in_csr<true>);
161+
static_assert(
162+
P::get_property<sycl::ext::intel::experimental::
163+
first_symbol_in_high_order_bits_key>() ==
164+
sycl::ext::intel::experimental::first_symbol_in_high_order_bits_off);
165+
static_assert(
166+
P::get_property<sycl::ext::intel::experimental::protocol_key>() ==
167+
sycl::ext::intel::experimental::protocol_avalon);
168+
}

0 commit comments

Comments
 (0)