lavfi: modify names of deshake_kernel.h and unsharp_kernel.h
authorhighgod0401 <highgod0401@gmail.com>
Wed, 8 May 2013 00:53:09 +0000 (08:53 +0800)
committerMichael Niedermayer <michaelni@gmx.at>
Fri, 10 May 2013 10:44:04 +0000 (12:44 +0200)
Reviewed-by: Stefano Sabatini <stefasab@gmail.com>
Signed-off-by: Michael Niedermayer <michaelni@gmx.at>
libavfilter/Makefile
libavfilter/deshake_kernel.h [deleted file]
libavfilter/deshake_opencl_kernel.h [new file with mode: 0644]
libavfilter/opencl_allkernels.c
libavfilter/unsharp_kernel.h [deleted file]
libavfilter/unsharp_opencl_kernel.h [new file with mode: 0644]

index 913f502..73a24dd 100644 (file)
@@ -232,7 +232,7 @@ OBJS-$(CONFIG_AMOVIE_FILTER)                 += src_movie.o
 OBJS-$(CONFIG_MOVIE_FILTER)                  += src_movie.o
 
 SKIPHEADERS-$(CONFIG_LIBVIDSTAB)             += vidstabutils.h
-SKIPHEADERS-$(CONFIG_OPENCL)                 += opencl_internal.h deshake_kernel.h unsharp_kernel.h
+SKIPHEADERS-$(CONFIG_OPENCL)                 += opencl_internal.h deshake_opencl_kernel.h unsharp_opencl_kernel.h
 
 TOOLS     = graph2dot
 TESTPROGS = drawutils filtfmts formats
