Skip to content

Commit 1a1fd8d

Browse files
tiwaria1GarveyJoe
andauthored
Add another parameter value for the pipelined kernel interface property (#6966)
The `pipelined` kernel property is defined to control the pipelining behavior. This PR updates the extension to say how a value of `-1` can be used to enforce a particular pipelining behavior. Co-authored-by: Joe Garvey <[email protected]>
1 parent f872048 commit 1a1fd8d

File tree

1 file changed

+20
-9
lines changed

1 file changed

+20
-9
lines changed

sycl/doc/extensions/proposed/sycl_ext_intel_fpga_kernel_interface_properties.asciidoc

Lines changed: 20 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -111,7 +111,7 @@ struct register_map_interface_key {
111111
};
112112

113113
struct pipelined_key {
114-
template <size_t pipeline_directive_or_initiation_interval>
114+
template <int pipeline_directive_or_initiation_interval>
115115
using value_t = sycl::ext::oneapi::properties::property_value<
116116
pipelined_key,
117117
std::integral_constant<size_t, pipeline_directive_or_initiation_interval>>;
@@ -140,7 +140,7 @@ inline constexpr register_map_interface_key::value_t<
140140
register_map_interface_options_enum::do_not_wait_for_done_write>
141141
register_map_interface_do_not_wait_for_done_write;
142142

143-
template<size_t pipeline_directive_or_initiation_interval>
143+
template<int pipeline_directive_or_initiation_interval>
144144
inline constexpr pipelined_key::value_t<
145145
pipeline_directive_or_initiation_interval> pipelined;
146146

@@ -192,22 +192,32 @@ inline constexpr pipelined_key::value_t<
192192
`register_map_interface_do_not_wait_for_done_write`.
193193

194194
|`pipelined`
195-
|An unsigned integer value is accepted as property parameter.
195+
|A signed integer value is accepted as property parameter.
196196

197197
When the parameter is set to a non zero value, the property directs the
198198
compiler to pipeline calls to the kernel such that multiple invocations of the
199-
kernel can be in flight simultaneously. The parameter value also specifies the
200-
minimum number of cycles between successive invocations of the kernel. Example:
199+
kernel can be in flight simultaneously.
200+
201+
When the parameter is a positive value, the value specifies the 'initiation
202+
interval' (II) of the kernel i.e., the minimum number of cycles between successive
203+
invocations. Example:
204+
205+
* `pipelined<N>` - For `N > 0`, the compiler will pipeline multiple kernel
206+
invocations such that an invocation can be launched every `N` cycles if one is
207+
available.
208+
209+
When the parameter is set to `-1`, the compiler will determine the II and
210+
pipeline the kernel invocations.
201211

202-
* `pipelined<N>` - The compiler will pipeline multiple kernel invocations such
203-
that an invocation can be launched every `N` cycles if one is available.
204-
205212
When the parameter is set to `0`, the compiler will not pipeline kernel
206213
invocations.
207214

208215
If the `pipelined` property is not specified, the default behavior is
209216
equivalent to a combination of the property parameter values listed above, but
210217
the choice is implementation defined.
218+
219+
If the property parameter (N) is not specified, the default value is `-1`.
220+
Valid values for `N` are values greater than or equal to `-1`.
211221
|===
212222

213223
Device compilers that do not support this extension may accept and ignore these
@@ -287,4 +297,5 @@ q.single_task(KernelFunctor{a, b, c}).wait();
287297
|========================================
288298
|Rev|Date|Author|Changes
289299
|1|2022-03-01|Abhishek Tiwari|*Initial public working draft*
290-
|========================================
300+
|2|2022-12-05|Abhishek Tiwari|*Make pipelined property parameter a signed int*
301+
|========================================

0 commit comments

Comments
 (0)