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

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...

Unclear.

I have rewrote this one, please check it.

static bool write(const dataT &Data) instead?

Thanks!

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?

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...

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.

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.

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.