channel deadlock

[Altera_Forum]

The below source code is used in opencl programming guide. It is told that without using "mem_fence(CLK_CHANNEL_MEM_FENCE)", the consumer module might end up with deadlock.  

 1) Consider a case where producer has written in channel c0 first and consumer has read from channel c1. Consumer is currently reading from empty channel. It would stall for couple of cycles but why would there be a deadlock?   2) Cant producer write in channel c0 and c1 in parallel?   __kernel void producer (__global const uint * src, const uint iterations) { for (int i = 0; i < iterations; i++) { write_channel_intel(c0, src[2*i]); write_channel_intel(c1, src[2*i+1]); } } __kernel void consumer (__global uint * dst, const uint iterations) { for (int i = 0; i < iterations; i++) { /*During compilation, the AOC might reorder the way the consumer kernel writes to memory to optimize memory access. Therefore, c1 might be read before c0, which is the reverse of what appears in code.*/ dst[2*i+1] = read_channel_intel(c0); dst[2*i] = read_channel_intel(c1); } }

[Altera_Forum]

You should take channel depth into account to understand why this could cause a deadlock. If both channels are empty, the process will never deadlock, but since the order and rate of producing and consuming might not be the same (due to other stallable operations like external memory accesses), a situation could happen in which c0 is full while c1 is empty. In this case, the producer will stall on trying to write to c0 that is full, while the consumer will stall on reading from c1 that is empty. This will obviously cause a deadlock. 

 In your specific snippet, it is unlikely for the code to deadlock, though, even if the compiler reorders the channel operations.  Regarding your second question, I guess this is not possible or else channel reordering sould never cause a deadlock. However, if you have an unrolled loop over an array of channels, then I believe those channels could be operated in parallel.