/*M///////////////////////////////////////////////////////////////////////////////////////
//
//  IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
//
//  By downloading, copying, installing or using the software you agree to this license.
//  If you do not agree to this license, do not download, install,
//  copy or use the software.
//
//
//                           License Agreement
//                For Open Source Computer Vision Library
//
// Copyright (C) 2010-2012, Institute Of Software Chinese Academy Of Science, all rights reserved.
// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved.
// Third party copyrights are property of their respective owners.
//
// @Authors
//    Jia Haipeng, jiahaipeng95@gmail.com
//
// Redistribution and use in source and binary forms, with or without modification,
// are permitted provided that the following conditions are met:
//
//   * Redistribution's of source code must retain the above copyright notice,
//     this list of conditions and the following disclaimer.
//
//   * Redistribution's in binary form must reproduce the above copyright notice,
//     this list of conditions and the following disclaimer in the documentation
//     and/or other materials provided with the distribution.
//
//   * The name of the copyright holders may not be used to endorse or promote products
//     derived from this software without specific prior written permission.
//
// This software is provided by the copyright holders and contributors as is and
// any express or implied warranties, including, but not limited to, the implied
// warranties of merchantability and fitness for a particular purpose are disclaimed.
// In no event shall the Intel Corporation or contributors be liable for any direct,
// indirect, incidental, special, exemplary, or consequential damages
// (including, but not limited to, procurement of substitute goods or services;
// loss of use, data, or profits; or business interruption) however caused
// and on any theory of liability, whether in contract, strict liability,
// or tort (including negligence or otherwise) arising in any way out of
// the use of this software, even if advised of the possibility of such damage.
//
//M*/

#ifdef DOUBLE_SUPPORT
#ifdef cl_amd_fp64
#pragma OPENCL EXTENSION cl_amd_fp64:enable
#elif defined (cl_khr_fp64)
#pragma OPENCL EXTENSION cl_khr_fp64:enable
#endif
#endif

///////////////////////////////////////////////////////////////////////////////////////////////
//////////////////////////////////optimized code using vector roi//////////////////////////
////////////vector fuction name format: merge_vector_C(channels number)D_(data type depth)//////
////////////////////////////////////////////////////////////////////////////////////////////////

__kernel void merge_vector_C2_D0(__global uchar *mat_dst,  int dst_step,  int dst_offset,
                                 __global uchar *mat_src0, int src0_step, int src0_offset,
                                 __global uchar *mat_src1, int src1_step, int src1_offset,
                                 int rows, int cols, int dst_step1)

{
    int x = get_global_id(0);
    int y = get_global_id(1);

    if ((x < cols) && (y < rows))
    {
        x = x << 1;

        #define dst_align  ((dst_offset & 3) >> 1)
        int src0_index = mad24(y, src0_step, src0_offset + x - dst_align);
        int src1_index = mad24(y, src1_step, src1_offset + x - dst_align);

        int dst_start  = mad24(y, dst_step, dst_offset);
        int dst_end    = mad24(y, dst_step, dst_offset + dst_step1);
        int dst_index  = mad24(y, dst_step, dst_offset + (x << 1) & (int)0xfffffffc);

        __global uchar4 * dst  = (__global uchar4 *)(mat_dst + dst_index);
        __global uchar  * src0 = mat_src0 + src0_index;
        __global uchar  * src1 = src0     + 1;
        __global uchar  * src2 = mat_src1 + src1_index;
        __global uchar  * src3 = src2     + 1;

        uchar4 dst_data = *dst;
        uchar  data_0   = *(src0);
        uchar  data_1   = *(src1);
        uchar  data_2   = *(src2);
        uchar  data_3   = *(src3);

        uchar4 tmp_data = (uchar4)(data_0, data_2, data_1, data_3);

        tmp_data.xy = dst_index + 0 >= dst_start ? tmp_data.xy : dst_data.xy;
        tmp_data.zw = dst_index + 2 <  dst_end   ? tmp_data.zw : dst_data.zw;

        *dst = tmp_data;
    }
}
__kernel void merge_vector_C2_D1(__global char *mat_dst,  int dst_step,  int dst_offset,
                                 __global char *mat_src0, int src0_step, int src0_offset,
                                 __global char *mat_src1, int src1_step, int src1_offset,
                                 int rows, int cols, int dst_step1)

{
    int x = get_global_id(0);
    int y = get_global_id(1);

    if ((x < cols) && (y < rows))
    {
        x = x << 1;

        #define dst_align  ((dst_offset & 3) >> 1)
        int src0_index = mad24(y, src0_step, src0_offset + x - dst_align);
        int src1_index = mad24(y, src1_step, src1_offset + x - dst_align);

        int dst_start  = mad24(y, dst_step, dst_offset);
        int dst_end    = mad24(y, dst_step, dst_offset + dst_step1);
        int dst_index  = mad24(y, dst_step, dst_offset + (x << 1) & (int)0xfffffffc);

        __global char4 * dst  = (__global char4 *)(mat_dst + dst_index);
        __global char  * src0 = mat_src0 + src0_index;
        __global char  * src1 = src0     + 1;
        __global char  * src2 = mat_src1 + src1_index;
        __global char  * src3 = src2     + 1;

        char4 dst_data = *dst;
        char  data_0   = *(src0);
        char  data_1   = *(src1);
        char  data_2   = *(src2);
        char  data_3   = *(src3);

        char4 tmp_data = (char4)(data_0, data_2, data_1, data_3);

        tmp_data.xy = dst_index + 0 >= dst_start ? tmp_data.xy : dst_data.xy;
        tmp_data.zw = dst_index + 2 <  dst_end   ? tmp_data.zw : dst_data.zw;

        *dst = tmp_data;
    }
}
__kernel void merge_vector_C2_D2(__global ushort *mat_dst,  int dst_step,  int dst_offset,
                                 __global ushort *mat_src0, int src0_step, int src0_offset,
                                 __global ushort *mat_src1, int src1_step, int src1_offset,
                                 int rows, int cols, int dst_step1)

{
    int x = get_global_id(0);
    int y = get_global_id(1);

    if ((x < cols) && (y < rows))
    {
        int src0_index = mad24(y, src0_step, src0_offset);
        int src1_index = mad24(y, src1_step, src1_offset);

        int dst_index  = mad24(y, dst_step , dst_offset);

        __global ushort*  src0 = (__global ushort * )((__global uchar *)mat_src0 + src0_index + (x << 1));
        __global ushort*  src1 = (__global ushort * )((__global uchar *)mat_src1 + src1_index + (x << 1));
        __global ushort2* dist = (__global ushort2 *)((__global uchar *)mat_dst  + dst_index  + (x << 2));

        ushort  src0_data = *src0;
        ushort  src1_data = *src1;

        *dist = (ushort2)(src0_data, src1_data);

    }
}
__kernel void merge_vector_C2_D3(__global short *mat_dst,  int dst_step,  int dst_offset,
                                 __global short *mat_src0, int src0_step, int src0_offset,
                                 __global short *mat_src1, int src1_step, int src1_offset,
                                 int rows, int cols, int dst_step1)
{
    int x = get_global_id(0);
    int y = get_global_id(1);

    if ((x < cols) && (y < rows))
    {
        int src0_index = mad24(y, src0_step, src0_offset);
        int src1_index = mad24(y, src1_step, src1_offset);

        int dst_index  = mad24(y, dst_step , dst_offset);

        __global short*  src0 = (__global short * )((__global uchar *)mat_src0 + src0_index + (x << 1));
        __global short*  src1 = (__global short * )((__global uchar *)mat_src1 + src1_index + (x << 1));
        __global short2* dist = (__global short2 *)((__global uchar *)mat_dst  + dst_index   + (x << 2));

        short  src0_data = *src0;
        short  src1_data = *src1;

        *dist = (short2)(src0_data, src1_data);
    }
}

