FFmpeg: libavfilter/vf_pad_cuda.c Source File

FFmpeg
vf_pad_cuda.c
Go to the documentation of this file.
1 /*
2  * This file is part of FFmpeg.
3  *
4  * FFmpeg is free software; you can redistribute it and/or
5  * modify it under the terms of the GNU Lesser General Public
6  * License as published by the Free Software Foundation; either
7  * version 2.1 of the License, or (at your option) any later version.
8  *
9  * FFmpeg is distributed in the hope that it will be useful,
10  * but WITHOUT ANY WARRANTY; without even the implied warranty of
11  * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU
12  * Lesser General Public License for more details.
13  *
14  * You should have received a copy of the GNU Lesser General Public
15  * License along with FFmpeg; if not, write to the Free Software
16  * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA
17  */
18 
19 /**
20  * @file
21  * CUDA video padding filter
22  */
23 
24 #include <float.h>
25 
26 #include "filters.h"
27 #include "libavutil/avstring.h"
28 #include "libavutil/common.h"
29 #include "libavutil/cuda_check.h"
30 #include "libavutil/eval.h"
31 #include "libavutil/hwcontext.h"
32 #include "libavutil/hwcontext_cuda_internal.h"
33 #include "libavutil/imgutils.h"
34 #include "libavutil/internal.h"
35 #include "libavutil/opt.h"
36 #include "libavutil/pixdesc.h"
37 #include "libavutil/colorspace.h"
38 
39 #include "cuda/load_helper.h"
40 
41  #define CHECK_CU(x) FF_CUDA_CHECK_DL(ctx, device_hwctx->internal->cuda_dl, x)
42  #define DIV_UP(a, b) ( ((a) + (b) - 1) / (b) )
43  #define BLOCK_X 32
44  #define BLOCK_Y 16
45 
46  static const enum AVPixelFormat supported_formats[] = {
47  AV_PIX_FMT_YUV420P,
48  AV_PIX_FMT_YUV444P,
49  AV_PIX_FMT_YUVA420P,
50  AV_PIX_FMT_YUVA444P,
51  AV_PIX_FMT_NV12,
52 };
53 
54  typedef struct CUDAPadContext {
55   const AVClass *class;
56 
57   AVBufferRef *frames_ctx;
58 
59   int w, h; ///< output dimensions, a value of 0 will result in the input size
60   int x, y; ///< offsets of the input area with respect to the padded area
61   int in_w, in_h; ///< width and height for the padded input video
62 
63   char *w_expr; ///< width expression
64   char *h_expr; ///< height expression
65   char *x_expr; ///< x offset expression
66   char *y_expr; ///< y offset expression
67 
68   uint8_t rgba_color[4]; ///< color for the padding area
69   uint8_t parsed_color[4];
70   AVRational aspect;
71 
72   int eval_mode;
73 
74   int last_out_w, last_out_h; ///< used to evaluate the prior output width and height with the incoming frame
75 
76   AVCUDADeviceContext *hwctx;
77   CUmodule cu_module;
78   CUfunction cu_func_uchar;
79   CUfunction cu_func_uchar2;
80 } CUDAPadContext;
81 
82  static const char *const var_names[] = {
83  "in_w", "iw",
84  "in_h", "ih",
85  "out_w", "ow",
86  "out_h", "oh",
87  "x",
88  "y",
89  "a",
90  "sar",
91  "dar",
92  "hsub",
93  "vsub",
94  NULL
95 };
96 
97 enum {
98   VAR_IN_W,
99   VAR_IW,
100   VAR_IN_H,
101   VAR_IH,
102   VAR_OUT_W,
103   VAR_OW,
104   VAR_OUT_H,
105   VAR_OH,
106   VAR_X,
107   VAR_Y,
108   VAR_A,
109   VAR_SAR,
110   VAR_DAR,
111   VAR_HSUB,
112   VAR_VSUB,
113   VARS_NB
114 };
115 
116  enum EvalMode {
117   EVAL_MODE_INIT,
118   EVAL_MODE_FRAME,
119   EVAL_MODE_NB
120 };
121 
122  static int eval_expr(AVFilterContext *ctx)
123 {
124  CUDAPadContext *s = ctx->priv;
125  AVFilterLink *inlink = ctx->inputs[0];
126  const AVPixFmtDescriptor *desc = av_pix_fmt_desc_get(inlink->format);
127 
128  double var_values[VARS_NB], res;
129  char *expr;
130  int ret;
131 
132  var_values[VAR_IN_W] = var_values[VAR_IW] = s->in_w;
133  var_values[VAR_IN_H] = var_values[VAR_IH] = s->in_h;
134  var_values[VAR_OUT_W] = var_values[VAR_OW] = NAN;
135  var_values[VAR_OUT_H] = var_values[VAR_OH] = NAN;
136  var_values[VAR_A] = (double)s->in_w / s->in_h;
137  var_values[VAR_SAR] = inlink->sample_aspect_ratio.num ?
138  (double)inlink->sample_aspect_ratio.num /
139  inlink->sample_aspect_ratio.den : 1;
140  var_values[VAR_DAR] = var_values[VAR_A] * var_values[VAR_SAR];
141  var_values[VAR_HSUB] = 1 << desc->log2_chroma_w;
142  var_values[VAR_VSUB] = 1 << desc->log2_chroma_h;
143 
144  expr = s->w_expr;
145  ret = av_expr_parse_and_eval(&res, expr, var_names, var_values, NULL, NULL, NULL, NULL, NULL, 0, ctx);
146  if (ret < 0)
147  goto fail;
148 
149  s->w = res;
150  if (s->w < 0) {
151  av_log(ctx, AV_LOG_ERROR, "Width expression is negative.\n");
152  ret = AVERROR(EINVAL);
153  goto fail;
154  }
155 
156  var_values[VAR_OUT_W] = var_values[VAR_OW] = s->w;
157 
158  expr = s->h_expr;
159  ret = av_expr_parse_and_eval(&res, expr, var_names, var_values, NULL, NULL, NULL, NULL, NULL, 0, ctx);
160  if (ret < 0)
161  goto fail;
162 
163  s->h = res;
164  if (s->h < 0) {
165  av_log(ctx, AV_LOG_ERROR, "Height expression is negative.\n");
166  ret = AVERROR(EINVAL);
167  goto fail;
168  }
169  var_values[VAR_OUT_H] = var_values[VAR_OH] = s->h;
170 
171  if (!s->h)
172  s->h = s->in_h;
173 
174  var_values[VAR_OUT_H] = var_values[VAR_OH] = s->h;
175 
176 
177  expr = s->w_expr;
178  ret = av_expr_parse_and_eval(&res, expr, var_names, var_values, NULL, NULL, NULL, NULL, NULL, 0, ctx);
179  if (ret < 0)
180  goto fail;
181 
182  s->w = res;
183  if (s->w < 0) {
184  av_log(ctx, AV_LOG_ERROR, "Width expression is negative.\n");
185  ret = AVERROR(EINVAL);
186  goto fail;
187  }
188  if (!s->w)
189  s->w = s->in_w;
190 
191  var_values[VAR_OUT_W] = var_values[VAR_OW] = s->w;
192 
193 
194  expr = s->x_expr;
195  ret = av_expr_parse_and_eval(&res, expr, var_names, var_values, NULL, NULL, NULL, NULL, NULL, 0, ctx);
196  if (ret < 0)
197  goto fail;
198 
199  s->x = res;
200 
201 
202  expr = s->y_expr;
203  ret = av_expr_parse_and_eval(&res, expr, var_names, var_values, NULL, NULL, NULL, NULL, NULL, 0, ctx);
204  if (ret < 0)
205  goto fail;
206 
207  s->y = res;
208 
209  if (s->x < 0 || s->x + s->in_w > s->w) {
210  s->x = (s->w - s->in_w) / 2;
211  av_log(ctx, AV_LOG_VERBOSE, "centering X offset.\n");
212  }
213 
214  if (s->y < 0 || s->y + s->in_h > s->h) {
215  s->y = (s->h - s->in_h) / 2;
216  av_log(ctx, AV_LOG_VERBOSE, "centering Y offset.\n");
217  }
218 
219  s->w = av_clip(s->w, 1, INT_MAX);
220  s->h = av_clip(s->h, 1, INT_MAX);
221 
222  if (s->w < s->in_w || s->h < s->in_h) {
223  av_log(ctx, AV_LOG_ERROR, "Padded size < input size.\n");
224  return AVERROR(EINVAL);
225  }
226 
227  av_log(ctx, AV_LOG_DEBUG,
228  "w:%d h:%d -> w:%d h:%d x:%d y:%d color:0x%02X%02X%02X%02X\n",
229  inlink->w, inlink->h, s->w, s->h, s->x, s->y, s->rgba_color[0],
230  s->rgba_color[1], s->rgba_color[2], s->rgba_color[3]);
231 
232  return 0;
233 
234 fail:
235  av_log(ctx, AV_LOG_ERROR, "Error evaluating '%s'\n", expr);
236  return ret;
237 }
238 
239  static int cuda_pad_alloc_out_frames_ctx(AVFilterContext *ctx, AVBufferRef **out_frames_ctx, const int width, const int height)
240 {
241  AVFilterLink *inlink = ctx->inputs[0];
242  FilterLink *inl = ff_filter_link(inlink);
243  AVHWFramesContext *in_frames_ctx = (AVHWFramesContext *)inl->hw_frames_ctx->data;
244  int ret;
245 
246  *out_frames_ctx = av_hwframe_ctx_alloc(in_frames_ctx->device_ref);
247  if (!*out_frames_ctx) {
248  return AVERROR(ENOMEM);
249  }
250 
251  AVHWFramesContext *out_fc = (AVHWFramesContext *)(*out_frames_ctx)->data;
252  out_fc->format = AV_PIX_FMT_CUDA;
253  out_fc->sw_format = in_frames_ctx->sw_format;
254 
255  out_fc->width = FFALIGN(width, 32);
256  out_fc->height = FFALIGN(height, 32);
257 
258  ret = av_hwframe_ctx_init(*out_frames_ctx);
259  if (ret < 0) {
260  av_log(ctx, AV_LOG_ERROR, "Failed to init output ctx\n");
261  av_buffer_unref(out_frames_ctx);
262  return ret;
263  }
264 
265  return 0;
266 }
267 
268  static av_cold int cuda_pad_init(AVFilterContext *ctx)
269 {
270  CUDAPadContext *s = ctx->priv;
271 
272  s->last_out_w = -1;
273  s->last_out_h = -1;
274 
275  return 0;
276 }
277 
278  static av_cold void cuda_pad_uninit(AVFilterContext *ctx)
279 {
280  CUDAPadContext *s = ctx->priv;
281  CUcontext dummy;
282 
283  av_buffer_unref(&s->frames_ctx);
284 
285  if (s->hwctx && s->cu_module) {
286  CudaFunctions *cu = s->hwctx->internal->cuda_dl;
287  AVCUDADeviceContext *device_hwctx = s->hwctx;
288  CHECK_CU(cu->cuCtxPushCurrent(s->hwctx->cuda_ctx));
289  CHECK_CU(cu->cuModuleUnload(s->cu_module));
290  CHECK_CU(cu->cuCtxPopCurrent(&dummy));
291  }
292 
293  s->cu_module = NULL;
294  s->hwctx = NULL;
295 }
296 
297  static av_cold int cuda_pad_load_functions(AVFilterContext *ctx)
298 {
299  CUDAPadContext *s = ctx->priv;
300  CudaFunctions *cu = s->hwctx->internal->cuda_dl;
301  CUcontext dummy_cu_ctx;
302  int ret;
303 
304  AVCUDADeviceContext *device_hwctx = s->hwctx;
305 
306  extern const unsigned char ff_vf_pad_cuda_ptx_data[];
307  extern const unsigned int ff_vf_pad_cuda_ptx_len;
308 
309  ret = CHECK_CU(cu->cuCtxPushCurrent(device_hwctx->cuda_ctx));
310  if (ret < 0)
311  return ret;
312 
313  ret = ff_cuda_load_module(ctx, device_hwctx, &s->cu_module,
314  ff_vf_pad_cuda_ptx_data, ff_vf_pad_cuda_ptx_len);
315  if (ret < 0) {
316  av_log(ctx, AV_LOG_ERROR, "Failed to load CUDA module\n");
317  goto end;
318  }
319 
320  ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_uchar, s->cu_module, "pad_uchar"));
321  if (ret < 0) {
322  av_log(ctx, AV_LOG_ERROR, "Failed to load pad_planar_cuda\n");
323  goto end;
324  }
325 
326  ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_uchar2, s->cu_module, "pad_uchar2"));
327  if (ret < 0)
328  av_log(ctx, AV_LOG_ERROR, "Failed to load pad_uv_cuda\n");
329 
330 end:
331  CHECK_CU(cu->cuCtxPopCurrent(&dummy_cu_ctx));
332 
333  return ret;
334 }
335 
336  static int cuda_pad_config_props(AVFilterLink *outlink)
337 {
338  AVFilterContext *ctx = outlink->src;
339  CUDAPadContext *s = ctx->priv;
340 
341  AVFilterLink *inlink = ctx->inputs[0];
342  FilterLink *inl = ff_filter_link(inlink);
343 
344  FilterLink *ol = ff_filter_link(outlink);
345 
346  AVHWFramesContext *in_frames_ctx;
347  int format_supported = 0;
348  int ret;
349 
350  s->in_w = inlink->w;
351  s->in_h = inlink->h;
352  ret = eval_expr(ctx);
353  if (ret < 0)
354  return ret;
355 
356  if (!inl->hw_frames_ctx) {
357  av_log(ctx, AV_LOG_ERROR, "No hw context provided on input\n");
358  return AVERROR(EINVAL);
359  }
360 
361  in_frames_ctx = (AVHWFramesContext *)inl->hw_frames_ctx->data;
362  s->hwctx = in_frames_ctx->device_ctx->hwctx;
363 
364  for (int i = 0; i < FF_ARRAY_ELEMS(supported_formats); i++) {
365  if (in_frames_ctx->sw_format == supported_formats[i]) {
366  format_supported = 1;
367  break;
368  }
369  }
370  if (!format_supported) {
371  av_log(ctx, AV_LOG_ERROR, "Unsupported input format.\n");
372  return AVERROR(EINVAL);
373  }
374 
375  s->parsed_color[0] = RGB_TO_Y_BT709(s->rgba_color[0], s->rgba_color[1], s->rgba_color[2]);
376  s->parsed_color[1] = RGB_TO_U_BT709(s->rgba_color[0], s->rgba_color[1], s->rgba_color[2], 0);
377  s->parsed_color[2] = RGB_TO_V_BT709(s->rgba_color[0], s->rgba_color[1], s->rgba_color[2], 0);
378  s->parsed_color[3] = s->rgba_color[3];
379 
380  ret = cuda_pad_alloc_out_frames_ctx(ctx, &s->frames_ctx, s->w, s->h);
381  if (ret < 0)
382  return ret;
383 
384  ol->hw_frames_ctx = av_buffer_ref(s->frames_ctx);
385  if (!ol->hw_frames_ctx)
386  return AVERROR(ENOMEM);
387 
388  outlink->w = s->w;
389  outlink->h = s->h;
390  outlink->time_base = inlink->time_base;
391  outlink->format = AV_PIX_FMT_CUDA;
392 
393  s->last_out_w = s->w;
394  s->last_out_h = s->h;
395 
396  ret = cuda_pad_load_functions(ctx);
397  if (ret < 0)
398  return ret;
399 
400  return 0;
401 }
402 
403  static int cuda_pad_pad(AVFilterContext *ctx, AVFrame *out, const AVFrame *in)
404 {
405  CUDAPadContext *s = ctx->priv;
406  FilterLink *inl = ff_filter_link(ctx->inputs[0]);
407 
408  AVHWFramesContext *in_frames_ctx = (AVHWFramesContext *)inl->hw_frames_ctx->data;
409  const AVPixFmtDescriptor *pixdesc = av_pix_fmt_desc_get(in_frames_ctx->sw_format);
410 
411  CudaFunctions *cu = s->hwctx->internal->cuda_dl;
412  AVCUDADeviceContext *device_hwctx = s->hwctx;
413  int ret;
414 
415 
416  const int nb_planes = av_pix_fmt_count_planes(in_frames_ctx->sw_format);
417  for (int plane = 0; plane < nb_planes; plane++) {
418  const AVComponentDescriptor *cur_comp = &pixdesc->comp[0];
419  for (int comp = 1; comp < pixdesc->nb_components && cur_comp->plane != plane; comp++)
420  cur_comp = &pixdesc->comp[comp];
421 
422  int hsub = (plane == 1 || plane == 2) ? pixdesc->log2_chroma_w : 0;
423  int vsub = (plane == 1 || plane == 2) ? pixdesc->log2_chroma_h : 0;
424 
425  int src_w = AV_CEIL_RSHIFT(s->in_w, hsub);
426  int src_h = AV_CEIL_RSHIFT(s->in_h, vsub);
427 
428  int dst_w = AV_CEIL_RSHIFT(s->w, hsub);
429  int dst_h = AV_CEIL_RSHIFT(s->h, vsub);
430 
431  int y_plane_offset = AV_CEIL_RSHIFT(s->y, vsub);
432  int x_plane_offset = AV_CEIL_RSHIFT(s->x, hsub);
433 
434  if (x_plane_offset + src_w > dst_w || y_plane_offset + src_h > dst_h) {
435  av_log(ctx, AV_LOG_ERROR,
436  "ROI out of bounds in plane %d: offset=(%d,%d) in=(%dx%d) "
437  "out=(%dx%d)\n",
438  plane, x_plane_offset, y_plane_offset, src_w, src_h, dst_w, dst_h);
439  return AVERROR(EINVAL);
440  }
441 
442  int dst_linesize = out->linesize[plane] / cur_comp->step;
443  int src_linesize = in->linesize[plane] / cur_comp->step;
444 
445  CUdeviceptr d_dst = (CUdeviceptr)out->data[plane];
446  CUdeviceptr d_src = (CUdeviceptr)in->data[plane];
447 
448  CUfunction cuda_func;
449 
450  if (cur_comp->step == 1 && cur_comp->depth == 8)
451  cuda_func = s->cu_func_uchar;
452  else if(cur_comp->step == 2 && cur_comp->depth == 8)
453  cuda_func = s->cu_func_uchar2;
454  else
455  return AVERROR_BUG;
456 
457  void *kernel_args[] = {
458  &d_dst, &dst_linesize, &dst_w, &dst_h,
459  &d_src, &src_linesize, &src_w, &src_h,
460  &x_plane_offset, &y_plane_offset, &s->parsed_color[plane]
461  };
462 
463  unsigned int grid_x = DIV_UP(dst_w, BLOCK_X);
464  unsigned int grid_y = DIV_UP(dst_h, BLOCK_Y);
465 
466  ret = CHECK_CU(cu->cuLaunchKernel(cuda_func, grid_x, grid_y, 1,
467  BLOCK_X, BLOCK_Y, 1,
468  0, s->hwctx->stream, kernel_args, NULL));
469 
470  if (ret < 0) {
471  av_log(ctx, AV_LOG_ERROR, "Failed to launch kernel for plane %d\n", plane);
472  return ret;
473  }
474  }
475 
476  return 0;
477 }
478 
479  static int cuda_pad_filter_frame(AVFilterLink *inlink, AVFrame *in)
480 {
481  AVFilterContext *ctx = inlink->dst;
482  CUDAPadContext *s = ctx->priv;
483  AVFilterLink *outlink = ctx->outputs[0];
484 
485  FilterLink *outl = ff_filter_link(outlink);
486 
487  AVHWFramesContext *out_frames_ctx = (AVHWFramesContext *)outl->hw_frames_ctx->data;
488  AVCUDADeviceContext *device_hwctx = out_frames_ctx->device_ctx->hwctx;
489 
490  int ret;
491 
492  if (s->eval_mode == EVAL_MODE_FRAME) {
493  s->in_w = in->width;
494  s->in_h = in->height;
495  s->aspect = in->sample_aspect_ratio;
496 
497  ret = eval_expr(ctx);
498  if (ret < 0) {
499  av_frame_free(&in);
500  return ret;
501  }
502  }
503 
504 
505  if (s->x == 0 && s->y == 0 &&
506  s->w == in->width && s->h == in->height) {
507  av_log(ctx, AV_LOG_DEBUG, "No border. Passing the frame unmodified.\n");
508  s->last_out_w = s->w;
509  s->last_out_h = s->h;
510  return ff_filter_frame(outlink, in);
511  }
512 
513 
514  if (s->w != s->last_out_w || s->h != s->last_out_h) {
515 
516  av_buffer_unref(&s->frames_ctx);
517 
518  ret = cuda_pad_alloc_out_frames_ctx(ctx, &s->frames_ctx, s->w, s->h);
519  if (ret < 0)
520  return ret;
521 
522  av_buffer_unref(&outl->hw_frames_ctx);
523  outl->hw_frames_ctx = av_buffer_ref(s->frames_ctx);
524  if (!outl->hw_frames_ctx) {
525  av_frame_free(&in);
526  av_log(ctx, AV_LOG_ERROR, "Failed to allocate output frame context.\n");
527  return AVERROR(ENOMEM);
528  }
529  outlink->w = s->w;
530  outlink->h = s->h;
531 
532  s->last_out_w = s->w;
533  s->last_out_h = s->h;
534  }
535 
536  AVFrame *out = av_frame_alloc();
537  if (!out) {
538  av_frame_free(&in);
539  av_log(ctx, AV_LOG_ERROR, "Failed to allocate output AVFrame.\n");
540  return AVERROR(ENOMEM);
541  }
542  ret = av_hwframe_get_buffer(outl->hw_frames_ctx, out, 0);
543  if (ret < 0) {
544  av_log(ctx, AV_LOG_ERROR, "Unable to get output buffer: %s\n",
545  av_err2str(ret));
546  av_frame_free(&out);
547  av_frame_free(&in);
548  return ret;
549  }
550 
551  CUcontext dummy;
552  ret = CHECK_CU(device_hwctx->internal->cuda_dl->cuCtxPushCurrent(
553  device_hwctx->cuda_ctx));
554  if (ret < 0) {
555  av_frame_free(&out);
556  av_frame_free(&in);
557  return ret;
558  }
559 
560  ret = cuda_pad_pad(ctx, out, in);
561 
562  CHECK_CU(device_hwctx->internal->cuda_dl->cuCtxPopCurrent(&dummy));
563 
564  if (ret < 0) {
565  av_frame_free(&out);
566  av_frame_free(&in);
567  return ret;
568  }
569 
570  av_frame_copy_props(out, in);
571  out->width = s->w;
572  out->height = s->h;
573 
574 
575  av_reduce(&out->sample_aspect_ratio.num, &out->sample_aspect_ratio.den,
576  (int64_t)in->sample_aspect_ratio.num * out->height * in->width,
577  (int64_t)in->sample_aspect_ratio.den * out->width * in->height,
578  INT_MAX);
579 
580  av_frame_free(&in);
581  return ff_filter_frame(outlink, out);
582 }
583 
584  #define OFFSET(x) offsetof(CUDAPadContext, x)
585  #define FLAGS (AV_OPT_FLAG_FILTERING_PARAM | AV_OPT_FLAG_VIDEO_PARAM)
586 
587  static const AVOption cuda_pad_options[] = {
588  { "width", "set the pad area width expression", OFFSET(w_expr), AV_OPT_TYPE_STRING, {.str = "iw"}, 0, 0, FLAGS },
589  { "w", "set the pad area width expression", OFFSET(w_expr), AV_OPT_TYPE_STRING, {.str = "iw"}, 0, 0, FLAGS },
590  { "height", "set the pad area height expression", OFFSET(h_expr), AV_OPT_TYPE_STRING, {.str = "ih"}, 0, 0, FLAGS },
591  { "h", "set the pad area height expression", OFFSET(h_expr), AV_OPT_TYPE_STRING, {.str = "ih"}, 0, 0, FLAGS },
592  { "x", "set the x offset expression for the input image position", OFFSET(x_expr), AV_OPT_TYPE_STRING, {.str = "0"}, 0, 0, FLAGS },
593  { "y", "set the y offset expression for the input image position", OFFSET(y_expr), AV_OPT_TYPE_STRING, {.str = "0"}, 0, 0, FLAGS },
594  { "color", "set the color of the padded area border", OFFSET(rgba_color), AV_OPT_TYPE_COLOR, {.str = "black"}, .flags = FLAGS },
595  { "eval", "specify when to evaluate expressions", OFFSET(eval_mode), AV_OPT_TYPE_INT, {.i64 = EVAL_MODE_INIT}, 0, EVAL_MODE_NB-1, FLAGS, .unit = "eval" },
596  { "init", "eval expressions once during initialization", 0, AV_OPT_TYPE_CONST, {.i64=EVAL_MODE_INIT}, .flags = FLAGS, .unit = "eval" },
597  { "frame", "eval expressions during initialization and per-frame", 0, AV_OPT_TYPE_CONST, {.i64=EVAL_MODE_FRAME}, .flags = FLAGS, .unit = "eval" },
598  { "aspect", "pad to fit an aspect instead of a resolution", OFFSET(aspect), AV_OPT_TYPE_RATIONAL, {.dbl = 0}, 0, DBL_MAX, FLAGS },
599  { NULL }
600 };
601 
602  static const AVClass cuda_pad_class = {
603  .class_name = "pad_cuda",
604  .item_name = av_default_item_name,
605  .option = cuda_pad_options,
606  .version = LIBAVUTIL_VERSION_INT,
607 };
608 
609  static const AVFilterPad cuda_pad_inputs[] = {{
610  .name = "default",
611  .type = AVMEDIA_TYPE_VIDEO,
612  .filter_frame = cuda_pad_filter_frame
613 }};
614 
615  static const AVFilterPad cuda_pad_outputs[] = {{
616  .name = "default",
617  .type = AVMEDIA_TYPE_VIDEO,
618  .config_props = cuda_pad_config_props,
619 }};
620 
621  const FFFilter ff_vf_pad_cuda = {
622  .p.name = "pad_cuda",
623  .p.description = NULL_IF_CONFIG_SMALL("CUDA-based GPU padding filter"),
624  .init = cuda_pad_init,
625  .uninit = cuda_pad_uninit,
626 
627  .p.priv_class = &cuda_pad_class,
628 
629  FILTER_INPUTS(cuda_pad_inputs),
630  FILTER_OUTPUTS(cuda_pad_outputs),
631 
632  FILTER_SINGLE_PIXFMT(AV_PIX_FMT_CUDA),
633 
634  .priv_size = sizeof(CUDAPadContext),
635  .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE,
636 };
VAR_X
@ VAR_X
Definition: vf_pad_cuda.c:106
CUDAPadContext::in_h
int in_h
width and height for the padded input video
Definition: vf_pad_cuda.c:61
AVHWDeviceContext::hwctx
void * hwctx
The format-specific data, allocated and freed by libavutil along with this context.
Definition: hwcontext.h:88
AV_PIX_FMT_CUDA
@ AV_PIX_FMT_CUDA
HW acceleration through CUDA.
Definition: pixfmt.h:260
AVPixelFormat
AVPixelFormat
Pixel format.
Definition: pixfmt.h:71
VAR_IH
@ VAR_IH
Definition: vf_pad_cuda.c:101
av_clip
#define av_clip
Definition: common.h:100
cuda_pad_outputs
static const AVFilterPad cuda_pad_outputs[]
Definition: vf_pad_cuda.c:615
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
hwcontext_cuda_internal.h
out
FILE * out
Definition: movenc.c:55
VAR_OUT_H
@ VAR_OUT_H
Definition: vf_pad_cuda.c:104
comp
static void comp(unsigned char *dst, ptrdiff_t dst_stride, unsigned char *src, ptrdiff_t src_stride, int add)
Definition: eamad.c:79
ff_filter_frame
int ff_filter_frame(AVFilterLink *link, AVFrame *frame)
Send a frame of data to the next filter.
Definition: avfilter.c:1067
av_pix_fmt_desc_get
const AVPixFmtDescriptor * av_pix_fmt_desc_get(enum AVPixelFormat pix_fmt)
Definition: pixdesc.c:3456
AVBufferRef::data
uint8_t * data
The data buffer.
Definition: buffer.h:90
AVHWFramesContext::format
enum AVPixelFormat format
The pixel format identifying the underlying HW surface type.
Definition: hwcontext.h:200
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
CUDAPadContext::w_expr
char * w_expr
width expression
Definition: vf_pad_cuda.c:63
CHECK_CU
#define CHECK_CU(x)
Definition: vf_pad_cuda.c:41
av_frame_free
void av_frame_free(AVFrame **frame)
Free the frame and any dynamically allocated objects in it, e.g.
Definition: frame.c:64
av_hwframe_ctx_init
int av_hwframe_ctx_init(AVBufferRef *ref)
Finalize the context before use.
Definition: hwcontext.c:337
FILTER_INPUTS
#define FILTER_INPUTS(array)
Definition: filters.h:263
AVFrame
This structure describes decoded (raw) audio or video data.
Definition: frame.h:427
pixdesc.h
AVFrame::width
int width
Definition: frame.h:499
av_hwframe_ctx_alloc
AVBufferRef * av_hwframe_ctx_alloc(AVBufferRef *device_ref_in)
Allocate an AVHWFramesContext tied to a given device context.
Definition: hwcontext.c:263
DIV_UP
#define DIV_UP(a, b)
Definition: vf_pad_cuda.c:42
AVComponentDescriptor::depth
int depth
Number of bits in the component.
Definition: pixdesc.h:57
CUDAPadContext::h
int h
output dimensions, a value of 0 will result in the input size
Definition: vf_pad_cuda.c:59
AVOption
AVOption.
Definition: opt.h:429
VAR_HSUB
@ VAR_HSUB
Definition: vf_pad_cuda.c:111
VAR_A
@ VAR_A
Definition: vf_pad_cuda.c:108
AVComponentDescriptor::step
int step
Number of elements between 2 horizontally consecutive pixels.
Definition: pixdesc.h:40
AV_LOG_VERBOSE
#define AV_LOG_VERBOSE
Detailed information.
Definition: log.h:226
float.h
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:220
AV_OPT_TYPE_RATIONAL
@ AV_OPT_TYPE_RATIONAL
Underlying C type is AVRational.
Definition: opt.h:280
AVHWFramesContext::width
int width
The allocated dimensions of the frames in this pool.
Definition: hwcontext.h:220
CUDAPadContext::x
int x
Definition: vf_pad_cuda.c:60
CUDAPadContext::h_expr
char * h_expr
height expression
Definition: vf_pad_cuda.c:64
cuda_pad_load_functions
static av_cold int cuda_pad_load_functions(AVFilterContext *ctx)
Definition: vf_pad_cuda.c:297
eval_expr
static int eval_expr(AVFilterContext *ctx)
Definition: vf_pad_cuda.c:122
AVFrame::data
uint8_t * data[AV_NUM_DATA_POINTERS]
pointer to the picture/channel planes.
Definition: frame.h:448
VAR_IN_W
@ VAR_IN_W
Definition: vf_pad_cuda.c:98
CUDAPadContext::y_expr
char * y_expr
y offset expression
Definition: vf_pad_cuda.c:66
hsub
static void hsub(htype *dst, const htype *src, int bins)
Definition: vf_median.c:74
av_pix_fmt_count_planes
int av_pix_fmt_count_planes(enum AVPixelFormat pix_fmt)
Definition: pixdesc.c:3496
CUDAPadContext::eval_mode
int eval_mode
Definition: vf_pad_cuda.c:72
fail
#define fail()
Definition: checkasm.h:207
dummy
int dummy
Definition: motion.c:66
cuda_pad_inputs
static const AVFilterPad cuda_pad_inputs[]
Definition: vf_pad_cuda.c:609
cuda_pad_options
static const AVOption cuda_pad_options[]
Definition: vf_pad_cuda.c:587
AVCUDADeviceContext::cuda_ctx
CUcontext cuda_ctx
Definition: hwcontext_cuda.h:43
av_reduce
int av_reduce(int *dst_num, int *dst_den, int64_t num, int64_t den, int64_t max)
Reduce a fraction.
Definition: rational.c:35
AVRational::num
int num
Numerator.
Definition: rational.h:59
cuda_pad_alloc_out_frames_ctx
static int cuda_pad_alloc_out_frames_ctx(AVFilterContext *ctx, AVBufferRef **out_frames_ctx, const int width, const int height)
Definition: vf_pad_cuda.c:239
AVFilterPad
A filter pad used for either input or output.
Definition: filters.h:39
VAR_SAR
@ VAR_SAR
Definition: vf_pad_cuda.c:109
av_frame_alloc
AVFrame * av_frame_alloc(void)
Allocate an AVFrame and set its fields to default values.
Definition: frame.c:52
colorspace.h
AV_LOG_ERROR
#define AV_LOG_ERROR
Something went wrong and cannot losslessly be recovered.
Definition: log.h:210
FF_ARRAY_ELEMS
#define FF_ARRAY_ELEMS(a)
Definition: sinewin_tablegen.c:29
av_cold
#define av_cold
Definition: attributes.h:106
cuda_pad_pad
static int cuda_pad_pad(AVFilterContext *ctx, AVFrame *out, const AVFrame *in)
Definition: vf_pad_cuda.c:403
AVHWFramesContext::height
int height
Definition: hwcontext.h:220
FFFilter
Definition: filters.h:266
CUDAPadContext::y
int y
offsets of the input area with respect to the padded area
Definition: vf_pad_cuda.c:60
RGB_TO_Y_BT709
#define RGB_TO_Y_BT709(r, g, b)
Definition: vf_pseudocolor.c:674
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_CEIL_RSHIFT
#define AV_CEIL_RSHIFT(a, b)
Definition: common.h:60
EVAL_MODE_NB
@ EVAL_MODE_NB
Definition: vf_pad_cuda.c:119
RGB_TO_U_BT709
#define RGB_TO_U_BT709(r1, g1, b1, max)
Definition: vf_pseudocolor.c:678
FLAGS
#define FLAGS
Definition: vf_pad_cuda.c:585
cuda_pad_class
static const AVClass cuda_pad_class
Definition: vf_pad_cuda.c:602
filters.h
AV_LOG_DEBUG
#define AV_LOG_DEBUG
Stuff which is only useful for libav* developers.
Definition: log.h:231
ctx
AVFormatContext * ctx
Definition: movenc.c:49
EVAL_MODE_FRAME
@ EVAL_MODE_FRAME
Definition: vf_pad_cuda.c:118
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
FILTER_OUTPUTS
#define FILTER_OUTPUTS(array)
Definition: filters.h:264
NAN
#define NAN
Definition: mathematics.h:115
BLOCK_X
#define BLOCK_X
Definition: vf_pad_cuda.c:43
if
if(ret)
Definition: filter_design.txt:179
cuda_pad_config_props
static int cuda_pad_config_props(AVFilterLink *outlink)
Definition: vf_pad_cuda.c:336
LIBAVUTIL_VERSION_INT
#define LIBAVUTIL_VERSION_INT
Definition: version.h:85
AVClass
Describe the class of an AVClass context structure.
Definition: log.h:76
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:213
av_frame_copy_props
int av_frame_copy_props(AVFrame *dst, const AVFrame *src)
Copy only "metadata" fields from src to dst.
Definition: frame.c:599
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
AVComponentDescriptor
Definition: pixdesc.h:30
AVRational
Rational number (pair of numerator and denominator).
Definition: rational.h:58
AVHWFramesContext::device_ref
AVBufferRef * device_ref
A reference to the parent AVHWDeviceContext.
Definition: hwcontext.h:129
EVAL_MODE_INIT
@ EVAL_MODE_INIT
Definition: vf_pad_cuda.c:117
AV_OPT_TYPE_COLOR
@ AV_OPT_TYPE_COLOR
Underlying C type is uint8_t[4].
Definition: opt.h:323
AVComponentDescriptor::plane
int plane
Which of the 4 planes contains the component.
Definition: pixdesc.h:34
av_default_item_name
const char * av_default_item_name(void *ptr)
Return the context name.
Definition: log.c:241
ff_vf_pad_cuda
const FFFilter ff_vf_pad_cuda
Definition: vf_pad_cuda.c:621
double
double
Definition: af_crystalizer.c:132
CUDAPadContext::last_out_w
int last_out_w
Definition: vf_pad_cuda.c:74
CUDAPadContext::rgba_color
uint8_t rgba_color[4]
color for the padding area
Definition: vf_pad_cuda.c:68
ff_filter_link
static FilterLink * ff_filter_link(AVFilterLink *link)
Definition: filters.h:198
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:207
eval.h
CUDAPadContext::last_out_h
int last_out_h
used to evaluate the prior output width and height with the incoming frame
Definition: vf_pad_cuda.c:74
RGB_TO_V_BT709
#define RGB_TO_V_BT709(r1, g1, b1, max)
Definition: vf_pseudocolor.c:682
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
av_expr_parse_and_eval
int av_expr_parse_and_eval(double *d, const char *s, const char *const *const_names, const double *const_values, const char *const *func1_names, double(*const *funcs1)(void *, double), const char *const *func2_names, double(*const *funcs2)(void *, double, double), void *opaque, int log_offset, void *log_ctx)
Parse and evaluate an expression.
Definition: eval.c:805
height
#define height
Definition: dsp.h:89
CUDAPadContext::w
int w
Definition: vf_pad_cuda.c:59
av_err2str
#define av_err2str(errnum)
Convenience macro, the return value should be used only directly in function arguments but never stan...
Definition: error.h:122
for
for(k=2;k<=8;++k)
Definition: h264pred_template.c:424
CUDAPadContext::in_w
int in_w
Definition: vf_pad_cuda.c:61
CUDAPadContext::cu_func_uchar2
CUfunction cu_func_uchar2
Definition: vf_pad_cuda.c:79
CUDAPadContext
Definition: vf_pad_cuda.c:54
VAR_OUT_W
@ VAR_OUT_W
Definition: vf_pad_cuda.c:102
AV_PIX_FMT_YUVA444P
@ AV_PIX_FMT_YUVA444P
planar YUV 4:4:4 32bpp, (1 Cr & Cb sample per 1x1 Y & A samples)
Definition: pixfmt.h:174
cuda_pad_filter_frame
static int cuda_pad_filter_frame(AVFilterLink *inlink, AVFrame *in)
Definition: vf_pad_cuda.c:479
supported_formats
static enum AVPixelFormat supported_formats[]
Definition: vf_pad_cuda.c:46
var_names
static const char *const var_names[]
Definition: vf_pad_cuda.c:82
CUDAPadContext::hwctx
AVCUDADeviceContext * hwctx
Definition: vf_pad_cuda.c:76
i
#define i(width, name, range_min, range_max)
Definition: cbs_h2645.c:256
internal.h
common.h
EvalMode
EvalMode
Definition: af_volume.h:39
CUDAPadContext::parsed_color
uint8_t parsed_color[4]
Definition: vf_pad_cuda.c:69
AVFilterPad::name
const char * name
Pad name.
Definition: filters.h:45
BLOCK_Y
#define BLOCK_Y
Definition: vf_pad_cuda.c:44
CUDAPadContext::frames_ctx
AVBufferRef * frames_ctx
Definition: vf_pad_cuda.c:57
AVHWFramesContext
This struct describes a set or pool of "hardware" frames (i.e.
Definition: hwcontext.h:118
AVCUDADeviceContext
This struct is allocated as AVHWDeviceContext.hwctx.
Definition: hwcontext_cuda.h:42
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
AVClass::class_name
const char * class_name
The name of the class; usually it is the same name as the context structure type to which the AVClass...
Definition: log.h:81
AVHWFramesContext::device_ctx
AVHWDeviceContext * device_ctx
The parent AVHWDeviceContext.
Definition: hwcontext.h:137
cuda_check.h
VAR_OW
@ VAR_OW
Definition: vf_pad_cuda.c:103
AVFrame::sample_aspect_ratio
AVRational sample_aspect_ratio
Sample aspect ratio for the video frame, 0/1 if unknown/unspecified.
Definition: frame.h:524
AVFrame::height
int height
Definition: frame.h:499
VAR_OH
@ VAR_OH
Definition: vf_pad_cuda.c:105
OFFSET
#define OFFSET(x)
Definition: vf_pad_cuda.c:584
AVRational::den
int den
Denominator.
Definition: rational.h:60
AV_OPT_TYPE_INT
@ AV_OPT_TYPE_INT
Underlying C type is int.
Definition: opt.h:259
VAR_IN_H
@ VAR_IN_H
Definition: vf_pad_cuda.c:100
CUDAPadContext::x_expr
char * x_expr
x offset expression
Definition: vf_pad_cuda.c:65
AV_PIX_FMT_YUV444P
@ AV_PIX_FMT_YUV444P
planar YUV 4:4:4, 24bpp, (1 Cr & Cb sample per 1x1 Y samples)
Definition: pixfmt.h:78
AVFilterContext
An instance of a filter.
Definition: avfilter.h:274
VARS_NB
@ VARS_NB
Definition: vf_pad_cuda.c:113
desc
const char * desc
Definition: libsvtav1.c:78
AVMEDIA_TYPE_VIDEO
@ AVMEDIA_TYPE_VIDEO
Definition: avutil.h:200
FFFilter::p
AVFilter p
The public AVFilter.
Definition: filters.h:270
AVBufferRef
A reference to a data buffer.
Definition: buffer.h:82
cuda_pad_uninit
static av_cold void cuda_pad_uninit(AVFilterContext *ctx)
Definition: vf_pad_cuda.c:278
AVPixFmtDescriptor
Descriptor that unambiguously describes how the bits of a pixel are stored in the up to 4 data planes...
Definition: pixdesc.h:69
VAR_VSUB
@ VAR_VSUB
Definition: vf_pad_cuda.c:112
FFALIGN
#define FFALIGN(x, a)
Definition: macros.h:78
cuda_pad_init
static av_cold int cuda_pad_init(AVFilterContext *ctx)
Definition: vf_pad_cuda.c:268
VAR_DAR
@ VAR_DAR
Definition: vf_pad_cuda.c:110
imgutils.h
CUDAPadContext::cu_func_uchar
CUfunction cu_func_uchar
Definition: vf_pad_cuda.c:78
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:472
av_log
#define av_log(a,...)
Definition: tableprint_vlc.h:27
avstring.h
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
width
#define width
Definition: dsp.h:89
VAR_Y
@ VAR_Y
Definition: vf_pad_cuda.c:107
av_hwframe_get_buffer
int av_hwframe_get_buffer(AVBufferRef *hwframe_ref, AVFrame *frame, int flags)
Allocate a new frame attached to the given AVHWFramesContext.
Definition: hwcontext.c:506
CUDAPadContext::aspect
AVRational aspect
Definition: vf_pad_cuda.c:70
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:253
CUDAPadContext::cu_module
CUmodule cu_module
Definition: vf_pad_cuda.c:77
VAR_IW
@ VAR_IW
Definition: vf_pad_cuda.c:99
Generated on Wed Nov 19 2025 19:23:01 for FFmpeg by   doxygen 1.8.17

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