Channel from many work items to single task

Altera_Forum · ‎09-21-2014

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?

Altera_Forum · ‎09-26-2014

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

Altera_Forum · ‎09-27-2014

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 · ‎09-30-2014

--- 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 · ‎10-01-2014

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 · ‎01-15-2015

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