__kernel void merge_vector_C2_D4(__global int *mat_dst,  int dst_step,  int dst_offset,
                                 __global int *mat_src0, int src0_step, int src0_offset,
                                 __global int *mat_src1, int src1_step, int src1_offset,
                                 int rows, int cols, int dst_step1)
{
    int x = get_global_id(0);
    int y = get_global_id(1);

    if ((x < cols) && (y < rows))
    {
        int src0_index = mad24(y, src0_step, src0_offset);
        int src1_index = mad24(y, src1_step, src1_offset);
        int dst_index  = mad24(y, dst_step , dst_offset);

        int src0 = *((__global int *)((__global uchar *)mat_src0 + src0_index + (x << 2)));
        int src1 = *((__global int *)((__global uchar *)mat_src1 + src1_index + (x << 2)));

        *((__global int2 *)((__global uchar *)mat_dst  + dst_index + (x << 3))) = (int2)(src0, src1);
    }
}
__kernel void merge_vector_C2_D5(__global float *mat_dst,  int dst_step,  int dst_offset,
                                 __global float *mat_src0, int src0_step, int src0_offset,
                                 __global float *mat_src1, int src1_step, int src1_offset,
                                 int rows, int cols, int dst_step1)
{
    int x = get_global_id(0);
    int y = get_global_id(1);

    if ((x < cols) && (y < rows))
    {
        int src0_index = mad24(y, src0_step, src0_offset);
        int src1_index = mad24(y, src1_step, src1_offset);
        int dst_index  = mad24(y, dst_step , dst_offset);

        float src0 = *((__global float *)((__global uchar *)mat_src0 + src0_index + (x << 2)));
        float src1 = *((__global float *)((__global uchar *)mat_src1 + src1_index + (x << 2)));

        *((__global float2 *)((__global uchar *)mat_dst  + dst_index + (x << 3))) = (float2)(src0, src1);
    }
}

#if defined (DOUBLE_SUPPORT)
__kernel void merge_vector_C2_D6(__global double *mat_dst,  int dst_step,  int dst_offset,
                                 __global double *mat_src0, int src0_step, int src0_offset,
                                 __global double *mat_src1, int src1_step, int src1_offset,
                                 int rows, int cols, int dst_step1)
{
    int x = get_global_id(0);
    int y = get_global_id(1);

    if ((x < cols) && (y < rows))
    {
        int src0_index = mad24(y, src0_step, src0_offset);
        int src1_index = mad24(y, src1_step, src1_offset);
        int dst_index  = mad24(y, dst_step , dst_offset);

        double src0 = *((__global double *)((__global uchar *)mat_src0 + src0_index + (x << 3)));
        double src1 = *((__global double *)((__global uchar *)mat_src1 + src1_index + (x << 3)));

        *((__global double2 *)((__global uchar *)mat_dst  + dst_index + (x << 4))) = (double2)(src0, src1);
    }
}
#endif

