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

ID 683342
Date 4/22/2019
Public
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.