1. ffmpeg
  2. Commit 7faa8d8b03832e9f1c67029c859db5af82897a48
Browse code

lavfi: Add OpenCL unsharp mask filter

Intended to replace existing opencl mode of the unsharp filter.
Supports many more pixel formats and works without immediate upload
and download of frame data. The options are compatible with the
existing filter.

Mark Thompson authored on 2017/11/15 04:47:27
Showing 7 changed files
configure
History View file @ 7faa8d8
... ... 
@@ -3283,6 +3283,7 @@ tinterlace_filter_deps="gpl"
3283 3283 
 tinterlace_merge_test_deps="tinterlace_filter"
3284 3284 
 tinterlace_pad_test_deps="tinterlace_filter"
3285 3285 
 tonemap_filter_deps="const_nan"
3286 
+unsharp_opencl_filter_deps="opencl"
3286 3287 
 uspp_filter_deps="gpl avcodec"
3287 3288 
 unsharp_filter_suggest="opencl"
3288 3289 
 vaguedenoiser_filter_deps="gpl"
libavfilter/Makefile
History View file @ 7faa8d8
... ... 
@@ -331,6 +331,8 @@ OBJS-$(CONFIG_TRANSPOSE_FILTER)              += vf_transpose.o
331 331 
 OBJS-$(CONFIG_TRIM_FILTER)                   += trim.o
332 332 
 OBJS-$(CONFIG_UNPREMULTIPLY_FILTER)          += vf_premultiply.o framesync.o
333 333 
 OBJS-$(CONFIG_UNSHARP_FILTER)                += vf_unsharp.o
334 
+OBJS-$(CONFIG_UNSHARP_OPENCL_FILTER)         += vf_unsharp_opencl.o opencl.o \
335 
+                                                opencl/unsharp.o
334 336 
 OBJS-$(CONFIG_USPP_FILTER)                   += vf_uspp.o
335 337 
 OBJS-$(CONFIG_VAGUEDENOISER_FILTER)          += vf_vaguedenoiser.o
336 338 
 OBJS-$(CONFIG_VECTORSCOPE_FILTER)            += vf_vectorscope.o
libavfilter/allfilters.c
History View file @ 7faa8d8
... ... 
@@ -339,6 +339,7 @@ static void register_all(void)
339 339 
     REGISTER_FILTER(TRIM,           trim,           vf);
340 340 
     REGISTER_FILTER(UNPREMULTIPLY,  unpremultiply,  vf);
341 341 
     REGISTER_FILTER(UNSHARP,        unsharp,        vf);
342 
+    REGISTER_FILTER(UNSHARP_OPENCL, unsharp_opencl, vf);
342 343 
     REGISTER_FILTER(USPP,           uspp,           vf);
343 344 
     REGISTER_FILTER(VAGUEDENOISER,  vaguedenoiser,  vf);
344 345 
     REGISTER_FILTER(VECTORSCOPE,    vectorscope,    vf);
