FFmpeg
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 "libavutil/avassert.h"
23 #include "libavutil/cuda_check.h"
24 #include "internal.h"
25 #include "yadif.h"
26 
27 #include "cuda/load_helper.h"
28 
29 extern const unsigned char ff_vf_yadif_cuda_ptx_data[];
30 extern const unsigned int ff_vf_yadif_cuda_ptx_len;
31 
32 typedef struct DeintCUDAContext {
34 
39 
40  CUcontext cu_ctx;
41  CUstream stream;
42  CUmodule cu_module;
43  CUfunction cu_func_uchar;
44  CUfunction cu_func_uchar2;
45  CUfunction cu_func_ushort;
46  CUfunction cu_func_ushort2;
48 
49 #define DIV_UP(a, b) ( ((a) + (b) - 1) / (b) )
50 #define ALIGN_UP(a, b) (((a) + (b) - 1) & ~((b) - 1))
51 #define BLOCKX 32
52 #define BLOCKY 16
53 
54 #define CHECK_CU(x) FF_CUDA_CHECK_DL(ctx, s->hwctx->internal->cuda_dl, x)
55 
56 static CUresult call_kernel(AVFilterContext *ctx, CUfunction func,
57  CUdeviceptr prev, CUdeviceptr cur, CUdeviceptr next,
58  CUarray_format format, int channels,
59  int src_width, // Width is pixels per channel
60  int src_height, // Height is pixels per channel
61  int src_pitch, // Pitch is bytes
62  CUdeviceptr dst,
63  int dst_width, // Width is pixels per channel
64  int dst_height, // Height is pixels per channel
65  int dst_pitch, // Pitch is pixels per channel
66  int parity, int tff)
67 {
68  DeintCUDAContext *s = ctx->priv;
69  CudaFunctions *cu = s->hwctx->internal->cuda_dl;
70  CUtexObject tex_prev = 0, tex_cur = 0, tex_next = 0;
71  int ret;
72  int skip_spatial_check = s->yadif.mode&2;
73 
74  void *args[] = { &dst, &tex_prev, &tex_cur, &tex_next,
75  &dst_width, &dst_height, &dst_pitch,
76  &src_width, &src_height, &parity, &tff,
77  &skip_spatial_check };
78 
79  CUDA_TEXTURE_DESC tex_desc = {
80  .filterMode = CU_TR_FILTER_MODE_POINT,
81  .flags = CU_TRSF_READ_AS_INTEGER,
82  };
83 
84  CUDA_RESOURCE_DESC res_desc = {
85  .resType = CU_RESOURCE_TYPE_PITCH2D,
86  .res.pitch2D.format = format,
87  .res.pitch2D.numChannels = channels,
88  .res.pitch2D.width = src_width,
89  .res.pitch2D.height = src_height,
90  .res.pitch2D.pitchInBytes = src_pitch,
91  };
92 
93  res_desc.res.pitch2D.devPtr = (CUdeviceptr)prev;
94  ret = CHECK_CU(cu->cuTexObjectCreate(&tex_prev, &res_desc, &tex_desc, NULL));
95  if (ret < 0)
96  goto exit;
97 
98  res_desc.res.pitch2D.devPtr = (CUdeviceptr)cur;
99  ret = CHECK_CU(cu->cuTexObjectCreate(&tex_cur, &res_desc, &tex_desc, NULL));
100  if (ret < 0)
101  goto exit;
102 
103  res_desc.res.pitch2D.devPtr = (CUdeviceptr)next;
104  ret = CHECK_CU(cu->cuTexObjectCreate(&tex_next, &res_desc, &tex_desc, NULL));
105  if (ret < 0)
106  goto exit;
107 
108  ret = CHECK_CU(cu->cuLaunchKernel(func,
109  DIV_UP(dst_width, BLOCKX), DIV_UP(dst_height, BLOCKY), 1,
110  BLOCKX, BLOCKY, 1,
111  0, s->stream, args, NULL));
112 
113 exit:
114  if (tex_prev)
115  CHECK_CU(cu->cuTexObjectDestroy(tex_prev));
116  if (tex_cur)
117  CHECK_CU(cu->cuTexObjectDestroy(tex_cur));
118  if (tex_next)
119  CHECK_CU(cu->cuTexObjectDestroy(tex_next));
120 
121  return ret;
122 }
123 
124 static void filter(AVFilterContext *ctx, AVFrame *dst,
125  int parity, int tff)
126 {
127  DeintCUDAContext *s = ctx->priv;
128  YADIFContext *y = &s->yadif;
129  CudaFunctions *cu = s->hwctx->internal->cuda_dl;
130  CUcontext dummy;
131  int i, ret;
132 
133  ret = CHECK_CU(cu->cuCtxPushCurrent(s->cu_ctx));
134  if (ret < 0)
135  return;
136 
137  for (i = 0; i < y->csp->nb_components; i++) {
138  CUfunction func;
139  CUarray_format format;
140  int pixel_size, channels;
141  const AVComponentDescriptor *comp = &y->csp->comp[i];
142 
143  if (comp->plane < i) {
144  // We process planes as a whole, so don't reprocess
145  // them for additional components
146  continue;
147  }
148 
149  pixel_size = (comp->depth + comp->shift) / 8;
150  channels = comp->step / pixel_size;
151  if (pixel_size > 2 || channels > 2) {
152  av_log(ctx, AV_LOG_ERROR, "Unsupported pixel format: %s\n", y->csp->name);
153  goto exit;
154  }
155  switch (pixel_size) {
156  case 1:
157  func = channels == 1 ? s->cu_func_uchar : s->cu_func_uchar2;
158  format = CU_AD_FORMAT_UNSIGNED_INT8;
159  break;
160  case 2:
161  func = channels == 1 ? s->cu_func_ushort : s->cu_func_ushort2;
162  format = CU_AD_FORMAT_UNSIGNED_INT16;
163  break;
164  default:
165  av_log(ctx, AV_LOG_ERROR, "Unsupported pixel format: %s\n", y->csp->name);
166  goto exit;
167  }
169  "Deinterlacing plane %d: pixel_size: %d channels: %d\n",
170  comp->plane, pixel_size, channels);
172  (CUdeviceptr)y->prev->data[i],
173  (CUdeviceptr)y->cur->data[i],
174  (CUdeviceptr)y->next->data[i],
175  format, channels,
176  AV_CEIL_RSHIFT(y->cur->width, i ? y->csp->log2_chroma_w : 0),
177  AV_CEIL_RSHIFT(y->cur->height, i ? y->csp->log2_chroma_h : 0),
178  y->cur->linesize[i],
179  (CUdeviceptr)dst->data[i],
180  AV_CEIL_RSHIFT(dst->width, i ? y->csp->log2_chroma_w : 0),
181  AV_CEIL_RSHIFT(dst->height, i ? y->csp->log2_chroma_h : 0),
182  dst->linesize[i] / comp->step,
183  parity, tff);
184  }
185 
186 exit:
187  CHECK_CU(cu->cuCtxPopCurrent(&dummy));
188  return;
189 }
190 
192 {
193  CUcontext dummy;
194  DeintCUDAContext *s = ctx->priv;
195  YADIFContext *y = &s->yadif;
196 
197  if (s->hwctx && s->cu_module) {
198  CudaFunctions *cu = s->hwctx->internal->cuda_dl;
199  CHECK_CU(cu->cuCtxPushCurrent(s->cu_ctx));
200  CHECK_CU(cu->cuModuleUnload(s->cu_module));
201  CHECK_CU(cu->cuCtxPopCurrent(&dummy));
202  }
203 
204  av_frame_free(&y->prev);
205  av_frame_free(&y->cur);
206  av_frame_free(&y->next);
207 
208  av_buffer_unref(&s->device_ref);
209  s->hwctx = NULL;
210  av_buffer_unref(&s->input_frames_ref);
211  s->input_frames = NULL;
212 }
213 
215 {
216  enum AVPixelFormat pix_fmts[] = {
218  };
219  int ret;
220 
222  &ctx->inputs[0]->outcfg.formats)) < 0)
223  return ret;
225  &ctx->outputs[0]->incfg.formats)) < 0)
226  return ret;
227 
228  return 0;
229 }
230 
232 {
233  AVFilterContext *ctx = inlink->dst;
234  DeintCUDAContext *s = ctx->priv;
235 
236  if (!inlink->hw_frames_ctx) {
237  av_log(ctx, AV_LOG_ERROR, "A hardware frames reference is "
238  "required to associate the processing device.\n");
239  return AVERROR(EINVAL);
240  }
241 
242  s->input_frames_ref = av_buffer_ref(inlink->hw_frames_ctx);
243  if (!s->input_frames_ref) {
244  av_log(ctx, AV_LOG_ERROR, "A input frames reference create "
245  "failed.\n");
246  return AVERROR(ENOMEM);
247  }
248  s->input_frames = (AVHWFramesContext*)s->input_frames_ref->data;
249 
250  return 0;
251 }
252 
254 {
255  AVHWFramesContext *output_frames;
256  AVFilterContext *ctx = link->src;
257  DeintCUDAContext *s = ctx->priv;
258  YADIFContext *y = &s->yadif;
259  CudaFunctions *cu;
260  int ret = 0;
261  CUcontext dummy;
262 
263  av_assert0(s->input_frames);
264  s->device_ref = av_buffer_ref(s->input_frames->device_ref);
265  if (!s->device_ref) {
266  av_log(ctx, AV_LOG_ERROR, "A device reference create "
267  "failed.\n");
268  return AVERROR(ENOMEM);
269  }
270  s->hwctx = ((AVHWDeviceContext*)s->device_ref->data)->hwctx;
271  s->cu_ctx = s->hwctx->cuda_ctx;
272  s->stream = s->hwctx->stream;
273  cu = s->hwctx->internal->cuda_dl;
274 
275  link->hw_frames_ctx = av_hwframe_ctx_alloc(s->device_ref);
276  if (!link->hw_frames_ctx) {
277  av_log(ctx, AV_LOG_ERROR, "Failed to create HW frame context "
278  "for output.\n");
279  ret = AVERROR(ENOMEM);
280  goto exit;
281  }
282 
283  output_frames = (AVHWFramesContext*)link->hw_frames_ctx->data;
284 
285  output_frames->format = AV_PIX_FMT_CUDA;
286  output_frames->sw_format = s->input_frames->sw_format;
287  output_frames->width = ctx->inputs[0]->w;
288  output_frames->height = ctx->inputs[0]->h;
289 
290  output_frames->initial_pool_size = 4;
291 
293  if (ret < 0)
294  goto exit;
295 
297  if (ret < 0) {
298  av_log(ctx, AV_LOG_ERROR, "Failed to initialise CUDA frame "
299  "context for output: %d\n", ret);
300  goto exit;
301  }
302 
303  link->time_base = av_mul_q(ctx->inputs[0]->time_base, (AVRational){1, 2});
304  link->w = ctx->inputs[0]->w;
305  link->h = ctx->inputs[0]->h;
306 
307  if(y->mode & 1)
308  link->frame_rate = av_mul_q(ctx->inputs[0]->frame_rate,
309  (AVRational){2, 1});
310 
311  if (link->w < 3 || link->h < 3) {
312  av_log(ctx, AV_LOG_ERROR, "Video of less than 3 columns or lines is not supported\n");
313  ret = AVERROR(EINVAL);
314  goto exit;
315  }
316 
317  y->csp = av_pix_fmt_desc_get(output_frames->sw_format);
318  y->filter = filter;
319 
320  ret = CHECK_CU(cu->cuCtxPushCurrent(s->cu_ctx));
321  if (ret < 0)
322  goto exit;
323 
325  if (ret < 0)
326  goto exit;
327 
328  ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_uchar, s->cu_module, "yadif_uchar"));
329  if (ret < 0)
330  goto exit;
331 
332  ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_uchar2, s->cu_module, "yadif_uchar2"));
333  if (ret < 0)
334  goto exit;
335 
336  ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_ushort, s->cu_module, "yadif_ushort"));
337  if (ret < 0)
338  goto exit;
339 
340  ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_ushort2, s->cu_module, "yadif_ushort2"));
341  if (ret < 0)
342  goto exit;
343 
344 exit:
345  CHECK_CU(cu->cuCtxPopCurrent(&dummy));
346 
347  return ret;
348 }
349 
350 static const AVClass yadif_cuda_class = {
351  .class_name = "yadif_cuda",
352  .item_name = av_default_item_name,
353  .option = ff_yadif_options,
354  .version = LIBAVUTIL_VERSION_INT,
355  .category = AV_CLASS_CATEGORY_FILTER,
356 };
357 
358 static const AVFilterPad deint_cuda_inputs[] = {
359  {
360  .name = "default",
361  .type = AVMEDIA_TYPE_VIDEO,
362  .filter_frame = ff_yadif_filter_frame,
363  .config_props = config_input,
364  },
365  { NULL }
366 };
367 
368 static const AVFilterPad deint_cuda_outputs[] = {
369  {
370  .name = "default",
371  .type = AVMEDIA_TYPE_VIDEO,
372  .request_frame = ff_yadif_request_frame,
373  .config_props = config_output,
374  },
375  { NULL }
376 };
377 
379  .name = "yadif_cuda",
380  .description = NULL_IF_CONFIG_SMALL("Deinterlace CUDA frames"),
381  .priv_size = sizeof(DeintCUDAContext),
382  .priv_class = &yadif_cuda_class,
388  .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE,
389 };
func
int(* func)(AVBPrint *dst, const char *in, const char *arg)
Definition: jacosubdec.c:68
AV_PIX_FMT_CUDA
@ AV_PIX_FMT_CUDA
HW acceleration through CUDA.
Definition: pixfmt.h:225
AVPixelFormat
AVPixelFormat
Pixel format.
Definition: pixfmt.h:64
DeintCUDAContext::input_frames
AVHWFramesContext * input_frames
Definition: vf_yadif_cuda.c:38
AVERROR
Filter the word “frame” indicates either a video frame or a group of audio as stored in an AVFrame structure Format for each input and each output the list of supported formats For video that means pixel format For audio that means channel sample they are references to shared objects When the negotiation mechanism computes the intersection of the formats supported at each end of a all references to both lists are replaced with a reference to the intersection And when a single format is eventually chosen for a link amongst the remaining all references to the list are updated That means that if a filter requires that its input and output have the same format amongst a supported all it has to do is use a reference to the same list of formats query_formats can leave some formats unset and return AVERROR(EAGAIN) to cause the negotiation mechanism toagain later. That can be used by filters with complex requirements to use the format negotiated on one link to set the formats supported on another. Frame references ownership and permissions
ff_make_format_list
AVFilterFormats * ff_make_format_list(const int *fmts)
Create a list of supported formats.
Definition: formats.c:286
hwcontext_cuda_internal.h
FF_FILTER_FLAG_HWFRAME_AWARE
#define FF_FILTER_FLAG_HWFRAME_AWARE
The filter is aware of hardware frames, and any hardware frame context should not be automatically pr...
Definition: internal.h:339
comp
static void comp(unsigned char *dst, ptrdiff_t dst_stride, unsigned char *src, ptrdiff_t src_stride, int add)
Definition: eamad.c:85
av_pix_fmt_desc_get
const AVPixFmtDescriptor * av_pix_fmt_desc_get(enum AVPixelFormat pix_fmt)
Definition: pixdesc.c:2540
AVBufferRef::data
uint8_t * data
The data buffer.
Definition: buffer.h:90
AVHWFramesContext::format
enum AVPixelFormat format
The pixel format identifying the underlying HW surface type.
Definition: hwcontext.h:209
ff_cuda_load_module
int ff_cuda_load_module(void *avctx, AVCUDADeviceContext *hwctx, CUmodule *cu_module, const unsigned char *data, const unsigned int length)
Loads a CUDA module and applies any decompression, if neccesary.
Definition: load_helper.c:34
inlink
The exact code depends on how similar the blocks are and how related they are to the and needs to apply these operations to the correct inlink or outlink if there are several Macros are available to factor that when no extra processing is inlink
Definition: filter_design.txt:212
YADIFContext::csp
const AVPixFmtDescriptor * csp
Definition: yadif.h:75
av_frame_free
void av_frame_free(AVFrame **frame)
Free the frame and any dynamically allocated objects in it, e.g.
Definition: frame.c:112
YADIFContext::mode
int mode
YADIFMode.
Definition: yadif.h:53
av_hwframe_ctx_init
int av_hwframe_ctx_init(AVBufferRef *ref)
Finalize the context before use.
Definition: hwcontext.c:333
AVFrame
This structure describes decoded (raw) audio or video data.
Definition: frame.h:303
AVFrame::width
int width
Definition: frame.h:361
av_hwframe_ctx_alloc
AVBufferRef * av_hwframe_ctx_alloc(AVBufferRef *device_ref_in)
Allocate an AVHWFramesContext tied to a given device context.
Definition: hwcontext.c:247
AVPixFmtDescriptor::name
const char * name
Definition: pixdesc.h:70
AVFilter::name
const char * name
Filter name.
Definition: avfilter.h:149
AVHWFramesContext::width
int width
The allocated dimensions of the frames in this pool.
Definition: hwcontext.h:229
DIV_UP
#define DIV_UP(a, b)
Definition: vf_yadif_cuda.c:49
AVFrame::data
uint8_t * data[AV_NUM_DATA_POINTERS]
pointer to the picture/channel planes.
Definition: frame.h:317
DeintCUDAContext::cu_func_uchar
CUfunction cu_func_uchar
Definition: vf_yadif_cuda.c:43
config_input
static int config_input(AVFilterLink *inlink)
Definition: vf_yadif_cuda.c:231
DeintCUDAContext::cu_func_uchar2
CUfunction cu_func_uchar2
Definition: vf_yadif_cuda.c:44
AVFilterPad
A filter pad used for either input or output.
Definition: internal.h:54
AVHWDeviceContext
This struct aggregates all the (hardware/vendor-specific) "high-level" state, i.e.
Definition: hwcontext.h:61
filter
static void filter(AVFilterContext *ctx, AVFrame *dst, int parity, int tff)
Definition: vf_yadif_cuda.c:124
avassert.h
AV_LOG_TRACE
#define AV_LOG_TRACE
Extremely verbose debugging, useful for libav* development.
Definition: log.h:206
AV_LOG_ERROR
#define AV_LOG_ERROR
Something went wrong and cannot losslessly be recovered.
Definition: log.h:180
deint_cuda_outputs
static const AVFilterPad deint_cuda_outputs[]
Definition: vf_yadif_cuda.c:368
av_cold
#define av_cold
Definition: attributes.h:90
inputs
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
Definition: filter_design.txt:243
AVHWFramesContext::height
int height
Definition: hwcontext.h:229
DeintCUDAContext::cu_ctx
CUcontext cu_ctx
Definition: vf_yadif_cuda.c:40
s
#define s(width, name)
Definition: cbs_vp9.c:257
DeintCUDAContext::cu_func_ushort2
CUfunction cu_func_ushort2
Definition: vf_yadif_cuda.c:46
AV_CEIL_RSHIFT
#define AV_CEIL_RSHIFT(a, b)
Definition: common.h:58
ff_formats_ref
int ff_formats_ref(AVFilterFormats *f, AVFilterFormats **ref)
Add *ref as a new reference to formats.
Definition: formats.c:459
ff_vf_yadif_cuda_ptx_data
const unsigned char ff_vf_yadif_cuda_ptx_data[]
deint_cuda_inputs
static const AVFilterPad deint_cuda_inputs[]
Definition: vf_yadif_cuda.c:358
BLOCKX
#define BLOCKX
Definition: vf_yadif_cuda.c:51
av_assert0
#define av_assert0(cond)
assert() equivalent, that is always enabled.
Definition: avassert.h:37
outputs
static const AVFilterPad outputs[]
Definition: af_acontrast.c:203
pix_fmts
static enum AVPixelFormat pix_fmts[]
Definition: libkvazaar.c:290
ctx
AVFormatContext * ctx
Definition: movenc.c:48
channels
channels
Definition: aptx.h:33
AVPixFmtDescriptor::log2_chroma_w
uint8_t log2_chroma_w
Amount to shift the luma width right to find the chroma width.
Definition: pixdesc.h:80
load_helper.h
link
Filter the word “frame” indicates either a video frame or a group of audio as stored in an AVFrame structure Format for each input and each output the list of supported formats For video that means pixel format For audio that means channel sample they are references to shared objects When the negotiation mechanism computes the intersection of the formats supported at each end of a link
Definition: filter_design.txt:23
LIBAVUTIL_VERSION_INT
#define LIBAVUTIL_VERSION_INT
Definition: version.h:85
yadif_cuda_class
static const AVClass yadif_cuda_class
Definition: vf_yadif_cuda.c:350
AVClass
Describe the class of an AVClass context structure.
Definition: log.h:66
DeintCUDAContext::cu_module
CUmodule cu_module
Definition: vf_yadif_cuda.c:42
NULL
#define NULL
Definition: coverity.c:32
AVHWFramesContext::sw_format
enum AVPixelFormat sw_format
The pixel format identifying the actual data layout of the hardware frames.
Definition: hwcontext.h:222
av_buffer_unref
void av_buffer_unref(AVBufferRef **buf)
Free a given reference and automatically free the buffer if there are no more references to it.
Definition: buffer.c:125
AVComponentDescriptor
Definition: pixdesc.h:30
AVPixFmtDescriptor::nb_components
uint8_t nb_components
The number of components each pixel has, (1-4)
Definition: pixdesc.h:71
AVRational
Rational number (pair of numerator and denominator).
Definition: rational.h:58
DeintCUDAContext::stream
CUstream stream
Definition: vf_yadif_cuda.c:41
av_default_item_name
const char * av_default_item_name(void *ptr)
Return the context name.
Definition: log.c:235
DeintCUDAContext::yadif
YADIFContext yadif
Definition: vf_yadif_cuda.c:33
deint_cuda_uninit
static av_cold void deint_cuda_uninit(AVFilterContext *ctx)
Definition: vf_yadif_cuda.c:191
CHECK_CU
#define CHECK_CU(x)
Definition: vf_yadif_cuda.c:54
AV_CLASS_CATEGORY_FILTER
@ AV_CLASS_CATEGORY_FILTER
Definition: log.h:36
ff_vf_yadif_cuda_ptx_len
const unsigned int ff_vf_yadif_cuda_ptx_len
ff_yadif_options
const AVOption ff_yadif_options[]
Definition: yadif_common.c:198
yadif.h
config_output
static int config_output(AVFilterLink *link)
Definition: vf_yadif_cuda.c:253
NULL_IF_CONFIG_SMALL
#define NULL_IF_CONFIG_SMALL(x)
Return NULL if CONFIG_SMALL is true, otherwise the argument without modification.
Definition: internal.h:116
YADIFContext::filter
void(* filter)(AVFilterContext *ctx, AVFrame *dstpic, int parity, int tff)
Definition: yadif.h:64
parity
mcdeint parity
Definition: vf_mcdeint.c:277
deint_cuda_query_formats
static int deint_cuda_query_formats(AVFilterContext *ctx)
Definition: vf_yadif_cuda.c:214
DeintCUDAContext
Definition: vf_yadif_cuda.c:32
YADIFContext::prev
AVFrame * prev
Definition: yadif.h:61
DeintCUDAContext::input_frames_ref
AVBufferRef * input_frames_ref
Definition: vf_yadif_cuda.c:37
internal.h
call_kernel
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:56
i
int i
Definition: input.c:406
YADIFContext
Definition: yadif.h:50
AVFilterPad::name
const char * name
Pad name.
Definition: internal.h:60
DeintCUDAContext::hwctx
AVCUDADeviceContext * hwctx
Definition: vf_yadif_cuda.c:35
BLOCKY
#define BLOCKY
Definition: vf_yadif_cuda.c:52
AVFilter
Filter definition.
Definition: avfilter.h:145
AVHWFramesContext
This struct describes a set or pool of "hardware" frames (i.e.
Definition: hwcontext.h:124
AVCUDADeviceContext
This struct is allocated as AVHWDeviceContext.hwctx.
Definition: hwcontext_cuda.h:42
ret
ret
Definition: filter_design.txt:187
AVClass::class_name
const char * class_name
The name of the class; usually it is the same name as the context structure type to which the AVClass...
Definition: log.h:71
cuda_check.h
ff_vf_yadif_cuda
const AVFilter ff_vf_yadif_cuda
Definition: vf_yadif_cuda.c:378
AVFrame::hw_frames_ctx
AVBufferRef * hw_frames_ctx
For hwaccel-format frames, this should be a reference to the AVHWFramesContext describing the frame.
Definition: frame.h:607
AVFrame::height
int height
Definition: frame.h:361
YADIFContext::next
AVFrame * next
Definition: yadif.h:60
AV_PIX_FMT_NONE
@ AV_PIX_FMT_NONE
Definition: pixfmt.h:65
ff_yadif_request_frame
int ff_yadif_request_frame(AVFilterLink *link)
Definition: yadif_common.c:159
dummy
int dummy
Definition: motion.c:65
AVPixFmtDescriptor::comp
AVComponentDescriptor comp[4]
Parameters that describe how pixels are packed.
Definition: pixdesc.h:105
av_mul_q
AVRational av_mul_q(AVRational b, AVRational c)
Multiply two rationals.
Definition: rational.c:80
av_buffer_ref
AVBufferRef * av_buffer_ref(AVBufferRef *buf)
Create a new reference to an AVBuffer.
Definition: buffer.c:93
AVFilterContext
An instance of a filter.
Definition: avfilter.h:333
AVHWFramesContext::initial_pool_size
int initial_pool_size
Initial size of the frame pool.
Definition: hwcontext.h:199
AVMEDIA_TYPE_VIDEO
@ AVMEDIA_TYPE_VIDEO
Definition: avutil.h:201
AVBufferRef
A reference to a data buffer.
Definition: buffer.h:82
DeintCUDAContext::device_ref
AVBufferRef * device_ref
Definition: vf_yadif_cuda.c:36
format
fg outputs[0] format
Definition: ffmpeg_filter.c:175
query_formats
static int query_formats(AVFilterContext *ctx)
Definition: aeval.c:243
AVFILTER_FLAG_SUPPORT_TIMELINE_INTERNAL
#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:134
flags
#define flags(name, subs,...)
Definition: cbs_av1.c:561
AVFrame::linesize
int linesize[AV_NUM_DATA_POINTERS]
For video, size in bytes of each picture line.
Definition: frame.h:334
av_log
#define av_log(a,...)
Definition: tableprint_vlc.h:28
DeintCUDAContext::cu_func_ushort
CUfunction cu_func_ushort
Definition: vf_yadif_cuda.c:45
uninit
static av_cold int uninit(AVCodecContext *avctx)
Definition: crystalhd.c:279
YADIFContext::cur
AVFrame * cur
Definition: yadif.h:59
AVPixFmtDescriptor::log2_chroma_h
uint8_t log2_chroma_h
Amount to shift the luma height right to find the chroma height.
Definition: pixdesc.h:89
ff_yadif_filter_frame
int ff_yadif_filter_frame(AVFilterLink *link, AVFrame *frame)
Definition: yadif_common.c:92
ff_filter_init_hw_frames
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:1535