/*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, Multicoreware, Inc., all rights reserved.
// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved.
// Third party copyrights are property of their respective owners.
//
// @Authors
//    Jin Ma jin@multicorewareinc.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 oclMaterials 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*/

__kernel void buildMotionMapsKernel(__global float* forwardMotionX,
    __global float* forwardMotionY,
    __global float* backwardMotionX,
    __global float* backwardMotionY,
    __global float* forwardMapX,
    __global float* forwardMapY,
    __global float* backwardMapX,
    __global float* backwardMapY,
    int forwardMotionX_row,
    int forwardMotionX_col,
    int forwardMotionX_step,
    int forwardMotionY_step,
    int backwardMotionX_step,
    int backwardMotionY_step,
    int forwardMapX_step,
    int forwardMapY_step,
    int backwardMapX_step,
    int backwardMapY_step
    )
{
    int x = get_global_id(0);
    int y = get_global_id(1);

    if(x < forwardMotionX_col && y < forwardMotionX_row)
    {
        float fx = forwardMotionX[y * forwardMotionX_step + x];
        float fy = forwardMotionY[y * forwardMotionY_step + x];

        float bx = backwardMotionX[y * backwardMotionX_step + x];
        float by = backwardMotionY[y * backwardMotionY_step + x];

        forwardMapX[y * forwardMapX_step + x] = x + bx;
        forwardMapY[y * forwardMapY_step + x] = y + by;

        backwardMapX[y * backwardMapX_step + x] = x + fx;
        backwardMapY[y * backwardMapY_step + x] = y + fy;
    }
}

__kernel void upscaleKernel(__global float* src,
    __global float* dst,
    int src_step,
    int dst_step,
    int src_row,
    int src_col,
    int scale,
    int channels
    )
{
    int x = get_global_id(0);
    int y = get_global_id(1);

    if(x < src_col && y < src_row)
    {
        if(channels == 1)
        {
            dst[y * scale * dst_step + x * scale] = src[y * src_step + x];
        }else if(channels == 3)
        {
            dst[y * channels * scale * dst_step + 3 * x * scale + 0] = src[y * channels * src_step + 3 * x + 0];
            dst[y * channels * scale * dst_step + 3 * x * scale + 1] = src[y * channels * src_step + 3 * x + 1];
            dst[y * channels * scale * dst_step + 3 * x * scale + 2] = src[y * channels * src_step + 3 * x + 2];
        }else
        {
            dst[y * channels * scale * dst_step + 4 * x * scale + 0] = src[y * channels * src_step + 4 * x + 0];
            dst[y * channels * scale * dst_step + 4 * x * scale + 1] = src[y * channels * src_step + 4 * x + 1];
            dst[y * channels * scale * dst_step + 4 * x * scale + 2] = src[y * channels * src_step + 4 * x + 2];
            dst[y * channels * scale * dst_step + 4 * x * scale + 3] = src[y * channels * src_step + 4 * x + 3];
        }
    }
}


float diffSign(float a, float b)
{
    return a > b ? 1.0f : a < b ? -1.0f : 0.0f;
}

float3 diffSign3(float3 a, float3 b)
{
    float3 pos;
    pos.x = a.x > b.x ? 1.0f : a.x < b.x ? -1.0f : 0.0f;
    pos.y = a.y > b.y ? 1.0f : a.y < b.y ? -1.0f : 0.0f;
    pos.z = a.z > b.z ? 1.0f : a.z < b.z ? -1.0f : 0.0f;
    return pos;
}

float4 diffSign4(float4 a, float4 b)
{
    float4 pos;
    pos.x = a.x > b.x ? 1.0f : a.x < b.x ? -1.0f : 0.0f;
    pos.y = a.y > b.y ? 1.0f : a.y < b.y ? -1.0f : 0.0f;
    pos.z = a.z > b.z ? 1.0f : a.z < b.z ? -1.0f : 0.0f;
    pos.w = 0.0f;
    return pos;
}

__kernel void diffSignKernel(__global float* src1,
    __global float* src2,
    __global float* dst,
    int src1_row,
    int src1_col,
    int dst_step,
    int src1_step,
    int src2_step)
{
    int x = get_global_id(0);
    int y = get_global_id(1);

    if(x < src1_col && y < src1_row)
    {
        dst[y * dst_step + x] = diffSign(src1[y * src1_step + x], src2[y * src2_step + x]);
    }
    barrier(CLK_LOCAL_MEM_FENCE);
}

