Forum Discussion

hiratz's avatar
hiratz
Icon for Occasional Contributor rankOccasional Contributor
6 years ago

Why does a blocked channel in one kernel also block other kernels' running?

Hi,

@HRZ​

I noticed the following scenario:

I have two kernels, say Producer and Consumer. I put them into two DIFFERENT command queues, so they can be launched and run concurrently. Meanwhile, there is a channel whose depth is set N between them. Each work item in the producer writes one value to the channel and the each work item in the consumer reads the channel. I use NDRange and the number of workitems is W.

Then

When W < N, everything is OK;

When W > N, the execution of Producer is blocked. But the Consumer also cannot be executed.

Since the Consumer is in a different queue from the Producer, why is it also blocked? Actually I guess it is even not launched.

I show a simple code sample here. In this example, N = 8 and W = 32. The "printf" in line 13 cannot be executed. Even if I commented line 16 (channel read), it still cannot be executed. So I conclude that the consumer is not launched. If I'm right, why? My original thought was: even if the producer is blocked, eventually it will be unblocked as long as the consumer is able to execute and keep reading the channel and making space for the channel.

Thank you!

channel ulong ch1 __attribute__ ((depth(8)));
 
__attribute__((max_work_group_size(32)))
__kernel void producer(buf_addr_space const int * restrict buf_in)
{
    size_t gid = get_global_id(0);
    write_channel_intel(ch1, buf_in[gid]);
}
 
__attribute__((max_work_group_size(32)))
__kernel void consumer(buf_addr_space int * restrict buf_out)
{
    printf("--------- test ------------\n");
    size_t gid = get_global_id(0);
 
    int val = 5 + read_channel_intel(ch1);
    buf_out[gid] = val;
}

9 Replies

  • HRZ's avatar
    HRZ
    Icon for Frequent Contributor rankFrequent Contributor
    6 years ago

    Is this with emulation or actual FPGA execution? Printf data is cached in Block RAMs and dumped after kernel execution on the actual FPGA (while it is dumped right away in the emulator) and hence, if the kernel runs into a deadlock, the printf output will never be displayed (printf is useless for debugging deadlocks on the FPGA). Other than that, I don't see any reason why your code above would deadlock unless there is some problem in your host code; e.g. one kernel is waiting for an event from the other or there is a clFinish on the producer kernel before the consumer kernel is launched. Your description very well matches a synchronization problem in the host code rather than the kernels.

    • hiratz's avatar
      hiratz
      Icon for Occasional Contributor rankOccasional Contributor
      6 years ago

      Thanks, HRZ.

      I just tested this with emulator (with the option "-emulator-channel-depth-model=strict"). My host code is very simple and no kernels are waiting for any events. I show the key code as follows (you can see I put Producer and Consumer into two different command queues, respectively). It is really weird.

      #define MAX_CMD_QUEUES 4
#define N_KERNEL 2
 
void run_kernel(cl_command_queue cmd_queue[MAX_CMD_QUEUES], cl_kernel (&kernel)[N_KERNEL], size_t n_thread)
{
    size_t global_work_size[1] = {(size_t)n_thread};
    size_t local_work_size[1] = {(size_t)n_thread};
 
    cl_event event_write[1], event_exec[2];
    cl_int status;
    cl_command_queue &rcmd_queue0 = cmd_queue[0], &rcmd_queue1 = cmd_queue[1], &rcmd_queue2 = cmd_queue[2];
 
    // Write data into buf_in from h_in
    status = clEnqueueWriteBuffer(rcmd_queue0, buf_in, CL_TRUE, 0, bufsize, h_in, 0, NULL, event_write);
    error_check(status, "Write buf_in failed!\n");
 
    // Launch the kernel Producer
    status = clEnqueueNDRangeKernel(rcmd_queue0, kernel[0], 1, NULL, global_work_size, local_work_size, 1, event_write, NULL);
    error_check(status, "Run kernel Producer error!\n");
 
    // Launch the kernel Consumer
    status = clEnqueueNDRangeKernel(rcmd_queue1, kernel[1], 1, NULL, global_work_size, local_work_size, 0, NULL, NULL);
    error_check(status, "Run kernel Consumer error!\n");
 
    // Read results back to h_out from buf_out
    read_back_results(rcmd_queue1, bufsize, buf_out, (char*)h_out, NULL);
 
    clReleaseEvent(event_write[0]);
    clReleaseEvent(event_exec[0]);
    clReleaseEvent(event_exec[1]);
}

      ​

      • HRZ's avatar
        HRZ
        Icon for Frequent Contributor rankFrequent Contributor
        6 years ago

        Have you tried commenting the error_check functions? Depending on its implementation, that function could be serializing the kernel launches. If you provide the full code for your example so that I can compile it on my own machine, I might be able to find the source of the problem.