lavfi: Add OpenCL avgblur filter
[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         global_work[0] = output->width;
174         global_work[1] = output->height;
175
176         av_log(avctx, AV_LOG_DEBUG, "Run kernel on plane %d "
177                "(%"SIZE_SPECIFIER"x%"SIZE_SPECIFIER").\n",
178                p, global_work[0], global_work[1]);
179
180         cle = clEnqueueNDRangeKernel(ctx->command_queue, ctx->kernel_horiz, 2, NULL,
181                                      global_work, NULL,
182                                      0, NULL, NULL);
183         if (cle != CL_SUCCESS) {
184             av_log(avctx, AV_LOG_ERROR, "Failed to enqueue kernel: %d.\n",
185                    cle);
186             err = AVERROR(EIO);
187             goto fail;
188         }
189
190         cle = clSetKernelArg(ctx->kernel_vert, 0, sizeof(cl_mem), &dst);
191         if (cle != CL_SUCCESS) {
192             av_log(avctx, AV_LOG_ERROR, "Failed to set kernel "
193                    "destination image argument: %d.\n", cle);
194             goto fail;
195         }
196         cle = clSetKernelArg(ctx->kernel_vert, 1, sizeof(cl_mem), &inter);
197         if (cle != CL_SUCCESS) {
198             av_log(avctx, AV_LOG_ERROR, "Failed to set kernel "
199                    "source image argument: %d.\n", cle);
200             goto fail;
201         }
202         cle = clSetKernelArg(ctx->kernel_vert, 2, sizeof(cl_int), &radius_y);
203         if (cle != CL_SUCCESS) {
204             av_log(avctx, AV_LOG_ERROR, "Failed to set kernel "
205                    "sizeY argument: %d.\n", cle);
206             goto fail;
207         }
208
209         global_work[0] = output->width;
210         global_work[1] = output->height;
211
212         av_log(avctx, AV_LOG_DEBUG, "Run kernel on plane %d "
213                "(%"SIZE_SPECIFIER"x%"SIZE_SPECIFIER").\n",
214                p, global_work[0], global_work[1]);
215
216         cle = clEnqueueNDRangeKernel(ctx->command_queue, ctx->kernel_vert, 2, NULL,
217                                      global_work, NULL,
218                                      0, NULL, NULL);
219         if (cle != CL_SUCCESS) {
220             av_log(avctx, AV_LOG_ERROR, "Failed to enqueue kernel: %d.\n",
221                    cle);
222             err = AVERROR(EIO);
223             goto fail;
224         }
225
226     }
227
228     cle = clFinish(ctx->command_queue);
229     if (cle != CL_SUCCESS) {
230         av_log(avctx, AV_LOG_ERROR, "Failed to finish command queue: %d.\n",
231                cle);
232         err = AVERROR(EIO);
233         goto fail;
234     }
235
236     err = av_frame_copy_props(output, input);
237     if (err < 0)
238         goto fail;
239
240     av_frame_free(&input);
241     av_frame_free(&intermediate);
242
243     av_log(ctx, AV_LOG_DEBUG, "Filter output: %s, %ux%u (%"PRId64").\n",
244            av_get_pix_fmt_name(output->format),
245            output->width, output->height, output->pts);
246
247     return ff_filter_frame(outlink, output);
248
249 fail:
250     clFinish(ctx->command_queue);
251     av_frame_free(&input);
252     av_frame_free(&output);
253     av_frame_free(&intermediate);
254     return err;
255 }
256
257 static av_cold void avgblur_opencl_uninit(AVFilterContext *avctx)
258 {
259     AverageBlurOpenCLContext *ctx = avctx->priv;
260     cl_int cle;
261
262
263     if (ctx->kernel_horiz) {
264         cle = clReleaseKernel(ctx->kernel_horiz);
265         if (cle != CL_SUCCESS)
266             av_log(avctx, AV_LOG_ERROR, "Failed to release "
267                    "kernel: %d.\n", cle);
268     }
269
270     if (ctx->kernel_vert) {
271         cle = clReleaseKernel(ctx->kernel_vert);
272         if (cle != CL_SUCCESS)
273             av_log(avctx, AV_LOG_ERROR, "Failed to release "
274                    "kernel: %d.\n", cle);
275     }
276
277     if (ctx->command_queue) {
278         cle = clReleaseCommandQueue(ctx->command_queue);
279         if (cle != CL_SUCCESS)
280             av_log(avctx, AV_LOG_ERROR, "Failed to release "
281                    "command queue: %d.\n", cle);
282     }
283
284     ff_opencl_filter_uninit(avctx);
285 }
286
287 #define OFFSET(x) offsetof(AverageBlurOpenCLContext, x)
288 #define FLAGS (AV_OPT_FLAG_FILTERING_PARAM | AV_OPT_FLAG_VIDEO_PARAM)
289 static const AVOption avgblur_opencl_options[] = {
290     { "sizeX",  "set horizontal size",  OFFSET(radius),  AV_OPT_TYPE_INT, {.i64=1},   1, 1024, FLAGS },
291     { "planes", "set planes to filter", OFFSET(planes),  AV_OPT_TYPE_INT, {.i64=0xF}, 0,  0xF, FLAGS },
292     { "sizeY",  "set vertical size",    OFFSET(radiusV), AV_OPT_TYPE_INT, {.i64=0},   0, 1024, FLAGS },
293     { NULL }
294 };
295
296 AVFILTER_DEFINE_CLASS(avgblur_opencl);
297
298 static const AVFilterPad avgblur_opencl_inputs[] = {
299     {
300         .name         = "default",
301         .type         = AVMEDIA_TYPE_VIDEO,
302         .filter_frame = &avgblur_opencl_filter_frame,
303         .config_props = &ff_opencl_filter_config_input,
304     },
305     { NULL }
306 };
307
308 static const AVFilterPad avgblur_opencl_outputs[] = {
309     {
310         .name         = "default",
311         .type         = AVMEDIA_TYPE_VIDEO,
312         .config_props = &ff_opencl_filter_config_output,
313     },
314     { NULL }
315 };
316
317 AVFilter ff_vf_avgblur_opencl = {
318     .name           = "avgblur_opencl",
319     .description    = NULL_IF_CONFIG_SMALL("Apply average blur filter"),
320     .priv_size      = sizeof(AverageBlurOpenCLContext),
321     .priv_class     = &avgblur_opencl_class,
322     .init           = &ff_opencl_filter_init,
323     .uninit         = &avgblur_opencl_uninit,
324     .query_formats  = &ff_opencl_filter_query_formats,
325     .inputs         = avgblur_opencl_inputs,
326     .outputs        = avgblur_opencl_outputs,
327     .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE,
328 };