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

ID 683846
Date 10/04/2021
Public

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

Document Table of Contents

5.5.3.1. Work-item Serial Execution of Pipes

Work-item serial execution refers to an ordered execution behavior where work-item sequential IDs determine their execution order in the compute unit.

When you implement pipes in a kernel, the Intel® FPGA SDK for OpenCL™ Offline Compiler enforces that kernel behavior is equivalent to having at most one work-group in flight. The offline compiler also ensures that the kernel executes pipes in work-item serial execution, where the kernel executes work items with smaller IDs first. A work-item has the identifier (x, y, z, group), where x, y, z are the local 3D identifiers, and group is the work-group identifier.

The work-item ID (x0, y0, z0, group0) is considered to be smaller than the ID (x1, y1, z1, group1) if one of the following conditions is true:

  • group0 < group1
  • group0 = group1 and z0 < z1
  • group0 = group1 and z0 = z1 and y0 < y1
  • group0 = group1 and z0 = z1 and y0 = y1 and x0 < x1

Work items with incremental IDs execute in a sequential order. For example, the work-item with an ID (x0, y0, z0, group0) executes the write channel call first. Then, the work-item with an ID (x1, y0, z0, group0) executes the call, and so on. Defining this order ensures that the system is verifiable with external models.

Pipe Execution in a Loop with Multiple Work Items

Only in kernels compiled as an NDRange kernel, when pipes exist in the body of a loop with multiple work items, as shown below, each loop iteration executes prior to subsequent iterations. This implies that loop iteration 0 of each work item in a work group executes before iteration 1 of each work item in a work group, and so on.

__kernel void ordering (__global int * data,
                        write_only pipe int __attribute__((blocking)) req)
{
    write_pipe (req, &data[get_global_id(0)]);
}