2 * This file is part of FFmpeg.
4 * FFmpeg is free software; you can redistribute it and/or
5 * modify it under the terms of the GNU Lesser General Public
6 * License as published by the Free Software Foundation; either
7 * version 2.1 of the License, or (at your option) any later version.
9 * FFmpeg is distributed in the hope that it will be useful,
10 * but WITHOUT ANY WARRANTY; without even the implied warranty of
11 * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU
12 * Lesser General Public License for more details.
14 * You should have received a copy of the GNU Lesser General Public
15 * License along with FFmpeg; if not, write to the Free Software
16 * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA
19 #include "libavutil/log.h"
20 #include "libavutil/mem.h"
21 #include "libavutil/opt.h"
22 #include "libavutil/pixdesc.h"
25 #include "framesync.h"
28 #include "opencl_source.h"
31 typedef struct OverlayOpenCLContext {
32 OpenCLFilterContext ocf;
36 cl_command_queue command_queue;
47 } OverlayOpenCLContext;
49 static int overlay_opencl_load(AVFilterContext *avctx,
50 enum AVPixelFormat main_format,
51 enum AVPixelFormat overlay_format)
53 OverlayOpenCLContext *ctx = avctx->priv;
55 const char *source = ff_opencl_source_overlay;
57 const AVPixFmtDescriptor *main_desc, *overlay_desc;
58 int err, i, main_planes, overlay_planes;
60 main_desc = av_pix_fmt_desc_get(main_format);
61 overlay_desc = av_pix_fmt_desc_get(overlay_format);
63 main_planes = overlay_planes = 0;
64 for (i = 0; i < main_desc->nb_components; i++)
65 main_planes = FFMAX(main_planes,
66 main_desc->comp[i].plane + 1);
67 for (i = 0; i < overlay_desc->nb_components; i++)
68 overlay_planes = FFMAX(overlay_planes,
69 overlay_desc->comp[i].plane + 1);
71 ctx->nb_planes = main_planes;
72 ctx->x_subsample = 1 << main_desc->log2_chroma_w;
73 ctx->y_subsample = 1 << main_desc->log2_chroma_h;
75 if (ctx->x_position % ctx->x_subsample ||
76 ctx->y_position % ctx->y_subsample) {
77 av_log(avctx, AV_LOG_WARNING, "Warning: overlay position (%d, %d) "
78 "does not match subsampling (%d, %d).\n",
79 ctx->x_position, ctx->y_position,
80 ctx->x_subsample, ctx->y_subsample);
83 if (main_planes == overlay_planes) {
84 if (main_desc->nb_components == overlay_desc->nb_components)
85 kernel = "overlay_no_alpha";
87 kernel = "overlay_internal_alpha";
88 ctx->alpha_separate = 0;
90 kernel = "overlay_external_alpha";
91 ctx->alpha_separate = 1;
94 av_log(avctx, AV_LOG_DEBUG, "Using kernel %s.\n", kernel);
96 err = ff_opencl_filter_load_program(avctx, &source, 1);
100 ctx->command_queue = clCreateCommandQueue(ctx->ocf.hwctx->context,
101 ctx->ocf.hwctx->device_id,
103 if (!ctx->command_queue) {
104 av_log(avctx, AV_LOG_ERROR, "Failed to create OpenCL "
105 "command queue: %d.\n", cle);
110 ctx->kernel = clCreateKernel(ctx->ocf.program, kernel, &cle);
112 av_log(avctx, AV_LOG_ERROR, "Failed to create kernel: %d.\n", cle);
117 ctx->initialised = 1;
121 if (ctx->command_queue)
122 clReleaseCommandQueue(ctx->command_queue);
124 clReleaseKernel(ctx->kernel);
128 static int overlay_opencl_blend(FFFrameSync *fs)
130 AVFilterContext *avctx = fs->parent;
131 AVFilterLink *outlink = avctx->outputs[0];
132 OverlayOpenCLContext *ctx = avctx->priv;
133 AVFrame *input_main, *input_overlay;
137 size_t global_work[2];
141 err = ff_framesync_get_frame(fs, 0, &input_main, 0);
144 err = ff_framesync_get_frame(fs, 1, &input_overlay, 0);
148 if (!ctx->initialised) {
149 AVHWFramesContext *main_fc =
150 (AVHWFramesContext*)input_main->hw_frames_ctx->data;
151 AVHWFramesContext *overlay_fc =
152 (AVHWFramesContext*)input_overlay->hw_frames_ctx->data;
154 err = overlay_opencl_load(avctx, main_fc->sw_format,
155 overlay_fc->sw_format);
160 output = ff_get_video_buffer(outlink, outlink->w, outlink->h);
162 err = AVERROR(ENOMEM);
166 for (plane = 0; plane < ctx->nb_planes; plane++) {
169 mem = (cl_mem)output->data[plane];
170 cle = clSetKernelArg(ctx->kernel, kernel_arg++, sizeof(cl_mem), &mem);
171 if (cle != CL_SUCCESS)
172 goto fail_kernel_arg;
174 mem = (cl_mem)input_main->data[plane];
175 cle = clSetKernelArg(ctx->kernel, kernel_arg++, sizeof(cl_mem), &mem);
176 if (cle != CL_SUCCESS)
177 goto fail_kernel_arg;
179 mem = (cl_mem)input_overlay->data[plane];
180 cle = clSetKernelArg(ctx->kernel, kernel_arg++, sizeof(cl_mem), &mem);
181 if (cle != CL_SUCCESS)
182 goto fail_kernel_arg;
184 if (ctx->alpha_separate) {
185 mem = (cl_mem)input_overlay->data[ctx->nb_planes];
186 cle = clSetKernelArg(ctx->kernel, kernel_arg++, sizeof(cl_mem), &mem);
187 if (cle != CL_SUCCESS)
188 goto fail_kernel_arg;
191 x = ctx->x_position / (plane == 0 ? 1 : ctx->x_subsample);
192 y = ctx->y_position / (plane == 0 ? 1 : ctx->y_subsample);
194 cle = clSetKernelArg(ctx->kernel, kernel_arg++, sizeof(cl_int), &x);
195 if (cle != CL_SUCCESS)
196 goto fail_kernel_arg;
197 cle = clSetKernelArg(ctx->kernel, kernel_arg++, sizeof(cl_int), &y);
198 if (cle != CL_SUCCESS)
199 goto fail_kernel_arg;
201 if (ctx->alpha_separate) {
202 cl_int alpha_adj_x = plane == 0 ? 1 : ctx->x_subsample;
203 cl_int alpha_adj_y = plane == 0 ? 1 : ctx->y_subsample;
205 cle = clSetKernelArg(ctx->kernel, kernel_arg++, sizeof(cl_int), &alpha_adj_x);
206 if (cle != CL_SUCCESS)
207 goto fail_kernel_arg;
208 cle = clSetKernelArg(ctx->kernel, kernel_arg++, sizeof(cl_int), &alpha_adj_y);
209 if (cle != CL_SUCCESS)
210 goto fail_kernel_arg;
213 err = ff_opencl_filter_work_size_from_image(avctx, global_work,
218 cle = clEnqueueNDRangeKernel(ctx->command_queue, ctx->kernel, 2, NULL,
219 global_work, NULL, 0, NULL, NULL);
220 if (cle != CL_SUCCESS) {
221 av_log(avctx, AV_LOG_ERROR, "Failed to enqueue "
222 "overlay kernel for plane %d: %d.\n", cle, plane);
228 cle = clFinish(ctx->command_queue);
229 if (cle != CL_SUCCESS) {
230 av_log(avctx, AV_LOG_ERROR, "Failed to finish "
231 "command queue: %d.\n", cle);
236 err = av_frame_copy_props(output, input_main);
238 av_log(avctx, AV_LOG_DEBUG, "Filter output: %s, %ux%u (%"PRId64").\n",
239 av_get_pix_fmt_name(output->format),
240 output->width, output->height, output->pts);
242 return ff_filter_frame(outlink, output);
245 av_log(avctx, AV_LOG_ERROR, "Failed to set kernel arg %d: %d.\n",
249 av_frame_free(&output);
253 static int overlay_opencl_config_output(AVFilterLink *outlink)
255 AVFilterContext *avctx = outlink->src;
256 OverlayOpenCLContext *ctx = avctx->priv;
259 err = ff_opencl_filter_config_output(outlink);
263 err = ff_framesync_init_dualinput(&ctx->fs, avctx);
267 return ff_framesync_configure(&ctx->fs);
270 static av_cold int overlay_opencl_init(AVFilterContext *avctx)
272 OverlayOpenCLContext *ctx = avctx->priv;
274 ctx->fs.on_event = &overlay_opencl_blend;
276 return ff_opencl_filter_init(avctx);
279 static int overlay_opencl_activate(AVFilterContext *avctx)
281 OverlayOpenCLContext *ctx = avctx->priv;
283 return ff_framesync_activate(&ctx->fs);
286 static av_cold void overlay_opencl_uninit(AVFilterContext *avctx)
288 OverlayOpenCLContext *ctx = avctx->priv;
292 cle = clReleaseKernel(ctx->kernel);
293 if (cle != CL_SUCCESS)
294 av_log(avctx, AV_LOG_ERROR, "Failed to release "
295 "kernel: %d.\n", cle);
298 if (ctx->command_queue) {
299 cle = clReleaseCommandQueue(ctx->command_queue);
300 if (cle != CL_SUCCESS)
301 av_log(avctx, AV_LOG_ERROR, "Failed to release "
302 "command queue: %d.\n", cle);
305 ff_opencl_filter_uninit(avctx);
307 ff_framesync_uninit(&ctx->fs);
310 #define OFFSET(x) offsetof(OverlayOpenCLContext, x)
311 #define FLAGS (AV_OPT_FLAG_FILTERING_PARAM | AV_OPT_FLAG_VIDEO_PARAM)
312 static const AVOption overlay_opencl_options[] = {
313 { "x", "Overlay x position",
314 OFFSET(x_position), AV_OPT_TYPE_INT, { .i64 = 0 }, 0, INT_MAX, .flags = FLAGS },
315 { "y", "Overlay y position",
316 OFFSET(y_position), AV_OPT_TYPE_INT, { .i64 = 0 }, 0, INT_MAX, .flags = FLAGS },
320 AVFILTER_DEFINE_CLASS(overlay_opencl);
322 static const AVFilterPad overlay_opencl_inputs[] = {
325 .type = AVMEDIA_TYPE_VIDEO,
326 .config_props = &ff_opencl_filter_config_input,
330 .type = AVMEDIA_TYPE_VIDEO,
331 .config_props = &ff_opencl_filter_config_input,
336 static const AVFilterPad overlay_opencl_outputs[] = {
339 .type = AVMEDIA_TYPE_VIDEO,
340 .config_props = &overlay_opencl_config_output,
345 AVFilter ff_vf_overlay_opencl = {
346 .name = "overlay_opencl",
347 .description = NULL_IF_CONFIG_SMALL("Overlay one video on top of another"),
348 .priv_size = sizeof(OverlayOpenCLContext),
349 .priv_class = &overlay_opencl_class,
350 .init = &overlay_opencl_init,
351 .uninit = &overlay_opencl_uninit,
352 .query_formats = &ff_opencl_filter_query_formats,
353 .activate = &overlay_opencl_activate,
354 .inputs = overlay_opencl_inputs,
355 .outputs = overlay_opencl_outputs,
356 .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE,