Visible to Intel only — GUID: zik1521227815263
Ixiasoft
Visible to Intel only — GUID: zik1521227815263
Ixiasoft
5.4.4. Restrictions in the Implementation of Intel® FPGA SDK for OpenCL™ Standard Edition Channels Extension
Multiple Channel Call Site
__kernel void k1() {
read_channel_intel (channel1);
read_channel_intel (channel1);
read_channel_intel (channel1);
}
__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.