romabo
New ContributorJoined 1 year ago
User Profile
New Contributor
User Widgets
Contributions
Re: HLS i++ style IP generation with oneAPI for RTL IP integration in Quartus
Sorry for the late reply I didn't receive email notifications. The code I used corresponds to the basic example shown in the video you linked that I adapted to single precision floating point multiplication addition. #include <iostream> // oneAPI headers #include <sycl/ext/intel/fpga_extensions.hpp> #include <sycl/sycl.hpp> #include "exception_handler.hpp" using PipeOutProps = decltype(sycl::ext::oneapi::experimental::properties( sycl::ext::intel::experimental::protocol< sycl::ext::intel::experimental::protocol_name::avalon_mm_uses_ready>)); class PipeOutResID; using PipeOutRes = sycl::ext::intel::experimental::pipe<PipeOutResID, float, 0, PipeOutProps>; class MulAddFp32ID; // Kernel name struct MulAddFp32 { float a; float b; float c; void operator()() const { float res = a * b + c; PipeOutRes::write(res); } auto get(sycl::ext::oneapi::experimental::properties_tag) { return sycl::ext::oneapi::experimental::properties{ sycl::ext::intel::experimental::streaming_interface_remove_downstream_stall}; } }; int main() { bool passed = false; try { // Use compile-time macros to select either: // - the FPGA emulator device (CPU emulation of the FPGA) // - the FPGA device (a real FPGA) // - the simulator device #if FPGA_SIMULATOR auto selector = sycl::ext::intel::fpga_simulator_selector_v; #elif FPGA_HARDWARE auto selector = sycl::ext::intel::fpga_selector_v; #else // #if FPGA_EMULATOR auto selector = sycl::ext::intel::fpga_emulator_selector_v; #endif sycl::queue q(selector, fpga_tools::exception_handler, sycl::property::queue::enable_profiling{}); auto device = q.get_device(); // make sure the device supports USM host allocations if (!device.has(sycl::aspect::usm_host_allocations)) { std::cerr << "This design must either target a board that supports USM " "Host/Shared allocations, or IP Component Authoring. " << std::endl; std::terminate(); } std::cout << "Running on device: " << device.get_info<sycl::info::device::name>().c_str() << std::endl; // Kernel inputs float a = 10.1; float b = 20.1; float c = 30.1; q.single_task<MulAddFp32ID>(MulAddFp32{a,b,c}).wait(); float res = PipeOutRes::read(q); // verify that VC is correct passed = true; float expected = a*b + c; if (res != expected) { std::cout << "result " << res << ", expected (" << expected << ") A=" << a << " + B=" << b << " + C=" << c << std::endl; passed = false; } std::cout << (passed ? "PASSED" : "FAILED") << std::endl; } catch (sycl::exception const &e) { // Catches exceptions in the host code. std::cerr << "Caught a SYCL host exception:\n" << e.what() << "\n"; // Most likely the runtime couldn't find FPGA hardware! if (e.code().value() == CL_DEVICE_NOT_FOUND) { std::cerr << "If you are targeting an FPGA, please ensure that your " "system has a correctly configured FPGA board.\n"; std::cerr << "Run sys_check in the oneAPI root directory to verify.\n"; std::cerr << "If you are targeting the FPGA emulator, compile with " "-DFPGA_EMULATOR.\n"; } std::terminate(); } return passed ? EXIT_SUCCESS : EXIT_FAILURE; } Which gives the following estimated resource utilization summary Name ALMs ALUTs FFs MLABs RAMs DSPs Pipe resources 1 2 33 0 0 0 MulAddFp32ID 40 80 146 0 1 1 * -Xstarget=Agilex5 -Xsclock=200MHz -Xsdsp-mode=prefer-dsp The RAM usage is apparently used for cluster logic and extra logic is used for computation as shown below: When running fpga compilation Quartus fitter I get the following resource utilization summary. Name ALMs ALUTs FFs MLABs RAMs DSPs Quartus Fitter: Device Image 129 180 565 0 0 0 MulAddFp32ID 119 165 510 0 0 0 * -Xstarget=Agilex5 -Xsclock=200MHz -Xsdsp-mode=prefer-dsp Let alone the fact that DSP is not inferred by the fitter, I don't really understand the schedule view and what the cluster exactly corresponds to along with the numerous write cycles. I also tried with I/O pipes which appeared to be more suited for my application but I also get a similar operation. // Pipe In struct io_pipe_read_a_id { static constexpr unsigned id = 0; }; struct io_pipe_read_b_id { static constexpr unsigned id = 0; }; struct io_pipe_read_c_id { static constexpr unsigned id = 0; }; using PipeReadA = sycl::ext::intel::kernel_readable_io_pipe<io_pipe_read_a_id, float, 0>; using PipeReadB = sycl::ext::intel::kernel_readable_io_pipe<io_pipe_read_b_id, float, 0>; using PipeReadC = sycl::ext::intel::kernel_readable_io_pipe<io_pipe_read_c_id, float, 0>; // Pipe Out struct io_pipe_write_res_id { static constexpr unsigned id = 1; }; using PipeWriteRes = sycl::ext::intel::kernel_writeable_io_pipe<io_pipe_write_res_id, float, 0>; class MulAddFp32ID; // Kernel name struct MulAddFp32 { auto get(sycl::ext::oneapi::experimental::properties_tag) { return sycl::ext::oneapi::experimental::properties{ sycl::ext::intel::experimental::streaming_interface<>}; } float a = PipeReadA::read(); float b = PipeReadB::read(); float c = PipeReadC::read(); void operator()() const { float res = a * b + c; PipeWriteRes::write(res); } }; @whitepau_altera wrote: If you want to get a similar IP with SYCL HLS, you specify a streaming invocation interface using a kernel property, and streaming data interfaces using pipes. Thank you for confirming. I was mostly interested in the details of what streaming interface/data I should use to get minimal "kernel" operation so as the generate IP would basically operate as an RTL module registers input/outputs with start/done signals. Thank you for your time.
