vf_avgblur_opencl: Don't run kernel on pixels outside the image
[ffmpeg.git] / libavfilter / vf_avgblur_opencl.c
1 /*
2  * Copyright (c) 2018 Dylan Fernando
3  *
4  * This file is part of FFmpeg.
5  *
6  * FFmpeg is free software; you can redistribute it and/or
7  * modify it under the terms of the GNU Lesser General Public
8  * License as published by the Free Software Foundation; either
9  * version 2.1 of the License, or (at your option) any later version.
10  *
11  * FFmpeg is distributed in the hope that it will be useful,
12  * but WITHOUT ANY WARRANTY; without even the implied warranty of
13  * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU
14  * Lesser General Public License for more details.
15  *
16  * You should have received a copy of the GNU Lesser General Public
17  * License along with FFmpeg; if not, write to the Free Software
18  * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA
19  */
20
21 #include "libavutil/common.h"
22 #include "libavutil/imgutils.h"
23 #include "libavutil/mem.h"
24 #include "libavutil/opt.h"
25 #include "libavutil/pixdesc.h"
26
27 #include "avfilter.h"
28 #include "internal.h"
29 #include "opencl.h"
30 #include "opencl_source.h"
31 #include "video.h"
32
33
34 typedef struct AverageBlurOpenCLContext {
35     OpenCLFilterContext ocf;
36
37     int              initialised;
38     cl_kernel        kernel_horiz;
39     cl_kernel        kernel_vert;
40     cl_command_queue command_queue;
41
42     int radius;
43     int radiusV;
44     int planes;
45
46 } AverageBlurOpenCLContext;
47
48
49 static int avgblur_opencl_init(AVFilterContext *avctx)
50 {
51     AverageBlurOpenCLContext *ctx = avctx->priv;
52     cl_int cle;
53     int err;
54
55     err = ff_opencl_filter_load_program(avctx, &ff_opencl_source_avgblur, 1);
56     if (err < 0)
57         goto fail;
58
59     ctx->command_queue = clCreateCommandQueue(ctx->ocf.hwctx->context,
60                                               ctx->ocf.hwctx->device_id,
61                                               0, &cle);
62     if (!ctx->command_queue) {
63         av_log(avctx, AV_LOG_ERROR, "Failed to create OpenCL "
64                "command queue: %d.\n", cle);
65         err = AVERROR(EIO);
66         goto fail;
67     }
68
69     ctx->kernel_horiz = clCreateKernel(ctx->ocf.program,"avgblur_horiz", &cle);
70     if (!ctx->kernel_horiz) {
71         av_log(avctx, AV_LOG_ERROR, "Failed to create kernel: %d.\n", cle);
72         err = AVERROR(EIO);
73         goto fail;
74     }
75
76     ctx->kernel_vert = clCreateKernel(ctx->ocf.program,"avgblur_vert", &cle);
77     if (!ctx->kernel_vert) {
78         av_log(avctx, AV_LOG_ERROR, "Failed to create kernel: %d.\n", cle);
79         err = AVERROR(EIO);
80         goto fail;
81     }
82
83     if (ctx->radiusV <= 0) {
84         ctx->radiusV = ctx->radius;
85     }
86
87     ctx->initialised = 1;
88     return 0;
89
90 fail:
91     if (ctx->command_queue)
92         clReleaseCommandQueue(ctx->command_queue);
93     if (ctx->kernel_horiz)
94         clReleaseKernel(ctx->kernel_horiz);
95     if (ctx->kernel_vert)
96         clReleaseKernel(ctx->kernel_vert);
97     return err;
98 }
99
100 static int avgblur_opencl_filter_frame(AVFilterLink *inlink, AVFrame *input)
101 {
102     AVFilterContext    *avctx = inlink->dst;
103     AVFilterLink     *outlink = avctx->outputs[0];
104     AverageBlurOpenCLContext *ctx = avctx->priv;
105     AVFrame *output = NULL;
106     AVFrame *intermediate = NULL;
107     cl_int cle;
108     size_t global_work[2];
109     cl_mem src, dst, inter;
110     int err, p, radius_x, radius_y;
111
112     av_log(ctx, AV_LOG_DEBUG, "Filter input: %s, %ux%u (%"PRId64").\n",
113            av_get_pix_fmt_name(input->format),
114            input->width, input->height, input->pts);
115
116     if (!input->hw_frames_ctx)
117         return AVERROR(EINVAL);
118
119     if (!ctx->initialised) {
120         err = avgblur_opencl_init(avctx);
121         if (err < 0)
122             goto fail;
123
124     }
125
126     output = ff_get_video_buffer(outlink, outlink->w, outlink->h);
127     if (!output) {
128         err = AVERROR(ENOMEM);
129         goto fail;
130     }
131
132     intermediate = ff_get_video_buffer(outlink, outlink->w, outlink->h);
133     if (!intermediate) {
134         err = AVERROR(ENOMEM);
135         goto fail;
136     }
137
138     for (p = 0; p < FF_ARRAY_ELEMS(output->data); p++) {
139         src = (cl_mem) input->data[p];
140         dst = (cl_mem)output->data[p];
141         inter = (cl_mem) intermediate->data[p];
142
143         if (!dst)
144             break;
145
146         radius_x = ctx->radius;
147         radius_y = ctx->radiusV;
148
149         if (!(ctx->planes & (1 << p))) {
150             radius_x = 0;
151             radius_y = 0;
152         }
153
154         cle = clSetKernelArg(ctx->kernel_horiz, 0, sizeof(cl_mem), &inter);
155         if (cle != CL_SUCCESS) {
156             av_log(avctx, AV_LOG_ERROR, "Failed to set kernel "
157                    "destination image argument: %d.\n", cle);
158             goto fail;
159         }
160         cle = clSetKernelArg(ctx->kernel_horiz, 1, sizeof(cl_mem), &src);
161         if (cle != CL_SUCCESS) {
162             av_log(avctx, AV_LOG_ERROR, "Failed to set kernel "
163                    "source image argument: %d.\n", cle);
164             goto fail;
165         }
166         cle = clSetKernelArg(ctx->kernel_horiz, 2, sizeof(cl_int), &radius_x);
167         if (cle != CL_SUCCESS) {
168             av_log(avctx, AV_LOG_ERROR, "Failed to set kernel "
169                    "sizeX argument: %d.\n", cle);
170             goto fail;
171         }
172
173         err = ff_opencl_filter_work_size_from_image(avctx, global_work,
174                                                     intermediate, p, 0);
175         if (err < 0)
176             goto fail;
177
178         av_log(avctx, AV_LOG_DEBUG, "Run kernel on plane %d "
179                "(%"SIZE_SPECIFIER"x%"SIZE_SPECIFIER").\n",
180                p, global_work[0], global_work[1]);
181
182         cle = clEnqueueNDRangeKernel(ctx->command_queue, ctx->kernel_horiz, 2, NULL,
183                                      global_work, NULL,
184                                      0, NULL, NULL);
185         if (cle != CL_SUCCESS) {
186             av_log(avctx, AV_LOG_ERROR, "Failed to enqueue kernel: %d.\n",
187                    cle);
188             err = AVERROR(EIO);
189             goto fail;
190         }
191
192         cle = clSetKernelArg(ctx->kernel_vert, 0, sizeof(cl_mem), &dst);
193         if (cle != CL_SUCCESS) {
194             av_log(avctx, AV_LOG_ERROR, "Failed to set kernel "
195                    "destination image argument: %d.\n", cle);
196             goto fail;
197         }
198         cle = clSetKernelArg(ctx->kernel_vert, 1, sizeof(cl_mem), &inter);
199         if (cle != CL_SUCCESS) {
200             av_log(avctx, AV_LOG_ERROR, "Failed to set kernel "
201                    "source image argument: %d.\n", cle);
202             goto fail;
203         }
204         cle = clSetKernelArg(ctx->kernel_vert, 2, sizeof(cl_int), &radius_y);
205         if (cle != CL_SUCCESS) {
206             av_log(avctx, AV_LOG_ERROR, "Failed to set kernel "
207                    "sizeY argument: %d.\n", cle);
208             goto fail;
209         }
210
211         err = ff_opencl_filter_work_size_from_image(avctx, global_work,
212                                                     output, p, 0);
213         if (err < 0)
214             goto fail;
215
216         av_log(avctx, AV_LOG_DEBUG, "Run kernel on plane %d "
217                "(%"SIZE_SPECIFIER"x%"SIZE_SPECIFIER").\n",
218                p, global_work[0], global_work[1]);
219
220         cle = clEnqueueNDRangeKernel(ctx->command_queue, ctx->kernel_vert, 2, NULL,
221                                      global_work, NULL,
222                                      0, NULL, NULL);
223         if (cle != CL_SUCCESS) {
224             av_log(avctx, AV_LOG_ERROR, "Failed to enqueue kernel: %d.\n",
225                    cle);
226             err = AVERROR(EIO);
227             goto fail;
228         }
229
230     }
231
232     cle = clFinish(ctx->command_queue);
233     if (cle != CL_SUCCESS) {
234         av_log(avctx, AV_LOG_ERROR, "Failed to finish command queue: %d.\n",
235                cle);
236         err = AVERROR(EIO);
237         goto fail;
238     }
239
240     err = av_frame_copy_props(output, input);
241     if (err < 0)
242         goto fail;
243
244     av_frame_free(&input);
245     av_frame_free(&intermediate);
246
247     av_log(ctx, AV_LOG_DEBUG, "Filter output: %s, %ux%u (%"PRId64").\n",
248            av_get_pix_fmt_name(output->format),
249            output->width, output->height, output->pts);
250
251     return ff_filter_frame(outlink, output);
252
253 fail:
254     clFinish(ctx->command_queue);
255     av_frame_free(&input);
256     av_frame_free(&output);
257     av_frame_free(&intermediate);
258     return err;
259 }
260
261 static av_cold void avgblur_opencl_uninit(AVFilterContext *avctx)
262 {
263     AverageBlurOpenCLContext *ctx = avctx->priv;
264     cl_int cle;
265
266
267     if (ctx->kernel_horiz) {
268         cle = clReleaseKernel(ctx->kernel_horiz);
269         if (cle != CL_SUCCESS)
270             av_log(avctx, AV_LOG_ERROR, "Failed to release "
271                    "kernel: %d.\n", cle);
272     }
273
274     if (ctx->kernel_vert) {
275         cle = clReleaseKernel(ctx->kernel_vert);
276         if (cle != CL_SUCCESS)
277             av_log(avctx, AV_LOG_ERROR, "Failed to release "
278                    "kernel: %d.\n", cle);
279     }
280
281     if (ctx->command_queue) {
282         cle = clReleaseCommandQueue(ctx->command_queue);
283         if (cle != CL_SUCCESS)
284             av_log(avctx, AV_LOG_ERROR, "Failed to release "
285                    "command queue: %d.\n", cle);
286     }
287
288     ff_opencl_filter_uninit(avctx);
289 }
290
291 #define OFFSET(x) offsetof(AverageBlurOpenCLContext, x)
292 #define FLAGS (AV_OPT_FLAG_FILTERING_PARAM | AV_OPT_FLAG_VIDEO_PARAM)
293 static const AVOption avgblur_opencl_options[] = {
294     { "sizeX",  "set horizontal size",  OFFSET(radius),  AV_OPT_TYPE_INT, {.i64=1},   1, 1024, FLAGS },
295     { "planes", "set planes to filter", OFFSET(planes),  AV_OPT_TYPE_INT, {.i64=0xF}, 0,  0xF, FLAGS },
296     { "sizeY",  "set vertical size",    OFFSET(radiusV), AV_OPT_TYPE_INT, {.i64=0},   0, 1024, FLAGS },
297     { NULL }
298 };
299
300 AVFILTER_DEFINE_CLASS(avgblur_opencl);
301
302 static const AVFilterPad avgblur_opencl_inputs[] = {
303     {
304         .name         = "default",
305         .type         = AVMEDIA_TYPE_VIDEO,
306         .filter_frame = &avgblur_opencl_filter_frame,
307         .config_props = &ff_opencl_filter_config_input,
308     },
309     { NULL }
310 };
311
312 static const AVFilterPad avgblur_opencl_outputs[] = {
313     {
314         .name         = "default",
315         .type         = AVMEDIA_TYPE_VIDEO,
316         .config_props = &ff_opencl_filter_config_output,
317     },
318     { NULL }
319 };
320
321 AVFilter ff_vf_avgblur_opencl = {
322     .name           = "avgblur_opencl",
323     .description    = NULL_IF_CONFIG_SMALL("Apply average blur filter"),
324     .priv_size      = sizeof(AverageBlurOpenCLContext),
325     .priv_class     = &avgblur_opencl_class,
326     .init           = &ff_opencl_filter_init,
327     .uninit         = &avgblur_opencl_uninit,
328     .query_formats  = &ff_opencl_filter_query_formats,
329     .inputs         = avgblur_opencl_inputs,
330     .outputs        = avgblur_opencl_outputs,
331     .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE,
332 };