libavfilter/opencl/unsharp.cl
History View file @ 7faa8d8
345 346 
new file mode 100644
... ... 
@@ -0,0 +1,99 @@
0 
+/*
1 
+ * This file is part of FFmpeg.
2 
+ *
3 
+ * FFmpeg is free software; you can redistribute it and/or
4 
+ * modify it under the terms of the GNU Lesser General Public
5 
+ * License as published by the Free Software Foundation; either
6 
+ * version 2.1 of the License, or (at your option) any later version.
7 
+ *
8 
+ * FFmpeg is distributed in the hope that it will be useful,
9 
+ * but WITHOUT ANY WARRANTY; without even the implied warranty of
10 
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU
11 
+ * Lesser General Public License for more details.
12 
+ *
13 
+ * You should have received a copy of the GNU Lesser General Public
14 
+ * License along with FFmpeg; if not, write to the Free Software
15 
+ * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA
16 
+ */
17 
+
18 
+__kernel void unsharp_global(__write_only image2d_t dst,
19 
+                             __read_only  image2d_t src,
20 
+                             int size_x,
21 
+                             int size_y,
22 
+                             float amount,
23 
+                             __constant float *coef_matrix)
24 
+{
25 
+    const sampler_t sampler = (CLK_NORMALIZED_COORDS_FALSE |
26 
+                               CLK_FILTER_NEAREST);
27 
+    int2 loc    = (int2)(get_global_id(0), get_global_id(1));
28 
+    int2 centre = (int2)(size_x / 2, size_y / 2);
29 
+
30 
+    float4 val = read_imagef(src, sampler, loc);
31 
+    float4 sum = 0.0f;
32 
+    int x, y;
33 
+
34 
+    for (y = 0; y < size_y; y++) {
35 
+        for (x = 0; x < size_x; x++) {
36 
+            int2 pos = loc + (int2)(x, y) - centre;
37 
+            sum += coef_matrix[y * size_x + x] *
38 
+                read_imagef(src, sampler, pos);
39 
+        }
40 
+    }
41 
+
42 
+    write_imagef(dst, loc, val + (val - sum) * amount);
43 
+}
44 
+
45 
+__kernel void unsharp_local(__write_only image2d_t dst,
46 
+                            __read_only  image2d_t src,
47 
+                            int size_x,
48 
+                            int size_y,
49 
+                            float amount,
50 
+                            __constant float *coef_x,
51 
+                            __constant float *coef_y)
52 
+{
53 
+    const sampler_t sampler = (CLK_NORMALIZED_COORDS_FALSE |
54 
+                               CLK_ADDRESS_CLAMP_TO_EDGE |
55 
+                               CLK_FILTER_NEAREST);
56 
+    int2 block = (int2)(get_group_id(0), get_group_id(1)) * 16;
57 
+    int2 pos   = (int2)(get_local_id(0), get_local_id(1));
58 
+
59 
+    __local float4 tmp[32][32];
60 
+
61 
+    int rad_x = size_x / 2;
62 
+    int rad_y = size_y / 2;
63 
+    int x, y;
64 
+
65 
+    for (y = 0; y <= 1; y++) {
66 
+        for (x = 0; x <= 1; x++) {
67 
+            tmp[pos.y + 16 * y][pos.x + 16 * x] =
68 
+                read_imagef(src, sampler, block + pos + (int2)(16 * x - 8, 16 * y - 8));
69 
+        }
70 
+    }
71 
+
72 
+    barrier(CLK_LOCAL_MEM_FENCE);
73 
+
74 
+    float4 val = tmp[pos.y + 8][pos.x + 8];
75 
+
76 
+    float4 horiz[2];
77 
+    for (y = 0; y <= 1; y++) {
78 
+        horiz[y] = 0.0f;
79 
+        for (x = 0; x < size_x; x++)
80 
+            horiz[y] += coef_x[x] * tmp[pos.y + y * 16][pos.x + 8 + x - rad_x];
81 
+    }
82 
+
83 
+    barrier(CLK_LOCAL_MEM_FENCE);
84 
+
85 
+    for (y = 0; y <= 1; y++) {
86 
+        tmp[pos.y + y * 16][pos.x + 8] = horiz[y];
87 
+    }
88 
+
89 
+    barrier(CLK_LOCAL_MEM_FENCE);
90 
+
91 
+    float4 sum = 0.0f;
92 
+    for (y = 0; y < size_y; y++)
93 
+        sum += coef_y[y] * tmp[pos.y + 8 + y - rad_y][pos.x + 8];
94 
+
95 
+    if (block.x + pos.x < get_image_width(dst) &&
96 
+        block.y + pos.y < get_image_height(dst))
97 
+        write_imagef(dst, block + pos, val + (val - sum) * amount);
98 
+}
libavfilter/opencl_source.h
History View file @ 7faa8d8
... ... 
@@ -20,5 +20,6 @@
20 20 
 #define AVFILTER_OPENCL_SOURCE_H
21 21
22 22 
 extern const char *ff_opencl_source_overlay;
23 
+extern const char *ff_opencl_source_unsharp;
23 24
24 25 
 #endif /* AVFILTER_OPENCL_SOURCE_H */
