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

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