Merging-in the OpenCL Scaling code from the OpenCL branch to trunk.

Patch originally by the Multicoreware Inc team, followed by improvements and fixes by Micheal Wootton from AMD Inc,

OpenCL:
This patch implements Bicubic Scaling in OpenCL.  
Note that HandBrake currently uses Lanczos so the performance difference appears to be much more significant. We may offer an option of BiCubic in software later.
Bicubic scaling may appear a bit sharper than the equivalent Lanczos encode and may increase file size a bit. Quality may be better or worse depending on the scaling and content and personal preference towards sharpness.
When comparing performance with a custom HandBrake build that runs Software Bicubic to OpenCL Bicubic, performance increase is about 5~7% on average on a modern GPU.


Hardware Decode via DXVA:
We also have optional DXVA decoding which may come in useful for slower/lower end systems that have a capable GPU.
This is only available on input sources that use the libav decode path.
Most GPU hardware for decoding is designed for playback, so if you are running on a high end CPU, it will bottleneck the encode process.

Requires OpenCL 1.1 or later supporting GPU.

Front end changes and testing framework are not included in this patch. This will be resolved later.

Patch will be revised further before the UI is implemented.

git-svn-id: svn://svn.handbrake.fr/HandBrake/trunk@5792 b64f7644-9d1e-0410-96f1-a4d463321fa5

1 files changed, 771 insertions, 0 deletions

diff --git a/libhb/openclkernels.h b/libhb/openclkernels.h
new file mode 100644
index 000000000..f324aafdd
--- /dev/null
+++ b/libhb/openclkernels.h
@@ -0,0 +1,771 @@
+/* openclkernels.h

+

+   Copyright (c) 2003-2012 HandBrake Team

+   This file is part of the HandBrake source code

+   Homepage: <http://handbrake.fr/>.

+   It may be used under the terms of the GNU General Public License v2.

+   For full terms see the file COPYING file or visit http://www.gnu.org/licenses/gpl-2.0.html

+   

+   Authors: Peng Gao <[email protected]> <http://www.multicorewareinc.com/>

+            Li   Cao <[email protected]> <http://www.multicorewareinc.com/>

+

+ */

+ 

+#ifndef USE_EXTERNAL_KERNEL

+

+#define KERNEL( ... )# __VA_ARGS__

+

+

+char *kernel_src_hscale = KERNEL (

+

+    typedef unsigned char fixed8;

+

+/*******************************************************************************************************

+dst:          Horizontal scale destination;

+src:          YUV content in opencl buf;

+hf_Y:         Horizontal filter coefficients for Y planes;

+hf_UV:        Horizontal filter coefficients for UV planes;

+hi_Y:         Horizontal filter index for Y planes;

+hi_UV:        Horizontal filter index for UV planes;

+stride:       Src width;

+filter_len:   Length of filter;

+********************************************************************************************************/

+    kernel void frame_h_scale (

+        global fixed8 *src,

+        global float *hf_Y,

+        global float *hf_UV,

+        global int *hi_Y,

+        global int *hi_UV,

+        global fixed8 *dst,

+        int stride, //src_width

+        int filter_len

+        )

+    {

+        int x = get_global_id( 0 );

+        int y = get_global_id( 1 );

+        int width = get_global_size( 0 );

+        int height = get_global_size( 1 );

+        float result_Y = 0, result_U = 0, result_V = 0;

+        int i = 0;

+

+        global fixed8 *src_Y = src;

+        global fixed8 *src_U = src_Y + stride * height;

+        global fixed8 *src_V = src_U + (stride >> 1) * (height >> 1);

+

+        global fixed8 *dst_Y = dst;

+        global fixed8 *dst_U = dst_Y + width * height;

+        global fixed8 *dst_V = dst_U + (width >> 1) * (height >> 1);

+

+        int xy = y * width + x;

+        global fixed8 *rowdata_Y = src_Y + (y * stride);

+        for( int i = 0; i < filter_len; i++ )

+        {

+            result_Y += ( hf_Y[x + i * width] * rowdata_Y[hi_Y[x] + i]);

+        }

+        dst_Y[xy] = result_Y;

+

+        if( y < (height >> 1) && x < (width >> 1) )

+        {

+            int xy = y * (width >> 1) + x;

+            global fixed8 *rowdata_U = src_U + (y * (stride >> 1));

+            global fixed8 *rowdata_V = src_V + (y * (stride >> 1));

+            for( i = 0; i < filter_len; i++ )

+            {

+                result_U += ( hf_UV[x + i * (width >> 1)] * rowdata_U[hi_UV[x] + i]);

+                result_V += ( hf_UV[x + i * (width >> 1)] * rowdata_V[hi_UV[x] + i]);

+            }

+            dst_U[xy] = result_U;

+            dst_V[xy] = result_V;

+        }

+    }

+    );