__kernel void merge_vector_C3_D0(__global uchar *mat_dst,  int dst_step,  int dst_offset,
                                 __global uchar *mat_src0, int src0_step, int src0_offset,
                                 __global uchar *mat_src1, int src1_step, int src1_offset,
                                 __global uchar *mat_src2, int src2_step, int src2_offset, int offset_cols,
                                 int rows, int cols, int dst_step1)
{
    int x = get_global_id(0);
    int y = get_global_id(1);

    if ((x < cols) && (y < rows))
    {
        x = x << 2;

        int src0_index = mad24(y, src0_step, x + src0_offset - offset_cols);
        int src1_index = mad24(y, src1_step, x + src1_offset - offset_cols);
        int src2_index = mad24(y, src2_step, x + src2_offset - offset_cols);

        int dst_start = mad24(y, dst_step, dst_offset);
        int dst_end   = mad24(y, dst_step, dst_offset + dst_step1);
        int dst_index = mad24(y, dst_step, dst_offset + 3 * x - offset_cols * 3);

        uchar data0_0 = *(mat_src0 + src0_index + 0);
        uchar data0_1 = *(mat_src0 + src0_index + 1);
        uchar data0_2 = *(mat_src0 + src0_index + 2);
        uchar data0_3 = *(mat_src0 + src0_index + 3);

        uchar data1_0 = *(mat_src1 + src1_index + 0);
        uchar data1_1 = *(mat_src1 + src1_index + 1);
        uchar data1_2 = *(mat_src1 + src1_index + 2);
        uchar data1_3 = *(mat_src1 + src1_index + 3);

        uchar data2_0 = *(mat_src2 + src2_index + 0);
        uchar data2_1 = *(mat_src2 + src2_index + 1);
        uchar data2_2 = *(mat_src2 + src2_index + 2);
        uchar data2_3 = *(mat_src2 + src2_index + 3);

        uchar4 tmp_data0 = (uchar4)(data0_0, data1_0, data2_0, data0_1);
        uchar4 tmp_data1 = (uchar4)(data1_1, data2_1, data0_2, data1_2);
        uchar4 tmp_data2 = (uchar4)(data2_2, data0_3, data1_3, data2_3);

        uchar4 dst_data0 = *((__global uchar4*)(mat_dst + dst_index + 0));
        uchar4 dst_data1 = *((__global uchar4*)(mat_dst + dst_index + 4));
        uchar4 dst_data2 = *((__global uchar4*)(mat_dst + dst_index + 8));

        tmp_data0.x = ((dst_index + 0  >= dst_start) && (dst_index + 0  < dst_end)) ? tmp_data0.x : dst_data0.x;
        tmp_data0.y = ((dst_index + 1  >= dst_start) && (dst_index + 1  < dst_end)) ? tmp_data0.y : dst_data0.y;
        tmp_data0.z = ((dst_index + 2  >= dst_start) && (dst_index + 2  < dst_end)) ? tmp_data0.z : dst_data0.z;
        tmp_data0.w = ((dst_index + 3  >= dst_start) && (dst_index + 3  < dst_end)) ? tmp_data0.w : dst_data0.w;

        tmp_data1.x = ((dst_index + 4  >= dst_start) && (dst_index + 4  < dst_end)) ? tmp_data1.x : dst_data1.x;
        tmp_data1.y = ((dst_index + 5  >= dst_start) && (dst_index + 5  < dst_end)) ? tmp_data1.y : dst_data1.y;
        tmp_data1.z = ((dst_index + 6  >= dst_start) && (dst_index + 6  < dst_end)) ? tmp_data1.z : dst_data1.z;
        tmp_data1.w = ((dst_index + 7  >= dst_start) && (dst_index + 7  < dst_end)) ? tmp_data1.w : dst_data1.w;

        tmp_data2.x = ((dst_index + 8  >= dst_start) && (dst_index + 8  < dst_end)) ? tmp_data2.x : dst_data2.x;
        tmp_data2.y = ((dst_index + 9  >= dst_start) && (dst_index + 9  < dst_end)) ? tmp_data2.y : dst_data2.y;
        tmp_data2.z = ((dst_index + 10 >= dst_start) && (dst_index + 10 < dst_end)) ? tmp_data2.z : dst_data2.z;
        tmp_data2.w = ((dst_index + 11 >= dst_start) && (dst_index + 11 < dst_end)) ? tmp_data2.w : dst_data2.w;

        *((__global uchar4*)(mat_dst + dst_index + 0)) = tmp_data0;
        *((__global uchar4*)(mat_dst + dst_index + 4)) = tmp_data1;
        *((__global uchar4*)(mat_dst + dst_index + 8)) = tmp_data2;
    }
}
__kernel void merge_vector_C3_D1(__global char *mat_dst,  int dst_step,  int dst_offset,
                                 __global char *mat_src0, int src0_step, int src0_offset,
                                 __global char *mat_src1, int src1_step, int src1_offset,
                                 __global char *mat_src2, int src2_step, int src2_offset, int offset_cols,
                                 int rows, int cols, int dst_step1)
{
    int x = get_global_id(0);
    int y = get_global_id(1);

    if ((x < cols) && (y < rows))
    {
        x = x << 2;

        int src0_index = mad24(y, src0_step, x + src0_offset - offset_cols);
        int src1_index = mad24(y, src1_step, x + src1_offset - offset_cols);
        int src2_index = mad24(y, src2_step, x + src2_offset - offset_cols);

        int dst_start = mad24(y, dst_step, dst_offset);
        int dst_end   = mad24(y, dst_step, dst_offset + dst_step1);
        int dst_index = mad24(y, dst_step, dst_offset + 3 * x - offset_cols * 3);

        char data0_0 = *(mat_src0 + src0_index + 0);
        char data0_1 = *(mat_src0 + src0_index + 1);
        char data0_2 = *(mat_src0 + src0_index + 2);
        char data0_3 = *(mat_src0 + src0_index + 3);

        char data1_0 = *(mat_src1 + src1_index + 0);
        char data1_1 = *(mat_src1 + src1_index + 1);
        char data1_2 = *(mat_src1 + src1_index + 2);
        char data1_3 = *(mat_src1 + src1_index + 3);

        char data2_0 = *(mat_src2 + src2_index + 0);
        char data2_1 = *(mat_src2 + src2_index + 1);
        char data2_2 = *(mat_src2 + src2_index + 2);
        char data2_3 = *(mat_src2 + src2_index + 3);

        char4 tmp_data0 = (char4)(data0_0, data1_0, data2_0, data0_1);
        char4 tmp_data1 = (char4)(data1_1, data2_1, data0_2, data1_2);
        char4 tmp_data2 = (char4)(data2_2, data0_3, data1_3, data2_3);

        char4 dst_data0 = *((__global char4*)(mat_dst + dst_index + 0));
        char4 dst_data1 = *((__global char4*)(mat_dst + dst_index + 4));
        char4 dst_data2 = *((__global char4*)(mat_dst + dst_index + 8));

        tmp_data0.x = ((dst_index + 0  >= dst_start) && (dst_index + 0  < dst_end)) ? tmp_data0.x : dst_data0.x;
        tmp_data0.y = ((dst_index + 1  >= dst_start) && (dst_index + 1  < dst_end)) ? tmp_data0.y : dst_data0.y;
        tmp_data0.z = ((dst_index + 2  >= dst_start) && (dst_index + 2  < dst_end)) ? tmp_data0.z : dst_data0.z;
        tmp_data0.w = ((dst_index + 3  >= dst_start) && (dst_index + 3  < dst_end)) ? tmp_data0.w : dst_data0.w;

        tmp_data1.x = ((dst_index + 4  >= dst_start) && (dst_index + 4  < dst_end)) ? tmp_data1.x : dst_data1.x;
        tmp_data1.y = ((dst_index + 5  >= dst_start) && (dst_index + 5  < dst_end)) ? tmp_data1.y : dst_data1.y;
        tmp_data1.z = ((dst_index + 6  >= dst_start) && (dst_index + 6  < dst_end)) ? tmp_data1.z : dst_data1.z;
        tmp_data1.w = ((dst_index + 7  >= dst_start) && (dst_index + 7  < dst_end)) ? tmp_data1.w : dst_data1.w;

        tmp_data2.x = ((dst_index + 8  >= dst_start) && (dst_index + 8  < dst_end)) ? tmp_data2.x : dst_data2.x;
        tmp_data2.y = ((dst_index + 9  >= dst_start) && (dst_index + 9  < dst_end)) ? tmp_data2.y : dst_data2.y;
        tmp_data2.z = ((dst_index + 10 >= dst_start) && (dst_index + 10 < dst_end)) ? tmp_data2.z : dst_data2.z;
        tmp_data2.w = ((dst_index + 11 >= dst_start) && (dst_index + 11 < dst_end)) ? tmp_data2.w : dst_data2.w;

        *((__global char4*)(mat_dst + dst_index + 0)) = tmp_data0;
        *((__global char4*)(mat_dst + dst_index + 4)) = tmp_data1;
        *((__global char4*)(mat_dst + dst_index + 8)) = tmp_data2;
    }
}
__kernel void merge_vector_C3_D2(__global ushort *mat_dst,  int dst_step,  int dst_offset,
                                 __global ushort *mat_src0, int src0_step, int src0_offset,
                                 __global ushort *mat_src1, int src1_step, int src1_offset,
                                 __global ushort *mat_src2, int src2_step, int src2_offset, int offset_cols,
                                 int rows, int cols, int dst_step1)
{
    int x = get_global_id(0);
    int y = get_global_id(1);

    if ((x < cols) && (y < rows))
    {
        x = x << 1;

        int src0_index = mad24(y, src0_step, (x << 1) + src0_offset - offset_cols);
        int src1_index = mad24(y, src1_step, (x << 1) + src1_offset - offset_cols);
        int src2_index = mad24(y, src2_step, (x << 1) + src2_offset - offset_cols);

        int dst_start = mad24(y, dst_step, dst_offset);
        int dst_end   = mad24(y, dst_step, dst_offset + dst_step1);
        int dst_index = mad24(y, dst_step, dst_offset + 6 * x - offset_cols * 6);

        ushort data0_0 = *((__global ushort *)((__global char *)mat_src0 + src0_index + 0));
        ushort data0_1 = *((__global ushort *)((__global char *)mat_src0 + src0_index + 2));

        ushort data1_0 = *((__global ushort *)((__global char *)mat_src1 + src1_index + 0));
        ushort data1_1 = *((__global ushort *)((__global char *)mat_src1 + src1_index + 2));

        ushort data2_0 = *((__global ushort *)((__global char *)mat_src2 + src2_index + 0));
        ushort data2_1 = *((__global ushort *)((__global char *)mat_src2 + src2_index + 2));

        ushort2 tmp_data0 = (ushort2)(data0_0, data1_0);
        ushort2 tmp_data1 = (ushort2)(data2_0, data0_1);
        ushort2 tmp_data2 = (ushort2)(data1_1, data2_1);

        ushort2 dst_data0 = *((__global ushort2*)((__global char *)mat_dst + dst_index + 0));
        ushort2 dst_data1 = *((__global ushort2*)((__global char *)mat_dst + dst_index + 4));
        ushort2 dst_data2 = *((__global ushort2*)((__global char *)mat_dst + dst_index + 8));

        tmp_data0.x = ((dst_index + 0  >= dst_start) && (dst_index + 0  < dst_end)) ? tmp_data0.x : dst_data0.x;
        tmp_data0.y = ((dst_index + 2  >= dst_start) && (dst_index + 2  < dst_end)) ? tmp_data0.y : dst_data0.y;

        tmp_data1.x = ((dst_index + 4  >= dst_start) && (dst_index + 4  < dst_end)) ? tmp_data1.x : dst_data1.x;
        tmp_data1.y = ((dst_index + 6  >= dst_start) && (dst_index + 6  < dst_end)) ? tmp_data1.y : dst_data1.y;

        tmp_data2.x = ((dst_index + 8  >= dst_start) && (dst_index + 8  < dst_end)) ? tmp_data2.x : dst_data2.x;
        tmp_data2.y = ((dst_index + 10 >= dst_start) && (dst_index + 10 < dst_end)) ? tmp_data2.y : dst_data2.y;

        *((__global ushort2*)((__global char *)mat_dst + dst_index + 0)) = tmp_data0;
        *((__global ushort2*)((__global char *)mat_dst + dst_index + 4)) = tmp_data1;
        *((__global ushort2*)((__global char *)mat_dst + dst_index + 8)) = tmp_data2;
    }
}
__kernel void merge_vector_C3_D3(__global short *mat_dst,  int dst_step,  int dst_offset,
                                 __global short *mat_src0, int src0_step, int src0_offset,
                                 __global short *mat_src1, int src1_step, int src1_offset,
                                 __global short *mat_src2, int src2_step, int src2_offset, int offset_cols,
                                 int rows, int cols, int dst_step1)
{
    int x = get_global_id(0);
    int y = get_global_id(1);

    if ((x < cols) && (y < rows))
    {
        x = x << 1;

        int src0_index = mad24(y, src0_step, (x << 1) + src0_offset - offset_cols);
        int src1_index = mad24(y, src1_step, (x << 1) + src1_offset - offset_cols);
        int src2_index = mad24(y, src2_step, (x << 1) + src2_offset - offset_cols);

        int dst_start = mad24(y, dst_step, dst_offset);
        int dst_end   = mad24(y, dst_step, dst_offset + dst_step1);
        int dst_index = mad24(y, dst_step, dst_offset + 6 * x - offset_cols * 6);

        short data0_0 = *((__global short *)((__global char *)mat_src0 + src0_index + 0));
        short data0_1 = *((__global short *)((__global char *)mat_src0 + src0_index + 2));

        short data1_0 = *((__global short *)((__global char *)mat_src1 + src1_index + 0));
        short data1_1 = *((__global short *)((__global char *)mat_src1 + src1_index + 2));

        short data2_0 = *((__global short *)((__global char *)mat_src2 + src2_index + 0));
        short data2_1 = *((__global short *)((__global char *)mat_src2 + src2_index + 2));

        short2 tmp_data0 = (short2)(data0_0, data1_0);
        short2 tmp_data1 = (short2)(data2_0, data0_1);
        short2 tmp_data2 = (short2)(data1_1, data2_1);

        short2 dst_data0 = *((__global short2*)((__global char *)mat_dst + dst_index + 0));
        short2 dst_data1 = *((__global short2*)((__global char *)mat_dst + dst_index + 4));
        short2 dst_data2 = *((__global short2*)((__global char *)mat_dst + dst_index + 8));

        tmp_data0.x = ((dst_index + 0  >= dst_start) && (dst_index + 0  < dst_end)) ? tmp_data0.x : dst_data0.x;
        tmp_data0.y = ((dst_index + 2  >= dst_start) && (dst_index + 2  < dst_end)) ? tmp_data0.y : dst_data0.y;

        tmp_data1.x = ((dst_index + 4  >= dst_start) && (dst_index + 4  < dst_end)) ? tmp_data1.x : dst_data1.x;
        tmp_data1.y = ((dst_index + 6  >= dst_start) && (dst_index + 6  < dst_end)) ? tmp_data1.y : dst_data1.y;

        tmp_data2.x = ((dst_index + 8  >= dst_start) && (dst_index + 8  < dst_end)) ? tmp_data2.x : dst_data2.x;
        tmp_data2.y = ((dst_index + 10 >= dst_start) && (dst_index + 10 < dst_end)) ? tmp_data2.y : dst_data2.y;

        *((__global short2*)((__global char *)mat_dst + dst_index + 0)) = tmp_data0;
        *((__global short2*)((__global char *)mat_dst + dst_index + 4)) = tmp_data1;
        *((__global short2*)((__global char *)mat_dst + dst_index + 8)) = tmp_data2;
    }
}
__kernel void merge_vector_C3_D4(__global int *mat_dst,  int dst_step,  int dst_offset,
                                 __global int *mat_src0, int src0_step, int src0_offset,
                                 __global int *mat_src1, int src1_step, int src1_offset,
                                 __global int *mat_src2, int src2_step, int src2_offset, int offset_cols,
                                 int rows, int cols, int dst_step1)
{
    int x = get_global_id(0);
    int y = get_global_id(1);

    if ((x < cols) && (y < rows))
    {
        int src0_index = mad24(y, src0_step, src0_offset);
        int src1_index = mad24(y, src1_step, src1_offset);
        int src2_index = mad24(y, src2_step, src2_offset);

        int dst_index  = mad24(y, dst_step , dst_offset);

        __global int* src0 = (__global int * )((__global uchar *)mat_src0 + src0_index + (x << 2));
        __global int* src1 = (__global int * )((__global uchar *)mat_src1 + src1_index + (x << 2));
        __global int* src2 = (__global int * )((__global uchar *)mat_src2 + src2_index + (x << 2));

        __global int* dist0 = (__global int *)((__global uchar *)mat_dst  + dst_index  + 3 * (x << 2));
        __global int* dist1 = dist0 + 1;
        __global int* dist2 = dist0 + 2;

        int  src0_data = *src0;
        int  src1_data = *src1;
        int  src2_data = *src2;

        *dist0 = src0_data;
        *dist1 = src1_data;
        *dist2 = src2_data;
    }
}
__kernel void merge_vector_C3_D5(__global float *mat_dst,  int dst_step,  int dst_offset,
                                 __global float *mat_src0, int src0_step, int src0_offset,
                                 __global float *mat_src1, int src1_step, int src1_offset,
                                 __global float *mat_src2, int src2_step, int src2_offset, int offset_cols,
                                 int rows, int cols, int dst_step1)
{
    int x = get_global_id(0);
    int y = get_global_id(1);

    if ((x < cols) && (y < rows))
    {
        int src0_index = mad24(y, src0_step, src0_offset);
        int src1_index = mad24(y, src1_step, src1_offset);
        int src2_index = mad24(y, src2_step, src2_offset);

        int dst_index  = mad24(y, dst_step , dst_offset);

        __global float* src0 = (__global float * )((__global uchar *)mat_src0 + src0_index + (x << 2));
        __global float* src1 = (__global float * )((__global uchar *)mat_src1 + src1_index + (x << 2));
        __global float* src2 = (__global float * )((__global uchar *)mat_src2 + src2_index + (x << 2));

        __global float* dist0 = (__global float *)((__global uchar *)mat_dst  + dst_index  + 3 * (x << 2));
        __global float* dist1 = dist0 + 1;
        __global float* dist2 = dist0 + 2;

        float  src0_data = *src0;
        float  src1_data = *src1;
        float  src2_data = *src2;

        *dist0 = src0_data;
        *dist1 = src1_data;
        *dist2 = src2_data;
    }
}

