Community
cancel
Showing results for 
Search instead for 
Did you mean: 
Altera_Forum
Honored Contributor I
1,237 Views

Data sharing among various kernels

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

 

}
0 Kudos
3 Replies
Altera_Forum
Honored Contributor I
32 Views

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
Honored Contributor I
32 Views

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
Honored Contributor I
32 Views

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.

Reply