+

+/*******************************************************************************************************

+dst:          Vertical scale destination;

+src:          YUV content in opencl buf;

+hf_Y:         Vertical filter coefficients for Y planes;

+hf_UV:        Vertical filter coefficients for UV planes;

+hi_Y:         Vertical filter index for Y planes;

+hi_UV:        Vertical filter index for UV planes;

+stride:       Src height;

+filter_len:   Length of filter;

+********************************************************************************************************/

+char *kernel_src_vscale = KERNEL (

+

+    kernel void frame_v_scale (

+        global fixed8 *src,

+        global float *vf_Y,

+        global float *vf_UV,

+        global int *vi_Y,

+        global int *vi_UV,

+        global fixed8 *dst,

+        int src_height,

+        int filter_len

+        )

+    {

+        int x = get_global_id( 0 );

+        int y = get_global_id( 1 );

+        int width = get_global_size( 0 );

+        int height = get_global_size( 1 );

+        float result_Y = 0, result_U = 0, result_V = 0;

+        int i = 0;

+

+        global fixed8 *src_Y = src;

+        global fixed8 *src_U = src_Y + src_height * width;

+        global fixed8 *src_V = src_U + (src_height >> 1) * (width >> 1);

+

+        global fixed8 *dst_Y = dst;

+        global fixed8 *dst_U = dst_Y + height * width;

+        global fixed8 *dst_V = dst_U + (height >> 1) * (width >> 1);

+

+        int xy = y * width + x;

+        for( i = 0; i < filter_len; i++ )

+        {

+            result_Y += vf_Y[y + i * height] * src_Y[(vi_Y[y] + i) * width + x];

+        }

+        dst_Y[xy] = result_Y;

+

+        if( y < (height >> 1) && x < (width >> 1) )

+        {

+            int xy = y * (width >> 1) + x;

+            for( i = 0; i < filter_len; i++ )

+            {

+                result_U += vf_UV[y + i * (height >> 1)] * src_U[(vi_UV[y] + i) * (width >> 1) + x];

+                result_V += vf_UV[y + i * (height >> 1)] * src_V[(vi_UV[y] + i) * (width >> 1) + x];

+            }

+            dst_U[xy] = result_U;

+            dst_V[xy] = result_V;

+        }

+    }

+    );

+

+/*******************************************************************************************************

+input:    Input buffer;

+output:   Output buffer;

+w:        Width of frame;

+h:        Height of frame;

+********************************************************************************************************/

+char *kernel_src_nvtoyuv = KERNEL (

+

+    kernel void nv12toyuv ( global char *input, global char* output, int w, int h )

+    {

+        int x = get_global_id( 0 );

+        int y = get_global_id( 1 );

+        int idx = y * (w >> 1) + x;

+        vstore4((vload4( 0, input + (idx << 2))), 0, output + (idx << 2)); //Y

+        char2 uv = vload2( 0, input + (idx << 1) + w * h );

+        output[idx + w * h] = uv.s0;

+        output[idx + w * h + ((w * h) >> 2)] = uv.s1;

+    }

+    );

+

+/*******************************************************************************************************

+dst:           Horizontal scale destination;

+src:           YUV content in opencl buf;

+yfilter:       Opencl memory of horizontal filter coefficients for luma/alpha planes;

+yfilterPos:    Opencl memory of horizontal filter starting positions for each dst[i] for luma/alpha planes;

+yfilterSize:   Horizontal filter size for luma/alpha pixels;

+cfilter:       Opencl memory of horizontal filter coefficients for chroma planes;

+cfilterPos:    Opencl memory of horizontal filter starting positions for each dst[i] for chroma planes;

+cfilterSize:   Horizontal filter size for chroma pixels;

+dstStride:     Width of destination luma/alpha planes;

+dstChrStride:  Width of destination chroma planes;

+********************************************************************************************************/

+