diff --git a/libavfilter/deshake_kernel.h b/libavfilter/deshake_kernel.h
deleted file mode 100644 (file)
index 1eb06fe..0000000
+++ /dev/null
@@ -1,217 +0,0 @@
-/*
- * Copyright (C) 2013 Wei Gao <weigao@multicorewareinc.com>
- *
- *
- * 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_DESHAKE_KERNEL_H
-#define AVFILTER_DESHAKE_KERNEL_H
-
-#include "libavutil/opencl.h"
-
-const char *ff_kernel_deshake_opencl = AV_OPENCL_KERNEL(
-
-inline unsigned char pixel(global const unsigned char *src, float x, float y,
-                           int w, int h,int stride, unsigned char def)
-{
-    return (x < 0 || y < 0 || x >= w || y >= h) ? def : src[(int)x + (int)y * stride];
-}
-unsigned char interpolate_nearest(float x, float y, global const unsigned char *src,
-                                  int width, int height, int stride, unsigned char def)
-{
-    return pixel(src, (int)(x + 0.5), (int)(y + 0.5), width, height, stride, def);
-}
-
-unsigned char interpolate_bilinear(float x, float y, global const unsigned char *src,
-                                   int width, int height, int stride, unsigned char def)
-{
-    int x_c, x_f, y_c, y_f;
-    int v1, v2, v3, v4;
-
-    if (x < -1 || x > width || y < -1 || y > height) {
-        return def;
-    } else {
-        x_f = (int)x;
-        x_c = x_f + 1;
-
-        y_f = (int)y;
-        y_c = y_f + 1;
-
-        v1 = pixel(src, x_c, y_c, width, height, stride, def);
-        v2 = pixel(src, x_c, y_f, width, height, stride, def);
-        v3 = pixel(src, x_f, y_c, width, height, stride, def);
-        v4 = pixel(src, x_f, y_f, width, height, stride, def);
-
-        return (v1*(x - x_f)*(y - y_f) + v2*((x - x_f)*(y_c - y)) +
-                v3*(x_c - x)*(y - y_f) + v4*((x_c - x)*(y_c - y)));
-    }
-}
-
-unsigned char interpolate_biquadratic(float x, float y, global const unsigned char *src,
-                                      int width, int height, int stride, unsigned char def)
-{
-    int     x_c, x_f, y_c, y_f;
-    unsigned char v1,  v2,  v3,  v4;
-    float   f1,  f2,  f3,  f4;
-
-    if (x < - 1 || x > width || y < -1 || y > height)
-        return def;
-    else {
-        x_f = (int)x;
-        x_c = x_f + 1;
-        y_f = (int)y;
-        y_c = y_f + 1;
-
-        v1 = pixel(src, x_c, y_c, width, height, stride, def);
-        v2 = pixel(src, x_c, y_f, width, height, stride, def);
-        v3 = pixel(src, x_f, y_c, width, height, stride, def);
-        v4 = pixel(src, x_f, y_f, width, height, stride, def);
-
-        f1 = 1 - sqrt((x_c - x) * (y_c - y));
-        f2 = 1 - sqrt((x_c - x) * (y - y_f));
-        f3 = 1 - sqrt((x - x_f) * (y_c - y));
-        f4 = 1 - sqrt((x - x_f) * (y - y_f));
-        return (v1 * f1 + v2 * f2 + v3 * f3 + v4 * f4) / (f1 + f2 + f3 + f4);
-    }
-}
-
-inline const float clipf(float a, float amin, float amax)
-{
-    if      (a < amin) return amin;
-    else if (a > amax) return amax;
-    else               return a;
-}
-
-inline int mirror(int v, int m)
-{
-    while ((unsigned)v > (unsigned)m) {
-        v = -v;
-        if (v < 0)
-            v += 2 * m;
-    }
-    return v;
-}
-
-kernel void avfilter_transform(global  unsigned char *src,
-                               global  unsigned char *dst,
-                               global          float *matrix,
-                               global          float *matrix2,
-                                                 int interpolate,
-                                                 int fillmethod,
-                                                 int src_stride_lu,
-                                                 int dst_stride_lu,
-                                                 int src_stride_ch,
-                                                 int dst_stride_ch,
-                                                 int height,
-                                                 int width,
-                                                 int ch,
-                                                 int cw)
-{
-     int global_id = get_global_id(0);
-
-     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 *tempdst;
-     global unsigned char *tempsrc;
-
-     int x;
-     int y;
-     float x_s;
-     float y_s;
-     int tempsrc_stride;
-     int tempdst_stride;
-     int temp_height;
-     int temp_width;
-     int curpos;
-     unsigned char def = 0;
-     if (global_id < width*height) {
-        y = global_id/width;
-        x = global_id%width;
-        x_s = x * matrix[0] + y * matrix[1] + matrix[2];
-        y_s = x * matrix[3] + y * matrix[4] + matrix[5];
-        tempdst = dst_y;
-        tempsrc = src_y;
-        tempsrc_stride = src_stride_lu;
-        tempdst_stride = dst_stride_lu;
-        temp_height = height;
-        temp_width = width;
-     } else if ((global_id >= width*height)&&(global_id < width*height + ch*cw)) {
-        y = (global_id - width*height)/cw;
-        x = (global_id - width*height)%cw;
-        x_s = x * matrix2[0] + y * matrix2[1] + matrix2[2];
-        y_s = x * matrix2[3] + y * matrix2[4] + matrix2[5];
-        tempdst = dst_u;
-        tempsrc = src_u;
-        tempsrc_stride = src_stride_ch;
-        tempdst_stride = dst_stride_ch;
-        temp_height = ch;
-        temp_width = cw;
-     } else {
-        y = (global_id - width*height - ch*cw)/cw;
-        x = (global_id - width*height - ch*cw)%cw;
-        x_s = x * matrix2[0] + y * matrix2[1] + matrix2[2];
-        y_s = x * matrix2[3] + y * matrix2[4] + matrix2[5];
-        tempdst = dst_v;
-        tempsrc = src_v;
-        tempsrc_stride = src_stride_ch;
-        tempdst_stride = dst_stride_ch;
-        temp_height = ch;
-        temp_width = cw;
-     }
-     curpos = y * tempdst_stride + x;
-     switch (fillmethod) {
-        case 0: //FILL_BLANK
-            def = 0;
-            break;
-        case 1: //FILL_ORIGINAL
-            def = tempsrc[y*tempsrc_stride+x];
-            break;
-        case 2: //FILL_CLAMP
-            y_s = clipf(y_s, 0, temp_height - 1);
-            x_s = clipf(x_s, 0, temp_width - 1);
-            def = tempsrc[(int)y_s * tempsrc_stride + (int)x_s];
-            break;
-        case 3: //FILL_MIRROR
-            y_s = mirror(y_s,temp_height - 1);
-            x_s = mirror(x_s,temp_width - 1);
-            def = tempsrc[(int)y_s * tempsrc_stride + (int)x_s];
-            break;
-    }
-    switch (interpolate) {
-        case 0: //INTERPOLATE_NEAREST
-            tempdst[curpos] = interpolate_nearest(x_s, y_s, tempsrc, temp_width, temp_height, tempsrc_stride, def);
-            break;
-        case 1: //INTERPOLATE_BILINEAR
-            tempdst[curpos] = interpolate_bilinear(x_s, y_s, tempsrc, temp_width, temp_height, tempsrc_stride, def);
-            break;
-        case 2: //INTERPOLATE_BIQUADRATIC
-            tempdst[curpos] = interpolate_biquadratic(x_s, y_s, tempsrc, temp_width, temp_height, tempsrc_stride, def);
-            break;
-        default:
-            return;
-    }
-}
-);
-
-#endif /* AVFILTER_DESHAKE_KERNEL_H */
diff --git a/libavfilter/deshake_opencl_kernel.h b/libavfilter/deshake_opencl_kernel.h
new file mode 100644 (file)
index 0000000..ca0bf83
--- /dev/null
@@ -0,0 +1,217 @@
+/*
+ * Copyright (C) 2013 Wei Gao <weigao@multicorewareinc.com>
+ *
+ *
+ * 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_DESHAKE_OPENCL_KERNEL_H
+#define AVFILTER_DESHAKE_OPENCL_KERNEL_H
+
+#include "libavutil/opencl.h"
+
+const char *ff_kernel_deshake_opencl = AV_OPENCL_KERNEL(
+
+inline unsigned char pixel(global const unsigned char *src, float x, float y,
+                           int w, int h,int stride, unsigned char def)
+{
+    return (x < 0 || y < 0 || x >= w || y >= h) ? def : src[(int)x + (int)y * stride];
+}
+unsigned char interpolate_nearest(float x, float y, global const unsigned char *src,
+                                  int width, int height, int stride, unsigned char def)
+{
+    return pixel(src, (int)(x + 0.5), (int)(y + 0.5), width, height, stride, def);
+}
+
+unsigned char interpolate_bilinear(float x, float y, global const unsigned char *src,
+                                   int width, int height, int stride, unsigned char def)
+{
+    int x_c, x_f, y_c, y_f;
+    int v1, v2, v3, v4;
+
+    if (x < -1 || x > width || y < -1 || y > height) {
+        return def;
+    } else {
+        x_f = (int)x;
+        x_c = x_f + 1;
+
+        y_f = (int)y;
+        y_c = y_f + 1;
+
+        v1 = pixel(src, x_c, y_c, width, height, stride, def);
+        v2 = pixel(src, x_c, y_f, width, height, stride, def);
+        v3 = pixel(src, x_f, y_c, width, height, stride, def);
+        v4 = pixel(src, x_f, y_f, width, height, stride, def);
+
+        return (v1*(x - x_f)*(y - y_f) + v2*((x - x_f)*(y_c - y)) +
+                v3*(x_c - x)*(y - y_f) + v4*((x_c - x)*(y_c - y)));
+    }
+}
+
+unsigned char interpolate_biquadratic(float x, float y, global const unsigned char *src,
+                                      int width, int height, int stride, unsigned char def)
+{
+    int     x_c, x_f, y_c, y_f;
+    unsigned char v1,  v2,  v3,  v4;
+    float   f1,  f2,  f3,  f4;
+
+    if (x < - 1 || x > width || y < -1 || y > height)
+        return def;
+    else {
+        x_f = (int)x;
+        x_c = x_f + 1;
+        y_f = (int)y;
+        y_c = y_f + 1;
+
+        v1 = pixel(src, x_c, y_c, width, height, stride, def);
+        v2 = pixel(src, x_c, y_f, width, height, stride, def);
+        v3 = pixel(src, x_f, y_c, width, height, stride, def);
+        v4 = pixel(src, x_f, y_f, width, height, stride, def);
+
+        f1 = 1 - sqrt((x_c - x) * (y_c - y));
+        f2 = 1 - sqrt((x_c - x) * (y - y_f));
+        f3 = 1 - sqrt((x - x_f) * (y_c - y));
+        f4 = 1 - sqrt((x - x_f) * (y - y_f));
+        return (v1 * f1 + v2 * f2 + v3 * f3 + v4 * f4) / (f1 + f2 + f3 + f4);
+    }
+}
+
+inline const float clipf(float a, float amin, float amax)
+{
+    if      (a < amin) return amin;
+    else if (a > amax) return amax;
+    else               return a;
+}
+
+inline int mirror(int v, int m)
+{
+    while ((unsigned)v > (unsigned)m) {
+        v = -v;
+        if (v < 0)
+            v += 2 * m;
+    }
+    return v;
+}
+
+kernel void avfilter_transform(global  unsigned char *src,
+                               global  unsigned char *dst,
+                               global          float *matrix,
+                               global          float *matrix2,
+                                                 int interpolate,
+                                                 int fillmethod,
+                                                 int src_stride_lu,
+                                                 int dst_stride_lu,
+                                                 int src_stride_ch,
+                                                 int dst_stride_ch,
+                                                 int height,
+                                                 int width,
+                                                 int ch,
+                                                 int cw)
+{
+     int global_id = get_global_id(0);
+
+     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 *tempdst;
+     global unsigned char *tempsrc;
+
+     int x;
+     int y;
+     float x_s;
+     float y_s;
+     int tempsrc_stride;
+     int tempdst_stride;
+     int temp_height;
+     int temp_width;
+     int curpos;
+     unsigned char def = 0;
+     if (global_id < width*height) {
+        y = global_id/width;
+        x = global_id%width;
+        x_s = x * matrix[0] + y * matrix[1] + matrix[2];
+        y_s = x * matrix[3] + y * matrix[4] + matrix[5];
+        tempdst = dst_y;
+        tempsrc = src_y;
+        tempsrc_stride = src_stride_lu;
+        tempdst_stride = dst_stride_lu;
+        temp_height = height;
+        temp_width = width;
+     } else if ((global_id >= width*height)&&(global_id < width*height + ch*cw)) {
+        y = (global_id - width*height)/cw;
+        x = (global_id - width*height)%cw;
+        x_s = x * matrix2[0] + y * matrix2[1] + matrix2[2];
+        y_s = x * matrix2[3] + y * matrix2[4] + matrix2[5];
+        tempdst = dst_u;
+        tempsrc = src_u;
+        tempsrc_stride = src_stride_ch;
+        tempdst_stride = dst_stride_ch;
+        temp_height = ch;
+        temp_width = cw;
+     } else {
+        y = (global_id - width*height - ch*cw)/cw;
+        x = (global_id - width*height - ch*cw)%cw;
+        x_s = x * matrix2[0] + y * matrix2[1] + matrix2[2];
+        y_s = x * matrix2[3] + y * matrix2[4] + matrix2[5];
+        tempdst = dst_v;
+        tempsrc = src_v;
+        tempsrc_stride = src_stride_ch;
+        tempdst_stride = dst_stride_ch;
+        temp_height = ch;
+        temp_width = cw;
+     }
+     curpos = y * tempdst_stride + x;
+     switch (fillmethod) {
+        case 0: //FILL_BLANK
+            def = 0;
+            break;
+        case 1: //FILL_ORIGINAL
+            def = tempsrc[y*tempsrc_stride+x];
+            break;
+        case 2: //FILL_CLAMP
+            y_s = clipf(y_s, 0, temp_height - 1);
+            x_s = clipf(x_s, 0, temp_width - 1);
+            def = tempsrc[(int)y_s * tempsrc_stride + (int)x_s];
+            break;
+        case 3: //FILL_MIRROR
+            y_s = mirror(y_s,temp_height - 1);
+            x_s = mirror(x_s,temp_width - 1);
+            def = tempsrc[(int)y_s * tempsrc_stride + (int)x_s];
+            break;
+    }
+    switch (interpolate) {
+        case 0: //INTERPOLATE_NEAREST
+            tempdst[curpos] = interpolate_nearest(x_s, y_s, tempsrc, temp_width, temp_height, tempsrc_stride, def);
+            break;
+        case 1: //INTERPOLATE_BILINEAR
+            tempdst[curpos] = interpolate_bilinear(x_s, y_s, tempsrc, temp_width, temp_height, tempsrc_stride, def);
+            break;
+        case 2: //INTERPOLATE_BIQUADRATIC
+            tempdst[curpos] = interpolate_biquadratic(x_s, y_s, tempsrc, temp_width, temp_height, tempsrc_stride, def);
+            break;
+        default:
+            return;
+    }
+}
+);
+
+#endif /* AVFILTER_DESHAKE_OPENCL_KERNEL_H */
index b4530c0..6d80fa8 100644 (file)
@@ -21,8 +21,8 @@
 #include "opencl_allkernels.h"
 #if CONFIG_OPENCL
 #include "libavutil/opencl.h"
-#include "deshake_kernel.h"
-#include "unsharp_kernel.h"
+#include "deshake_opencl_kernel.h"
+#include "unsharp_opencl_kernel.h"
 #endif
 
 #define OPENCL_REGISTER_KERNEL_CODE(X, x)                                              \
diff --git a/libavfilter/unsharp_kernel.h b/libavfilter/unsharp_kernel.h
deleted file mode 100644 (file)
index dd65bc0..0000000
+++ /dev/null
@@ -1,137 +0,0 @@
-/*
- * Copyright (C) 2013 Wei Gao <weigao@multicorewareinc.com>
- *
- * 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_KERNEL_H
-#define AVFILTER_UNSHARP_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(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_KERNEL_H */
diff --git a/libavfilter/unsharp_opencl_kernel.h b/libavfilter/unsharp_opencl_kernel.h
new file mode 100644 (file)
index 0000000..0cc8e90
--- /dev/null
@@ -0,0 +1,137 @@
+/*
+ * Copyright (C) 2013 Wei Gao <weigao@multicorewareinc.com>
+ *
+ * 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(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 */