Results 1 to 2 of 2

Thread: channel deadlock

  1. #1
    Join Date
    May 2018
    Posts
    4
    Rep Power
    1

    Question channel deadlock

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

  2. #2
    Join Date
    Jan 2017
    Posts
    692
    Rep Power
    1

    Default Re: channel deadlock

    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.

Similar Threads

  1. Channel Deadlock Debugging
    By amaier17 in forum OpenCL
    Replies: 9
    Last Post: July 21st, 2017, 04:38 AM
  2. Replies: 3
    Last Post: January 10th, 2017, 02:05 AM
  3. Replies: 1
    Last Post: July 15th, 2013, 11:10 PM
  4. Input decay and state machine deadlock
    By ironmoose in forum General Altera Discussion
    Replies: 6
    Last Post: February 23rd, 2011, 07:48 AM
  5. Replies: 6
    Last Post: July 15th, 2010, 07:53 PM

Tags for this Thread

Bookmarks

Posting Permissions

  • You may not post new threads
  • You may not post replies
  • You may not post attachments
  • You may not edit your posts
  •