/*
  * Copyright (C) 2013 Wei Gao <weigao@multicorewareinc.com>

  * Copyright (C) 2013 Lenny Wang

  *
  * This file is part of FFmpeg.
  *
  * FFmpeg is free software; you can redistribute it and/or
  * modify it under the terms of the GNU Lesser General Public
  * License as published by the Free Software Foundation; either
  * version 2.1 of the License, or (at your option) any later version.
  *
  * FFmpeg is distributed in the hope that it will be useful,
  * but WITHOUT ANY WARRANTY; without even the implied warranty of
  * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU
  * Lesser General Public License for more details.
  *
  * You should have received a copy of the GNU Lesser General Public
  * License along with FFmpeg; if not, write to the Free Software
  * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA
  */
 

 #ifndef AVFILTER_UNSHARP_OPENCL_KERNEL_H
 #define AVFILTER_UNSHARP_OPENCL_KERNEL_H

 
 #include "libavutil/opencl.h"
 
 const char *ff_kernel_unsharp_opencl = AV_OPENCL_KERNEL(
 inline unsigned char clip_uint8(int a)
 {
     if (a & (~0xFF))
         return (-a)>>31;
     else
         return a;
 }
 

 kernel void unsharp_luma(
                     global unsigned char *src,
                     global unsigned char *dst,

                     global int *mask_x,
                     global int *mask_y,

                     int amount,
                     int scalebits,
                     int halfscale,
                     int src_stride,
                     int dst_stride,
                     int width,
                     int height)
 {
     int2 threadIdx, blockIdx, globalIdx;
     threadIdx.x = get_local_id(0);
     threadIdx.y = get_local_id(1);
     blockIdx.x = get_group_id(0);
     blockIdx.y = get_group_id(1);
     globalIdx.x = get_global_id(0);
     globalIdx.y = get_global_id(1);
 
     if (!amount) {
         if (globalIdx.x < width && globalIdx.y < height)
             dst[globalIdx.x + globalIdx.y*dst_stride] = src[globalIdx.x + globalIdx.y*src_stride];
         return;
     }
 

     local unsigned int l[32][32];
     local unsigned int lcx[LU_RADIUS_X];
     local unsigned int lcy[LU_RADIUS_Y];

     int indexIx, indexIy, i, j;
 

     //load up tile: actual workspace + halo of 8 points in x and y \n

     for(i = 0; i <= 1; i++) {
         indexIy = -8 + (blockIdx.y + i) * 16 + threadIdx.y;
         indexIy = indexIy < 0 ? 0 : indexIy;
         indexIy = indexIy >= height ? height - 1: indexIy;
         for(j = 0; j <= 1; j++) {
             indexIx = -8 + (blockIdx.x + j) * 16 + threadIdx.x;
             indexIx = indexIx < 0 ? 0 : indexIx;
             indexIx = indexIx >= width ? width - 1: indexIx;
             l[i*16 + threadIdx.y][j*16 + threadIdx.x] = src[indexIy*src_stride + indexIx];
         }
     }
 
     int indexL = threadIdx.y*16 + threadIdx.x;

     if (indexL < LU_RADIUS_X)
         lcx[indexL] = mask_x[indexL];
     if (indexL < LU_RADIUS_Y)
         lcy[indexL] = mask_y[indexL];

     barrier(CLK_LOCAL_MEM_FENCE);
 

     //needed for unsharp mask application in the end \n
     int orig_value = (int)l[threadIdx.y + 8][threadIdx.x + 8];

     int idx, idy, maskIndex;
     int temp[2] = {0};
     int steps_x = (LU_RADIUS_X-1)/2;
     int steps_y = (LU_RADIUS_Y-1)/2;
 
     // compute the actual workspace + left&right halos \n
       \n#pragma unroll\n
     for (j = 0; j <=1; j++) {
       //extra work to cover left and right halos \n
       idx = 16*j + threadIdx.x;
       \n#pragma unroll\n
         for (i = -steps_y; i <= steps_y; i++) {
           idy = 8 + i + threadIdx.y;
           maskIndex = (i + steps_y);
           temp[j] += (int)l[idy][idx] * lcy[maskIndex];

         }
     }

     barrier(CLK_LOCAL_MEM_FENCE);
     //save results from the vertical filter in local memory \n
     idy = 8 + threadIdx.y;
       \n#pragma unroll\n
     for (j = 0; j <=1; j++) {
       idx = 16*j + threadIdx.x;
       l[idy][idx] = temp[j];
     }
     barrier(CLK_LOCAL_MEM_FENCE);
 
     //compute results with the horizontal filter \n
     int sum = 0;
     idy = 8 + threadIdx.y;
     \n#pragma unroll\n
       for (j = -steps_x; j <= steps_x; j++) {
         idx = 8 + j + threadIdx.x;
         maskIndex = j + steps_x;
         sum += (int)l[idy][idx] * lcx[maskIndex];
       }
 
     int res = orig_value + (((orig_value - (int)((sum + halfscale) >> scalebits)) * amount) >> 16);
 

     if (globalIdx.x < width && globalIdx.y < height)
         dst[globalIdx.x + globalIdx.y*dst_stride] = clip_uint8(res);
 }
 
 kernel void unsharp_chroma(
                     global unsigned char *src_y,
                     global unsigned char *dst_y,

                     global int *mask_x,
                     global int *mask_y,

                     int amount,
                     int scalebits,
                     int halfscale,
                     int src_stride_lu,
                     int src_stride_ch,
                     int dst_stride_lu,
                     int dst_stride_ch,
                     int width,
                     int height,
                     int cw,
                     int ch)
 {
     global unsigned char *dst_u = dst_y + height * dst_stride_lu;
     global unsigned char *dst_v = dst_u + ch * dst_stride_ch;
     global unsigned char *src_u = src_y + height * src_stride_lu;
     global unsigned char *src_v = src_u + ch * src_stride_ch;
     int2 threadIdx, blockIdx, globalIdx;
     threadIdx.x = get_local_id(0);
     threadIdx.y = get_local_id(1);
     blockIdx.x = get_group_id(0);
     blockIdx.y = get_group_id(1);
     globalIdx.x = get_global_id(0);
     globalIdx.y = get_global_id(1);
     int padch = get_global_size(1)/2;
     global unsigned char *src = globalIdx.y>=padch ? src_v : src_u;
     global unsigned char *dst = globalIdx.y>=padch ? dst_v : dst_u;
 
     blockIdx.y = globalIdx.y>=padch ? blockIdx.y - get_num_groups(1)/2 : blockIdx.y;
     globalIdx.y = globalIdx.y>=padch ? globalIdx.y - padch : globalIdx.y;
 
     if (!amount) {
         if (globalIdx.x < cw && globalIdx.y < ch)
             dst[globalIdx.x + globalIdx.y*dst_stride_ch] = src[globalIdx.x + globalIdx.y*src_stride_ch];
         return;
     }
 

     local unsigned int l[32][32];
     local unsigned int lcx[CH_RADIUS_X];
     local unsigned int lcy[CH_RADIUS_Y];

     int indexIx, indexIy, i, j;
     for(i = 0; i <= 1; i++) {
         indexIy = -8 + (blockIdx.y + i) * 16 + threadIdx.y;
         indexIy = indexIy < 0 ? 0 : indexIy;
         indexIy = indexIy >= ch ? ch - 1: indexIy;
         for(j = 0; j <= 1; j++) {
             indexIx = -8 + (blockIdx.x + j) * 16 + threadIdx.x;
             indexIx = indexIx < 0 ? 0 : indexIx;
             indexIx = indexIx >= cw ? cw - 1: indexIx;
             l[i*16 + threadIdx.y][j*16 + threadIdx.x] = src[indexIy * src_stride_ch + indexIx];
         }
     }
 
     int indexL = threadIdx.y*16 + threadIdx.x;

     if (indexL < CH_RADIUS_X)
         lcx[indexL] = mask_x[indexL];
     if (indexL < CH_RADIUS_Y)
         lcy[indexL] = mask_y[indexL];

     barrier(CLK_LOCAL_MEM_FENCE);
 

     int orig_value = (int)l[threadIdx.y + 8][threadIdx.x + 8];
 

     int idx, idy, maskIndex;
     int steps_x = CH_RADIUS_X/2;
     int steps_y = CH_RADIUS_Y/2;

     int temp[2] = {0,0};

 
     \n#pragma unroll\n

       for (j = 0; j <= 1; j++) {
         idx = 16*j + threadIdx.x;

         \n#pragma unroll\n

           for (i = -steps_y; i <= steps_y; i++) {
             idy = 8 + i + threadIdx.y;
             maskIndex = i + steps_y;
             temp[j] += (int)l[idy][idx] * lcy[maskIndex];
           }
       }
 
     barrier(CLK_LOCAL_MEM_FENCE);
     idy = 8 + threadIdx.y;
     \n#pragma unroll\n
     for (j = 0; j <= 1; j++) {
       idx = 16*j + threadIdx.x;
       l[idy][idx] = temp[j];

     }

     barrier(CLK_LOCAL_MEM_FENCE);
 
     //compute results with the horizontal filter \n
     int sum = 0;
     idy = 8 + threadIdx.y;
     \n#pragma unroll\n
       for (j = -steps_x; j <= steps_x; j++) {
         idx = 8 + j + threadIdx.x;
         maskIndex = j + steps_x;
         sum += (int)l[idy][idx] * lcx[maskIndex];
       }
 
     int res = orig_value + (((orig_value - (int)((sum + halfscale) >> scalebits)) * amount) >> 16);
 

     if (globalIdx.x < cw && globalIdx.y < ch)
         dst[globalIdx.x + globalIdx.y*dst_stride_ch] = clip_uint8(res);
 }
 
 kernel void unsharp_default(global  unsigned char *src,

                     global  unsigned char *dst,
                     const global  unsigned int *mask_lu,
                     const global  unsigned int *mask_ch,
                     int amount_lu,
                     int amount_ch,
                     int step_x_lu,
                     int step_y_lu,
                     int step_x_ch,
                     int step_y_ch,
                     int scalebits_lu,
                     int scalebits_ch,
                     int halfscale_lu,
                     int halfscale_ch,
                     int src_stride_lu,
                     int src_stride_ch,
                     int dst_stride_lu,
                     int dst_stride_ch,
                     int height,
                     int width,
                     int ch,
                     int cw)
 {
     global unsigned char *dst_y = dst;
     global unsigned char *dst_u = dst_y + height * dst_stride_lu;
     global unsigned char *dst_v = dst_u + ch * dst_stride_ch;
 
     global unsigned char *src_y = src;
     global unsigned char *src_u = src_y + height * src_stride_lu;
     global unsigned char *src_v = src_u + ch * src_stride_ch;
 
     global unsigned char *temp_dst;
     global unsigned char *temp_src;
     const global unsigned int *temp_mask;
     int global_id = get_global_id(0);
     int i, j, x, y, temp_src_stride, temp_dst_stride, temp_height, temp_width, temp_steps_x, temp_steps_y,
         temp_amount, temp_scalebits, temp_halfscale, sum, idx_x, idx_y, temp, res;
     if (global_id < width * height) {
         y = global_id / width;
         x = global_id % width;
         temp_dst = dst_y;
         temp_src = src_y;
         temp_src_stride = src_stride_lu;
         temp_dst_stride = dst_stride_lu;
         temp_height = height;
         temp_width = width;
         temp_steps_x = step_x_lu;
         temp_steps_y = step_y_lu;
         temp_mask = mask_lu;
         temp_amount = amount_lu;
         temp_scalebits = scalebits_lu;
         temp_halfscale = halfscale_lu;
     } else if ((global_id >= width * height) && (global_id < width * height + ch * cw)) {
         y = (global_id - width * height) / cw;
         x = (global_id - width * height) % cw;
         temp_dst = dst_u;
         temp_src = src_u;
         temp_src_stride = src_stride_ch;
         temp_dst_stride = dst_stride_ch;
         temp_height = ch;
         temp_width = cw;
         temp_steps_x = step_x_ch;
         temp_steps_y = step_y_ch;
         temp_mask = mask_ch;
         temp_amount = amount_ch;
         temp_scalebits = scalebits_ch;
         temp_halfscale = halfscale_ch;
     } else {
         y = (global_id - width * height - ch * cw) / cw;
         x = (global_id - width * height - ch * cw) % cw;
         temp_dst = dst_v;
         temp_src = src_v;
         temp_src_stride = src_stride_ch;
         temp_dst_stride = dst_stride_ch;
         temp_height = ch;
         temp_width = cw;
         temp_steps_x = step_x_ch;
         temp_steps_y = step_y_ch;
         temp_mask = mask_ch;
         temp_amount = amount_ch;
         temp_scalebits = scalebits_ch;
         temp_halfscale = halfscale_ch;
     }
     if (temp_amount) {
         sum = 0;
         for (j = 0; j <= 2 * temp_steps_y; j++) {
             idx_y = (y - temp_steps_y + j) <= 0 ? 0 : (y - temp_steps_y + j) >= temp_height ? temp_height-1 : y - temp_steps_y + j;
             for (i = 0; i <= 2 * temp_steps_x; i++) {
                 idx_x = (x - temp_steps_x + i) <= 0 ? 0 : (x - temp_steps_x + i) >= temp_width ? temp_width-1 : x - temp_steps_x + i;
                 sum += temp_mask[i + j * (2 * temp_steps_x + 1)] * temp_src[idx_x + idx_y * temp_src_stride];
             }
         }
         temp = (int)temp_src[x + y * temp_src_stride];
         res = temp + (((temp - (int)((sum + temp_halfscale) >> temp_scalebits)) * temp_amount) >> 16);
         temp_dst[x + y * temp_dst_stride] = clip_uint8(res);
     } else {
         temp_dst[x + y * temp_dst_stride] = temp_src[x + y * temp_src_stride];
     }
 }
 );
 

 #endif /* AVFILTER_UNSHARP_OPENCL_KERNEL_H */