#if defined (DOUBLE_SUPPORT)
__kernel void merge_vector_C3_D6(__global double *mat_dst,  int dst_step,  int dst_offset,
                                 __global double *mat_src0, int src0_step, int src0_offset,
                                 __global double *mat_src1, int src1_step, int src1_offset,
                                 __global double *mat_src2, int src2_step, int src2_offset, int offset_cols,
                                 int rows, int cols, int dst_step1)
{
    int x = get_global_id(0);
    int y = get_global_id(1);

    if ((x < cols) && (y < rows))
    {
        int src0_index = mad24(y, src0_step, src0_offset);
        int src1_index = mad24(y, src1_step, src1_offset);
        int src2_index = mad24(y, src2_step, src2_offset);

        int dst_index  = mad24(y, dst_step , dst_offset);

        __global double* src0 = (__global double * )((__global uchar *)mat_src0 + src0_index + (x << 3));
        __global double* src1 = (__global double * )((__global uchar *)mat_src1 + src1_index + (x << 3));
        __global double* src2 = (__global double * )((__global uchar *)mat_src2 + src2_index + (x << 3));

        __global double* dist0 = (__global double *)((__global uchar *)mat_dst  + dst_index  + 3 * (x << 3));
        __global double* dist1 = dist0 + 1;
        __global double* dist2 = dist0 + 2;

        double  src0_data = *src0;
        double  src1_data = *src1;
        double  src2_data = *src2;

        *dist0 = src0_data;
        *dist1 = src1_data;
        *dist2 = src2_data;
    }
}
#endif
__kernel void merge_vector_C4_D0(__global uchar *mat_dst,  int dst_step,  int dst_offset,
                                 __global uchar *mat_src0, int src0_step, int src0_offset,
                                 __global uchar *mat_src1, int src1_step, int src1_offset,
                                 __global uchar *mat_src2, int src2_step, int src2_offset,
                                 __global uchar *mat_src3, int src3_step, int src3_offset,
                                 int rows, int cols, int dst_step1)
{
    int x = get_global_id(0);
    int y = get_global_id(1);

    if ((x < cols) && (y < rows))
    {
        int src0_index = mad24(y, src0_step, src0_offset);
        int src1_index = mad24(y, src1_step, src1_offset);
        int src2_index = mad24(y, src2_step, src2_offset);
        int src3_index = mad24(y, src3_step, src3_offset);
        int dst_index  = mad24(y, dst_step , dst_offset);

        uchar src0 = *(mat_src0 + src0_index + x );
        uchar src1 = *(mat_src1 + src1_index + x);
        uchar src2 = *(mat_src2 + src2_index + x);
        uchar src3 = *(mat_src3 + src3_index + x);

        *((__global uchar4 *)(mat_dst  + dst_index + (x << 2))) = (uchar4)(src0, src1, src2, src3);
    }
}
__kernel void merge_vector_C4_D1(__global char *mat_dst,  int dst_step,  int dst_offset,
                                 __global char *mat_src0, int src0_step, int src0_offset,
                                 __global char *mat_src1, int src1_step, int src1_offset,
                                 __global char *mat_src2, int src2_step, int src2_offset,
                                 __global char *mat_src3, int src3_step, int src3_offset,
                                 int rows, int cols, int dst_step1)
{
    int x = get_global_id(0);
    int y = get_global_id(1);

    if ((x < cols) && (y < rows))
    {
        int src0_index = mad24(y, src0_step, src0_offset);
        int src1_index = mad24(y, src1_step, src1_offset);
        int src2_index = mad24(y, src2_step, src2_offset);
        int src3_index = mad24(y, src3_step, src3_offset);
        int dst_index  = mad24(y, dst_step , dst_offset);

        char src0 = *(mat_src0 + src0_index + x );
        char src1 = *(mat_src1 + src1_index + x);
        char src2 = *(mat_src2 + src2_index + x);
        char src3 = *(mat_src3 + src3_index + x);

        *((__global char4 *)(mat_dst  + dst_index + (x << 2))) = (char4)(src0, src1, src2, src3);
    }
}
__kernel void merge_vector_C4_D2(__global ushort *mat_dst,  int dst_step,  int dst_offset,
                                 __global ushort *mat_src0, int src0_step, int src0_offset,
                                 __global ushort *mat_src1, int src1_step, int src1_offset,
                                 __global ushort *mat_src2, int src2_step, int src2_offset,
                                 __global ushort *mat_src3, int src3_step, int src3_offset,
                                 int rows, int cols, int dst_step1)
{
    int x = get_global_id(0);
    int y = get_global_id(1);

    if ((x < cols) && (y < rows))
    {
        int src0_index = mad24(y, src0_step, src0_offset);
        int src1_index = mad24(y, src1_step, src1_offset);
        int src2_index = mad24(y, src2_step, src2_offset);
        int src3_index = mad24(y, src3_step, src3_offset);
        int dst_index  = mad24(y, dst_step , dst_offset);

        ushort src0 = *((__global ushort *)((__global uchar *)mat_src0 + src0_index + (x << 1)));
        ushort src1 = *((__global ushort *)((__global uchar *)mat_src1 + src1_index + (x << 1)));
        ushort src2 = *((__global ushort *)((__global uchar *)mat_src2 + src2_index + (x << 1)));
        ushort src3 = *((__global ushort *)((__global uchar *)mat_src3 + src3_index + (x << 1)));

        *((__global ushort4 *)((__global uchar *)mat_dst  + dst_index + (x << 3))) = (ushort4)(src0, src1, src2, src3);
    }
}
__kernel void merge_vector_C4_D3(__global short *mat_dst,  int dst_step,  int dst_offset,
                                 __global short *mat_src0, int src0_step, int src0_offset,
                                 __global short *mat_src1, int src1_step, int src1_offset,
                                 __global short *mat_src2, int src2_step, int src2_offset,
                                 __global short *mat_src3, int src3_step, int src3_offset,
                                 int rows, int cols, int dst_step1)
{
    int x = get_global_id(0);
    int y = get_global_id(1);

    if ((x < cols) && (y < rows))
    {
        int src0_index = mad24(y, src0_step, src0_offset);
        int src1_index = mad24(y, src1_step, src1_offset);
        int src2_index = mad24(y, src2_step, src2_offset);
        int src3_index = mad24(y, src3_step, src3_offset);
        int dst_index  = mad24(y, dst_step , dst_offset);

        short src0 = *((__global short *)((__global uchar *)mat_src0 + src0_index + (x << 1)));
        short src1 = *((__global short *)((__global uchar *)mat_src1 + src1_index + (x << 1)));
        short src2 = *((__global short *)((__global uchar *)mat_src2 + src2_index + (x << 1)));
        short src3 = *((__global short *)((__global uchar *)mat_src3 + src3_index + (x << 1)));

        *((__global short4 *)((__global uchar *)mat_dst  + dst_index + (x << 3))) = (short4)(src0, src1, src2, src3);
    }
}
__kernel void merge_vector_C4_D4(__global int *mat_dst,  int dst_step,  int dst_offset,
                                 __global int *mat_src0, int src0_step, int src0_offset,
                                 __global int *mat_src1, int src1_step, int src1_offset,
                                 __global int *mat_src2, int src2_step, int src2_offset,
                                 __global int *mat_src3, int src3_step, int src3_offset,
                                 int rows, int cols, int dst_step1)
{
    int x = get_global_id(0);
    int y = get_global_id(1);

    if ((x < cols) && (y < rows))
    {
        int src0_index = mad24(y, src0_step, src0_offset);
        int src1_index = mad24(y, src1_step, src1_offset);
        int src2_index = mad24(y, src2_step, src2_offset);
        int src3_index = mad24(y, src3_step, src3_offset);
        int dst_index  = mad24(y, dst_step , dst_offset);

        int src0 = *((__global int *)((__global uchar *)mat_src0 + src0_index + (x << 2)));
        int src1 = *((__global int *)((__global uchar *)mat_src1 + src1_index + (x << 2)));
        int src2 = *((__global int *)((__global uchar *)mat_src2 + src2_index + (x << 2)));
        int src3 = *((__global int *)((__global uchar *)mat_src3 + src3_index + (x << 2)));

        *((__global int4 *)((__global uchar *)mat_dst  + dst_index + (x << 4))) = (int4)(src0, src1, src2, src3);
    }
}
__kernel void merge_vector_C4_D5(__global float *mat_dst,  int dst_step,  int dst_offset,
                                 __global float *mat_src0, int src0_step, int src0_offset,
                                 __global float *mat_src1, int src1_step, int src1_offset,
                                 __global float *mat_src2, int src2_step, int src2_offset,
                                 __global float *mat_src3, int src3_step, int src3_offset,
                                 int rows, int cols, int dst_step1)
{
    int x = get_global_id(0);
    int y = get_global_id(1);

    if ((x < cols) && (y < rows))
    {
        int src0_index = mad24(y, src0_step, src0_offset);
        int src1_index = mad24(y, src1_step, src1_offset);
        int src2_index = mad24(y, src2_step, src2_offset);
        int src3_index = mad24(y, src3_step, src3_offset);
        int dst_index  = mad24(y, dst_step , dst_offset);

        float src0 = *((__global float *)((__global uchar *)mat_src0 + src0_index + (x << 2)));
        float src1 = *((__global float *)((__global uchar *)mat_src1 + src1_index + (x << 2)));
        float src2 = *((__global float *)((__global uchar *)mat_src2 + src2_index + (x << 2)));
        float src3 = *((__global float *)((__global uchar *)mat_src3 + src3_index + (x << 2)));

        *((__global float4 *)((__global uchar *)mat_dst  + dst_index + (x << 4))) = (float4)(src0, src1, src2, src3);
    }
}