libavfilter/version.h
History View file @ 7faa8d8
... ... 
@@ -31,7 +31,7 @@
31 31
32 32 
 #define LIBAVFILTER_VERSION_MAJOR   7
33 33 
 #define LIBAVFILTER_VERSION_MINOR   2
34 
-#define LIBAVFILTER_VERSION_MICRO 101
34 
+#define LIBAVFILTER_VERSION_MICRO 102
35 35
36 36 
 #define LIBAVFILTER_VERSION_INT AV_VERSION_INT(LIBAVFILTER_VERSION_MAJOR, \
37 37 
                                                LIBAVFILTER_VERSION_MINOR, \
libavfilter/vf_unsharp_opencl.c
History View file @ 7faa8d8
38 38 
new file mode 100644
... ... 
@@ -0,0 +1,482 @@
0 
+/*
1 
+ * This file is part of FFmpeg.
2 
+ *
3 
+ * FFmpeg is free software; you can redistribute it and/or
4 
+ * modify it under the terms of the GNU Lesser General Public
5 
+ * License as published by the Free Software Foundation; either
6 
+ * version 2.1 of the License, or (at your option) any later version.
7 
+ *
8 
+ * FFmpeg is distributed in the hope that it will be useful,
9 
+ * but WITHOUT ANY WARRANTY; without even the implied warranty of
10 
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU
11 
+ * Lesser General Public License for more details.
12 
+ *
13 
+ * You should have received a copy of the GNU Lesser General Public
14 
+ * License along with FFmpeg; if not, write to the Free Software
15 
+ * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA
16 
+ */
17 
+
18 
+#include "libavutil/common.h"
19 
+#include "libavutil/imgutils.h"
20 
+#include "libavutil/mem.h"
21 
+#include "libavutil/opt.h"
22 
+#include "libavutil/pixdesc.h"
23 
+
24 
+#include "avfilter.h"
25 
+#include "internal.h"
26 
+#include "opencl.h"
27 
+#include "opencl_source.h"
28 
+#include "video.h"
29 
+
30 
+#define MAX_DIAMETER 23
31 
+
32 
+typedef struct UnsharpOpenCLContext {
33 
+    OpenCLFilterContext ocf;
34 
+
35 
+    int              initialised;
36 
+    cl_kernel        kernel;
37 
+    cl_command_queue command_queue;
38 
+
39 
+    float luma_size_x;
40 
+    float luma_size_y;
41 
+    float luma_amount;
42 
+    float chroma_size_x;
43 
+    float chroma_size_y;
44 
+    float chroma_amount;
45 
+
46 
+    int global;
47 
+
48 
+    int nb_planes;
49 
+    struct {
50 
+        float blur_x[MAX_DIAMETER];
51 
+        float blur_y[MAX_DIAMETER];
52 
+
53 
+        cl_mem   matrix;
54 
+        cl_mem   coef_x;
55 
+        cl_mem   coef_y;
56 
+
57 
+        cl_int   size_x;
58 
+        cl_int   size_y;
59 
+        cl_float amount;
60 
+        cl_float threshold;
61 
+    } plane[4];
62 
+} UnsharpOpenCLContext;
63 
+
64 
+
65 
+static int unsharp_opencl_init(AVFilterContext *avctx)
66 
+{
67 
+    UnsharpOpenCLContext *ctx = avctx->priv;
68 
+    cl_int cle;
69 
+    int err;
70 
+
71 
+    err = ff_opencl_filter_load_program(avctx, &ff_opencl_source_unsharp, 1);
72 
+    if (err < 0)
73 
+        goto fail;
74 
+
75 
+    ctx->command_queue = clCreateCommandQueue(ctx->ocf.hwctx->context,
76 
+                                              ctx->ocf.hwctx->device_id,
77 
+                                              0, &cle);
78 
+    if (!ctx->command_queue) {
79 
+        av_log(avctx, AV_LOG_ERROR, "Failed to create OpenCL "
80 
+               "command queue: %d.\n", cle);
81 
+        err = AVERROR(EIO);
82 
+        goto fail;
83 
+    }
84 
+
85 
+    // Use global kernel if mask size will be too big for the local store..
86 
+    ctx->global = (ctx->luma_size_x   > 17.0f ||
87 
+                   ctx->luma_size_y   > 17.0f ||
88 
+                   ctx->chroma_size_x > 17.0f ||
89 
+                   ctx->chroma_size_y > 17.0f);
90 
+
91 
+    ctx->kernel = clCreateKernel(ctx->ocf.program,
92 
+                                 ctx->global ? "unsharp_global"
93 
+                                             : "unsharp_local", &cle);
94 
+    if (!ctx->kernel) {
95 
+        av_log(avctx, AV_LOG_ERROR, "Failed to create kernel: %d.\n", cle);
96 
+        err = AVERROR(EIO);
97 
+        goto fail;
98 
+    }
99 
+
100 
+    ctx->initialised = 1;
101 
+    return 0;
102 
+
103 
+fail:
104 
+    if (ctx->command_queue)
105 
+        clReleaseCommandQueue(ctx->command_queue);
106 
+    if (ctx->kernel)
107 
+        clReleaseKernel(ctx->kernel);
108 
+    return err;
109 
+}
110 
+
111 
+static int unsharp_opencl_make_filter_params(AVFilterContext *avctx)
112 
+{
113 
+    UnsharpOpenCLContext *ctx = avctx->priv;
114 
+    const AVPixFmtDescriptor *desc;
115 
+    float *matrix;
116 
+    double val, sum;
117 
+    cl_int cle;
118 
+    cl_mem buffer;
119 
+    size_t matrix_bytes;
120 
+    float diam_x, diam_y, amount;
121 
+    int err, p, x, y, size_x, size_y;
122 
+
123 
+    desc = av_pix_fmt_desc_get(ctx->ocf.output_format);
124 
+
125 
+    ctx->nb_planes = 0;
126 
+    for (p = 0; p < desc->nb_components; p++)
127 
+        ctx->nb_planes = FFMAX(ctx->nb_planes, desc->comp[p].plane + 1);
128 
+
129 
+    for (p = 0; p < ctx->nb_planes; p++) {
130 
+        if (p == 0 || (desc->flags & AV_PIX_FMT_FLAG_RGB)) {
131 
+            diam_x = ctx->luma_size_x;
132 
+            diam_y = ctx->luma_size_y;
133 
+            amount = ctx->luma_amount;
134 
+        } else {
135 
+            diam_x = ctx->chroma_size_x;
136 
+            diam_y = ctx->chroma_size_y;
137 
+            amount = ctx->chroma_amount;
138 
+        }
139 
+        size_x = (int)ceil(diam_x) | 1;
140 
+        size_y = (int)ceil(diam_y) | 1;
141 
+        matrix_bytes = size_x * size_y * sizeof(float);
142 
+
143 
+        matrix = av_malloc(matrix_bytes);
144 
+        if (!matrix) {
145 
+            err = AVERROR(ENOMEM);
146 
+            goto fail;
147 
+        }
148 
+
149 
+        sum = 0.0;
150 
+        for (x = 0; x < size_x; x++) {
151 
+            double dx = (double)(x - size_x / 2) / diam_x;
152 
+            sum += ctx->plane[p].blur_x[x] = exp(-16.0 * (dx * dx));
153 
+        }
154 
+        for (x = 0; x < size_x; x++)
155 
+            ctx->plane[p].blur_x[x] /= sum;
156 
+
157 
+        sum = 0.0;
158 
+        for (y = 0; y < size_y; y++) {
159 
+            double dy = (double)(y - size_y / 2) / diam_y;
160 
+            sum += ctx->plane[p].blur_y[y] = exp(-16.0 * (dy * dy));
161 
+        }
162 
+        for (y = 0; y < size_y; y++)
163 
+            ctx->plane[p].blur_y[y] /= sum;
164 
+
165 
+        for (y = 0; y < size_y; y++) {
166 
+            for (x = 0; x < size_x; x++) {
167 
+                val = ctx->plane[p].blur_x[x] * ctx->plane[p].blur_y[y];
168 
+                matrix[y * size_x + x] = val;
169 
+            }
170 
+        }
171 
+
172 
+        if (ctx->global) {
173 
+            buffer = clCreateBuffer(ctx->ocf.hwctx->context,
174 
+                                    CL_MEM_READ_ONLY     |
175 
+                                    CL_MEM_COPY_HOST_PTR |
176 
+                                    CL_MEM_HOST_NO_ACCESS,
177 
+                                    matrix_bytes, matrix, &cle);
178 
+            if (!buffer) {
179 
+                av_log(avctx, AV_LOG_ERROR, "Failed to create matrix buffer: "
180 
+                       "%d.\n", cle);
181 
+                err = AVERROR(EIO);
182 
+                goto fail;
183 
+            }
184 
+            ctx->plane[p].matrix = buffer;
185 
+        } else {
186 
+            buffer = clCreateBuffer(ctx->ocf.hwctx->context,
187 
+                                    CL_MEM_READ_ONLY     |
188 
+                                    CL_MEM_COPY_HOST_PTR |
189 
+                                    CL_MEM_HOST_NO_ACCESS,
190 
+                                    sizeof(ctx->plane[p].blur_x),
191 
+                                    ctx->plane[p].blur_x, &cle);
192 
+            if (!buffer) {
193 
+                av_log(avctx, AV_LOG_ERROR, "Failed to create x-coef buffer: "
194 
+                       "%d.\n", cle);
195 
+                err = AVERROR(EIO);
196 
+                goto fail;
197 
+            }
198 
+            ctx->plane[p].coef_x = buffer;
199 
+
200 
+            buffer = clCreateBuffer(ctx->ocf.hwctx->context,
201 
+                                    CL_MEM_READ_ONLY     |
202 
+                                    CL_MEM_COPY_HOST_PTR |
203 
+                                    CL_MEM_HOST_NO_ACCESS,
204 
+                                    sizeof(ctx->plane[p].blur_y),
205 
+                                    ctx->plane[p].blur_y, &cle);
206 
+            if (!buffer) {
207 
+                av_log(avctx, AV_LOG_ERROR, "Failed to create y-coef buffer: "
208 
+                       "%d.\n", cle);
209 
+                err = AVERROR(EIO);
210 
+                goto fail;
211 
+            }
212 
+            ctx->plane[p].coef_y = buffer;
213 
+        }
214 
+
215 
+        av_freep(&matrix);
216 
+
217 
+        ctx->plane[p].size_x = size_x;
218 
+        ctx->plane[p].size_y = size_y;
219 
+        ctx->plane[p].amount = amount;
220 
+    }
221 
+
222 
+    err = 0;
223 
+fail:
224 
+    av_freep(&matrix);
225 
+    return err;
226 
+}
227 
+
228 
+static int unsharp_opencl_filter_frame(AVFilterLink *inlink, AVFrame *input)
229 
+{
230 
+    AVFilterContext    *avctx = inlink->dst;
231 
+    AVFilterLink     *outlink = avctx->outputs[0];
232 
+    UnsharpOpenCLContext *ctx = avctx->priv;
233 
+    AVFrame *output = NULL;
234 
+    cl_int cle;
235 
+    size_t global_work[2];
236 
+    size_t local_work[2];
237 
+    cl_mem src, dst;
238 
+    int err, p;
239 
+
240 
+    av_log(ctx, AV_LOG_DEBUG, "Filter input: %s, %ux%u (%"PRId64").\n",
241 
+           av_get_pix_fmt_name(input->format),
242 
+           input->width, input->height, input->pts);
243 
+
244 
+    if (!input->hw_frames_ctx)
245 
+        return AVERROR(EINVAL);
246 
+
247 
+    if (!ctx->initialised) {
248 
+        err = unsharp_opencl_init(avctx);
249 
+        if (err < 0)
250 
+            goto fail;
251 
+
252 
+        err = unsharp_opencl_make_filter_params(avctx);
253 
+        if (err < 0)
254 
+            goto fail;
255 
+    }
256 
+
257 
+    output = ff_get_video_buffer(outlink, outlink->w, outlink->h);
258 
+    if (!output) {
259 
+        err = AVERROR(ENOMEM);
260 
+        goto fail;
261 
+    }
262 
+
263 
+    for (p = 0; p < FF_ARRAY_ELEMS(output->data); p++) {
264 
+        src = (cl_mem) input->data[p];
265 
+        dst = (cl_mem)output->data[p];
266 
+
267 
+        if (!dst)
268 
+            break;
269 
+
270 
+        cle = clSetKernelArg(ctx->kernel, 0, sizeof(cl_mem), &dst);
271 
+        if (cle != CL_SUCCESS) {
272 
+            av_log(avctx, AV_LOG_ERROR, "Failed to set kernel "
273 
+                   "destination image argument: %d.\n", cle);
274 
+            goto fail;
275 
+        }
276 
+        cle = clSetKernelArg(ctx->kernel, 1, sizeof(cl_mem), &src);
277 
+        if (cle != CL_SUCCESS) {
278 
+            av_log(avctx, AV_LOG_ERROR, "Failed to set kernel "
279 
+                   "source image argument: %d.\n", cle);
280 
+            goto fail;
281 
+        }
282 
+        cle = clSetKernelArg(ctx->kernel, 2, sizeof(cl_int), &ctx->plane[p].size_x);
283 
+        if (cle != CL_SUCCESS) {
284 
+            av_log(avctx, AV_LOG_ERROR, "Failed to set kernel "
285 
+                   "matrix size argument: %d.\n", cle);
286 
+            goto fail;
287 
+        }
288 
+        cle = clSetKernelArg(ctx->kernel, 3, sizeof(cl_int), &ctx->plane[p].size_y);
289 
+        if (cle != CL_SUCCESS) {
290 
+            av_log(avctx, AV_LOG_ERROR, "Failed to set kernel "
291 
+                   "matrix size argument: %d.\n", cle);
292 
+            goto fail;
293 
+        }
294 
+        cle = clSetKernelArg(ctx->kernel, 4, sizeof(cl_float), &ctx->plane[p].amount);
295 
+        if (cle != CL_SUCCESS) {
296 
+            av_log(avctx, AV_LOG_ERROR, "Failed to set kernel "
297 
+                   "amount argument: %d.\n", cle);
298 
+            goto fail;
299 
+        }
300 
+        if (ctx->global) {
301 
+            cle = clSetKernelArg(ctx->kernel, 5, sizeof(cl_mem), &ctx->plane[p].matrix);
302 
+            if (cle != CL_SUCCESS) {
303 
+                av_log(avctx, AV_LOG_ERROR, "Failed to set kernel "
304 
+                       "matrix argument: %d.\n", cle);
305 
+                goto fail;
306 
+            }
307 
+        } else {
308 
+            cle = clSetKernelArg(ctx->kernel, 5, sizeof(cl_mem), &ctx->plane[p].coef_x);
309 
+            if (cle != CL_SUCCESS) {
310 
+                av_log(avctx, AV_LOG_ERROR, "Failed to set kernel "
311 
+                       "x-coef argument: %d.\n", cle);
312 
+                goto fail;
313 
+            }
314 
+            cle = clSetKernelArg(ctx->kernel, 6, sizeof(cl_mem), &ctx->plane[p].coef_y);
315 
+            if (cle != CL_SUCCESS) {
316 
+                av_log(avctx, AV_LOG_ERROR, "Failed to set kernel "
317 
+                       "y-coef argument: %d.\n", cle);
318 
+                goto fail;
319 
+            }
320 
+        }
321 
+
322 
+        if (ctx->global) {
323 
+            global_work[0] = output->width;
324 
+            global_work[1] = output->height;
325 
+        } else {
326 
+            global_work[0] = FFALIGN(output->width,  16);
327 
+            global_work[1] = FFALIGN(output->height, 16);
328 
+            local_work[0]  = 16;
329 
+            local_work[1]  = 16;
330 
+        }
331 
+
332 
+        av_log(avctx, AV_LOG_DEBUG, "Run kernel on plane %d "
333 
+               "(%zux%zu).\n", p, global_work[0], global_work[1]);
334 
+
335 
+        cle = clEnqueueNDRangeKernel(ctx->command_queue, ctx->kernel, 2, NULL,
336 
+                                     global_work, ctx->global ? NULL : local_work,
337 
+                                     0, NULL, NULL);
338 
+        if (cle != CL_SUCCESS) {
339 
+            av_log(avctx, AV_LOG_ERROR, "Failed to enqueue kernel: %d.\n",
340 
+                   cle);
341 
+            err = AVERROR(EIO);
342 
+            goto fail;
343 
+        }
344 
+    }
345 
+
346 
+    cle = clFinish(ctx->command_queue);
347 
+    if (cle != CL_SUCCESS) {
348 
+        av_log(avctx, AV_LOG_ERROR, "Failed to finish command queue: %d.\n",
349 
+               cle);
350 
+        err = AVERROR(EIO);
351 
+        goto fail;
352 
+    }
353 
+
354 
+    err = av_frame_copy_props(output, input);
355 
+    if (err < 0)
356 
+        goto fail;
357 
+
358 
+    av_frame_free(&input);
359 
+
360 
+    av_log(ctx, AV_LOG_DEBUG, "Filter output: %s, %ux%u (%"PRId64").\n",
361 
+           av_get_pix_fmt_name(output->format),
362 
+           output->width, output->height, output->pts);
363 
+
364 
+    return ff_filter_frame(outlink, output);
365 
+
366 
+fail:
367 
+    clFinish(ctx->command_queue);
368 
+    av_frame_free(&input);
369 
+    av_frame_free(&output);
370 
+    return err;
371 
+}
372 
+
373 
+static av_cold void unsharp_opencl_uninit(AVFilterContext *avctx)
374 
+{
375 
+    UnsharpOpenCLContext *ctx = avctx->priv;
376 
+    cl_int cle;
377 
+    int i;
378 
+
379 
+    for (i = 0; i < ctx->nb_planes; i++) {
380 
+        if (ctx->plane[i].matrix)
381 
+            clReleaseMemObject(ctx->plane[i].matrix);
382 
+        if (ctx->plane[i].coef_x)
383 
+            clReleaseMemObject(ctx->plane[i].coef_x);
384 
+        if (ctx->plane[i].coef_y)
385 
+            clReleaseMemObject(ctx->plane[i].coef_y);
386 
+    }
387 
+
388 
+    if (ctx->kernel) {
389 
+        cle = clReleaseKernel(ctx->kernel);
390 
+        if (cle != CL_SUCCESS)
391 
+            av_log(avctx, AV_LOG_ERROR, "Failed to release "
392 
+                   "kernel: %d.\n", cle);
393 
+    }
394 
+
395 
+    if (ctx->command_queue) {
396 
+        cle = clReleaseCommandQueue(ctx->command_queue);
397 
+        if (cle != CL_SUCCESS)
398 
+            av_log(avctx, AV_LOG_ERROR, "Failed to release "
399 
+                   "command queue: %d.\n", cle);
400 
+    }
401 
+
402 
+    ff_opencl_filter_uninit(avctx);
403 
+}
404 
+
405 
+#define OFFSET(x) offsetof(UnsharpOpenCLContext, x)
406 
+#define FLAGS (AV_OPT_FLAG_FILTERING_PARAM | AV_OPT_FLAG_VIDEO_PARAM)
407 
+static const AVOption unsharp_opencl_options[] = {
408 
+    { "luma_msize_x",     "Set luma mask horizontal diameter (pixels)",
409 
+      OFFSET(luma_size_x),     AV_OPT_TYPE_FLOAT,
410 
+      { .dbl = 5.0 },   1, MAX_DIAMETER, FLAGS },
411 
+    { "lx",               "Set luma mask horizontal diameter (pixels)",
412 
+      OFFSET(luma_size_x),     AV_OPT_TYPE_FLOAT,
413 
+      { .dbl = 5.0 },   1, MAX_DIAMETER, FLAGS },
414 
+    { "luma_msize_y",     "Set luma mask vertical diameter (pixels)",
415 
+      OFFSET(luma_size_y),     AV_OPT_TYPE_FLOAT,
416 
+      { .dbl = 5.0 },   1, MAX_DIAMETER, FLAGS },
417 
+    { "ly",               "Set luma mask vertical diameter (pixels)",
418 
+      OFFSET(luma_size_y),     AV_OPT_TYPE_FLOAT,
419 
+      { .dbl = 5.0 },   1, MAX_DIAMETER, FLAGS },
420 
+    { "luma_amount",      "Set luma amount (multiplier)",
421 
+      OFFSET(luma_amount),     AV_OPT_TYPE_FLOAT,
422 
+      { .dbl = 1.0 }, -10, 10, FLAGS },
423 
+    { "la",               "Set luma amount (multiplier)",
424 
+      OFFSET(luma_amount),     AV_OPT_TYPE_FLOAT,
425 
+      { .dbl = 1.0 }, -10, 10, FLAGS },
426 
+
427 
+    { "chroma_msize_x",   "Set chroma mask horizontal diameter (pixels after subsampling)",
428 
+      OFFSET(chroma_size_x),   AV_OPT_TYPE_FLOAT,
429 
+      { .dbl = 5.0 },   1, MAX_DIAMETER, FLAGS },
430 
+    { "cx",               "Set chroma mask horizontal diameter (pixels after subsampling)",
431 
+      OFFSET(chroma_size_x),   AV_OPT_TYPE_FLOAT,
432 
+      { .dbl = 5.0 },   1, MAX_DIAMETER, FLAGS },
433 
+    { "chroma_msize_y",   "Set chroma mask vertical diameter (pixels after subsampling)",
434 
+      OFFSET(chroma_size_y),   AV_OPT_TYPE_FLOAT,
435 
+      { .dbl = 5.0 },   1, MAX_DIAMETER, FLAGS },
436 
+    { "cy",               "Set chroma mask vertical diameter (pixels after subsampling)",
437 
+      OFFSET(chroma_size_y),   AV_OPT_TYPE_FLOAT,
438 
+      { .dbl = 5.0 },   1, MAX_DIAMETER, FLAGS },
439 
+    { "chroma_amount",    "Set chroma amount (multiplier)",
440 
+      OFFSET(chroma_amount),   AV_OPT_TYPE_FLOAT,
441 
+      { .dbl = 0.0 }, -10, 10, FLAGS },
442 
+    { "ca",               "Set chroma amount (multiplier)",
443 
+      OFFSET(chroma_amount),   AV_OPT_TYPE_FLOAT,
444 
+      { .dbl = 0.0 }, -10, 10, FLAGS },
445 
+
446 
+    { NULL }
447 
+};
448 
+
449 
+AVFILTER_DEFINE_CLASS(unsharp_opencl);
450 
+
451 
+static const AVFilterPad unsharp_opencl_inputs[] = {
452 
+    {
453 
+        .name         = "default",
454 
+        .type         = AVMEDIA_TYPE_VIDEO,
455 
+        .filter_frame = &unsharp_opencl_filter_frame,
456 
+        .config_props = &ff_opencl_filter_config_input,
457 
+    },
458 
+    { NULL }
459 
+};
460 
+
461 
+static const AVFilterPad unsharp_opencl_outputs[] = {
462 
+    {
463 
+        .name         = "default",
464 
+        .type         = AVMEDIA_TYPE_VIDEO,
465 
+        .config_props = &ff_opencl_filter_config_output,
466 
+    },
467 
+    { NULL }
468 
+};
469 
+
470 
+AVFilter ff_vf_unsharp_opencl = {
471 
+    .name           = "unsharp_opencl",
472 
+    .description    = NULL_IF_CONFIG_SMALL("Apply unsharp mask to input video"),
473 
+    .priv_size      = sizeof(UnsharpOpenCLContext),
474 
+    .priv_class     = &unsharp_opencl_class,
475 
+    .init           = &ff_opencl_filter_init,
476 
+    .uninit         = &unsharp_opencl_uninit,
477 
+    .query_formats  = &ff_opencl_filter_query_formats,
478 
+    .inputs         = unsharp_opencl_inputs,
479 
+    .outputs        = unsharp_opencl_outputs,
480 
+    .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE,
481 
+};