opencv/modules/imgproc/src/opencl/pyramid_up.cl

// 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);
}