[SYCL][FPGA] Initial implementation of pipes feature #292

MrSidims · 2019-07-08T14:47:28Z

No description provided.

keryell

This looks like a super-useful feature on FPGA.
I have nevertheless a few remarks.

clang/lib/Sema/SemaType.cpp

sycl/doc/extensions/SYCL_pipes_lowering_to_SPIRV.rst

keryell · 2019-07-08T21:32:01Z

sycl/doc/extensions/SYCL_pipes_lowering_to_SPIRV.rst

+OpenCL 2.2 program pipe representation in SPIR-V
+================================================
+
+The SPIR-V program pipe representation is leveraged, to be an underlying


Spurious ,.
I do not know if "leveraged" is good here. Do you improve SPIR-V or you just recycle it for another purpose?

SPIR-V representation is re-used.

Then clarify...

sycl/doc/extensions/SYCL_pipes_lowering_to_SPIRV.rst

keryell · 2019-07-09T00:03:08Z

sycl/doc/extensions/SYCL_pipes_lowering_to_SPIRV.rst

+
+Global pipe creation itself is being done by SPIR-V built-in
+__spirv_CreatePipeFromPipeStorage_{read|write} which has no OpenCL
+representation and therefore stays in IR before and after SPIR-V tool-chain as:


I have rewrote this one, please check it.

sycl/include/CL/sycl/pipes.hpp

keryell · 2019-07-09T00:11:56Z

sycl/include/CL/sycl/pipes.hpp

+
+  // Writing to pipe is lowered to SPIR-V instruction OpWritePipe via SPIR-V
+  // friendly LLVM IR.
+  static void write(dataT Data, bool &Success) {


static bool write(const dataT &Data) instead?

The const reference is good. The status return code move to returned bool makes the signatures of read and write differ more than I would like. Read returns the data type in both cases. How strongly do you think that write should return bool success status, while read returns the data and takes a status ref arg?

sycl/include/CL/sycl/pipes.hpp

Ruyk · 2019-07-09T09:15:27Z

sycl/test/fpga_tests/fpga_pipes.cpp

+        while (!SuccessCode)
+          write_acc[idx] = Pipe::read(SuccessCode);
+      });
+    });


Is there any user-friendly documentation on how this Pipe extension works?
One of the things that confuses me is that, in the test above, there is nothing preventing a SYCL runtime to execute goo_nb before foo_nb, since the intersection of the requirements is empty and SYCL queues are out-of-order.

Is that relevant in your case? I figure for an FPGA it doesn't matter, but there is nothing in this extension that suggest it shouldn't work on anything that is not an FPGA but supports pipes.

Is there any user-friendly documentation on how this Pipe extension works?

It's going to be published. Don't really know when.
There is SYCL 2.2 spec (which is deprecated, but yet can provide some information). Also this implementation doesn't follows it.

I figure for an FPGA it doesn't matter, but there is nothing in this extension that suggest it shouldn't work on anything that is not an FPGA but supports pipes.

Yeah, I shall put some code in device information quarrying.

We'll publish the extension document soon in this repo

The problem is to come with a terse API and at the same time not giving up any performance on the table...
Perhaps a more explicit naming in the case of (non-)blocking? Adding a tag to pick another overload? Adding a template parameter?...
I understand that relying on an accessor to encode the blocking behavior like in SYCL 2.2 might be a too heavy syntax....

The pipes extension at SYCL level is described in the document, currently as a pending PR, at:
#635

Ah thank you for the cross-reference. It prevents me from looking for it...

sycl/test/fpga_tests/fpga_pipes.cpp

sycl/include/CL/sycl/pipes.hpp

asavonic

Current documentation mostly covers SPIR-V and LLVM IR levels.
Please describe how this feature is represented and used (with examples) at SYCL level.

clang/include/clang/Basic/TokenKinds.def

keryell

Interesting!
Still a few remarks.

clang/lib/Parse/ParseDecl.cpp

sycl/doc/extensions/SYCL_pipes_usage.rst

keryell · 2019-07-11T18:09:38Z

sycl/doc/extensions/SYCL_pipes_usage.rst