+char *kernel_src_hscaleall = KERNEL (

+

+    kernel void hscale_all_opencl (

+        global short *dst,

+        const global unsigned char *src,

+        const global short *yfilter,

+        const global int *yfilterPos,

+        int yfilterSize,

+        const global short *cfilter,

+        const global int *cfilterPos,

+        int cfilterSize,

+        int dstWidth,

+        int dstHeight,

+        int srcWidth,

+        int srcHeight,

+        int dstStride,

+        int dstChrStride,

+        int srcStride,

+        int srcChrStride)

+    {

+        int w = get_global_id(0);

+        int h = get_global_id(1);

+

+        int chrWidth = get_global_size(0);

+        int chrHeight = get_global_size(1);

+

+        int srcPos1 = h * srcStride + yfilterPos[w];

+        int srcPos2 = h * srcStride + yfilterPos[w + chrWidth];

+        int srcPos3 = (h + (srcHeight >> 1)) * srcStride + yfilterPos[w];

+        int srcPos4 = (h + (srcHeight >> 1)) * srcStride + yfilterPos[w + chrWidth];

+        int srcc1Pos = srcStride * srcHeight + (h) * (srcChrStride) + cfilterPos[w];

+        int srcc2Pos = srcc1Pos + ((srcChrStride)*(chrHeight));

+

+        int val1 = 0;

+        int val2 = 0;

+        int val3 = 0;

+        int val4 = 0;

+        int val5 = 0;

+        int val6 = 0;

+

+        int filterPos1 = yfilterSize * w;

+        int filterPos2 = yfilterSize * (w + chrWidth);

+        int cfilterPos1 = cfilterSize * w;

+

+        int j;

+        for (j = 0; j < yfilterSize; j++)

+        {

+            val1 += src[srcPos1 + j] * yfilter[filterPos1+ j];

+            val2 += src[srcPos2 + j] * yfilter[filterPos2 + j];

+            val3 += src[srcPos3 + j] * yfilter[filterPos1 + j];

+            val4 += src[srcPos4 + j] * yfilter[filterPos2 + j];

+            val5 += src[srcc1Pos+j] * cfilter[cfilterPos1 + j];

+            val6 += src[srcc2Pos+j] * cfilter[cfilterPos1 + j];

+        }

+        int dstPos1 = h *dstStride;

+        int dstPos2 = (h + chrHeight) * dstStride;

+

+        dst[dstPos1 + w] = ((val1 >> 7) > ((1 << 15) - 1) ? ((1 << 15) - 1) : (val1 >> 7));

+        dst[dstPos1 + w + chrWidth] = ((val2 >> 7) > ((1 << 15) - 1) ? ((1 << 15) - 1) : (val2 >> 7));

+        dst[dstPos2 + w] = ((val3 >> 7) > ((1 << 15) - 1) ? ((1 << 15) - 1) : (val3 >> 7));

+        dst[dstPos2 + w + chrWidth] = ((val4 >> 7) > ((1 << 15) - 1) ? ((1 << 15) - 1) : (val4 >> 7));

+

+        int dstPos3 = h * (dstChrStride) + w + dstStride * dstHeight;

+        int dstPos4 = h * (dstChrStride) + w + dstStride * dstHeight + ((dstChrStride) * chrHeight);

+        dst[dstPos3] = ((val5 >> 7) > ((1 << 15) - 1) ? ((1 << 15) - 1) : (val5 >> 7));

+        dst[dstPos4] = ((val6 >> 7) > ((1 << 15) - 1) ? ((1 << 15) - 1) : (val6 >> 7));

+    }

+    );

+

+char *kernel_src_hscalefast = KERNEL (

+

+    kernel void hscale_fast_opencl (

+        global short *dst,

+        const global unsigned char *src,

+        int xInc,

+        int chrXInc,

+        int dstWidth,

+        int dstHeight,

+        int srcWidth,

+        int srcHeight,

+        int dstStride,

+        int dstChrStride,

+        int srcStride,

+        int srcChrStride)

+    {

+

+        int w = get_global_id(0);

+        int h = get_global_id(1);

+

+        int chrWidth = get_global_size(0);

+        int chrHeight = get_global_size(1);

+        int xpos1 = 0;

+        int xpos2 = 0;

+        int xx = xpos1 >> 16;

+        int xalpha = (xpos1 & 0xFFFF) >> 9;

+        dst[h * dstStride + w] = (src[h * srcStride + xx] << 7) + (src[h * srcStride + xx + 1] -src[h * srcStride + xx]) * xalpha;

+        int lowpart = h + (chrHeight);

+        dst[lowpart * dstStride + w] = (src[lowpart * srcStride + xx] << 7) + (src[lowpart * srcStride + xx + 1] - src[lowpart * srcStride + xx]) * xalpha;

+

+        int inv_i = w * xInc >> 16;

+        if( inv_i >= srcWidth - 1)

+        {

+            dst[h*dstStride + w] = src[h*srcStride + srcWidth-1]*128;

+            dst[lowpart*dstStride + w] = src[lowpart*srcStride + srcWidth - 1] * 128;

+        }

+

+        int rightpart = w + (chrWidth);

+        xx = xpos2 >> 16;

+        xalpha = (xpos2 & 0xFFFF) >> 9;

+        dst[h * dstStride + rightpart] = (src[h *srcStride + xx] << 7) + (src[h * srcStride + xx + 1] - src[h * srcStride + xx]) * xalpha;

+        dst[lowpart * dstStride + rightpart] = (src[lowpart * srcStride + xx] << 7) + (src[lowpart * srcStride + xx + 1] - src[lowpart * srcStride + xx]) * xalpha;

+        inv_i = rightpart * xInc >> 16;

+        if( inv_i >= srcWidth - 1)

+        {

+            dst[h * dstStride + rightpart] = src[h * srcStride + srcWidth - 1] * 128;

+            dst[lowpart * dstStride + rightpart] = src[lowpart * srcStride + srcWidth - 1] * 128;

+        }

+

+        int xpos = 0;

+        xpos = chrXInc * w;

+        xx = xpos >> 16;

+        xalpha = (xpos & 0xFFFF) >> 9;

+        src += srcStride * srcHeight;

+        dst += dstStride * dstHeight;

+        dst[h * (dstChrStride) + w] = (src[h * (srcChrStride) + xx] * (xalpha^127) + src[h * (srcChrStride) + xx + 1] * xalpha);

+        inv_i = w * xInc >> 16;

+        if( inv_i >= (srcWidth >> 1) - 1)

+        {

+            dst[h * (dstChrStride) + w] = src[h * (srcChrStride) + (srcWidth >> 1) -1]*128;

+        }

+

+        xpos = chrXInc * (w);

+        xx = xpos >> 16;

+        src += srcChrStride * srcHeight >> 1;

+        dst += (dstChrStride * chrHeight);

+        dst[h * (dstChrStride) + w] = (src[h * (srcChrStride) + xx] * (xalpha^127) + src[h * (srcChrStride) + xx + 1 ] * xalpha);

+

+        if( inv_i >= (srcWidth >> 1) - 1)

+        {

+            //v channel:

+            dst[h * (dstChrStride) + w] = src[h * (srcChrStride) + (srcWidth >> 1) -1] * 128;

+        }

+    }

+    );

