FFmpeg: libavfilter/vf_overlay_cuda.c Source File

FFmpeg
vf_overlay_cuda.c
Go to the documentation of this file.
1 /*
2  * Copyright (c) 2020 Yaroslav Pogrebnyak <yyyaroslav@gmail.com>
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 /**
22  * @file
23  * Overlay one video on top of another using cuda hardware acceleration
24  */
25 
26 #include "libavutil/log.h"
27 #include "libavutil/opt.h"
28 #include "libavutil/pixdesc.h"
29 #include "libavutil/hwcontext.h"
30 #include "libavutil/hwcontext_cuda_internal.h"
31 #include "libavutil/cuda_check.h"
32 #include "libavutil/eval.h"
33 
34 #include "avfilter.h"
35 #include "filters.h"
36 #include "framesync.h"
37 
38 #include "cuda/load_helper.h"
39 
40  #define CHECK_CU(x) FF_CUDA_CHECK_DL(ctx, ctx->hwctx->internal->cuda_dl, x)
41  #define DIV_UP(a, b) ( ((a) + (b) - 1) / (b) )
42 
43  #define BLOCK_X 32
44  #define BLOCK_Y 16
45 
46  #define MAIN 0
47  #define OVERLAY 1
48 
49  static const enum AVPixelFormat supported_main_formats[] = {
50  AV_PIX_FMT_NV12,
51  AV_PIX_FMT_YUV420P,
52  AV_PIX_FMT_NONE,
53 };
54 
55  static const enum AVPixelFormat supported_overlay_formats[] = {
56  AV_PIX_FMT_NV12,
57  AV_PIX_FMT_YUV420P,
58  AV_PIX_FMT_YUVA420P,
59  AV_PIX_FMT_NONE,
60 };
61 
62  enum var_name {
63   VAR_MAIN_W, VAR_MW,
64   VAR_MAIN_H, VAR_MH,
65   VAR_OVERLAY_W, VAR_OW,
66   VAR_OVERLAY_H, VAR_OH,
67   VAR_X,
68   VAR_Y,
69   VAR_N,
70 #if FF_API_FRAME_PKT
71  VAR_POS,
72 #endif
73   VAR_T,
74   VAR_VARS_NB
75 };
76 
77  enum EvalMode {
78   EVAL_MODE_INIT,
79   EVAL_MODE_FRAME,
80   EVAL_MODE_NB
81 };
82 
83  static const char *const var_names[] = {
84  "main_w", "W", ///< width of the main video
85  "main_h", "H", ///< height of the main video
86  "overlay_w", "w", ///< width of the overlay video
87  "overlay_h", "h", ///< height of the overlay video
88  "x",
89  "y",
90  "n", ///< number of frame
91 #if FF_API_FRAME_PKT
92  "pos", ///< position in the file
93 #endif
94  "t", ///< timestamp expressed in seconds
95  NULL
96 };
97 
98 /**
99  * OverlayCUDAContext
100  */
101  typedef struct OverlayCUDAContext {
102   const AVClass *class;
103 
104   enum AVPixelFormat in_format_overlay;
105   enum AVPixelFormat in_format_main;
106 
107   AVBufferRef *hw_device_ctx;
108   AVCUDADeviceContext *hwctx;
109 
110   CUcontext cu_ctx;
111   CUmodule cu_module;
112   CUfunction cu_func;
113   CUstream cu_stream;
114 
115   FFFrameSync fs;
116 
117   int eval_mode;
118   int x_position;
119   int y_position;
120 
121   double var_values[VAR_VARS_NB];
122   char *x_expr, *y_expr;
123 
124   AVExpr *x_pexpr, *y_pexpr;
125 } OverlayCUDAContext;
126 
127 /**
128  * Helper to find out if provided format is supported by filter
129  */
130  static int format_is_supported(const enum AVPixelFormat formats[], enum AVPixelFormat fmt)
131 {
132  for (int i = 0; formats[i] != AV_PIX_FMT_NONE; i++)
133  if (formats[i] == fmt)
134  return 1;
135  return 0;
136 }
137 
138  static inline int normalize_xy(double d, int chroma_sub)
139 {
140  if (isnan(d))
141  return INT_MAX;
142  return (int)d & ~((1 << chroma_sub) - 1);
143 }
144 
145  static void eval_expr(AVFilterContext *ctx)
146 {
147  OverlayCUDAContext *s = ctx->priv;
148 
149  s->var_values[VAR_X] = av_expr_eval(s->x_pexpr, s->var_values, NULL);
150  s->var_values[VAR_Y] = av_expr_eval(s->y_pexpr, s->var_values, NULL);
151  /* necessary if x is expressed from y */
152  s->var_values[VAR_X] = av_expr_eval(s->x_pexpr, s->var_values, NULL);
153 
154  s->x_position = normalize_xy(s->var_values[VAR_X], 1);
155 
156  /* the cuda pixel format is using hwaccel, normalizing y is unnecessary */
157  s->y_position = s->var_values[VAR_Y];
158 }
159 
160  static int set_expr(AVExpr **pexpr, const char *expr, const char *option, void *log_ctx)
161 {
162  int ret;
163  AVExpr *old = NULL;
164 
165  if (*pexpr)
166  old = *pexpr;
167  ret = av_expr_parse(pexpr, expr, var_names,
168  NULL, NULL, NULL, NULL, 0, log_ctx);
169  if (ret < 0) {
170  av_log(log_ctx, AV_LOG_ERROR,
171  "Error when evaluating the expression '%s' for %s\n",
172  expr, option);
173  *pexpr = old;
174  return ret;
175  }
176 
177  av_expr_free(old);
178  return 0;
179 }
180 
181 /**
182  * Helper checks if we can process main and overlay pixel formats
183  */
184  static int formats_match(const enum AVPixelFormat format_main, const enum AVPixelFormat format_overlay) {
185  switch(format_main) {
186  case AV_PIX_FMT_NV12:
187  return format_overlay == AV_PIX_FMT_NV12;
188  case AV_PIX_FMT_YUV420P:
189  return format_overlay == AV_PIX_FMT_YUV420P ||
190  format_overlay == AV_PIX_FMT_YUVA420P;
191  default:
192  return 0;
193  }
194 }
195 
196 /**
197  * Call overlay kernell for a plane
198  */
199  static int overlay_cuda_call_kernel(
200  OverlayCUDAContext *ctx,
201  int x_position, int y_position,
202  uint8_t* main_data, int main_linesize,
203  int main_width, int main_height,
204  uint8_t* overlay_data, int overlay_linesize,
205  int overlay_width, int overlay_height,
206  uint8_t* alpha_data, int alpha_linesize,
207  int alpha_adj_x, int alpha_adj_y) {
208 
209  CudaFunctions *cu = ctx->hwctx->internal->cuda_dl;
210 
211  void* kernel_args[] = {
212  &x_position, &y_position,
213  &main_data, &main_linesize,
214  &overlay_data, &overlay_linesize,
215  &overlay_width, &overlay_height,
216  &alpha_data, &alpha_linesize,
217  &alpha_adj_x, &alpha_adj_y,
218  };
219 
220  return CHECK_CU(cu->cuLaunchKernel(
221  ctx->cu_func,
222  DIV_UP(main_width, BLOCK_X), DIV_UP(main_height, BLOCK_Y), 1,
223  BLOCK_X, BLOCK_Y, 1,
224  0, ctx->cu_stream, kernel_args, NULL));
225 }
226 
227 /**
228  * Perform blend overlay picture over main picture
229  */
230  static int overlay_cuda_blend(FFFrameSync *fs)
231 {
232  int ret;
233 
234  AVFilterContext *avctx = fs->parent;
235  OverlayCUDAContext *ctx = avctx->priv;
236  AVFilterLink *outlink = avctx->outputs[0];
237  AVFilterLink *inlink = avctx->inputs[0];
238  FilterLink *inl = ff_filter_link(inlink);
239 
240  CudaFunctions *cu = ctx->hwctx->internal->cuda_dl;
241  CUcontext dummy, cuda_ctx = ctx->hwctx->cuda_ctx;
242 
243  AVFrame *input_main, *input_overlay;
244 
245  ctx->cu_ctx = cuda_ctx;
246 
247  // read main and overlay frames from inputs
248  ret = ff_framesync_dualinput_get(fs, &input_main, &input_overlay);
249  if (ret < 0)
250  return ret;
251 
252  if (!input_main)
253  return AVERROR_BUG;
254 
255  if (!input_overlay)
256  return ff_filter_frame(outlink, input_main);
257 
258  ret = ff_inlink_make_frame_writable(inlink, &input_main);
259  if (ret < 0) {
260  av_frame_free(&input_main);
261  return ret;
262  }
263 
264  // push cuda context
265 
266  ret = CHECK_CU(cu->cuCtxPushCurrent(cuda_ctx));
267  if (ret < 0) {
268  av_frame_free(&input_main);
269  return ret;
270  }
271 
272  if (ctx->eval_mode == EVAL_MODE_FRAME) {
273  ctx->var_values[VAR_N] = inl->frame_count_out;
274  ctx->var_values[VAR_T] = input_main->pts == AV_NOPTS_VALUE ?
275  NAN : input_main->pts * av_q2d(inlink->time_base);
276 
277 #if FF_API_FRAME_PKT
278 FF_DISABLE_DEPRECATION_WARNINGS
279  {
280  int64_t pos = input_main->pkt_pos;
281  ctx->var_values[VAR_POS] = pos == -1 ? NAN : pos;
282  }
283 FF_ENABLE_DEPRECATION_WARNINGS
284 #endif
285 
286  ctx->var_values[VAR_OVERLAY_W] = ctx->var_values[VAR_OW] = input_overlay->width;
287  ctx->var_values[VAR_OVERLAY_H] = ctx->var_values[VAR_OH] = input_overlay->height;
288  ctx->var_values[VAR_MAIN_W ] = ctx->var_values[VAR_MW] = input_main->width;
289  ctx->var_values[VAR_MAIN_H ] = ctx->var_values[VAR_MH] = input_main->height;
290 
291  eval_expr(avctx);
292 
293  av_log(avctx, AV_LOG_DEBUG, "n:%f t:%f x:%f xi:%d y:%f yi:%d\n",
294  ctx->var_values[VAR_N], ctx->var_values[VAR_T],
295  ctx->var_values[VAR_X], ctx->x_position,
296  ctx->var_values[VAR_Y], ctx->y_position);
297  }
298 
299  // overlay first plane
300 
301  overlay_cuda_call_kernel(ctx,
302  ctx->x_position, ctx->y_position,
303  input_main->data[0], input_main->linesize[0],
304  input_main->width, input_main->height,
305  input_overlay->data[0], input_overlay->linesize[0],
306  input_overlay->width, input_overlay->height,
307  input_overlay->data[3], input_overlay->linesize[3], 1, 1);
308 
309  // overlay rest planes depending on pixel format
310 
311  switch(ctx->in_format_overlay) {
312  case AV_PIX_FMT_NV12:
313  overlay_cuda_call_kernel(ctx,
314  ctx->x_position, ctx->y_position / 2,
315  input_main->data[1], input_main->linesize[1],
316  input_main->width, input_main->height / 2,
317  input_overlay->data[1], input_overlay->linesize[1],
318  input_overlay->width, input_overlay->height / 2,
319  0, 0, 0, 0);
320  break;
321  case AV_PIX_FMT_YUV420P:
322  case AV_PIX_FMT_YUVA420P:
323  overlay_cuda_call_kernel(ctx,
324  ctx->x_position / 2 , ctx->y_position / 2,
325  input_main->data[1], input_main->linesize[1],
326  input_main->width / 2, input_main->height / 2,
327  input_overlay->data[1], input_overlay->linesize[1],
328  input_overlay->width / 2, input_overlay->height / 2,
329  input_overlay->data[3], input_overlay->linesize[3], 2, 2);
330 
331  overlay_cuda_call_kernel(ctx,
332  ctx->x_position / 2 , ctx->y_position / 2,
333  input_main->data[2], input_main->linesize[2],
334  input_main->width / 2, input_main->height / 2,
335  input_overlay->data[2], input_overlay->linesize[2],
336  input_overlay->width / 2, input_overlay->height / 2,
337  input_overlay->data[3], input_overlay->linesize[3], 2, 2);
338  break;
339  default:
340  av_log(ctx, AV_LOG_ERROR, "Passed unsupported overlay pixel format\n");
341  av_frame_free(&input_main);
342  CHECK_CU(cu->cuCtxPopCurrent(&dummy));
343  return AVERROR_BUG;
344  }
345 
346  CHECK_CU(cu->cuCtxPopCurrent(&dummy));
347 
348  return ff_filter_frame(outlink, input_main);
349 }
350 
351  static int config_input_overlay(AVFilterLink *inlink)
352 {
353  AVFilterContext *ctx = inlink->dst;
354  OverlayCUDAContext *s = inlink->dst->priv;
355  int ret;
356 
357 
358  /* Finish the configuration by evaluating the expressions
359  now when both inputs are configured. */
360  s->var_values[VAR_MAIN_W ] = s->var_values[VAR_MW] = ctx->inputs[MAIN ]->w;
361  s->var_values[VAR_MAIN_H ] = s->var_values[VAR_MH] = ctx->inputs[MAIN ]->h;
362  s->var_values[VAR_OVERLAY_W] = s->var_values[VAR_OW] = ctx->inputs[OVERLAY]->w;
363  s->var_values[VAR_OVERLAY_H] = s->var_values[VAR_OH] = ctx->inputs[OVERLAY]->h;
364  s->var_values[VAR_X] = NAN;
365  s->var_values[VAR_Y] = NAN;
366  s->var_values[VAR_N] = 0;
367  s->var_values[VAR_T] = NAN;
368 #if FF_API_FRAME_PKT
369  s->var_values[VAR_POS] = NAN;
370 #endif
371 
372  if ((ret = set_expr(&s->x_pexpr, s->x_expr, "x", ctx)) < 0 ||
373  (ret = set_expr(&s->y_pexpr, s->y_expr, "y", ctx)) < 0)
374  return ret;
375 
376  if (s->eval_mode == EVAL_MODE_INIT) {
377  eval_expr(ctx);
378  av_log(ctx, AV_LOG_VERBOSE, "x:%f xi:%d y:%f yi:%d\n",
379  s->var_values[VAR_X], s->x_position,
380  s->var_values[VAR_Y], s->y_position);
381  }
382 
383  return 0;
384 }
385 
386 /**
387  * Initialize overlay_cuda
388  */
389  static av_cold int overlay_cuda_init(AVFilterContext *avctx)
390 {
391  OverlayCUDAContext* ctx = avctx->priv;
392  ctx->fs.on_event = &overlay_cuda_blend;
393 
394  return 0;
395 }
396 
397 /**
398  * Uninitialize overlay_cuda
399  */
400  static av_cold void overlay_cuda_uninit(AVFilterContext *avctx)
401 {
402  OverlayCUDAContext* ctx = avctx->priv;
403 
404  ff_framesync_uninit(&ctx->fs);
405 
406  if (ctx->hwctx && ctx->cu_module) {
407  CUcontext dummy;
408  CudaFunctions *cu = ctx->hwctx->internal->cuda_dl;
409  CHECK_CU(cu->cuCtxPushCurrent(ctx->cu_ctx));
410  CHECK_CU(cu->cuModuleUnload(ctx->cu_module));
411  CHECK_CU(cu->cuCtxPopCurrent(&dummy));
412  }
413 
414  av_expr_free(ctx->x_pexpr); ctx->x_pexpr = NULL;
415  av_expr_free(ctx->y_pexpr); ctx->y_pexpr = NULL;
416  av_buffer_unref(&ctx->hw_device_ctx);
417  ctx->hwctx = NULL;
418 }
419 
420 /**
421  * Activate overlay_cuda
422  */
423  static int overlay_cuda_activate(AVFilterContext *avctx)
424 {
425  OverlayCUDAContext *ctx = avctx->priv;
426 
427  return ff_framesync_activate(&ctx->fs);
428 }
429 
430 /**
431  * Configure output
432  */
433  static int overlay_cuda_config_output(AVFilterLink *outlink)
434 {
435  extern const unsigned char ff_vf_overlay_cuda_ptx_data[];
436  extern const unsigned int ff_vf_overlay_cuda_ptx_len;
437 
438  int err;
439  FilterLink *outl = ff_filter_link(outlink);
440  AVFilterContext* avctx = outlink->src;
441  OverlayCUDAContext* ctx = avctx->priv;
442 
443  AVFilterLink *inlink = avctx->inputs[0];
444  FilterLink *inl = ff_filter_link(inlink);
445  AVHWFramesContext *frames_ctx = (AVHWFramesContext*)inl->hw_frames_ctx->data;
446 
447  AVFilterLink *inlink_overlay = avctx->inputs[1];
448  FilterLink *inl_overlay = ff_filter_link(inlink_overlay);
449  AVHWFramesContext *frames_ctx_overlay = (AVHWFramesContext*)inl_overlay->hw_frames_ctx->data;
450 
451  CUcontext dummy, cuda_ctx;
452  CudaFunctions *cu;
453 
454  // check main input formats
455 
456  if (!frames_ctx) {
457  av_log(ctx, AV_LOG_ERROR, "No hw context provided on main input\n");
458  return AVERROR(EINVAL);
459  }
460 
461  ctx->in_format_main = frames_ctx->sw_format;
462  if (!format_is_supported(supported_main_formats, ctx->in_format_main)) {
463  av_log(ctx, AV_LOG_ERROR, "Unsupported main input format: %s\n",
464  av_get_pix_fmt_name(ctx->in_format_main));
465  return AVERROR(ENOSYS);
466  }
467 
468  // check overlay input formats
469 
470  if (!frames_ctx_overlay) {
471  av_log(ctx, AV_LOG_ERROR, "No hw context provided on overlay input\n");
472  return AVERROR(EINVAL);
473  }
474 
475  ctx->in_format_overlay = frames_ctx_overlay->sw_format;
476  if (!format_is_supported(supported_overlay_formats, ctx->in_format_overlay)) {
477  av_log(ctx, AV_LOG_ERROR, "Unsupported overlay input format: %s\n",
478  av_get_pix_fmt_name(ctx->in_format_overlay));
479  return AVERROR(ENOSYS);
480  }
481 
482  // check we can overlay pictures with those pixel formats
483 
484  if (!formats_match(ctx->in_format_main, ctx->in_format_overlay)) {
485  av_log(ctx, AV_LOG_ERROR, "Can't overlay %s on %s \n",
486  av_get_pix_fmt_name(ctx->in_format_overlay), av_get_pix_fmt_name(ctx->in_format_main));
487  return AVERROR(EINVAL);
488  }
489 
490  // initialize
491 
492  ctx->hw_device_ctx = av_buffer_ref(frames_ctx->device_ref);
493  if (!ctx->hw_device_ctx)
494  return AVERROR(ENOMEM);
495  ctx->hwctx = ((AVHWDeviceContext*)ctx->hw_device_ctx->data)->hwctx;
496 
497  cuda_ctx = ctx->hwctx->cuda_ctx;
498  ctx->fs.time_base = inlink->time_base;
499 
500  ctx->cu_stream = ctx->hwctx->stream;
501 
502  outl->hw_frames_ctx = av_buffer_ref(inl->hw_frames_ctx);
503  if (!outl->hw_frames_ctx)
504  return AVERROR(ENOMEM);
505 
506  // load functions
507 
508  cu = ctx->hwctx->internal->cuda_dl;
509 
510  err = CHECK_CU(cu->cuCtxPushCurrent(cuda_ctx));
511  if (err < 0) {
512  return err;
513  }
514 
515  err = ff_cuda_load_module(ctx, ctx->hwctx, &ctx->cu_module, ff_vf_overlay_cuda_ptx_data, ff_vf_overlay_cuda_ptx_len);
516  if (err < 0) {
517  CHECK_CU(cu->cuCtxPopCurrent(&dummy));
518  return err;
519  }
520 
521  err = CHECK_CU(cu->cuModuleGetFunction(&ctx->cu_func, ctx->cu_module, "Overlay_Cuda"));
522  if (err < 0) {
523  CHECK_CU(cu->cuCtxPopCurrent(&dummy));
524  return err;
525  }
526 
527  CHECK_CU(cu->cuCtxPopCurrent(&dummy));
528 
529  // init dual input
530 
531  err = ff_framesync_init_dualinput(&ctx->fs, avctx);
532  if (err < 0) {
533  return err;
534  }
535 
536  return ff_framesync_configure(&ctx->fs);
537 }
538 
539 
540  #define OFFSET(x) offsetof(OverlayCUDAContext, x)
541  #define FLAGS (AV_OPT_FLAG_FILTERING_PARAM | AV_OPT_FLAG_VIDEO_PARAM)
542 
543  static const AVOption overlay_cuda_options[] = {
544  { "x", "set the x expression of overlay", OFFSET(x_expr), AV_OPT_TYPE_STRING, { .str = "0" }, 0, 0, FLAGS },
545  { "y", "set the y expression of overlay", OFFSET(y_expr), AV_OPT_TYPE_STRING, { .str = "0" }, 0, 0, FLAGS },
546  { "eof_action", "Action to take when encountering EOF from secondary input ",
547  OFFSET(fs.opt_eof_action), AV_OPT_TYPE_INT, { .i64 = EOF_ACTION_REPEAT },
548  EOF_ACTION_REPEAT, EOF_ACTION_PASS, .flags = FLAGS, .unit = "eof_action" },
549  { "repeat", "Repeat the previous frame.", 0, AV_OPT_TYPE_CONST, { .i64 = EOF_ACTION_REPEAT }, .flags = FLAGS, .unit = "eof_action" },
550  { "endall", "End both streams.", 0, AV_OPT_TYPE_CONST, { .i64 = EOF_ACTION_ENDALL }, .flags = FLAGS, .unit = "eof_action" },
551  { "pass", "Pass through the main input.", 0, AV_OPT_TYPE_CONST, { .i64 = EOF_ACTION_PASS }, .flags = FLAGS, .unit = "eof_action" },
552  { "eval", "specify when to evaluate expressions", OFFSET(eval_mode), AV_OPT_TYPE_INT, { .i64 = EVAL_MODE_FRAME }, 0, EVAL_MODE_NB - 1, FLAGS, .unit = "eval" },
553  { "init", "eval expressions once during initialization", 0, AV_OPT_TYPE_CONST, { .i64=EVAL_MODE_INIT }, .flags = FLAGS, .unit = "eval" },
554  { "frame", "eval expressions per-frame", 0, AV_OPT_TYPE_CONST, { .i64=EVAL_MODE_FRAME }, .flags = FLAGS, .unit = "eval" },
555  { "shortest", "force termination when the shortest input terminates", OFFSET(fs.opt_shortest), AV_OPT_TYPE_BOOL, { .i64 = 0 }, 0, 1, FLAGS },
556  { "repeatlast", "repeat overlay of the last overlay frame", OFFSET(fs.opt_repeatlast), AV_OPT_TYPE_BOOL, {.i64=1}, 0, 1, FLAGS },
557  { NULL },
558 };
559 
560 FRAMESYNC_DEFINE_CLASS(overlay_cuda, OverlayCUDAContext, fs);
561 
562  static const AVFilterPad overlay_cuda_inputs[] = {
563  {
564  .name = "main",
565  .type = AVMEDIA_TYPE_VIDEO,
566  },
567  {
568  .name = "overlay",
569  .type = AVMEDIA_TYPE_VIDEO,
570  .config_props = config_input_overlay,
571  },
572 };
573 
574  static const AVFilterPad overlay_cuda_outputs[] = {
575  {
576  .name = "default",
577  .type = AVMEDIA_TYPE_VIDEO,
578  .config_props = &overlay_cuda_config_output,
579  },
580 };
581 
582  const AVFilter ff_vf_overlay_cuda = {
583  .name = "overlay_cuda",
584  .description = NULL_IF_CONFIG_SMALL("Overlay one video on top of another using CUDA"),
585  .priv_size = sizeof(OverlayCUDAContext),
586  .priv_class = &overlay_cuda_class,
587  .init = &overlay_cuda_init,
588  .uninit = &overlay_cuda_uninit,
589  .activate = &overlay_cuda_activate,
590  FILTER_INPUTS(overlay_cuda_inputs),
591  FILTER_OUTPUTS(overlay_cuda_outputs),
592  FILTER_SINGLE_PIXFMT(AV_PIX_FMT_CUDA),
593  .preinit = overlay_cuda_framesync_preinit,
594  .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE,
595 };
formats
formats
Definition: signature.h:47
FF_ENABLE_DEPRECATION_WARNINGS
#define FF_ENABLE_DEPRECATION_WARNINGS
Definition: internal.h:73
ff_framesync_configure
int ff_framesync_configure(FFFrameSync *fs)
Configure a frame sync structure.
Definition: framesync.c:137
AV_PIX_FMT_CUDA
@ AV_PIX_FMT_CUDA
HW acceleration through CUDA.
Definition: pixfmt.h:260
AVPixelFormat
AVPixelFormat
Pixel format.
Definition: pixfmt.h:71
formats_match
static int formats_match(const enum AVPixelFormat format_main, const enum AVPixelFormat format_overlay)
Helper checks if we can process main and overlay pixel formats.
Definition: vf_overlay_cuda.c:184
AVERROR
Filter the word "frame" indicates either a video frame or a group of audio as stored in an AVFrame structure Format for each input and each output the list of supported formats For video that means pixel format For audio that means channel sample they are references to shared objects When the negotiation mechanism computes the intersection of the formats supported at each end of a all references to both lists are replaced with a reference to the intersection And when a single format is eventually chosen for a link amongst the remaining all references to the list are updated That means that if a filter requires that its input and output have the same format amongst a supported all it has to do is use a reference to the same list of formats query_formats can leave some formats unset and return AVERROR(EAGAIN) to cause the negotiation mechanism toagain later. That can be used by filters with complex requirements to use the format negotiated on one link to set the formats supported on another. Frame references ownership and permissions
opt.h
overlay_cuda_call_kernel
static int overlay_cuda_call_kernel(OverlayCUDAContext *ctx, int x_position, int y_position, uint8_t *main_data, int main_linesize, int main_width, int main_height, uint8_t *overlay_data, int overlay_linesize, int overlay_width, int overlay_height, uint8_t *alpha_data, int alpha_linesize, int alpha_adj_x, int alpha_adj_y)
Call overlay kernell for a plane.
Definition: vf_overlay_cuda.c:199
var_name
var_name
Definition: noise.c:47
hwcontext_cuda_internal.h
ff_framesync_uninit
void ff_framesync_uninit(FFFrameSync *fs)
Free all memory currently allocated.
Definition: framesync.c:301
ff_filter_frame
int ff_filter_frame(AVFilterLink *link, AVFrame *frame)
Send a frame of data to the next filter.
Definition: avfilter.c:1023
AVBufferRef::data
uint8_t * data
The data buffer.
Definition: buffer.h:90
overlay_cuda_inputs
static const AVFilterPad overlay_cuda_inputs[]
Definition: vf_overlay_cuda.c:562
OverlayCUDAContext::y_position
int y_position
Definition: vf_overlay_cuda.c:119
OverlayCUDAContext::cu_ctx
CUcontext cu_ctx
Definition: vf_overlay_cuda.c:110
EVAL_MODE_INIT
@ EVAL_MODE_INIT
Definition: vf_overlay_cuda.c:78
ff_cuda_load_module
int ff_cuda_load_module(void *avctx, AVCUDADeviceContext *hwctx, CUmodule *cu_module, const unsigned char *data, const unsigned int length)
Loads a CUDA module and applies any decompression, if necessary.
Definition: load_helper.c:35
int64_t
long long int64_t
Definition: coverity.c:34
inlink
The exact code depends on how similar the blocks are and how related they are to the and needs to apply these operations to the correct inlink or outlink if there are several Macros are available to factor that when no extra processing is inlink
Definition: filter_design.txt:212
av_frame_free
void av_frame_free(AVFrame **frame)
Free the frame and any dynamically allocated objects in it, e.g.
Definition: frame.c:162
FILTER_INPUTS
#define FILTER_INPUTS(array)
Definition: filters.h:262
AVFrame
This structure describes decoded (raw) audio or video data.
Definition: frame.h:389
pixdesc.h
AVFrame::pts
int64_t pts
Presentation timestamp in time_base units (time when frame should be shown to user).
Definition: frame.h:501
AVFrame::width
int width
Definition: frame.h:461
OverlayCUDAContext::hw_device_ctx
AVBufferRef * hw_device_ctx
Definition: vf_overlay_cuda.c:107
OverlayCUDAContext::in_format_overlay
enum AVPixelFormat in_format_overlay
Definition: vf_overlay_cuda.c:104
AVOption
AVOption.
Definition: opt.h:429
normalize_xy
static int normalize_xy(double d, int chroma_sub)
Definition: vf_overlay_cuda.c:138
EOF_ACTION_ENDALL
@ EOF_ACTION_ENDALL
Definition: framesync.h:28
overlay_cuda_outputs
static const AVFilterPad overlay_cuda_outputs[]
Definition: vf_overlay_cuda.c:574
AV_LOG_VERBOSE
#define AV_LOG_VERBOSE
Detailed information.
Definition: log.h:196
OverlayCUDAContext::in_format_main
enum AVPixelFormat in_format_main
Definition: vf_overlay_cuda.c:105
OVERLAY
#define OVERLAY
Definition: vf_overlay_cuda.c:47
av_buffer_ref
AVBufferRef * av_buffer_ref(const AVBufferRef *buf)
Create a new reference to an AVBuffer.
Definition: buffer.c:103
AVFilter::name
const char * name
Filter name.
Definition: avfilter.h:205
FFFrameSync
Frame sync structure.
Definition: framesync.h:168
AVFrame::data
uint8_t * data[AV_NUM_DATA_POINTERS]
pointer to the picture/channel planes.
Definition: frame.h:410
av_expr_parse
int av_expr_parse(AVExpr **expr, const char *s, const char *const *const_names, const char *const *func1_names, double(*const *funcs1)(void *, double), const char *const *func2_names, double(*const *funcs2)(void *, double, double), int log_offset, void *log_ctx)
Parse an expression.
Definition: eval.c:710
ff_vf_overlay_cuda
const AVFilter ff_vf_overlay_cuda
Definition: vf_overlay_cuda.c:582
AVFilterContext::priv
void * priv
private data for use by the filter
Definition: avfilter.h:472
VAR_MW
@ VAR_MW
Definition: vf_overlay_cuda.c:63
dummy
int dummy
Definition: motion.c:66
OverlayCUDAContext::fs
FFFrameSync fs
Definition: vf_overlay_cuda.c:115
av_expr_free
void av_expr_free(AVExpr *e)
Free a parsed expression previously created with av_expr_parse().
Definition: eval.c:358
AVFilterPad
A filter pad used for either input or output.
Definition: filters.h:38
AVHWDeviceContext
This struct aggregates all the (hardware/vendor-specific) "high-level" state, i.e.
Definition: hwcontext.h:60
AV_LOG_ERROR
#define AV_LOG_ERROR
Something went wrong and cannot losslessly be recovered.
Definition: log.h:180
av_cold
#define av_cold
Definition: attributes.h:90
s
#define s(width, name)
Definition: cbs_vp9.c:198
AV_PIX_FMT_YUVA420P
@ AV_PIX_FMT_YUVA420P
planar YUV 4:2:0, 20bpp, (1 Cr & Cb sample per 2x2 Y & A samples)
Definition: pixfmt.h:108
av_q2d
static double av_q2d(AVRational a)
Convert an AVRational to a double.
Definition: rational.h:104
OverlayCUDAContext::x_position
int x_position
Definition: vf_overlay_cuda.c:118
OFFSET
#define OFFSET(x)
Definition: vf_overlay_cuda.c:540
filters.h
AV_LOG_DEBUG
#define AV_LOG_DEBUG
Stuff which is only useful for libav* developers.
Definition: log.h:201
ctx
AVFormatContext * ctx
Definition: movenc.c:49
av_expr_eval
double av_expr_eval(AVExpr *e, const double *const_values, void *opaque)
Evaluate a previously parsed expression.
Definition: eval.c:792
AVExpr
Definition: eval.c:158
VAR_OH
@ VAR_OH
Definition: vf_overlay_cuda.c:66
load_helper.h
AV_PIX_FMT_YUV420P
@ AV_PIX_FMT_YUV420P
planar YUV 4:2:0, 12bpp, (1 Cr & Cb sample per 2x2 Y samples)
Definition: pixfmt.h:73
EOF_ACTION_PASS
@ EOF_ACTION_PASS
Definition: framesync.h:29
FILTER_OUTPUTS
#define FILTER_OUTPUTS(array)
Definition: filters.h:263
overlay_cuda_config_output
static int overlay_cuda_config_output(AVFilterLink *outlink)
Configure output.
Definition: vf_overlay_cuda.c:433
NAN
#define NAN
Definition: mathematics.h:115
ff_inlink_make_frame_writable
int ff_inlink_make_frame_writable(AVFilterLink *link, AVFrame **rframe)
Make sure a frame is writable.
Definition: avfilter.c:1498
if
if(ret)
Definition: filter_design.txt:179
option
option
Definition: libkvazaar.c:320
CHECK_CU
#define CHECK_CU(x)
Definition: vf_overlay_cuda.c:40
set_expr
static int set_expr(AVExpr **pexpr, const char *expr, const char *option, void *log_ctx)
Definition: vf_overlay_cuda.c:160
config_input_overlay
static int config_input_overlay(AVFilterLink *inlink)
Definition: vf_overlay_cuda.c:351
AVClass
Describe the class of an AVClass context structure.
Definition: log.h:66
BLOCK_Y
#define BLOCK_Y
Definition: vf_overlay_cuda.c:44
NULL
#define NULL
Definition: coverity.c:32
AVHWFramesContext::sw_format
enum AVPixelFormat sw_format
The pixel format identifying the actual data layout of the hardware frames.
Definition: hwcontext.h:210
VAR_POS
@ VAR_POS
Definition: noise.c:56
OverlayCUDAContext::x_expr
char * x_expr
Definition: vf_overlay_cuda.c:122
av_buffer_unref
void av_buffer_unref(AVBufferRef **buf)
Free a given reference and automatically free the buffer if there are no more references to it.
Definition: buffer.c:139
EVAL_MODE_NB
@ EVAL_MODE_NB
Definition: vf_overlay_cuda.c:80
fs
#define fs(width, name, subs,...)
Definition: cbs_vp9.c:200
AVHWFramesContext::device_ref
AVBufferRef * device_ref
A reference to the parent AVHWDeviceContext.
Definition: hwcontext.h:126
isnan
#define isnan(x)
Definition: libm.h:340
activate
filter_frame For filters that do not use the activate() callback
AVFilterContext::inputs
AVFilterLink ** inputs
array of pointers to input links
Definition: avfilter.h:465
MAIN
#define MAIN
Definition: vf_overlay_cuda.c:46
ff_filter_link
static FilterLink * ff_filter_link(AVFilterLink *link)
Definition: filters.h:197
FF_FILTER_FLAG_HWFRAME_AWARE
#define FF_FILTER_FLAG_HWFRAME_AWARE
The filter is aware of hardware frames, and any hardware frame context should not be automatically pr...
Definition: filters.h:206
overlay_cuda_init
static av_cold int overlay_cuda_init(AVFilterContext *avctx)
Initialize overlay_cuda.
Definition: vf_overlay_cuda.c:389
eval.h
FRAMESYNC_DEFINE_CLASS
FRAMESYNC_DEFINE_CLASS(overlay_cuda, OverlayCUDAContext, fs)
init
int(* init)(AVBSFContext *ctx)
Definition: dts2pts.c:368
NULL_IF_CONFIG_SMALL
#define NULL_IF_CONFIG_SMALL(x)
Return NULL if CONFIG_SMALL is true, otherwise the argument without modification.
Definition: internal.h:94
ff_framesync_init_dualinput
int ff_framesync_init_dualinput(FFFrameSync *fs, AVFilterContext *parent)
Initialize a frame sync structure for dualinput.
Definition: framesync.c:372
VAR_Y
@ VAR_Y
Definition: vf_overlay_cuda.c:68
OverlayCUDAContext::x_pexpr
AVExpr * x_pexpr
Definition: vf_overlay_cuda.c:124
VAR_MAIN_H
@ VAR_MAIN_H
Definition: vf_overlay_cuda.c:64
AV_NOPTS_VALUE
#define AV_NOPTS_VALUE
Undefined timestamp value.
Definition: avutil.h:248
FLAGS
#define FLAGS
Definition: vf_overlay_cuda.c:541
overlay_cuda_options
static const AVOption overlay_cuda_options[]
Definition: vf_overlay_cuda.c:543
VAR_T
@ VAR_T
Definition: vf_overlay_cuda.c:73
AVFrame::pkt_pos
attribute_deprecated int64_t pkt_pos
reordered pos from the last AVPacket that has been input into the decoder
Definition: frame.h:699
EVAL_MODE_FRAME
@ EVAL_MODE_FRAME
Definition: vf_overlay_cuda.c:79
OverlayCUDAContext
OverlayCUDAContext.
Definition: vf_overlay_cuda.c:101
overlay_cuda_uninit
static av_cold void overlay_cuda_uninit(AVFilterContext *avctx)
Uninitialize overlay_cuda.
Definition: vf_overlay_cuda.c:400
format_is_supported
static int format_is_supported(const enum AVPixelFormat formats[], enum AVPixelFormat fmt)
Helper to find out if provided format is supported by filter.
Definition: vf_overlay_cuda.c:130
VAR_X
@ VAR_X
Definition: vf_overlay_cuda.c:67
uninit
static void uninit(AVBSFContext *ctx)
Definition: pcm_rechunk.c:68
overlay_cuda_blend
static int overlay_cuda_blend(FFFrameSync *fs)
Perform blend overlay picture over main picture.
Definition: vf_overlay_cuda.c:230
log.h
i
#define i(width, name, range_min, range_max)
Definition: cbs_h2645.c:256
VAR_OVERLAY_H
@ VAR_OVERLAY_H
Definition: vf_overlay_cuda.c:66
OverlayCUDAContext::hwctx
AVCUDADeviceContext * hwctx
Definition: vf_overlay_cuda.c:108
EvalMode
EvalMode
Definition: af_volume.h:39
BLOCK_X
#define BLOCK_X
Definition: vf_overlay_cuda.c:43
OverlayCUDAContext::cu_stream
CUstream cu_stream
Definition: vf_overlay_cuda.c:113
OverlayCUDAContext::cu_func
CUfunction cu_func
Definition: vf_overlay_cuda.c:112
AVFilterPad::name
const char * name
Pad name.
Definition: filters.h:44
AVFilter
Filter definition.
Definition: avfilter.h:201
AVHWFramesContext
This struct describes a set or pool of "hardware" frames (i.e.
Definition: hwcontext.h:115
AVCUDADeviceContext
This struct is allocated as AVHWDeviceContext.hwctx.
Definition: hwcontext_cuda.h:42
supported_overlay_formats
static enum AVPixelFormat supported_overlay_formats[]
Definition: vf_overlay_cuda.c:55
ret
ret
Definition: filter_design.txt:187
AV_PIX_FMT_NV12
@ AV_PIX_FMT_NV12
planar YUV 4:2:0, 12bpp, 1 plane for Y and 1 plane for the UV components, which are interleaved (firs...
Definition: pixfmt.h:96
cuda_check.h
OverlayCUDAContext::cu_module
CUmodule cu_module
Definition: vf_overlay_cuda.c:111
pos
unsigned int pos
Definition: spdifenc.c:414
EOF_ACTION_REPEAT
@ EOF_ACTION_REPEAT
Definition: framesync.h:27
AVFrame::height
int height
Definition: frame.h:461
OverlayCUDAContext::y_pexpr
AVExpr * y_pexpr
Definition: vf_overlay_cuda.c:124
supported_main_formats
static enum AVPixelFormat supported_main_formats[]
Definition: vf_overlay_cuda.c:49
framesync.h
DIV_UP
#define DIV_UP(a, b)
Definition: vf_overlay_cuda.c:41
VAR_MH
@ VAR_MH
Definition: vf_overlay_cuda.c:64
VAR_OW
@ VAR_OW
Definition: vf_overlay_cuda.c:65
OverlayCUDAContext::y_expr
char * y_expr
Definition: vf_overlay_cuda.c:122
AV_PIX_FMT_NONE
@ AV_PIX_FMT_NONE
Definition: pixfmt.h:72
AV_OPT_TYPE_INT
@ AV_OPT_TYPE_INT
Underlying C type is int.
Definition: opt.h:259
avfilter.h
VAR_OVERLAY_W
@ VAR_OVERLAY_W
Definition: vf_overlay_cuda.c:65
AVFilterContext
An instance of a filter.
Definition: avfilter.h:457
FF_DISABLE_DEPRECATION_WARNINGS
#define FF_DISABLE_DEPRECATION_WARNINGS
Definition: internal.h:72
var_names
static const char *const var_names[]
Definition: vf_overlay_cuda.c:83
AVMEDIA_TYPE_VIDEO
@ AVMEDIA_TYPE_VIDEO
Definition: avutil.h:201
AVBufferRef
A reference to a data buffer.
Definition: buffer.h:82
VAR_MAIN_W
@ VAR_MAIN_W
Definition: vf_overlay_cuda.c:63
overlay_cuda_activate
static int overlay_cuda_activate(AVFilterContext *avctx)
Activate overlay_cuda.
Definition: vf_overlay_cuda.c:423
AV_OPT_TYPE_BOOL
@ AV_OPT_TYPE_BOOL
Underlying C type is int.
Definition: opt.h:327
VAR_VARS_NB
@ VAR_VARS_NB
Definition: vf_overlay_cuda.c:74
VAR_N
@ VAR_N
Definition: vf_overlay_cuda.c:69
hwcontext.h
AVERROR_BUG
#define AVERROR_BUG
Internal bug, also see AVERROR_BUG2.
Definition: error.h:52
AVFrame::linesize
int linesize[AV_NUM_DATA_POINTERS]
For video, a positive or negative value, which is typically indicating the size in bytes of each pict...
Definition: frame.h:434
av_log
#define av_log(a,...)
Definition: tableprint_vlc.h:27
eval_expr
static void eval_expr(AVFilterContext *ctx)
Definition: vf_overlay_cuda.c:145
ff_framesync_activate
int ff_framesync_activate(FFFrameSync *fs)
Examine the frames in the filter's input and try to produce output.
Definition: framesync.c:352
ff_framesync_dualinput_get
int ff_framesync_dualinput_get(FFFrameSync *fs, AVFrame **f0, AVFrame **f1)
Definition: framesync.c:390
AV_OPT_TYPE_STRING
@ AV_OPT_TYPE_STRING
Underlying C type is a uint8_t* that is either NULL or points to a C string allocated with the av_mal...
Definition: opt.h:276
OverlayCUDAContext::eval_mode
int eval_mode
Definition: vf_overlay_cuda.c:117
AV_OPT_TYPE_CONST
@ AV_OPT_TYPE_CONST
Special option type for declaring named constants.
Definition: opt.h:299
FILTER_SINGLE_PIXFMT
#define FILTER_SINGLE_PIXFMT(pix_fmt_)
Definition: filters.h:252
av_get_pix_fmt_name
const char * av_get_pix_fmt_name(enum AVPixelFormat pix_fmt)
Return the short name for a pixel format, NULL in case pix_fmt is unknown.
Definition: pixdesc.c:2885
AVFilterContext::outputs
AVFilterLink ** outputs
array of pointers to output links
Definition: avfilter.h:469
OverlayCUDAContext::var_values
double var_values[VAR_VARS_NB]
Definition: vf_overlay_cuda.c:121
Generated on Fri Aug 22 2025 13:59:18 for FFmpeg by   doxygen 1.8.17

AltStyle によって変換されたページ (->オリジナル) /