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.4.5.3. Implementing Blocking Channel Reads

The read_channel_intel API call allows you to receive data across a channel.
To implement a blocking channel read, include the following read_channel_intel function signature:
<type> read_channel_intel(channel <type> channel_id);

Where:

channel_id identifies the buffer to which the channel connects, and it must match the channel_id of the corresponding write channel (write_channel_intel).

<type> defines a channel data width. Ensure that the variable the kernel assigns to read the channel data is convertible from <type>.

The following code snippet demonstrates the implementation of the read_channel_intel API call:
//Defines chan, a kernel file-scope channel variable.
channel long chan;

/*Defines the kernel, which reads eight bytes (size of long) from the channel and writes it back to global memory.*/
__kernel void kernel_read_channel (__global long * dst); {
  for (int i = 0; i < N; i++) {
     //Reads the eight bytes from the channel.
     dst[i] = read_channel_intel(chan);
  }
}
CAUTION:
If the channel is empty (that is, if the FIFO buffer is empty), you cannot receive data across a read channel using the read_channel_intel API call. Doing so causes your kernel to stall until at least one element of data becomes available from the FIFO buffer.

Implementing Nonblocking Channel Reads

Perform nonblocking reads to facilitate applications where data is not always available and the operation should not wait for data to become available.
The nonblocking read signature is similar to a blocking read. However, it populates the address pointed to by the bool pointer valid indicating whether a read operation successfully read data from the channel.

On a successful read (valid set to true), the value read from the channel is returned by the read_channel_nb_intel function. On a failed read (valid set to false), the return value of the read_channel_nb_intel function is not defined.

To implement a blocking channel read, use the following read_channel_nb_intel function signature:
<type> read_channel_nb_intel(channel <type> channel_id, bool * valid);
The following code snippet demonstrates the use of the nonblocking channel read extension:
channel long chan;

__kernel void kernel_read_channel (__global long * dst) {
  int i = 0;
  while (i < N) {
    bool valid0, valid1;
    long data0 = read_channel_nb_intel(chan, &valid0);
    long data1 = read_channel_nb_intel(chan, &valid1);
    if (valid0) {
       process(data0);
    }
    if (valid1) {
       process(data1); 
    } 
  }
}