Forum Discussion

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

Channel from many work items to single task

Hi,

I am trying to compile something that is computing something with many work items. They are spit out to a channel into a single work item. There is no way to know how many at compile time but i get this error:

Channel Pairing Type Does not match

Channel c0 needs one chan_read_altera and one chan_write_altera.

Multiple chan_read_altera/chan_write_altera with the same channel ID may cause this problem.

system_integrator: custom_ic_impl.cpp:4422: void custom_ic::System::build_channel_system(): Assertion `0' failed.

There is only one of each channel read/write call in the code and i the different kernels (one in the multiple work item, the other in a single task)

Any ideas?

Thanks,

Stephen

Edit: I think i realized why. Is it that the number of of calls has to be defined to be the same? (ie static?) If so that sucks ... i obviously want to read and write N times in each kernel but can i vary N?

5 Replies

  • Altera_Forum's avatar
    Altera_Forum
    Icon for Honored Contributor rankHonored Contributor
    12 years ago

    Your channel c0 has multiple chan_write_altera but one chan_read_altera to channel c0. It is not allowed.

  • Altera_Forum's avatar
    Altera_Forum
    Icon for Honored Contributor rankHonored Contributor
    12 years ago

    Hi,

    I wished to have N work-items in the write kernel, and call the single task kernel reading N times.

    I have found that with the static (defined) N, this seems to be faster.

    I now want to be able to vary the work group size of the writer, and pass N as a parameter to the single task kernel.

    The error above is what i got. Any way around it?

    Thanks,

    Stephen
  • Altera_Forum's avatar
    Altera_Forum
    Icon for Honored Contributor rankHonored Contributor
    12 years ago

    --- Quote Start ---

    Hi,

    I wished to have N work-items in the write kernel, and call the single task kernel reading N times.

    I have found that with the static (defined) N, this seems to be faster.

    Stephen

    --- Quote End ---

    Hi,

    I actually am interested in doing so, too. You mean that with a static global size, it does work?

    How do you define a static global size at compile time? I thought we only can specify it in the call to enqueueNDRangeKernel().
  • Altera_Forum's avatar
    Altera_Forum
    Icon for Honored Contributor rankHonored Contributor
    12 years ago

    Hmm i think it was just with the -O3 optimization flag. Prob just a bug then. Heres some examples that work 

    __kernel void producer (__global int * in_buf, int N)
{
for (int index=0; index < N; index++)
{
write_channel_altera( c0, in_buf );
}
}
__kernel void consumer (__global int * ret_buf, int N)
{
for(int index=0; index < N; index++)
{
ret_buf = read_channel_altera( c0 );
}
}

    This is what i meant by static ... 

    #pragma OPENCL EXTENSION cl_altera_channels : enable
channel int c0;
# define N 256
__attribute__((reqd_work_group_size(N,1,1)))
__kernel void producer (__global int * in_buf)
{
int gid = get_global_id(0);
write_channel_altera( c0, in_buf );
}
__kernel void consumer (__global int * ret_buf)
{
for(int index=0; index < N; index++)
{
ret_buf = read_channel_altera( c0 );
}
}

    This is what i was trying to do. Didn't try it without -O3 before but i guess this is solved then :D 

    #pragma OPENCL EXTENSION cl_altera_channels : enable
channel int c0;
__attribute__((max_work_group_size(256)))
__kernel void producer (__global int * in_buf)
{
int gid = get_global_id(0);
write_channel_altera( c0, in_buf );
}
__kernel void consumer (__global int * ret_buf, int N)
{
for(int index=0; index < N; index++)
{
ret_buf = read_channel_altera( c0 );
}
}

    edit: This is not exactly my code n cbf compiling so idk which of these is faster, i had 16 channels in parallel in mine
  • Altera_Forum's avatar
    Altera_Forum
    Icon for Honored Contributor rankHonored Contributor
    11 years ago

    I have just started working with channels so my knowledge is limited. I do recall reading that you can't use "num_simd_wk_items" when implementing channels. I wonder if -O3 is trying to use that attribute. I don't know for sure, but if it was, I could see you getting that error.

    Rudy