Visible to Intel only — GUID: hnj1521227876209
Ixiasoft
Visible to Intel only — GUID: hnj1521227876209
Ixiasoft
5.4.5.7. Implementing Buffered Channels Using the depth Channels Attribute
You may use a buffered channel to control data traffic, such as limiting throughput or synchronizing accesses to shared memory. In an unbuffered channel, a write operation cannot proceed until the read operation reads a data value. In a buffered channel, a write operation cannot proceed until the data value is copied to the buffer. If the buffer is full, the operation cannot proceed until the read operation reads a piece of data and removes it from the channel.
channel int c __attribute__((depth(10)));
__kernel void producer (__global int * in_data)
{
for (int i = 0; i < N; i++)
{
if (in_data[i])
{
write_channel_intel(c, in_data[i]);
}
}
}
__kernel void consumer (__global int * restrict check_data,
__global int * restrict out_data)
{
int last_val = 0;
for (int i = 0; i < N, i++)
{
if (check_data[i])
{
last_val = read_channel_intel(c);
}
out_data[i] = last_val;
}
}
In this example, the write operation can write ten data values to the channel without blocking. Once the channel is full, the write operation cannot proceed until an associated read operation to the channel occurs.
Because the channel read and write calls are conditional statements, the channel might experience an imbalance between read and write calls. You may add a buffer capacity to the channel to ensure that the producer and consumer kernels are decoupled. This step is particularly important if the producer kernel is writing data to the channel when the consumer kernel is not reading from it.