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 <stdio.h>
24 #include <string.h>
25 
26 #include "libavutil/avstring.h"
27 #include "libavutil/common.h"
28 #include "libavutil/hwcontext.h"
30 #include "libavutil/cuda_check.h"
31 #include "libavutil/internal.h"
32 #include "libavutil/opt.h"
33 #include "libavutil/pixdesc.h"
34 
35 #include "avfilter.h"
36 #include "formats.h"
37 #include "internal.h"
38 #include "scale_eval.h"
39 #include "video.h"
40 
41 static const enum AVPixelFormat supported_formats[] = {
48 };
49 
50 #define DIV_UP(a, b) ( ((a) + (b) - 1) / (b) )
51 #define ALIGN_UP(a, b) (((a) + (b) - 1) & ~((b) - 1))
52 #define NUM_BUFFERS 2
53 #define BLOCKX 32
54 #define BLOCKY 16
55 
56 #define CHECK_CU(x) FF_CUDA_CHECK_DL(ctx, s->hwctx->internal->cuda_dl, x)
57 
58 typedef struct CUDAScaleContext {
59  const AVClass *class;
60 
62 
65 
66  struct {
67  int width;
68  int height;
69  } planes_in[3], planes_out[3];
70 
73 
76 
77  /**
78  * Output sw format. AV_PIX_FMT_NONE for no conversion.
79  */
81 
82  char *w_expr; ///< width expression string
83  char *h_expr; ///< height expression string
84 
87 
88  CUcontext cu_ctx;
89  CUmodule cu_module;
90  CUfunction cu_func_uchar;
91  CUfunction cu_func_uchar2;
92  CUfunction cu_func_uchar4;
93  CUfunction cu_func_ushort;
94  CUfunction cu_func_ushort2;
95  CUfunction cu_func_ushort4;
96  CUstream cu_stream;
97 
98  CUdeviceptr srcBuffer;
99  CUdeviceptr dstBuffer;
102 
104 {
105  CUDAScaleContext *s = ctx->priv;
106 
107  s->format = AV_PIX_FMT_NONE;
108  s->frame = av_frame_alloc();
109  if (!s->frame)
110  return AVERROR(ENOMEM);
111 
112  s->tmp_frame = av_frame_alloc();
113  if (!s->tmp_frame)
114  return AVERROR(ENOMEM);
115 
116  return 0;
117 }
118 
120 {
121  CUDAScaleContext *s = ctx->priv;
122 
123  av_frame_free(&s->frame);
126 }
127 
129 {
130  static const enum AVPixelFormat pixel_formats[] = {
132  };
133  AVFilterFormats *pix_fmts = ff_make_format_list(pixel_formats);
134 
135  return ff_set_common_formats(ctx, pix_fmts);
136 }
137 
139 {
140  AVBufferRef *out_ref = NULL;
141  AVHWFramesContext *out_ctx;
142  int in_sw, in_sh, out_sw, out_sh;
143  int ret, i;
144 
145  av_pix_fmt_get_chroma_sub_sample(s->in_fmt, &in_sw, &in_sh);
146  av_pix_fmt_get_chroma_sub_sample(s->out_fmt, &out_sw, &out_sh);
147  if (!s->planes_out[0].width) {
148  s->planes_out[0].width = s->planes_in[0].width;
149  s->planes_out[0].height = s->planes_in[0].height;
150  }
151 
152  for (i = 1; i < FF_ARRAY_ELEMS(s->planes_in); i++) {
153  s->planes_in[i].width = s->planes_in[0].width >> in_sw;
154  s->planes_in[i].height = s->planes_in[0].height >> in_sh;
155  s->planes_out[i].width = s->planes_out[0].width >> out_sw;
156  s->planes_out[i].height = s->planes_out[0].height >> out_sh;
157  }
158 
159  out_ref = av_hwframe_ctx_alloc(device_ctx);
160  if (!out_ref)
161  return AVERROR(ENOMEM);
162  out_ctx = (AVHWFramesContext*)out_ref->data;
163 
164  out_ctx->format = AV_PIX_FMT_CUDA;
165  out_ctx->sw_format = s->out_fmt;
166  out_ctx->width = FFALIGN(s->planes_out[0].width, 32);
167  out_ctx->height = FFALIGN(s->planes_out[0].height, 32);
168 
169  ret = av_hwframe_ctx_init(out_ref);
170  if (ret < 0)
171  goto fail;
172 
173  av_frame_unref(s->frame);
174  ret = av_hwframe_get_buffer(out_ref, s->frame, 0);
175  if (ret < 0)
176  goto fail;
177 
178  s->frame->width = s->planes_out[0].width;
179  s->frame->height = s->planes_out[0].height;
180 
182  s->frames_ctx = out_ref;
183 
184  return 0;
185 fail:
186  av_buffer_unref(&out_ref);
187  return ret;
188 }
189 
191 {
192  int i;
193 
194  for (i = 0; i < FF_ARRAY_ELEMS(supported_formats); i++)
195  if (supported_formats[i] == fmt)
196  return 1;
197  return 0;
198 }
199 
200 static av_cold int init_processing_chain(AVFilterContext *ctx, int in_width, int in_height,
201  int out_width, int out_height)
202 {
203  CUDAScaleContext *s = ctx->priv;
204 
205  AVHWFramesContext *in_frames_ctx;
206 
207  enum AVPixelFormat in_format;
208  enum AVPixelFormat out_format;
209  int ret;
210 
211  /* check that we have a hw context */
212  if (!ctx->inputs[0]->hw_frames_ctx) {
213  av_log(ctx, AV_LOG_ERROR, "No hw context provided on input\n");
214  return AVERROR(EINVAL);
215  }
216  in_frames_ctx = (AVHWFramesContext*)ctx->inputs[0]->hw_frames_ctx->data;
217  in_format = in_frames_ctx->sw_format;
218  out_format = (s->format == AV_PIX_FMT_NONE) ? in_format : s->format;
219 
220  if (!format_is_supported(in_format)) {
221  av_log(ctx, AV_LOG_ERROR, "Unsupported input format: %s\n",
222  av_get_pix_fmt_name(in_format));
223  return AVERROR(ENOSYS);
224  }
225  if (!format_is_supported(out_format)) {
226  av_log(ctx, AV_LOG_ERROR, "Unsupported output format: %s\n",
227  av_get_pix_fmt_name(out_format));
228  return AVERROR(ENOSYS);
229  }
230 
231  if (in_width == out_width && in_height == out_height)
232  s->passthrough = 1;
233 
234  s->in_fmt = in_format;
235  s->out_fmt = out_format;
236 
237  s->planes_in[0].width = in_width;
238  s->planes_in[0].height = in_height;
239  s->planes_out[0].width = out_width;
240  s->planes_out[0].height = out_height;
241 
242  ret = init_stage(s, in_frames_ctx->device_ref);
243  if (ret < 0)
244  return ret;
245 
247  if (!ctx->outputs[0]->hw_frames_ctx)
248  return AVERROR(ENOMEM);
249 
250  return 0;
251 }
252 
254 {
255  AVFilterContext *ctx = outlink->src;
256  AVFilterLink *inlink = outlink->src->inputs[0];
257  CUDAScaleContext *s = ctx->priv;
259  AVCUDADeviceContext *device_hwctx = frames_ctx->device_ctx->hwctx;
260  CUcontext dummy, cuda_ctx = device_hwctx->cuda_ctx;
261  CudaFunctions *cu = device_hwctx->internal->cuda_dl;
262  int w, h;
263  int ret;
264 
265  extern char vf_scale_cuda_ptx[];
266 
267  s->hwctx = device_hwctx;
268  s->cu_stream = s->hwctx->stream;
269 
270  ret = CHECK_CU(cu->cuCtxPushCurrent(cuda_ctx));
271  if (ret < 0)
272  goto fail;
273 
274  ret = CHECK_CU(cu->cuModuleLoadData(&s->cu_module, vf_scale_cuda_ptx));
275  if (ret < 0)
276  goto fail;
277 
278  CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_uchar, s->cu_module, "Subsample_Bilinear_uchar"));
279  if (ret < 0)
280  goto fail;
281 
282  CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_uchar2, s->cu_module, "Subsample_Bilinear_uchar2"));
283  if (ret < 0)
284  goto fail;
285 
286  CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_uchar4, s->cu_module, "Subsample_Bilinear_uchar4"));
287  if (ret < 0)
288  goto fail;
289 
290  CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_ushort, s->cu_module, "Subsample_Bilinear_ushort"));
291  if (ret < 0)
292  goto fail;
293 
294  CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_ushort2, s->cu_module, "Subsample_Bilinear_ushort2"));
295  if (ret < 0)
296  goto fail;
297 
298  CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_ushort4, s->cu_module, "Subsample_Bilinear_ushort4"));
299  if (ret < 0)
300  goto fail;
301 
302 
303  CHECK_CU(cu->cuCtxPopCurrent(&dummy));
304 
305  if ((ret = ff_scale_eval_dimensions(s,
306  s->w_expr, s->h_expr,
307  inlink, outlink,
308  &w, &h)) < 0)
309  goto fail;
310 
311  ff_scale_adjust_dimensions(inlink, &w, &h,
313 
314  if (((int64_t)h * inlink->w) > INT_MAX ||
315  ((int64_t)w * inlink->h) > INT_MAX)
316  av_log(ctx, AV_LOG_ERROR, "Rescaled value for width or height is too big.\n");
317 
318  outlink->w = w;
319  outlink->h = h;
320 
321  ret = init_processing_chain(ctx, inlink->w, inlink->h, w, h);
322  if (ret < 0)
323  return ret;
324 
325  av_log(ctx, AV_LOG_VERBOSE, "w:%d h:%d -> w:%d h:%d\n",
326  inlink->w, inlink->h, outlink->w, outlink->h);
327 
328  if (inlink->sample_aspect_ratio.num) {
329  outlink->sample_aspect_ratio = av_mul_q((AVRational){outlink->h*inlink->w,
330  outlink->w*inlink->h},
331  inlink->sample_aspect_ratio);
332  } else {
333  outlink->sample_aspect_ratio = inlink->sample_aspect_ratio;
334  }
335 
336  return 0;
337 
338 fail:
339  return ret;
340 }
341 
342 static int call_resize_kernel(AVFilterContext *ctx, CUfunction func, int channels,
343  uint8_t *src_dptr, int src_width, int src_height, int src_pitch,
344  uint8_t *dst_dptr, int dst_width, int dst_height, int dst_pitch,
345  int pixel_size)
346 {
347  CUDAScaleContext *s = ctx->priv;
348  CudaFunctions *cu = s->hwctx->internal->cuda_dl;
349  CUdeviceptr dst_devptr = (CUdeviceptr)dst_dptr;
350  CUtexObject tex = 0;
351  void *args_uchar[] = { &tex, &dst_devptr, &dst_width, &dst_height, &dst_pitch, &src_width, &src_height };
352  int ret;
353 
354  CUDA_TEXTURE_DESC tex_desc = {
355  .filterMode = CU_TR_FILTER_MODE_LINEAR,
356  .flags = CU_TRSF_READ_AS_INTEGER,
357  };
358 
359  CUDA_RESOURCE_DESC res_desc = {
360  .resType = CU_RESOURCE_TYPE_PITCH2D,
361  .res.pitch2D.format = pixel_size == 1 ?
362  CU_AD_FORMAT_UNSIGNED_INT8 :
363  CU_AD_FORMAT_UNSIGNED_INT16,
364  .res.pitch2D.numChannels = channels,
365  .res.pitch2D.width = src_width,
366  .res.pitch2D.height = src_height,
367  .res.pitch2D.pitchInBytes = src_pitch * pixel_size,
368  .res.pitch2D.devPtr = (CUdeviceptr)src_dptr,
369  };
370 
371  ret = CHECK_CU(cu->cuTexObjectCreate(&tex, &res_desc, &tex_desc, NULL));
372  if (ret < 0)
373  goto exit;
374 
375  ret = CHECK_CU(cu->cuLaunchKernel(func,
376  DIV_UP(dst_width, BLOCKX), DIV_UP(dst_height, BLOCKY), 1,
377  BLOCKX, BLOCKY, 1, 0, s->cu_stream, args_uchar, NULL));
378 
379 exit:
380  if (tex)
381  CHECK_CU(cu->cuTexObjectDestroy(tex));
382 
383  return ret;
384 }
385 
387  AVFrame *out, AVFrame *in)
388 {
389  AVHWFramesContext *in_frames_ctx = (AVHWFramesContext*)in->hw_frames_ctx->data;
390  CUDAScaleContext *s = ctx->priv;
391 
392  switch (in_frames_ctx->sw_format) {
393  case AV_PIX_FMT_YUV420P:
394  call_resize_kernel(ctx, s->cu_func_uchar, 1,
395  in->data[0], in->width, in->height, in->linesize[0],
396  out->data[0], out->width, out->height, out->linesize[0],
397  1);
398  call_resize_kernel(ctx, s->cu_func_uchar, 1,
399  in->data[1], in->width/2, in->height/2, in->linesize[0]/2,
400  out->data[1], out->width/2, out->height/2, out->linesize[0]/2,
401  1);
402  call_resize_kernel(ctx, s->cu_func_uchar, 1,
403  in->data[2], in->width/2, in->height/2, in->linesize[0]/2,
404  out->data[2], out->width/2, out->height/2, out->linesize[0]/2,
405  1);
406  break;
407  case AV_PIX_FMT_YUV444P:
408  call_resize_kernel(ctx, s->cu_func_uchar, 1,
409  in->data[0], in->width, in->height, in->linesize[0],
410  out->data[0], out->width, out->height, out->linesize[0],
411  1);
412  call_resize_kernel(ctx, s->cu_func_uchar, 1,
413  in->data[1], in->width, in->height, in->linesize[0],
414  out->data[1], out->width, out->height, out->linesize[0],
415  1);
416  call_resize_kernel(ctx, s->cu_func_uchar, 1,
417  in->data[2], in->width, in->height, in->linesize[0],
418  out->data[2], out->width, out->height, out->linesize[0],
419  1);
420  break;
422  call_resize_kernel(ctx, s->cu_func_ushort, 1,
423  in->data[0], in->width, in->height, in->linesize[0] / 2,
424  out->data[0], out->width, out->height, out->linesize[0] / 2,
425  2);
426  call_resize_kernel(ctx, s->cu_func_ushort, 1,
427  in->data[1], in->width, in->height, in->linesize[1] / 2,
428  out->data[1], out->width, out->height, out->linesize[1] / 2,
429  2);
430  call_resize_kernel(ctx, s->cu_func_ushort, 1,
431  in->data[2], in->width, in->height, in->linesize[2] / 2,
432  out->data[2], out->width, out->height, out->linesize[2] / 2,
433  2);
434  break;
435  case AV_PIX_FMT_NV12:
436  call_resize_kernel(ctx, s->cu_func_uchar, 1,
437  in->data[0], in->width, in->height, in->linesize[0],
438  out->data[0], out->width, out->height, out->linesize[0],
439  1);
440  call_resize_kernel(ctx, s->cu_func_uchar2, 2,
441  in->data[1], in->width/2, in->height/2, in->linesize[1],
442  out->data[1], out->width/2, out->height/2, out->linesize[1]/2,
443  1);
444  break;
445  case AV_PIX_FMT_P010LE:
446  call_resize_kernel(ctx, s->cu_func_ushort, 1,
447  in->data[0], in->width, in->height, in->linesize[0]/2,
448  out->data[0], out->width, out->height, out->linesize[0]/2,
449  2);
450  call_resize_kernel(ctx, s->cu_func_ushort2, 2,
451  in->data[1], in->width / 2, in->height / 2, in->linesize[1]/2,
452  out->data[1], out->width / 2, out->height / 2, out->linesize[1] / 4,
453  2);
454  break;
455  case AV_PIX_FMT_P016LE:
456  call_resize_kernel(ctx, s->cu_func_ushort, 1,
457  in->data[0], in->width, in->height, in->linesize[0] / 2,
458  out->data[0], out->width, out->height, out->linesize[0] / 2,
459  2);
460  call_resize_kernel(ctx, s->cu_func_ushort2, 2,
461  in->data[1], in->width / 2, in->height / 2, in->linesize[1] / 2,
462  out->data[1], out->width / 2, out->height / 2, out->linesize[1] / 4,
463  2);
464  break;
465  default:
466  return AVERROR_BUG;
467  }
468 
469  return 0;
470 }
471 
473 {
474  CUDAScaleContext *s = ctx->priv;
475  AVFrame *src = in;
476  int ret;
477 
478  ret = scalecuda_resize(ctx, s->frame, src);
479  if (ret < 0)
480  return ret;
481 
482  src = s->frame;
483  ret = av_hwframe_get_buffer(src->hw_frames_ctx, s->tmp_frame, 0);
484  if (ret < 0)
485  return ret;
486 
487  av_frame_move_ref(out, s->frame);
489 
490  s->frame->width = s->planes_out[0].width;
491  s->frame->height = s->planes_out[0].height;
492 
493  ret = av_frame_copy_props(out, in);
494  if (ret < 0)
495  return ret;
496 
497  return 0;
498 }
499 
501 {
502  AVFilterContext *ctx = link->dst;
503  CUDAScaleContext *s = ctx->priv;
504  AVFilterLink *outlink = ctx->outputs[0];
505  CudaFunctions *cu = s->hwctx->internal->cuda_dl;
506 
507  AVFrame *out = NULL;
508  CUcontext dummy;
509  int ret = 0;
510 
511  out = av_frame_alloc();
512  if (!out) {
513  ret = AVERROR(ENOMEM);
514  goto fail;
515  }
516 
517  ret = CHECK_CU(cu->cuCtxPushCurrent(s->hwctx->cuda_ctx));
518  if (ret < 0)
519  goto fail;
520 
521  ret = cudascale_scale(ctx, out, in);
522 
523  CHECK_CU(cu->cuCtxPopCurrent(&dummy));
524  if (ret < 0)
525  goto fail;
526 
528  (int64_t)in->sample_aspect_ratio.num * outlink->h * link->w,
529  (int64_t)in->sample_aspect_ratio.den * outlink->w * link->h,
530  INT_MAX);
531 
532  av_frame_free(&in);
533  return ff_filter_frame(outlink, out);
534 fail:
535  av_frame_free(&in);
536  av_frame_free(&out);
537  return ret;
538 }
539 
540 #define OFFSET(x) offsetof(CUDAScaleContext, x)
541 #define FLAGS (AV_OPT_FLAG_FILTERING_PARAM|AV_OPT_FLAG_VIDEO_PARAM)
542 static const AVOption options[] = {
543  { "w", "Output video width", OFFSET(w_expr), AV_OPT_TYPE_STRING, { .str = "iw" }, .flags = FLAGS },
544  { "h", "Output video height", OFFSET(h_expr), AV_OPT_TYPE_STRING, { .str = "ih" }, .flags = FLAGS },
545  { "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" },
546  { "disable", NULL, 0, AV_OPT_TYPE_CONST, {.i64 = 0 }, 0, 0, FLAGS, "force_oar" },
547  { "decrease", NULL, 0, AV_OPT_TYPE_CONST, {.i64 = 1 }, 0, 0, FLAGS, "force_oar" },
548  { "increase", NULL, 0, AV_OPT_TYPE_CONST, {.i64 = 2 }, 0, 0, FLAGS, "force_oar" },
549  { "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 },
550  { NULL },
551 };
552 
553 static const AVClass cudascale_class = {
554  .class_name = "cudascale",
555  .item_name = av_default_item_name,
556  .option = options,
557  .version = LIBAVUTIL_VERSION_INT,
558 };
559 
560 static const AVFilterPad cudascale_inputs[] = {
561  {
562  .name = "default",
563  .type = AVMEDIA_TYPE_VIDEO,
564  .filter_frame = cudascale_filter_frame,
565  },
566  { NULL }
567 };
568 
569 static const AVFilterPad cudascale_outputs[] = {
570  {
571  .name = "default",
572  .type = AVMEDIA_TYPE_VIDEO,
573  .config_props = cudascale_config_props,
574  },
575  { NULL }
576 };
577 
579  .name = "scale_cuda",
580  .description = NULL_IF_CONFIG_SMALL("GPU accelerated video resizer"),
581 
582  .init = cudascale_init,
583  .uninit = cudascale_uninit,
584  .query_formats = cudascale_query_formats,
585 
586  .priv_size = sizeof(CUDAScaleContext),
587  .priv_class = &cudascale_class,
588 
589  .inputs = cudascale_inputs,
590  .outputs = cudascale_outputs,
591 
592  .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE,
593 };
int force_original_aspect_ratio
Definition: vf_scale_cuda.c:85
#define NULL
Definition: coverity.c:32
#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:385
enum AVPixelFormat format
Output sw format.
Definition: vf_scale_cuda.c:80
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:125
This structure describes decoded (raw) audio or video data.
Definition: frame.h:295
BYTE int const BYTE int src_pitch
Definition: avisynth_c.h:908
AVOption.
Definition: opt.h:246
const char * fmt
Definition: avisynth_c.h:861
CUfunction cu_func_uchar
Definition: vf_scale_cuda.c:90
static int call_resize_kernel(AVFilterContext *ctx, CUfunction func, int channels, uint8_t *src_dptr, int src_width, int src_height, int src_pitch, uint8_t *dst_dptr, int dst_width, int dst_height, int dst_pitch, int pixel_size)
planar YUV 4:4:4, 24bpp, (1 Cr & Cb sample per 1x1 Y samples)
Definition: pixfmt.h:71
#define LIBAVUTIL_VERSION_INT
Definition: version.h:85
Main libavfilter public API header.
channels
Definition: aptx.c:30
AVCUDADeviceContextInternal * internal
int num
Numerator.
Definition: rational.h:59
const char * av_default_item_name(void *ptr)
Return the context name.
Definition: log.c:191
int width
The allocated dimensions of the frames in this pool.
Definition: hwcontext.h:228
enum AVPixelFormat format
The pixel format identifying the underlying HW surface type.
Definition: hwcontext.h:208
AVFrame * tmp_frame
Definition: vf_scale_cuda.c:74
CUcontext cu_ctx
Definition: vf_scale_cuda.c:88
void av_frame_move_ref(AVFrame *dst, AVFrame *src)
Move everything contained in src to dst and reset src.
Definition: frame.c:582
char * w_expr
width expression string
Definition: vf_scale_cuda.c:82
AVCUDADeviceContext * hwctx
Definition: vf_scale_cuda.c:61
#define src
Definition: vp8dsp.c:254
#define BLOCKX
Definition: vf_scale_cuda.c:53
CUdeviceptr srcBuffer
Definition: vf_scale_cuda.c:98
struct CUDAScaleContext::@247 planes_out[3]
CUfunction cu_func_ushort
Definition: vf_scale_cuda.c:93
#define FLAGS
struct CUDAScaleContext::@247 planes_in[3]
AVFilterFormats * ff_make_format_list(const int *fmts)
Create a list of supported formats.
Definition: formats.c:283
AVBufferRef * frames_ctx
Definition: vf_scale_cuda.c:71
AVBufferRef * hw_frames_ctx
For hwaccel-format frames, this should be a reference to the AVHWFramesContext describing the frame...
Definition: frame.h:634
#define AV_PIX_FMT_P016
Definition: pixfmt.h:437
CUstream cu_stream
Definition: vf_scale_cuda.c:96
const char * name
Pad name.
Definition: internal.h:60
#define AV_PIX_FMT_P010
Definition: pixfmt.h:436
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:72
AVFilterLink ** inputs
array of pointers to input links
Definition: avfilter.h:346
static int scalecuda_resize(AVFilterContext *ctx, AVFrame *out, AVFrame *in)
int ff_filter_frame(AVFilterLink *link, AVFrame *frame)
Send a frame of data to the next filter.
Definition: avfilter.c:1093
uint8_t
#define av_cold
Definition: attributes.h:82
#define BLOCKY
Definition: vf_scale_cuda.c:54
AVFrame * av_frame_alloc(void)
Allocate an AVFrame and set its fields to default values.
Definition: frame.c:189
AVOptions.
static enum AVPixelFormat supported_formats[]
Definition: vf_scale_cuda.c:41
static const AVFilterPad cudascale_outputs[]
CUfunction cu_func_ushort4
Definition: vf_scale_cuda.c:95
void * hwctx
The format-specific data, allocated and freed by libavutil along with this context.
Definition: hwcontext.h:91
static int cudascale_scale(AVFilterContext *ctx, AVFrame *out, AVFrame *in)
#define AV_LOG_VERBOSE
Detailed information.
Definition: log.h:192
#define AV_PIX_FMT_YUV444P16
Definition: pixfmt.h:400
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
static av_cold int init_processing_chain(AVFilterContext *ctx, int in_width, int in_height, int out_width, int out_height)
#define FFALIGN(x, a)
Definition: macros.h:48
#define av_log(a,...)
A filter pad used for either input or output.
Definition: internal.h:54
#define i(width, name, range_min, range_max)
Definition: cbs_h2645.c:259
int width
Definition: frame.h:353
#define AV_LOG_ERROR
Something went wrong and cannot losslessly be recovered.
Definition: log.h:176
int ff_set_common_formats(AVFilterContext *ctx, AVFilterFormats *formats)
A helper for query_formats() which sets all links to the same list of formats.
Definition: formats.c:569
static int format_is_supported(enum AVPixelFormat fmt)
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:106
#define DIV_UP(a, b)
Definition: vf_scale_cuda.c:50
void av_frame_free(AVFrame **frame)
Free the frame and any dynamically allocated objects in it, e.g.
Definition: frame.c:202
int av_pix_fmt_get_chroma_sub_sample(enum AVPixelFormat pix_fmt, int *h_shift, int *v_shift)
Utility function to access log2_chroma_w log2_chroma_h from the pixel format AVPixFmtDescriptor.
Definition: pixdesc.c:2550
#define NULL_IF_CONFIG_SMALL(x)
Return NULL if CONFIG_SMALL is true, otherwise the argument without modification. ...
Definition: internal.h:186
void * priv
private data for use by the filter
Definition: avfilter.h:353
char * h_expr
height expression string
Definition: vf_scale_cuda.c:83
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
like NV12, with 16bpp per component, little-endian
Definition: pixfmt.h:300
int av_hwframe_ctx_init(AVBufferRef *ref)
Finalize the context before use.
Definition: hwcontext.c:329
int av_hwframe_get_buffer(AVBufferRef *hwframe_ref, AVFrame *frame, int flags)
Allocate a new frame attached to the given AVHWFramesContext.
Definition: hwcontext.c:465
#define fail()
Definition: checkasm.h:122
static int cudascale_filter_frame(AVFilterLink *link, AVFrame *in)
common internal API header
#define OFFSET(x)
AVHWDeviceContext * device_ctx
The parent AVHWDeviceContext.
Definition: hwcontext.h:148
CUmodule cu_module
Definition: vf_scale_cuda.c:89
uint8_t w
Definition: llviddspenc.c:38
AVFormatContext * ctx
Definition: movenc.c:48
static const AVClass cudascale_class
#define s(width, name)
Definition: cbs_vp9.c:257
FFmpeg internal API for CUDA.
CUfunction cu_func_uchar2
Definition: vf_scale_cuda.c:91
int dummy
Definition: motion.c:64
like NV12, with 10bpp per component, data in the high bits, zeros in the low bits, little-endian
Definition: pixfmt.h:284
CUdeviceptr dstBuffer
Definition: vf_scale_cuda.c:99
HW acceleration through CUDA.
Definition: pixfmt.h:235
static const AVOption options[]
static const AVFilterPad outputs[]
Definition: af_acontrast.c:203
#define FF_ARRAY_ELEMS(a)
int linesize[AV_NUM_DATA_POINTERS]
For video, size in bytes of each picture line.
Definition: frame.h:326
AVFilter ff_vf_scale_cuda
uint8_t * data
The data buffer.
Definition: buffer.h:89
AVRational sample_aspect_ratio
Sample aspect ratio for the video frame, 0/1 if unknown/unspecified.
Definition: frame.h:383
these buffered frames must be flushed immediately if a new input produces new the filter must not call request_frame to get more It must just process the frame or queue it The task of requesting more frames is left to the filter s request_frame method or the application If a filter has several inputs
#define AVERROR_BUG
Internal bug, also see AVERROR_BUG2.
Definition: error.h:50
uint8_t pi<< 24) CONV_FUNC_GROUP(AV_SAMPLE_FMT_FLT, float, AV_SAMPLE_FMT_U8, uint8_t,(*(const uint8_t *) pi-0x80)*(1.0f/(1<< 7))) CONV_FUNC_GROUP(AV_SAMPLE_FMT_DBL, double, AV_SAMPLE_FMT_U8, uint8_t,(*(const uint8_t *) pi-0x80)*(1.0/(1<< 7))) CONV_FUNC_GROUP(AV_SAMPLE_FMT_U8, uint8_t, AV_SAMPLE_FMT_S16, int16_t,(*(const int16_t *) pi >> 8)+0x80) CONV_FUNC_GROUP(AV_SAMPLE_FMT_FLT, float, AV_SAMPLE_FMT_S16, int16_t,*(const int16_t *) pi *(1.0f/(1<< 15))) CONV_FUNC_GROUP(AV_SAMPLE_FMT_DBL, double, AV_SAMPLE_FMT_S16, int16_t,*(const int16_t *) pi *(1.0/(1<< 15))) CONV_FUNC_GROUP(AV_SAMPLE_FMT_U8, uint8_t, AV_SAMPLE_FMT_S32, int32_t,(*(const int32_t *) pi >> 24)+0x80) CONV_FUNC_GROUP(AV_SAMPLE_FMT_FLT, float, AV_SAMPLE_FMT_S32, int32_t,*(const int32_t *) pi *(1.0f/(1U<< 31))) CONV_FUNC_GROUP(AV_SAMPLE_FMT_DBL, double, AV_SAMPLE_FMT_S32, int32_t,*(const int32_t *) pi *(1.0/(1U<< 31))) CONV_FUNC_GROUP(AV_SAMPLE_FMT_U8, uint8_t, AV_SAMPLE_FMT_FLT, float, av_clip_uint8(lrintf(*(const float *) pi *(1<< 7))+0x80)) CONV_FUNC_GROUP(AV_SAMPLE_FMT_S16, int16_t, AV_SAMPLE_FMT_FLT, float, av_clip_int16(lrintf(*(const float *) pi *(1<< 15)))) CONV_FUNC_GROUP(AV_SAMPLE_FMT_S32, int32_t, AV_SAMPLE_FMT_FLT, float, av_clipl_int32(llrintf(*(const float *) pi *(1U<< 31)))) CONV_FUNC_GROUP(AV_SAMPLE_FMT_U8, uint8_t, AV_SAMPLE_FMT_DBL, double, av_clip_uint8(lrint(*(const double *) pi *(1<< 7))+0x80)) CONV_FUNC_GROUP(AV_SAMPLE_FMT_S16, int16_t, AV_SAMPLE_FMT_DBL, double, av_clip_int16(lrint(*(const double *) pi *(1<< 15)))) CONV_FUNC_GROUP(AV_SAMPLE_FMT_S32, int32_t, AV_SAMPLE_FMT_DBL, double, av_clipl_int32(llrint(*(const double *) pi *(1U<< 31))))#define SET_CONV_FUNC_GROUP(ofmt, ifmt) static void set_generic_function(AudioConvert *ac){}void ff_audio_convert_free(AudioConvert **ac){if(!*ac) return;ff_dither_free(&(*ac) ->dc);av_freep(ac);}AudioConvert *ff_audio_convert_alloc(AVAudioResampleContext *avr, enum AVSampleFormat out_fmt, enum AVSampleFormat in_fmt, int channels, int sample_rate, int apply_map){AudioConvert *ac;int in_planar, out_planar;ac=av_mallocz(sizeof(*ac));if(!ac) return NULL;ac->avr=avr;ac->out_fmt=out_fmt;ac->in_fmt=in_fmt;ac->channels=channels;ac->apply_map=apply_map;if(avr->dither_method!=AV_RESAMPLE_DITHER_NONE &&av_get_packed_sample_fmt(out_fmt)==AV_SAMPLE_FMT_S16 &&av_get_bytes_per_sample(in_fmt) > 2){ac->dc=ff_dither_alloc(avr, out_fmt, in_fmt, channels, sample_rate, apply_map);if(!ac->dc){av_free(ac);return NULL;}return ac;}in_planar=ff_sample_fmt_is_planar(in_fmt, channels);out_planar=ff_sample_fmt_is_planar(out_fmt, channels);if(in_planar==out_planar){ac->func_type=CONV_FUNC_TYPE_FLAT;ac->planes=in_planar?ac->channels:1;}else if(in_planar) ac->func_type=CONV_FUNC_TYPE_INTERLEAVE;else ac->func_type=CONV_FUNC_TYPE_DEINTERLEAVE;set_generic_function(ac);if(ARCH_AARCH64) ff_audio_convert_init_aarch64(ac);if(ARCH_ARM) ff_audio_convert_init_arm(ac);if(ARCH_X86) ff_audio_convert_init_x86(ac);return ac;}int ff_audio_convert(AudioConvert *ac, AudioData *out, AudioData *in){int use_generic=1;int len=in->nb_samples;int p;if(ac->dc){av_log(ac->avr, AV_LOG_TRACE,"%d samples - audio_convert: %s to %s (dithered)\n", len, av_get_sample_fmt_name(ac->in_fmt), av_get_sample_fmt_name(ac->out_fmt));return ff_convert_dither(ac-> in
static av_cold void cudascale_uninit(AVFilterContext *ctx)
This struct is allocated as AVHWDeviceContext.hwctx.
int(* func)(AVBPrint *dst, const char *in, const char *arg)
Definition: jacosubdec.c:67
Describe the class of an AVClass context structure.
Definition: log.h:67
Filter definition.
Definition: avfilter.h:144
Rational number (pair of numerator and denominator).
Definition: rational.h:58
This struct describes a set or pool of "hardware" frames (i.e.
Definition: hwcontext.h:123
static av_cold int cudascale_config_props(AVFilterLink *outlink)
const char * name
Filter name.
Definition: avfilter.h:148
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
AVFilterLink ** outputs
array of pointers to output links
Definition: avfilter.h:350
static enum AVPixelFormat pix_fmts[]
Definition: libkvazaar.c:275
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:186
static const AVFilterPad cudascale_inputs[]
void av_frame_unref(AVFrame *frame)
Unreference all the buffers referenced by frame and reset the frame fields.
Definition: frame.c:553
uint8_t * data[AV_NUM_DATA_POINTERS]
pointer to the picture/channel planes.
Definition: frame.h:309
CUfunction cu_func_ushort2
Definition: vf_scale_cuda.c:94
AVBufferRef * device_ref
A reference to the parent AVHWDeviceContext.
Definition: hwcontext.h:140
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
A reference to a data buffer.
Definition: buffer.h:81
planar YUV 4:2:0, 12bpp, (1 Cr & Cb sample per 2x2 Y samples)
Definition: pixfmt.h:66
common internal and external API header
static av_cold int cudascale_init(AVFilterContext *ctx)
AVBufferRef * av_hwframe_ctx_alloc(AVBufferRef *device_ref_in)
Allocate an AVHWFramesContext tied to a given device context.
Definition: hwcontext.c:243
static int cudascale_query_formats(AVFilterContext *ctx)
AVBufferRef * av_buffer_ref(AVBufferRef *buf)
Create a new reference to an AVBuffer.
Definition: buffer.c:93
#define CHECK_CU(x)
Definition: vf_scale_cuda.c:56
enum AVPixelFormat in_fmt
Definition: vf_scale_cuda.c:63
int den
Denominator.
Definition: rational.h:60
CUfunction cu_func_uchar4
Definition: vf_scale_cuda.c:92
A list of supported formats for one end of a filter link.
Definition: formats.h:64
An instance of a filter.
Definition: avfilter.h:338
AVRational av_mul_q(AVRational b, AVRational c)
Multiply two rationals.
Definition: rational.c:80
int height
Definition: frame.h:353
FILE * out
Definition: movenc.c:54
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:2438
BYTE int dst_pitch
Definition: avisynth_c.h:908
internal API functions
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
enum AVPixelFormat sw_format
The pixel format identifying the actual data layout of the hardware frames.
Definition: hwcontext.h:221
AVPixelFormat
Pixel format.
Definition: pixfmt.h:64
static av_cold int init_stage(CUDAScaleContext *s, AVBufferRef *device_ctx)
enum AVPixelFormat out_fmt
Definition: vf_scale_cuda.c:64
int av_frame_copy_props(AVFrame *dst, const AVFrame *src)
Copy only "metadata" fields from src to dst.
Definition: frame.c:654