+    cgh.parallel_for<class foo_nb>(range<1> { dataSize }, [=](id<1> idx) {
+      bool SuccessCode = false;
+      while (!SuccessCode)
+        Pipe::write(42, SuccessCode); // Write '42' into a some_pipe allocated


What is the memory model of a pipe in the case you are executing it in a parallel_for and not a single_task?

Do you expect some independent-forward-progress here in the case dataSize > 1 (the capacity of the pipe...)?

Yes, I expect that. I should re-write the example code into:
cgh.parallel_for<class bar>(range<1> { 1024 }, [=](id<1> idx) { write_acc[idx] = my_pipe::read(); });

OpenCL analog of the code above is:
__kernel void foo(__read_only pipe uint in_pipe, __global uint *dst) { int gid = get_global_id(0); read_pipe(in_pipe, &dst[gid]); }

But it is illegal OpenCL code in general. You have to use pipe reservation for this and and I am not sure we want this anyway on FPGA...
So the assumptions are probably to be described in the document describing the extension.
What happens if there are multiple work-group on different CU using the same pipe input or output?

But it is illegal OpenCL code in general. You have to use pipe reservation for this and and I am not sure we want this anyway on FPGA...

You are right.

What happens if there are multiple work-group on different CU using the same pipe input or output?

... to be described in the document describing the extension :) I added a quote from there:

Accesses to a single pipe within two work-items of the same kernel, and/or threads of the host program: No ordering guarantees are made on the order of pipe operations across device work-items or host threads. For example, if two work-items executing a kernel write to a pipe, there are no guarantees that the work-item with lower id (for any definition of id) executes before the pipe write from a higher id. The execution order of work-items executing a kernel are not defined by SYCL, may be dynamically reordered, and may not be deterministic. If ordering guarantees are required across work-items and/or host threads, synchronization mechanisms such as atomics or barriers must be used.

@keryell, are you okay to merge this PR?

Yes I am fine with merging.
It looks like in #318 there are some corrections that will be done later.

keryell · 2019-07-11T18:11:31Z

sycl/doc/extensions/SYCL_pipes_usage.rst

+Simple example of SYCL program
+==============================
+.. code:: cpp
+  using Pipe = pipe<class some_pipe, int, 1>;


Is it legal to have a local class some_pipe declaration like here?

If I don't try to use it after - yes. It's a forward declaration that adds appropriate mangling for fields and methods of the template class pipe.

sycl/doc/extensions/SYCL_pipes_usage.rst

sycl/test/fpga_tests/fpga_pipes.cpp

clang/lib/Parse/ParseTentative.cpp

bader · 2019-07-12T11:00:48Z

clang/lib/Parse/ParseDecl.cpp

        // OpenCL 2.0 defined this keyword. OpenCL 1.2 and earlier should
        // support the "pipe" word as identifier.
        Tok.getIdentifierInfo()->revertTokenIDToIdentifier();
        goto DoneWithDeclSpec;
      }
      isInvalid = DS.SetTypePipe(true, Loc, PrevSpec, DiagID, Policy);
      break;
+    case tok::kw___pipe:
+      // __pipe is defined only for SYCL kernel language


Shouldn't we add a check for that?

No. If this keyword appears in any other language, clang already will report it be unknown keyword.

Yes, but if there is __pipe used in an existing library, unrelated to SYCL?
But do we need such a keyword?

Fixed.

But do we need such a keyword?

Will be reworked as attribute. Or you suggest to do some forward declaration expecting the translator to do its work?

I think the attribute is fine.

clang/lib/Sema/SemaDecl.cpp

sycl/include/CL/sycl/pipes.hpp

clang/lib/Sema/SemaDecl.cpp

Documentation has been added.

sycl/include/CL/sycl/pipes.hpp

bader · 2019-07-14T19:43:44Z

@keryell, @Ruyk, @asavonic, please, "resolve conversations", which do not require follow-up actions.
I'd like to merge this pull request, but it's hard to say if all issues are resolved or not - there are 97 comment at the moment.

I think there are at least three issues left to resolve (updated by @keryell last Friday).

Pipe types construction rely on keywords mechanism. Consider following code: __read_only __pipe int p; It declares a read only pipe with integer layout. In LLVM IR this type is seen as 'pipe_ro_t'. Write only pipe type will be looking like 'pipe_wo_t'. Signed-off-by: Dmitry Sidorov <dmitry.sidorov@intel.com>

Signed-off-by: Dmitry Sidorov <dmitry.sidorov@intel.com>

