Skip to content

Conversation

@tiwaria1
Copy link
Contributor

@tiwaria1 tiwaria1 commented Mar 3, 2022

SYCL extension contains the following new kernel properties

  • streaming_interface<...>
  • register_map_interface<...>
    The first two properties take enum arguments that provide the compiler information about whether the logic downstream to the kernel will back-pressure the kernel or not.
  • pipelined<N>
    Takes an integer, non-zero values specify minimum cycles between kernel invocations, and 0 specifies that pipelining should be disabled.

@tiwaria1 tiwaria1 requested a review from GarveyJoe March 3, 2022 00:42
|`ip_interface_streaming`
|The `ip_interface_streaming` property adds the requirement that the kernel must
have dedicated ports for input / output signals in the generated
Register Transfer Level (RTL) module.
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I would drop the mention of RTL. It confuses things to refer to the description of the hardware as being separate from the hardware itself.
Also, I think your intent is for this to apply to both control and data signals. Can you state the explicitly? And assuming that this should apply to data signals as well, should it apply to all inputs/outputs from the kernel, i.e. pipes as well, or only kernel arguments?

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Added explicit statement. Specified that it should apply to kernel arguments and control flow signals only. Including pipes seems to make sense. I need to think a bit more about including them.

*a = *b + *c;
}

auto get(sycl::ext::oneapi::experimental::properties_tag) {
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Is this a magic function that encode the properties? How it is connected to the single_task below?

|`pipelined_kernel`
| The `pipelined_kernel` property directs the compiler to pipeline calls to the
kernel so that multiple invocations of the kernel can be in flight
simultaneously.
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Do you mean several instances of a single_task or this is pipelining at the work-item level?

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This refers to several instances of a kernel, most commonly a single_task kernel. It does not refer to pipelining of work-items. However, there's no reason this couldn't be applied to a parallel_for kernel to request that successive launches of that kernel be pipelined in addition to the usual work item pipelining. In such a kernel, the work items could be injected into the datapath one after another and, after the last work item from a given kernel invocation has been injected, the first work-item of the next invocation could be injected without waiting for the pipeline to drain.

I don't see much value in a control to enable or disable pipelining at the work-item level because that's nearly always beneficial. This invocation pipelining is more situational though as it only provides value if the user intends to launch the same kernel multiple times.

@tiwaria1 tiwaria1 changed the title SYCL extension for FPGA kernel interface properties [SYCL][DOC] Add extension for FPGA kernel interface properties Apr 12, 2022
@tiwaria1 tiwaria1 marked this pull request as ready for review April 19, 2022 03:31
@tiwaria1 tiwaria1 requested a review from a team as a code owner April 19, 2022 03:31
@tiwaria1
Copy link
Contributor Author

ping @intel/dpcpp-specification-reviewers

steffenlarsen
steffenlarsen previously approved these changes Jun 6, 2022
Copy link
Contributor

@steffenlarsen steffenlarsen left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

LGTM!

Copy link
Contributor

@keryell keryell left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

pipelined could be simplified to take a single parameter with value false, true or some II value.
The current cartesian product with 2 parameters is useless since false and an II is meaningless.

@keryell
Copy link
Contributor

keryell commented Jun 11, 2022

Also having a default true value for the pipelined when used seems an interesting shortcut.

the following extensions:

- link:../experimental/sycl_ext_oneapi_properties.asciidoc[sycl_ext_oneapi_properties]
- link:sycl_ext_oneapi_kernel_properties.asciidoc[sycl_ext_oneapi_kernel_properties]
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Links are broken

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

They seem to be working now.
oneapi properties is in experimental folder: https://github.com/intel/llvm/tree/sycl/sycl/doc/extensions/experimental
and oneapi kernel properties will be in the proposed folder after 5755 PR goes in : https://github.com/intel/llvm/pull/5755/files

@tiwaria1 tiwaria1 requested a review from steffenlarsen August 11, 2022 19:15
@tiwaria1
Copy link
Contributor Author

@steffenlarsen I addressed the remaining comments. Please review and merge this PR. Based on your namespace comment on a different PR I have updated the namespace here to be sycl::ext::intel::experimental since this is a FPGA specific extension.

Copy link
Contributor

@steffenlarsen steffenlarsen left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Just a few more comments, then I think I'm happy!

@tiwaria1
Copy link
Contributor Author

tiwaria1 commented Sep 7, 2022

Just a few more comments, then I think I'm happy!

Hey Steffen, I addressed the other comments. Please take a look.

Copy link
Contributor

@steffenlarsen steffenlarsen left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

LGTM!

@steffenlarsen steffenlarsen merged commit 4b6bd14 into intel:sycl Sep 14, 2022
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

4 participants