Data sharing among various kernels

[Altera_Forum]

Hello, 

 I have tried working with read and write channels to pass data from one kernel to other. I wanted to know, whether a global variable can be shared by various kernels to pass the data. I have written a small program below indicating channel transfer. But is possible to transfer using shared global variables? # pragma OPENCL EXTENSION cl_intel_channels : enable  channel float data_ch __attribute__((depth(0)));  __kernel void Pipe_in(__global float *restrict x, __global float *restrict check) {  // Get index of the work item   uint count;  for (count = 0; count<10; count++){   write_channel_intel(data_ch, x[count]);    check[count] = x[count]; }  }  __kernel void Pipe_out(__global float *restrict y) {  // Get index of the work item   uint count;  for (count = 0; count<10; count++)   y[count] = read_channel_intel(data_ch);  }

[Altera_Forum]

You can share data between two kernels using global buffers; however, this is only possible if the kernels run sequentially, and one finishes before the other one starts. If the kernels are supposed to run in parallel, this is not possible since OpenCL guarantees global memory consistency only AFTER kernel execution.

[Altera_Forum]

Should the global buffer be defined as the argument in one of the kernel? Because when I define __global float share[10] in Pipe_in kernel (below program), i end up with errors. 

 # pragma OPENCL EXTENSION cl_intel_channels : enable  channel float data_ch __attribute__((depth(0)));  __kernel void Pipe_in(__global float *restrict x, __global float *restrict check) { // Get index of the work item  __global float share[10]; uint count;  for (count = 0; count<10; count++){  write_channel_intel(data_ch, x[count]);  share[count] = x[count]; }  }  __kernel void Pipe_out(__global float *restrict y) { // Get index of the work item   uint count;  for (count = 0; count<10; count++)  y[count] = share[count];  }

[Altera_Forum]

That is not how you share global buffers. You cannot define global buffers inside the kernel code. They must be defined and allocated in the host code. You can just add one extra global argument to each kernel and after allocating the shared buffer in the host code, pass the pointer to the shared buffer to both kernels.