FFmpeg
vf_bilateral_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"
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 
39 #include "cuda/load_helper.h"
40 
41 static const enum AVPixelFormat supported_formats[] = {
45 };
46 
47 #define DIV_UP(a, b) ( ((a) + (b) - 1) / (b) )
48 #define BLOCKX 32
49 #define BLOCKY 16
50 
51 #define CHECK_CU(x) FF_CUDA_CHECK_DL(ctx, s->hwctx->internal->cuda_dl, x)
52 
53 
54 typedef struct CUDABilateralContext {
55  const AVClass *class;
57 
58  enum AVPixelFormat in_fmt, out_fmt;
63 
65  float sigmaS;
66  float sigmaR;
67 
71 
72  CUcontext cu_ctx;
73  CUmodule cu_module;
74  CUfunction cu_func;
75  CUfunction cu_func_uv;
76  CUstream cu_stream;
78 
80 {
81  CUDABilateralContext *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 
95 {
96  CUDABilateralContext *s = ctx->priv;
97 
98  if (s->hwctx && s->cu_module) {
99  CudaFunctions *cu = s->hwctx->internal->cuda_dl;
100  CUcontext bilateral;
101 
102  CHECK_CU(cu->cuCtxPushCurrent(s->hwctx->cuda_ctx));
103  CHECK_CU(cu->cuModuleUnload(s->cu_module));
104  s->cu_module = NULL;
105  CHECK_CU(cu->cuCtxPopCurrent(&bilateral));
106  }
107 
108  av_frame_free(&s->frame);
109  av_buffer_unref(&s->frames_ctx);
110  av_frame_free(&s->tmp_frame);
111 }
112 
114 {
115  AVBufferRef *out_ref = NULL;
116  AVHWFramesContext *out_ctx;
117  int ret;
118 
119  out_ref = av_hwframe_ctx_alloc(device_ctx);
120  if (!out_ref)
121  return AVERROR(ENOMEM);
122  out_ctx = (AVHWFramesContext*)out_ref->data;
123 
124  out_ctx->format = AV_PIX_FMT_CUDA;
125  out_ctx->sw_format = s->out_fmt;
126  out_ctx->width = width;
127  out_ctx->height = height;
128 
129  ret = av_hwframe_ctx_init(out_ref);
130  if (ret < 0)
131  goto fail;
132 
133  av_frame_unref(s->frame);
134  ret = av_hwframe_get_buffer(out_ref, s->frame, 0);
135  if (ret < 0)
136  goto fail;
137 
138  av_buffer_unref(&s->frames_ctx);
139  s->frames_ctx = out_ref;
140 
141  return 0;
142 fail:
143  av_buffer_unref(&out_ref);
144  return ret;
145 }
146 
147 static int format_is_supported(enum AVPixelFormat fmt)
148 {
149  int i;
150 
151  for (i = 0; i < FF_ARRAY_ELEMS(supported_formats); i++)
152  if (supported_formats[i] == fmt)
153  return 1;
154  return 0;
155 }
156 
157 static av_cold void set_format_info(AVFilterContext *ctx, enum AVPixelFormat in_format, enum AVPixelFormat out_format)
158 {
159  CUDABilateralContext *s = ctx->priv;
160  int i, p, d;
161 
162  s->in_fmt = in_format;
163  s->out_fmt = out_format;
164 
165  s->in_desc = av_pix_fmt_desc_get(s->in_fmt);
166  s->out_desc = av_pix_fmt_desc_get(s->out_fmt);
167  s->in_planes = av_pix_fmt_count_planes(s->in_fmt);
168  s->out_planes = av_pix_fmt_count_planes(s->out_fmt);
169 
170  // find maximum step of each component of each plane
171  // For our subset of formats, this should accurately tell us how many channels CUDA needs
172  // i.e. 1 for Y plane, 2 for UV plane of NV12, 4 for single plane of RGB0 formats
173 
174  for (i = 0; i < s->in_desc->nb_components; i++) {
175  d = (s->in_desc->comp[i].depth + 7) / 8;
176  p = s->in_desc->comp[i].plane;
177  s->in_plane_channels[p] = FFMAX(s->in_plane_channels[p], s->in_desc->comp[i].step / d);
178 
179  s->in_plane_depths[p] = s->in_desc->comp[i].depth;
180  }
181 }
182 
184 {
185  CUDABilateralContext *s = ctx->priv;
186  AVHWFramesContext *in_frames_ctx;
187  int ret;
188 
189  /* check that we have a hw context */
190  if (!ctx->inputs[0]->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*)ctx->inputs[0]->hw_frames_ctx->data;
195 
196  if (!format_is_supported(in_frames_ctx->sw_format)) {
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, in_frames_ctx->sw_format);
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 
215 {
216  CUDABilateralContext *s = ctx->priv;
217  CUcontext bilateral, 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_bilateral_cuda_ptx_data[];
222  extern const unsigned int ff_vf_bilateral_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_bilateral_cuda_ptx_data, ff_vf_bilateral_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  av_log(ctx, AV_LOG_FATAL, "Failed loading Process_uchar\n");
236  goto fail;
237  }
238 
239  ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_uv, s->cu_module, "Process_uchar2"));
240  if (ret < 0) {
241  av_log(ctx, AV_LOG_FATAL, "Failed loading Process_uchar2\n");
242  goto fail;
243  }
244 
245 fail:
246  CHECK_CU(cu->cuCtxPopCurrent(&bilateral));
247 
248  return ret;
249 }
250 
252 {
253  AVFilterContext *ctx = outlink->src;
254  AVFilterLink *inlink = outlink->src->inputs[0];
255  CUDABilateralContext *s = ctx->priv;
256  AVHWFramesContext *frames_ctx = (AVHWFramesContext*)inlink->hw_frames_ctx->data;
257  AVCUDADeviceContext *device_hwctx = frames_ctx->device_ctx->hwctx;
258  int ret;
259 
260  s->hwctx = device_hwctx;
261  s->cu_stream = s->hwctx->stream;
262 
264  if (ret < 0)
265  return ret;
266 
267  outlink->sample_aspect_ratio = inlink->sample_aspect_ratio;
268 
269  // the window_size makes more sense when it is odd, so add 1 if even
270  s->window_size= (s->window_size%2) ? s->window_size : s->window_size+1;
271 
273  if (ret < 0)
274  return ret;
275 
276  return 0;
277 }
278 
279 static int call_cuda_kernel(AVFilterContext *ctx, CUfunction func,
280  CUtexObject src_tex[3], AVFrame *out_frame,
281  int width, int height, int pitch,
282  int width_uv, int height_uv, int pitch_uv,
283  int window_size, float sigmaS, float sigmaR)
284 {
285  CUDABilateralContext *s = ctx->priv;
286  CudaFunctions *cu = s->hwctx->internal->cuda_dl;
287  int ret;
288 
289  CUdeviceptr dst_devptr[3] = {
290  (CUdeviceptr)out_frame->data[0], (CUdeviceptr)out_frame->data[1], (CUdeviceptr)out_frame->data[2]
291  };
292 
293  void *args_uchar[] = {
294  &src_tex[0], &src_tex[1], &src_tex[2],
295  &dst_devptr[0], &dst_devptr[1], &dst_devptr[2],
296  &width, &height, &pitch,
297  &width_uv, &height_uv, &pitch_uv,
298  &window_size, &sigmaS, &sigmaR
299  };
300 
301  ret = CHECK_CU(cu->cuLaunchKernel(func,
303  BLOCKX, BLOCKY, 1, 0, s->cu_stream, args_uchar, NULL));
304  if (ret < 0)
305  return ret;
306 
307  return ret;
308 }
309 
311  AVFrame *out, AVFrame *in)
312 {
313  CUDABilateralContext *s = ctx->priv;
314  CudaFunctions *cu = s->hwctx->internal->cuda_dl;
315  CUcontext bilateral, cuda_ctx = s->hwctx->cuda_ctx;
316  int i, ret;
317 
318  CUtexObject tex[3] = { 0, 0, 0 };
319 
320  ret = CHECK_CU(cu->cuCtxPushCurrent(cuda_ctx));
321  if (ret < 0)
322  return ret;
323 
324  for (i = 0; i < s->in_planes; i++) {
325  CUDA_TEXTURE_DESC tex_desc = {
326  .filterMode = CU_TR_FILTER_MODE_LINEAR,
327  .flags = 0, // CU_TRSF_READ_AS_INTEGER to get raw ints instead of normalized floats from tex2D
328  };
329 
330  CUDA_RESOURCE_DESC res_desc = {
331  .resType = CU_RESOURCE_TYPE_PITCH2D,
332  .res.pitch2D.format = CU_AD_FORMAT_UNSIGNED_INT8,
333  .res.pitch2D.numChannels = s->in_plane_channels[i],
334  .res.pitch2D.pitchInBytes = in->linesize[i],
335  .res.pitch2D.devPtr = (CUdeviceptr)in->data[i],
336  };
337 
338  if (i == 1 || i == 2) {
339  res_desc.res.pitch2D.width = AV_CEIL_RSHIFT(in->width, s->in_desc->log2_chroma_w);
340  res_desc.res.pitch2D.height = AV_CEIL_RSHIFT(in->height, s->in_desc->log2_chroma_h);
341  } else {
342  res_desc.res.pitch2D.width = in->width;
343  res_desc.res.pitch2D.height = in->height;
344  }
345 
346  ret = CHECK_CU(cu->cuTexObjectCreate(&tex[i], &res_desc, &tex_desc, NULL));
347  if (ret < 0)
348  goto exit;
349  }
350 
351  ret = call_cuda_kernel(ctx, (s->in_plane_channels[1] > 1) ? s->cu_func_uv : s->cu_func,
352  tex, out,
353  out->width, out->height, out->linesize[0],
354  AV_CEIL_RSHIFT(out->width, s->out_desc->log2_chroma_w),
355  AV_CEIL_RSHIFT(out->height, s->out_desc->log2_chroma_h),
356  out->linesize[1] >> ((s->in_plane_channels[1] > 1) ? 1 : 0),
357  s->window_size, s->sigmaS, s->sigmaR);
358  if (ret < 0)
359  goto exit;
360 
361 exit:
362  for (i = 0; i < s->in_planes; i++)
363  if (tex[i])
364  CHECK_CU(cu->cuTexObjectDestroy(tex[i]));
365 
366  CHECK_CU(cu->cuCtxPopCurrent(&bilateral));
367 
368  return ret;
369 }
370 
372 {
373  CUDABilateralContext *s = ctx->priv;
374  AVFrame *src = in;
375  int ret;
376 
378  if (ret < 0)
379  return ret;
380 
381  src = s->frame;
382  ret = av_hwframe_get_buffer(src->hw_frames_ctx, s->tmp_frame, 0);
383  if (ret < 0)
384  return ret;
385 
386  av_frame_move_ref(out, s->frame);
387  av_frame_move_ref(s->frame, s->tmp_frame);
388 
389  ret = av_frame_copy_props(out, in);
390  if (ret < 0)
391  return ret;
392 
393  return 0;
394 }
395 
397 {
398  AVFilterContext *ctx = link->dst;
399  CUDABilateralContext *s = ctx->priv;
400  AVFilterLink *outlink = ctx->outputs[0];
401  CudaFunctions *cu = s->hwctx->internal->cuda_dl;
402 
403  AVFrame *out = NULL;
404  CUcontext bilateral;
405  int ret = 0;
406 
407  out = av_frame_alloc();
408  if (!out) {
409  ret = AVERROR(ENOMEM);
410  goto fail;
411  }
412 
413  ret = CHECK_CU(cu->cuCtxPushCurrent(s->hwctx->cuda_ctx));
414  if (ret < 0)
415  goto fail;
416 
418 
419  CHECK_CU(cu->cuCtxPopCurrent(&bilateral));
420  if (ret < 0)
421  goto fail;
422 
423  av_frame_free(&in);
424  return ff_filter_frame(outlink, out);
425 fail:
426  av_frame_free(&in);
427  av_frame_free(&out);
428  return ret;
429 }
430 
431 #define OFFSET(x) offsetof(CUDABilateralContext, x)
432 #define FLAGS (AV_OPT_FLAG_FILTERING_PARAM|AV_OPT_FLAG_VIDEO_PARAM)
433 static const AVOption options[] = {
434  { "sigmaS", "set spatial sigma", OFFSET(sigmaS), AV_OPT_TYPE_FLOAT, {.dbl=0.1}, 0.1, 512, FLAGS },
435  { "sigmaR", "set range sigma", OFFSET(sigmaR), AV_OPT_TYPE_FLOAT, {.dbl=0.1}, 0.1, 512, FLAGS },
436  { "window_size", "set neighbours window_size", OFFSET(window_size), AV_OPT_TYPE_INT, {.i64=1}, 1, 255, FLAGS },
437  { NULL }
438 };
439 
441  .class_name = "cudabilateral",
442  .item_name = av_default_item_name,
443  .option = options,
444  .version = LIBAVUTIL_VERSION_INT,
445 };
446 
448  {
449  .name = "default",
450  .type = AVMEDIA_TYPE_VIDEO,
451  .filter_frame = cuda_bilateral_filter_frame,
452  },
453 };
454 
456  {
457  .name = "default",
458  .type = AVMEDIA_TYPE_VIDEO,
459  .config_props = cuda_bilateral_config_props,
460  },
461 };
462 
464  .name = "bilateral_cuda",
465  .description = NULL_IF_CONFIG_SMALL("GPU accelerated bilateral filter"),
466 
467  .init = cudabilateral_init,
468  .uninit = cudabilateral_uninit,
469 
470  .priv_size = sizeof(CUDABilateralContext),
471  .priv_class = &cuda_bilateral_class,
472 
475 
477 
478  .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE,
479 };
cuda_bilateral_process
static int cuda_bilateral_process(AVFilterContext *ctx, AVFrame *out, AVFrame *in)
Definition: vf_bilateral_cuda.c:371
func
int(* func)(AVBPrint *dst, const char *in, const char *arg)
Definition: jacosubdec.c:68
AVHWDeviceContext::hwctx
void * hwctx
The format-specific data, allocated and freed by libavutil along with this context.
Definition: hwcontext.h:92
CUDABilateralContext::in_desc
const AVPixFmtDescriptor * in_desc
Definition: vf_bilateral_cuda.c:59
AV_PIX_FMT_CUDA
@ AV_PIX_FMT_CUDA
HW acceleration through CUDA.
Definition: pixfmt.h:253
AVPixelFormat
AVPixelFormat
Pixel format.
Definition: pixfmt.h:64
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
CUDABilateralContext::frames_ctx
AVBufferRef * frames_ctx
Definition: vf_bilateral_cuda.c:68
opt.h
hwcontext_cuda_internal.h
out
FILE * out
Definition: movenc.c:54
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:370
options
static const AVOption options[]
Definition: vf_bilateral_cuda.c:433
ff_filter_frame
int ff_filter_frame(AVFilterLink *link, AVFrame *frame)
Send a frame of data to the next filter.
Definition: avfilter.c:1009
CUDABilateralContext::hwctx
AVCUDADeviceContext * hwctx
Definition: vf_bilateral_cuda.c:56
av_pix_fmt_desc_get
const AVPixFmtDescriptor * av_pix_fmt_desc_get(enum AVPixelFormat pix_fmt)
Definition: pixdesc.c:2858
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
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, int window_size, float sigmaS, float sigmaR)
Definition: vf_bilateral_cuda.c:279
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
CUDABilateralContext::cu_stream
CUstream cu_stream
Definition: vf_bilateral_cuda.c:76
av_frame_free
void av_frame_free(AVFrame **frame)
Free the frame and any dynamically allocated objects in it, e.g.
Definition: frame.c:116
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:325
pixdesc.h
AVFrame::width
int width
Definition: frame.h:397
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
BLOCKY
#define BLOCKY
Definition: vf_bilateral_cuda.c:49
AVOption
AVOption.
Definition: opt.h:251
cuda_bilateral_filter_frame
static int cuda_bilateral_filter_frame(AVFilterLink *link, AVFrame *in)
Definition: vf_bilateral_cuda.c:396
float.h
init_processing_chain
static av_cold int init_processing_chain(AVFilterContext *ctx, int width, int height)
Definition: vf_bilateral_cuda.c:183
CUDABilateralContext::in_plane_depths
int in_plane_depths[4]
Definition: vf_bilateral_cuda.c:61
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:175
CUDABilateralContext::out_planes
int out_planes
Definition: vf_bilateral_cuda.c:60
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_bilateral_cuda.c:47
cuda_bilateral_load_functions
static av_cold int cuda_bilateral_load_functions(AVFilterContext *ctx)
Definition: vf_bilateral_cuda.c:214
AVFrame::data
uint8_t * data[AV_NUM_DATA_POINTERS]
pointer to the picture/channel planes.
Definition: frame.h:346
formats.h
av_pix_fmt_count_planes
int av_pix_fmt_count_planes(enum AVPixelFormat pix_fmt)
Definition: pixdesc.c:2898
fail
#define fail()
Definition: checkasm.h:134
CUDABilateralContext::frame
AVFrame * frame
Definition: vf_bilateral_cuda.c:69
cuda_bilateral_inputs
static const AVFilterPad cuda_bilateral_inputs[]
Definition: vf_bilateral_cuda.c:447
CUDABilateralContext::tmp_frame
AVFrame * tmp_frame
Definition: vf_bilateral_cuda.c:70
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:104
CUDABilateralContext::out_fmt
enum AVPixelFormat in_fmt out_fmt
Definition: vf_bilateral_cuda.c:58
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
CUDABilateralContext::cu_ctx
CUcontext cu_ctx
Definition: vf_bilateral_cuda.c:72
AVHWFramesContext::height
int height
Definition: hwcontext.h:229
CUDABilateralContext
Definition: vf_bilateral_cuda.c:54
width
#define width
s
#define s(width, name)
Definition: cbs_vp9.c:256
AV_CEIL_RSHIFT
#define AV_CEIL_RSHIFT(a, b)
Definition: common.h:50
ctx
AVFormatContext * ctx
Definition: movenc.c:48
format_is_supported
static int format_is_supported(enum AVPixelFormat fmt)
Definition: vf_bilateral_cuda.c:147
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:190
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
CUDABilateralContext::in_plane_channels
int in_plane_channels[4]
Definition: vf_bilateral_cuda.c:62
if
if(ret)
Definition: filter_design.txt:179
LIBAVUTIL_VERSION_INT
#define LIBAVUTIL_VERSION_INT
Definition: version.h:85
set_format_info
static av_cold void set_format_info(AVFilterContext *ctx, enum AVPixelFormat in_format, enum AVPixelFormat out_format)
Definition: vf_bilateral_cuda.c:157
AVClass
Describe the class of an AVClass context structure.
Definition: log.h:66
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:603
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
FLAGS
#define FLAGS
Definition: vf_bilateral_cuda.c:432
AVHWFramesContext::device_ref
AVBufferRef * device_ref
A reference to the parent AVHWDeviceContext.
Definition: hwcontext.h:141
AVFilterContext::inputs
AVFilterLink ** inputs
array of pointers to input links
Definition: avfilter.h:423
av_default_item_name
const char * av_default_item_name(void *ptr)
Return the context name.
Definition: log.c:237
ff_vf_bilateral_cuda
const AVFilter ff_vf_bilateral_cuda
Definition: vf_bilateral_cuda.c:463
CUDABilateralContext::cu_func_uv
CUfunction cu_func_uv
Definition: vf_bilateral_cuda.c:75
cuda_bilateral_process_internal
static int cuda_bilateral_process_internal(AVFilterContext *ctx, AVFrame *out, AVFrame *in)
Definition: vf_bilateral_cuda.c:310
CUDABilateralContext::window_size
int window_size
Definition: vf_bilateral_cuda.c:64
cudabilateral_init
static av_cold int cudabilateral_init(AVFilterContext *ctx)
Definition: vf_bilateral_cuda.c:79
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
CUDABilateralContext::cu_module
CUmodule cu_module
Definition: vf_bilateral_cuda.c:73
cuda_bilateral_outputs
static const AVFilterPad cuda_bilateral_outputs[]
Definition: vf_bilateral_cuda.c:455
cudabilateral_uninit
static av_cold void cudabilateral_uninit(AVFilterContext *ctx)
Definition: vf_bilateral_cuda.c:94
init_hwframe_ctx
static av_cold int init_hwframe_ctx(CUDABilateralContext *s, AVBufferRef *device_ctx, int width, int height)
Definition: vf_bilateral_cuda.c:113
height
#define height
supported_formats
static enum AVPixelFormat supported_formats[]
Definition: vf_bilateral_cuda.c:41
cuda_bilateral_config_props
static av_cold int cuda_bilateral_config_props(AVFilterLink *outlink)
Definition: vf_bilateral_cuda.c:251
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:180
i
#define i(width, name, range_min, range_max)
Definition: cbs_h2645.c:269
CUDABilateralContext::sigmaS
float sigmaS
Definition: vf_bilateral_cuda.c:65
internal.h
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:516
av_frame_unref
void av_frame_unref(AVFrame *frame)
Unreference all the buffers referenced by frame and reset the frame fields.
Definition: frame.c:487
CUDABilateralContext::sigmaR
float sigmaR
Definition: vf_bilateral_cuda.c:66
AVFilterPad::name
const char * name
Pad name.
Definition: internal.h:55
BLOCKX
#define BLOCKX
Definition: vf_bilateral_cuda.c:48
AVFilter
Filter definition.
Definition: avfilter.h:171
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
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
cuda_bilateral_class
static const AVClass cuda_bilateral_class
Definition: vf_bilateral_cuda.c:440
AVFrame::height
int height
Definition: frame.h:397
AV_OPT_TYPE_INT
@ AV_OPT_TYPE_INT
Definition: opt.h:225
avfilter.h
AV_PIX_FMT_YUV444P
@ AV_PIX_FMT_YUV444P
planar YUV 4:4:4, 24bpp, (1 Cr & Cb sample per 1x1 Y samples)
Definition: pixfmt.h:71
AVFilterContext
An instance of a filter.
Definition: avfilter.h:415
AVMEDIA_TYPE_VIDEO
@ AVMEDIA_TYPE_VIDEO
Definition: avutil.h:201
OFFSET
#define OFFSET(x)
Definition: vf_bilateral_cuda.c:431
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
FILTER_OUTPUTS
#define FILTER_OUTPUTS(array)
Definition: internal.h:191
src
INIT_CLIP pixel * src
Definition: h264pred_template.c:418
CUDABilateralContext::out_desc
const AVPixFmtDescriptor * out_desc
Definition: vf_bilateral_cuda.c:59
d
d
Definition: ffmpeg_filter.c:156
CUDABilateralContext::cu_func
CUfunction cu_func
Definition: vf_bilateral_cuda.c:74
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:370
av_log
#define av_log(a,...)
Definition: tableprint_vlc.h:27
avstring.h
CHECK_CU
#define CHECK_CU(x)
Definition: vf_bilateral_cuda.c:51
CUDABilateralContext::in_planes
int in_planes
Definition: vf_bilateral_cuda.c:60
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
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:2778