Visible to Intel only — GUID: mtu1521225134364
Ixiasoft
Visible to Intel only — GUID: mtu1521225134364
Ixiasoft
5.5.5.5. Implementing Buffered Pipes Using the depth Attribute
You may use a buffered pipe to control data traffic, such as limiting throughput or synchronizing accesses to shared memory. In an unbuffered pipe, a write operation can only proceed when the read operation is expecting to read data. Use unbuffered pipes in conjunction with blocking read and write behaviors in kernels that execute concurrently. The unbuffered pipes provide self-synchronizing data transfers efficiently.
In a buffered pipe, a write operation can only proceed if there is capacity in the pipe to hold the incoming packet. A read operation can only proceed if there is at least one packet in the pipe.
Use buffered pipes if pipe calls are predicated differently in the writer and reader kernels, and the kernels do not execute concurrently.
__kernel void
producer (__global int *in_data,
write_only pipe int __attribute__((blocking))
__attribute__((depth(10))) c)
{
for (i = 0; i < N; i++)
{
if (in_data[i])
{
write_pipe( c, &in_data[i] );
}
}
}
__kernel void
consumer (__global int *check_data,
__global int *out_data,
read_only pipe int __attribute__((blocking)) c )
{
int last_val = 0;
for (i = 0; i < N; i++)
{
if (check_data[i])
{
read_pipe( c, &last_val );
}
out_data[i] = last_val;
}
}
In this example, the write operation can write ten data values to the pipe successfully. After the pipe is full, the write kernel returns failure until a read kernel consumes some of the data in the pipe.
Because the pipe read and write calls are conditional statements, the pipe might experience an imbalance between read and write calls. You may add a buffer capacity to the pipe to ensure that the producer and consumer kernels are decoupled. This step is particularly important if the producer kernel is writing data to the pipe when the consumer kernel is not reading from it.