5.4.5.2. Implementing Blocking Channel Writes

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

Download PDF

ID 683846

Date 3/28/2022

Version

Public

A newer version of this document is available. Customers should click here to go to the newest version.

Visible to Intel only — Ixiasoft

View Details

5.4.5.2. Implementing Blocking Channel Writes

The write_channel_intel API call allows you to send data across a channel.

To implement a blocking channel write, use the following write_channel_intel function signature:

void write_channel_intel (channel <type> channel_id, const <type> data);

Where:

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

data is the data that the channel write operation writes to the channel.

<type> defines a channel data width. Follow the OpenCL™ conversion rules to ensure that data the kernel writes to a channel is convertible to <type>.

The following code snippet demonstrates the implementation of the write_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 global memory, and passes this data to the channel.*/ 
__kernel void kernel_write_channel( __global const long * src ) {
  for (int i = 0; i < N; i++) {
     //Writes the eight bytes to the channel.
     write_channel_intel(chan, src[i]);
  }
}

CAUTION:

When you send data across a channel using the write_channel_intel API call, keep in mind that if the channel is full (that is, if the FIFO buffer is full of data), your kernel stalls and waits until at least one data slot becomes available in the FIFO buffer. Use the Intel® FPGA dynamic profiler for OpenCL™ to check for channel stalls.

Implementing Nonblocking Channel Writes

Perform nonblocking channel writes to facilitate applications where writes to a full FIFO buffer should not cause the kernel to stall until a slot in the FIFO buffer becomes free. A nonblocking channel write returns a boolean value that indicates whether data was written successfully to the channel (that is, the channel was not full).

Consider a scenario where your application has one data producer with two identical workers. Assume the time each worker takes to process a message varies depending on the contents of the data. In this case, there might be situations where one worker is busy while the other is free. A nonblocking write can facilitate work distribution such that both workers are busy.

To implement a nonblocking channel write, include the following write_channel_nb_intel function signature:

bool write_channel_nb_intel(channel <type> channel_id, const <type> data);

The following code snippet of the kernel producer facilitates work distribution using the nonblocking channel write extension:

channel long worker0, worker1;
__kernel void producer( __global const long * src ) {
  for(int i = 0; i < N; i++)  {
        
    bool success = false;
    do {
      success = write_channel_nb_intel(worker0, src[i]);
      if(!success) {
        success = write_channel_nb_intel(worker1, src[i]);
      }
    }
    while(!success);
  }
}

Select Your Language

Using Intel.com Search

Quick Links

Recent Searches

Advanced Search

Only search in

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

5.4.5.2. Implementing Blocking Channel Writes

Implementing Nonblocking Channel Writes