/*
 *
 *  Simple kernel example to demonstrate packing 2 128 words into each IO channel word
 *
 */
#pragma OPENCL EXTENSION cl_intel_channels : enable

channel short16 serial_in0    __attribute__((io("kernel_input_ch0")));    //Channel 0 Rx
channel short16 serial_out0   __attribute__((io("kernel_output_ch0")));   //Channel 0 Tx

/*
        Test transmit kernel 0.
        Using short16 for testing 256 bit interface.
*/
__kernel void TX_Kernel_0(__global short8 *restrict input,
						unsigned int batch_size)
{
	unsigned int out_index = 0;
	unsigned char word_count = 0;
	short8 prev_word;
	
	while (out_index < batch_size)
	{
		short8 word = input[out_index];
		if (word_count == 1)
		{
			union
			{
				short8 in[2];
				short16 out;
			}u;
			
			u.in[0] = prev_word;
			u.in[1] = word;
			write_channel_intel(serial_out0, u.out);
		}
		prev_word = word;
		word_count = word_count == 0? 1: 0;
		out_index++;
	}
}

__kernel void RX_Kernel_0(__global short8 *restrict output, 
						unsigned int batch_size) 
{
	unsigned int out_index = 0;
	unsigned char word_count = 0;
	
	union
	{
		short8 out[2];
		short16 in;
	}u;
	
	while (out_index < batch_size)
	{
		if (word_count == 0)
			u.in = read_channel_intel(serial_in0);
		output[out_index] = word_count==0? u.out[0]: u.out[1];
		out_index++;
		word_count = word_count == 0? 1: 0;
	}
}