Forum Discussion
Occasional ContributorWhy 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
HRZFrequent Contributor
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.
hiratzThanks, 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
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.
