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

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

5.4.4. Restrictions in the Implementation of Intel® FPGA SDK for OpenCL™ Standard Edition Channels Extension

There are certain design restrictions to the implementation of channels in your OpenCL™ application.

Multiple Channel Call Site

A kernel can read from the same channel multiple times. However, multiple kernels cannot read from the same channel. Similarly, a kernel can write to the same channel multiple times but multiple kernels cannot write to the same channel.
__kernel void k1() {
     read_channel_intel (channel1);
     read_channel_intel (channel1);
     read_channel_intel (channel1);
}
The Intel® FPGA SDK for OpenCL™ Offline Compiler cannot compile the following code and will issue an error:
__kernel void k1(){
     write_channel_intel (channel1, 1);
}

__kernel void k2() {
     write_channel_intel (channel1, 2);
}

Feedback and Feed-Forward Channels

Channels within a kernel can be either read_only or write_only. Performance of a kernel that reads and writes to the same channel might be poor.

Static Indexing

The Intel® FPGA SDK for OpenCL™ channels extension does not support dynamic indexing into arrays of channel IDs because it leads to inefficient hardware.

Consider the following example:

channel int ch[WORKGROUP_SIZE];

__kernel void consumer()
{

    int gid = get_global_id(0);
    int value = read_channel_intel(ch[gid]);

    //statements
}

Compilation of this example kernel fails with the following error message:

Compiler Error: Indexing into channel array ch could not be resolved to all constant

To avoid this compilation error, index into arrays of channel IDs statically, as shown below:

channel int ch[WORKGROUP_SIZE];
 
__kernel void consumer() {
 
    int gid = get_global_id(0);
    int value;

    switch(gid)
    {
        case 0: value = read_channel_intel(ch[0]); break;
        case 1: value = read_channel_intel(ch[1]); break;
        case 2: value = read_channel_intel(ch[2]); break;
        case 3: value = read_channel_intel(ch[3]); break;
        //statements
	
	    case WORKGROUP_SIZE-1:read_channel_intel(ch[WORKGROUP_SIZE-1]); break;
    }

    //statements
}

Kernel Vectorization Support

You cannot vectorize kernels that use channels; that is, do not include the num_simd_work_items kernel attribute in your kernel code. Vectorizing a kernel that uses channels creates multiple channel accesses inside the same kernel and requires arbitration, which negates the advantages of vectorization. As a result, the SDK's channel extension does not support kernel vectorization.

Instruction-Level Parallelism on read_channel_intel and write_channel_intel Calls

If no data dependencies exist between read_channel_intel and write_channel_intel calls, the offline compiler attempts to execute these instructions in parallel. As a result, the offline compiler might execute these read_channel_intel and write_channel_intel calls in an order that does not follow the sequence expressed in the OpenCL kernel code.

Consider the following code sequence:

in_data1 = read_channel_intel(channel1);
in_data2 = read_channel_intel(channel2);
in_data3 = read_channel_intel(channel3);

Because there are no data dependencies between the read_channel_intel calls, the offline compiler can execute them in any order.