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"
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 "filters.h"
34 #include "cuda/load_helper.h"
35 
36 static const enum AVPixelFormat supported_formats[] = {
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 
51 
52  enum AVPixelFormat in_fmt, out_fmt;
57 
58  uint8_t chromakey_rgba[4];
59  uint16_t chromakey_uv[2];
60  int is_yuv;
61  float similarity;
62  float blend;
63 
67 
68  CUcontext cu_ctx;
69  CUmodule cu_module;
70  CUfunction cu_func;
71  CUfunction cu_func_uv;
72  CUstream cu_stream;
74 
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 
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 
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 
182 {
183  FilterLink *inl = ff_filter_link(ctx->inputs[0]);
184  FilterLink *outl = ff_filter_link(ctx->outputs[0]);
185  ChromakeyCUDAContext *s = ctx->priv;
186  AVHWFramesContext *in_frames_ctx;
187  int ret;
188 
189  /* check that we have a hw context */
190  if (!inl->hw_frames_ctx) {
191  av_log(ctx, AV_LOG_ERROR, "No hw context provided on input\n");
192  return AVERROR(EINVAL);
193  }
194  in_frames_ctx = (AVHWFramesContext *)inl->hw_frames_ctx->data;
195 
196  if (!format_is_supported(in_frames_ctx->sw_format))
197  {
198  av_log(ctx, AV_LOG_ERROR, "Unsupported format: %s\n", av_get_pix_fmt_name(in_frames_ctx->sw_format));
199  return AVERROR(ENOSYS);
200  }
201 
203 
204  ret = init_hwframe_ctx(s, in_frames_ctx->device_ref, width, height);
205  if (ret < 0)
206  return ret;
207 
208  outl->hw_frames_ctx = av_buffer_ref(s->frames_ctx);
209  if (!outl->hw_frames_ctx)
210  return AVERROR(ENOMEM);
211 
212  return 0;
213 }
214 
216 {
217  ChromakeyCUDAContext *s = ctx->priv;
218  CUcontext context, cuda_ctx = s->hwctx->cuda_ctx;
219  CudaFunctions *cu = s->hwctx->internal->cuda_dl;
220  int ret;
221 
222  extern const unsigned char ff_vf_chromakey_cuda_ptx_data[];
223  extern const unsigned int ff_vf_chromakey_cuda_ptx_len;
224 
225  ret = CHECK_CU(cu->cuCtxPushCurrent(cuda_ctx));
226  if (ret < 0)
227  return ret;
228 
229  ret = ff_cuda_load_module(ctx, s->hwctx, &s->cu_module,
230  ff_vf_chromakey_cuda_ptx_data, ff_vf_chromakey_cuda_ptx_len);
231  if (ret < 0)
232  goto fail;
233 
234  ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_func, s->cu_module, "Process_uchar"));
235  if (ret < 0)
236  {
237  av_log(ctx, AV_LOG_FATAL, "Failed loading Process_uchar\n");
238  goto fail;
239  }
240 
241  ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_uv, s->cu_module, "Process_uchar2"));
242  if (ret < 0)
243  {
244  av_log(ctx, AV_LOG_FATAL, "Failed loading Process_uchar2\n");
245  goto fail;
246  }
247 
248 fail:
249  CHECK_CU(cu->cuCtxPopCurrent(&context));
250 
251  return ret;
252 }
253 
254 #define FIXNUM(x) lrint((x) * (1 << 10))
255 #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)
256 #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)
257 
259 {
260  AVFilterContext *ctx = outlink->src;
261  AVFilterLink *inlink = outlink->src->inputs[0];
263  ChromakeyCUDAContext *s = ctx->priv;
264  AVHWFramesContext *frames_ctx;
265  AVCUDADeviceContext *device_hwctx;
266  int ret;
267 
268  if (s->is_yuv) {
269  s->chromakey_uv[0] = s->chromakey_rgba[1];
270  s->chromakey_uv[1] = s->chromakey_rgba[2];
271  } else {
272  s->chromakey_uv[0] = RGB_TO_U(s->chromakey_rgba);
273  s->chromakey_uv[1] = RGB_TO_V(s->chromakey_rgba);
274  }
275 
277  if (ret < 0)
278  return ret;
279 
280  frames_ctx = (AVHWFramesContext *)inl->hw_frames_ctx->data;
281  device_hwctx = frames_ctx->device_ctx->hwctx;
282 
283  s->hwctx = device_hwctx;
284  s->cu_stream = s->hwctx->stream;
285 
286  outlink->sample_aspect_ratio = inlink->sample_aspect_ratio;
287 
289  if (ret < 0)
290  return ret;
291 
292  return 0;
293 }
294 
295 static int call_cuda_kernel(AVFilterContext *ctx, CUfunction func,
296  CUtexObject src_tex[3], AVFrame *out_frame,
297  int width, int height, int pitch,
298  int width_uv, int height_uv, int pitch_uv,
299  float u_key, float v_key, float similarity,
300  float blend)
301 {
302  ChromakeyCUDAContext *s = ctx->priv;
303  CudaFunctions *cu = s->hwctx->internal->cuda_dl;
304 
305  CUdeviceptr dst_devptr[4] = {
306  (CUdeviceptr)out_frame->data[0], (CUdeviceptr)out_frame->data[1],
307  (CUdeviceptr)out_frame->data[2], (CUdeviceptr)out_frame->data[3]
308  };
309 
310  void *args_uchar[] = {
311  &src_tex[0], &src_tex[1], &src_tex[2],
312  &dst_devptr[0], &dst_devptr[1], &dst_devptr[2], &dst_devptr[3],
313  &width, &height, &pitch,
314  &width_uv, &height_uv, &pitch_uv,
315  &u_key, &v_key, &similarity, &blend
316  };
317 
318  return CHECK_CU(cu->cuLaunchKernel(func,
320  BLOCKX, BLOCKY, 1, 0, s->cu_stream, args_uchar, NULL));
321 }
322 
324  AVFrame *out, AVFrame *in)
325 {
326  ChromakeyCUDAContext *s = ctx->priv;
327  CudaFunctions *cu = s->hwctx->internal->cuda_dl;
328  CUcontext context, cuda_ctx = s->hwctx->cuda_ctx;
329  int i, ret;
330 
331  CUtexObject tex[3] = {0, 0, 0};
332 
333  ret = CHECK_CU(cu->cuCtxPushCurrent(cuda_ctx));
334  if (ret < 0)
335  return ret;
336 
337  for (i = 0; i < s->in_planes; i++)
338  {
339  CUDA_TEXTURE_DESC tex_desc = {
340  .filterMode = CU_TR_FILTER_MODE_LINEAR,
341  .flags = 0, // CU_TRSF_READ_AS_INTEGER to get raw ints instead of normalized floats from tex2D
342  };
343 
344  CUDA_RESOURCE_DESC res_desc = {
345  .resType = CU_RESOURCE_TYPE_PITCH2D,
346  .res.pitch2D.format = CU_AD_FORMAT_UNSIGNED_INT8,
347  .res.pitch2D.numChannels = s->in_plane_channels[i],
348  .res.pitch2D.pitchInBytes = in->linesize[i],
349  .res.pitch2D.devPtr = (CUdeviceptr)in->data[i],
350  };
351 
352  if (i == 1 || i == 2)
353  {
354  res_desc.res.pitch2D.width = AV_CEIL_RSHIFT(in->width, s->in_desc->log2_chroma_w);
355  res_desc.res.pitch2D.height = AV_CEIL_RSHIFT(in->height, s->in_desc->log2_chroma_h);
356  }
357  else
358  {
359  res_desc.res.pitch2D.width = in->width;
360  res_desc.res.pitch2D.height = in->height;
361  }
362 
363  ret = CHECK_CU(cu->cuTexObjectCreate(&tex[i], &res_desc, &tex_desc, NULL));
364  if (ret < 0)
365  goto exit;
366  }
367 
368  ret = call_cuda_kernel(ctx, (s->in_plane_channels[1] > 1) ? s->cu_func_uv : s->cu_func,
369  tex, out,
370  out->width, out->height, out->linesize[0],
371  AV_CEIL_RSHIFT(out->width, s->out_desc->log2_chroma_w),
372  AV_CEIL_RSHIFT(out->height, s->out_desc->log2_chroma_h),
373  out->linesize[1],
374  s->chromakey_uv[0], s->chromakey_uv[1],
375  s->similarity, s->blend);
376  if (ret < 0)
377  goto exit;
378 
379 exit:
380  for (i = 0; i < s->in_planes; i++)
381  if (tex[i])
382  CHECK_CU(cu->cuTexObjectDestroy(tex[i]));
383 
384  CHECK_CU(cu->cuCtxPopCurrent(&context));
385 
386  return ret;
387 }
388 
390 {
391  ChromakeyCUDAContext *s = ctx->priv;
392  AVFrame *src = in;
393  int ret;
394 
396  if (ret < 0)
397  return ret;
398 
399  src = s->frame;
400  ret = av_hwframe_get_buffer(src->hw_frames_ctx, s->tmp_frame, 0);
401  if (ret < 0)
402  return ret;
403 
404  av_frame_move_ref(out, s->frame);
405  av_frame_move_ref(s->frame, s->tmp_frame);
406 
407  ret = av_frame_copy_props(out, in);
408  if (ret < 0)
409  return ret;
410 
411  return 0;
412 }
413 
415 {
416  AVFilterContext *ctx = link->dst;
417  ChromakeyCUDAContext *s = ctx->priv;
418  AVFilterLink *outlink = ctx->outputs[0];
419  CudaFunctions *cu = s->hwctx->internal->cuda_dl;
420 
421  AVFrame *out = NULL;
422  CUcontext context;
423  int ret = 0;
424 
425  out = av_frame_alloc();
426  if (!out)
427  {
428  ret = AVERROR(ENOMEM);
429  goto fail;
430  }
431 
432  ret = CHECK_CU(cu->cuCtxPushCurrent(s->hwctx->cuda_ctx));
433  if (ret < 0)
434  goto fail;
435 
437 
438  CHECK_CU(cu->cuCtxPopCurrent(&context));
439  if (ret < 0)
440  goto fail;
441 
442  av_frame_free(&in);
443  return ff_filter_frame(outlink, out);
444 fail:
445  av_frame_free(&in);
446  av_frame_free(&out);
447  return ret;
448 }
449 
450 #define OFFSET(x) offsetof(ChromakeyCUDAContext, x)
451 #define FLAGS (AV_OPT_FLAG_FILTERING_PARAM | AV_OPT_FLAG_VIDEO_PARAM)
452 static const AVOption options[] = {
453  {"color", "set the chromakey key color", OFFSET(chromakey_rgba), AV_OPT_TYPE_COLOR, {.str = "black"}, 0, 0, FLAGS},
454  {"similarity", "set the chromakey similarity value", OFFSET(similarity), AV_OPT_TYPE_FLOAT, {.dbl = 0.01}, 0.01, 1.0, FLAGS},
455  {"blend", "set the chromakey key blend value", OFFSET(blend), AV_OPT_TYPE_FLOAT, {.dbl = 0.0}, 0.0, 1.0, FLAGS},
456  {"yuv", "color parameter is in yuv instead of rgb", OFFSET(is_yuv), AV_OPT_TYPE_BOOL, {.i64 = 0}, 0, 1, FLAGS},
457  {NULL},
458 };
459 
460 static const AVClass cudachromakey_class = {
461  .class_name = "cudachromakey",
462  .item_name = av_default_item_name,
463  .option = options,
464  .version = LIBAVUTIL_VERSION_INT,
465 };
466 
468  {
469  .name = "default",
470  .type = AVMEDIA_TYPE_VIDEO,
471  .filter_frame = cudachromakey_filter_frame,
472  },
473 };
474 
476  {
477  .name = "default",
478  .type = AVMEDIA_TYPE_VIDEO,
479  .config_props = cudachromakey_config_props,
480  },
481 };
482 
484  .name = "chromakey_cuda",
485  .description = NULL_IF_CONFIG_SMALL("GPU accelerated chromakey filter"),
486 
487  .init = cudachromakey_init,
488  .uninit = cudachromakey_uninit,
489 
490  .priv_size = sizeof(ChromakeyCUDAContext),
491  .priv_class = &cudachromakey_class,
492 
495 
497 
498  .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE,
499 };
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:295
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:55
ChromakeyCUDAContext::in_desc
const AVPixFmtDescriptor * in_desc
Definition: vf_chromakey_cuda.c:53
ff_filter_frame
int ff_filter_frame(AVFilterLink *link, AVFrame *frame)
Send a frame of data to the next filter.
Definition: avfilter.c:1062
av_pix_fmt_desc_get
const AVPixFmtDescriptor * av_pix_fmt_desc_get(enum AVPixelFormat pix_fmt)
Definition: pixdesc.c:2965
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:35
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:162
av_hwframe_ctx_init
int av_hwframe_ctx_init(AVBufferRef *ref)
Finalize the context before use.
Definition: hwcontext.c:322
FILTER_INPUTS
#define FILTER_INPUTS(array)
Definition: filters.h:262
AVFrame
This structure describes decoded (raw) audio or video data.
Definition: frame.h:389
pixdesc.h
AVFrame::width
int width
Definition: frame.h:461
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:429
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:205
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:410
av_pix_fmt_count_planes
int av_pix_fmt_count_planes(enum AVPixelFormat pix_fmt)
Definition: pixdesc.c:3005
fail
#define fail()
Definition: checkasm.h:188
cudachromakey_load_functions
static av_cold int cudachromakey_load_functions(AVFilterContext *ctx)
Definition: vf_chromakey_cuda.c:215
ChromakeyCUDAContext::is_yuv
int is_yuv
Definition: vf_chromakey_cuda.c:60
AVFilterPad
A filter pad used for either input or output.
Definition: filters.h:38
av_frame_alloc
AVFrame * av_frame_alloc(void)
Allocate an AVFrame and set its fields to default values.
Definition: frame.c:150
ChromakeyCUDAContext::cu_ctx
CUcontext cu_ctx
Definition: vf_chromakey_cuda.c:68
FLAGS
#define FLAGS
Definition: vf_chromakey_cuda.c:451
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:209
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
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:60
ff_vf_chromakey_cuda
const AVFilter ff_vf_chromakey_cuda
Definition: vf_chromakey_cuda.c:483
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:255
ChromakeyCUDAContext::blend
float blend
Definition: vf_chromakey_cuda.c:62
filters.h
options
static const AVOption options[]
Definition: vf_chromakey_cuda.c:452
ctx
AVFormatContext * ctx
Definition: movenc.c:49
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:450
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_OUTPUTS
#define FILTER_OUTPUTS(array)
Definition: filters.h:263
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:75
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:713
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
Underlying C type is uint8_t[4].
Definition: opt.h:323
AVFilterContext::inputs
AVFilterLink ** inputs
array of pointers to input links
Definition: avfilter.h:465
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
ff_filter_link
static FilterLink * ff_filter_link(AVFilterLink *link)
Definition: filters.h:197
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: filters.h:206
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:94
height
#define height
Definition: dsp.h:85
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
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:256
ChromakeyCUDAContext::chromakey_uv
uint16_t chromakey_uv[2]
Definition: vf_chromakey_cuda.c:59
AV_OPT_TYPE_FLOAT
@ AV_OPT_TYPE_FLOAT
Underlying C type is float.
Definition: opt.h:271
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:256
internal.h
cudachromakey_class
static const AVClass cudachromakey_class
Definition: vf_chromakey_cuda.c:460
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:637
av_frame_unref
void av_frame_unref(AVFrame *frame)
Unreference all the buffers referenced by frame and reset the frame fields.
Definition: frame.c:610
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: filters.h:44
AVFilter
Filter definition.
Definition: avfilter.h:201
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:203
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:80
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:475
AVFrame::height
int height
Definition: frame.h:461
avfilter.h
cudachromakey_filter_frame
static int cudachromakey_filter_frame(AVFilterLink *link, AVFrame *in)
Definition: vf_chromakey_cuda.c:414
cudachromakey_inputs
static const AVFilterPad cudachromakey_inputs[]
Definition: vf_chromakey_cuda.c:467
BLOCKX
#define BLOCKX
Definition: vf_chromakey_cuda.c:43
AVFilterContext
An instance of a filter.
Definition: avfilter.h:457
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:323
cudachromakey_config_props
static av_cold int cudachromakey_config_props(AVFilterLink *outlink)
Definition: vf_chromakey_cuda.c:258
CHECK_CU
#define CHECK_CU(x)
Definition: vf_chromakey_cuda.c:45
AV_OPT_TYPE_BOOL
@ AV_OPT_TYPE_BOOL
Underlying C type is int.
Definition: opt.h:327
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:434
av_log
#define av_log(a,...)
Definition: tableprint_vlc.h:27
width
#define width
Definition: dsp.h:85
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
FILTER_SINGLE_PIXFMT
#define FILTER_SINGLE_PIXFMT(pix_fmt_)
Definition: filters.h:252
cudachromakey_process
static int cudachromakey_process(AVFilterContext *ctx, AVFrame *out, AVFrame *in)
Definition: vf_chromakey_cuda.c:389
src
#define src
Definition: vp8dsp.c:248
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:2885