FFmpeg
 All Data Structures Namespaces Files Functions Variables Typedefs Enumerations Enumerator Macros Groups Pages
vf_yadif_cuda.c
Go to the documentation of this file.
1 /*
2  * Copyright (C) 2018 Philip Langdale <philipl@overt.org>
3  *
4  * This file is part of FFmpeg.
5  *
6  * FFmpeg is free software; you can redistribute it and/or
7  * modify it under the terms of the GNU Lesser General Public
8  * License as published by the Free Software Foundation; either
9  * version 2.1 of the License, or (at your option) any later version.
10  *
11  * FFmpeg is distributed in the hope that it will be useful,
12  * but WITHOUT ANY WARRANTY; without even the implied warranty of
13  * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU
14  * Lesser General Public License for more details.
15  *
16  * You should have received a copy of the GNU Lesser General Public
17  * License along with FFmpeg; if not, write to the Free Software
18  * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA
19  */
20 
21 #include <cuda.h>
22 #include "libavutil/avassert.h"
24 #include "internal.h"
25 #include "yadif.h"
26 
27 extern char vf_yadif_cuda_ptx[];
28 
29 typedef struct DeintCUDAContext {
31 
36 
37  CUcontext cu_ctx;
38  CUstream stream;
39  CUmodule cu_module;
40  CUfunction cu_func_uchar;
41  CUfunction cu_func_uchar2;
42  CUfunction cu_func_ushort;
43  CUfunction cu_func_ushort2;
45 
46 #define DIV_UP(a, b) ( ((a) + (b) - 1) / (b) )
47 #define ALIGN_UP(a, b) (((a) + (b) - 1) & ~((b) - 1))
48 #define BLOCKX 32
49 #define BLOCKY 16
50 
51 static int check_cu(AVFilterContext *avctx, CUresult err, const char *func)
52 {
53  const char *err_name;
54  const char *err_string;
55 
56  av_log(avctx, AV_LOG_TRACE, "Calling %s\n", func);
57 
58  if (err == CUDA_SUCCESS)
59  return 0;
60 
61  cuGetErrorName(err, &err_name);
62  cuGetErrorString(err, &err_string);
63 
64  av_log(avctx, AV_LOG_ERROR, "%s failed", func);
65  if (err_name && err_string)
66  av_log(avctx, AV_LOG_ERROR, " -> %s: %s", err_name, err_string);
67  av_log(avctx, AV_LOG_ERROR, "\n");
68 
69  return AVERROR_EXTERNAL;
70 }
71 
72 #define CHECK_CU(x) check_cu(ctx, (x), #x)
73 
74 static CUresult call_kernel(AVFilterContext *ctx, CUfunction func,
75  CUdeviceptr prev, CUdeviceptr cur, CUdeviceptr next,
76  CUarray_format format, int channels,
77  int src_width, // Width is pixels per channel
78  int src_height, // Height is pixels per channel
79  int src_pitch, // Pitch is bytes
80  CUdeviceptr dst,
81  int dst_width, // Width is pixels per channel
82  int dst_height, // Height is pixels per channel
83  int dst_pitch, // Pitch is pixels per channel
84  int parity, int tff)
85 {
86  DeintCUDAContext *s = ctx->priv;
87  CUtexObject tex_prev = 0, tex_cur = 0, tex_next = 0;
88  CUresult err;
89  int skip_spatial_check = s->yadif.mode&2;
90 
91  void *args[] = { &dst, &tex_prev, &tex_cur, &tex_next,
92  &dst_width, &dst_height, &dst_pitch,
93  &src_width, &src_height, &parity, &tff,
94  &skip_spatial_check };
95 
96  CUDA_TEXTURE_DESC tex_desc = {
97  .filterMode = CU_TR_FILTER_MODE_POINT,
98  .flags = CU_TRSF_READ_AS_INTEGER,
99  };
100 
101  CUDA_RESOURCE_DESC res_desc = {
102  .resType = CU_RESOURCE_TYPE_PITCH2D,
103  .res.pitch2D.format = format,
104  .res.pitch2D.numChannels = channels,
105  .res.pitch2D.width = src_width,
106  .res.pitch2D.height = src_height,
107  .res.pitch2D.pitchInBytes = src_pitch,
108  };
109 
110  res_desc.res.pitch2D.devPtr = (CUdeviceptr)prev;
111  err = CHECK_CU(cuTexObjectCreate(&tex_prev, &res_desc, &tex_desc, NULL));
112  if (err != CUDA_SUCCESS) {
113  goto exit;
114  }
115 
116  res_desc.res.pitch2D.devPtr = (CUdeviceptr)cur;
117  err = CHECK_CU(cuTexObjectCreate(&tex_cur, &res_desc, &tex_desc, NULL));
118  if (err != CUDA_SUCCESS) {
119  goto exit;
120  }
121 
122  res_desc.res.pitch2D.devPtr = (CUdeviceptr)next;
123  err = CHECK_CU(cuTexObjectCreate(&tex_next, &res_desc, &tex_desc, NULL));
124  if (err != CUDA_SUCCESS) {
125  goto exit;
126  }
127 
128  err = CHECK_CU(cuLaunchKernel(func,
129  DIV_UP(dst_width, BLOCKX), DIV_UP(dst_height, BLOCKY), 1,
130  BLOCKX, BLOCKY, 1,
131  0, s->stream, args, NULL));
132 
133 exit:
134  if (tex_prev)
135  CHECK_CU(cuTexObjectDestroy(tex_prev));
136  if (tex_cur)
137  CHECK_CU(cuTexObjectDestroy(tex_cur));
138  if (tex_next)
139  CHECK_CU(cuTexObjectDestroy(tex_next));
140 
141  return err;
142 }
143 
144 static void filter(AVFilterContext *ctx, AVFrame *dst,
145  int parity, int tff)
146 {
147  DeintCUDAContext *s = ctx->priv;
148  YADIFContext *y = &s->yadif;
149  CUcontext dummy;
150  CUresult err;
151  int i;
152 
153  err = CHECK_CU(cuCtxPushCurrent(s->cu_ctx));
154  if (err != CUDA_SUCCESS) {
155  goto exit;
156  }
157 
158  for (i = 0; i < y->csp->nb_components; i++) {
159  CUfunction func;
160  CUarray_format format;
161  int pixel_size, channels;
162  const AVComponentDescriptor *comp = &y->csp->comp[i];
163 
164  if (comp->plane < i) {
165  // We process planes as a whole, so don't reprocess
166  // them for additional components
167  continue;
168  }
169 
170  pixel_size = (comp->depth + comp->shift) / 8;
171  channels = comp->step / pixel_size;
172  if (pixel_size > 2 || channels > 2) {
173  av_log(ctx, AV_LOG_ERROR, "Unsupported pixel format: %s\n", y->csp->name);
174  goto exit;
175  }
176  switch (pixel_size) {
177  case 1:
178  func = channels == 1 ? s->cu_func_uchar : s->cu_func_uchar2;
179  format = CU_AD_FORMAT_UNSIGNED_INT8;
180  break;
181  case 2:
182  func = channels == 1 ? s->cu_func_ushort : s->cu_func_ushort2;
183  format = CU_AD_FORMAT_UNSIGNED_INT16;
184  break;
185  default:
186  av_log(ctx, AV_LOG_ERROR, "Unsupported pixel format: %s\n", y->csp->name);
187  goto exit;
188  }
189  av_log(ctx, AV_LOG_TRACE,
190  "Deinterlacing plane %d: pixel_size: %d channels: %d\n",
191  comp->plane, pixel_size, channels);
192  call_kernel(ctx, func,
193  (CUdeviceptr)y->prev->data[i],
194  (CUdeviceptr)y->cur->data[i],
195  (CUdeviceptr)y->next->data[i],
196  format, channels,
197  AV_CEIL_RSHIFT(y->cur->width, i ? y->csp->log2_chroma_w : 0),
198  AV_CEIL_RSHIFT(y->cur->height, i ? y->csp->log2_chroma_h : 0),
199  y->cur->linesize[i],
200  (CUdeviceptr)dst->data[i],
201  AV_CEIL_RSHIFT(dst->width, i ? y->csp->log2_chroma_w : 0),
202  AV_CEIL_RSHIFT(dst->height, i ? y->csp->log2_chroma_h : 0),
203  dst->linesize[i] / comp->step,
204  parity, tff);
205  }
206 
207  err = CHECK_CU(cuStreamSynchronize(s->stream));
208  if (err != CUDA_SUCCESS) {
209  goto exit;
210  }
211 
212 exit:
213  CHECK_CU(cuCtxPopCurrent(&dummy));
214  return;
215 }
216 
218 {
219  CUcontext dummy;
220  DeintCUDAContext *s = ctx->priv;
221  YADIFContext *y = &s->yadif;
222 
223  if (s->cu_module) {
224  CHECK_CU(cuCtxPushCurrent(s->cu_ctx));
225  CHECK_CU(cuModuleUnload(s->cu_module));
226  CHECK_CU(cuCtxPopCurrent(&dummy));
227  }
228 
229  av_frame_free(&y->prev);
230  av_frame_free(&y->cur);
231  av_frame_free(&y->next);
232 
234  s->hwctx = NULL;
236  s->input_frames = NULL;
237 }
238 
240 {
241  enum AVPixelFormat pix_fmts[] = {
243  };
244  int ret;
245 
246  if ((ret = ff_formats_ref(ff_make_format_list(pix_fmts),
247  &ctx->inputs[0]->out_formats)) < 0)
248  return ret;
249  if ((ret = ff_formats_ref(ff_make_format_list(pix_fmts),
250  &ctx->outputs[0]->in_formats)) < 0)
251  return ret;
252 
253  return 0;
254 }
255 
256 static int config_input(AVFilterLink *inlink)
257 {
258  AVFilterContext *ctx = inlink->dst;
259  DeintCUDAContext *s = ctx->priv;
260 
261  if (!inlink->hw_frames_ctx) {
262  av_log(ctx, AV_LOG_ERROR, "A hardware frames reference is "
263  "required to associate the processing device.\n");
264  return AVERROR(EINVAL);
265  }
266 
268  if (!s->input_frames_ref) {
269  av_log(ctx, AV_LOG_ERROR, "A input frames reference create "
270  "failed.\n");
271  return AVERROR(ENOMEM);
272  }
274 
275  return 0;
276 }
277 
278 static int config_output(AVFilterLink *link)
279 {
280  AVHWFramesContext *output_frames;
281  AVFilterContext *ctx = link->src;
282  DeintCUDAContext *s = ctx->priv;
283  YADIFContext *y = &s->yadif;
284  int ret = 0;
285  CUcontext dummy;
286  CUresult err;
287 
290  if (!s->device_ref) {
291  av_log(ctx, AV_LOG_ERROR, "A device reference create "
292  "failed.\n");
293  return AVERROR(ENOMEM);
294  }
295  s->hwctx = ((AVHWDeviceContext*)s->device_ref->data)->hwctx;
296  s->cu_ctx = s->hwctx->cuda_ctx;
297  s->stream = s->hwctx->stream;
298 
300  if (!link->hw_frames_ctx) {
301  av_log(ctx, AV_LOG_ERROR, "Failed to create HW frame context "
302  "for output.\n");
303  ret = AVERROR(ENOMEM);
304  goto exit;
305  }
306 
307  output_frames = (AVHWFramesContext*)link->hw_frames_ctx->data;
308 
309  output_frames->format = AV_PIX_FMT_CUDA;
310  output_frames->sw_format = s->input_frames->sw_format;
311  output_frames->width = ctx->inputs[0]->w;
312  output_frames->height = ctx->inputs[0]->h;
313 
314  output_frames->initial_pool_size = 4;
315 
316  ret = ff_filter_init_hw_frames(ctx, link, 10);
317  if (ret < 0)
318  goto exit;
319 
320  ret = av_hwframe_ctx_init(link->hw_frames_ctx);
321  if (ret < 0) {
322  av_log(ctx, AV_LOG_ERROR, "Failed to initialise CUDA frame "
323  "context for output: %d\n", ret);
324  goto exit;
325  }
326 
327  link->time_base.num = ctx->inputs[0]->time_base.num;
328  link->time_base.den = ctx->inputs[0]->time_base.den * 2;
329  link->w = ctx->inputs[0]->w;
330  link->h = ctx->inputs[0]->h;
331 
332  if(y->mode & 1)
333  link->frame_rate = av_mul_q(ctx->inputs[0]->frame_rate,
334  (AVRational){2, 1});
335 
336  if (link->w < 3 || link->h < 3) {
337  av_log(ctx, AV_LOG_ERROR, "Video of less than 3 columns or lines is not supported\n");
338  ret = AVERROR(EINVAL);
339  goto exit;
340  }
341 
342  y->csp = av_pix_fmt_desc_get(output_frames->sw_format);
343  y->filter = filter;
344 
345  err = CHECK_CU(cuCtxPushCurrent(s->cu_ctx));
346  if (err != CUDA_SUCCESS) {
347  ret = AVERROR_EXTERNAL;
348  goto exit;
349  }
350 
351  err = CHECK_CU(cuModuleLoadData(&s->cu_module, vf_yadif_cuda_ptx));
352  if (err != CUDA_SUCCESS) {
353  ret = AVERROR_INVALIDDATA;
354  goto exit;
355  }
356 
357  err = CHECK_CU(cuModuleGetFunction(&s->cu_func_uchar, s->cu_module, "yadif_uchar"));
358  if (err != CUDA_SUCCESS) {
359  ret = AVERROR_INVALIDDATA;
360  goto exit;
361  }
362 
363  err = CHECK_CU(cuModuleGetFunction(&s->cu_func_uchar2, s->cu_module, "yadif_uchar2"));
364  if (err != CUDA_SUCCESS) {
365  ret = AVERROR_INVALIDDATA;
366  goto exit;
367  }
368 
369  err= CHECK_CU(cuModuleGetFunction(&s->cu_func_ushort, s->cu_module, "yadif_ushort"));
370  if (err != CUDA_SUCCESS) {
371  ret = AVERROR_INVALIDDATA;
372  goto exit;
373  }
374 
375  err = CHECK_CU(cuModuleGetFunction(&s->cu_func_ushort2, s->cu_module, "yadif_ushort2"));
376  if (err != CUDA_SUCCESS) {
377  ret = AVERROR_INVALIDDATA;
378  goto exit;
379  }
380 
381 exit:
382  CHECK_CU(cuCtxPopCurrent(&dummy));
383 
384  return ret;
385 }
386 
387 static const AVClass yadif_cuda_class = {
388  .class_name = "yadif_cuda",
389  .item_name = av_default_item_name,
390  .option = ff_yadif_options,
391  .version = LIBAVUTIL_VERSION_INT,
392  .category = AV_CLASS_CATEGORY_FILTER,
393 };
394 
395 static const AVFilterPad deint_cuda_inputs[] = {
396  {
397  .name = "default",
398  .type = AVMEDIA_TYPE_VIDEO,
399  .filter_frame = ff_yadif_filter_frame,
400  .config_props = config_input,
401  },
402  { NULL }
403 };
404 
405 static const AVFilterPad deint_cuda_outputs[] = {
406  {
407  .name = "default",
408  .type = AVMEDIA_TYPE_VIDEO,
409  .request_frame = ff_yadif_request_frame,
410  .config_props = config_output,
411  },
412  { NULL }
413 };
414 
416  .name = "yadif_cuda",
417  .description = NULL_IF_CONFIG_SMALL("Deinterlace CUDA frames"),
418  .priv_size = sizeof(DeintCUDAContext),
419  .priv_class = &yadif_cuda_class,
422  .inputs = deint_cuda_inputs,
423  .outputs = deint_cuda_outputs,
425  .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE,
426 };
This struct aggregates all the (hardware/vendor-specific) "high-level" state, i.e.
Definition: hwcontext.h:60
int plane
Which of the 4 planes contains the component.
Definition: pixdesc.h:35
#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
#define AVERROR_INVALIDDATA
Invalid data found when processing input.
Definition: error.h:59
static const char * format[]
Definition: af_aiir.c:330
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
const AVPixFmtDescriptor * av_pix_fmt_desc_get(enum AVPixelFormat pix_fmt)
Definition: pixdesc.c:2446
This structure describes decoded (raw) audio or video data.
Definition: frame.h:226
BYTE int const BYTE int src_pitch
Definition: avisynth_c.h:813
static CUresult call_kernel(AVFilterContext *ctx, CUfunction func, CUdeviceptr prev, CUdeviceptr cur, CUdeviceptr next, CUarray_format format, int channels, int src_width, int src_height, int src_pitch, CUdeviceptr dst, int dst_width, int dst_height, int dst_pitch, int parity, int tff)
Definition: vf_yadif_cuda.c:74
#define LIBAVUTIL_VERSION_INT
Definition: version.h:85
static av_cold void deint_cuda_uninit(AVFilterContext *ctx)
channels
Definition: aptx.c:30
int num
Numerator.
Definition: rational.h:59
int ff_yadif_request_frame(AVFilterLink *link)
Definition: yadif_common.c:154
const char * av_default_item_name(void *ptr)
Return the context name.
Definition: log.c:191
#define CHECK_CU(x)
Definition: vf_yadif_cuda.c:72
int width
The allocated dimensions of the frames in this pool.
Definition: hwcontext.h:228
static int config_input(AVFilterLink *inlink)
enum AVPixelFormat format
The pixel format identifying the underlying HW surface type.
Definition: hwcontext.h:208
static void filter(AVFilterContext *ctx, AVFrame *dst, int parity, int tff)
AVHWFramesContext * input_frames
Definition: vf_yadif_cuda.c:35
uint8_t log2_chroma_w
Amount to shift the luma width right to find the chroma width.
Definition: pixdesc.h:92
AVFilterFormats * ff_make_format_list(const int *fmts)
Create a list of supported formats.
Definition: formats.c:283
CUcontext cu_ctx
Definition: vf_yadif_cuda.c:37
const char * name
Pad name.
Definition: internal.h:60
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
#define av_assert0(cond)
assert() equivalent, that is always enabled.
Definition: avassert.h:37
AVFrame * cur
Definition: yadif.h:53
AVComponentDescriptor comp[4]
Parameters that describe how pixels are packed.
Definition: pixdesc.h:117
#define av_cold
Definition: attributes.h:82
static av_cold int uninit(AVCodecContext *avctx)
Definition: crystalhd.c:279
#define AV_LOG_TRACE
Extremely verbose debugging, useful for libav* development.
Definition: log.h:202
AVFrame * next
Definition: yadif.h:54
AVFrame * prev
Definition: yadif.h:55
#define av_log(a,...)
const char * name
Definition: pixdesc.h:82
An API-specific header for AV_HWDEVICE_TYPE_CUDA.
A filter pad used for either input or output.
Definition: internal.h:54
int width
Definition: frame.h:284
#define AV_LOG_ERROR
Something went wrong and cannot losslessly be recovered.
Definition: log.h:176
uint8_t log2_chroma_h
Amount to shift the luma height right to find the chroma height.
Definition: pixdesc.h:101
#define AVERROR(e)
Definition: error.h:43
void av_frame_free(AVFrame **frame)
Free the frame and any dynamically allocated objects in it, e.g.
Definition: frame.c:202
#define NULL_IF_CONFIG_SMALL(x)
Return NULL if CONFIG_SMALL is true, otherwise the argument without modification. ...
Definition: internal.h:186
#define DIV_UP(a, b)
Definition: vf_yadif_cuda.c:46
void * priv
private data for use by the filter
Definition: avfilter.h:353
#define BLOCKY
Definition: vf_yadif_cuda.c:49
simple assert() macros that are a bit more flexible than ISO C assert().
int av_hwframe_ctx_init(AVBufferRef *ref)
Finalize the context before use.
Definition: hwcontext.c:329
char vf_yadif_cuda_ptx[]
int initial_pool_size
Initial size of the frame pool.
Definition: hwcontext.h:198
uint8_t nb_components
The number of components each pixel has, (1-4)
Definition: pixdesc.h:83
static const AVClass yadif_cuda_class
#define BLOCKX
Definition: vf_yadif_cuda.c:48
int ff_formats_ref(AVFilterFormats *f, AVFilterFormats **ref)
Add *ref as a new reference to formats.
Definition: formats.c:440
AVFormatContext * ctx
Definition: movenc.c:48
CUfunction cu_func_uchar2
Definition: vf_yadif_cuda.c:41
#define s(width, name)
Definition: cbs_vp9.c:257
mcdeint parity
Definition: vf_mcdeint.c:274
int dummy
Definition: motion.c:64
AVCUDADeviceContext * hwctx
Definition: vf_yadif_cuda.c:32
static const AVFilterPad inputs[]
Definition: af_acontrast.c:193
HW acceleration through CUDA.
Definition: pixfmt.h:235
static const AVFilterPad outputs[]
Definition: af_acontrast.c:203
AVBufferRef * device_ref
Definition: vf_yadif_cuda.c:33
static void comp(unsigned char *dst, ptrdiff_t dst_stride, unsigned char *src, ptrdiff_t src_stride, int add)
Definition: eamad.c:83
CUfunction cu_func_ushort2
Definition: vf_yadif_cuda.c:43
CUmodule cu_module
Definition: vf_yadif_cuda.c:39
int linesize[AV_NUM_DATA_POINTERS]
For video, size in bytes of each picture line.
Definition: frame.h:257
AVFilter ff_vf_yadif_cuda
uint8_t * data
The data buffer.
Definition: buffer.h:89
This struct is allocated as AVHWDeviceContext.hwctx.
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
int ff_yadif_filter_frame(AVFilterLink *link, AVFrame *frame)
Definition: yadif_common.c:90
This struct describes a set or pool of "hardware" frames (i.e.
Definition: hwcontext.h:123
int(* func)(AVBPrint *dst, const char *in, const char *arg)
Definition: jacosubdec.c:67
static const AVFilterPad deint_cuda_outputs[]
const char * name
Filter name.
Definition: avfilter.h:148
YADIFContext yadif
Definition: vf_yadif_cuda.c:30
#define AVFILTER_FLAG_SUPPORT_TIMELINE_INTERNAL
Same as AVFILTER_FLAG_SUPPORT_TIMELINE_GENERIC, except that the filter will have its filter_frame() c...
Definition: avfilter.h:133
AVFilterLink ** outputs
array of pointers to output links
Definition: avfilter.h:350
int shift
Number of least significant bits that must be shifted away to get the value.
Definition: pixdesc.h:53
static enum AVPixelFormat pix_fmts[]
Definition: libkvazaar.c:266
CUfunction cu_func_ushort
Definition: vf_yadif_cuda.c:42
#define flags(name, subs,...)
Definition: cbs_av1.c:596
uint8_t * data[AV_NUM_DATA_POINTERS]
pointer to the picture/channel planes.
Definition: frame.h:240
AVBufferRef * device_ref
A reference to the parent AVHWDeviceContext.
Definition: hwcontext.h:140
const AVPixFmtDescriptor * csp
Definition: yadif.h:69
static int config_output(AVFilterLink *link)
AVBufferRef * input_frames_ref
Definition: vf_yadif_cuda.c:34
static const AVFilterPad deint_cuda_inputs[]
A reference to a data buffer.
Definition: buffer.h:81
static int check_cu(AVFilterContext *avctx, CUresult err, const char *func)
Definition: vf_yadif_cuda.c:51
static int query_formats(AVFilterContext *ctx)
Definition: aeval.c:244
AVBufferRef * av_hwframe_ctx_alloc(AVBufferRef *device_ref_in)
Allocate an AVHWFramesContext tied to a given device context.
Definition: hwcontext.c:243
const AVOption ff_yadif_options[]
Definition: yadif_common.c:192
AVBufferRef * av_buffer_ref(AVBufferRef *buf)
Create a new reference to an AVBuffer.
Definition: buffer.c:93
int ff_filter_init_hw_frames(AVFilterContext *avctx, AVFilterLink *link, int default_pool_size)
Perform any additional setup required for hardware frames.
Definition: avfilter.c:1640
int den
Denominator.
Definition: rational.h:60
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:284
static int deint_cuda_query_formats(AVFilterContext *ctx)
BYTE int dst_pitch
Definition: avisynth_c.h:813
internal API functions
int depth
Number of bits in the component.
Definition: pixdesc.h:58
enum AVPixelFormat sw_format
The pixel format identifying the actual data layout of the hardware frames.
Definition: hwcontext.h:221
#define AVERROR_EXTERNAL
Generic error in an external library.
Definition: error.h:57
AVPixelFormat
Pixel format.
Definition: pixfmt.h:64
CUfunction cu_func_uchar
Definition: vf_yadif_cuda.c:40
int mode
YADIFMode.
Definition: yadif.h:47
int step
Number of elements between 2 horizontally consecutive pixels.
Definition: pixdesc.h:41
#define AV_CEIL_RSHIFT(a, b)
Definition: common.h:58