Inter-Kernel and Kernel-Host Pipes in one design
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
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!