Intel® FPGA SDK for OpenCL™ Pro Edition: Programming Guide

ID 683846
Date 9/26/2022

A newer version of this document is available. Customers should click here to go to the newest version.

Document Table of Contents

7.27. Pipelining Loops in Non-task Kernels (-auto-pipeline)

To direct the Intel® FPGA SDK for OpenCL™ Offline Compiler to compile your design and pipeline loops in non-task (NDRange) kernels, include the -auto-pipeline option in your aoc command. The host program invokes non-task kernels through the kernel execution function clEnqueueNDRangeKernel.

At a command prompt, invoke the aoc -auto-pipeline <your_kernel_filename1>.cl [<your_kernel_filename2>.cl ...] command, where [ <your_kernel_filename2>.cl ...] are the optional space-delimited file names of kernels that you can compile in addition to <your_kernel_filename1>.cl.

Without the -auto-pipeline option, the compiler does not pipeline the loops in non-task kernels, but instead executes different work items in parallel. With the -auto-pipeline option, the compiler attempts to pipeline the loops in your design, but the pipelining is not guaranteed.

Note: The -auto-pipeline option might improve or degrade performance depending on the memory access pattern in your design.
  • If the auto-pipelining is successful, the Loop Analysis report displays the message ND-Range as task. The compiler-generated loops appear marked as Compiler generated NDRange-as-task loop in the report.
  • If the compiler chooses not to auto-pipeline the loops, the Loop Analysis report displays a message for the kernel. The reasons for not auto-pipelining a loop can be one of the following:
    • A barrier in the function is not at the top-level function scope.
    • Kernel uses a local or private memory.
    • Kernel uses a volatile or atomic memory or channels.
    • Kernel uses the num_compute_units(X,Y,Z) attribute.
Tip: If you do not want the compiler to pipeline some infrequently used loops while allowing other loops to be auto-pipelined, use the disable_loop_pipelining pragma on specific loops when using the -auto-pipeline option. This pragma disables the loop pipelining.