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

ID 683342
Date 4/22/2019
Public
Document Table of Contents

5.5.5.3. Implementing Pipe Writes

The write_pipe API call allows you to send data across a pipe.

Intel® only supports the convenience version of the write_pipe function. By default, write_pipe calls are nonblocking. Pipe write operations are successful only if there is capacity in the pipe to hold the incoming packet.

To implement a pipe write, include the following write_pipe function signature:
int write_pipe (write_only pipe <type> pipe_id, const <type> *data);

Where:

pipe_id identifies the buffer to which the pipe connects, and it must match the pipe_id of the corresponding read pipe (read_pipe).

data is the data that the pipe write operation writes to the pipe. It is a pointer to the packet type of the pipe. Note that writing to the pipe might lead to a global or local memory load, depending on the source address space of the data pointer.

<type> defines a pipe data width. The return value indicates whether the pipe write operation is successful. If successful, the return value is 0. If pipe write is unsuccessful, the return value is -1.

The following code snippet demonstrates the implementation of the write_pipe API call:
/*Declares the writable nonblocking pipe, p, which contains packets of type int*/
__kernel void kernel_write_pipe (__global const long *src,
                                 write_only pipe int p)
{
    for (int i = 0; i < N; i++)
    {
        //Performs the actual writing
		//Emulates blocking behavior via the use of a while loop
		while (write_pipe(p, &src[i]) < 0) { }
    }
}

The while loop is unnecessary if you specify a blocking attribute. To facilitate better hardware implementations, Intel® provides facility for blocking write_pipe calls by specifying the blocking attribute (that is, __attribute__((blocking))) on the pipe argument declaration for the kernel. Blocking write_pipe calls always return success.

CAUTION:
When you send data across a blocking write pipe using the write_pipe API call, keep in mind that if the pipe is full (that is, if the FIFO buffer is full of data), your kernel will stall. Use the Intel® FPGA dynamic profiler for OpenCL™ to check for pipe stalls.