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

ID 683846
Date 10/04/2021
Public

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

Document Table of Contents

11.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 (int 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 (int i = 0; i < size; ++i) {
		data_out[i] = read_channel_intel(chain_channels[N]);;
	}
}
Figure 38. 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].