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