Forum Discussion
The weird aggressive aocl optimization "removing unnecessary storage to local memory"
- 7 years ago
With respect to functional verification, what I do is that I construct my host code in a way that both run-time and offline compilation are supported, the latter for FPGAs and the former for other devices, and I use AMD's OpenCL SDK for other devices. In this case, as long as the run-time OpenCL driver is installed, the same host code can then be used to execute the same kernel on any type of CPU, GPU or FPGA. You can take a look at the host code/makefiles of the optimized benchmarks in the following repository as example of achieving this:
https://github.com/fpga-opencl-benchmarks/rodinia_fpga
I emulated all of those kernels on CPUs/GPUs using the same host and kernel codes. What I would tell you is that if an NDRange kernel with sufficiently large local and global size performs correctly on a GPU, it should also perform correctly on an FPGA (unless there is a bug in the FPGA compiler). A CPU should also work fine even if the whole kernel runs on one core, since there will still be multiple threads (work-items) running on that core that could be issued out of order and this is usually enough to show concurrency issues but a GPU would likely be more trustworthy in this case.
With respect to, let's say HDL vs. OpenCL, many old-school HDL programmers tend to think that OpenCL or HLS tools in general are insufficient and it is possible to achieve better results using HDL. This is indeed true in some cases like latency-sensitive or low-power applications where clock-by-clock control over the code is required, or applications that are limited by logic resources, but I would not say this is the case for high-throughput applications where limitation is Memory/PCI-E bandwidth or DSP count since these limitations are independent of the programming language. With respect to the particular case of unpipelinable nested loops, HDL or OpenCL would not make a difference. If you have a regular outer loop with an irregular inner loop, the outer loop cannot be pipelined; it doesn't matter how you "describe" the code. There are two ways to approach such loops on FPGAs:
1- Use NDRange and let the run-time work-item scheduler do its best in maximizing pipeline efficiency and minimizing the average loop II.
2- Collapse the nested loop as long as it is not too irregular and get an II of one at the cost of a noticeable Fmax hit. Though by "collapse" I mean manual collapse and not the compiler's "coalesce" pragma. Take a look at Section 3.2.4.3 in this document:
https://arxiv.org/abs/1810.09773
Even though the provided example involves collapsing a regular nested loop, this optimization also sometimes applies to irregular nested loops. I such case, the condition inside the collapsed loop that is used to increment the variable of the original outer loop will have more than one statement (which complicates the critical path and reduces the Fmax). Indeed the possibility also exists to implement parts of your application in HDL and use it as an HDL library in an OpenCL kernel but you are going to run into complications if your HDL library does not have a fixed latency and I highly doubt you would be able to achieve much better results in the end.
Finally, with respect to NDRange vs. Single Work-item, I recommend reading Section 3.1 (and particularly 3.1.4) of the document I posted above.
Then, you try on 1.1 first.
I've been running it for more than 15 minutes and the errors have not yet appeared, which seems a good sign! Previously the errors happened within 5 minutes after I launched it.
I will let you know once the compilation is finished successfully. For ndrange (#define attr_setup attr_max_wg), the compilation usually needs 4 ~ 5 hours under 17.1.1; for task (single work item) (#define attr_setup attr_task), the compilation usually is much faster and needs 2 hours under 17.1.1.
Here is my configuration script:
sys_quartus_dir=/export/quartus_pro
my_quartus_dir=$HOME
export QSYS_ROOTDIR=${sys_quartus_dir}/17.1.1/qsys/bin
export QUARTUS_ROOTDIR=${sys_quartus_dir}/17.1.1/quartus/bin
export QUARTUS_ROOTDIR_OVERRIDE=${sys_quartus_dir}/17.1.1/quartus
export PATH="${sys_quartus_dir}/17.1.1/quartus/bin/:$PATH"
export PATH="${sys_quartus_dir}/17.1.1/qsys/bin:$PATH"
# PAC_A10 BSP version 1.1
export AOCL_BOARD_PACKAGE_ROOT=/export/fpga/release/a10_gx_pac_ias_1_1_pv/opencl/opencl_bsp
# export PATH="$my_quartus_dir/intelFPGA_pro/19.1/hld/board/a10_ref/:$PATH"
# set OpenCL version 19.1
export ALTERAOCLSDKROOT=$my_quartus_dir/intelFPGA_pro/19.1/hld
export INTELFPGAOCLSDKROOT=$my_quartus_dir/intelFPGA_pro/19.1/hld
export PATH=/homes/hiratz/intelFPGA_pro/19.1/hld/bin:$PATH
#run the OpenCL Setup script in 19.1
source $my_quartus_dir/intelFPGA_pro/19.1/hld/init_opencl.sh
/export/fpga/bin/qsub-fpga- KennyT_altera7 years ago
Super Contributor
Good to hear that, btw, how do you access the Intel Harp machine? I will look into requesting upgrade it to 1.2 version
- hiratz7 years ago
Occasional Contributor
The Intel's Harp environment is documented at: https://wiki.intel-research.net/FPGA.html
To apply for access, you can fill out the form at: https://registration.intel-research.net/register
Once you get an account, you can send emails to [email protected] for help requests.
(Actually I already asked about the bsp 1.2's installation on Intel's Harp community and they replied they probably will update it soon)
Update about my progress:
I successfully compiled the code for both ndrange type and task type. The former took 1:42 hours (1 hour, 42 mins) and the latter took 1:29 hours.
Please note: for the task type compilation, __attribute__((task)) is not anymore supported in 19.1. You have to replace it with __attribute__((max_global_work_dim(0))). In addition, though 17.1.1 allows the code like "get_global_id()" or "get_global_size()" to stay in the code when compiling a task type version, 19.1 does not allow doing so. If you forcibly do so, you will get an obviously abnormal resource utilization report:
Logic utilization (423226%), ALUTs (502536%), Dedicated logic registers (82%), Memory blocks(65%), DSP blocks(2%)
So you may want to remove them and replace the variable "gid" with 0 and the variable "gsize" with 1.
However, the binary (aocx file) cannot work correctly for both emulation and hardware execution.
The emulation shows the binary has some problems:
Using AOCX: decom_comp.aocx Error: Malformed program interface definition found in binary: Error: FAILED to read auto-discovery string at byte 44. Full auto-discovery string value is 19 emulatorhash0000000000000000000000000000 Emul atorDevice 0 0 0 0 3 decomp 0 0 0 0 0 0 0 1 0 10 2 1 8 2 1 8 2 1 8 2 1 8 0 0 4 0 0 4 2 1 8 2 1 8 0 0 4 0 0 4 0 0 0 0 0 1 2147483647 0 1 comp ress 0 0 0 0 0 0 0 1 0 13 2 1 8 2 1 8 0 0 4 0 0 4 0 0 4 2 1 8 2 1 8 0 0 4 0 0 4 2 1 8 2 1 8 0 0 4 0 0 4 0 0 0 0 0 1 2147483647 0 1 merge_str eams 0 0 0 0 0 0 0 1 0 8 2 1 8 2 1 8 2 1 8 2 1 8 2 1 8 0 0 4 0 0 4 0 0 4 0 0 0 0 0 1 2147483647 0 1 Error: Invalid binary ERROR: CL_INVALID_BINARY Location: ../common-fpga/src/AOCLUtils/opencl.cpp:392 Failed to create program with binaryThe hardware version's launching shows similar errors. Some additional messages are:
acl_hal_mmd.c:1393:assert failure: Failed to initialize kernel interface acl_hal_mmd.c:1393: l_try_device: Assertion `0' failed.So I think there may be some mismatch across Quartus 17.1.1, OpenCL RTE 19.1 and the underlying bsp 1.1
So for now, it is better for me to stick with 17.1.1