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"
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 #include "internal.h"
38 
39 #include "cuda/load_helper.h"
40 
41 #define CHECK_CU(x) FF_CUDA_CHECK_DL(ctx, ctx->hwctx->internal->cuda_dl, x)
42 #define DIV_UP(a, b) ( ((a) + (b) - 1) / (b) )
43 
44 #define BLOCK_X 32
45 #define BLOCK_Y 16
46 
47 #define MAIN 0
48 #define OVERLAY 1
49 
50 static const enum AVPixelFormat supported_main_formats[] = {
54 };
55 
61 };
62 
63 enum var_name {
71 #if FF_API_FRAME_PKT
72  VAR_POS,
73 #endif
76 };
77 
78 enum EvalMode {
82 };
83 
84 static const char *const var_names[] = {
85  "main_w", "W", ///< width of the main video
86  "main_h", "H", ///< height of the main video
87  "overlay_w", "w", ///< width of the overlay video
88  "overlay_h", "h", ///< height of the overlay video
89  "x",
90  "y",
91  "n", ///< number of frame
92 #if FF_API_FRAME_PKT
93  "pos", ///< position in the file
94 #endif
95  "t", ///< timestamp expressed in seconds
96  NULL
97 };
98 
99 /**
100  * OverlayCUDAContext
101  */
102 typedef struct OverlayCUDAContext {
103  const AVClass *class;
104 
107 
110 
111  CUcontext cu_ctx;
112  CUmodule cu_module;
113  CUfunction cu_func;
114  CUstream cu_stream;
115 
117 
121 
123  char *x_expr, *y_expr;
124 
127 
128 /**
129  * Helper to find out if provided format is supported by filter
130  */
131 static int format_is_supported(const enum AVPixelFormat formats[], enum AVPixelFormat fmt)
132 {
133  for (int i = 0; formats[i] != AV_PIX_FMT_NONE; i++)
134  if (formats[i] == fmt)
135  return 1;
136  return 0;
137 }
138 
139 static inline int normalize_xy(double d, int chroma_sub)
140 {
141  if (isnan(d))
142  return INT_MAX;
143  return (int)d & ~((1 << chroma_sub) - 1);
144 }
145 
147 {
148  OverlayCUDAContext *s = ctx->priv;
149 
150  s->var_values[VAR_X] = av_expr_eval(s->x_pexpr, s->var_values, NULL);
151  s->var_values[VAR_Y] = av_expr_eval(s->y_pexpr, s->var_values, NULL);
152  /* necessary if x is expressed from y */
153  s->var_values[VAR_X] = av_expr_eval(s->x_pexpr, s->var_values, NULL);
154 
155  s->x_position = normalize_xy(s->var_values[VAR_X], 1);
156 
157  /* the cuda pixel format is using hwaccel, normalizing y is unnecessary */
158  s->y_position = s->var_values[VAR_Y];
159 }
160 
161 static int set_expr(AVExpr **pexpr, const char *expr, const char *option, void *log_ctx)
162 {
163  int ret;
164  AVExpr *old = NULL;
165 
166  if (*pexpr)
167  old = *pexpr;
168  ret = av_expr_parse(pexpr, expr, var_names,
169  NULL, NULL, NULL, NULL, 0, log_ctx);
170  if (ret < 0) {
171  av_log(log_ctx, AV_LOG_ERROR,
172  "Error when evaluating the expression '%s' for %s\n",
173  expr, option);
174  *pexpr = old;
175  return ret;
176  }
177 
178  av_expr_free(old);
179  return 0;
180 }
181 
182 /**
183  * Helper checks if we can process main and overlay pixel formats
184  */
185 static int formats_match(const enum AVPixelFormat format_main, const enum AVPixelFormat format_overlay) {
186  switch(format_main) {
187  case AV_PIX_FMT_NV12:
188  return format_overlay == AV_PIX_FMT_NV12;
189  case AV_PIX_FMT_YUV420P:
190  return format_overlay == AV_PIX_FMT_YUV420P ||
191  format_overlay == AV_PIX_FMT_YUVA420P;
192  default:
193  return 0;
194  }
195 }
196 
197 /**
198  * Call overlay kernell for a plane
199  */
202  int x_position, int y_position,
203  uint8_t* main_data, int main_linesize,
204  int main_width, int main_height,
205  uint8_t* overlay_data, int overlay_linesize,
206  int overlay_width, int overlay_height,
207  uint8_t* alpha_data, int alpha_linesize,
208  int alpha_adj_x, int alpha_adj_y) {
209 
210  CudaFunctions *cu = ctx->hwctx->internal->cuda_dl;
211 
212  void* kernel_args[] = {
213  &x_position, &y_position,
214  &main_data, &main_linesize,
215  &overlay_data, &overlay_linesize,
216  &overlay_width, &overlay_height,
217  &alpha_data, &alpha_linesize,
218  &alpha_adj_x, &alpha_adj_y,
219  };
220 
221  return CHECK_CU(cu->cuLaunchKernel(
222  ctx->cu_func,
223  DIV_UP(main_width, BLOCK_X), DIV_UP(main_height, BLOCK_Y), 1,
224  BLOCK_X, BLOCK_Y, 1,
225  0, ctx->cu_stream, kernel_args, NULL));
226 }
227 
228 /**
229  * Perform blend overlay picture over main picture
230  */
232 {
233  int ret;
234 
235  AVFilterContext *avctx = fs->parent;
236  OverlayCUDAContext *ctx = avctx->priv;
237  AVFilterLink *outlink = avctx->outputs[0];
238  AVFilterLink *inlink = avctx->inputs[0];
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] = inlink->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
279  {
280  int64_t pos = input_main->pkt_pos;
281  ctx->var_values[VAR_POS] = pos == -1 ? NAN : pos;
282  }
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 
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:
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:
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 
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 
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  */
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  */
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  */
424 {
425  OverlayCUDAContext *ctx = avctx->priv;
426 
427  return ff_framesync_activate(&ctx->fs);
428 }
429 
430 /**
431  * Configure output
432  */
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  AVFilterContext* avctx = outlink->src;
440  OverlayCUDAContext* ctx = avctx->priv;
441 
442  AVFilterLink *inlink = avctx->inputs[0];
443  AVHWFramesContext *frames_ctx = (AVHWFramesContext*)inlink->hw_frames_ctx->data;
444 
445  AVFilterLink *inlink_overlay = avctx->inputs[1];
446  AVHWFramesContext *frames_ctx_overlay = (AVHWFramesContext*)inlink_overlay->hw_frames_ctx->data;
447 
448  CUcontext dummy, cuda_ctx;
449  CudaFunctions *cu;
450 
451  // check main input formats
452 
453  if (!frames_ctx) {
454  av_log(ctx, AV_LOG_ERROR, "No hw context provided on main input\n");
455  return AVERROR(EINVAL);
456  }
457 
458  ctx->in_format_main = frames_ctx->sw_format;
459  if (!format_is_supported(supported_main_formats, ctx->in_format_main)) {
460  av_log(ctx, AV_LOG_ERROR, "Unsupported main input format: %s\n",
461  av_get_pix_fmt_name(ctx->in_format_main));
462  return AVERROR(ENOSYS);
463  }
464 
465  // check overlay input formats
466 
467  if (!frames_ctx_overlay) {
468  av_log(ctx, AV_LOG_ERROR, "No hw context provided on overlay input\n");
469  return AVERROR(EINVAL);
470  }
471 
472  ctx->in_format_overlay = frames_ctx_overlay->sw_format;
473  if (!format_is_supported(supported_overlay_formats, ctx->in_format_overlay)) {
474  av_log(ctx, AV_LOG_ERROR, "Unsupported overlay input format: %s\n",
475  av_get_pix_fmt_name(ctx->in_format_overlay));
476  return AVERROR(ENOSYS);
477  }
478 
479  // check we can overlay pictures with those pixel formats
480 
481  if (!formats_match(ctx->in_format_main, ctx->in_format_overlay)) {
482  av_log(ctx, AV_LOG_ERROR, "Can't overlay %s on %s \n",
483  av_get_pix_fmt_name(ctx->in_format_overlay), av_get_pix_fmt_name(ctx->in_format_main));
484  return AVERROR(EINVAL);
485  }
486 
487  // initialize
488 
489  ctx->hw_device_ctx = av_buffer_ref(frames_ctx->device_ref);
490  if (!ctx->hw_device_ctx)
491  return AVERROR(ENOMEM);
492  ctx->hwctx = ((AVHWDeviceContext*)ctx->hw_device_ctx->data)->hwctx;
493 
494  cuda_ctx = ctx->hwctx->cuda_ctx;
495  ctx->fs.time_base = inlink->time_base;
496 
497  ctx->cu_stream = ctx->hwctx->stream;
498 
499  outlink->hw_frames_ctx = av_buffer_ref(inlink->hw_frames_ctx);
500  if (!outlink->hw_frames_ctx)
501  return AVERROR(ENOMEM);
502 
503  // load functions
504 
505  cu = ctx->hwctx->internal->cuda_dl;
506 
507  err = CHECK_CU(cu->cuCtxPushCurrent(cuda_ctx));
508  if (err < 0) {
509  return err;
510  }
511 
512  err = ff_cuda_load_module(ctx, ctx->hwctx, &ctx->cu_module, ff_vf_overlay_cuda_ptx_data, ff_vf_overlay_cuda_ptx_len);
513  if (err < 0) {
514  CHECK_CU(cu->cuCtxPopCurrent(&dummy));
515  return err;
516  }
517 
518  err = CHECK_CU(cu->cuModuleGetFunction(&ctx->cu_func, ctx->cu_module, "Overlay_Cuda"));
519  if (err < 0) {
520  CHECK_CU(cu->cuCtxPopCurrent(&dummy));
521  return err;
522  }
523 
524  CHECK_CU(cu->cuCtxPopCurrent(&dummy));
525 
526  // init dual input
527 
528  err = ff_framesync_init_dualinput(&ctx->fs, avctx);
529  if (err < 0) {
530  return err;
531  }
532 
533  return ff_framesync_configure(&ctx->fs);
534 }
535 
536 
537 #define OFFSET(x) offsetof(OverlayCUDAContext, x)
538 #define FLAGS (AV_OPT_FLAG_FILTERING_PARAM | AV_OPT_FLAG_VIDEO_PARAM)
539 
540 static const AVOption overlay_cuda_options[] = {
541  { "x", "set the x expression of overlay", OFFSET(x_expr), AV_OPT_TYPE_STRING, { .str = "0" }, 0, 0, FLAGS },
542  { "y", "set the y expression of overlay", OFFSET(y_expr), AV_OPT_TYPE_STRING, { .str = "0" }, 0, 0, FLAGS },
543  { "eof_action", "Action to take when encountering EOF from secondary input ",
544  OFFSET(fs.opt_eof_action), AV_OPT_TYPE_INT, { .i64 = EOF_ACTION_REPEAT },
545  EOF_ACTION_REPEAT, EOF_ACTION_PASS, .flags = FLAGS, .unit = "eof_action" },
546  { "repeat", "Repeat the previous frame.", 0, AV_OPT_TYPE_CONST, { .i64 = EOF_ACTION_REPEAT }, .flags = FLAGS, .unit = "eof_action" },
547  { "endall", "End both streams.", 0, AV_OPT_TYPE_CONST, { .i64 = EOF_ACTION_ENDALL }, .flags = FLAGS, .unit = "eof_action" },
548  { "pass", "Pass through the main input.", 0, AV_OPT_TYPE_CONST, { .i64 = EOF_ACTION_PASS }, .flags = FLAGS, .unit = "eof_action" },
549  { "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" },
550  { "init", "eval expressions once during initialization", 0, AV_OPT_TYPE_CONST, { .i64=EVAL_MODE_INIT }, .flags = FLAGS, .unit = "eval" },
551  { "frame", "eval expressions per-frame", 0, AV_OPT_TYPE_CONST, { .i64=EVAL_MODE_FRAME }, .flags = FLAGS, .unit = "eval" },
552  { "shortest", "force termination when the shortest input terminates", OFFSET(fs.opt_shortest), AV_OPT_TYPE_BOOL, { .i64 = 0 }, 0, 1, FLAGS },
553  { "repeatlast", "repeat overlay of the last overlay frame", OFFSET(fs.opt_repeatlast), AV_OPT_TYPE_BOOL, {.i64=1}, 0, 1, FLAGS },
554  { NULL },
555 };
556 
558 
560  {
561  .name = "main",
562  .type = AVMEDIA_TYPE_VIDEO,
563  },
564  {
565  .name = "overlay",
566  .type = AVMEDIA_TYPE_VIDEO,
567  .config_props = config_input_overlay,
568  },
569 };
570 
572  {
573  .name = "default",
574  .type = AVMEDIA_TYPE_VIDEO,
575  .config_props = &overlay_cuda_config_output,
576  },
577 };
578 
580  .name = "overlay_cuda",
581  .description = NULL_IF_CONFIG_SMALL("Overlay one video on top of another using CUDA"),
582  .priv_size = sizeof(OverlayCUDAContext),
583  .priv_class = &overlay_cuda_class,
590  .preinit = overlay_cuda_framesync_preinit,
591  .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE,
592 };
formats
formats
Definition: signature.h:48
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:134
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:185
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:200
var_name
var_name
Definition: noise.c:46
hwcontext_cuda_internal.h
ff_framesync_uninit
void ff_framesync_uninit(FFFrameSync *fs)
Free all memory currently allocated.
Definition: framesync.c:304
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: internal.h:343
ff_filter_frame
int ff_filter_frame(AVFilterLink *link, AVFrame *frame)
Send a frame of data to the next filter.
Definition: avfilter.c:1007
overlay_cuda_inputs
static const AVFilterPad overlay_cuda_inputs[]
Definition: vf_overlay_cuda.c:559
OverlayCUDAContext::y_position
int y_position
Definition: vf_overlay_cuda.c:120
OverlayCUDAContext::cu_ctx
CUcontext cu_ctx
Definition: vf_overlay_cuda.c:111
EVAL_MODE_INIT
@ EVAL_MODE_INIT
Definition: vf_overlay_cuda.c:79
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: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:100
AVFrame
This structure describes decoded (raw) audio or video data.
Definition: frame.h:340
pixdesc.h
AVFrame::pts
int64_t pts
Presentation timestamp in time_base units (time when frame should be shown to user).
Definition: frame.h:452
AVFrame::width
int width
Definition: frame.h:412
OverlayCUDAContext::hw_device_ctx
AVBufferRef * hw_device_ctx
Definition: vf_overlay_cuda.c:108
OverlayCUDAContext::in_format_overlay
enum AVPixelFormat in_format_overlay
Definition: vf_overlay_cuda.c:105
AVOption
AVOption.
Definition: opt.h:294
normalize_xy
static int normalize_xy(double d, int chroma_sub)
Definition: vf_overlay_cuda.c:139
EOF_ACTION_ENDALL
@ EOF_ACTION_ENDALL
Definition: framesync.h:28
overlay_cuda_outputs
static const AVFilterPad overlay_cuda_outputs[]
Definition: vf_overlay_cuda.c:571
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:106
OVERLAY
#define OVERLAY
Definition: vf_overlay_cuda.c:48
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:170
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:361
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:711
ff_vf_overlay_cuda
const AVFilter ff_vf_overlay_cuda
Definition: vf_overlay_cuda.c:579
AVFilterContext::priv
void * priv
private data for use by the filter
Definition: avfilter.h:424
VAR_MW
@ VAR_MW
Definition: vf_overlay_cuda.c:64
dummy
int dummy
Definition: motion.c:66
OverlayCUDAContext::fs
FFFrameSync fs
Definition: vf_overlay_cuda.c:116
av_expr_free
void av_expr_free(AVExpr *e)
Free a parsed expression previously created with av_expr_parse().
Definition: eval.c:359
AVFilterPad
A filter pad used for either input or output.
Definition: internal.h:33
AVHWDeviceContext
This struct aggregates all the (hardware/vendor-specific) "high-level" state, i.e.
Definition: hwcontext.h:62
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:119
OFFSET
#define OFFSET(x)
Definition: vf_overlay_cuda.c:537
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:48
av_expr_eval
double av_expr_eval(AVExpr *e, const double *const_values, void *opaque)
Evaluate a previously parsed expression.
Definition: eval.c:793
AVExpr
Definition: eval.c:159
VAR_OH
@ VAR_OH
Definition: vf_overlay_cuda.c:67
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
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
FILTER_INPUTS
#define FILTER_INPUTS(array)
Definition: internal.h:182
ff_inlink_make_frame_writable
int ff_inlink_make_frame_writable(AVFilterLink *link, AVFrame **rframe)
Make sure a frame is writable.
Definition: avfilter.c:1485
option
option
Definition: libkvazaar.c:320
CHECK_CU
#define CHECK_CU(x)
Definition: vf_overlay_cuda.c:41
set_expr
static int set_expr(AVExpr **pexpr, const char *expr, const char *option, void *log_ctx)
Definition: vf_overlay_cuda.c:161
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:45
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:226
VAR_POS
@ VAR_POS
Definition: noise.c:55
OverlayCUDAContext::x_expr
char * x_expr
Definition: vf_overlay_cuda.c:123
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:81
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:142
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:417
MAIN
#define MAIN
Definition: vf_overlay_cuda.c:47
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:365
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:106
ff_framesync_init_dualinput
int ff_framesync_init_dualinput(FFFrameSync *fs, AVFilterContext *parent)
Initialize a frame sync structure for dualinput.
Definition: framesync.c:375
VAR_Y
@ VAR_Y
Definition: vf_overlay_cuda.c:69
OverlayCUDAContext::x_pexpr
AVExpr * x_pexpr
Definition: vf_overlay_cuda.c:125
VAR_MAIN_H
@ VAR_MAIN_H
Definition: vf_overlay_cuda.c:65
AV_NOPTS_VALUE
#define AV_NOPTS_VALUE
Undefined timestamp value.
Definition: avutil.h:248
FLAGS
#define FLAGS
Definition: vf_overlay_cuda.c:538
overlay_cuda_options
static const AVOption overlay_cuda_options[]
Definition: vf_overlay_cuda.c:540
VAR_T
@ VAR_T
Definition: vf_overlay_cuda.c:74
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:687
EVAL_MODE_FRAME
@ EVAL_MODE_FRAME
Definition: vf_overlay_cuda.c:80
OverlayCUDAContext
OverlayCUDAContext.
Definition: vf_overlay_cuda.c:102
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:131
VAR_X
@ VAR_X
Definition: vf_overlay_cuda.c:68
internal.h
FILTER_SINGLE_PIXFMT
#define FILTER_SINGLE_PIXFMT(pix_fmt_)
Definition: internal.h:172
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:231
log.h
i
#define i(width, name, range_min, range_max)
Definition: cbs_h2645.c:255
VAR_OVERLAY_H
@ VAR_OVERLAY_H
Definition: vf_overlay_cuda.c:67
OverlayCUDAContext::hwctx
AVCUDADeviceContext * hwctx
Definition: vf_overlay_cuda.c:109
EvalMode
EvalMode
Definition: af_volume.h:39
BLOCK_X
#define BLOCK_X
Definition: vf_overlay_cuda.c:44
OverlayCUDAContext::cu_stream
CUstream cu_stream
Definition: vf_overlay_cuda.c:114
OverlayCUDAContext::cu_func
CUfunction cu_func
Definition: vf_overlay_cuda.c:113
AVFilterPad::name
const char * name
Pad name.
Definition: internal.h:39
AVFilter
Filter definition.
Definition: avfilter.h:166
AVHWFramesContext
This struct describes a set or pool of "hardware" frames (i.e.
Definition: hwcontext.h:125
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:56
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:112
pos
unsigned int pos
Definition: spdifenc.c:413
EOF_ACTION_REPEAT
@ EOF_ACTION_REPEAT
Definition: framesync.h:27
AVFrame::height
int height
Definition: frame.h:412
OverlayCUDAContext::y_pexpr
AVExpr * y_pexpr
Definition: vf_overlay_cuda.c:125
supported_main_formats
static enum AVPixelFormat supported_main_formats[]
Definition: vf_overlay_cuda.c:50
framesync.h
DIV_UP
#define DIV_UP(a, b)
Definition: vf_overlay_cuda.c:42
VAR_MH
@ VAR_MH
Definition: vf_overlay_cuda.c:65
VAR_OW
@ VAR_OW
Definition: vf_overlay_cuda.c:66
OverlayCUDAContext::y_expr
char * y_expr
Definition: vf_overlay_cuda.c:123
AV_PIX_FMT_NONE
@ AV_PIX_FMT_NONE
Definition: pixfmt.h:72
AV_OPT_TYPE_INT
@ AV_OPT_TYPE_INT
Definition: opt.h:225
avfilter.h
VAR_OVERLAY_W
@ VAR_OVERLAY_W
Definition: vf_overlay_cuda.c:66
AVFilterContext
An instance of a filter.
Definition: avfilter.h:409
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:84
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:64
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
Definition: opt.h:244
FILTER_OUTPUTS
#define FILTER_OUTPUTS(array)
Definition: internal.h:183
VAR_VARS_NB
@ VAR_VARS_NB
Definition: vf_overlay_cuda.c:75
VAR_N
@ VAR_N
Definition: vf_overlay_cuda.c:70
d
d
Definition: ffmpeg_filter.c:424
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:385
av_log
#define av_log(a,...)
Definition: tableprint_vlc.h:27
eval_expr
static void eval_expr(AVFilterContext *ctx)
Definition: vf_overlay_cuda.c:146
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:355
ff_framesync_dualinput_get
int ff_framesync_dualinput_get(FFFrameSync *fs, AVFrame **f0, AVFrame **f1)
Definition: framesync.c:393
AV_OPT_TYPE_STRING
@ AV_OPT_TYPE_STRING
Definition: opt.h:229
OverlayCUDAContext::eval_mode
int eval_mode
Definition: vf_overlay_cuda.c:118
AV_OPT_TYPE_CONST
@ AV_OPT_TYPE_CONST
Definition: opt.h:234
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:2888
AVFilterContext::outputs
AVFilterLink ** outputs
array of pointers to output links
Definition: avfilter.h:421
OverlayCUDAContext::var_values
double var_values[VAR_VARS_NB]
Definition: vf_overlay_cuda.c:122