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