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  if (s->hwctx && s->cu_module) {
124  CudaFunctions *cu = s->hwctx->internal->cuda_dl;
125  CUcontext dummy;
126 
127  CHECK_CU(cu->cuCtxPushCurrent(s->hwctx->cuda_ctx));
128  CHECK_CU(cu->cuModuleUnload(s->cu_module));
129  s->cu_module = NULL;
130  CHECK_CU(cu->cuCtxPopCurrent(&dummy));
131  }
132 
133  av_frame_free(&s->frame);
136 }
137 
139 {
140  static const enum AVPixelFormat pixel_formats[] = {
142  };
143  AVFilterFormats *pix_fmts = ff_make_format_list(pixel_formats);
144 
145  return ff_set_common_formats(ctx, pix_fmts);
146 }
147 
149 {
150  AVBufferRef *out_ref = NULL;
151  AVHWFramesContext *out_ctx;
152  int in_sw, in_sh, out_sw, out_sh;
153  int ret, i;
154 
155  av_pix_fmt_get_chroma_sub_sample(s->in_fmt, &in_sw, &in_sh);
156  av_pix_fmt_get_chroma_sub_sample(s->out_fmt, &out_sw, &out_sh);
157  if (!s->planes_out[0].width) {
158  s->planes_out[0].width = s->planes_in[0].width;
159  s->planes_out[0].height = s->planes_in[0].height;
160  }
161 
162  for (i = 1; i < FF_ARRAY_ELEMS(s->planes_in); i++) {
163  s->planes_in[i].width = s->planes_in[0].width >> in_sw;
164  s->planes_in[i].height = s->planes_in[0].height >> in_sh;
165  s->planes_out[i].width = s->planes_out[0].width >> out_sw;
166  s->planes_out[i].height = s->planes_out[0].height >> out_sh;
167  }
168 
169  out_ref = av_hwframe_ctx_alloc(device_ctx);
170  if (!out_ref)
171  return AVERROR(ENOMEM);
172  out_ctx = (AVHWFramesContext*)out_ref->data;
173 
174  out_ctx->format = AV_PIX_FMT_CUDA;
175  out_ctx->sw_format = s->out_fmt;
176  out_ctx->width = FFALIGN(s->planes_out[0].width, 32);
177  out_ctx->height = FFALIGN(s->planes_out[0].height, 32);
178 
179  ret = av_hwframe_ctx_init(out_ref);
180  if (ret < 0)
181  goto fail;
182 
183  av_frame_unref(s->frame);
184  ret = av_hwframe_get_buffer(out_ref, s->frame, 0);
185  if (ret < 0)
186  goto fail;
187 
188  s->frame->width = s->planes_out[0].width;
189  s->frame->height = s->planes_out[0].height;
190 
192  s->frames_ctx = out_ref;
193 
194  return 0;
195 fail:
196  av_buffer_unref(&out_ref);
197  return ret;
198 }
199 
200 static int format_is_supported(enum AVPixelFormat fmt)
201 {
202  int i;
203 
204  for (i = 0; i < FF_ARRAY_ELEMS(supported_formats); i++)
205  if (supported_formats[i] == fmt)
206  return 1;
207  return 0;
208 }
209 
210 static av_cold int init_processing_chain(AVFilterContext *ctx, int in_width, int in_height,
211  int out_width, int out_height)
212 {
213  CUDAScaleContext *s = ctx->priv;
214 
215  AVHWFramesContext *in_frames_ctx;
216 
217  enum AVPixelFormat in_format;
218  enum AVPixelFormat out_format;
219  int ret;
220 
221  /* check that we have a hw context */
222  if (!ctx->inputs[0]->hw_frames_ctx) {
223  av_log(ctx, AV_LOG_ERROR, "No hw context provided on input\n");
224  return AVERROR(EINVAL);
225  }
226  in_frames_ctx = (AVHWFramesContext*)ctx->inputs[0]->hw_frames_ctx->data;
227  in_format = in_frames_ctx->sw_format;
228  out_format = (s->format == AV_PIX_FMT_NONE) ? in_format : s->format;
229 
230  if (!format_is_supported(in_format)) {
231  av_log(ctx, AV_LOG_ERROR, "Unsupported input format: %s\n",
232  av_get_pix_fmt_name(in_format));
233  return AVERROR(ENOSYS);
234  }
235  if (!format_is_supported(out_format)) {
236  av_log(ctx, AV_LOG_ERROR, "Unsupported output format: %s\n",
237  av_get_pix_fmt_name(out_format));
238  return AVERROR(ENOSYS);
239  }
240 
241  if (in_width == out_width && in_height == out_height)
242  s->passthrough = 1;
243 
244  s->in_fmt = in_format;
245  s->out_fmt = out_format;
246 
247  s->planes_in[0].width = in_width;
248  s->planes_in[0].height = in_height;
249  s->planes_out[0].width = out_width;
250  s->planes_out[0].height = out_height;
251 
252  ret = init_stage(s, in_frames_ctx->device_ref);
253  if (ret < 0)
254  return ret;
255 
257  if (!ctx->outputs[0]->hw_frames_ctx)
258  return AVERROR(ENOMEM);
259 
260  return 0;
261 }
262 
264 {
265  AVFilterContext *ctx = outlink->src;
266  AVFilterLink *inlink = outlink->src->inputs[0];
267  CUDAScaleContext *s = ctx->priv;
269  AVCUDADeviceContext *device_hwctx = frames_ctx->device_ctx->hwctx;
270  CUcontext dummy, cuda_ctx = device_hwctx->cuda_ctx;
271  CudaFunctions *cu = device_hwctx->internal->cuda_dl;
272  int w, h;
273  int ret;
274 
275  extern char vf_scale_cuda_ptx[];
276 
277  s->hwctx = device_hwctx;
278  s->cu_stream = s->hwctx->stream;
279 
280  ret = CHECK_CU(cu->cuCtxPushCurrent(cuda_ctx));
281  if (ret < 0)
282  goto fail;
283 
284  ret = CHECK_CU(cu->cuModuleLoadData(&s->cu_module, vf_scale_cuda_ptx));
285  if (ret < 0)
286  goto fail;
287 
288  CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_uchar, s->cu_module, "Subsample_Bilinear_uchar"));
289  if (ret < 0)
290  goto fail;
291 
292  CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_uchar2, s->cu_module, "Subsample_Bilinear_uchar2"));
293  if (ret < 0)
294  goto fail;
295 
296  CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_uchar4, s->cu_module, "Subsample_Bilinear_uchar4"));
297  if (ret < 0)
298  goto fail;
299 
300  CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_ushort, s->cu_module, "Subsample_Bilinear_ushort"));
301  if (ret < 0)
302  goto fail;
303 
304  CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_ushort2, s->cu_module, "Subsample_Bilinear_ushort2"));
305  if (ret < 0)
306  goto fail;
307 
308  CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_ushort4, s->cu_module, "Subsample_Bilinear_ushort4"));
309  if (ret < 0)
310  goto fail;
311 
312 
313  CHECK_CU(cu->cuCtxPopCurrent(&dummy));
314 
315  if ((ret = ff_scale_eval_dimensions(s,
316  s->w_expr, s->h_expr,
317  inlink, outlink,
318  &w, &h)) < 0)
319  goto fail;
320 
321  ff_scale_adjust_dimensions(inlink, &w, &h,
323 
324  if (((int64_t)h * inlink->w) > INT_MAX ||
325  ((int64_t)w * inlink->h) > INT_MAX)
326  av_log(ctx, AV_LOG_ERROR, "Rescaled value for width or height is too big.\n");
327 
328  outlink->w = w;
329  outlink->h = h;
330 
331  ret = init_processing_chain(ctx, inlink->w, inlink->h, w, h);
332  if (ret < 0)
333  return ret;
334 
335  av_log(ctx, AV_LOG_VERBOSE, "w:%d h:%d -> w:%d h:%d\n",
336  inlink->w, inlink->h, outlink->w, outlink->h);
337 
338  if (inlink->sample_aspect_ratio.num) {
339  outlink->sample_aspect_ratio = av_mul_q((AVRational){outlink->h*inlink->w,
340  outlink->w*inlink->h},
341  inlink->sample_aspect_ratio);
342  } else {
343  outlink->sample_aspect_ratio = inlink->sample_aspect_ratio;
344  }
345 
346  return 0;
347 
348 fail:
349  return ret;
350 }
351 
352 static int call_resize_kernel(AVFilterContext *ctx, CUfunction func, int channels,
353  uint8_t *src_dptr, int src_width, int src_height, int src_pitch,
354  uint8_t *dst_dptr, int dst_width, int dst_height, int dst_pitch,
355  int pixel_size)
356 {
357  CUDAScaleContext *s = ctx->priv;
358  CudaFunctions *cu = s->hwctx->internal->cuda_dl;
359  CUdeviceptr dst_devptr = (CUdeviceptr)dst_dptr;
360  CUtexObject tex = 0;
361  void *args_uchar[] = { &tex, &dst_devptr, &dst_width, &dst_height, &dst_pitch, &src_width, &src_height };
362  int ret;
363 
364  CUDA_TEXTURE_DESC tex_desc = {
365  .filterMode = CU_TR_FILTER_MODE_LINEAR,
366  .flags = CU_TRSF_READ_AS_INTEGER,
367  };
368 
369  CUDA_RESOURCE_DESC res_desc = {
370  .resType = CU_RESOURCE_TYPE_PITCH2D,
371  .res.pitch2D.format = pixel_size == 1 ?
372  CU_AD_FORMAT_UNSIGNED_INT8 :
373  CU_AD_FORMAT_UNSIGNED_INT16,
374  .res.pitch2D.numChannels = channels,
375  .res.pitch2D.width = src_width,
376  .res.pitch2D.height = src_height,
377  .res.pitch2D.pitchInBytes = src_pitch * pixel_size,
378  .res.pitch2D.devPtr = (CUdeviceptr)src_dptr,
379  };
380 
381  ret = CHECK_CU(cu->cuTexObjectCreate(&tex, &res_desc, &tex_desc, NULL));
382  if (ret < 0)
383  goto exit;
384 
385  ret = CHECK_CU(cu->cuLaunchKernel(func,
386  DIV_UP(dst_width, BLOCKX), DIV_UP(dst_height, BLOCKY), 1,
387  BLOCKX, BLOCKY, 1, 0, s->cu_stream, args_uchar, NULL));
388 
389 exit:
390  if (tex)
391  CHECK_CU(cu->cuTexObjectDestroy(tex));
392 
393  return ret;
394 }
395 
397  AVFrame *out, AVFrame *in)
398 {
399  AVHWFramesContext *in_frames_ctx = (AVHWFramesContext*)in->hw_frames_ctx->data;
400  CUDAScaleContext *s = ctx->priv;
401 
402  switch (in_frames_ctx->sw_format) {
403  case AV_PIX_FMT_YUV420P:
404  call_resize_kernel(ctx, s->cu_func_uchar, 1,
405  in->data[0], in->width, in->height, in->linesize[0],
406  out->data[0], out->width, out->height, out->linesize[0],
407  1);
408  call_resize_kernel(ctx, s->cu_func_uchar, 1,
409  in->data[1], in->width/2, in->height/2, in->linesize[0]/2,
410  out->data[1], out->width/2, out->height/2, out->linesize[0]/2,
411  1);
412  call_resize_kernel(ctx, s->cu_func_uchar, 1,
413  in->data[2], in->width/2, in->height/2, in->linesize[0]/2,
414  out->data[2], out->width/2, out->height/2, out->linesize[0]/2,
415  1);
416  break;
417  case AV_PIX_FMT_YUV444P:
418  call_resize_kernel(ctx, s->cu_func_uchar, 1,
419  in->data[0], in->width, in->height, in->linesize[0],
420  out->data[0], out->width, out->height, out->linesize[0],
421  1);
422  call_resize_kernel(ctx, s->cu_func_uchar, 1,
423  in->data[1], in->width, in->height, in->linesize[0],
424  out->data[1], out->width, out->height, out->linesize[0],
425  1);
426  call_resize_kernel(ctx, s->cu_func_uchar, 1,
427  in->data[2], in->width, in->height, in->linesize[0],
428  out->data[2], out->width, out->height, out->linesize[0],
429  1);
430  break;
432  call_resize_kernel(ctx, s->cu_func_ushort, 1,
433  in->data[0], in->width, in->height, in->linesize[0] / 2,
434  out->data[0], out->width, out->height, out->linesize[0] / 2,
435  2);
436  call_resize_kernel(ctx, s->cu_func_ushort, 1,
437  in->data[1], in->width, in->height, in->linesize[1] / 2,
438  out->data[1], out->width, out->height, out->linesize[1] / 2,
439  2);
440  call_resize_kernel(ctx, s->cu_func_ushort, 1,
441  in->data[2], in->width, in->height, in->linesize[2] / 2,
442  out->data[2], out->width, out->height, out->linesize[2] / 2,
443  2);
444  break;
445  case AV_PIX_FMT_NV12:
446  call_resize_kernel(ctx, s->cu_func_uchar, 1,
447  in->data[0], in->width, in->height, in->linesize[0],
448  out->data[0], out->width, out->height, out->linesize[0],
449  1);
450  call_resize_kernel(ctx, s->cu_func_uchar2, 2,
451  in->data[1], in->width/2, in->height/2, in->linesize[1],
452  out->data[1], out->width/2, out->height/2, out->linesize[1]/2,
453  1);
454  break;
455  case AV_PIX_FMT_P010LE:
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  case AV_PIX_FMT_P016LE:
466  call_resize_kernel(ctx, s->cu_func_ushort, 1,
467  in->data[0], in->width, in->height, in->linesize[0] / 2,
468  out->data[0], out->width, out->height, out->linesize[0] / 2,
469  2);
470  call_resize_kernel(ctx, s->cu_func_ushort2, 2,
471  in->data[1], in->width / 2, in->height / 2, in->linesize[1] / 2,
472  out->data[1], out->width / 2, out->height / 2, out->linesize[1] / 4,
473  2);
474  break;
475  default:
476  return AVERROR_BUG;
477  }
478 
479  return 0;
480 }
481 
483 {
484  CUDAScaleContext *s = ctx->priv;
485  AVFrame *src = in;
486  int ret;
487 
488  ret = scalecuda_resize(ctx, s->frame, src);
489  if (ret < 0)
490  return ret;
491 
492  src = s->frame;
493  ret = av_hwframe_get_buffer(src->hw_frames_ctx, s->tmp_frame, 0);
494  if (ret < 0)
495  return ret;
496 
497  av_frame_move_ref(out, s->frame);
499 
500  s->frame->width = s->planes_out[0].width;
501  s->frame->height = s->planes_out[0].height;
502 
503  ret = av_frame_copy_props(out, in);
504  if (ret < 0)
505  return ret;
506 
507  return 0;
508 }
509 
511 {
512  AVFilterContext *ctx = link->dst;
513  CUDAScaleContext *s = ctx->priv;
514  AVFilterLink *outlink = ctx->outputs[0];
515  CudaFunctions *cu = s->hwctx->internal->cuda_dl;
516 
517  AVFrame *out = NULL;
518  CUcontext dummy;
519  int ret = 0;
520 
521  out = av_frame_alloc();
522  if (!out) {
523  ret = AVERROR(ENOMEM);
524  goto fail;
525  }
526 
527  ret = CHECK_CU(cu->cuCtxPushCurrent(s->hwctx->cuda_ctx));
528  if (ret < 0)
529  goto fail;
530 
531  ret = cudascale_scale(ctx, out, in);
532 
533  CHECK_CU(cu->cuCtxPopCurrent(&dummy));
534  if (ret < 0)
535  goto fail;
536 
538  (int64_t)in->sample_aspect_ratio.num * outlink->h * link->w,
539  (int64_t)in->sample_aspect_ratio.den * outlink->w * link->h,
540  INT_MAX);
541 
542  av_frame_free(&in);
543  return ff_filter_frame(outlink, out);
544 fail:
545  av_frame_free(&in);
546  av_frame_free(&out);
547  return ret;
548 }
549 
550 #define OFFSET(x) offsetof(CUDAScaleContext, x)
551 #define FLAGS (AV_OPT_FLAG_FILTERING_PARAM|AV_OPT_FLAG_VIDEO_PARAM)
552 static const AVOption options[] = {
553  { "w", "Output video width", OFFSET(w_expr), AV_OPT_TYPE_STRING, { .str = "iw" }, .flags = FLAGS },
554  { "h", "Output video height", OFFSET(h_expr), AV_OPT_TYPE_STRING, { .str = "ih" }, .flags = FLAGS },
555  { "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" },
556  { "disable", NULL, 0, AV_OPT_TYPE_CONST, {.i64 = 0 }, 0, 0, FLAGS, "force_oar" },
557  { "decrease", NULL, 0, AV_OPT_TYPE_CONST, {.i64 = 1 }, 0, 0, FLAGS, "force_oar" },
558  { "increase", NULL, 0, AV_OPT_TYPE_CONST, {.i64 = 2 }, 0, 0, FLAGS, "force_oar" },
559  { "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 },
560  { NULL },
561 };
562 
563 static const AVClass cudascale_class = {
564  .class_name = "cudascale",
565  .item_name = av_default_item_name,
566  .option = options,
567  .version = LIBAVUTIL_VERSION_INT,
568 };
569 
570 static const AVFilterPad cudascale_inputs[] = {
571  {
572  .name = "default",
573  .type = AVMEDIA_TYPE_VIDEO,
574  .filter_frame = cudascale_filter_frame,
575  },
576  { NULL }
577 };
578 
579 static const AVFilterPad cudascale_outputs[] = {
580  {
581  .name = "default",
582  .type = AVMEDIA_TYPE_VIDEO,
583  .config_props = cudascale_config_props,
584  },
585  { NULL }
586 };
587 
589  .name = "scale_cuda",
590  .description = NULL_IF_CONFIG_SMALL("GPU accelerated video resizer"),
591 
592  .init = cudascale_init,
593  .uninit = cudascale_uninit,
594  .query_formats = cudascale_query_formats,
595 
596  .priv_size = sizeof(CUDAScaleContext),
597  .priv_class = &cudascale_class,
598 
599  .inputs = cudascale_inputs,
600  .outputs = cudascale_outputs,
601 
602  .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE,
603 };
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:339
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:308
AVOption.
Definition: opt.h:248
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.
AVCUDADeviceContextInternal * internal
struct CUDAScaleContext::@236 planes_out[3]
int num
Numerator.
Definition: rational.h:59
const char * av_default_item_name(void *ptr)
Return the context name.
Definition: log.c:235
int width
The allocated dimensions of the frames in this pool.
Definition: hwcontext.h:229
enum AVPixelFormat format
The pixel format identifying the underlying HW surface type.
Definition: hwcontext.h:209
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:595
char * w_expr
width expression string
Definition: vf_scale_cuda.c:82
AVCUDADeviceContext * hwctx
Definition: vf_scale_cuda.c:61
#define BLOCKX
Definition: vf_scale_cuda.c:53
CUdeviceptr srcBuffer
Definition: vf_scale_cuda.c:98
CUfunction cu_func_ushort
Definition: vf_scale_cuda.c:93
#define FLAGS
AVFilterFormats * ff_make_format_list(const int *fmts)
Create a list of supported formats.
Definition: formats.c:287
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:647
#define AV_PIX_FMT_P016
Definition: pixfmt.h:449
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:448
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:349
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:1091
uint8_t
#define av_cold
Definition: attributes.h:88
#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:190
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:92
static int cudascale_scale(AVFilterContext *ctx, AVFrame *out, AVFrame *in)
#define AV_LOG_VERBOSE
Detailed information.
Definition: log.h:210
#define AV_PIX_FMT_YUV444P16
Definition: pixfmt.h:412
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
channels
Definition: aptx.h:33
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 src
Definition: vp8dsp.c:254
int width
Definition: frame.h:366
#define AV_LOG_ERROR
Something went wrong and cannot losslessly be recovered.
Definition: log.h:194
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:588
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:57
#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:203
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:2601
#define NULL_IF_CONFIG_SMALL(x)
Return NULL if CONFIG_SMALL is true, otherwise the argument without modification. ...
Definition: internal.h:153
void * priv
private data for use by the filter
Definition: avfilter.h:356
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:333
int av_hwframe_get_buffer(AVBufferRef *hwframe_ref, AVFrame *frame, int flags)
Allocate a new frame attached to the given AVHWFramesContext.
Definition: hwcontext.c:502
#define fail()
Definition: checkasm.h:123
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:149
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)
struct CUDAScaleContext::@236 planes_in[3]
int linesize[AV_NUM_DATA_POINTERS]
For video, size in bytes of each picture line.
Definition: frame.h:339
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:396
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:145
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:124
static av_cold int cudascale_config_props(AVFilterLink *outlink)
const char * name
Filter name.
Definition: avfilter.h:149
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:353
static enum AVPixelFormat pix_fmts[]
Definition: libkvazaar.c:300
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
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:566
uint8_t * data[AV_NUM_DATA_POINTERS]
pointer to the picture/channel planes.
Definition: frame.h:322
CUfunction cu_func_ushort2
Definition: vf_scale_cuda.c:94
AVBufferRef * device_ref
A reference to the parent AVHWDeviceContext.
Definition: hwcontext.h:141
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:247
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:65
An instance of a filter.
Definition: avfilter.h:341
AVRational av_mul_q(AVRational b, AVRational c)
Multiply two rationals.
Definition: rational.c:80
int height
Definition: frame.h:366
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:2489
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:222
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:671
int i
Definition: input.c:407