#if defined (DOUBLE_SUPPORT)
__kernel void merge_vector_C4_D6(__global double *mat_dst,  int dst_step,  int dst_offset,
                                 __global double *mat_src0, int src0_step, int src0_offset,
                                 __global double *mat_src1, int src1_step, int src1_offset,
                                 __global double *mat_src2, int src2_step, int src2_offset,
                                 __global double *mat_src3, int src3_step, int src3_offset,
                                 int rows, int cols, int dst_step1)
{
    int x = get_global_id(0);
    int y = get_global_id(1);

    if ((x < cols) && (y < rows))
    {
        int src0_index = mad24(y, src0_step, src0_offset);
        int src1_index = mad24(y, src1_step, src1_offset);
        int src2_index = mad24(y, src2_step, src2_offset);
        int src3_index = mad24(y, src3_step, src3_offset);
        int dst_index  = mad24(y, dst_step , dst_offset);

        double src0 = *((__global double *)((__global uchar *)mat_src0 + src0_index + (x << 3)));
        double src1 = *((__global double *)((__global uchar *)mat_src1 + src1_index + (x << 3)));
        double src2 = *((__global double *)((__global uchar *)mat_src2 + src2_index + (x << 3)));
        double src3 = *((__global double *)((__global uchar *)mat_src3 + src3_index + (x << 3)));

        *((__global double4 *)((__global uchar *)mat_dst  + dst_index + (x << 5))) = (double4)(src0, src1, src2, src3);
    }
}
#endif
///////////////////////////////////////////////////////////////////////////////////////////////
//////////////////////////////////optimized code using vector  no roi//////////////////////////
////////////vector fuction name format: merge_vector_C(channels number)D_(data type depth)//////
////////////////////////////////////////////////////////////////////////////////////////////////
__kernel void merge_vector_C2_D0_1(int rows, int cols,
                                   __global uchar *mat_dst,  int dst_step,
                                   __global uchar *mat_src0, int src0_step,
                                   __global uchar *mat_src1, int src1_step)
{
    int x = get_global_id(0);
    int y = get_global_id(1);

    if ((x < cols) && (y < rows))
    {
        __global uchar4  *src0_y = (__global uchar4 * )(mat_src0 + y * src0_step);
        __global uchar4  *src1_y = (__global uchar4 * )(mat_src1 + y * src1_step);
        __global uchar8 *dst_y  = (__global uchar8 *)(mat_dst  + y * dst_step);

        uchar4 value1 = src0_y[x];
        uchar4 value2 = src1_y[x];

        uchar8 value;
        value.even = value1;
        value.odd = value2;

        dst_y[x] = value;
    }
}
__kernel void merge_vector_C2_D1_1(int rows, int cols,
                                   __global char *mat_dst,  int dst_step,
                                   __global char *mat_src0, int src0_step,
                                   __global char *mat_src1, int src1_step)
{
    int x = get_global_id(0);
    int y = get_global_id(1);

    if ((x < cols) && (y < rows))
    {
        __global char4  *src0_y = (__global char4 * )(mat_src0 + y * src0_step);
        __global char4  *src1_y = (__global char4 * )(mat_src1 + y * src1_step);
        __global char8 *dst_y  = (__global char8 *)(mat_dst  + y * dst_step);

        char4 value1 = src0_y[x];
        char4 value2 = src1_y[x];

        char8 value;
        value.even = value1;
        value.odd = value2;

        dst_y[x] = value;
    }
}
__kernel void merge_vector_C2_D2_1(int rows, int cols,
                                   __global ushort *mat_dst,  int dst_step,
                                   __global ushort *mat_src0, int src0_step,
                                   __global ushort *mat_src1, int src1_step)
{
    int x = get_global_id(0);
    int y = get_global_id(1);

    if ((x < cols) && (y < rows))
    {
        __global ushort2  *src0_y = (__global ushort2 *)((__global uchar *)mat_src0 + y * src0_step);
        __global ushort2  *src1_y = (__global ushort2 *)((__global uchar *)mat_src1 + y * src1_step);
        __global ushort4  *dst_y  = (__global ushort4 *)((__global uchar *)mat_dst  + y * dst_step);

        ushort2 value1 = src0_y[x];
        ushort2 value2 = src1_y[x];

        ushort4 value;
        value.even = value1;
        value.odd = value2;

        dst_y[x] = value;
    }
}
__kernel void merge_vector_C2_D3_1(int rows, int cols,
                                   __global short *mat_dst,  int dst_step,
                                   __global short *mat_src0, int src0_step,
                                   __global short *mat_src1, int src1_step)
{
    int x = get_global_id(0);
    int y = get_global_id(1);

    if ((x < cols) && (y < rows))
    {
        __global short2  *src0_y = (__global short2 *)((__global uchar *)mat_src0 + y * src0_step);
        __global short2  *src1_y = (__global short2 *)((__global uchar *)mat_src1 + y * src1_step);
        __global short4 *dst_y   = (__global short4 *)((__global uchar *)mat_dst  + y * dst_step);

        short2 value1 = src0_y[x];
        short2 value2 = src1_y[x];

        short4 value;
        value.even = value1;
        value.odd = value2;

        dst_y[x] = value;
    }
}