Pipe are supported via generation SPIR-V friendly IR and lowering it directly to SPIR-V code. non-blocking pipes are supported; blocking pipes are supported as a workaround. Signed-off-by: Dmitry Sidorov <dmitry.sidorov@intel.com>

bader · 2019-07-15T17:22:02Z

@keryell, are you okay to merge this PR?

Signed-off-by: Michael Kinsner <michael.kinsner@intel.com>

keryell reviewed Jul 9, 2019

View reviewed changes

Ruyk reviewed Jul 9, 2019

View reviewed changes

sycl/test/fpga_tests/fpga_pipes.cpp Outdated Show resolved Hide resolved

AlexeySachkov reviewed Jul 9, 2019

View reviewed changes

sycl/include/CL/sycl/pipes.hpp Outdated Show resolved Hide resolved

asavonic previously requested changes Jul 10, 2019

View reviewed changes

Fznamznon reviewed Jul 11, 2019

View reviewed changes

clang/include/clang/Basic/TokenKinds.def Outdated Show resolved Hide resolved

MrSidims force-pushed the private/MrSidims/GlobalPipes_WOCL branch 2 times, most recently from f97d84a to 30ae85f Compare July 11, 2019 15:16

keryell reviewed Jul 11, 2019

View reviewed changes

MrSidims force-pushed the private/MrSidims/GlobalPipes_WOCL branch from fffd69f to c627b78 Compare July 12, 2019 08:57

MrSidims requested a review from asavonic July 12, 2019 10:01

bader reviewed Jul 12, 2019

View reviewed changes

bader self-assigned this Jul 12, 2019

MrSidims force-pushed the private/MrSidims/GlobalPipes_WOCL branch 2 times, most recently from 81c2eff to ec15072 Compare July 12, 2019 12:10

bader reviewed Jul 12, 2019

View reviewed changes

clang/lib/Sema/SemaDecl.cpp Outdated Show resolved Hide resolved

MrSidims force-pushed the private/MrSidims/GlobalPipes_WOCL branch 2 times, most recently from fdacf2e to ad598c1 Compare July 12, 2019 12:50

bader previously approved these changes Jul 12, 2019

View reviewed changes

asavonic reviewed Jul 12, 2019

View reviewed changes

sycl/include/CL/sycl/pipes.hpp Show resolved Hide resolved

MrSidims added 2 commits July 15, 2019 11:22

[SYCL][FPGA] Declare pipe types and SPIR-V friendly functions

63c738d

Signed-off-by: Dmitry Sidorov <dmitry.sidorov@intel.com>

MrSidims dismissed bader’s stale review via a5f9027 July 15, 2019 08:36

MrSidims force-pushed the private/MrSidims/GlobalPipes_WOCL branch from ad598c1 to a5f9027 Compare July 15, 2019 08:36

bader previously approved these changes Jul 15, 2019

View reviewed changes

[SYCL][FPGA] Initial implementation of pipe feature

cb9c774

Pipe are supported via generation SPIR-V friendly IR and lowering it directly to SPIR-V code. non-blocking pipes are supported; blocking pipes are supported as a workaround. Signed-off-by: Dmitry Sidorov <dmitry.sidorov@intel.com>

MrSidims dismissed bader’s stale review via cb9c774 July 15, 2019 09:52

MrSidims force-pushed the private/MrSidims/GlobalPipes_WOCL branch from a5f9027 to cb9c774 Compare July 15, 2019 09:52

bader approved these changes Jul 15, 2019

View reviewed changes

bader mentioned this pull request Jul 15, 2019

[SYCL][FPGA] Rework blocking pipes #318

Merged

bader merged commit 74527fc into intel:sycl Jul 15, 2019

mkinsner pushed a commit to mkinsner/llvm that referenced this pull request Nov 13, 2019

Const ref to data in write methods, as suggested in intel#292 (comment)

589224c

Signed-off-by: Michael Kinsner <michael.kinsner@intel.com>

[SYCL][FPGA] Initial implementation of pipes feature #292

[SYCL][FPGA] Initial implementation of pipes feature #292

Uh oh!

Conversation

MrSidims commented Jul 8, 2019

Uh oh!

keryell left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

asavonic left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

keryell left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

MrSidims Jul 15, 2019 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

MrSidims Jul 15, 2019 •

edited

Loading

MrSidims Jul 15, 2019 •

edited

Loading