avfilter/af_afir: do not over allocate fft buffer
[ffmpeg.git] / libavfilter / vf_neighbor_opencl.c
1 /*
2  * Copyright (c) 2018 Danil Iashchenko
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 #include "libavutil/avstring.h"
27
28
29 #include "avfilter.h"
30 #include "internal.h"
31 #include "opencl.h"
32 #include "opencl_source.h"
33 #include "video.h"
34
35 typedef struct NeighborOpenCLContext {
36     OpenCLFilterContext ocf;
37
38     int              initialised;
39     cl_kernel        kernel;
40     cl_command_queue command_queue;
41
42     char *matrix_str[4];
43
44     cl_float threshold[4];
45     cl_int coordinates;
46     cl_mem coord;
47
48 } NeighborOpenCLContext;
49
50 static int neighbor_opencl_init(AVFilterContext *avctx)
51 {
52     NeighborOpenCLContext *ctx = avctx->priv;
53     const char *kernel_name;
54     cl_int cle;
55     int err;
56
57     err = ff_opencl_filter_load_program(avctx, &ff_opencl_source_neighbor, 1);
58     if (err < 0)
59         goto fail;
60
61     ctx->command_queue = clCreateCommandQueue(ctx->ocf.hwctx->context,
62                                               ctx->ocf.hwctx->device_id,
63                                               0, &cle);
64     CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create OpenCL "
65                      "command queue %d.\n", cle);
66
67     if (!strcmp(avctx->filter->name, "erosion_opencl")){
68         kernel_name = "erosion_global";
69     } else if (!strcmp(avctx->filter->name, "dilation_opencl")){
70         kernel_name = "dilation_global";
71     }
72     ctx->kernel = clCreateKernel(ctx->ocf.program, kernel_name, &cle);
73     CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create "
74                      "kernel %d.\n", cle);
75
76     ctx->initialised = 1;
77     return 0;
78
79 fail:
80     if (ctx->command_queue)
81         clReleaseCommandQueue(ctx->command_queue);
82     if (ctx->kernel)
83         clReleaseKernel(ctx->kernel);
84     return err;
85 }
86
87 static int neighbor_opencl_make_filter_params(AVFilterContext *avctx)
88 {
89     NeighborOpenCLContext *ctx = avctx->priv;
90     cl_int matrix[9];
91     cl_mem buffer;
92     cl_int cle;
93     int i;
94
95     for (i = 0; i < 4; i++) {
96         ctx->threshold[i] /= 255.0;
97     }
98
99     matrix[4] = 0;
100     for (i = 0; i < 8; i++) {
101         if (ctx->coordinates & (1 << i)) {
102             matrix[i > 3 ? i + 1: i] = 1;
103         }
104     }
105     buffer = clCreateBuffer(ctx->ocf.hwctx->context,
106                             CL_MEM_READ_ONLY |
107                             CL_MEM_COPY_HOST_PTR |
108                             CL_MEM_HOST_NO_ACCESS,
109                             9 * sizeof(cl_int), matrix, &cle);
110     if (!buffer) {
111         av_log(avctx, AV_LOG_ERROR, "Failed to create matrix buffer: "
112                "%d.\n", cle);
113         return AVERROR(EIO);
114     }
115     ctx->coord = buffer;
116
117     return 0;
118 }
119
120
121 static int neighbor_opencl_filter_frame(AVFilterLink *inlink, AVFrame *input)
122 {
123     AVFilterContext *avctx = inlink->dst;
124     AVFilterLink *outlink = avctx->outputs[0];
125     NeighborOpenCLContext *ctx = avctx->priv;
126     AVFrame *output = NULL;
127     cl_int cle;
128     size_t global_work[2];
129     cl_mem src, dst;
130     int err, p;
131     size_t origin[3] = {0, 0, 0};
132     size_t region[3] = {0, 0, 1};
133
134     av_log(ctx, AV_LOG_DEBUG, "Filter input: %s, %ux%u (%"PRId64").\n",
135            av_get_pix_fmt_name(input->format),
136            input->width, input->height, input->pts);
137
138     if (!input->hw_frames_ctx)
139         return AVERROR(EINVAL);
140
141     if (!ctx->initialised) {
142         err = neighbor_opencl_init(avctx);
143         if (err < 0)
144             goto fail;
145
146         err = neighbor_opencl_make_filter_params(avctx);
147         if (err < 0)
148             goto fail;
149
150     }
151
152     output = ff_get_video_buffer(outlink, outlink->w, outlink->h);
153     if (!output) {
154         err = AVERROR(ENOMEM);
155         goto fail;
156     }
157
158     for (p = 0; p < FF_ARRAY_ELEMS(output->data); p++) {
159         src = (cl_mem) input->data[p];
160         dst = (cl_mem)output->data[p];
161
162         if (!dst)
163             break;
164
165         if (ctx->threshold[p] == 0) {
166             err = ff_opencl_filter_work_size_from_image(avctx, region, output, p, 0);
167             if (err < 0)
168                 goto fail;
169
170             cle = clEnqueueCopyImage(ctx->command_queue, src, dst,
171                                      origin, origin, region, 0, NULL, NULL);
172             CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to copy plane %d: %d.\n",
173                              p, cle);
174         } else {
175             CL_SET_KERNEL_ARG(ctx->kernel, 0, cl_mem,   &dst);
176             CL_SET_KERNEL_ARG(ctx->kernel, 1, cl_mem,   &src);
177             CL_SET_KERNEL_ARG(ctx->kernel, 2, cl_float, &ctx->threshold[p]);
178             CL_SET_KERNEL_ARG(ctx->kernel, 3, cl_mem,   &ctx->coord);
179
180             err = ff_opencl_filter_work_size_from_image(avctx, global_work, output, p, 0);
181             if (err < 0)
182                 goto fail;
183
184             av_log(avctx, AV_LOG_DEBUG, "Run kernel on plane %d "
185                    "(%"SIZE_SPECIFIER"x%"SIZE_SPECIFIER").\n",
186                    p, global_work[0], global_work[1]);
187
188             cle = clEnqueueNDRangeKernel(ctx->command_queue, ctx->kernel, 2, NULL,
189                                          global_work, NULL,
190                                          0, NULL, NULL);
191             CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to enqueue "
192                              "kernel: %d.\n", cle);
193         }
194     }
195
196     cle = clFinish(ctx->command_queue);
197     CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to finish command queue: %d.\n", cle);
198
199     err = av_frame_copy_props(output, input);
200     if (err < 0)
201         goto fail;
202
203     av_frame_free(&input);
204
205     av_log(ctx, AV_LOG_DEBUG, "Filter output: %s, %ux%u (%"PRId64").\n",
206            av_get_pix_fmt_name(output->format),
207            output->width, output->height, output->pts);
208
209     return ff_filter_frame(outlink, output);
210
211 fail:
212     clFinish(ctx->command_queue);
213     av_frame_free(&input);
214     av_frame_free(&output);
215     return err;
216 }
217
218 static av_cold void neighbor_opencl_uninit(AVFilterContext *avctx)
219 {
220     NeighborOpenCLContext *ctx = avctx->priv;
221     cl_int cle;
222
223     clReleaseMemObject(ctx->coord);
224
225     if (ctx->kernel) {
226         cle = clReleaseKernel(ctx->kernel);
227         if (cle != CL_SUCCESS)
228             av_log(avctx, AV_LOG_ERROR, "Failed to release "
229                    "kernel: %d.\n", cle);
230     }
231
232     if (ctx->command_queue) {
233         cle = clReleaseCommandQueue(ctx->command_queue);
234         if (cle != CL_SUCCESS)
235             av_log(avctx, AV_LOG_ERROR, "Failed to release "
236                    "command queue: %d.\n", cle);
237     }
238
239     ff_opencl_filter_uninit(avctx);
240 }
241
242 static const AVFilterPad neighbor_opencl_inputs[] = {
243     {
244         .name         = "default",
245         .type         = AVMEDIA_TYPE_VIDEO,
246         .filter_frame = &neighbor_opencl_filter_frame,
247         .config_props = &ff_opencl_filter_config_input,
248     },
249     { NULL }
250 };
251
252 static const AVFilterPad neighbor_opencl_outputs[] = {
253     {
254         .name         = "default",
255         .type         = AVMEDIA_TYPE_VIDEO,
256         .config_props = &ff_opencl_filter_config_output,
257     },
258     { NULL }
259 };
260
261 #define OFFSET(x) offsetof(NeighborOpenCLContext, x)
262 #define FLAGS (AV_OPT_FLAG_FILTERING_PARAM | AV_OPT_FLAG_VIDEO_PARAM)
263
264 #if CONFIG_EROSION_OPENCL_FILTER
265
266 static const AVOption erosion_opencl_options[] = {
267     { "threshold0",  "set threshold for 1st plane",   OFFSET(threshold[0]),   AV_OPT_TYPE_FLOAT, {.dbl=65535.0}, 0.0, 65535, FLAGS },
268     { "threshold1",  "set threshold for 2nd plane",   OFFSET(threshold[1]),   AV_OPT_TYPE_FLOAT, {.dbl=65535.0}, 0.0, 65535, FLAGS },
269     { "threshold2",  "set threshold for 3rd plane",   OFFSET(threshold[2]),   AV_OPT_TYPE_FLOAT, {.dbl=65535.0}, 0.0, 65535, FLAGS },
270     { "threshold3",  "set threshold for 4th plane",   OFFSET(threshold[3]),   AV_OPT_TYPE_FLOAT, {.dbl=65535.0}, 0.0, 65535, FLAGS },
271     { "coordinates", "set coordinates",               OFFSET(coordinates),    AV_OPT_TYPE_INT,   {.i64=255},     0,   255,   FLAGS },
272     { NULL }
273 };
274
275 AVFILTER_DEFINE_CLASS(erosion_opencl);
276
277 AVFilter ff_vf_erosion_opencl = {
278     .name           = "erosion_opencl",
279     .description    = NULL_IF_CONFIG_SMALL("Apply erosion effect"),
280     .priv_size      = sizeof(NeighborOpenCLContext),
281     .priv_class     = &erosion_opencl_class,
282     .init           = &ff_opencl_filter_init,
283     .uninit         = &neighbor_opencl_uninit,
284     .query_formats  = &ff_opencl_filter_query_formats,
285     .inputs         = neighbor_opencl_inputs,
286     .outputs        = neighbor_opencl_outputs,
287     .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE,
288 };
289
290 #endif /* CONFIG_EROSION_OPENCL_FILTER */
291
292 #if CONFIG_DILATION_OPENCL_FILTER
293
294 static const AVOption dilation_opencl_options[] = {
295     { "threshold0",  "set threshold for 1st plane",   OFFSET(threshold[0]),   AV_OPT_TYPE_FLOAT, {.dbl=65535.0}, 0.0, 65535, FLAGS },
296     { "threshold1",  "set threshold for 2nd plane",   OFFSET(threshold[1]),   AV_OPT_TYPE_FLOAT, {.dbl=65535.0}, 0.0, 65535, FLAGS },
297     { "threshold2",  "set threshold for 3rd plane",   OFFSET(threshold[2]),   AV_OPT_TYPE_FLOAT, {.dbl=65535.0}, 0.0, 65535, FLAGS },
298     { "threshold3",  "set threshold for 4th plane",   OFFSET(threshold[3]),   AV_OPT_TYPE_FLOAT, {.dbl=65535.0}, 0.0, 65535, FLAGS },
299     { "coordinates", "set coordinates",               OFFSET(coordinates),    AV_OPT_TYPE_INT,   {.i64=255},     0,   255,   FLAGS },
300     { NULL }
301 };
302
303 AVFILTER_DEFINE_CLASS(dilation_opencl);
304
305 AVFilter ff_vf_dilation_opencl = {
306     .name           = "dilation_opencl",
307     .description    = NULL_IF_CONFIG_SMALL("Apply dilation effect"),
308     .priv_size      = sizeof(NeighborOpenCLContext),
309     .priv_class     = &dilation_opencl_class,
310     .init           = &ff_opencl_filter_init,
311     .uninit         = &neighbor_opencl_uninit,
312     .query_formats  = &ff_opencl_filter_query_formats,
313     .inputs         = neighbor_opencl_inputs,
314     .outputs        = neighbor_opencl_outputs,
315     .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE,
316 };
317
318 #endif /* CONFIG_DILATION_OPENCL_FILTER */