__kernel void merge_vector_C2_D4_1(int rows, int cols,
                                   __global int *mat_dst,  int dst_step,
                                   __global int *mat_src0, int src0_step,
                                   __global int *mat_src1, int src1_step)
{
    int x = get_global_id(0);
    int y = get_global_id(1);

    if ((x < cols) && (y < rows))
    {
        __global int  *src0_y = (__global int *)((__global uchar *)mat_src0 + y * src0_step);
        __global int  *src1_y = (__global int *)((__global uchar *)mat_src1 + y * src1_step);
        __global int2  *dst_y  = (__global int2 *)((__global uchar *)mat_dst  + y * dst_step);

        int value1 = src0_y[x];
        int value2 = src1_y[x];

        int2 value;
        value.even = value1;
        value.odd = value2;

        dst_y[x] = value;
    }
}
__kernel void merge_vector_C2_D5_1(int rows, int cols,
                                   __global float *mat_dst,  int dst_step,
                                   __global float *mat_src0, int src0_step,
                                   __global float *mat_src1, int src1_step)
{
    int x = get_global_id(0);
    int y = get_global_id(1);

    if ((x < cols) && (y < rows))
    {
        __global float  *src0_y = (__global float *)((__global uchar *)mat_src0 + y * src0_step);
        __global float  *src1_y = (__global float *)((__global uchar *)mat_src1 + y * src1_step);
        __global float2  *dst_y  = (__global float2 *)((__global uchar *)mat_dst  + y * dst_step);

        float value1 = src0_y[x];
        float value2 = src1_y[x];

        dst_y[x] = (float2)(value1, value2);
    }
}

#if defined (DOUBLE_SUPPORT)
__kernel void merge_vector_C2_D6_1(int rows, int cols,
                                   __global double *mat_dst,  int dst_step,
                                   __global double *mat_src0, int src0_step,
                                   __global double *mat_src1, int src1_step)
{
    int x = get_global_id(0);
    int y = get_global_id(1);

    if ((x < cols) && (y < rows))
    {
        __global double  *src0_y = (__global double *)((__global uchar *)mat_src0 + y * src0_step);
        __global double  *src1_y = (__global double *)((__global uchar *)mat_src1 + y * src1_step);
        __global double2 *dst_y  = (__global double2 *)((__global uchar *)mat_dst  + y * dst_step);

        double value1 = src0_y[x];
        double value2 = src1_y[x];

        dst_y[x] = (double2)(value1, value2);
    }
}
#endif

