Intel® Quartus® Prime Software
Intel® Quartus® Prime Design Software, Design Entry, Synthesis, Simulation, Verification, Timing Analysis, System Design (Platform Designer, formerly Qsys)
Announcements
Intel Support hours are Monday-Fridays, 8am-5pm PST, except Holidays. Thanks to our community members who provide support during our down time or before we get to your questions. We appreciate you!

Need Forum Guidance? Click here
Search our FPGA Knowledge Articles here.

channel deadlock

Altera_Forum
Honored Contributor II
815 Views

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); 

}
0 Kudos
1 Reply
Altera_Forum
Honored Contributor II
88 Views

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.
Reply