5.4.2. Channel Data Behavior

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

Download PDF

ID 683342

Date 4/22/2019

Version

Public

Visible to Intel only — GUID: gom1521227791495

Ixiasoft

View Details

5.4.2. Channel Data Behavior

Data written to a channel remains in a channel as long as the kernel program remains loaded on the FPGA device.

In other words, data written to a channel persists across multiple work-groups and NDRange invocations. However, data is not persistent across multiple or different invocations of kernel programs that lead to FPGA device reprogramming.

Data in channels does not persist between context, program, device, kernel, or platform releases, even if the OpenCL implementation performs optimizations that avoid reprogramming operations on a device. For example, if you run a host program twice using the same .aocx file, or if a host program releases and reacquires a context, the data in the channel might or might not persist across the operation. FPGA device reset operations might happen behind the scenes on object releases that purge data in any channels

Consider the following code example:

channel int c0;

__kernel void producer() {
    for (int i = 0; i < 10; i++) {
        write_channel_intel (c0, i);
    }
}

__kernel void consumer (__global uint * restrict dst) {
    for (int i = 0; i < 5; i++) {
        dst[i] = read_channel_intel(c0);
    }
}

Figure 7. Channel Data FIFO Ordering

The kernel producer writes ten elements ([0, 9]) to the channel. The kernel consumer does not contain any work-item identifier queries; therefore, it will receive an implicit reqd_work_group_size attribute of (1,1,1). The implied reqd_work_group_size(1,1,1) attribute means that consumer must be launched as a single work-item kernel. In the example above, consumer reads five elements from the channel per invocation. During the first invocation, the kernel consumer reads values 0 to 4 from the channel. Because the data persists across NDRange invocations, the second time you execute the kernel consumer, it reads values 5 to 9.

For this example, to avoid a deadlock from occurring, you need to invoke the kernel consumer twice for every invocation of the kernel producer. If you call consumer less than twice, producer stalls because the channel becomes full. If you call consumer more than twice, consumer stalls because there is insufficient data in the channel.

Select Your Language

Using Intel.com Search

Quick Links

Recent Searches

Advanced Search

Only search in

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

5.4.2. Channel Data Behavior