FFmpeg: libavfilter/vf_scale_cuda.c Source File

FFmpeg
vf_scale_cuda.c
Go to the documentation of this file.
1 /*
2 * Copyright (c) 2017, NVIDIA CORPORATION. All rights reserved.
3 *
4 * Permission is hereby granted, free of charge, to any person obtaining a
5 * copy of this software and associated documentation files (the "Software"),
6 * to deal in the Software without restriction, including without limitation
7 * the rights to use, copy, modify, merge, publish, distribute, sublicense,
8 * and/or sell copies of the Software, and to permit persons to whom the
9 * Software is furnished to do so, subject to the following conditions:
10 *
11 * The above copyright notice and this permission notice shall be included in
12 * all copies or substantial portions of the Software.
13 *
14 * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
15 * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
16 * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL
17 * THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
18 * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
19 * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER
20 * DEALINGS IN THE SOFTWARE.
21 */
22 
23 #include <float.h>
24 #include <stdio.h>
25 
26 #include "libavutil/common.h"
27 #include "libavutil/hwcontext.h"
28 #include "libavutil/hwcontext_cuda_internal.h"
29 #include "libavutil/cuda_check.h"
30 #include "libavutil/internal.h"
31 #include "libavutil/opt.h"
32 #include "libavutil/pixdesc.h"
33 
34 #include "avfilter.h"
35 #include "internal.h"
36 #include "scale_eval.h"
37 #include "video.h"
38 
39 #include "cuda/load_helper.h"
40 #include "vf_scale_cuda.h"
41 
42  static const enum AVPixelFormat supported_formats[] = {
43  AV_PIX_FMT_YUV420P,
44  AV_PIX_FMT_NV12,
45  AV_PIX_FMT_YUV444P,
46  AV_PIX_FMT_P010,
47  AV_PIX_FMT_P016,
48  AV_PIX_FMT_YUV444P16,
49  AV_PIX_FMT_0RGB32,
50  AV_PIX_FMT_0BGR32,
51  AV_PIX_FMT_RGB32,
52  AV_PIX_FMT_BGR32,
53 };
54 
55  #define DIV_UP(a, b) ( ((a) + (b) - 1) / (b) )
56  #define BLOCKX 32
57  #define BLOCKY 16
58 
59  #define CHECK_CU(x) FF_CUDA_CHECK_DL(ctx, s->hwctx->internal->cuda_dl, x)
60 
61 enum {
62   INTERP_ALGO_DEFAULT,
63 
64   INTERP_ALGO_NEAREST,
65   INTERP_ALGO_BILINEAR,
66   INTERP_ALGO_BICUBIC,
67   INTERP_ALGO_LANCZOS,
68 
69   INTERP_ALGO_COUNT
70 };
71 
72  typedef struct CUDAScaleContext {
73   const AVClass *class;
74 
75   AVCUDADeviceContext *hwctx;
76 
77   enum AVPixelFormat in_fmt, out_fmt;
78   const AVPixFmtDescriptor *in_desc, *out_desc;
79   int in_planes, out_planes;
80   int in_plane_depths[4];
81   int in_plane_channels[4];
82 
83   AVBufferRef *frames_ctx;
84   AVFrame *frame;
85 
86   AVFrame *tmp_frame;
87   int passthrough;
88 
89  /**
90  * Output sw format. AV_PIX_FMT_NONE for no conversion.
91  */
92   enum AVPixelFormat format;
93 
94   char *w_expr; ///< width expression string
95   char *h_expr; ///< height expression string
96 
97   int force_original_aspect_ratio;
98   int force_divisible_by;
99 
100   CUcontext cu_ctx;
101   CUmodule cu_module;
102   CUfunction cu_func;
103   CUfunction cu_func_uv;
104   CUstream cu_stream;
105 
106   int interp_algo;
107   int interp_use_linear;
108   int interp_as_integer;
109 
110   float param;
111 } CUDAScaleContext;
112 
113  static av_cold int cudascale_init(AVFilterContext *ctx)
114 {
115  CUDAScaleContext *s = ctx->priv;
116 
117  s->frame = av_frame_alloc();
118  if (!s->frame)
119  return AVERROR(ENOMEM);
120 
121  s->tmp_frame = av_frame_alloc();
122  if (!s->tmp_frame)
123  return AVERROR(ENOMEM);
124 
125  return 0;
126 }
127 
128  static av_cold void cudascale_uninit(AVFilterContext *ctx)
129 {
130  CUDAScaleContext *s = ctx->priv;
131 
132  if (s->hwctx && s->cu_module) {
133  CudaFunctions *cu = s->hwctx->internal->cuda_dl;
134  CUcontext dummy;
135 
136  CHECK_CU(cu->cuCtxPushCurrent(s->hwctx->cuda_ctx));
137  CHECK_CU(cu->cuModuleUnload(s->cu_module));
138  s->cu_module = NULL;
139  CHECK_CU(cu->cuCtxPopCurrent(&dummy));
140  }
141 
142  av_frame_free(&s->frame);
143  av_buffer_unref(&s->frames_ctx);
144  av_frame_free(&s->tmp_frame);
145 }
146 
147  static av_cold int init_hwframe_ctx(CUDAScaleContext *s, AVBufferRef *device_ctx, int width, int height)
148 {
149  AVBufferRef *out_ref = NULL;
150  AVHWFramesContext *out_ctx;
151  int ret;
152 
153  out_ref = av_hwframe_ctx_alloc(device_ctx);
154  if (!out_ref)
155  return AVERROR(ENOMEM);
156  out_ctx = (AVHWFramesContext*)out_ref->data;
157 
158  out_ctx->format = AV_PIX_FMT_CUDA;
159  out_ctx->sw_format = s->out_fmt;
160  out_ctx->width = FFALIGN(width, 32);
161  out_ctx->height = FFALIGN(height, 32);
162 
163  ret = av_hwframe_ctx_init(out_ref);
164  if (ret < 0)
165  goto fail;
166 
167  av_frame_unref(s->frame);
168  ret = av_hwframe_get_buffer(out_ref, s->frame, 0);
169  if (ret < 0)
170  goto fail;
171 
172  s->frame->width = width;
173  s->frame->height = height;
174 
175  av_buffer_unref(&s->frames_ctx);
176  s->frames_ctx = out_ref;
177 
178  return 0;
179 fail:
180  av_buffer_unref(&out_ref);
181  return ret;
182 }
183 
184  static int format_is_supported(enum AVPixelFormat fmt)
185 {
186  int i;
187 
188  for (i = 0; i < FF_ARRAY_ELEMS(supported_formats); i++)
189  if (supported_formats[i] == fmt)
190  return 1;
191  return 0;
192 }
193 
194  static av_cold void set_format_info(AVFilterContext *ctx, enum AVPixelFormat in_format, enum AVPixelFormat out_format)
195 {
196  CUDAScaleContext *s = ctx->priv;
197  int i, p, d;
198 
199  s->in_fmt = in_format;
200  s->out_fmt = out_format;
201 
202  s->in_desc = av_pix_fmt_desc_get(s->in_fmt);
203  s->out_desc = av_pix_fmt_desc_get(s->out_fmt);
204  s->in_planes = av_pix_fmt_count_planes(s->in_fmt);
205  s->out_planes = av_pix_fmt_count_planes(s->out_fmt);
206 
207  // find maximum step of each component of each plane
208  // For our subset of formats, this should accurately tell us how many channels CUDA needs
209  // i.e. 1 for Y plane, 2 for UV plane of NV12, 4 for single plane of RGB0 formats
210 
211  for (i = 0; i < s->in_desc->nb_components; i++) {
212  d = (s->in_desc->comp[i].depth + 7) / 8;
213  p = s->in_desc->comp[i].plane;
214  s->in_plane_channels[p] = FFMAX(s->in_plane_channels[p], s->in_desc->comp[i].step / d);
215 
216  s->in_plane_depths[p] = s->in_desc->comp[i].depth;
217  }
218 }
219 
220  static av_cold int init_processing_chain(AVFilterContext *ctx, int in_width, int in_height,
221  int out_width, int out_height)
222 {
223  CUDAScaleContext *s = ctx->priv;
224 
225  AVHWFramesContext *in_frames_ctx;
226 
227  enum AVPixelFormat in_format;
228  enum AVPixelFormat out_format;
229  int ret;
230 
231  /* check that we have a hw context */
232  if (!ctx->inputs[0]->hw_frames_ctx) {
233  av_log(ctx, AV_LOG_ERROR, "No hw context provided on input\n");
234  return AVERROR(EINVAL);
235  }
236  in_frames_ctx = (AVHWFramesContext*)ctx->inputs[0]->hw_frames_ctx->data;
237  in_format = in_frames_ctx->sw_format;
238  out_format = (s->format == AV_PIX_FMT_NONE) ? in_format : s->format;
239 
240  if (!format_is_supported(in_format)) {
241  av_log(ctx, AV_LOG_ERROR, "Unsupported input format: %s\n",
242  av_get_pix_fmt_name(in_format));
243  return AVERROR(ENOSYS);
244  }
245  if (!format_is_supported(out_format)) {
246  av_log(ctx, AV_LOG_ERROR, "Unsupported output format: %s\n",
247  av_get_pix_fmt_name(out_format));
248  return AVERROR(ENOSYS);
249  }
250 
251  set_format_info(ctx, in_format, out_format);
252 
253  if (s->passthrough && in_width == out_width && in_height == out_height && in_format == out_format) {
254  s->frames_ctx = av_buffer_ref(ctx->inputs[0]->hw_frames_ctx);
255  if (!s->frames_ctx)
256  return AVERROR(ENOMEM);
257  } else {
258  s->passthrough = 0;
259 
260  ret = init_hwframe_ctx(s, in_frames_ctx->device_ref, out_width, out_height);
261  if (ret < 0)
262  return ret;
263 
264  if (in_width == out_width && in_height == out_height &&
265  in_format == out_format && s->interp_algo == INTERP_ALGO_DEFAULT)
266  s->interp_algo = INTERP_ALGO_NEAREST;
267  }
268 
269  ctx->outputs[0]->hw_frames_ctx = av_buffer_ref(s->frames_ctx);
270  if (!ctx->outputs[0]->hw_frames_ctx)
271  return AVERROR(ENOMEM);
272 
273  return 0;
274 }
275 
276  static av_cold int cudascale_load_functions(AVFilterContext *ctx)
277 {
278  CUDAScaleContext *s = ctx->priv;
279  CUcontext dummy, cuda_ctx = s->hwctx->cuda_ctx;
280  CudaFunctions *cu = s->hwctx->internal->cuda_dl;
281  char buf[128];
282  int ret;
283 
284  const char *in_fmt_name = av_get_pix_fmt_name(s->in_fmt);
285  const char *out_fmt_name = av_get_pix_fmt_name(s->out_fmt);
286 
287  const char *function_infix = "";
288 
289  extern const unsigned char ff_vf_scale_cuda_ptx_data[];
290  extern const unsigned int ff_vf_scale_cuda_ptx_len;
291 
292  switch(s->interp_algo) {
293  case INTERP_ALGO_NEAREST:
294  function_infix = "Nearest";
295  s->interp_use_linear = 0;
296  s->interp_as_integer = 1;
297  break;
298  case INTERP_ALGO_BILINEAR:
299  function_infix = "Bilinear";
300  s->interp_use_linear = 1;
301  s->interp_as_integer = 1;
302  break;
303  case INTERP_ALGO_DEFAULT:
304  case INTERP_ALGO_BICUBIC:
305  function_infix = "Bicubic";
306  s->interp_use_linear = 0;
307  s->interp_as_integer = 0;
308  break;
309  case INTERP_ALGO_LANCZOS:
310  function_infix = "Lanczos";
311  s->interp_use_linear = 0;
312  s->interp_as_integer = 0;
313  break;
314  default:
315  av_log(ctx, AV_LOG_ERROR, "Unknown interpolation algorithm\n");
316  return AVERROR_BUG;
317  }
318 
319  ret = CHECK_CU(cu->cuCtxPushCurrent(cuda_ctx));
320  if (ret < 0)
321  return ret;
322 
323  ret = ff_cuda_load_module(ctx, s->hwctx, &s->cu_module,
324  ff_vf_scale_cuda_ptx_data, ff_vf_scale_cuda_ptx_len);
325  if (ret < 0)
326  goto fail;
327 
328  snprintf(buf, sizeof(buf), "Subsample_%s_%s_%s", function_infix, in_fmt_name, out_fmt_name);
329  ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_func, s->cu_module, buf));
330  if (ret < 0) {
331  av_log(ctx, AV_LOG_FATAL, "Unsupported conversion: %s -> %s\n", in_fmt_name, out_fmt_name);
332  ret = AVERROR(ENOSYS);
333  goto fail;
334  }
335 
336  snprintf(buf, sizeof(buf), "Subsample_%s_%s_%s_uv", function_infix, in_fmt_name, out_fmt_name);
337  ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_uv, s->cu_module, buf));
338  if (ret < 0)
339  goto fail;
340 
341 fail:
342  CHECK_CU(cu->cuCtxPopCurrent(&dummy));
343 
344  return ret;
345 }
346 
347  static av_cold int cudascale_config_props(AVFilterLink *outlink)
348 {
349  AVFilterContext *ctx = outlink->src;
350  AVFilterLink *inlink = outlink->src->inputs[0];
351  CUDAScaleContext *s = ctx->priv;
352  AVHWFramesContext *frames_ctx = (AVHWFramesContext*)inlink->hw_frames_ctx->data;
353  AVCUDADeviceContext *device_hwctx = frames_ctx->device_ctx->hwctx;
354  int w, h;
355  int ret;
356 
357  s->hwctx = device_hwctx;
358  s->cu_stream = s->hwctx->stream;
359 
360  if ((ret = ff_scale_eval_dimensions(s,
361  s->w_expr, s->h_expr,
362  inlink, outlink,
363  &w, &h)) < 0)
364  goto fail;
365 
366  ff_scale_adjust_dimensions(inlink, &w, &h,
367  s->force_original_aspect_ratio, s->force_divisible_by);
368 
369  if (((int64_t)h * inlink->w) > INT_MAX ||
370  ((int64_t)w * inlink->h) > INT_MAX)
371  av_log(ctx, AV_LOG_ERROR, "Rescaled value for width or height is too big.\n");
372 
373  outlink->w = w;
374  outlink->h = h;
375 
376  ret = init_processing_chain(ctx, inlink->w, inlink->h, w, h);
377  if (ret < 0)
378  return ret;
379 
380  if (inlink->sample_aspect_ratio.num) {
381  outlink->sample_aspect_ratio = av_mul_q((AVRational){outlink->h*inlink->w,
382  outlink->w*inlink->h},
383  inlink->sample_aspect_ratio);
384  } else {
385  outlink->sample_aspect_ratio = inlink->sample_aspect_ratio;
386  }
387 
388  av_log(ctx, AV_LOG_VERBOSE, "w:%d h:%d fmt:%s -> w:%d h:%d fmt:%s%s\n",
389  inlink->w, inlink->h, av_get_pix_fmt_name(s->in_fmt),
390  outlink->w, outlink->h, av_get_pix_fmt_name(s->out_fmt),
391  s->passthrough ? " (passthrough)" : "");
392 
393  ret = cudascale_load_functions(ctx);
394  if (ret < 0)
395  return ret;
396 
397  return 0;
398 
399 fail:
400  return ret;
401 }
402 
403  static int call_resize_kernel(AVFilterContext *ctx, CUfunction func,
404  CUtexObject src_tex[4], int src_width, int src_height,
405  AVFrame *out_frame, int dst_width, int dst_height, int dst_pitch)
406 {
407  CUDAScaleContext *s = ctx->priv;
408  CudaFunctions *cu = s->hwctx->internal->cuda_dl;
409 
410  CUdeviceptr dst_devptr[4] = {
411  (CUdeviceptr)out_frame->data[0], (CUdeviceptr)out_frame->data[1],
412  (CUdeviceptr)out_frame->data[2], (CUdeviceptr)out_frame->data[3]
413  };
414 
415  void *args_uchar[] = {
416  &src_tex[0], &src_tex[1], &src_tex[2], &src_tex[3],
417  &dst_devptr[0], &dst_devptr[1], &dst_devptr[2], &dst_devptr[3],
418  &dst_width, &dst_height, &dst_pitch,
419  &src_width, &src_height, &s->param
420  };
421 
422  return CHECK_CU(cu->cuLaunchKernel(func,
423  DIV_UP(dst_width, BLOCKX), DIV_UP(dst_height, BLOCKY), 1,
424  BLOCKX, BLOCKY, 1, 0, s->cu_stream, args_uchar, NULL));
425 }
426 
427  static int scalecuda_resize(AVFilterContext *ctx,
428  AVFrame *out, AVFrame *in)
429 {
430  CUDAScaleContext *s = ctx->priv;
431  CudaFunctions *cu = s->hwctx->internal->cuda_dl;
432  CUcontext dummy, cuda_ctx = s->hwctx->cuda_ctx;
433  int i, ret;
434 
435  CUtexObject tex[4] = { 0, 0, 0, 0 };
436 
437  ret = CHECK_CU(cu->cuCtxPushCurrent(cuda_ctx));
438  if (ret < 0)
439  return ret;
440 
441  for (i = 0; i < s->in_planes; i++) {
442  CUDA_TEXTURE_DESC tex_desc = {
443  .filterMode = s->interp_use_linear ?
444  CU_TR_FILTER_MODE_LINEAR :
445  CU_TR_FILTER_MODE_POINT,
446  .flags = s->interp_as_integer ? CU_TRSF_READ_AS_INTEGER : 0,
447  };
448 
449  CUDA_RESOURCE_DESC res_desc = {
450  .resType = CU_RESOURCE_TYPE_PITCH2D,
451  .res.pitch2D.format = s->in_plane_depths[i] <= 8 ?
452  CU_AD_FORMAT_UNSIGNED_INT8 :
453  CU_AD_FORMAT_UNSIGNED_INT16,
454  .res.pitch2D.numChannels = s->in_plane_channels[i],
455  .res.pitch2D.pitchInBytes = in->linesize[i],
456  .res.pitch2D.devPtr = (CUdeviceptr)in->data[i],
457  };
458 
459  if (i == 1 || i == 2) {
460  res_desc.res.pitch2D.width = AV_CEIL_RSHIFT(in->width, s->in_desc->log2_chroma_w);
461  res_desc.res.pitch2D.height = AV_CEIL_RSHIFT(in->height, s->in_desc->log2_chroma_h);
462  } else {
463  res_desc.res.pitch2D.width = in->width;
464  res_desc.res.pitch2D.height = in->height;
465  }
466 
467  ret = CHECK_CU(cu->cuTexObjectCreate(&tex[i], &res_desc, &tex_desc, NULL));
468  if (ret < 0)
469  goto exit;
470  }
471 
472  // scale primary plane(s). Usually Y (and A), or single plane of RGB frames.
473  ret = call_resize_kernel(ctx, s->cu_func,
474  tex, in->width, in->height,
475  out, out->width, out->height, out->linesize[0]);
476  if (ret < 0)
477  goto exit;
478 
479  if (s->out_planes > 1) {
480  // scale UV plane. Scale function sets both U and V plane, or singular interleaved plane.
481  ret = call_resize_kernel(ctx, s->cu_func_uv, tex,
482  AV_CEIL_RSHIFT(in->width, s->in_desc->log2_chroma_w),
483  AV_CEIL_RSHIFT(in->height, s->in_desc->log2_chroma_h),
484  out,
485  AV_CEIL_RSHIFT(out->width, s->out_desc->log2_chroma_w),
486  AV_CEIL_RSHIFT(out->height, s->out_desc->log2_chroma_h),
487  out->linesize[1]);
488  if (ret < 0)
489  goto exit;
490  }
491 
492 exit:
493  for (i = 0; i < s->in_planes; i++)
494  if (tex[i])
495  CHECK_CU(cu->cuTexObjectDestroy(tex[i]));
496 
497  CHECK_CU(cu->cuCtxPopCurrent(&dummy));
498 
499  return ret;
500 }
501 
502  static int cudascale_scale(AVFilterContext *ctx, AVFrame *out, AVFrame *in)
503 {
504  CUDAScaleContext *s = ctx->priv;
505  AVFilterLink *outlink = ctx->outputs[0];
506  AVFrame *src = in;
507  int ret;
508 
509  ret = scalecuda_resize(ctx, s->frame, src);
510  if (ret < 0)
511  return ret;
512 
513  src = s->frame;
514  ret = av_hwframe_get_buffer(src->hw_frames_ctx, s->tmp_frame, 0);
515  if (ret < 0)
516  return ret;
517 
518  av_frame_move_ref(out, s->frame);
519  av_frame_move_ref(s->frame, s->tmp_frame);
520 
521  s->frame->width = outlink->w;
522  s->frame->height = outlink->h;
523 
524  ret = av_frame_copy_props(out, in);
525  if (ret < 0)
526  return ret;
527 
528  return 0;
529 }
530 
531  static int cudascale_filter_frame(AVFilterLink *link, AVFrame *in)
532 {
533  AVFilterContext *ctx = link->dst;
534  CUDAScaleContext *s = ctx->priv;
535  AVFilterLink *outlink = ctx->outputs[0];
536  CudaFunctions *cu = s->hwctx->internal->cuda_dl;
537 
538  AVFrame *out = NULL;
539  CUcontext dummy;
540  int ret = 0;
541 
542  if (s->passthrough)
543  return ff_filter_frame(outlink, in);
544 
545  out = av_frame_alloc();
546  if (!out) {
547  ret = AVERROR(ENOMEM);
548  goto fail;
549  }
550 
551  ret = CHECK_CU(cu->cuCtxPushCurrent(s->hwctx->cuda_ctx));
552  if (ret < 0)
553  goto fail;
554 
555  ret = cudascale_scale(ctx, out, in);
556 
557  CHECK_CU(cu->cuCtxPopCurrent(&dummy));
558  if (ret < 0)
559  goto fail;
560 
561  av_reduce(&out->sample_aspect_ratio.num, &out->sample_aspect_ratio.den,
562  (int64_t)in->sample_aspect_ratio.num * outlink->h * link->w,
563  (int64_t)in->sample_aspect_ratio.den * outlink->w * link->h,
564  INT_MAX);
565 
566  av_frame_free(&in);
567  return ff_filter_frame(outlink, out);
568 fail:
569  av_frame_free(&in);
570  av_frame_free(&out);
571  return ret;
572 }
573 
574  static AVFrame *cudascale_get_video_buffer(AVFilterLink *inlink, int w, int h)
575 {
576  CUDAScaleContext *s = inlink->dst->priv;
577 
578  return s->passthrough ?
579  ff_null_get_video_buffer (inlink, w, h) :
580  ff_default_get_video_buffer(inlink, w, h);
581 }
582 
583  #define OFFSET(x) offsetof(CUDAScaleContext, x)
584  #define FLAGS (AV_OPT_FLAG_FILTERING_PARAM|AV_OPT_FLAG_VIDEO_PARAM)
585  static const AVOption options[] = {
586  { "w", "Output video width", OFFSET(w_expr), AV_OPT_TYPE_STRING, { .str = "iw" }, .flags = FLAGS },
587  { "h", "Output video height", OFFSET(h_expr), AV_OPT_TYPE_STRING, { .str = "ih" }, .flags = FLAGS },
588  { "interp_algo", "Interpolation algorithm used for resizing", OFFSET(interp_algo), AV_OPT_TYPE_INT, { .i64 = INTERP_ALGO_DEFAULT }, 0, INTERP_ALGO_COUNT - 1, FLAGS, .unit = "interp_algo" },
589  { "nearest", "nearest neighbour", 0, AV_OPT_TYPE_CONST, { .i64 = INTERP_ALGO_NEAREST }, 0, 0, FLAGS, .unit = "interp_algo" },
590  { "bilinear", "bilinear", 0, AV_OPT_TYPE_CONST, { .i64 = INTERP_ALGO_BILINEAR }, 0, 0, FLAGS, .unit = "interp_algo" },
591  { "bicubic", "bicubic", 0, AV_OPT_TYPE_CONST, { .i64 = INTERP_ALGO_BICUBIC }, 0, 0, FLAGS, .unit = "interp_algo" },
592  { "lanczos", "lanczos", 0, AV_OPT_TYPE_CONST, { .i64 = INTERP_ALGO_LANCZOS }, 0, 0, FLAGS, .unit = "interp_algo" },
593  { "format", "Output video pixel format", OFFSET(format), AV_OPT_TYPE_PIXEL_FMT, { .i64 = AV_PIX_FMT_NONE }, INT_MIN, INT_MAX, .flags=FLAGS },
594  { "passthrough", "Do not process frames at all if parameters match", OFFSET(passthrough), AV_OPT_TYPE_BOOL, { .i64 = 1 }, 0, 1, FLAGS },
595  { "param", "Algorithm-Specific parameter", OFFSET(param), AV_OPT_TYPE_FLOAT, { .dbl = SCALE_CUDA_PARAM_DEFAULT }, -FLT_MAX, FLT_MAX, FLAGS },
596  { "force_original_aspect_ratio", "decrease or increase w/h if necessary to keep the original AR", OFFSET(force_original_aspect_ratio), AV_OPT_TYPE_INT, { .i64 = 0 }, 0, 2, FLAGS, .unit = "force_oar" },
597  { "disable", NULL, 0, AV_OPT_TYPE_CONST, {.i64 = 0 }, 0, 0, FLAGS, .unit = "force_oar" },
598  { "decrease", NULL, 0, AV_OPT_TYPE_CONST, {.i64 = 1 }, 0, 0, FLAGS, .unit = "force_oar" },
599  { "increase", NULL, 0, AV_OPT_TYPE_CONST, {.i64 = 2 }, 0, 0, FLAGS, .unit = "force_oar" },
600  { "force_divisible_by", "enforce that the output resolution is divisible by a defined integer when force_original_aspect_ratio is used", OFFSET(force_divisible_by), AV_OPT_TYPE_INT, { .i64 = 1 }, 1, 256, FLAGS },
601  { NULL },
602 };
603 
604  static const AVClass cudascale_class = {
605  .class_name = "cudascale",
606  .item_name = av_default_item_name,
607  .option = options,
608  .version = LIBAVUTIL_VERSION_INT,
609 };
610 
611  static const AVFilterPad cudascale_inputs[] = {
612  {
613  .name = "default",
614  .type = AVMEDIA_TYPE_VIDEO,
615  .filter_frame = cudascale_filter_frame,
616  .get_buffer.video = cudascale_get_video_buffer,
617  },
618 };
619 
620  static const AVFilterPad cudascale_outputs[] = {
621  {
622  .name = "default",
623  .type = AVMEDIA_TYPE_VIDEO,
624  .config_props = cudascale_config_props,
625  },
626 };
627 
628  const AVFilter ff_vf_scale_cuda = {
629  .name = "scale_cuda",
630  .description = NULL_IF_CONFIG_SMALL("GPU accelerated video resizer"),
631 
632  .init = cudascale_init,
633  .uninit = cudascale_uninit,
634 
635  .priv_size = sizeof(CUDAScaleContext),
636  .priv_class = &cudascale_class,
637 
638  FILTER_INPUTS(cudascale_inputs),
639  FILTER_OUTPUTS(cudascale_outputs),
640 
641  FILTER_SINGLE_PIXFMT(AV_PIX_FMT_CUDA),
642 
643  .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE,
644 };
options
static const AVOption options[]
Definition: vf_scale_cuda.c:585
func
int(* func)(AVBPrint *dst, const char *in, const char *arg)
Definition: jacosubdec.c:68
AVHWDeviceContext::hwctx
void * hwctx
The format-specific data, allocated and freed by libavutil along with this context.
Definition: hwcontext.h:85
AV_PIX_FMT_CUDA
@ AV_PIX_FMT_CUDA
HW acceleration through CUDA.
Definition: pixfmt.h:260
AVPixelFormat
AVPixelFormat
Pixel format.
Definition: pixfmt.h:71
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
INTERP_ALGO_LANCZOS
@ INTERP_ALGO_LANCZOS
Definition: vf_scale_cuda.c:67
opt.h
CUDAScaleContext::frames_ctx
AVBufferRef * frames_ctx
Definition: vf_scale_cuda.c:83
hwcontext_cuda_internal.h
cudascale_init
static av_cold int cudascale_init(AVFilterContext *ctx)
Definition: vf_scale_cuda.c:113
out
FILE * out
Definition: movenc.c:54
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:351
AV_PIX_FMT_BGR32
#define AV_PIX_FMT_BGR32
Definition: pixfmt.h:453
ff_filter_frame
int ff_filter_frame(AVFilterLink *link, AVFrame *frame)
Send a frame of data to the next filter.
Definition: avfilter.c:1018
av_pix_fmt_desc_get
const AVPixFmtDescriptor * av_pix_fmt_desc_get(enum AVPixelFormat pix_fmt)
Definition: pixdesc.c:2962
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:197
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
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
CUDAScaleContext::tmp_frame
AVFrame * tmp_frame
Definition: vf_scale_cuda.c:86
CUDAScaleContext::passthrough
int passthrough
Definition: vf_scale_cuda.c:87
cudascale_uninit
static av_cold void cudascale_uninit(AVFilterContext *ctx)
Definition: vf_scale_cuda.c:128
av_frame_free
void av_frame_free(AVFrame **frame)
Free the frame and any dynamically allocated objects in it, e.g.
Definition: frame.c:130
av_hwframe_ctx_init
int av_hwframe_ctx_init(AVBufferRef *ref)
Finalize the context before use.
Definition: hwcontext.c:322
CUDAScaleContext::w_expr
char * w_expr
width expression string
Definition: vf_scale_cuda.c:94
AVFrame
This structure describes decoded (raw) audio or video data.
Definition: frame.h:344
pixdesc.h
AVFrame::width
int width
Definition: frame.h:416
w
uint8_t w
Definition: llviddspenc.c:38
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:248
AVOption
AVOption.
Definition: opt.h:346
init_hwframe_ctx
static av_cold int init_hwframe_ctx(CUDAScaleContext *s, AVBufferRef *device_ctx, int width, int height)
Definition: vf_scale_cuda.c:147
CUDAScaleContext::interp_use_linear
int interp_use_linear
Definition: vf_scale_cuda.c:107
FLAGS
#define FLAGS
Definition: vf_scale_cuda.c:584
AV_LOG_VERBOSE
#define AV_LOG_VERBOSE
Detailed information.
Definition: log.h:196
ff_scale_eval_dimensions
int ff_scale_eval_dimensions(void *log_ctx, const char *w_expr, const char *h_expr, AVFilterLink *inlink, AVFilterLink *outlink, int *ret_w, int *ret_h)
Parse and evaluate string expressions for width and height.
Definition: scale_eval.c:57
float.h
DIV_UP
#define DIV_UP(a, b)
Definition: vf_scale_cuda.c:55
FFMAX
#define FFMAX(a, b)
Definition: macros.h: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:170
AVHWFramesContext::width
int width
The allocated dimensions of the frames in this pool.
Definition: hwcontext.h:217
video.h
AVFrame::data
uint8_t * data[AV_NUM_DATA_POINTERS]
pointer to the picture/channel planes.
Definition: frame.h:365
CUDAScaleContext::frame
AVFrame * frame
Definition: vf_scale_cuda.c:84
ff_default_get_video_buffer
AVFrame * ff_default_get_video_buffer(AVFilterLink *link, int w, int h)
Definition: video.c:107
av_pix_fmt_count_planes
int av_pix_fmt_count_planes(enum AVPixelFormat pix_fmt)
Definition: pixdesc.c:3002
CUDAScaleContext::cu_func_uv
CUfunction cu_func_uv
Definition: vf_scale_cuda.c:103
vf_scale_cuda.h
fail
#define fail()
Definition: checkasm.h:179
CHECK_CU
#define CHECK_CU(x)
Definition: vf_scale_cuda.c:59
CUDAScaleContext::in_planes
int in_planes
Definition: vf_scale_cuda.c:79
dummy
int dummy
Definition: motion.c:66
scalecuda_resize
static int scalecuda_resize(AVFilterContext *ctx, AVFrame *out, AVFrame *in)
Definition: vf_scale_cuda.c:427
CUDAScaleContext::in_plane_channels
int in_plane_channels[4]
Definition: vf_scale_cuda.c:81
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
cudascale_load_functions
static av_cold int cudascale_load_functions(AVFilterContext *ctx)
Definition: vf_scale_cuda.c:276
AVFilterPad
A filter pad used for either input or output.
Definition: internal.h:33
av_frame_alloc
AVFrame * av_frame_alloc(void)
Allocate an AVFrame and set its fields to default values.
Definition: frame.c:118
cudascale_class
static const AVClass cudascale_class
Definition: vf_scale_cuda.c:604
cudascale_config_props
static av_cold int cudascale_config_props(AVFilterLink *outlink)
Definition: vf_scale_cuda.c:347
CUDAScaleContext::hwctx
AVCUDADeviceContext * hwctx
Definition: vf_scale_cuda.c:75
INTERP_ALGO_NEAREST
@ INTERP_ALGO_NEAREST
Definition: vf_scale_cuda.c:64
AV_LOG_ERROR
#define AV_LOG_ERROR
Something went wrong and cannot losslessly be recovered.
Definition: log.h:180
set_format_info
static av_cold void set_format_info(AVFilterContext *ctx, enum AVPixelFormat in_format, enum AVPixelFormat out_format)
Definition: vf_scale_cuda.c:194
FF_ARRAY_ELEMS
#define FF_ARRAY_ELEMS(a)
Definition: sinewin_tablegen.c:29
av_cold
#define av_cold
Definition: attributes.h:90
AVHWFramesContext::height
int height
Definition: hwcontext.h:217
CUDAScaleContext::interp_as_integer
int interp_as_integer
Definition: vf_scale_cuda.c:108
width
#define width
s
#define s(width, name)
Definition: cbs_vp9.c:198
AV_PIX_FMT_YUV444P16
#define AV_PIX_FMT_YUV444P16
Definition: pixfmt.h:491
AV_CEIL_RSHIFT
#define AV_CEIL_RSHIFT(a, b)
Definition: common.h:58
format
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 format(the sample packing is implied by the sample format) and sample rate. The lists are not just lists
AV_PIX_FMT_0BGR32
#define AV_PIX_FMT_0BGR32
Definition: pixfmt.h:456
CUDAScaleContext::cu_stream
CUstream cu_stream
Definition: vf_scale_cuda.c:104
ctx
AVFormatContext * ctx
Definition: movenc.c:48
INTERP_ALGO_BICUBIC
@ INTERP_ALGO_BICUBIC
Definition: vf_scale_cuda.c:66
load_helper.h
AV_PIX_FMT_YUV420P
@ AV_PIX_FMT_YUV420P
planar YUV 4:2:0, 12bpp, (1 Cr & Cb sample per 2x2 Y samples)
Definition: pixfmt.h:73
FILTER_INPUTS
#define FILTER_INPUTS(array)
Definition: internal.h:182
format_is_supported
static int format_is_supported(enum AVPixelFormat fmt)
Definition: vf_scale_cuda.c:184
link
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 link
Definition: filter_design.txt:23
CUDAScaleContext::param
float param
Definition: vf_scale_cuda.c:110
if
if(ret)
Definition: filter_design.txt:179
CUDAScaleContext::force_divisible_by
int force_divisible_by
Definition: vf_scale_cuda.c:98
INTERP_ALGO_BILINEAR
@ INTERP_ALGO_BILINEAR
Definition: vf_scale_cuda.c:65
INTERP_ALGO_COUNT
@ INTERP_ALGO_COUNT
Definition: vf_scale_cuda.c:69
LIBAVUTIL_VERSION_INT
#define LIBAVUTIL_VERSION_INT
Definition: version.h:85
CUDAScaleContext::interp_algo
int interp_algo
Definition: vf_scale_cuda.c:106
AVClass
Describe the class of an AVClass context structure.
Definition: log.h:66
OFFSET
#define OFFSET(x)
Definition: vf_scale_cuda.c:583
NULL
#define NULL
Definition: coverity.c:32
AVHWFramesContext::sw_format
enum AVPixelFormat sw_format
The pixel format identifying the actual data layout of the hardware frames.
Definition: hwcontext.h:210
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:679
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
CUDAScaleContext::out_desc
const AVPixFmtDescriptor * out_desc
Definition: vf_scale_cuda.c:78
CUDAScaleContext::h_expr
char * h_expr
height expression string
Definition: vf_scale_cuda.c:95
BLOCKY
#define BLOCKY
Definition: vf_scale_cuda.c:57
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:126
AVFilterContext::inputs
AVFilterLink ** inputs
array of pointers to input links
Definition: avfilter.h:415
av_default_item_name
const char * av_default_item_name(void *ptr)
Return the context name.
Definition: log.c:237
cudascale_inputs
static const AVFilterPad cudascale_inputs[]
Definition: vf_scale_cuda.c:611
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
CUDAScaleContext::cu_func
CUfunction cu_func
Definition: vf_scale_cuda.c:102
scale_eval.h
height
#define height
AV_PIX_FMT_RGB32
#define AV_PIX_FMT_RGB32
Definition: pixfmt.h:451
CUDAScaleContext::cu_module
CUmodule cu_module
Definition: vf_scale_cuda.c:101
ff_null_get_video_buffer
AVFrame * ff_null_get_video_buffer(AVFilterLink *link, int w, int h)
Definition: video.c:44
internal.h
CUDAScaleContext
Definition: vf_scale_cuda.c:72
AV_OPT_TYPE_FLOAT
@ AV_OPT_TYPE_FLOAT
Definition: opt.h:238
FILTER_SINGLE_PIXFMT
#define FILTER_SINGLE_PIXFMT(pix_fmt_)
Definition: internal.h:172
supported_formats
static enum AVPixelFormat supported_formats[]
Definition: vf_scale_cuda.c:42
INTERP_ALGO_DEFAULT
@ INTERP_ALGO_DEFAULT
Definition: vf_scale_cuda.c:62
i
#define i(width, name, range_min, range_max)
Definition: cbs_h2645.c:255
internal.h
common.h
call_resize_kernel
static int call_resize_kernel(AVFilterContext *ctx, CUfunction func, CUtexObject src_tex[4], int src_width, int src_height, AVFrame *out_frame, int dst_width, int dst_height, int dst_pitch)
Definition: vf_scale_cuda.c:403
av_frame_move_ref
void av_frame_move_ref(AVFrame *dst, AVFrame *src)
Move everything contained in src to dst and reset src.
Definition: frame.c:603
av_frame_unref
void av_frame_unref(AVFrame *frame)
Unreference all the buffers referenced by frame and reset the frame fields.
Definition: frame.c:576
AVFilterPad::name
const char * name
Pad name.
Definition: internal.h:39
AV_PIX_FMT_P016
#define AV_PIX_FMT_P016
Definition: pixfmt.h:530
CUDAScaleContext::in_plane_depths
int in_plane_depths[4]
Definition: vf_scale_cuda.c:80
AVFilter
Filter definition.
Definition: avfilter.h:166
AVHWFramesContext
This struct describes a set or pool of "hardware" frames (i.e.
Definition: hwcontext.h:115
AVCUDADeviceContext
This struct is allocated as AVHWDeviceContext.hwctx.
Definition: hwcontext_cuda.h:42
BLOCKX
#define BLOCKX
Definition: vf_scale_cuda.c:56
CUDAScaleContext::out_fmt
enum AVPixelFormat in_fmt out_fmt
Definition: vf_scale_cuda.c:77
ret
ret
Definition: filter_design.txt:187
AV_LOG_FATAL
#define AV_LOG_FATAL
Something went wrong and recovery is not possible.
Definition: log.h:174
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:71
AV_PIX_FMT_0RGB32
#define AV_PIX_FMT_0RGB32
Definition: pixfmt.h:455
AVHWFramesContext::device_ctx
AVHWDeviceContext * device_ctx
The parent AVHWDeviceContext.
Definition: hwcontext.h:134
cuda_check.h
AVFrame::sample_aspect_ratio
AVRational sample_aspect_ratio
Sample aspect ratio for the video frame, 0/1 if unknown/unspecified.
Definition: frame.h:451
ff_scale_adjust_dimensions
int ff_scale_adjust_dimensions(AVFilterLink *inlink, int *ret_w, int *ret_h, int force_original_aspect_ratio, int force_divisible_by)
Transform evaluated width and height obtained from ff_scale_eval_dimensions into actual target width ...
Definition: scale_eval.c:113
AVFrame::height
int height
Definition: frame.h:416
cudascale_get_video_buffer
static AVFrame * cudascale_get_video_buffer(AVFilterLink *inlink, int w, int h)
Definition: vf_scale_cuda.c:574
AVRational::den
int den
Denominator.
Definition: rational.h:60
AV_PIX_FMT_NONE
@ AV_PIX_FMT_NONE
Definition: pixfmt.h:72
AV_OPT_TYPE_INT
@ AV_OPT_TYPE_INT
Definition: opt.h:235
avfilter.h
AV_OPT_TYPE_PIXEL_FMT
@ AV_OPT_TYPE_PIXEL_FMT
Definition: opt.h:246
av_mul_q
AVRational av_mul_q(AVRational b, AVRational c)
Multiply two rationals.
Definition: rational.c:80
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:407
AV_PIX_FMT_P010
#define AV_PIX_FMT_P010
Definition: pixfmt.h:528
cudascale_outputs
static const AVFilterPad cudascale_outputs[]
Definition: vf_scale_cuda.c:620
AVMEDIA_TYPE_VIDEO
@ AVMEDIA_TYPE_VIDEO
Definition: avutil.h:201
AVBufferRef
A reference to a data buffer.
Definition: buffer.h:82
cudascale_scale
static int cudascale_scale(AVFilterContext *ctx, AVFrame *out, AVFrame *in)
Definition: vf_scale_cuda.c:502
cudascale_filter_frame
static int cudascale_filter_frame(AVFilterLink *link, AVFrame *in)
Definition: vf_scale_cuda.c:531
AVPixFmtDescriptor
Descriptor that unambiguously describes how the bits of a pixel are stored in the up to 4 data planes...
Definition: pixdesc.h:69
CUDAScaleContext::force_original_aspect_ratio
int force_original_aspect_ratio
Definition: vf_scale_cuda.c:97
CUDAScaleContext::format
enum AVPixelFormat format
Output sw format.
Definition: vf_scale_cuda.c:92
FFALIGN
#define FFALIGN(x, a)
Definition: macros.h:78
AV_OPT_TYPE_BOOL
@ AV_OPT_TYPE_BOOL
Definition: opt.h:251
FILTER_OUTPUTS
#define FILTER_OUTPUTS(array)
Definition: internal.h:183
src
INIT_CLIP pixel * src
Definition: h264pred_template.c:418
SCALE_CUDA_PARAM_DEFAULT
#define SCALE_CUDA_PARAM_DEFAULT
Definition: vf_scale_cuda.h:26
d
d
Definition: ffmpeg_filter.c:409
CUDAScaleContext::in_desc
const AVPixFmtDescriptor * in_desc
Definition: vf_scale_cuda.c:78
init_processing_chain
static av_cold int init_processing_chain(AVFilterContext *ctx, int in_width, int in_height, int out_width, int out_height)
Definition: vf_scale_cuda.c:220
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:389
av_log
#define av_log(a,...)
Definition: tableprint_vlc.h:27
h
h
Definition: vp9dsp_template.c:2038
AV_OPT_TYPE_STRING
@ AV_OPT_TYPE_STRING
Definition: opt.h:239
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:491
AV_OPT_TYPE_CONST
@ AV_OPT_TYPE_CONST
Definition: opt.h:244
snprintf
#define snprintf
Definition: snprintf.h:34
ff_vf_scale_cuda
const AVFilter ff_vf_scale_cuda
Definition: vf_scale_cuda.c:628
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:2882
CUDAScaleContext::out_planes
int out_planes
Definition: vf_scale_cuda.c:79
CUDAScaleContext::cu_ctx
CUcontext cu_ctx
Definition: vf_scale_cuda.c:100
Generated on Thu Sep 26 2024 23:15:40 for FFmpeg by   doxygen 1.8.17

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