Inter-Kernel and Kernel-Host Pipes in one design

[Ric2]

Hi, we are transfering a HLS IP core design to the new sycl HLS tool flow. During this migration several issues came up. We have one of them isolated here and generated a reproducer (please see attached file). The issue comes while mixing Inter-Kernel and Kernel-Host Pipes with DIFFERENT data types. The compiler generates EMU when datatype is the same but crashes (icpx: error: fpga compiler command failed with exit code 245) when datatype is different. Please notice that this is only for EMU. SIM and FPGA targets run through without issues and gererate expected results.
The reproducer is based on Pipes Sample and modified to add Host Pipes.
Any suggestion how to get different data types emulated? Thanks!

oneAPI DPC++/C++ Compiler 2024.1.0 (2024.1.0.20240308), Ubuntu 22.04.4 LTS

[Ric2]

Corrected attachement!
Sorry for the confusion, I forgot to redo some test code changes. Corrected now. Issue still as described above.

[BoonBengT_Altera]

Hi @Ric2,

Thank you for posting in Intel community forum and hope all is well.

Noted on the attached example, please do allow us some time to compile the mention example and will get back with you on updates or clarification.

Appreciate the patients.

Best Wishes

BB

[BoonBengT_Altera]

Hi @Ric2,

Some updates, we hit into some issues in compiling the provided design however was able to compiled them after.

Hence we are working internally on this and will get back to you as soon as possible.

Please rest assure your request has been looked at.

Thank you for the patients.

Best Wishes

BB

[whitepau_altera · Accepted solution]

This issue was successfully reproduced.

You can use the sycl::ext::intel::experimental::pipe as an inter-kernel pipe or as a host pipe. I was able to get your code to compile by changing ProducerToConsumerPipe as so:

using ProducerToConsumerPipe = ext::intel::experimental::pipe< // Defined in the SYCL headers.
    class ProducerConsumerPipeId,                // An identifier for the pipe.
    INTRA_KERNEL_TYPE,                           // The type of data in the pipe.
    4>;                                          // The capacity of the pipe.

It's odd that you can't have a mixture of experimental pipes and non-experimental pipes; we've opened a bug report internally on this. Thanks for reporting!

[BoonBengT_Altera]

Hi @Ric2,

Great! Good to know that your doubts has been clarify, with no further clarification on this thread, it will be transitioned to community support for further help on doubts in this thread. Please login to ‘https://supporttickets.intel.com’, view details of the desire request, and post a feed/response within the next 15 days to allow me to continue to support you. After 15 days, this thread will be transitioned to community support.

Thank you for the questions and as always pleasure having you here.

Best Wishes

BB

p/s: If any answer from the community or Intel Support are helpful, please feel free to give best answer or rate 4/5 survey.