Multiple read/write from channels

The Fmax bottleneck only happens if one of the channel loops is not unrolled. I commented out both conditions in the channel loops and left the one in the compute loop and all loops where pipelined with an II of one without any Fmax bottleneck. This is the code:

#define W 32
 
#pragma OPENCL EXTENSION cl_intel_channels : enable;
 
channel float channel_float __attribute__((depth(W)));
 
__kernel void producer(const int N)
{
	int outer_loop_limit=(int)(N/(W))+1;
	for(int i=0;i<outer_loop_limit;i++)
	{
		const int it=i*W;
		#pragma unroll
		for(int j=0;j<W;j++)
		{
			//if(it+j<N)	//handle boundary conditions
				write_channel_intel(channel_float,(float)(2.0));
		}
	}
}
 
__kernel void consumer(const int N, __global float* res)
{
	const int outer_loop_limit=(int)(N/(W))+1;
	float acc_o=0, acc_i=0;
	float mult[W], x[W];
 
	for(int i=0; i<outer_loop_limit; i++)
	{
		const int it=i*W;
		#pragma unroll
		for(int j=0; j<W; j++)
		{
			//if(it+j<N)	//handle boundary conditions
				x[j]=read_channel_intel(channel_float);
		}
		acc_i=0;
 
		#pragma unroll
		for(int j=0; j<W; j++)
		{
			if(it+j<N)
				acc_i+=x[j];
		}
		acc_o+=acc_i;
	}
	*res=acc_o;
}

Regarding the problem with portability, I would personally recommend against having multiple reads/writes from/to channels. In fact, up until v17.1 (or maybe v17.0), doing so was not even allowed. If you want to completely avoid extra channel reads/writes while having pipelineable loops, you might be able to achieve this by using W channels instead of one channel that is unrolled W times. Something like this:

#pragma unroll
for (int j = 0; j < W; j++)
{
	if(it + j < N)
	{
		x[j] = read_channel_intel(channel_float[j])
	}
}

In this case the compiler should not fail to unroll the loop due to conditional execution anymore, since now you have W channels instead of one.

You might be able to avoid the issue by merging both channel reads and both channel writes into one read and one write and instead use an "if" condition inside of the unrolled read/write loops to avoid reading/writing the out-of-bound data.

I've tried it, loops are pipelined with II=1. But this comes at the expenses of a lower Fmax (as the report states)

The new code look likes this:

__kernel void producer(const int N)
{
	int outer_loop_limit=(int)(N/(W))+1;
	for(int i=0;i<outer_loop_limit;i++)
	{
		const int it=i*W;
		#pragma unroll
		for(int j=0;j<W;j++)
		{
			if(it+j<N)	//handle boundary conditions
				write_channel_intel(channel_float,(float)(2.0));
		}
	}
}
 
__kernel void consumer(const int N, __global float* res)
{
	const int outer_loop_limit=(int)(N/(W))+1;
	float acc_o=0, acc_i=0;
	float mult[W], x[W];
 
	for(int i=0; i<outer_loop_limit; i++)
	{
		const int it=i*W;
		#pragma unroll
		for(int j=0; j<W; j++)
		{
			if(it+j<N)	//handle boundary conditions
				x[j]=read_channel_intel(channel_float);
		}
		acc_i=0;
 
		#pragma unroll
		for(int j=0; j<W; j++)
		{
			if(it+j<N)
				acc_i+=x[j];
		}
		acc_o+=acc_i;
	}
}

The report states that a bottleneck occurs at line 34 (Fmax) due to data dependencies (it indicates the comparisons in the two loop guards of consumer and in the definition of the outer_loop_limit constant)

I am not getting any messages from the compiler regarding an Fmax bottleneck in your second code; however, on my environment, the compiler fails to unroll the loops over the channel operations now due to "conditional channel execution". Looks like my suggestion made things worse instead of better.

Is there any reason why you are trying to avoid passing the out-of-bound data via channels? Since W is small, the overhead of passing the few extra indexes will be extremely small. I suggest that you send the extra data via channels anyway to avoid the "conditional channel execution", but instead avoid processing the extra data. You already have the same condition in the second loop in the consumer kernel and that loop is correctly unrolled and pipelined; this should be enough to generate correct output even if you do send the out-of-bound data through the channels.

Sorry,

in the second posted code, I forgot to copy the last line in the consumer kernel, to save the data into memory (otherwise the compiler remove compilation)

    //...
   *res=acc_o;
}

Apart from this, even in my case it complains about conditional write, but then, according the report, loops are unrolled but the with lower Fmax.

I would like to avoid generates more data than needed, for the sake of code portability (e.g. the generator is implemented by some other code). Apart from this, even if I try to do it, the problem remains (bottleneck is still Fmax) and a new compiler warning appears "Cannot unroll loop for.body3 in producer because channel endpoints would undergo different amounts of unrolling"

Ok, i see.

Multiple channels may work, but I have to keep their number low (at least, this is what Intel documentation states).

Thanks