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

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

12.4.2. Using Channels with Kernel Copies

To implement channels within compute units (that is, replicated kernel copies), create an array of channels and then index into that array using the return value of get_compute_id().

The example code below implements channels within multiple compute units.

#define N 4
channel int chain_channels[N+1];

__attribute__((max_global_work_dim(0)))
__kernel void reader(global int *data_in, int size) {
	for (i = 0; i < size; ++i) {
		write_channel_intel(chain_channels[0], data_in[i]);
	}
}

__attribute__((max_global_work_dim(0)))
__attribute__((autorun))	
__attribute__((num_compute_units(N)))					
__kernel void plusOne() {
	int compute_id = get_compute_id(0);
	int input = read_channel_intel(chain_channels[compute_id]);
	write_channel_intel(chain_channels[compute_id+1], input + 1);
}

__attribute__((max_global_work_dim(0)))
__kernel void writer(global int *data_out, int size) {
	for (i = 0; i < size; ++i) {
		data_out[i] = read_channel_intel(chain_channels[N]);;
	}
}
Figure 22. Example Topology of Kernel Copies that Implement ChannelsThis figure illustrates the topology of the group of kernels that the OpenCL™ application code above generates.
Note: The implementation of kernel copies is functionally equivalent to defining four separate kernels in your source code and then hard-coding unique indexes for the accesses to chain_channels[N].