FFmpeg: libavfilter/vf_chromakey_cuda.c Source File

FFmpeg
vf_chromakey_cuda.c
Go to the documentation of this file.
1 /*
2  * Copyright (c) 2022 Mohamed Khaled <Mohamed_Khaled_Kamal@outlook.com>
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 <float.h>
22 #include <stdio.h>
23 
24 #include "libavutil/common.h"
25 #include "libavutil/hwcontext.h"
26 #include "libavutil/hwcontext_cuda_internal.h"
27 #include "libavutil/cuda_check.h"
28 #include "libavutil/internal.h"
29 #include "libavutil/opt.h"
30 #include "libavutil/pixdesc.h"
31 
32 #include "avfilter.h"
33 #include "internal.h"
34 #include "cuda/load_helper.h"
35 
36  static const enum AVPixelFormat supported_formats[] = {
37  AV_PIX_FMT_YUV420P,
38  AV_PIX_FMT_NV12,
39  AV_PIX_FMT_YUVA420P
40 };
41 
42  #define DIV_UP(a, b) (((a) + (b)-1) / (b))
43  #define BLOCKX 32
44  #define BLOCKY 16
45  #define CHECK_CU(x) FF_CUDA_CHECK_DL(ctx, s->hwctx->internal->cuda_dl, x)
46 
47  typedef struct ChromakeyCUDAContext {
48   const AVClass *class;
49 
50   AVCUDADeviceContext *hwctx;
51 
52   enum AVPixelFormat in_fmt, out_fmt;
53   const AVPixFmtDescriptor *in_desc, *out_desc;
54   int in_planes, out_planes;
55   int in_plane_depths[4];
56   int in_plane_channels[4];
57 
58   uint8_t chromakey_rgba[4];
59   uint16_t chromakey_uv[2];
60   int is_yuv;
61   float similarity;
62   float blend;
63 
64   AVBufferRef *frames_ctx;
65   AVFrame *frame;
66   AVFrame *tmp_frame;
67 
68   CUcontext cu_ctx;
69   CUmodule cu_module;
70   CUfunction cu_func;
71   CUfunction cu_func_uv;
72   CUstream cu_stream;
73 } ChromakeyCUDAContext;
74 
75  static av_cold int cudachromakey_init(AVFilterContext *ctx)
76 {
77  ChromakeyCUDAContext *s = ctx->priv;
78 
79  s->frame = av_frame_alloc();
80  if (!s->frame)
81  return AVERROR(ENOMEM);
82 
83  s->tmp_frame = av_frame_alloc();
84  if (!s->tmp_frame)
85  return AVERROR(ENOMEM);
86 
87  return 0;
88 }
89 
90  static av_cold void cudachromakey_uninit(AVFilterContext *ctx)
91 {
92  ChromakeyCUDAContext *s = ctx->priv;
93 
94  if (s->hwctx && s->cu_module)
95  {
96  CudaFunctions *cu = s->hwctx->internal->cuda_dl;
97  CUcontext context;
98 
99  CHECK_CU(cu->cuCtxPushCurrent(s->hwctx->cuda_ctx));
100  CHECK_CU(cu->cuModuleUnload(s->cu_module));
101  s->cu_module = NULL;
102  CHECK_CU(cu->cuCtxPopCurrent(&context));
103  }
104 
105  av_frame_free(&s->frame);
106  av_buffer_unref(&s->frames_ctx);
107  av_frame_free(&s->tmp_frame);
108 }
109 
110  static av_cold int init_hwframe_ctx(ChromakeyCUDAContext *s, AVBufferRef *device_ctx, int width, int height)
111 {
112  AVBufferRef *out_ref = NULL;
113  AVHWFramesContext *out_ctx;
114  int ret;
115 
116  out_ref = av_hwframe_ctx_alloc(device_ctx);
117  if (!out_ref)
118  return AVERROR(ENOMEM);
119  out_ctx = (AVHWFramesContext *)out_ref->data;
120 
121  out_ctx->format = AV_PIX_FMT_CUDA;
122  out_ctx->sw_format = s->out_fmt;
123  out_ctx->width = width;
124  out_ctx->height = height;
125 
126  ret = av_hwframe_ctx_init(out_ref);
127  if (ret < 0)
128  goto fail;
129 
130  av_frame_unref(s->frame);
131  ret = av_hwframe_get_buffer(out_ref, s->frame, 0);
132  if (ret < 0)
133  goto fail;
134 
135  av_buffer_unref(&s->frames_ctx);
136  s->frames_ctx = out_ref;
137 
138  return 0;
139 fail:
140  av_buffer_unref(&out_ref);
141  return ret;
142 }
143 
144  static int format_is_supported(enum AVPixelFormat fmt)
145 {
146  int i;
147 
148  for (i = 0; i < FF_ARRAY_ELEMS(supported_formats); i++)
149  if (supported_formats[i] == fmt)
150  return 1;
151  return 0;
152 }
153 
154  static av_cold void set_format_info(AVFilterContext *ctx, enum AVPixelFormat in_format, enum AVPixelFormat out_format)
155 {
156  ChromakeyCUDAContext *s = ctx->priv;
157  int i, p, d;
158 
159  s->in_fmt = in_format;
160  s->out_fmt = out_format;
161 
162  s->in_desc = av_pix_fmt_desc_get(s->in_fmt);
163  s->out_desc = av_pix_fmt_desc_get(s->out_fmt);
164  s->in_planes = av_pix_fmt_count_planes(s->in_fmt);
165  s->out_planes = av_pix_fmt_count_planes(s->out_fmt);
166 
167  // find maximum step of each component of each plane
168  // For our subset of formats, this should accurately tell us how many channels CUDA needs
169  // i.e. 1 for Y plane, 2 for UV plane of NV12, 4 for single plane of RGB0 formats
170 
171  for (i = 0; i < s->in_desc->nb_components; i++)
172  {
173  d = (s->in_desc->comp[i].depth + 7) / 8;
174  p = s->in_desc->comp[i].plane;
175  s->in_plane_channels[p] = FFMAX(s->in_plane_channels[p], s->in_desc->comp[i].step / d);
176 
177  s->in_plane_depths[p] = s->in_desc->comp[i].depth;
178  }
179 }
180 
181  static av_cold int init_processing_chain(AVFilterContext *ctx, int width, int height)
182 {
183  ChromakeyCUDAContext *s = ctx->priv;
184  AVHWFramesContext *in_frames_ctx;
185  int ret;
186 
187  /* check that we have a hw context */
188  if (!ctx->inputs[0]->hw_frames_ctx)
189  {
190  av_log(ctx, AV_LOG_ERROR, "No hw context provided on input\n");
191  return AVERROR(EINVAL);
192  }
193  in_frames_ctx = (AVHWFramesContext *)ctx->inputs[0]->hw_frames_ctx->data;
194 
195  if (!format_is_supported(in_frames_ctx->sw_format))
196  {
197  av_log(ctx, AV_LOG_ERROR, "Unsupported format: %s\n", av_get_pix_fmt_name(in_frames_ctx->sw_format));
198  return AVERROR(ENOSYS);
199  }
200 
201  set_format_info(ctx, in_frames_ctx->sw_format, AV_PIX_FMT_YUVA420P);
202 
203  ret = init_hwframe_ctx(s, in_frames_ctx->device_ref, width, height);
204  if (ret < 0)
205  return ret;
206 
207  ctx->outputs[0]->hw_frames_ctx = av_buffer_ref(s->frames_ctx);
208  if (!ctx->outputs[0]->hw_frames_ctx)
209  return AVERROR(ENOMEM);
210 
211  return 0;
212 }
213 
214  static av_cold int cudachromakey_load_functions(AVFilterContext *ctx)
215 {
216  ChromakeyCUDAContext *s = ctx->priv;
217  CUcontext context, cuda_ctx = s->hwctx->cuda_ctx;
218  CudaFunctions *cu = s->hwctx->internal->cuda_dl;
219  int ret;
220 
221  extern const unsigned char ff_vf_chromakey_cuda_ptx_data[];
222  extern const unsigned int ff_vf_chromakey_cuda_ptx_len;
223 
224  ret = CHECK_CU(cu->cuCtxPushCurrent(cuda_ctx));
225  if (ret < 0)
226  return ret;
227 
228  ret = ff_cuda_load_module(ctx, s->hwctx, &s->cu_module,
229  ff_vf_chromakey_cuda_ptx_data, ff_vf_chromakey_cuda_ptx_len);
230  if (ret < 0)
231  goto fail;
232 
233  ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_func, s->cu_module, "Process_uchar"));
234  if (ret < 0)
235  {
236  av_log(ctx, AV_LOG_FATAL, "Failed loading Process_uchar\n");
237  goto fail;
238  }
239 
240  ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_uv, s->cu_module, "Process_uchar2"));
241  if (ret < 0)
242  {
243  av_log(ctx, AV_LOG_FATAL, "Failed loading Process_uchar2\n");
244  goto fail;
245  }
246 
247 fail:
248  CHECK_CU(cu->cuCtxPopCurrent(&context));
249 
250  return ret;
251 }
252 
253  #define FIXNUM(x) lrint((x) * (1 << 10))
254  #define RGB_TO_U(rgb) (((-FIXNUM(0.16874) * rgb[0] - FIXNUM(0.33126) * rgb[1] + FIXNUM(0.50000) * rgb[2] + (1 << 9) - 1) >> 10) + 128)
255  #define RGB_TO_V(rgb) (((FIXNUM(0.50000) * rgb[0] - FIXNUM(0.41869) * rgb[1] - FIXNUM(0.08131) * rgb[2] + (1 << 9) - 1) >> 10) + 128)
256 
257  static av_cold int cudachromakey_config_props(AVFilterLink *outlink)
258 {
259  AVFilterContext *ctx = outlink->src;
260  AVFilterLink *inlink = outlink->src->inputs[0];
261  ChromakeyCUDAContext *s = ctx->priv;
262  AVHWFramesContext *frames_ctx = (AVHWFramesContext *)inlink->hw_frames_ctx->data;
263  AVCUDADeviceContext *device_hwctx = frames_ctx->device_ctx->hwctx;
264  int ret;
265 
266  s->hwctx = device_hwctx;
267  s->cu_stream = s->hwctx->stream;
268 
269  if (s->is_yuv) {
270  s->chromakey_uv[0] = s->chromakey_rgba[1];
271  s->chromakey_uv[1] = s->chromakey_rgba[2];
272  } else {
273  s->chromakey_uv[0] = RGB_TO_U(s->chromakey_rgba);
274  s->chromakey_uv[1] = RGB_TO_V(s->chromakey_rgba);
275  }
276 
277  ret = init_processing_chain(ctx, inlink->w, inlink->h);
278  if (ret < 0)
279  return ret;
280 
281  outlink->sample_aspect_ratio = inlink->sample_aspect_ratio;
282 
283  ret = cudachromakey_load_functions(ctx);
284  if (ret < 0)
285  return ret;
286 
287  return 0;
288 }
289 
290  static int call_cuda_kernel(AVFilterContext *ctx, CUfunction func,
291  CUtexObject src_tex[3], AVFrame *out_frame,
292  int width, int height, int pitch,
293  int width_uv, int height_uv, int pitch_uv,
294  float u_key, float v_key, float similarity,
295  float blend)
296 {
297  ChromakeyCUDAContext *s = ctx->priv;
298  CudaFunctions *cu = s->hwctx->internal->cuda_dl;
299 
300  CUdeviceptr dst_devptr[4] = {
301  (CUdeviceptr)out_frame->data[0], (CUdeviceptr)out_frame->data[1],
302  (CUdeviceptr)out_frame->data[2], (CUdeviceptr)out_frame->data[3]
303  };
304 
305  void *args_uchar[] = {
306  &src_tex[0], &src_tex[1], &src_tex[2],
307  &dst_devptr[0], &dst_devptr[1], &dst_devptr[2], &dst_devptr[3],
308  &width, &height, &pitch,
309  &width_uv, &height_uv, &pitch_uv,
310  &u_key, &v_key, &similarity, &blend
311  };
312 
313  return CHECK_CU(cu->cuLaunchKernel(func,
314  DIV_UP(width, BLOCKX), DIV_UP(height, BLOCKY), 1,
315  BLOCKX, BLOCKY, 1, 0, s->cu_stream, args_uchar, NULL));
316 }
317 
318  static int cudachromakey_process_internal(AVFilterContext *ctx,
319  AVFrame *out, AVFrame *in)
320 {
321  ChromakeyCUDAContext *s = ctx->priv;
322  CudaFunctions *cu = s->hwctx->internal->cuda_dl;
323  CUcontext context, cuda_ctx = s->hwctx->cuda_ctx;
324  int i, ret;
325 
326  CUtexObject tex[3] = {0, 0, 0};
327 
328  ret = CHECK_CU(cu->cuCtxPushCurrent(cuda_ctx));
329  if (ret < 0)
330  return ret;
331 
332  for (i = 0; i < s->in_planes; i++)
333  {
334  CUDA_TEXTURE_DESC tex_desc = {
335  .filterMode = CU_TR_FILTER_MODE_LINEAR,
336  .flags = 0, // CU_TRSF_READ_AS_INTEGER to get raw ints instead of normalized floats from tex2D
337  };
338 
339  CUDA_RESOURCE_DESC res_desc = {
340  .resType = CU_RESOURCE_TYPE_PITCH2D,
341  .res.pitch2D.format = CU_AD_FORMAT_UNSIGNED_INT8,
342  .res.pitch2D.numChannels = s->in_plane_channels[i],
343  .res.pitch2D.pitchInBytes = in->linesize[i],
344  .res.pitch2D.devPtr = (CUdeviceptr)in->data[i],
345  };
346 
347  if (i == 1 || i == 2)
348  {
349  res_desc.res.pitch2D.width = AV_CEIL_RSHIFT(in->width, s->in_desc->log2_chroma_w);
350  res_desc.res.pitch2D.height = AV_CEIL_RSHIFT(in->height, s->in_desc->log2_chroma_h);
351  }
352  else
353  {
354  res_desc.res.pitch2D.width = in->width;
355  res_desc.res.pitch2D.height = in->height;
356  }
357 
358  ret = CHECK_CU(cu->cuTexObjectCreate(&tex[i], &res_desc, &tex_desc, NULL));
359  if (ret < 0)
360  goto exit;
361  }
362 
363  ret = call_cuda_kernel(ctx, (s->in_plane_channels[1] > 1) ? s->cu_func_uv : s->cu_func,
364  tex, out,
365  out->width, out->height, out->linesize[0],
366  AV_CEIL_RSHIFT(out->width, s->out_desc->log2_chroma_w),
367  AV_CEIL_RSHIFT(out->height, s->out_desc->log2_chroma_h),
368  out->linesize[1],
369  s->chromakey_uv[0], s->chromakey_uv[1],
370  s->similarity, s->blend);
371  if (ret < 0)
372  goto exit;
373 
374 exit:
375  for (i = 0; i < s->in_planes; i++)
376  if (tex[i])
377  CHECK_CU(cu->cuTexObjectDestroy(tex[i]));
378 
379  CHECK_CU(cu->cuCtxPopCurrent(&context));
380 
381  return ret;
382 }
383 
384  static int cudachromakey_process(AVFilterContext *ctx, AVFrame *out, AVFrame *in)
385 {
386  ChromakeyCUDAContext *s = ctx->priv;
387  AVFrame *src = in;
388  int ret;
389 
390  ret = cudachromakey_process_internal(ctx, s->frame, src);
391  if (ret < 0)
392  return ret;
393 
394  src = s->frame;
395  ret = av_hwframe_get_buffer(src->hw_frames_ctx, s->tmp_frame, 0);
396  if (ret < 0)
397  return ret;
398 
399  av_frame_move_ref(out, s->frame);
400  av_frame_move_ref(s->frame, s->tmp_frame);
401 
402  ret = av_frame_copy_props(out, in);
403  if (ret < 0)
404  return ret;
405 
406  return 0;
407 }
408 
409  static int cudachromakey_filter_frame(AVFilterLink *link, AVFrame *in)
410 {
411  AVFilterContext *ctx = link->dst;
412  ChromakeyCUDAContext *s = ctx->priv;
413  AVFilterLink *outlink = ctx->outputs[0];
414  CudaFunctions *cu = s->hwctx->internal->cuda_dl;
415 
416  AVFrame *out = NULL;
417  CUcontext context;
418  int ret = 0;
419 
420  out = av_frame_alloc();
421  if (!out)
422  {
423  ret = AVERROR(ENOMEM);
424  goto fail;
425  }
426 
427  ret = CHECK_CU(cu->cuCtxPushCurrent(s->hwctx->cuda_ctx));
428  if (ret < 0)
429  goto fail;
430 
431  ret = cudachromakey_process(ctx, out, in);
432 
433  CHECK_CU(cu->cuCtxPopCurrent(&context));
434  if (ret < 0)
435  goto fail;
436 
437  av_frame_free(&in);
438  return ff_filter_frame(outlink, out);
439 fail:
440  av_frame_free(&in);
441  av_frame_free(&out);
442  return ret;
443 }
444 
445  #define OFFSET(x) offsetof(ChromakeyCUDAContext, x)
446  #define FLAGS (AV_OPT_FLAG_FILTERING_PARAM | AV_OPT_FLAG_VIDEO_PARAM)
447  static const AVOption options[] = {
448  {"color", "set the chromakey key color", OFFSET(chromakey_rgba), AV_OPT_TYPE_COLOR, {.str = "black"}, 0, 0, FLAGS},
449  {"similarity", "set the chromakey similarity value", OFFSET(similarity), AV_OPT_TYPE_FLOAT, {.dbl = 0.01}, 0.01, 1.0, FLAGS},
450  {"blend", "set the chromakey key blend value", OFFSET(blend), AV_OPT_TYPE_FLOAT, {.dbl = 0.0}, 0.0, 1.0, FLAGS},
451  {"yuv", "color parameter is in yuv instead of rgb", OFFSET(is_yuv), AV_OPT_TYPE_BOOL, {.i64 = 0}, 0, 1, FLAGS},
452  {NULL},
453 };
454 
455  static const AVClass cudachromakey_class = {
456  .class_name = "cudachromakey",
457  .item_name = av_default_item_name,
458  .option = options,
459  .version = LIBAVUTIL_VERSION_INT,
460 };
461 
462  static const AVFilterPad cudachromakey_inputs[] = {
463  {
464  .name = "default",
465  .type = AVMEDIA_TYPE_VIDEO,
466  .filter_frame = cudachromakey_filter_frame,
467  },
468 };
469 
470  static const AVFilterPad cudachromakey_outputs[] = {
471  {
472  .name = "default",
473  .type = AVMEDIA_TYPE_VIDEO,
474  .config_props = cudachromakey_config_props,
475  },
476 };
477 
478  const AVFilter ff_vf_chromakey_cuda = {
479  .name = "chromakey_cuda",
480  .description = NULL_IF_CONFIG_SMALL("GPU accelerated chromakey filter"),
481 
482  .init = cudachromakey_init,
483  .uninit = cudachromakey_uninit,
484 
485  .priv_size = sizeof(ChromakeyCUDAContext),
486  .priv_class = &cudachromakey_class,
487 
488  FILTER_INPUTS(cudachromakey_inputs),
489  FILTER_OUTPUTS(cudachromakey_outputs),
490 
491  FILTER_SINGLE_PIXFMT(AV_PIX_FMT_CUDA),
492 
493  .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE,
494 };
func
int(* func)(AVBPrint *dst, const char *in, const char *arg)
Definition: jacosubdec.c:68
ChromakeyCUDAContext
Definition: vf_chromakey_cuda.c:47
AVHWDeviceContext::hwctx
void * hwctx
The format-specific data, allocated and freed by libavutil along with this context.
Definition: hwcontext.h:85
call_cuda_kernel
static int call_cuda_kernel(AVFilterContext *ctx, CUfunction func, CUtexObject src_tex[3], AVFrame *out_frame, int width, int height, int pitch, int width_uv, int height_uv, int pitch_uv, float u_key, float v_key, float similarity, float blend)
Definition: vf_chromakey_cuda.c:290
AV_PIX_FMT_CUDA
@ AV_PIX_FMT_CUDA
HW acceleration through CUDA.
Definition: pixfmt.h:260
AVPixelFormat
AVPixelFormat
Pixel format.
Definition: pixfmt.h:71
supported_formats
static enum AVPixelFormat supported_formats[]
Definition: vf_chromakey_cuda.c:36
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
opt.h
hwcontext_cuda_internal.h
out
FILE * out
Definition: movenc.c:54
ChromakeyCUDAContext::in_desc
const AVPixFmtDescriptor * in_desc
Definition: vf_chromakey_cuda.c:53
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:351
ff_filter_frame
int ff_filter_frame(AVFilterLink *link, AVFrame *frame)
Send a frame of data to the next filter.
Definition: avfilter.c:1018
av_pix_fmt_desc_get
const AVPixFmtDescriptor * av_pix_fmt_desc_get(enum AVPixelFormat pix_fmt)
Definition: pixdesc.c:2962
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:197
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 necessary.
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
av_frame_free
void av_frame_free(AVFrame **frame)
Free the frame and any dynamically allocated objects in it, e.g.
Definition: frame.c:130
av_hwframe_ctx_init
int av_hwframe_ctx_init(AVBufferRef *ref)
Finalize the context before use.
Definition: hwcontext.c:322
AVFrame
This structure describes decoded (raw) audio or video data.
Definition: frame.h:344
pixdesc.h
AVFrame::width
int width
Definition: frame.h:416
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:248
AVOption
AVOption.
Definition: opt.h:346
ChromakeyCUDAContext::cu_func_uv
CUfunction cu_func_uv
Definition: vf_chromakey_cuda.c:71
ChromakeyCUDAContext::chromakey_rgba
uint8_t chromakey_rgba[4]
Definition: vf_chromakey_cuda.c:58
float.h
FFMAX
#define FFMAX(a, b)
Definition: macros.h:47
av_buffer_ref
AVBufferRef * av_buffer_ref(const AVBufferRef *buf)
Create a new reference to an AVBuffer.
Definition: buffer.c:103
AVFilter::name
const char * name
Filter name.
Definition: avfilter.h:170
AVHWFramesContext::width
int width
The allocated dimensions of the frames in this pool.
Definition: hwcontext.h:217
DIV_UP
#define DIV_UP(a, b)
Definition: vf_chromakey_cuda.c:42
ChromakeyCUDAContext::tmp_frame
AVFrame * tmp_frame
Definition: vf_chromakey_cuda.c:66
AVFrame::data
uint8_t * data[AV_NUM_DATA_POINTERS]
pointer to the picture/channel planes.
Definition: frame.h:365
av_pix_fmt_count_planes
int av_pix_fmt_count_planes(enum AVPixelFormat pix_fmt)
Definition: pixdesc.c:3002
fail
#define fail()
Definition: checkasm.h:179
cudachromakey_load_functions
static av_cold int cudachromakey_load_functions(AVFilterContext *ctx)
Definition: vf_chromakey_cuda.c:214
ChromakeyCUDAContext::is_yuv
int is_yuv
Definition: vf_chromakey_cuda.c:60
AVFilterPad
A filter pad used for either input or output.
Definition: internal.h:33
av_frame_alloc
AVFrame * av_frame_alloc(void)
Allocate an AVFrame and set its fields to default values.
Definition: frame.c:118
ChromakeyCUDAContext::cu_ctx
CUcontext cu_ctx
Definition: vf_chromakey_cuda.c:68
FLAGS
#define FLAGS
Definition: vf_chromakey_cuda.c:446
BLOCKY
#define BLOCKY
Definition: vf_chromakey_cuda.c:44
AV_LOG_ERROR
#define AV_LOG_ERROR
Something went wrong and cannot losslessly be recovered.
Definition: log.h:180
FF_ARRAY_ELEMS
#define FF_ARRAY_ELEMS(a)
Definition: sinewin_tablegen.c:29
av_cold
#define av_cold
Definition: attributes.h:90
AVHWFramesContext::height
int height
Definition: hwcontext.h:217
ChromakeyCUDAContext::frames_ctx
AVBufferRef * frames_ctx
Definition: vf_chromakey_cuda.c:64
width
#define width
s
#define s(width, name)
Definition: cbs_vp9.c:198
AV_PIX_FMT_YUVA420P
@ AV_PIX_FMT_YUVA420P
planar YUV 4:2:0, 20bpp, (1 Cr & Cb sample per 2x2 Y & A samples)
Definition: pixfmt.h:108
AV_CEIL_RSHIFT
#define AV_CEIL_RSHIFT(a, b)
Definition: common.h:58
ff_vf_chromakey_cuda
const AVFilter ff_vf_chromakey_cuda
Definition: vf_chromakey_cuda.c:478
ChromakeyCUDAContext::cu_stream
CUstream cu_stream
Definition: vf_chromakey_cuda.c:72
ChromakeyCUDAContext::in_plane_depths
int in_plane_depths[4]
Definition: vf_chromakey_cuda.c:55
RGB_TO_U
#define RGB_TO_U(rgb)
Definition: vf_chromakey_cuda.c:254
ChromakeyCUDAContext::blend
float blend
Definition: vf_chromakey_cuda.c:62
options
static const AVOption options[]
Definition: vf_chromakey_cuda.c:447
ctx
AVFormatContext * ctx
Definition: movenc.c:48
cudachromakey_uninit
static av_cold void cudachromakey_uninit(AVFilterContext *ctx)
Definition: vf_chromakey_cuda.c:90
OFFSET
#define OFFSET(x)
Definition: vf_chromakey_cuda.c:445
load_helper.h
AV_PIX_FMT_YUV420P
@ AV_PIX_FMT_YUV420P
planar YUV 4:2:0, 12bpp, (1 Cr & Cb sample per 2x2 Y samples)
Definition: pixfmt.h:73
FILTER_INPUTS
#define FILTER_INPUTS(array)
Definition: internal.h:182
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
if
if(ret)
Definition: filter_design.txt:179
context
it s the only field you need to keep assuming you have a context There is some magic you don t need to care about around this just let it vf default minimum maximum flags name is the option keep it simple and lowercase description are in without and describe what they for example set the foo of the bar offset is the offset of the field in your context
Definition: writing_filters.txt:91
LIBAVUTIL_VERSION_INT
#define LIBAVUTIL_VERSION_INT
Definition: version.h:85
AVClass
Describe the class of an AVClass context structure.
Definition: log.h:66
init_processing_chain
static av_cold int init_processing_chain(AVFilterContext *ctx, int width, int height)
Definition: vf_chromakey_cuda.c:181
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:210
av_frame_copy_props
int av_frame_copy_props(AVFrame *dst, const AVFrame *src)
Copy only "metadata" fields from src to dst.
Definition: frame.c:679
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:139
AVHWFramesContext::device_ref
AVBufferRef * device_ref
A reference to the parent AVHWDeviceContext.
Definition: hwcontext.h:126
ChromakeyCUDAContext::cu_func
CUfunction cu_func
Definition: vf_chromakey_cuda.c:70
AV_OPT_TYPE_COLOR
@ AV_OPT_TYPE_COLOR
Definition: opt.h:250
AVFilterContext::inputs
AVFilterLink ** inputs
array of pointers to input links
Definition: avfilter.h:415
av_default_item_name
const char * av_default_item_name(void *ptr)
Return the context name.
Definition: log.c:237
ChromakeyCUDAContext::hwctx
AVCUDADeviceContext * hwctx
Definition: vf_chromakey_cuda.c:50
ChromakeyCUDAContext::out_fmt
enum AVPixelFormat in_fmt out_fmt
Definition: vf_chromakey_cuda.c:52
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:106
ChromakeyCUDAContext::out_desc
const AVPixFmtDescriptor * out_desc
Definition: vf_chromakey_cuda.c:53
init_hwframe_ctx
static av_cold int init_hwframe_ctx(ChromakeyCUDAContext *s, AVBufferRef *device_ctx, int width, int height)
Definition: vf_chromakey_cuda.c:110
cudachromakey_init
static av_cold int cudachromakey_init(AVFilterContext *ctx)
Definition: vf_chromakey_cuda.c:75
ChromakeyCUDAContext::similarity
float similarity
Definition: vf_chromakey_cuda.c:61
height
#define height
ChromakeyCUDAContext::in_plane_channels
int in_plane_channels[4]
Definition: vf_chromakey_cuda.c:56
ChromakeyCUDAContext::out_planes
int out_planes
Definition: vf_chromakey_cuda.c:54
RGB_TO_V
#define RGB_TO_V(rgb)
Definition: vf_chromakey_cuda.c:255
ChromakeyCUDAContext::chromakey_uv
uint16_t chromakey_uv[2]
Definition: vf_chromakey_cuda.c:59
internal.h
AV_OPT_TYPE_FLOAT
@ AV_OPT_TYPE_FLOAT
Definition: opt.h:238
FILTER_SINGLE_PIXFMT
#define FILTER_SINGLE_PIXFMT(pix_fmt_)
Definition: internal.h:172
ChromakeyCUDAContext::cu_module
CUmodule cu_module
Definition: vf_chromakey_cuda.c:69
i
#define i(width, name, range_min, range_max)
Definition: cbs_h2645.c:255
internal.h
cudachromakey_class
static const AVClass cudachromakey_class
Definition: vf_chromakey_cuda.c:455
common.h
av_frame_move_ref
void av_frame_move_ref(AVFrame *dst, AVFrame *src)
Move everything contained in src to dst and reset src.
Definition: frame.c:603
av_frame_unref
void av_frame_unref(AVFrame *frame)
Unreference all the buffers referenced by frame and reset the frame fields.
Definition: frame.c:576
set_format_info
static av_cold void set_format_info(AVFilterContext *ctx, enum AVPixelFormat in_format, enum AVPixelFormat out_format)
Definition: vf_chromakey_cuda.c:154
AVFilterPad::name
const char * name
Pad name.
Definition: internal.h:39
AVFilter
Filter definition.
Definition: avfilter.h:166
AVHWFramesContext
This struct describes a set or pool of "hardware" frames (i.e.
Definition: hwcontext.h:115
AVCUDADeviceContext
This struct is allocated as AVHWDeviceContext.hwctx.
Definition: hwcontext_cuda.h:42
format_is_supported
static int format_is_supported(enum AVPixelFormat fmt)
Definition: vf_chromakey_cuda.c:144
ret
ret
Definition: filter_design.txt:187
AV_LOG_FATAL
#define AV_LOG_FATAL
Something went wrong and recovery is not possible.
Definition: log.h:174
AV_PIX_FMT_NV12
@ AV_PIX_FMT_NV12
planar YUV 4:2:0, 12bpp, 1 plane for Y and 1 plane for the UV components, which are interleaved (firs...
Definition: pixfmt.h:96
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
AVHWFramesContext::device_ctx
AVHWDeviceContext * device_ctx
The parent AVHWDeviceContext.
Definition: hwcontext.h:134
cuda_check.h
cudachromakey_outputs
static const AVFilterPad cudachromakey_outputs[]
Definition: vf_chromakey_cuda.c:470
AVFrame::height
int height
Definition: frame.h:416
avfilter.h
cudachromakey_filter_frame
static int cudachromakey_filter_frame(AVFilterLink *link, AVFrame *in)
Definition: vf_chromakey_cuda.c:409
cudachromakey_inputs
static const AVFilterPad cudachromakey_inputs[]
Definition: vf_chromakey_cuda.c:462
BLOCKX
#define BLOCKX
Definition: vf_chromakey_cuda.c:43
AVFilterContext
An instance of a filter.
Definition: avfilter.h:407
ChromakeyCUDAContext::in_planes
int in_planes
Definition: vf_chromakey_cuda.c:54
AVMEDIA_TYPE_VIDEO
@ AVMEDIA_TYPE_VIDEO
Definition: avutil.h:201
AVBufferRef
A reference to a data buffer.
Definition: buffer.h:82
AVPixFmtDescriptor
Descriptor that unambiguously describes how the bits of a pixel are stored in the up to 4 data planes...
Definition: pixdesc.h:69
ChromakeyCUDAContext::frame
AVFrame * frame
Definition: vf_chromakey_cuda.c:65
cudachromakey_process_internal
static int cudachromakey_process_internal(AVFilterContext *ctx, AVFrame *out, AVFrame *in)
Definition: vf_chromakey_cuda.c:318
cudachromakey_config_props
static av_cold int cudachromakey_config_props(AVFilterLink *outlink)
Definition: vf_chromakey_cuda.c:257
CHECK_CU
#define CHECK_CU(x)
Definition: vf_chromakey_cuda.c:45
AV_OPT_TYPE_BOOL
@ AV_OPT_TYPE_BOOL
Definition: opt.h:251
FILTER_OUTPUTS
#define FILTER_OUTPUTS(array)
Definition: internal.h:183
src
INIT_CLIP pixel * src
Definition: h264pred_template.c:418
d
d
Definition: ffmpeg_filter.c:409
hwcontext.h
AVFrame::linesize
int linesize[AV_NUM_DATA_POINTERS]
For video, a positive or negative value, which is typically indicating the size in bytes of each pict...
Definition: frame.h:389
av_log
#define av_log(a,...)
Definition: tableprint_vlc.h:27
av_hwframe_get_buffer
int av_hwframe_get_buffer(AVBufferRef *hwframe_ref, AVFrame *frame, int flags)
Allocate a new frame attached to the given AVHWFramesContext.
Definition: hwcontext.c:491
cudachromakey_process
static int cudachromakey_process(AVFilterContext *ctx, AVFrame *out, AVFrame *in)
Definition: vf_chromakey_cuda.c:384
av_get_pix_fmt_name
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:2882
Generated on Thu Sep 26 2024 23:15:35 for FFmpeg by   doxygen 1.8.17

AltStyle によって変換されたページ (->オリジナル) /