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 "framesync.h"
36 #include "internal.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[] = {
53 };
54 
60 };
61 
62 enum var_name {
73 };
74 
75 enum EvalMode {
79 };
80 
81 static const char *const var_names[] = {
82  "main_w", "W", ///< width of the main video
83  "main_h", "H", ///< height of the main video
84  "overlay_w", "w", ///< width of the overlay video
85  "overlay_h", "h", ///< height of the overlay video
86  "x",
87  "y",
88  "n", ///< number of frame
89  "pos", ///< position in the file
90  "t", ///< timestamp expressed in seconds
91  NULL
92 };
93 
94 /**
95  * OverlayCUDAContext
96  */
97 typedef struct OverlayCUDAContext {
98  const AVClass *class;
99 
102 
105 
106  CUcontext cu_ctx;
107  CUmodule cu_module;
108  CUfunction cu_func;
109  CUstream cu_stream;
110 
112 
116 
118  char *x_expr, *y_expr;
119 
122 
123 /**
124  * Helper to find out if provided format is supported by filter
125  */
126 static int format_is_supported(const enum AVPixelFormat formats[], enum AVPixelFormat fmt)
127 {
128  for (int i = 0; formats[i] != AV_PIX_FMT_NONE; i++)
129  if (formats[i] == fmt)
130  return 1;
131  return 0;
132 }
133 
134 static inline int normalize_xy(double d, int chroma_sub)
135 {
136  if (isnan(d))
137  return INT_MAX;
138  return (int)d & ~((1 << chroma_sub) - 1);
139 }
140 
142 {
143  OverlayCUDAContext *s = ctx->priv;
144 
145  s->var_values[VAR_X] = av_expr_eval(s->x_pexpr, s->var_values, NULL);
146  s->var_values[VAR_Y] = av_expr_eval(s->y_pexpr, s->var_values, NULL);
147  /* necessary if x is expressed from y */
148  s->var_values[VAR_X] = av_expr_eval(s->x_pexpr, s->var_values, NULL);
149 
150  s->x_position = normalize_xy(s->var_values[VAR_X], 1);
151 
152  /* the cuda pixel format is using hwaccel, normalizing y is unnecessary */
153  s->y_position = s->var_values[VAR_Y];
154 }
155 
156 static int set_expr(AVExpr **pexpr, const char *expr, const char *option, void *log_ctx)
157 {
158  int ret;
159  AVExpr *old = NULL;
160 
161  if (*pexpr)
162  old = *pexpr;
163  ret = av_expr_parse(pexpr, expr, var_names,
164  NULL, NULL, NULL, NULL, 0, log_ctx);
165  if (ret < 0) {
166  av_log(log_ctx, AV_LOG_ERROR,
167  "Error when evaluating the expression '%s' for %s\n",
168  expr, option);
169  *pexpr = old;
170  return ret;
171  }
172 
173  av_expr_free(old);
174  return 0;
175 }
176 
177 /**
178  * Helper checks if we can process main and overlay pixel formats
179  */
180 static int formats_match(const enum AVPixelFormat format_main, const enum AVPixelFormat format_overlay) {
181  switch(format_main) {
182  case AV_PIX_FMT_NV12:
183  return format_overlay == AV_PIX_FMT_NV12;
184  case AV_PIX_FMT_YUV420P:
185  return format_overlay == AV_PIX_FMT_YUV420P ||
186  format_overlay == AV_PIX_FMT_YUVA420P;
187  default:
188  return 0;
189  }
190 }
191 
192 /**
193  * Call overlay kernell for a plane
194  */
197  int x_position, int y_position,
198  uint8_t* main_data, int main_linesize,
199  int main_width, int main_height,
200  uint8_t* overlay_data, int overlay_linesize,
201  int overlay_width, int overlay_height,
202  uint8_t* alpha_data, int alpha_linesize,
203  int alpha_adj_x, int alpha_adj_y) {
204 
205  CudaFunctions *cu = ctx->hwctx->internal->cuda_dl;
206 
207  void* kernel_args[] = {
208  &x_position, &y_position,
209  &main_data, &main_linesize,
210  &overlay_data, &overlay_linesize,
211  &overlay_width, &overlay_height,
212  &alpha_data, &alpha_linesize,
213  &alpha_adj_x, &alpha_adj_y,
214  };
215 
216  return CHECK_CU(cu->cuLaunchKernel(
217  ctx->cu_func,
218  DIV_UP(main_width, BLOCK_X), DIV_UP(main_height, BLOCK_Y), 1,
219  BLOCK_X, BLOCK_Y, 1,
220  0, ctx->cu_stream, kernel_args, NULL));
221 }
222 
223 /**
224  * Perform blend overlay picture over main picture
225  */
227 {
228  int ret;
229 
230  AVFilterContext *avctx = fs->parent;
231  OverlayCUDAContext *ctx = avctx->priv;
232  AVFilterLink *outlink = avctx->outputs[0];
233  AVFilterLink *inlink = avctx->inputs[0];
234 
235  CudaFunctions *cu = ctx->hwctx->internal->cuda_dl;
236  CUcontext dummy, cuda_ctx = ctx->hwctx->cuda_ctx;
237 
238  AVFrame *input_main, *input_overlay;
239 
240  int pos = 0;
241 
242  ctx->cu_ctx = cuda_ctx;
243 
244  // read main and overlay frames from inputs
245  ret = ff_framesync_dualinput_get(fs, &input_main, &input_overlay);
246  if (ret < 0)
247  return ret;
248 
249  if (!input_main)
250  return AVERROR_BUG;
251 
252  if (!input_overlay)
253  return ff_filter_frame(outlink, input_main);
254 
255  ret = av_frame_make_writable(input_main);
256  if (ret < 0) {
257  av_frame_free(&input_main);
258  return ret;
259  }
260 
261  // push cuda context
262 
263  ret = CHECK_CU(cu->cuCtxPushCurrent(cuda_ctx));
264  if (ret < 0) {
265  av_frame_free(&input_main);
266  return ret;
267  }
268 
269  if (ctx->eval_mode == EVAL_MODE_FRAME) {
270  pos = input_main->pkt_pos;
271  ctx->var_values[VAR_N] = inlink->frame_count_out;
272  ctx->var_values[VAR_T] = input_main->pts == AV_NOPTS_VALUE ?
273  NAN : input_main->pts * av_q2d(inlink->time_base);
274  ctx->var_values[VAR_POS] = pos == -1 ? NAN : pos;
275  ctx->var_values[VAR_OVERLAY_W] = ctx->var_values[VAR_OW] = input_overlay->width;
276  ctx->var_values[VAR_OVERLAY_H] = ctx->var_values[VAR_OH] = input_overlay->height;
277  ctx->var_values[VAR_MAIN_W ] = ctx->var_values[VAR_MW] = input_main->width;
278  ctx->var_values[VAR_MAIN_H ] = ctx->var_values[VAR_MH] = input_main->height;
279 
280  eval_expr(avctx);
281 
282  av_log(avctx, AV_LOG_DEBUG, "n:%f t:%f pos:%f x:%f xi:%d y:%f yi:%d\n",
283  ctx->var_values[VAR_N], ctx->var_values[VAR_T], ctx->var_values[VAR_POS],
284  ctx->var_values[VAR_X], ctx->x_position,
285  ctx->var_values[VAR_Y], ctx->y_position);
286  }
287 
288  // overlay first plane
289 
291  ctx->x_position, ctx->y_position,
292  input_main->data[0], input_main->linesize[0],
293  input_main->width, input_main->height,
294  input_overlay->data[0], input_overlay->linesize[0],
295  input_overlay->width, input_overlay->height,
296  input_overlay->data[3], input_overlay->linesize[3], 1, 1);
297 
298  // overlay rest planes depending on pixel format
299 
300  switch(ctx->in_format_overlay) {
301  case AV_PIX_FMT_NV12:
303  ctx->x_position, ctx->y_position / 2,
304  input_main->data[1], input_main->linesize[1],
305  input_main->width, input_main->height / 2,
306  input_overlay->data[1], input_overlay->linesize[1],
307  input_overlay->width, input_overlay->height / 2,
308  0, 0, 0, 0);
309  break;
310  case AV_PIX_FMT_YUV420P:
311  case AV_PIX_FMT_YUVA420P:
313  ctx->x_position / 2 , ctx->y_position / 2,
314  input_main->data[1], input_main->linesize[1],
315  input_main->width / 2, input_main->height / 2,
316  input_overlay->data[1], input_overlay->linesize[1],
317  input_overlay->width / 2, input_overlay->height / 2,
318  input_overlay->data[3], input_overlay->linesize[3], 2, 2);
319 
321  ctx->x_position / 2 , ctx->y_position / 2,
322  input_main->data[2], input_main->linesize[2],
323  input_main->width / 2, input_main->height / 2,
324  input_overlay->data[2], input_overlay->linesize[2],
325  input_overlay->width / 2, input_overlay->height / 2,
326  input_overlay->data[3], input_overlay->linesize[3], 2, 2);
327  break;
328  default:
329  av_log(ctx, AV_LOG_ERROR, "Passed unsupported overlay pixel format\n");
330  av_frame_free(&input_main);
331  CHECK_CU(cu->cuCtxPopCurrent(&dummy));
332  return AVERROR_BUG;
333  }
334 
335  CHECK_CU(cu->cuCtxPopCurrent(&dummy));
336 
337  return ff_filter_frame(outlink, input_main);
338 }
339 
341 {
342  AVFilterContext *ctx = inlink->dst;
343  OverlayCUDAContext *s = inlink->dst->priv;
344  int ret;
345 
346 
347  /* Finish the configuration by evaluating the expressions
348  now when both inputs are configured. */
349  s->var_values[VAR_MAIN_W ] = s->var_values[VAR_MW] = ctx->inputs[MAIN ]->w;
350  s->var_values[VAR_MAIN_H ] = s->var_values[VAR_MH] = ctx->inputs[MAIN ]->h;
351  s->var_values[VAR_OVERLAY_W] = s->var_values[VAR_OW] = ctx->inputs[OVERLAY]->w;
352  s->var_values[VAR_OVERLAY_H] = s->var_values[VAR_OH] = ctx->inputs[OVERLAY]->h;
353  s->var_values[VAR_X] = NAN;
354  s->var_values[VAR_Y] = NAN;
355  s->var_values[VAR_N] = 0;
356  s->var_values[VAR_T] = NAN;
357  s->var_values[VAR_POS] = NAN;
358 
359  if ((ret = set_expr(&s->x_pexpr, s->x_expr, "x", ctx)) < 0 ||
360  (ret = set_expr(&s->y_pexpr, s->y_expr, "y", ctx)) < 0)
361  return ret;
362 
363  if (s->eval_mode == EVAL_MODE_INIT) {
364  eval_expr(ctx);
365  av_log(ctx, AV_LOG_VERBOSE, "x:%f xi:%d y:%f yi:%d\n",
366  s->var_values[VAR_X], s->x_position,
367  s->var_values[VAR_Y], s->y_position);
368  }
369 
370  return 0;
371 }
372 
373 /**
374  * Initialize overlay_cuda
375  */
377 {
378  OverlayCUDAContext* ctx = avctx->priv;
379  ctx->fs.on_event = &overlay_cuda_blend;
380 
381  return 0;
382 }
383 
384 /**
385  * Uninitialize overlay_cuda
386  */
388 {
389  OverlayCUDAContext* ctx = avctx->priv;
390 
391  ff_framesync_uninit(&ctx->fs);
392 
393  if (ctx->hwctx && ctx->cu_module) {
394  CUcontext dummy;
395  CudaFunctions *cu = ctx->hwctx->internal->cuda_dl;
396  CHECK_CU(cu->cuCtxPushCurrent(ctx->cu_ctx));
397  CHECK_CU(cu->cuModuleUnload(ctx->cu_module));
398  CHECK_CU(cu->cuCtxPopCurrent(&dummy));
399  }
400 
401  av_expr_free(ctx->x_pexpr); ctx->x_pexpr = NULL;
402  av_expr_free(ctx->y_pexpr); ctx->y_pexpr = NULL;
403  av_buffer_unref(&ctx->hw_device_ctx);
404  ctx->hwctx = NULL;
405 }
406 
407 /**
408  * Activate overlay_cuda
409  */
411 {
412  OverlayCUDAContext *ctx = avctx->priv;
413 
414  return ff_framesync_activate(&ctx->fs);
415 }
416 
417 /**
418  * Configure output
419  */
421 {
422  extern const unsigned char ff_vf_overlay_cuda_ptx_data[];
423  extern const unsigned int ff_vf_overlay_cuda_ptx_len;
424 
425  int err;
426  AVFilterContext* avctx = outlink->src;
427  OverlayCUDAContext* ctx = avctx->priv;
428 
429  AVFilterLink *inlink = avctx->inputs[0];
430  AVHWFramesContext *frames_ctx = (AVHWFramesContext*)inlink->hw_frames_ctx->data;
431 
432  AVFilterLink *inlink_overlay = avctx->inputs[1];
433  AVHWFramesContext *frames_ctx_overlay = (AVHWFramesContext*)inlink_overlay->hw_frames_ctx->data;
434 
435  CUcontext dummy, cuda_ctx;
436  CudaFunctions *cu;
437 
438  // check main input formats
439 
440  if (!frames_ctx) {
441  av_log(ctx, AV_LOG_ERROR, "No hw context provided on main input\n");
442  return AVERROR(EINVAL);
443  }
444 
445  ctx->in_format_main = frames_ctx->sw_format;
446  if (!format_is_supported(supported_main_formats, ctx->in_format_main)) {
447  av_log(ctx, AV_LOG_ERROR, "Unsupported main input format: %s\n",
448  av_get_pix_fmt_name(ctx->in_format_main));
449  return AVERROR(ENOSYS);
450  }
451 
452  // check overlay input formats
453 
454  if (!frames_ctx_overlay) {
455  av_log(ctx, AV_LOG_ERROR, "No hw context provided on overlay input\n");
456  return AVERROR(EINVAL);
457  }
458 
459  ctx->in_format_overlay = frames_ctx_overlay->sw_format;
460  if (!format_is_supported(supported_overlay_formats, ctx->in_format_overlay)) {
461  av_log(ctx, AV_LOG_ERROR, "Unsupported overlay input format: %s\n",
462  av_get_pix_fmt_name(ctx->in_format_overlay));
463  return AVERROR(ENOSYS);
464  }
465 
466  // check we can overlay pictures with those pixel formats
467 
468  if (!formats_match(ctx->in_format_main, ctx->in_format_overlay)) {
469  av_log(ctx, AV_LOG_ERROR, "Can't overlay %s on %s \n",
470  av_get_pix_fmt_name(ctx->in_format_overlay), av_get_pix_fmt_name(ctx->in_format_main));
471  return AVERROR(EINVAL);
472  }
473 
474  // initialize
475 
476  ctx->hw_device_ctx = av_buffer_ref(frames_ctx->device_ref);
477  if (!ctx->hw_device_ctx)
478  return AVERROR(ENOMEM);
479  ctx->hwctx = ((AVHWDeviceContext*)ctx->hw_device_ctx->data)->hwctx;
480 
481  cuda_ctx = ctx->hwctx->cuda_ctx;
482  ctx->fs.time_base = inlink->time_base;
483 
484  ctx->cu_stream = ctx->hwctx->stream;
485 
486  outlink->hw_frames_ctx = av_buffer_ref(inlink->hw_frames_ctx);
487  if (!outlink->hw_frames_ctx)
488  return AVERROR(ENOMEM);
489 
490  // load functions
491 
492  cu = ctx->hwctx->internal->cuda_dl;
493 
494  err = CHECK_CU(cu->cuCtxPushCurrent(cuda_ctx));
495  if (err < 0) {
496  return err;
497  }
498 
499  err = ff_cuda_load_module(ctx, ctx->hwctx, &ctx->cu_module, ff_vf_overlay_cuda_ptx_data, ff_vf_overlay_cuda_ptx_len);
500  if (err < 0) {
501  CHECK_CU(cu->cuCtxPopCurrent(&dummy));
502  return err;
503  }
504 
505  err = CHECK_CU(cu->cuModuleGetFunction(&ctx->cu_func, ctx->cu_module, "Overlay_Cuda"));
506  if (err < 0) {
507  CHECK_CU(cu->cuCtxPopCurrent(&dummy));
508  return err;
509  }
510 
511  CHECK_CU(cu->cuCtxPopCurrent(&dummy));
512 
513  // init dual input
514 
515  err = ff_framesync_init_dualinput(&ctx->fs, avctx);
516  if (err < 0) {
517  return err;
518  }
519 
520  return ff_framesync_configure(&ctx->fs);
521 }
522 
523 
524 #define OFFSET(x) offsetof(OverlayCUDAContext, x)
525 #define FLAGS (AV_OPT_FLAG_FILTERING_PARAM | AV_OPT_FLAG_VIDEO_PARAM)
526 
527 static const AVOption overlay_cuda_options[] = {
528  { "x", "set the x expression of overlay", OFFSET(x_expr), AV_OPT_TYPE_STRING, { .str = "0" }, 0, 0, FLAGS },
529  { "y", "set the y expression of overlay", OFFSET(y_expr), AV_OPT_TYPE_STRING, { .str = "0" }, 0, 0, FLAGS },
530  { "eof_action", "Action to take when encountering EOF from secondary input ",
531  OFFSET(fs.opt_eof_action), AV_OPT_TYPE_INT, { .i64 = EOF_ACTION_REPEAT },
532  EOF_ACTION_REPEAT, EOF_ACTION_PASS, .flags = FLAGS, "eof_action" },
533  { "repeat", "Repeat the previous frame.", 0, AV_OPT_TYPE_CONST, { .i64 = EOF_ACTION_REPEAT }, .flags = FLAGS, "eof_action" },
534  { "endall", "End both streams.", 0, AV_OPT_TYPE_CONST, { .i64 = EOF_ACTION_ENDALL }, .flags = FLAGS, "eof_action" },
535  { "pass", "Pass through the main input.", 0, AV_OPT_TYPE_CONST, { .i64 = EOF_ACTION_PASS }, .flags = FLAGS, "eof_action" },
536  { "eval", "specify when to evaluate expressions", OFFSET(eval_mode), AV_OPT_TYPE_INT, { .i64 = EVAL_MODE_FRAME }, 0, EVAL_MODE_NB - 1, FLAGS, "eval" },
537  { "init", "eval expressions once during initialization", 0, AV_OPT_TYPE_CONST, { .i64=EVAL_MODE_INIT }, .flags = FLAGS, .unit = "eval" },
538  { "frame", "eval expressions per-frame", 0, AV_OPT_TYPE_CONST, { .i64=EVAL_MODE_FRAME }, .flags = FLAGS, .unit = "eval" },
539  { "shortest", "force termination when the shortest input terminates", OFFSET(fs.opt_shortest), AV_OPT_TYPE_BOOL, { .i64 = 0 }, 0, 1, FLAGS },
540  { "repeatlast", "repeat overlay of the last overlay frame", OFFSET(fs.opt_repeatlast), AV_OPT_TYPE_BOOL, {.i64=1}, 0, 1, FLAGS },
541  { NULL },
542 };
543 
545 
547  {
548  .name = "main",
549  .type = AVMEDIA_TYPE_VIDEO,
550  },
551  {
552  .name = "overlay",
553  .type = AVMEDIA_TYPE_VIDEO,
554  .config_props = config_input_overlay,
555  },
556 };
557 
559  {
560  .name = "default",
561  .type = AVMEDIA_TYPE_VIDEO,
562  .config_props = &overlay_cuda_config_output,
563  },
564 };
565 
567  .name = "overlay_cuda",
568  .description = NULL_IF_CONFIG_SMALL("Overlay one video on top of another using CUDA"),
569  .priv_size = sizeof(OverlayCUDAContext),
570  .priv_class = &overlay_cuda_class,
577  .preinit = overlay_cuda_framesync_preinit,
578  .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE,
579 };
formats
formats
Definition: signature.h:48
ff_framesync_configure
int ff_framesync_configure(FFFrameSync *fs)
Configure a frame sync structure.
Definition: framesync.c:119
AV_PIX_FMT_CUDA
@ AV_PIX_FMT_CUDA
HW acceleration through CUDA.
Definition: pixfmt.h:225
AVPixelFormat
AVPixelFormat
Pixel format.
Definition: pixfmt.h:64
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:180
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:195
VAR_POS
@ VAR_POS
Definition: vf_overlay_cuda.c:70
hwcontext_cuda_internal.h
ff_framesync_uninit
void ff_framesync_uninit(FFFrameSync *fs)
Free all memory currently allocated.
Definition: framesync.c:285
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:370
ff_filter_frame
int ff_filter_frame(AVFilterLink *link, AVFrame *frame)
Send a frame of data to the next filter.
Definition: avfilter.c:999
overlay_cuda_inputs
static const AVFilterPad overlay_cuda_inputs[]
Definition: vf_overlay_cuda.c:546
OverlayCUDAContext::y_position
int y_position
Definition: vf_overlay_cuda.c:115
OverlayCUDAContext::cu_ctx
CUcontext cu_ctx
Definition: vf_overlay_cuda.c:106
EVAL_MODE_INIT
@ EVAL_MODE_INIT
Definition: vf_overlay_cuda.c:76
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 neccesary.
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:111
AVFrame
This structure describes decoded (raw) audio or video data.
Definition: frame.h:325
av_frame_make_writable
int av_frame_make_writable(AVFrame *frame)
Ensure that the frame data is writable, avoiding data copy if possible.
Definition: frame.c:540
pixdesc.h
AVFrame::pts
int64_t pts
Presentation timestamp in time_base units (time when frame should be shown to user).
Definition: frame.h:432
AVFrame::width
int width
Definition: frame.h:397
OverlayCUDAContext::hw_device_ctx
AVBufferRef * hw_device_ctx
Definition: vf_overlay_cuda.c:103
OverlayCUDAContext::in_format_overlay
enum AVPixelFormat in_format_overlay
Definition: vf_overlay_cuda.c:100
AVOption
AVOption.
Definition: opt.h:251
normalize_xy
static int normalize_xy(double d, int chroma_sub)
Definition: vf_overlay_cuda.c:134
EOF_ACTION_ENDALL
@ EOF_ACTION_ENDALL
Definition: framesync.h:28
overlay_cuda_outputs
static const AVFilterPad overlay_cuda_outputs[]
Definition: vf_overlay_cuda.c:558
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:101
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:175
FFFrameSync
Frame sync structure.
Definition: framesync.h:146
AVFrame::data
uint8_t * data[AV_NUM_DATA_POINTERS]
pointer to the picture/channel planes.
Definition: frame.h:346
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:685
init
static int init
Definition: av_tx.c:47
ff_vf_overlay_cuda
const AVFilter ff_vf_overlay_cuda
Definition: vf_overlay_cuda.c:566
AVFilterContext::priv
void * priv
private data for use by the filter
Definition: avfilter.h:423
VAR_MW
@ VAR_MW
Definition: vf_overlay_cuda.c:63
dummy
int dummy
Definition: motion.c:65
OverlayCUDAContext::fs
FFFrameSync fs
Definition: vf_overlay_cuda.c:111
av_expr_free
void av_expr_free(AVExpr *e)
Free a parsed expression previously created with av_expr_parse().
Definition: eval.c:336
AVFilterPad
A filter pad used for either input or output.
Definition: internal.h:49
AVHWDeviceContext
This struct aggregates all the (hardware/vendor-specific) "high-level" state, i.e.
Definition: hwcontext.h:61
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:256
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:101
AVFrame::pkt_pos
int64_t pkt_pos
reordered pos from the last AVPacket that has been input into the decoder
Definition: frame.h:605
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:114
OFFSET
#define OFFSET(x)
Definition: vf_overlay_cuda.c:524
var_name
var_name
Definition: noise_bsf.c:46
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:766
AVExpr
Definition: eval.c:157
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:66
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:420
NAN
#define NAN
Definition: mathematics.h:64
FILTER_INPUTS
#define FILTER_INPUTS(array)
Definition: internal.h:190
option
option
Definition: libkvazaar.c:312
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:156
config_input_overlay
static int config_input_overlay(AVFilterLink *inlink)
Definition: vf_overlay_cuda.c:340
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:222
OverlayCUDAContext::x_expr
char * x_expr
Definition: vf_overlay_cuda.c:118
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:78
fs
#define fs(width, name, subs,...)
Definition: cbs_vp9.c:258
AVHWFramesContext::device_ref
AVBufferRef * device_ref
A reference to the parent AVHWDeviceContext.
Definition: hwcontext.h:141
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:416
MAIN
#define MAIN
Definition: vf_overlay_cuda.c:46
overlay_cuda_init
static av_cold int overlay_cuda_init(AVFilterContext *avctx)
Initialize overlay_cuda.
Definition: vf_overlay_cuda.c:376
eval.h
FRAMESYNC_DEFINE_CLASS
FRAMESYNC_DEFINE_CLASS(overlay_cuda, OverlayCUDAContext, fs)
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:117
ff_framesync_init_dualinput
int ff_framesync_init_dualinput(FFFrameSync *fs, AVFilterContext *parent)
Initialize a frame sync structure for dualinput.
Definition: framesync.c:353
VAR_Y
@ VAR_Y
Definition: vf_overlay_cuda.c:68
OverlayCUDAContext::x_pexpr
AVExpr * x_pexpr
Definition: vf_overlay_cuda.c:120
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:525
overlay_cuda_options
static const AVOption overlay_cuda_options[]
Definition: vf_overlay_cuda.c:527
VAR_T
@ VAR_T
Definition: vf_overlay_cuda.c:71
EVAL_MODE_FRAME
@ EVAL_MODE_FRAME
Definition: vf_overlay_cuda.c:77
OverlayCUDAContext
OverlayCUDAContext.
Definition: vf_overlay_cuda.c:97
overlay_cuda_uninit
static av_cold void overlay_cuda_uninit(AVFilterContext *avctx)
Uninitialize overlay_cuda.
Definition: vf_overlay_cuda.c:387
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:126
VAR_X
@ VAR_X
Definition: vf_overlay_cuda.c:67
internal.h
FILTER_SINGLE_PIXFMT
#define FILTER_SINGLE_PIXFMT(pix_fmt_)
Definition: internal.h:180
overlay_cuda_blend
static int overlay_cuda_blend(FFFrameSync *fs)
Perform blend overlay picture over main picture.
Definition: vf_overlay_cuda.c:226
log.h
i
#define i(width, name, range_min, range_max)
Definition: cbs_h2645.c:269
VAR_OVERLAY_H
@ VAR_OVERLAY_H
Definition: vf_overlay_cuda.c:66
OverlayCUDAContext::hwctx
AVCUDADeviceContext * hwctx
Definition: vf_overlay_cuda.c:104
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:109
OverlayCUDAContext::cu_func
CUfunction cu_func
Definition: vf_overlay_cuda.c:108
AVFilterPad::name
const char * name
Pad name.
Definition: internal.h:55
AVFilter
Filter definition.
Definition: avfilter.h:171
AVHWFramesContext
This struct describes a set or pool of "hardware" frames (i.e.
Definition: hwcontext.h:124
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:89
cuda_check.h
OverlayCUDAContext::cu_module
CUmodule cu_module
Definition: vf_overlay_cuda.c:107
pos
unsigned int pos
Definition: spdifenc.c:412
EOF_ACTION_REPEAT
@ EOF_ACTION_REPEAT
Definition: framesync.h:27
AVFrame::height
int height
Definition: frame.h:397
OverlayCUDAContext::y_pexpr
AVExpr * y_pexpr
Definition: vf_overlay_cuda.c:120
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:118
AV_PIX_FMT_NONE
@ AV_PIX_FMT_NONE
Definition: pixfmt.h:65
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:65
AVFilterContext
An instance of a filter.
Definition: avfilter.h:408
var_names
static const char *const var_names[]
Definition: vf_overlay_cuda.c:81
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:410
AV_OPT_TYPE_BOOL
@ AV_OPT_TYPE_BOOL
Definition: opt.h:244
FILTER_OUTPUTS
#define FILTER_OUTPUTS(array)
Definition: internal.h:191
VAR_VARS_NB
@ VAR_VARS_NB
Definition: vf_overlay_cuda.c:72
VAR_N
@ VAR_N
Definition: vf_overlay_cuda.c:69
d
d
Definition: ffmpeg_filter.c:153
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:370
av_log
#define av_log(a,...)
Definition: tableprint_vlc.h:27
eval_expr
static void eval_expr(AVFilterContext *ctx)
Definition: vf_overlay_cuda.c:141
uninit
static av_cold int uninit(AVCodecContext *avctx)
Definition: crystalhd.c:285
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:336
ff_framesync_dualinput_get
int ff_framesync_dualinput_get(FFFrameSync *fs, AVFrame **f0, AVFrame **f1)
Definition: framesync.c:371
AV_OPT_TYPE_STRING
@ AV_OPT_TYPE_STRING
Definition: opt.h:229
OverlayCUDAContext::eval_mode
int eval_mode
Definition: vf_overlay_cuda.c:113
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:2582
AVFilterContext::outputs
AVFilterLink ** outputs
array of pointers to output links
Definition: avfilter.h:420
OverlayCUDAContext::var_values
double var_values[VAR_VARS_NB]
Definition: vf_overlay_cuda.c:117