Forum Discussion

Altera_Forum's avatar
Altera_Forum
Icon for Honored Contributor rankHonored Contributor
7 years ago

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

}

3 Replies

  • Altera_Forum's avatar
    Altera_Forum
    Icon for Honored Contributor rankHonored Contributor

    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's avatar
    Altera_Forum
    Icon for Honored Contributor rankHonored Contributor

    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's avatar
    Altera_Forum
    Icon for Honored Contributor rankHonored Contributor

    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.