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.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 
85  CUcontext cu_ctx;
86  CUmodule cu_module;
87  CUfunction cu_func_uchar;
88  CUfunction cu_func_uchar2;
89  CUfunction cu_func_uchar4;
90  CUfunction cu_func_ushort;
91  CUfunction cu_func_ushort2;
92  CUfunction cu_func_ushort4;
93  CUstream cu_stream;
94 
95  CUdeviceptr srcBuffer;
96  CUdeviceptr dstBuffer;
99 
101 {
102  CUDAScaleContext *s = ctx->priv;
103 
104  s->format = AV_PIX_FMT_NONE;
105  s->frame = av_frame_alloc();
106  if (!s->frame)
107  return AVERROR(ENOMEM);
108 
109  s->tmp_frame = av_frame_alloc();
110  if (!s->tmp_frame)
111  return AVERROR(ENOMEM);
112 
113  return 0;
114 }
115 
117 {
118  CUDAScaleContext *s = ctx->priv;
119 
120  av_frame_free(&s->frame);
123 }
124 
126 {
127  static const enum AVPixelFormat pixel_formats[] = {
129  };
130  AVFilterFormats *pix_fmts = ff_make_format_list(pixel_formats);
131 
132  return ff_set_common_formats(ctx, pix_fmts);
133 }
134 
136 {
137  AVBufferRef *out_ref = NULL;
138  AVHWFramesContext *out_ctx;
139  int in_sw, in_sh, out_sw, out_sh;
140  int ret, i;
141 
142  av_pix_fmt_get_chroma_sub_sample(s->in_fmt, &in_sw, &in_sh);
143  av_pix_fmt_get_chroma_sub_sample(s->out_fmt, &out_sw, &out_sh);
144  if (!s->planes_out[0].width) {
145  s->planes_out[0].width = s->planes_in[0].width;
146  s->planes_out[0].height = s->planes_in[0].height;
147  }
148 
149  for (i = 1; i < FF_ARRAY_ELEMS(s->planes_in); i++) {
150  s->planes_in[i].width = s->planes_in[0].width >> in_sw;
151  s->planes_in[i].height = s->planes_in[0].height >> in_sh;
152  s->planes_out[i].width = s->planes_out[0].width >> out_sw;
153  s->planes_out[i].height = s->planes_out[0].height >> out_sh;
154  }
155 
156  out_ref = av_hwframe_ctx_alloc(device_ctx);
157  if (!out_ref)
158  return AVERROR(ENOMEM);
159  out_ctx = (AVHWFramesContext*)out_ref->data;
160 
161  out_ctx->format = AV_PIX_FMT_CUDA;
162  out_ctx->sw_format = s->out_fmt;
163  out_ctx->width = FFALIGN(s->planes_out[0].width, 32);
164  out_ctx->height = FFALIGN(s->planes_out[0].height, 32);
165 
166  ret = av_hwframe_ctx_init(out_ref);
167  if (ret < 0)
168  goto fail;
169 
170  av_frame_unref(s->frame);
171  ret = av_hwframe_get_buffer(out_ref, s->frame, 0);
172  if (ret < 0)
173  goto fail;
174 
175  s->frame->width = s->planes_out[0].width;
176  s->frame->height = s->planes_out[0].height;
177 
179  s->frames_ctx = out_ref;
180 
181  return 0;
182 fail:
183  av_buffer_unref(&out_ref);
184  return ret;
185 }
186 
188 {
189  int i;
190 
191  for (i = 0; i < FF_ARRAY_ELEMS(supported_formats); i++)
192  if (supported_formats[i] == fmt)
193  return 1;
194  return 0;
195 }
196 
197 static av_cold int init_processing_chain(AVFilterContext *ctx, int in_width, int in_height,
198  int out_width, int out_height)
199 {
200  CUDAScaleContext *s = ctx->priv;
201 
202  AVHWFramesContext *in_frames_ctx;
203 
204  enum AVPixelFormat in_format;
205  enum AVPixelFormat out_format;
206  int ret;
207 
208  /* check that we have a hw context */
209  if (!ctx->inputs[0]->hw_frames_ctx) {
210  av_log(ctx, AV_LOG_ERROR, "No hw context provided on input\n");
211  return AVERROR(EINVAL);
212  }
213  in_frames_ctx = (AVHWFramesContext*)ctx->inputs[0]->hw_frames_ctx->data;
214  in_format = in_frames_ctx->sw_format;
215  out_format = (s->format == AV_PIX_FMT_NONE) ? in_format : s->format;
216 
217  if (!format_is_supported(in_format)) {
218  av_log(ctx, AV_LOG_ERROR, "Unsupported input format: %s\n",
219  av_get_pix_fmt_name(in_format));
220  return AVERROR(ENOSYS);
221  }
222  if (!format_is_supported(out_format)) {
223  av_log(ctx, AV_LOG_ERROR, "Unsupported output format: %s\n",
224  av_get_pix_fmt_name(out_format));
225  return AVERROR(ENOSYS);
226  }
227 
228  if (in_width == out_width && in_height == out_height)
229  s->passthrough = 1;
230 
231  s->in_fmt = in_format;
232  s->out_fmt = out_format;
233 
234  s->planes_in[0].width = in_width;
235  s->planes_in[0].height = in_height;
236  s->planes_out[0].width = out_width;
237  s->planes_out[0].height = out_height;
238 
239  ret = init_stage(s, in_frames_ctx->device_ref);
240  if (ret < 0)
241  return ret;
242 
244  if (!ctx->outputs[0]->hw_frames_ctx)
245  return AVERROR(ENOMEM);
246 
247  return 0;
248 }
249 
251 {
252  AVFilterContext *ctx = outlink->src;
253  AVFilterLink *inlink = outlink->src->inputs[0];
254  CUDAScaleContext *s = ctx->priv;
256  AVCUDADeviceContext *device_hwctx = frames_ctx->device_ctx->hwctx;
257  CUcontext dummy, cuda_ctx = device_hwctx->cuda_ctx;
258  CudaFunctions *cu = device_hwctx->internal->cuda_dl;
259  int w, h;
260  int ret;
261 
262  extern char vf_scale_cuda_ptx[];
263 
264  s->hwctx = device_hwctx;
265  s->cu_stream = s->hwctx->stream;
266 
267  ret = CHECK_CU(cu->cuCtxPushCurrent(cuda_ctx));
268  if (ret < 0)
269  goto fail;
270 
271  ret = CHECK_CU(cu->cuModuleLoadData(&s->cu_module, vf_scale_cuda_ptx));
272  if (ret < 0)
273  goto fail;
274 
275  CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_uchar, s->cu_module, "Subsample_Bilinear_uchar"));
276  if (ret < 0)
277  goto fail;
278 
279  CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_uchar2, s->cu_module, "Subsample_Bilinear_uchar2"));
280  if (ret < 0)
281  goto fail;
282 
283  CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_uchar4, s->cu_module, "Subsample_Bilinear_uchar4"));
284  if (ret < 0)
285  goto fail;
286 
287  CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_ushort, s->cu_module, "Subsample_Bilinear_ushort"));
288  if (ret < 0)
289  goto fail;
290 
291  CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_ushort2, s->cu_module, "Subsample_Bilinear_ushort2"));
292  if (ret < 0)
293  goto fail;
294 
295  CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_ushort4, s->cu_module, "Subsample_Bilinear_ushort4"));
296  if (ret < 0)
297  goto fail;
298 
299 
300  CHECK_CU(cu->cuCtxPopCurrent(&dummy));
301 
302  if ((ret = ff_scale_eval_dimensions(s,
303  s->w_expr, s->h_expr,
304  inlink, outlink,
305  &w, &h)) < 0)
306  goto fail;
307 
308  if (((int64_t)h * inlink->w) > INT_MAX ||
309  ((int64_t)w * inlink->h) > INT_MAX)
310  av_log(ctx, AV_LOG_ERROR, "Rescaled value for width or height is too big.\n");
311 
312  outlink->w = w;
313  outlink->h = h;
314 
315  ret = init_processing_chain(ctx, inlink->w, inlink->h, w, h);
316  if (ret < 0)
317  return ret;
318 
319  av_log(ctx, AV_LOG_VERBOSE, "w:%d h:%d -> w:%d h:%d\n",
320  inlink->w, inlink->h, outlink->w, outlink->h);
321 
322  if (inlink->sample_aspect_ratio.num) {
323  outlink->sample_aspect_ratio = av_mul_q((AVRational){outlink->h*inlink->w,
324  outlink->w*inlink->h},
325  inlink->sample_aspect_ratio);
326  } else {
327  outlink->sample_aspect_ratio = inlink->sample_aspect_ratio;
328  }
329 
330  return 0;
331 
332 fail:
333  return ret;
334 }
335 
336 static int call_resize_kernel(AVFilterContext *ctx, CUfunction func, int channels,
337  uint8_t *src_dptr, int src_width, int src_height, int src_pitch,
338  uint8_t *dst_dptr, int dst_width, int dst_height, int dst_pitch,
339  int pixel_size)
340 {
341  CUDAScaleContext *s = ctx->priv;
342  CudaFunctions *cu = s->hwctx->internal->cuda_dl;
343  CUdeviceptr dst_devptr = (CUdeviceptr)dst_dptr;
344  CUtexObject tex = 0;
345  void *args_uchar[] = { &tex, &dst_devptr, &dst_width, &dst_height, &dst_pitch, &src_width, &src_height };
346  int ret;
347 
348  CUDA_TEXTURE_DESC tex_desc = {
349  .filterMode = CU_TR_FILTER_MODE_LINEAR,
350  .flags = CU_TRSF_READ_AS_INTEGER,
351  };
352 
353  CUDA_RESOURCE_DESC res_desc = {
354  .resType = CU_RESOURCE_TYPE_PITCH2D,
355  .res.pitch2D.format = pixel_size == 1 ?
356  CU_AD_FORMAT_UNSIGNED_INT8 :
357  CU_AD_FORMAT_UNSIGNED_INT16,
358  .res.pitch2D.numChannels = channels,
359  .res.pitch2D.width = src_width,
360  .res.pitch2D.height = src_height,
361  .res.pitch2D.pitchInBytes = src_pitch * pixel_size,
362  .res.pitch2D.devPtr = (CUdeviceptr)src_dptr,
363  };
364 
365  ret = CHECK_CU(cu->cuTexObjectCreate(&tex, &res_desc, &tex_desc, NULL));
366  if (ret < 0)
367  goto exit;
368 
369  ret = CHECK_CU(cu->cuLaunchKernel(func,
370  DIV_UP(dst_width, BLOCKX), DIV_UP(dst_height, BLOCKY), 1,
371  BLOCKX, BLOCKY, 1, 0, s->cu_stream, args_uchar, NULL));
372 
373 exit:
374  if (tex)
375  CHECK_CU(cu->cuTexObjectDestroy(tex));
376 
377  return ret;
378 }
379 
381  AVFrame *out, AVFrame *in)
382 {
383  AVHWFramesContext *in_frames_ctx = (AVHWFramesContext*)in->hw_frames_ctx->data;
384  CUDAScaleContext *s = ctx->priv;
385 
386  switch (in_frames_ctx->sw_format) {
387  case AV_PIX_FMT_YUV420P:
388  call_resize_kernel(ctx, s->cu_func_uchar, 1,
389  in->data[0], in->width, in->height, in->linesize[0],
390  out->data[0], out->width, out->height, out->linesize[0],
391  1);
392  call_resize_kernel(ctx, s->cu_func_uchar, 1,
393  in->data[1], in->width/2, in->height/2, in->linesize[0]/2,
394  out->data[1], out->width/2, out->height/2, out->linesize[0]/2,
395  1);
396  call_resize_kernel(ctx, s->cu_func_uchar, 1,
397  in->data[2], in->width/2, in->height/2, in->linesize[0]/2,
398  out->data[2], out->width/2, out->height/2, out->linesize[0]/2,
399  1);
400  break;
401  case AV_PIX_FMT_YUV444P:
402  call_resize_kernel(ctx, s->cu_func_uchar, 1,
403  in->data[0], in->width, in->height, in->linesize[0],
404  out->data[0], out->width, out->height, out->linesize[0],
405  1);
406  call_resize_kernel(ctx, s->cu_func_uchar, 1,
407  in->data[1], in->width, in->height, in->linesize[0],
408  out->data[1], out->width, out->height, out->linesize[0],
409  1);
410  call_resize_kernel(ctx, s->cu_func_uchar, 1,
411  in->data[2], in->width, in->height, in->linesize[0],
412  out->data[2], out->width, out->height, out->linesize[0],
413  1);
414  break;
416  call_resize_kernel(ctx, s->cu_func_ushort, 1,
417  in->data[0], in->width, in->height, in->linesize[0] / 2,
418  out->data[0], out->width, out->height, out->linesize[0] / 2,
419  2);
420  call_resize_kernel(ctx, s->cu_func_ushort, 1,
421  in->data[1], in->width, in->height, in->linesize[1] / 2,
422  out->data[1], out->width, out->height, out->linesize[1] / 2,
423  2);
424  call_resize_kernel(ctx, s->cu_func_ushort, 1,
425  in->data[2], in->width, in->height, in->linesize[2] / 2,
426  out->data[2], out->width, out->height, out->linesize[2] / 2,
427  2);
428  break;
429  case AV_PIX_FMT_NV12:
430  call_resize_kernel(ctx, s->cu_func_uchar, 1,
431  in->data[0], in->width, in->height, in->linesize[0],
432  out->data[0], out->width, out->height, out->linesize[0],
433  1);
434  call_resize_kernel(ctx, s->cu_func_uchar2, 2,
435  in->data[1], in->width/2, in->height/2, in->linesize[1],
436  out->data[1], out->width/2, out->height/2, out->linesize[1]/2,
437  1);
438  break;
439  case AV_PIX_FMT_P010LE:
440  call_resize_kernel(ctx, s->cu_func_ushort, 1,
441  in->data[0], in->width, in->height, in->linesize[0]/2,
442  out->data[0], out->width, out->height, out->linesize[0]/2,
443  2);
444  call_resize_kernel(ctx, s->cu_func_ushort2, 2,
445  in->data[1], in->width / 2, in->height / 2, in->linesize[1]/2,
446  out->data[1], out->width / 2, out->height / 2, out->linesize[1] / 4,
447  2);
448  break;
449  case AV_PIX_FMT_P016LE:
450  call_resize_kernel(ctx, s->cu_func_ushort, 1,
451  in->data[0], in->width, in->height, in->linesize[0] / 2,
452  out->data[0], out->width, out->height, out->linesize[0] / 2,
453  2);
454  call_resize_kernel(ctx, s->cu_func_ushort2, 2,
455  in->data[1], in->width / 2, in->height / 2, in->linesize[1] / 2,
456  out->data[1], out->width / 2, out->height / 2, out->linesize[1] / 4,
457  2);
458  break;
459  default:
460  return AVERROR_BUG;
461  }
462 
463  return 0;
464 }
465 
467 {
468  CUDAScaleContext *s = ctx->priv;
469  AVFrame *src = in;
470  int ret;
471 
472  ret = scalecuda_resize(ctx, s->frame, src);
473  if (ret < 0)
474  return ret;
475 
476  src = s->frame;
477  ret = av_hwframe_get_buffer(src->hw_frames_ctx, s->tmp_frame, 0);
478  if (ret < 0)
479  return ret;
480 
481  av_frame_move_ref(out, s->frame);
483 
484  s->frame->width = s->planes_out[0].width;
485  s->frame->height = s->planes_out[0].height;
486 
487  ret = av_frame_copy_props(out, in);
488  if (ret < 0)
489  return ret;
490 
491  return 0;
492 }
493 
495 {
496  AVFilterContext *ctx = link->dst;
497  CUDAScaleContext *s = ctx->priv;
498  AVFilterLink *outlink = ctx->outputs[0];
499  CudaFunctions *cu = s->hwctx->internal->cuda_dl;
500 
501  AVFrame *out = NULL;
502  CUcontext dummy;
503  int ret = 0;
504 
505  out = av_frame_alloc();
506  if (!out) {
507  ret = AVERROR(ENOMEM);
508  goto fail;
509  }
510 
511  ret = CHECK_CU(cu->cuCtxPushCurrent(s->hwctx->cuda_ctx));
512  if (ret < 0)
513  goto fail;
514 
515  ret = cudascale_scale(ctx, out, in);
516 
517  CHECK_CU(cu->cuCtxPopCurrent(&dummy));
518  if (ret < 0)
519  goto fail;
520 
522  (int64_t)in->sample_aspect_ratio.num * outlink->h * link->w,
523  (int64_t)in->sample_aspect_ratio.den * outlink->w * link->h,
524  INT_MAX);
525 
526  av_frame_free(&in);
527  return ff_filter_frame(outlink, out);
528 fail:
529  av_frame_free(&in);
530  av_frame_free(&out);
531  return ret;
532 }
533 
534 #define OFFSET(x) offsetof(CUDAScaleContext, x)
535 #define FLAGS (AV_OPT_FLAG_FILTERING_PARAM|AV_OPT_FLAG_VIDEO_PARAM)
536 static const AVOption options[] = {
537  { "w", "Output video width", OFFSET(w_expr), AV_OPT_TYPE_STRING, { .str = "iw" }, .flags = FLAGS },
538  { "h", "Output video height", OFFSET(h_expr), AV_OPT_TYPE_STRING, { .str = "ih" }, .flags = FLAGS },
539  { NULL },
540 };
541 
542 static const AVClass cudascale_class = {
543  .class_name = "cudascale",
544  .item_name = av_default_item_name,
545  .option = options,
546  .version = LIBAVUTIL_VERSION_INT,
547 };
548 
549 static const AVFilterPad cudascale_inputs[] = {
550  {
551  .name = "default",
552  .type = AVMEDIA_TYPE_VIDEO,
553  .filter_frame = cudascale_filter_frame,
554  },
555  { NULL }
556 };
557 
558 static const AVFilterPad cudascale_outputs[] = {
559  {
560  .name = "default",
561  .type = AVMEDIA_TYPE_VIDEO,
562  .config_props = cudascale_config_props,
563  },
564  { NULL }
565 };
566 
568  .name = "scale_cuda",
569  .description = NULL_IF_CONFIG_SMALL("GPU accelerated video resizer"),
570 
571  .init = cudascale_init,
572  .uninit = cudascale_uninit,
573  .query_formats = cudascale_query_formats,
574 
575  .priv_size = sizeof(CUDAScaleContext),
576  .priv_class = &cudascale_class,
577 
578  .inputs = cudascale_inputs,
579  .outputs = cudascale_outputs,
580 
581  .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE,
582 };
#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
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)
Definition: scale.c:106
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:87
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:85
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:95
struct CUDAScaleContext::@247 planes_out[3]
CUfunction cu_func_ushort
Definition: vf_scale_cuda.c:90
#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:93
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:92
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)
#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:86
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:88
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:96
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
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:91
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:89
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