__kernel void merge_vector_C3_D0_1(int rows, int cols,
                                   __global uchar *mat_dst,  int dst_step,
                                   __global uchar *mat_src0, int src0_step,
                                   __global uchar *mat_src1, int src1_step,
                                   __global uchar *mat_src2, int src2_step)
{
    int x = get_global_id(0);
    int y = get_global_id(1);

    if ((x < cols) && (y < rows))
    {
        __global uchar4  *src0_y = (__global uchar4 * )(mat_src0 + y * src0_step);
        __global uchar4  *src1_y = (__global uchar4 * )(mat_src1 + y * src1_step);
        __global uchar4  *src2_y = (__global uchar4 * )(mat_src2 + y * src0_step);

        __global uchar4 *dst_y  = (__global uchar4 *)(mat_dst  + y * dst_step);

        uchar4 value0 = src0_y[x];
        uchar4 value1 = src1_y[x];
        uchar4 value2 = src2_y[x];

        dst_y[3 * x + 0] = (uchar4)(value0.s0, value1.s0, value2.s0,
                                    value0.s1);

        dst_y[3 * x + 1] = (uchar4)(value1.s1, value2.s1,
                                    value0.s2, value1.s2);

        dst_y[3 * x + 2] = (uchar4)(value2.s2,
                                    value0.s3, value1.s3, value2.s3);

    }
}
__kernel void merge_vector_C3_D1_1(int rows, int cols,
                                   __global char *mat_dst,  int dst_step,
                                   __global char *mat_src0, int src0_step,
                                   __global char *mat_src1, int src1_step,
                                   __global char *mat_src2, int src2_step)
{
    int x = get_global_id(0);
    int y = get_global_id(1);

    if ((x < cols) && (y < rows))
    {
        __global char4  *src0_y = (__global char4 * )(mat_src0 + y * src0_step);
        __global char4  *src1_y = (__global char4 * )(mat_src1 + y * src1_step);
        __global char4  *src2_y = (__global char4 * )(mat_src2 + y * src0_step);

        __global char4 *dst_y  = (__global char4 *)(mat_dst  + y * dst_step);

        char4 value0 = src0_y[x];
        char4 value1 = src1_y[x];
        char4 value2 = src2_y[x];

        dst_y[3 * x + 0] = (char4)(value0.s0, value1.s0, value2.s0,
                                   value0.s1);

        dst_y[3 * x + 1] = (char4)(value1.s1, value2.s1,
                                     value0.s2, value1.s2);

        dst_y[3 * x + 2] = (char4)(value2.s2,
                                     value0.s3, value1.s3, value2.s3);

        /* for test do not delete
        dst_y[3 * x + 0] = (char8)(value0.s0, value1.s0, value2.s0,
                                    value0.s1, value1.s1, value2.s1,
                                    value0.s2, value1.s2);

        dst_y[3 * x + 1] = (char8)(value2.s2,
                                    value0.s3, value1.s3, value2.s3,
                                    value0.s4, value1.s4, value2.s4,
                                    value0.s5);

        dst_y[3 * x + 2] = (char8)(value1.s5, value2.s5,
                                    value0.s6, value1.s6, value2.s6,
                                    value0.s7, value1.s7, value2.s7);
                                    */
    }
}
__kernel void merge_vector_C3_D2_1(int rows, int cols,
                                   __global ushort *mat_dst,  int dst_step,
                                   __global ushort *mat_src0, int src0_step,
                                   __global ushort *mat_src1, int src1_step,
                                   __global ushort *mat_src2, int src2_step)
{
    int x = get_global_id(0);
    int y = get_global_id(1);

    if ((x < cols) && (y < rows))
    {
        __global ushort2  *src0_y = (__global ushort2 * )((__global char *)mat_src0 + y * src0_step);
        __global ushort2  *src1_y = (__global ushort2 * )((__global char *)mat_src1 + y * src1_step);
        __global ushort2  *src2_y = (__global ushort2 * )((__global char *)mat_src2 + y * src0_step);

        __global ushort2 *dst_y  = (__global ushort2 *)((__global char *)mat_dst  + y * dst_step);

        ushort2 value0 = src0_y[x];
        ushort2 value1 = src1_y[x];
        ushort2 value2 = src2_y[x];

        dst_y[3 * x + 0] = (ushort2)(value0.x, value1.x);
        dst_y[3 * x + 1] = (ushort2)(value2.x, value0.y);
        dst_y[3 * x + 2] = (ushort2)(value1.y, value2.y);

    }
}
__kernel void merge_vector_C3_D3_1(int rows, int cols,
                                   __global short *mat_dst,  int dst_step,
                                   __global short *mat_src0, int src0_step,
                                   __global short *mat_src1, int src1_step,
                                   __global short *mat_src2, int src2_step)
{
    int x = get_global_id(0);
    int y = get_global_id(1);

    if ((x < cols) && (y < rows))
    {
        __global short2  *src0_y = (__global short2 * )((__global char *)mat_src0 + y * src0_step);
        __global short2  *src1_y = (__global short2 * )((__global char *)mat_src1 + y * src1_step);
        __global short2  *src2_y = (__global short2 * )((__global char *)mat_src2 + y * src0_step);

        __global short2 *dst_y  = (__global short2 *)((__global char *)mat_dst  + y * dst_step);

        short2 value0 = src0_y[x];
        short2 value1 = src1_y[x];
        short2 value2 = src2_y[x];

        dst_y[3 * x + 0] = (short2)(value0.x, value1.x);
        dst_y[3 * x + 1] = (short2)(value2.x, value0.y);
        dst_y[3 * x + 2] = (short2)(value1.y, value2.y);

        /*
        dst_y[3 * x + 0] = (short4)(value0.s0, value1.s0, value2.s0,
                                    value0.s1);

        dst_y[3 * x + 1] = (short4)(value1.s1, value2.s1,
                                    value0.s2, value1.s2);

        dst_y[3 * x + 2] = (short4)(value2.s2,
                                    value0.s3, value1.s3, value2.s3);
                                    */
    }
}
__kernel void merge_vector_C3_D4_1(int rows, int cols,
                                   __global int *mat_dst,  int dst_step,
                                   __global int *mat_src0, int src0_step,
                                   __global int *mat_src1, int src1_step,
                                   __global int *mat_src2, int src2_step)
{
    int x = get_global_id(0);
    int y = get_global_id(1);

    if ((x < cols) && (y < rows))
    {
        __global int  *src0_y = (__global int * )((__global char *)mat_src0 + y * src0_step);
        __global int  *src1_y = (__global int * )((__global char *)mat_src1 + y * src1_step);
        __global int  *src2_y = (__global int * )((__global char *)mat_src2 + y * src0_step);

        __global int *dst_y  = (__global int *)((__global char *)mat_dst  + y * dst_step);

        int value0 = src0_y[x];
        int value1 = src1_y[x];
        int value2 = src2_y[x];

        dst_y[3 * x + 0] = value0;
        dst_y[3 * x + 1] = value1;
        dst_y[3 * x + 2] = value2;

        /*for test do not delete
        dst_y[3 * x + 0] = (int2)(value0.x, value1.x);
        dst_y[3 * x + 1] = (int2)(value2.x, value0.y);
        dst_y[3 * x + 2] = (int2)(value1.y, value2.y);
        */
    }
}
__kernel void merge_vector_C3_D5_1(int rows, int cols,
                                   __global float *mat_dst,  int dst_step,
                                   __global float *mat_src0, int src0_step,
                                   __global float *mat_src1, int src1_step,
                                   __global float *mat_src2, int src2_step)
{
    int x = get_global_id(0);
    int y = get_global_id(1);

    if ((x < cols) && (y < rows))
    {
        __global float  *src0_y = (__global float * )((__global char *)mat_src0 + y * src0_step);
        __global float  *src1_y = (__global float * )((__global char *)mat_src1 + y * src1_step);
        __global float  *src2_y = (__global float * )((__global char *)mat_src2 + y * src0_step);

        __global float *dst_y  = (__global float *)((__global char *)mat_dst  + y * dst_step);

        float value0 = src0_y[x];
        float value1 = src1_y[x];
        float value2 = src2_y[x];

        dst_y[3 * x + 0] = value0;
        dst_y[3 * x + 1] = value1;
        dst_y[3 * x + 2] = value2;
    }
}

#if defined (DOUBLE_SUPPORT)
__kernel void merge_vector_C3_D6_1(int rows, int cols,
                                   __global double *mat_dst,  int dst_step,
                                   __global double *mat_src0, int src0_step,
                                   __global double *mat_src1, int src1_step,
                                   __global double *mat_src2, int src2_step)
{
    int x = get_global_id(0);
    int y = get_global_id(1);

    if ((x < cols) && (y < rows))
    {
        __global double  *src0_y = (__global double * )((__global char *)mat_src0 + y * src0_step);
        __global double  *src1_y = (__global double * )((__global char *)mat_src1 + y * src1_step);
        __global double  *src2_y = (__global double * )((__global char *)mat_src2 + y * src0_step);

        __global double *dst_y  = (__global double *)((__global char *)mat_dst  + y * dst_step);

        double value0 = src0_y[x];
        double value1 = src1_y[x];
        double value2 = src2_y[x];

        dst_y[3 * x + 0] = value0;
        dst_y[3 * x + 1] = value1;
        dst_y[3 * x + 2] = value2;
    }
}
#endif
__kernel void merge_vector_C4_D0_1(int rows, int cols,
                                   __global uchar *mat_dst,  int dst_step,
                                   __global uchar *mat_src0, int src0_step,
                                   __global uchar *mat_src1, int src1_step,
                                   __global uchar *mat_src2, int src2_step,
                                   __global uchar *mat_src3, int src3_step)
{
    int x = get_global_id(0);
    int y = get_global_id(1);

    if ((x < cols) && (y < rows))
    {
        __global uchar4  *src0_y = (__global uchar4 * )(mat_src0 + y * src0_step);
        __global uchar4  *src1_y = (__global uchar4 * )(mat_src1 + y * src1_step);
        __global uchar4  *src2_y = (__global uchar4 * )(mat_src2 + y * src0_step);
        __global uchar4  *src3_y = (__global uchar4 * )(mat_src3 + y * src1_step);

        __global uchar16 *dst_y  = (__global uchar16 *)(mat_dst  + y * dst_step);

        uchar4 value0 = src0_y[x];
        uchar4 value1 = src1_y[x];
        uchar4 value2 = src2_y[x];
        uchar4 value3 = src3_y[x];

        dst_y[x] = (uchar16)(value0.x, value1.x, value2.x, value3.x,
                             value0.y, value1.y, value2.y, value3.y,
                             value0.z, value1.z, value2.z, value3.z,
                             value0.w, value1.w, value2.w, value3.w);
    }
}

