Intel® FPGA SDK for OpenCL™ Standard Edition: Best Practices Guide

ID 683176
Date 9/24/2018
Public
Document Table of Contents

3.1.2. Execution Order for Channels and Pipes

Each channel or pipe call in a kernel program translates into an instruction executed in the FPGA pipeline. The execution of a channel call or a pipe call occurs if a valid work-item executes through the pipeline. However, even if there is no control or data dependence between channel or pipe calls, their execution might not achieve perfect instruction-level parallelism in the kernel pipeline.

Consider the following code examples:

Table 6.  Kernel with Two Read Channel or Pipe Calls
Kernel with Two Read Channel Calls Kernel with Two Read Pipe Calls
__kernel void
consumer (__global uint*restrict dst) {
  for (int i = 0; i < 5; i++) {
    dst[2*i] = read_channel_intel(c0);
    dst[2*i+2] = read_channel_intel(c1);
  }
}
__kernel void
consumer (__global uint*restrict dst,
  read_only pipe uint 
    __attribute__((blocking)) c0,
  read_only pipe uint
    __attribute__((blocking)) c1)
{
  for (int i = 0; i < 5; i++) {
    read_pipe (c0, &dst[2*i]);
    read_pipe (c1, &dst[2*i+2]);
  }
}

The code example on the left makes two read channel calls. The code example on the right makes two read pipe calls. In most cases, the kernel executes these channel or pipe calls in parallel; however, channel and pipe call executions might occur out of sequence. Out-of-sequence execution means that the read operation from c1 can occur and complete before the read operation from c0.