--- /dev/null
+// This file is part of OpenCV project.
+// It is subject to the license terms in the LICENSE file found in the top-level directory
+// of this distribution and at http://opencv.org/license.html.
+
+__constant float kx[] = { 0.125, 0.5, 0.75, 0.5, 0.125 };
+__constant float ky[] = { 0.125, 0.5, 0.75, 0.5, 0.125 };
+
+#define OP(delta, y, x) (convert_float4(arr[(y + delta) * 5 + x]) * ky[y] * kx[x])
+
+__kernel void pyrUp_cols2(__global const uchar * src, int src_step, int src_offset, int src_rows, int src_cols,
+ __global uchar * dst, int dst_step, int dst_offset, int dst_rows, int dst_cols)
+{
+ int block_x = get_global_id(0);
+ int y = get_global_id(1) * 2;
+
+ if ((block_x * 4) >= dst_cols || y >= dst_rows) return;
+
+ uchar8 line[6];
+ uchar4 line_out;
+
+ int offset, src_index;
+ src_index = block_x * 2 + (y / 2 - 1) * src_step - 1 + src_offset;
+
+ uchar4 tmp;
+
+ line[0] = line[2] = line[4] = (uchar8)0;
+ line[1] = line[3] = line[5] = (uchar8)0;
+
+ offset = max(0, src_index + 1 * src_step);
+ tmp = vload4(0, src + offset);
+ if (offset == 0) tmp = (uchar4)(0, tmp.s012);
+ line[2].even = tmp;
+
+ offset = max(0, src_index + ((y == 0) ? 2 : 0) * src_step);
+ tmp = vload4(0, src + offset);
+ if (offset == 0) tmp = (uchar4)(0, tmp.s012);
+ line[0].even = tmp;
+
+ if (y == (dst_rows - 2))
+ line[4] = line[2];
+ else
+ line[4].even = vload4(0, src + src_index + 2 * src_step);
+
+ bool row_s = (block_x == 0);
+ bool row_e = ((block_x + 1) * 4 == dst_cols);
+ uchar4 arr[30];
+ uchar s, e;
+
+ s = line[0].s4;
+ e = line[0].s3;
+ arr[0] = row_s ? (uchar4)(s, e, line[0].s23) : (uchar4)(line[0].s0123);
+ arr[1] = row_s ? (uchar4)(e, line[0].s234) : (uchar4)(line[0].s1234);
+ arr[2] = (uchar4)(line[0].s2345);
+ arr[3] = row_e ? (uchar4)(line[0].s345, s) : (uchar4)(line[0].s3456);
+ arr[4] = row_e ? (uchar4)(line[0].s45, s, e) : (uchar4)(line[0].s4567);
+
+ s = line[1].s4;
+ e = line[1].s3;
+ arr[5] = row_s ? (uchar4)(s, e, line[1].s23) : (uchar4)(line[1].s0123);
+ arr[6] = row_s ? (uchar4)(e, line[1].s234) : (uchar4)(line[1].s1234);
+ arr[7] = (uchar4)(line[1].s2345);
+ arr[8] = row_e ? (uchar4)(line[1].s345, s) : (uchar4)(line[1].s3456);
+ arr[9] = row_e ? (uchar4)(line[1].s45, s, e) : (uchar4)(line[1].s4567);
+
+ s = line[2].s4;
+ e = line[2].s3;
+ arr[10] = row_s ? (uchar4)(s, e, line[2].s23) : (uchar4)(line[2].s0123);
+ arr[11] = row_s ? (uchar4)(e, line[2].s234) : (uchar4)(line[2].s1234);
+ arr[12] = (uchar4)(line[2].s2345);
+ arr[13] = row_e ? (uchar4)(line[2].s345, s) : (uchar4)(line[2].s3456);
+ arr[14] = row_e ? (uchar4)(line[2].s45, s, e) : (uchar4)(line[2].s4567);
+
+ s = line[3].s4;
+ e = line[3].s3;
+ arr[15] = row_s ? (uchar4)(s, e, line[3].s23) : (uchar4)(line[3].s0123);
+ arr[16] = row_s ? (uchar4)(e, line[3].s234) : (uchar4)(line[3].s1234);
+ arr[17] = (uchar4)(line[3].s2345);
+ arr[18] = row_e ? (uchar4)(line[3].s345, s) : (uchar4)(line[3].s3456);
+ arr[19] = row_e ? (uchar4)(line[3].s45, s, e) : (uchar4)(line[3].s4567);
+
+ s = line[4].s4;
+ e = line[4].s3;
+ arr[20] = row_s ? (uchar4)(s, e, line[4].s23) : (uchar4)(line[4].s0123);
+ arr[21] = row_s ? (uchar4)(e, line[4].s234) : (uchar4)(line[4].s1234);
+ arr[22] = (uchar4)(line[4].s2345);
+ arr[23] = row_e ? (uchar4)(line[4].s345, s) : (uchar4)(line[4].s3456);
+ arr[24] = row_e ? (uchar4)(line[4].s45, s, e) : (uchar4)(line[4].s4567);
+
+ s = line[5].s4;
+ e = line[5].s3;
+ arr[25] = row_s ? (uchar4)(s, e, line[5].s23) : (uchar4)(line[5].s0123);
+ arr[26] = row_s ? (uchar4)(e, line[5].s234) : (uchar4)(line[5].s1234);
+ arr[27] = (uchar4)(line[5].s2345);
+ arr[28] = row_e ? (uchar4)(line[5].s345, s) : (uchar4)(line[5].s3456);
+ arr[29] = row_e ? (uchar4)(line[5].s45, s, e) : (uchar4)(line[5].s4567);
+
+ float4 sum[2];
+
+ sum[0] = OP(0, 0, 0) + OP(0, 0, 1) + OP(0, 0, 2) + OP(0, 0, 3) + OP(0, 0, 4) +
+ OP(0, 1, 0) + OP(0, 1, 1) + OP(0, 1, 2) + OP(0, 1, 3) + OP(0, 1, 4) +
+ OP(0, 2, 0) + OP(0, 2, 1) + OP(0, 2, 2) + OP(0, 2, 3) + OP(0, 2, 4) +
+ OP(0, 3, 0) + OP(0, 3, 1) + OP(0, 3, 2) + OP(0, 3, 3) + OP(0, 3, 4) +
+ OP(0, 4, 0) + OP(0, 4, 1) + OP(0, 4, 2) + OP(0, 4, 3) + OP(0, 4, 4);
+
+ sum[1] = OP(1, 0, 0) + OP(1, 0, 1) + OP(1, 0, 2) + OP(1, 0, 3) + OP(1, 0, 4) +
+ OP(1, 1, 0) + OP(1, 1, 1) + OP(1, 1, 2) + OP(1, 1, 3) + OP(1, 1, 4) +
+ OP(1, 2, 0) + OP(1, 2, 1) + OP(1, 2, 2) + OP(1, 2, 3) + OP(1, 2, 4) +
+ OP(1, 3, 0) + OP(1, 3, 1) + OP(1, 3, 2) + OP(1, 3, 3) + OP(1, 3, 4) +
+ OP(1, 4, 0) + OP(1, 4, 1) + OP(1, 4, 2) + OP(1, 4, 3) + OP(1, 4, 4);
+
+ int dst_index = block_x * 4 + y * dst_step + dst_offset;
+ vstore4(convert_uchar4_sat_rte(sum[0]), 0, dst + dst_index);
+ vstore4(convert_uchar4_sat_rte(sum[1]), 0, dst + dst_index + dst_step);
+}