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

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

5.5.5.4. Implementing Pipe Reads

The read_pipe API call allows you to receive data across a pipe.

Intel® only supports the convenience version of the read_pipe function. By default, read_pipe calls are non-blocking.

To implement a pipe read, include the following read_pipe function signature:
int read_pipe (read_only_pipe <type> pipe_id, <type> *data);

Where:

  • pipe_id identifies the buffer to which the pipe connects, and it must match the pipe_id of the corresponding pipe write operation (write_pipe).
  • data is the data that the pipe read operation reads from the pipe. It is a pointer to the location of the data.
    Note: read_pipe call might lead to a global or local memory load, depending on the source address space of the data pointer.
  • <type> defines the packet size of the data.
The following code snippet demonstrates the implementation of the read_pipe API call:
/*Declares the read_only_pipe that contains packets
of type long.*/
/*Declares that read_pipe calls within the kernel will exhibit
blocking behavior*/
__kernel void kernel_read_pipe(__global long *dst,
                               read_only pipe long __attribute__((blocking)) p)
{
    for (int i = 0; i < N; i++)
    {
        /*Reads from a long from the pipe and stores it
        into global memory at the specified location*/
        read_pipe(p, &dst[i]);
    }
}

To facilitate better hardware implementations, Intel® provides facility for blocking read_pipe calls by specifying the blocking attribute (that is, __attribute__((blocking))) on the pipe argument declaration for the kernel. Blocking read_pipe calls always return success.

CAUTION:
If the pipe is empty (that is, if the FIFO buffer is empty), you cannot receive data across a blocking read pipe using the read_pipe API call. Doing so causes your kernel to stall.