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

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