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

ID 683846
Date 3/28/2022
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. Multiple Work-Item Ordering for Pipes

The OpenCL™ specification does not define a work-item ordering. The Intel® FPGA SDK for OpenCL™ enforces a work-item order to maintain the consistency in pipe read and write operations.

Multiple work-item accesses to a pipe can be useful in some scenarios. For example, they are useful when data words in the pipe are independent, or when the pipe is implemented for control logic. The main concern regarding multiple work-item accesses to a pipe is the order in which the kernel writes data to and reads data from the pipe. If possible, the OpenCL pipes process work-items read and write operations to a pipe in a deterministic order. As such, the read and write operations remain consistent across kernel invocations.

Requirements for Deterministic Multiple Work-Item Ordering

To guarantee deterministic ordering, the SDK checks that the pipe call is work-item invariant based on the following characteristics:

  • All paths through the kernel must execute the pipe call.
  • If the first requirement is not satisfied, none of the branch conditions that reach the pipe call should execute in a work-item-dependent manner.

If the SDK cannot guarantee deterministic ordering of multiple work-item accesses to a pipe, it warns you that the pipes might not have well-defined ordering with nondeterministic execution. Primarily, the SDK fails to provide deterministic ordering if you have work-item-variant code on loop executions with pipe calls, as illustrated below:

__kernel void
ordering (__global int * check, global int * data,
          write_only pipe int __attribute__((blocking)) req)
{
    int condition = check[get_global_id(0)];

    if (condition)
    {
        for (int i = 0; i < N; i++)
        {
            process(data);
            write_pipe (req, &data[i]);
        }
    }
	else
    {
        process(data);
    }
}