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

ID 683846
Date 12/19/2022
Document Table of Contents

5.5.2. Pipe Data Behavior

Data written to a pipe remains in a pipe provided that the kernel program remains loaded on the FPGA device. In other words, data written to a pipe persists across multiple work-groups and NDRange invocations. However, data is not persistent across multiple or different invocations of kernel programs that result in FPGA reprogramming operations.

Consider the following code example:

__kernel void producer (write_only pipe uint __attribute__((blocking)) c0)
    for (uint i = 0; i < 10; i++)
        write_pipe (c0, &i);

__kernel void consumer (__global uint * restrict dst,
                        read_only pipe uint __attribute__((blocking))
                        __attribute__((depth(10))) c0)
    for (int i = 0; i < 5; i++)
        read_pipe (c0, &dst[i]);

A read operation to a pipe reads the least recent piece of data written to the pipe first. Pipe data maintains a FIFO ordering within the pipe.

Figure 12. Pipe Data FIFO Ordering

The kernel producer writes ten elements ([0, 9]) to the pipe. The kernel consumer reads five elements from the pipe per NDRange invocation. During the first invocation, the kernel consumer reads values 0 to 4 from the pipe. Because the data persists across NDRange invocations, the second time you execute the kernel consumer, it reads values 5 to 9.

For this example, to avoid a deadlock from occurring, you need to invoke the kernel consumer twice for every invocation of the kernel producer. If you call consumer less than twice, producer stalls because the pipe becomes full. If you call consumer more than twice, consumer stalls because there is insufficient data in the pipe.