+

+char *kernel_src_vscalealldither = KERNEL (

+

+    kernel void vscale_all_dither_opencl (

+        global unsigned char *dst,

+        const global short *src,

+        const global short *yfilter,

+        int yfilterSize,

+        const global short *cfilter,

+        int cfilterSize,

+        const global int *yfilterPos,

+        const global int *cfilterPos,

+        int dstWidth,

+        int dstHeight,

+        int srcWidth,

+        int srcHeight,

+        int dstStride,

+        int dstChrStride,

+        int srcStride,

+        int srcChrStride)

+    {

+        const unsigned char hb_dither_8x8_128[8][8] = {

+            {  36, 68,  60, 92,  34, 66,  58, 90, },

+            { 100,  4, 124, 28,  98,  2, 122, 26, },

+            {  52, 84,  44, 76,  50, 82,  42, 74, },

+            { 116, 20, 108, 12, 114, 18, 106, 10, },

+            {  32, 64,  56, 88,  38, 70,  62, 94, },

+            {  96,  0, 120, 24, 102,  6, 126, 30, },

+            {  48, 80,  40, 72,  54, 86,  46, 78, },

+            { 112, 16, 104,  8, 118, 22, 110, 14, },

+        };

+

+

+        int w = get_global_id(0);

+        int h = get_global_id(1);

+

+        int chrWidth = get_global_size(0);

+        int chrHeight = get_global_size(1);

+        const unsigned char *local_up_dither;

+        const unsigned char *local_down_dither;

+

+        local_up_dither = hb_dither_8x8_128[h & 7];

+        local_down_dither = hb_dither_8x8_128[(h + chrHeight) & 7];

+

+        //yscale;

+        int srcPos1 = (yfilterPos[h]) * srcStride + w;

+        int srcPos2 = (yfilterPos[h]) * srcStride + w + (chrWidth);

+        int srcPos3 = (yfilterPos[h + chrHeight]) * srcStride + w;

+        int srcPos4 = (yfilterPos[h + chrHeight]) * srcStride + w + chrWidth;

+        int src1Pos = dstStride * srcHeight + (cfilterPos[h]) * dstChrStride + (w);

+        int src2Pos = dstStride * srcHeight + (dstChrStride*(srcHeight>>1)) + (cfilterPos[h]) * dstChrStride + w;

+

+        int val1 = (local_up_dither[w & 7] << 12); //y offset is 0;

+        int val2 = (local_up_dither[(w + chrWidth) & 7] << 12);

+        int val3 = (local_down_dither[w &7] << 12);

+        int val4 = (local_down_dither[(w + chrWidth) & 7] << 12);

+        int val5 = (local_up_dither[w & 7] << 12);

+        int val6 = (local_up_dither[(w + 3) & 7] << 12);   // 3 is offset of the chrome channel.

+

+        int j;

+        int filterPos1 = h * yfilterSize;

+        int filterPos2 = ( h + chrHeight ) * yfilterSize;

+        for(j = 0; j < yfilterSize; j++)

+        {

+            val1 += src[srcPos1] * yfilter[filterPos1 + j];

+            srcPos1 += srcStride;

+            val2 += src[srcPos2] * yfilter[filterPos1 + j];

+            srcPos2 += srcStride;

+            val3 += src[srcPos3] * yfilter[filterPos2 + j];

+            srcPos3 += srcStride;

+            val4 += src[srcPos4] * yfilter[filterPos2 + j];

+            srcPos4 += srcStride;

+            val5 += src[src1Pos] * cfilter[filterPos1 + j];

+            val6 += src[src2Pos] * cfilter[filterPos1 + j];

+            src1Pos += dstChrStride;

+            src2Pos += dstChrStride;

+        }

+        dst[h * dstStride + w] = (((val1 >> 19)&(~0xFF)) ? ((-(val1 >> 19)) >> 31) : (val1 >> 19));

+        dst[h * dstStride + w + chrWidth] = (((val2 >> 19)&(~0xFF)) ? ((-(val2 >> 19)) >> 31) : (val2 >> 19));

+        dst[(h + chrHeight) * dstStride + w] = (((val3 >> 19)&(~0xFF)) ? ((-(val3 >> 19)) >> 31) : (val3 >> 19));

+        dst[(h + chrHeight) * dstStride + w + chrWidth] = (((val4 >> 19)&(~0xFF)) ? ((-(val4 >> 19)) >> 31) : (val4 >> 19));

+

+        int dst1Pos = dstStride * dstHeight + h*(dstChrStride)+(w);

+        int dst2Pos = (dstChrStride * chrHeight) + dst1Pos;

+        dst[dst1Pos] = (((val5 >> 19)&(~0xFF)) ? ((-(val5 >> 19)) >> 31) : (val5 >> 19));

+        dst[dst2Pos] = (((val6 >> 19)&(~0xFF)) ? ((-(val6 >> 19)) >> 31) : (val6 >> 19));

+    }

+    );

