Visible to Intel only — GUID: kok1521225088340
Ixiasoft
Visible to Intel only — GUID: kok1521225088340
Ixiasoft
5.5.2. Pipe Data Behavior
Consider the following code example:
__kernel void producer (write_only pipe uint __attribute__((blocking)) c0)
{
for (uint i = 0; i < 10; i++)
{
write_pipe (c0, &i);
}
}
__kernel void consumer (__global uint * restrict dst,
read_only pipe uint __attribute__((blocking))
__attribute__((depth(10))) c0)
{
for (int i = 0; i < 5; i++)
{
read_pipe (c0, &dst[i]);
}
}
A read operation to a pipe reads the least recent piece of data written to the pipe first. Pipe data maintains a FIFO ordering within the pipe.
The kernel producer writes ten elements ([0, 9]) to the pipe. The kernel consumer reads five elements from the pipe per NDRange invocation. During the first invocation, the kernel consumer reads values 0 to 4 from the pipe. 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 pipe becomes full. If you call consumer more than twice, consumer stalls because there is insufficient data in the pipe.