mem_fence() not working for channels

[Altera_Forum]

Hi all, 

 I am testing the function of the feed-forward model(ping-pong buffer) mentioned in the programming guide. And I found the mem_fence function is not working. Here's the code I used for testing:   channel int c_id __attribute__((depth(100)));  __attribute__((reqd_work_group_size(10,1,1))) __kernel void producer (__global int *restrict x, __global volatile int *restrict producer_data){  int global_x = get_global_id(0);  producer_data[global_x] = global_x;  mem_fence(CLK_CHANNEL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);  write_channel_altera(c_id, global_x); } __attribute__((reqd_work_group_size(10,1,1))) __kernel void consumer (__global int *restrict y, __global volatile int *restrict producer_data, __global int *restrict output) {  int global_x = read_channel_altera(c_id);  int sum2 = producer_data[global_x] + global_x;  output[global_x] = sum2; }   The producer_data[] is initialized to zero. The producer kernel writes it's global id to the producer_data[global_x] and channel. After the consumer kernel reads from the channel, it writes producer_data[global_x] + global_x to output[global_x]. The value of output[global_x] should always be global_x * 2.  However in my experiment, the output array is not always global_x * 2, sometimes output[global_x] = global_x. The consumer kernel reads the data before producer writes to global memory. mem_fence() here is not working.  I guess the problem lies with the way to create a shared buffer. I put the producer and consumer kernels in different command queues for the concurrent execution and use clCreateBuffer(...,CL_MEM_READ_WRITE,...) to create a shared buffer. The programming guide mentioned clCreateBuffer(...,CL_MEM_READ_WRITE,...) allocates memory to nonshared DDR memory banks and shared memory should be allocated by using clCreateBuffer(...,CL_MEM_ALLOC_HOST_PTR,...). However, when I am using the clCreateBuffer(...,CL_MEM_ALLOC_HOST_PTR,...) function, these two kernels cannot execute concurrently. Consumer kernel will wait until producer kernel finishes.  How to allocate the shared buffer for two concurrent kernels? Any example host code for the feed-forward model(ping-pong buffer) will greatly help.  Thanks in advance.

[Altera_Forum]

--- Quote Start ---  I guess the problem lies with the way to create a shared buffer. I put the producer and consumer kernels in different command queues for the concurrent execution and use clCreateBuffer(...,CL_MEM_READ_WRITE,...) to create a shared buffer. The programming guide mentioned clCreateBuffer(...,CL_MEM_READ_WRITE,...) allocates memory to nonshared DDR memory banks and shared memory should be allocated by using clCreateBuffer(...,CL_MEM_ALLOC_HOST_PTR,...). However, when I am using the clCreateBuffer(...,CL_MEM_ALLOC_HOST_PTR,...) function, these two kernels cannot execute concurrently. Consumer kernel will wait until producer kernel finishes.  --- Quote End ---    You are mixing two different concepts. The type of "shared memory" that the guide recommends to be allocated using CL_MEM_ALLOC_HOST_PTR is for FPGA SoCs which share the same "physical memory" between the ARM processor and the FPGA; it does not apply to PCI-E-attached FPGA boards and is not related to whether a "global buffer" is shared between two or more kernels.  Regarding your problem, OpenCL does NOT guarantee global memory consistency unless when kernel execution has finished. Hence, by default, trying to share a global buffer between to concurrent kernels with one writing to it and the other reading from it can (and likely will) lead to undefined behavior. Altera's guide claims using mem_fence(CLK_CHANNEL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE) can allow two such kernels to synchronize a shared global memory buffer through tokens passed via a channel as you are trying to do in this case; however, I have seen multiple people trying this and reporting that it doesn't work in the forum. In fact, the programming guide has an example of this in "Section 5.4.5.6 Use Models of Intel FPGA SDK for OpenCL Channels Implementation" but I wouldn't be surprised if that doesn't work either. Maybe you should consider using a Single Work-item implementation as is done in Altera's example and then it might work.