+

+char *kernel_src_vscaleallnodither = KERNEL (

+

+    kernel void vscale_all_nodither_opencl (

+        global unsigned char *dst,

+        const global short *src,

+        const global short *yfilter,

+        int yfilterSize,

+        const global short *cfilter,

+        int cfilterSize,

+        const global int *yfilterPos,

+        const global int *cfilterPos,

+        int dstWidth,

+        int dstHeight,

+        int srcWidth,

+        int srcHeight,

+        int dstStride,

+        int dstChrStride,

+        int srcStride,

+        int srcChrStride)

+    {

+        const unsigned char hb_sws_pb_64[8] = {

+            64, 64, 64, 64, 64, 64, 64, 64

+        };

+

+        int w = get_global_id(0);

+        int h = get_global_id(1);

+

+        int chrWidth = get_global_size(0);

+        int chrHeight = get_global_size(1);

+        const unsigned char *local_up_dither;

+        const unsigned char *local_down_dither;

+

+        local_up_dither = hb_sws_pb_64;

+        local_down_dither = hb_sws_pb_64;

+

+

+        //yscale;

+        int srcPos1 = (yfilterPos[h]) * srcStride + w;

+        int srcPos2 = (yfilterPos[h]) * srcStride + w + (chrWidth);

+        int srcPos3 = (yfilterPos[h + chrHeight]) * srcStride + w;

+        int srcPos4 = (yfilterPos[h + chrHeight]) * srcStride + w + chrWidth;

+        int src1Pos = dstStride * srcHeight + (cfilterPos[h]) * dstChrStride + (w);

+        int src2Pos = dstStride * srcHeight + (dstChrStride*(srcHeight>>1)) + (cfilterPos[h]) * dstChrStride + w;

+

+        int val1 = (local_up_dither[w & 7] << 12); //y offset is 0;

+        int val2 = (local_up_dither[(w + chrWidth) & 7] << 12);

+        int val3 = (local_down_dither[w &7] << 12);

+        int val4 = (local_down_dither[(w + chrWidth) & 7] << 12);

+        int val5 = (local_up_dither[w & 7] << 12);

+        int val6 = (local_up_dither[(w + 3) & 7] << 12);   // 3 is offset of the chrome channel.

+

+

+        int j;

+        int filterPos1 = h * yfilterSize;

+        int filterPos2 = ( h + chrHeight ) * yfilterSize;

+        for(j = 0; j < yfilterSize; j++)

+        {

+            val1 += src[srcPos1] * yfilter[filterPos1 + j];

+            srcPos1 += srcStride;

+            val2 += src[srcPos2] * yfilter[filterPos1 + j];

+            srcPos2 += srcStride;

+            val3 += src[srcPos3] * yfilter[filterPos2 + j];

+            srcPos3 += srcStride;

+            val4 += src[srcPos4] * yfilter[filterPos2 + j];

+            srcPos4 += srcStride;

+            val5 += src[src1Pos] * cfilter[filterPos1 + j];

+            val6 += src[src2Pos] * cfilter[filterPos1 + j];

+            src1Pos += dstChrStride;

+            src2Pos += dstChrStride;

+        }

+        dst[h * dstStride + w] = (((val1 >> 19)&(~0xFF)) ? ((-(val1 >> 19)) >> 31) : (val1 >> 19));

+        dst[h * dstStride + w + chrWidth] = (((val2 >> 19)&(~0xFF)) ? ((-(val2 >> 19)) >> 31) : (val2 >> 19));

+        dst[(h + chrHeight) * dstStride + w] = (((val3 >> 19)&(~0xFF)) ? ((-(val3 >> 19)) >> 31) : (val3 >> 19));

+        dst[(h + chrHeight) * dstStride + w + chrWidth] = (((val4 >> 19)&(~0xFF)) ? ((-(val4 >> 19)) >> 31) : (val4 >> 19));;

+

+        int dst1Pos = dstStride * dstHeight + h * (dstChrStride) + (w);

+        int dst2Pos = (dstChrStride * chrHeight) + dst1Pos;

+        dst[dst1Pos] = (((val5 >> 19)&(~0xFF)) ? ((-(val5 >> 19)) >> 31) : (val5 >> 19));

+        dst[dst2Pos] = (((val6 >> 19)&(~0xFF)) ? ((-(val6 >> 19)) >> 31) : (val6 >> 19));

+    }

+    );