__kernel void merge_vector_C4_D1_1(int rows, int cols,
                                   __global char *mat_dst,  int dst_step,
                                   __global char *mat_src0, int src0_step,
                                   __global char *mat_src1, int src1_step,
                                   __global char *mat_src2, int src2_step,
                                   __global char *mat_src3, int src3_step)
{
    int x = get_global_id(0);
    int y = get_global_id(1);

    if ((x < cols) && (y < rows))
    {
        __global char4  *src0_y = (__global char4 * )(mat_src0 + y * src0_step);
        __global char4  *src1_y = (__global char4 * )(mat_src1 + y * src1_step);
        __global char4  *src2_y = (__global char4 * )(mat_src2 + y * src0_step);
        __global char4  *src3_y = (__global char4 * )(mat_src3 + y * src1_step);

        __global char16 *dst_y  = (__global char16 *)(mat_dst  + y * dst_step);

        char4 value0 = src0_y[x];
        char4 value1 = src1_y[x];
        char4 value2 = src2_y[x];
        char4 value3 = src3_y[x];

        dst_y[x] = (char16)(value0.x, value1.x, value2.x, value3.x,
                            value0.y, value1.y, value2.y, value3.y,
                            value0.z, value1.z, value2.z, value3.z,
                            value0.w, value1.w, value2.w, value3.w);
    }
}
__kernel void merge_vector_C4_D2_1(int rows, int cols,
                                   __global ushort *mat_dst,  int dst_step,
                                   __global ushort *mat_src0, int src0_step,
                                   __global ushort *mat_src1, int src1_step,
                                   __global ushort *mat_src2, int src2_step,
                                   __global ushort *mat_src3, int src3_step)
{
    int x = get_global_id(0);
    int y = get_global_id(1);

    if ((x < cols) && (y < rows))
    {
        __global ushort2  *src0_y = (__global ushort2 * )((__global uchar*)mat_src0 + y * src0_step);
        __global ushort2  *src1_y = (__global ushort2 * )((__global uchar*)mat_src1 + y * src1_step);
        __global ushort2  *src2_y = (__global ushort2 * )((__global uchar*)mat_src2 + y * src0_step);
        __global ushort2  *src3_y = (__global ushort2 * )((__global uchar*)mat_src3 + y * src1_step);

        __global ushort8 *dst_y  = (__global ushort8 *)((__global uchar*)mat_dst  + y * dst_step);

        ushort2 value0 = src0_y[x];
        ushort2 value1 = src1_y[x];
        ushort2 value2 = src2_y[x];
        ushort2 value3 = src3_y[x];

        dst_y[x] = (ushort8)(value0.x, value1.x, value2.x, value3.x,
                             value0.y, value1.y, value2.y, value3.y);
    }
}
__kernel void merge_vector_C4_D3_1(int rows, int cols,
                                   __global short *mat_dst,  int dst_step,
                                   __global short *mat_src0, int src0_step,
                                   __global short *mat_src1, int src1_step,
                                   __global short *mat_src2, int src2_step,
                                   __global short *mat_src3, int src3_step)
{
    int x = get_global_id(0);
    int y = get_global_id(1);

    if ((x < cols) && (y < rows))
    {
        __global short2  *src0_y = (__global short2 * )((__global uchar*)mat_src0 + y * src0_step);
        __global short2  *src1_y = (__global short2 * )((__global uchar*)mat_src1 + y * src1_step);
        __global short2  *src2_y = (__global short2 * )((__global uchar*)mat_src2 + y * src0_step);
        __global short2  *src3_y = (__global short2 * )((__global uchar*)mat_src3 + y * src1_step);

        __global short8 *dst_y  = (__global short8 *)((__global uchar*)mat_dst  + y * dst_step);

        short2 value0 = src0_y[x];
        short2 value1 = src1_y[x];
        short2 value2 = src2_y[x];
        short2 value3 = src3_y[x];

        dst_y[x] = (short8)(value0.x, value1.x, value2.x, value3.x,
                            value0.y, value1.y, value2.y, value3.y);
    }
}
__kernel void merge_vector_C4_D4_1(int rows, int cols,
                                   __global int *mat_dst,  int dst_step,
                                   __global int *mat_src0, int src0_step,
                                   __global int *mat_src1, int src1_step,
                                   __global int *mat_src2, int src2_step,
                                   __global int *mat_src3, int src3_step)
{
    int x = get_global_id(0);
    int y = get_global_id(1);

    if ((x < cols) && (y < rows))
    {
        __global int *src0_y = (__global int * )((__global uchar*)mat_src0 + y * src0_step);
        __global int *src1_y = (__global int * )((__global uchar*)mat_src1 + y * src1_step);
        __global int *src2_y = (__global int * )((__global uchar*)mat_src2 + y * src0_step);
        __global int *src3_y = (__global int * )((__global uchar*)mat_src3 + y * src1_step);

        __global int4 *dst_y  = (__global int4 *)((__global uchar*)mat_dst  + y * dst_step);

        int value0 = src0_y[x];
        int value1 = src1_y[x];
        int value2 = src2_y[x];
        int value3 = src3_y[x];

        dst_y[x] = (int4)(value0, value1, value2, value3);
    }
}
__kernel void merge_vector_C4_D5_1(int rows, int cols,
                                   __global float *mat_dst,  int dst_step,
                                   __global float *mat_src0, int src0_step,
                                   __global float *mat_src1, int src1_step,
                                   __global float *mat_src2, int src2_step,
                                   __global float *mat_src3, int src3_step)
{
    int x = get_global_id(0);
    int y = get_global_id(1);

    if ((x < cols) && (y < rows))
    {
        __global float *src0_y = (__global float * )((__global uchar*)mat_src0 + y * src0_step);
        __global float *src1_y = (__global float * )((__global uchar*)mat_src1 + y * src1_step);
        __global float *src2_y = (__global float * )((__global uchar*)mat_src2 + y * src0_step);
        __global float *src3_y = (__global float * )((__global uchar*)mat_src3 + y * src1_step);

        __global float4 *dst_y  = (__global float4 *)((__global uchar*)mat_dst  + y * dst_step);

        float value0 = src0_y[x];
        float value1 = src1_y[x];
        float value2 = src2_y[x];
        float value3 = src3_y[x];

        dst_y[x] = (float4)(value0, value1, value2, value3);
    }
}

#if defined (DOUBLE_SUPPORT)
__kernel void merge_vector_C4_D6_1(int rows, int cols,
                                   __global double *mat_dst,  int dst_step,
                                   __global double *mat_src0, int src0_step,
                                   __global double *mat_src1, int src1_step,
                                   __global double *mat_src2, int src2_step,
                                   __global double *mat_src3, int src3_step)
{
    int x = get_global_id(0);
    int y = get_global_id(1);

    if ((x < cols) && (y < rows))
    {
        __global double *src0_y = (__global double * )((__global uchar*)mat_src0 + y * src0_step);
        __global double *src1_y = (__global double * )((__global uchar*)mat_src1 + y * src1_step);
        __global double *src2_y = (__global double * )((__global uchar*)mat_src2 + y * src0_step);
        __global double *src3_y = (__global double * )((__global uchar*)mat_src3 + y * src1_step);

        __global double4 *dst_y  = (__global double4 *)((__global uchar*)mat_dst  + y * dst_step);

        double value0 = src0_y[x];
        double value1 = src1_y[x];
        double value2 = src2_y[x];
        double value3 = src3_y[x];

        dst_y[x] = (double4)(value0, value1, value2, value3);
    }
}
#endif