__kernel void calcBtvRegularizationKernel(__global float* src,
    __global float* dst,
    int src_step,
    int dst_step,
    int src_row,
    int src_col,
    int ksize,
    int channels,
    __global float* c_btvRegWeights
    )
{
    int x = get_global_id(0) + ksize;
    int y = get_global_id(1) + ksize;

    if ((y < src_row - ksize) && (x < src_col - ksize))
    {
        if(channels == 1)
        {
            const float srcVal = src[y * src_step + x];
            float dstVal = 0.0f;

            for (int m = 0, count = 0; m <= ksize; ++m)
            {
                for (int l = ksize; l + m >= 0; --l, ++count)
                    dstVal = dstVal + c_btvRegWeights[count] * (diffSign(srcVal, src[(y + m) * src_step + (x + l)]) - diffSign(src[(y - m) * src_step + (x - l)], srcVal));
            }
            dst[y * dst_step + x] = dstVal;
        }else if(channels == 3)
        {
            float3 srcVal;
            srcVal.x = src[y * src_step + 3 * x + 0];
            srcVal.y = src[y * src_step + 3 * x + 1];
            srcVal.z = src[y * src_step + 3 * x + 2];

            float3 dstVal;
            dstVal.x = 0.0f;
            dstVal.y = 0.0f;
            dstVal.z = 0.0f;

            for (int m = 0, count = 0; m <= ksize; ++m)
            {
                for (int l = ksize; l + m >= 0; --l, ++count)
                {
                    float3 src1;
                    src1.x = src[(y + m) * src_step + 3 * (x + l) + 0];
                    src1.y = src[(y + m) * src_step + 3 * (x + l) + 1];
                    src1.z = src[(y + m) * src_step + 3 * (x + l) + 2];

                    float3 src2;
                    src2.x = src[(y - m) * src_step + 3 * (x - l) + 0];
                    src2.y = src[(y - m) * src_step + 3 * (x - l) + 1];
                    src2.z = src[(y - m) * src_step + 3 * (x - l) + 2];

                    dstVal = dstVal + c_btvRegWeights[count] * (diffSign3(srcVal, src1) - diffSign3(src2, srcVal));
                }
            }
            dst[y * dst_step + 3 * x + 0] = dstVal.x;
            dst[y * dst_step + 3 * x + 1] = dstVal.y;
            dst[y * dst_step + 3 * x + 2] = dstVal.z;
        }else
        {
            float4 srcVal;
            srcVal.x = src[y * src_step + 4 * x + 0];//r type =float
            srcVal.y = src[y * src_step + 4 * x + 1];//g
            srcVal.z = src[y * src_step + 4 * x + 2];//b
            srcVal.w = src[y * src_step + 4 * x + 3];//a

            float4 dstVal;
            dstVal.x = 0.0f;
            dstVal.y = 0.0f;
            dstVal.z = 0.0f;
            dstVal.w = 0.0f;

            for (int m = 0, count = 0; m <= ksize; ++m)
            {
                for (int l = ksize; l + m >= 0; --l, ++count)
                {
                    float4 src1;
                    src1.x = src[(y + m) * src_step + 4 * (x + l) + 0];
                    src1.y = src[(y + m) * src_step + 4 * (x + l) + 1];
                    src1.z = src[(y + m) * src_step + 4 * (x + l) + 2];
                    src1.w = src[(y + m) * src_step + 4 * (x + l) + 3];

                    float4 src2;
                    src2.x = src[(y - m) * src_step + 4 * (x - l) + 0];
                    src2.y = src[(y - m) * src_step + 4 * (x - l) + 1];
                    src2.z = src[(y - m) * src_step + 4 * (x - l) + 2];
                    src2.w = src[(y - m) * src_step + 4 * (x - l) + 3];

                    dstVal = dstVal + c_btvRegWeights[count] * (diffSign4(srcVal, src1) - diffSign4(src2, srcVal));

                }
            }
            dst[y * dst_step + 4 * x + 0] = dstVal.x;
            dst[y * dst_step + 4 * x + 1] = dstVal.y;
            dst[y * dst_step + 4 * x + 2] = dstVal.z;
            dst[y * dst_step + 4 * x + 3] = dstVal.w;
        }
    }
}