+

+char *kernel_src_vscalefast = KERNEL (

+

+    kernel void vscale_fast_opencl (

+        global unsigned char *dst,

+        const global short *src,

+        const global int *yfilterPos,

+        const global int *cfilterPos,

+        int dstWidth,

+        int dstHeight,

+        int srcWidth,

+        int srcHeight,

+        int dstStride,

+        int dstChrStride,

+        int srcStride,

+        int srcChrStride)

+    {

+        const unsigned char hb_sws_pb_64[8] = {

+            64, 64, 64, 64, 64, 64, 64, 64

+        };

+

+        int w = get_global_id(0);

+        int h = get_global_id(1);

+

+        int chrWidth = get_global_size(0);

+        int chrHeight = get_global_size(1);

+

+        const unsigned char *local_up_dither;

+        const unsigned char *local_down_dither;

+

+        local_up_dither = hb_sws_pb_64;

+        local_down_dither = hb_sws_pb_64;

+

+

+        int rightpart = w + chrWidth;

+        int bh = h + chrHeight; // bottom part

+        short val1 = (src[(yfilterPos[h]) * dstStride + w] + local_up_dither[(w + 0) & 7]) >> 7; //lum offset is 0;

+        short val2 = (src[(yfilterPos[h]) * dstStride + rightpart] + local_up_dither[rightpart & 7]) >> 7;

+        short val3 = (src[(yfilterPos[bh]) * dstStride + w] + local_down_dither[w & 7]) >> 7;

+        short val4 = (src[(yfilterPos[bh]) * dstStride + rightpart] + local_down_dither[rightpart & 7]) >> 7;

+        dst[h * dstStride + w] = ((val1&(~0xFF)) ? ((-val1) >> 31) : (val1));

+        dst[h * dstStride + rightpart] = ((val2&(~0xFF)) ? ((-val2) >> 31) : (val2));

+        dst[bh * dstStride + w] = ((val3&(~0xFF)) ? ((-val3) >> 31) : (val3));

+        dst[bh * dstStride + rightpart] = ((val4&(~0xFF)) ? ((-val4) >> 31) : (val4));

+

+        src += dstStride * srcHeight;

+        dst += dstStride * dstHeight;

+        val1 = (src[cfilterPos[h] * (dstChrStride) + w] + local_up_dither[ w & 7]) >> 7;

+        dst[h * (dstChrStride) + w] = ((val1&(~0xFF)) ? ((-val1) >> 31) : (val1));

+

+        src += dstChrStride * (srcHeight >> 1);

+        dst += dstChrStride * chrHeight;

+        val1 = (src[cfilterPos[h] * dstChrStride + w] + local_up_dither[ (w + 3) & 7] ) >> 7;

+        dst[h * dstChrStride + w] = ((val1&(~0xFF)) ? ((-val1) >> 31) : (val1));

+

+    }
+    );
+
+char *kernel_src_scale = KERNEL (
+
+__kernel __attribute__((reqd_work_group_size(64, 1, 1))) void frame_scale(__global uchar *dst,
+                          __global const uchar *src,
+                          const float xscale,
+                          const float yscale,
+                          const int srcPlaneOffset0,
+                          const int srcPlaneOffset1,
+                          const int srcPlaneOffset2,
+                          const int dstPlaneOffset0,
+                          const int dstPlaneOffset1,
+                          const int dstPlaneOffset2,
+                          const int srcRowWords0,
+                          const int srcRowWords1,
+                          const int srcRowWords2,
+                          const int dstRowWords0,
+                          const int dstRowWords1,
+                          const int dstRowWords2,
+                          const int srcWidth,
+                          const int srcHeight,
+                          const int dstWidth,
+                          const int dstHeight,
+                          __global const float4* restrict xweights,
+                          __global const float4* restrict yweights
+                          )
+{
+    const int x = get_global_id(0);
+    const int y = get_global_id(1);
+    const int z = get_global_id(2);
+
+    // Abort work items outside the dst image bounds.
+
+    if ((get_group_id(0) * 64 >= (dstWidth >> ((z == 0) ? 0 : 1))) || (get_group_id(1) * 16 >= (dstHeight >> ((z == 0) ? 0 : 1))))
+      return;
+
+    const int srcPlaneOffset = (z == 0) ? srcPlaneOffset0 : ((z == 1) ? srcPlaneOffset1 : srcPlaneOffset2);
+    const int dstPlaneOffset = (z == 0) ? dstPlaneOffset0 : ((z == 1) ? dstPlaneOffset1 : dstPlaneOffset2);
+    const int srcRowWords = (z == 0) ? srcRowWords0: ((z == 1) ? srcRowWords1 : srcRowWords2);
+    const int dstRowWords = (z == 0) ? dstRowWords0: ((z == 1) ? dstRowWords1 : dstRowWords2);
+
+    __local uchar pixels[64 * 36];
+    const int localRowPixels = 64;
+    const int groupHeight = 16;     // src pixel height output by the workgroup
+    const int ypad = 2;
+    const int localx = get_local_id(0);
+
+    const int globalStartRow = floor((get_group_id(1) * groupHeight) / yscale);
+    const int globalRowCount = ceil(groupHeight / yscale) + 2 * ypad;
+
+    float4 weights = xweights[x];
+    int4 woffs = floor(x / xscale);
+    woffs += (int4)(-1, 0, 1, 2);
+    woffs = clamp(woffs, 0, (srcWidth >> ((z == 0) ? 0 : 1)) - 1);
+    const int maxy = (srcHeight >> ((z == 0) ? 0 : 1)) - 1;
+
+    // Scale x from global into LDS
+
+    for (int i = 0; i <= globalRowCount; ++i) {
+        int4 offs = srcPlaneOffset + clamp(globalStartRow - ypad + i, 0, maxy) * srcRowWords;
+        offs += woffs;
+        pixels[localx + i * localRowPixels] = convert_uchar(clamp(round(dot(weights,
+                       (float4)(src[offs.x], src[offs.y], src[offs.z], src[offs.w]))), 0.0f, 255.0f));
+    }
+ 
+    barrier(CLK_LOCAL_MEM_FENCE);
+
+    // Scale y from LDS into global
+
+    if (x >= dstWidth >> ((z == 0) ? 0 : 1))
+        return;
+
+    int off = dstPlaneOffset + x + (get_group_id(1) * groupHeight) * dstRowWords;
+
+    for (int i = 0; i < groupHeight; ++i) {
+        if (y >= dstHeight >> ((z == 0) ? 0 : 1))
+          break;
+        int localy = floor((get_group_id(1) * groupHeight + i) / yscale);
+        localy = localy - globalStartRow + ypad;
+        int loff = localx + localy * localRowPixels;
+        dst[off] = convert_uchar(clamp(round(dot(yweights[get_group_id(1) * groupHeight + i],
+                                (float4)(pixels[loff - localRowPixels], pixels[loff], pixels[loff + localRowPixels]
+                                , pixels[loff + localRowPixels * 2]))), 0.0f, 255.0f));
+        off += dstRowWords;
+    }
+}
+);
+
+
+char *kernel_src_yadif_filter = KERNEL(
+    void filter_v6(
+        global unsigned char *dst,
+        global unsigned char *prev, 

+        global unsigned char *cur, 

+        global unsigned char *next,

+        int x,

+        int y,

+        int width,

+        int height,

+        int parity,

+        int inlinesize,

+        int outlinesize,

+        int inmode,

+        int uvflag

+    )

+    {

+

+        int flag = uvflag * (y >=height) * height;  

+        int prefs = select(-(inlinesize), inlinesize,((y+1) - flag) <height);

+        int mrefs = select(inlinesize, -(inlinesize),y - flag);

+        int mode  = select(inmode,2,(y - flag==1) || (y - flag + 2==height));

+        int score;

+    

+        global unsigned char *prev2 = parity ? prev : cur ;

+        global unsigned char *next2 = parity ? cur  : next;

+        int index = x + y * inlinesize;

+        int outindex = x + y * outlinesize;

+        int c = cur[index + mrefs]; 

+        int d = (prev2[index] + next2[index])>>1; 

+        int e = cur[index + prefs]; 

+        int temporal_diff0 = abs((prev2[index]) - (next2[index])); 

+        int temporal_diff1 =(abs(prev[index + mrefs] - c) + abs(prev[index + prefs] - e) )>>1; 

+        int temporal_diff2 =(abs(next[index + mrefs] - c) + abs(next[index + prefs] - e) )>>1; 

+        int diff = max(max(temporal_diff0>>1, temporal_diff1), temporal_diff2); 

+        int spatial_pred = (c+e)>>1; 

+        int spatial_score = abs(cur[index + mrefs-1] - cur[index + prefs-1]) + abs(c-e) + abs(cur[index + mrefs+1] - cur[index + prefs+1]) - 1; 

+        //check -1

+        score = abs(cur[index + mrefs-2] - cur[index + prefs])

+            + abs(cur[index + mrefs-1] - cur[index + prefs+1])

+            + abs(cur[index + mrefs] - cur[index + prefs+2]);

+        if (score < spatial_score)

+        {

+            spatial_score= score;

+            spatial_pred= (cur[index + mrefs-1] + cur[index + prefs+1])>>1;

+        }

+        //check -2

+        score = abs(cur[index + mrefs-3] - cur[index + prefs+1])

+            + abs(cur[index + mrefs-2] - cur[index + prefs+2])

+            + abs(cur[index + mrefs-1] - cur[index + prefs+3]);

+        if (score < spatial_score)

+        {

+            spatial_score= score;

+            spatial_pred= (cur[index + mrefs-2] + cur[index + prefs+2])>>1;

+        }

+        //check 1

+        score = abs(cur[index + mrefs] - cur[index + prefs-2])

+            + abs(cur[index + mrefs+1] - cur[index + prefs-1])

+            + abs(cur[index + mrefs+2] - cur[index + prefs]);

+        if (score < spatial_score)

+        {

+            spatial_score= score;

+            spatial_pred= (cur[index + mrefs+1] + cur[index + prefs-1])>>1;

+        }

+        //check 2

+        score = abs(cur[index + mrefs+1] - cur[index + prefs-3])

+            + abs(cur[index + mrefs+2] - cur[index + prefs-2])

+            + abs(cur[index + mrefs+3] - cur[index + prefs-1]);

+        if (score < spatial_score)

+        {

+            spatial_score= score;

+            spatial_pred= (cur[index + mrefs+2] + cur[index + prefs-2])>>1;

+        }

+        if (mode < 2)

+        { 

+            int b = (prev2[index + (mrefs<<1)] + next2[index + (mrefs<<1)])>>1; 

+            int f = (prev2[index + (prefs<<1)] + next2[index + (prefs<<1)])>>1; 

+            int diffmax = max(max(d-e, d-c), min(b-c, f-e)); 

+            int diffmin = min(min(d-e, d-c), max(b-c, f-e)); 

+

+            diff = max(max(diff, diffmin), -diffmax); 

+        } 

+        if (spatial_pred > d + diff) 

+        {

+            spatial_pred = d + diff; 

+        }

+        else if (spatial_pred < d - diff) 

+        {

+            spatial_pred = d - diff; 

+        }

+

+        dst[outindex] = spatial_pred; 

+    }

+

+    kernel void yadif_filter(

+        global unsigned char *dst,

+        global unsigned char *prev,

+        global unsigned char *cur,

+        global unsigned char *next,

+        int parity,

+        int inlinesizeY,

+        int inlinesizeUV,

+        int outlinesizeY,

+        int outlinesizeUV,

+        int mode)

+    {

+        int x=get_global_id(0);

+        int y=(get_global_id(1)<<1) + (!parity);

+        int width=(get_global_size(0)<<1)/3;

+        int height=get_global_size(1)<<1;

+    

+

+        global unsigned char *dst_Y=dst;

+        global unsigned char *dst_U=dst_Y+height*outlinesizeY;

+

+        global unsigned char *prev_Y=prev;

+        global unsigned char *prev_U=prev_Y+height*inlinesizeY;

+

+        global unsigned char *cur_Y=cur;

+        global unsigned char *cur_U=cur_Y+height*inlinesizeY;

+

+        global unsigned char *next_Y=next;

+        global unsigned char *next_U=next_Y+height*inlinesizeY;

+

+        if(x < width)

+        {

+            filter_v6(dst_Y,prev_Y,cur_Y,next_Y,x,y,width,height,parity,inlinesizeY,outlinesizeY,mode,0);

+        }

+        else

+        {

+            x = x - width;

+            filter_v6(dst_U,prev_U,cur_U,next_U,x,y,width>>1,height>>1,parity,inlinesizeUV,outlinesizeUV,mode,1);

+        }

+    }

+    );

+

+#endif