search

triSYCL / sycl

SYCL for Vitis: Experimental fusion of triSYCL with Intel SYCL oneAPI DPC++ up-streaming effort into Clang/LLVM
Other
107 stars 19 forks source link

Is barrier in kernel supported or going to be supported? #180

Open fxzjshm opened 2 years ago

fxzjshm commented 2 years ago

barrier() of sycl::nd_item seems not working, the log says:

===>The following messages were generated while  performing high-level synthesis for kernel: geIXT1_EEERKT0_EUlvE_85uG06fiV Log file: /tmp/dedisp-1337dez7lrhv1u/vxx_comp_tmp/geIXT1_EEERKT0_EUlvE_85uG06fiV/geIXT1_EEERKT0_EUlvE_85uG06fiV/vitis_hls.log :
ERROR: [v++ 214-194] in function 'geIXT1_EEERKT0_EUlvE_85uG06fiV': Undefined function ControlBarrier
ERROR: [v++ 214-135] Syn check fail!
ERROR: [v++ 200-1715] Encountered problem during source synthesis
ERROR: [v++ 60-300] Failed to build kernel(ip) geIXT1_EEERKT0_EUlvE_85uG06fiV, see log for details: /tmp/dedisp-1337dez7lrhv1u/vxx_comp_tmp/geIXT1_EEERKT0_EUlvE_85uG06fiV/geIXT1_EEERKT0_EUlvE_85uG06fiV/vitis_hls.log
ERROR: [v++ 60-773] In '/tmp/dedisp-1337dez7lrhv1u/vxx_comp_tmp/geIXT1_EEERKT0_EUlvE_85uG06fiV/geIXT1_EEERKT0_EUlvE_85uG06fiV/vitis_hls.log', caught Tcl error: ERROR: [HLS 214-194] in function 'geIXT1_EEERKT0_EUlvE_85uG06fiV': Undefined function ControlBarrier
ERROR: [v++ 60-773] In '/tmp/dedisp-1337dez7lrhv1u/vxx_comp_tmp/geIXT1_EEERKT0_EUlvE_85uG06fiV/geIXT1_EEERKT0_EUlvE_85uG06fiV/vitis_hls.log', caught Tcl error: ERROR: [HLS 214-135] Syn check fail!
ERROR: [v++ 60-773] In '/tmp/dedisp-1337dez7lrhv1u/vxx_comp_tmp/geIXT1_EEERKT0_EUlvE_85uG06fiV/geIXT1_EEERKT0_EUlvE_85uG06fiV/vitis_hls.log', caught Tcl error: ERROR: [HLS 200-1715] Encountered problem during source synthesis
ERROR: [v++ 60-599] Kernel compilation failed to complete
ERROR: [v++ 60-592] Failed to finish compilation

Environment: Vitis 2021.2 & 2022.1, compiling with -fsycl-targets=fpga64_hls_hw_emu

keryell commented 2 years ago

It used to work with the OpenCL/SPIR back-end since it was implemented by Vitis OpenCL but we deprecated it and we have moved to HLS back-end which is really single-task oriented, more suitable for FPGA. The problem with parallel_for and barriers is that at some point you need a GPU thread emulator, requiring some memory to store the thread states and emulate the context switches. We could implement the compiler transformations to do it, but at the end this kind of very GPU-friendly code will probably work well on GPU but not with good performance on FPGA. So this is not our priority, in comparison to implementing features really suitable for FPGA. But what is the kind of applications are trying to run on FPGA?

fxzjshm commented 2 years ago

Well, I know barriers and mem_fences isn't designed for FPGA, so I'm just asking for a possibility.
I was porting some CUDA-based scientific research program into SYCL (mainly on GPU and CPU), and my teacher let me try if it can be ported to FPGA as well. I will re-design it for FPGA then (well, if a can).
Thank for explanation!