2 * Copyright (c) 2018 Dylan Fernando
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"
30 #include "opencl_source.h"
34 typedef struct AverageBlurOpenCLContext {
35 OpenCLFilterContext ocf;
38 cl_kernel kernel_horiz;
39 cl_kernel kernel_vert;
40 cl_command_queue command_queue;
46 } AverageBlurOpenCLContext;
49 static int avgblur_opencl_init(AVFilterContext *avctx)
51 AverageBlurOpenCLContext *ctx = avctx->priv;
55 err = ff_opencl_filter_load_program(avctx, &ff_opencl_source_avgblur, 1);
59 ctx->command_queue = clCreateCommandQueue(ctx->ocf.hwctx->context,
60 ctx->ocf.hwctx->device_id,
62 if (!ctx->command_queue) {
63 av_log(avctx, AV_LOG_ERROR, "Failed to create OpenCL "
64 "command queue: %d.\n", cle);
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);
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);
83 if (ctx->radiusV <= 0) {
84 ctx->radiusV = ctx->radius;
91 if (ctx->command_queue)
92 clReleaseCommandQueue(ctx->command_queue);
93 if (ctx->kernel_horiz)
94 clReleaseKernel(ctx->kernel_horiz);
96 clReleaseKernel(ctx->kernel_vert);
100 static int avgblur_opencl_filter_frame(AVFilterLink *inlink, AVFrame *input)
102 AVFilterContext *avctx = inlink->dst;
103 AVFilterLink *outlink = avctx->outputs[0];
104 AverageBlurOpenCLContext *ctx = avctx->priv;
105 AVFrame *output = NULL;
106 AVFrame *intermediate = NULL;
108 size_t global_work[2];
109 cl_mem src, dst, inter;
110 int err, p, radius_x, radius_y;
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);
116 if (!input->hw_frames_ctx)
117 return AVERROR(EINVAL);
119 if (!ctx->initialised) {
120 err = avgblur_opencl_init(avctx);
126 output = ff_get_video_buffer(outlink, outlink->w, outlink->h);
128 err = AVERROR(ENOMEM);
132 intermediate = ff_get_video_buffer(outlink, outlink->w, outlink->h);
134 err = AVERROR(ENOMEM);
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];
146 radius_x = ctx->radius;
147 radius_y = ctx->radiusV;
149 if (!(ctx->planes & (1 << p))) {
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;
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;
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;
176 err = ff_opencl_filter_work_size_from_image(avctx, global_work,
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]);
185 cle = clEnqueueNDRangeKernel(ctx->command_queue, ctx->kernel_horiz, 2, NULL,
188 if (cle != CL_SUCCESS) {
189 av_log(avctx, AV_LOG_ERROR, "Failed to enqueue kernel: %d.\n",
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;
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;
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;
217 err = ff_opencl_filter_work_size_from_image(avctx, global_work,
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]);
226 cle = clEnqueueNDRangeKernel(ctx->command_queue, ctx->kernel_vert, 2, NULL,
229 if (cle != CL_SUCCESS) {
230 av_log(avctx, AV_LOG_ERROR, "Failed to enqueue kernel: %d.\n",
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",
246 err = av_frame_copy_props(output, input);
250 av_frame_free(&input);
251 av_frame_free(&intermediate);
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);
257 return ff_filter_frame(outlink, output);
260 clFinish(ctx->command_queue);
261 av_frame_free(&input);
262 av_frame_free(&output);
263 av_frame_free(&intermediate);
267 static av_cold void avgblur_opencl_uninit(AVFilterContext *avctx)
269 AverageBlurOpenCLContext *ctx = avctx->priv;
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);
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);
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);
294 ff_opencl_filter_uninit(avctx);
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 },
306 AVFILTER_DEFINE_CLASS(avgblur_opencl);
308 static const AVFilterPad avgblur_opencl_inputs[] = {
311 .type = AVMEDIA_TYPE_VIDEO,
312 .filter_frame = &avgblur_opencl_filter_frame,
313 .config_props = &ff_opencl_filter_config_input,
318 static const AVFilterPad avgblur_opencl_outputs[] = {
321 .type = AVMEDIA_TYPE_VIDEO,
322 .config_props = &ff_opencl_filter_config_output,
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,