2 * Copyright (c) 2018 Danil Iashchenko
4 * This file is part of FFmpeg.
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.
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.
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
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"
32 #include "opencl_source.h"
35 typedef struct NeighborOpenCLContext {
36 OpenCLFilterContext ocf;
40 cl_command_queue command_queue;
44 cl_float threshold[4];
48 } NeighborOpenCLContext;
50 static int neighbor_opencl_init(AVFilterContext *avctx)
52 NeighborOpenCLContext *ctx = avctx->priv;
53 const char *kernel_name;
57 err = ff_opencl_filter_load_program(avctx, &ff_opencl_source_neighbor, 1);
61 ctx->command_queue = clCreateCommandQueue(ctx->ocf.hwctx->context,
62 ctx->ocf.hwctx->device_id,
64 CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create OpenCL "
65 "command queue %d.\n", cle);
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";
72 ctx->kernel = clCreateKernel(ctx->ocf.program, kernel_name, &cle);
73 CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create "
80 if (ctx->command_queue)
81 clReleaseCommandQueue(ctx->command_queue);
83 clReleaseKernel(ctx->kernel);
87 static int neighbor_opencl_make_filter_params(AVFilterContext *avctx)
89 NeighborOpenCLContext *ctx = avctx->priv;
95 for (i = 0; i < 4; i++) {
96 ctx->threshold[i] /= 255.0;
100 for (i = 0; i < 8; i++) {
101 if (ctx->coordinates & (1 << i)) {
102 matrix[i > 3 ? i + 1: i] = 1;
105 buffer = clCreateBuffer(ctx->ocf.hwctx->context,
107 CL_MEM_COPY_HOST_PTR |
108 CL_MEM_HOST_NO_ACCESS,
109 9 * sizeof(cl_int), matrix, &cle);
111 av_log(avctx, AV_LOG_ERROR, "Failed to create matrix buffer: "
121 static int neighbor_opencl_filter_frame(AVFilterLink *inlink, AVFrame *input)
123 AVFilterContext *avctx = inlink->dst;
124 AVFilterLink *outlink = avctx->outputs[0];
125 NeighborOpenCLContext *ctx = avctx->priv;
126 AVFrame *output = NULL;
128 size_t global_work[2];
131 size_t origin[3] = {0, 0, 0};
132 size_t region[3] = {0, 0, 1};
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);
138 if (!input->hw_frames_ctx)
139 return AVERROR(EINVAL);
141 if (!ctx->initialised) {
142 err = neighbor_opencl_init(avctx);
146 err = neighbor_opencl_make_filter_params(avctx);
152 output = ff_get_video_buffer(outlink, outlink->w, outlink->h);
154 err = AVERROR(ENOMEM);
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];
165 if (ctx->threshold[p] == 0) {
166 err = ff_opencl_filter_work_size_from_image(avctx, region, output, p, 0);
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",
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);
180 err = ff_opencl_filter_work_size_from_image(avctx, global_work, output, p, 0);
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]);
188 cle = clEnqueueNDRangeKernel(ctx->command_queue, ctx->kernel, 2, NULL,
191 CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to enqueue "
192 "kernel: %d.\n", cle);
196 cle = clFinish(ctx->command_queue);
197 CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to finish command queue: %d.\n", cle);
199 err = av_frame_copy_props(output, input);
203 av_frame_free(&input);
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);
209 return ff_filter_frame(outlink, output);
212 clFinish(ctx->command_queue);
213 av_frame_free(&input);
214 av_frame_free(&output);
218 static av_cold void neighbor_opencl_uninit(AVFilterContext *avctx)
220 NeighborOpenCLContext *ctx = avctx->priv;
223 clReleaseMemObject(ctx->coord);
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);
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);
239 ff_opencl_filter_uninit(avctx);
242 static const AVFilterPad neighbor_opencl_inputs[] = {
245 .type = AVMEDIA_TYPE_VIDEO,
246 .filter_frame = &neighbor_opencl_filter_frame,
247 .config_props = &ff_opencl_filter_config_input,
252 static const AVFilterPad neighbor_opencl_outputs[] = {
255 .type = AVMEDIA_TYPE_VIDEO,
256 .config_props = &ff_opencl_filter_config_output,
261 #define OFFSET(x) offsetof(NeighborOpenCLContext, x)
262 #define FLAGS (AV_OPT_FLAG_FILTERING_PARAM | AV_OPT_FLAG_VIDEO_PARAM)
264 #if CONFIG_EROSION_OPENCL_FILTER
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 },
275 AVFILTER_DEFINE_CLASS(erosion_opencl);
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,
290 #endif /* CONFIG_EROSION_OPENCL_FILTER */
292 #if CONFIG_DILATION_OPENCL_FILTER
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 },
303 AVFILTER_DEFINE_CLASS(dilation_opencl);
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,
318 #endif /* CONFIG_DILATION_OPENCL_FILTER */