FFmpeg: libavfilter/vf_bilateral_cuda.c Source File

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 
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 "filters.h"
34 #include "video.h"
35 
36 #include "cuda/load_helper.h"
37 
38  static const enum AVPixelFormat supported_formats[] = {
39  AV_PIX_FMT_YUV420P,
40  AV_PIX_FMT_NV12,
41  AV_PIX_FMT_YUV444P
42 };
43 
44  #define DIV_UP(a, b) ( ((a) + (b) - 1) / (b) )
45  #define BLOCKX 32
46  #define BLOCKY 16
47 
48  #define CHECK_CU(x) FF_CUDA_CHECK_DL(ctx, s->hwctx->internal->cuda_dl, x)
49 
50 
51  typedef struct CUDABilateralContext {
52   const AVClass *class;
53   AVCUDADeviceContext *hwctx;
54 
55   const AVPixFmtDescriptor *in_desc;
56   int in_planes;
57   int in_plane_channels[4];
58 
59   int window_size;
60   float sigmaS;
61   float sigmaR;
62 
63   CUcontext cu_ctx;
64   CUmodule cu_module;
65   CUfunction cu_func;
66   CUfunction cu_func_uv;
67   CUstream cu_stream;
68 } CUDABilateralContext;
69 
70  static av_cold void cudabilateral_uninit(AVFilterContext *ctx)
71 {
72  CUDABilateralContext *s = ctx->priv;
73 
74  if (s->hwctx && s->cu_module) {
75  CudaFunctions *cu = s->hwctx->internal->cuda_dl;
76  CUcontext bilateral;
77 
78  CHECK_CU(cu->cuCtxPushCurrent(s->hwctx->cuda_ctx));
79  CHECK_CU(cu->cuModuleUnload(s->cu_module));
80  s->cu_module = NULL;
81  CHECK_CU(cu->cuCtxPopCurrent(&bilateral));
82  }
83 }
84 
85  static int format_is_supported(enum AVPixelFormat fmt)
86 {
87  int i;
88 
89  for (i = 0; i < FF_ARRAY_ELEMS(supported_formats); i++)
90  if (supported_formats[i] == fmt)
91  return 1;
92  return 0;
93 }
94 
95  static av_cold void set_format_info(AVFilterContext *ctx, enum AVPixelFormat format)
96 {
97  CUDABilateralContext *s = ctx->priv;
98  int i, p, d;
99 
100  s->in_desc = av_pix_fmt_desc_get(format);
101  s->in_planes = av_pix_fmt_count_planes(format);
102 
103  // find maximum step of each component of each plane
104  // For our subset of formats, this should accurately tell us how many channels CUDA needs
105  // i.e. 1 for Y plane, 2 for UV plane of NV12, 4 for single plane of RGB0 formats
106 
107  for (i = 0; i < s->in_desc->nb_components; i++) {
108  d = (s->in_desc->comp[i].depth + 7) / 8;
109  p = s->in_desc->comp[i].plane;
110  s->in_plane_channels[p] = FFMAX(s->in_plane_channels[p], s->in_desc->comp[i].step / d);
111  }
112 }
113 
114  static av_cold int init_processing_chain(AVFilterContext *ctx, int width, int height)
115 {
116  FilterLink *inl = ff_filter_link(ctx->inputs[0]);
117  FilterLink *outl = ff_filter_link(ctx->outputs[0]);
118  AVHWFramesContext *in_frames_ctx;
119 
120  /* check that we have a hw context */
121  if (!inl->hw_frames_ctx) {
122  av_log(ctx, AV_LOG_ERROR, "No hw context provided on input\n");
123  return AVERROR(EINVAL);
124  }
125  in_frames_ctx = (AVHWFramesContext*)inl->hw_frames_ctx->data;
126 
127  if (!format_is_supported(in_frames_ctx->sw_format)) {
128  av_log(ctx, AV_LOG_ERROR, "Unsupported format: %s\n", av_get_pix_fmt_name(in_frames_ctx->sw_format));
129  return AVERROR(ENOSYS);
130  }
131 
132  set_format_info(ctx, in_frames_ctx->sw_format);
133 
134  outl->hw_frames_ctx = av_buffer_ref(inl->hw_frames_ctx);
135  if (!outl->hw_frames_ctx)
136  return AVERROR(ENOMEM);
137 
138  return 0;
139 }
140 
141  static av_cold int cuda_bilateral_load_functions(AVFilterContext *ctx)
142 {
143  CUDABilateralContext *s = ctx->priv;
144  CUcontext bilateral, cuda_ctx = s->hwctx->cuda_ctx;
145  CudaFunctions *cu = s->hwctx->internal->cuda_dl;
146  int ret;
147 
148  extern const unsigned char ff_vf_bilateral_cuda_ptx_data[];
149  extern const unsigned int ff_vf_bilateral_cuda_ptx_len;
150 
151  ret = CHECK_CU(cu->cuCtxPushCurrent(cuda_ctx));
152  if (ret < 0)
153  return ret;
154 
155  ret = ff_cuda_load_module(ctx, s->hwctx, &s->cu_module,
156  ff_vf_bilateral_cuda_ptx_data, ff_vf_bilateral_cuda_ptx_len);
157  if (ret < 0)
158  goto fail;
159 
160  ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_func, s->cu_module, "Process_uchar"));
161  if (ret < 0) {
162  av_log(ctx, AV_LOG_FATAL, "Failed loading Process_uchar\n");
163  goto fail;
164  }
165 
166  ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_uv, s->cu_module, "Process_uchar2"));
167  if (ret < 0) {
168  av_log(ctx, AV_LOG_FATAL, "Failed loading Process_uchar2\n");
169  goto fail;
170  }
171 
172 fail:
173  CHECK_CU(cu->cuCtxPopCurrent(&bilateral));
174 
175  return ret;
176 }
177 
178  static av_cold int cuda_bilateral_config_props(AVFilterLink *outlink)
179 {
180  AVFilterContext *ctx = outlink->src;
181  AVFilterLink *inlink = outlink->src->inputs[0];
182  FilterLink *inl = ff_filter_link(inlink);
183  CUDABilateralContext *s = ctx->priv;
184  AVHWFramesContext *frames_ctx;
185  AVCUDADeviceContext *device_hwctx;
186  int ret;
187 
188  ret = init_processing_chain(ctx, inlink->w, inlink->h);
189  if (ret < 0)
190  return ret;
191 
192  frames_ctx = (AVHWFramesContext*)inl->hw_frames_ctx->data;
193  device_hwctx = frames_ctx->device_ctx->hwctx;
194 
195  s->hwctx = device_hwctx;
196  s->cu_stream = s->hwctx->stream;
197 
198  outlink->sample_aspect_ratio = inlink->sample_aspect_ratio;
199 
200  // the window_size makes more sense when it is odd, so add 1 if even
201  s->window_size |= 1;
202 
203  ret = cuda_bilateral_load_functions(ctx);
204  if (ret < 0)
205  return ret;
206 
207  return 0;
208 }
209 
210  static int call_cuda_kernel(AVFilterContext *ctx, CUfunction func,
211  CUtexObject src_tex[3], AVFrame *out_frame,
212  int width, int height, int pitch,
213  int width_uv, int height_uv, int pitch_uv,
214  int window_size, float sigmaS, float sigmaR)
215 {
216  CUDABilateralContext *s = ctx->priv;
217  CudaFunctions *cu = s->hwctx->internal->cuda_dl;
218  int ret;
219 
220  CUdeviceptr dst_devptr[3] = {
221  (CUdeviceptr)out_frame->data[0], (CUdeviceptr)out_frame->data[1], (CUdeviceptr)out_frame->data[2]
222  };
223 
224  void *args_uchar[] = {
225  &src_tex[0], &src_tex[1], &src_tex[2],
226  &dst_devptr[0], &dst_devptr[1], &dst_devptr[2],
227  &width, &height, &pitch,
228  &width_uv, &height_uv, &pitch_uv,
229  &window_size, &sigmaS, &sigmaR
230  };
231 
232  ret = CHECK_CU(cu->cuLaunchKernel(func,
233  DIV_UP(width, BLOCKX), DIV_UP(height, BLOCKY), 1,
234  BLOCKX, BLOCKY, 1, 0, s->cu_stream, args_uchar, NULL));
235  if (ret < 0)
236  return ret;
237 
238  return ret;
239 }
240 
241  static int cuda_bilateral_process_internal(AVFilterContext *ctx,
242  AVFrame *out, AVFrame *in)
243 {
244  CUDABilateralContext *s = ctx->priv;
245  CudaFunctions *cu = s->hwctx->internal->cuda_dl;
246  CUcontext bilateral, cuda_ctx = s->hwctx->cuda_ctx;
247  int i, ret;
248 
249  CUtexObject tex[3] = { 0, 0, 0 };
250 
251  ret = CHECK_CU(cu->cuCtxPushCurrent(cuda_ctx));
252  if (ret < 0)
253  return ret;
254 
255  for (i = 0; i < s->in_planes; i++) {
256  CUDA_TEXTURE_DESC tex_desc = {
257  .filterMode = CU_TR_FILTER_MODE_LINEAR,
258  .flags = 0, // CU_TRSF_READ_AS_INTEGER to get raw ints instead of normalized floats from tex2D
259  };
260 
261  CUDA_RESOURCE_DESC res_desc = {
262  .resType = CU_RESOURCE_TYPE_PITCH2D,
263  .res.pitch2D.format = CU_AD_FORMAT_UNSIGNED_INT8,
264  .res.pitch2D.numChannels = s->in_plane_channels[i],
265  .res.pitch2D.pitchInBytes = in->linesize[i],
266  .res.pitch2D.devPtr = (CUdeviceptr)in->data[i],
267  };
268 
269  if (i == 1 || i == 2) {
270  res_desc.res.pitch2D.width = AV_CEIL_RSHIFT(in->width, s->in_desc->log2_chroma_w);
271  res_desc.res.pitch2D.height = AV_CEIL_RSHIFT(in->height, s->in_desc->log2_chroma_h);
272  } else {
273  res_desc.res.pitch2D.width = in->width;
274  res_desc.res.pitch2D.height = in->height;
275  }
276 
277  ret = CHECK_CU(cu->cuTexObjectCreate(&tex[i], &res_desc, &tex_desc, NULL));
278  if (ret < 0)
279  goto exit;
280  }
281 
282  ret = call_cuda_kernel(ctx, (s->in_plane_channels[1] > 1) ? s->cu_func_uv : s->cu_func,
283  tex, out,
284  out->width, out->height, out->linesize[0],
285  AV_CEIL_RSHIFT(out->width, s->in_desc->log2_chroma_w),
286  AV_CEIL_RSHIFT(out->height, s->in_desc->log2_chroma_h),
287  out->linesize[1] >> ((s->in_plane_channels[1] > 1) ? 1 : 0),
288  s->window_size, s->sigmaS, s->sigmaR);
289 
290 exit:
291  for (i = 0; i < s->in_planes; i++)
292  if (tex[i])
293  CHECK_CU(cu->cuTexObjectDestroy(tex[i]));
294 
295  CHECK_CU(cu->cuCtxPopCurrent(&bilateral));
296 
297  return ret;
298 }
299 
300  static int cuda_bilateral_filter_frame(AVFilterLink *link, AVFrame *in)
301 {
302  AVFilterContext *ctx = link->dst;
303  CUDABilateralContext *s = ctx->priv;
304  AVFilterLink *outlink = ctx->outputs[0];
305  CudaFunctions *cu = s->hwctx->internal->cuda_dl;
306 
307  AVFrame *out;
308  CUcontext bilateral;
309  int ret = 0;
310 
311  out = ff_get_video_buffer(outlink, outlink->w, outlink->h);
312  if (!out) {
313  ret = AVERROR(ENOMEM);
314  goto fail;
315  }
316 
317  ret = CHECK_CU(cu->cuCtxPushCurrent(s->hwctx->cuda_ctx));
318  if (ret < 0)
319  goto fail;
320 
321  ret = cuda_bilateral_process_internal(ctx, out, in);
322  if (ret < 0)
323  goto fail;
324 
325  ret = av_frame_copy_props(out, in);
326  if (ret < 0)
327  goto fail;
328 
329  ret = CHECK_CU(cu->cuCtxPopCurrent(&bilateral));
330  if (ret < 0)
331  goto fail;
332 
333  av_frame_free(&in);
334  return ff_filter_frame(outlink, out);
335 fail:
336  av_frame_free(&in);
337  av_frame_free(&out);
338  return ret;
339 }
340 
341  #define OFFSET(x) offsetof(CUDABilateralContext, x)
342  #define FLAGS (AV_OPT_FLAG_FILTERING_PARAM|AV_OPT_FLAG_VIDEO_PARAM)
343  static const AVOption options[] = {
344  { "sigmaS", "set spatial sigma", OFFSET(sigmaS), AV_OPT_TYPE_FLOAT, {.dbl=0.1}, 0.1, 512, FLAGS },
345  { "sigmaR", "set range sigma", OFFSET(sigmaR), AV_OPT_TYPE_FLOAT, {.dbl=0.1}, 0.1, 512, FLAGS },
346  { "window_size", "set neighbours window_size", OFFSET(window_size), AV_OPT_TYPE_INT, {.i64=1}, 1, 255, FLAGS },
347  { NULL }
348 };
349 
350  static const AVClass cuda_bilateral_class = {
351  .class_name = "cudabilateral",
352  .item_name = av_default_item_name,
353  .option = options,
354  .version = LIBAVUTIL_VERSION_INT,
355 };
356 
357  static const AVFilterPad cuda_bilateral_inputs[] = {
358  {
359  .name = "default",
360  .type = AVMEDIA_TYPE_VIDEO,
361  .filter_frame = cuda_bilateral_filter_frame,
362  },
363 };
364 
365  static const AVFilterPad cuda_bilateral_outputs[] = {
366  {
367  .name = "default",
368  .type = AVMEDIA_TYPE_VIDEO,
369  .config_props = cuda_bilateral_config_props,
370  },
371 };
372 
373  const FFFilter ff_vf_bilateral_cuda = {
374  .p.name = "bilateral_cuda",
375  .p.description = NULL_IF_CONFIG_SMALL("GPU accelerated bilateral filter"),
376  .p.priv_class = &cuda_bilateral_class,
377  .uninit = cudabilateral_uninit,
378  .priv_size = sizeof(CUDABilateralContext),
379  FILTER_INPUTS(cuda_bilateral_inputs),
380  FILTER_OUTPUTS(cuda_bilateral_outputs),
381  FILTER_SINGLE_PIXFMT(AV_PIX_FMT_CUDA),
382  .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE,
383 };
ff_get_video_buffer
AVFrame * ff_get_video_buffer(AVFilterLink *link, int w, int h)
Request a picture buffer with a specific set of permissions.
Definition: video.c:117
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:88
CUDABilateralContext::in_desc
const AVPixFmtDescriptor * in_desc
Definition: vf_bilateral_cuda.c:55
AV_PIX_FMT_CUDA
@ AV_PIX_FMT_CUDA
HW acceleration through CUDA.
Definition: pixfmt.h:260
AVPixelFormat
AVPixelFormat
Pixel format.
Definition: pixfmt.h:71
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
options
static const AVOption options[]
Definition: vf_bilateral_cuda.c:343
ff_filter_frame
int ff_filter_frame(AVFilterLink *link, AVFrame *frame)
Send a frame of data to the next filter.
Definition: avfilter.c:1067
CUDABilateralContext::hwctx
AVCUDADeviceContext * hwctx
Definition: vf_bilateral_cuda.c:53
av_pix_fmt_desc_get
const AVPixFmtDescriptor * av_pix_fmt_desc_get(enum AVPixelFormat pix_fmt)
Definition: pixdesc.c:3456
AVBufferRef::data
uint8_t * data
The data buffer.
Definition: buffer.h:90
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:210
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
CUDABilateralContext::cu_stream
CUstream cu_stream
Definition: vf_bilateral_cuda.c:67
av_frame_free
void av_frame_free(AVFrame **frame)
Free the frame and any dynamically allocated objects in it, e.g.
Definition: frame.c:64
FILTER_INPUTS
#define FILTER_INPUTS(array)
Definition: filters.h:263
AVFrame
This structure describes decoded (raw) audio or video data.
Definition: frame.h:427
pixdesc.h
AVFrame::width
int width
Definition: frame.h:499
BLOCKY
#define BLOCKY
Definition: vf_bilateral_cuda.c:46
AVOption
AVOption.
Definition: opt.h:429
cuda_bilateral_filter_frame
static int cuda_bilateral_filter_frame(AVFilterLink *link, AVFrame *in)
Definition: vf_bilateral_cuda.c:300
float.h
init_processing_chain
static av_cold int init_processing_chain(AVFilterContext *ctx, int width, int height)
Definition: vf_bilateral_cuda.c:114
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:220
video.h
DIV_UP
#define DIV_UP(a, b)
Definition: vf_bilateral_cuda.c:44
cuda_bilateral_load_functions
static av_cold int cuda_bilateral_load_functions(AVFilterContext *ctx)
Definition: vf_bilateral_cuda.c:141
AVFrame::data
uint8_t * data[AV_NUM_DATA_POINTERS]
pointer to the picture/channel planes.
Definition: frame.h:448
av_pix_fmt_count_planes
int av_pix_fmt_count_planes(enum AVPixelFormat pix_fmt)
Definition: pixdesc.c:3496
fail
#define fail()
Definition: checkasm.h:207
cuda_bilateral_inputs
static const AVFilterPad cuda_bilateral_inputs[]
Definition: vf_bilateral_cuda.c:357
AVFilterPad
A filter pad used for either input or output.
Definition: filters.h:39
AV_LOG_ERROR
#define AV_LOG_ERROR
Something went wrong and cannot losslessly be recovered.
Definition: log.h:210
FF_ARRAY_ELEMS
#define FF_ARRAY_ELEMS(a)
Definition: sinewin_tablegen.c:29
av_cold
#define av_cold
Definition: attributes.h:106
CUDABilateralContext::cu_ctx
CUcontext cu_ctx
Definition: vf_bilateral_cuda.c:63
FFFilter
Definition: filters.h:266
CUDABilateralContext
Definition: vf_bilateral_cuda.c:51
s
#define s(width, name)
Definition: cbs_vp9.c:198
AV_CEIL_RSHIFT
#define AV_CEIL_RSHIFT(a, b)
Definition: common.h:60
filters.h
ctx
AVFormatContext * ctx
Definition: movenc.c:49
format_is_supported
static int format_is_supported(enum AVPixelFormat fmt)
Definition: vf_bilateral_cuda.c:85
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:264
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:57
if
if(ret)
Definition: filter_design.txt:179
LIBAVUTIL_VERSION_INT
#define LIBAVUTIL_VERSION_INT
Definition: version.h:85
AVClass
Describe the class of an AVClass context structure.
Definition: log.h:76
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:213
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:599
format
New swscale design to change SwsGraph is what coordinates multiple passes These can include cascaded scaling error diffusion and so on Or we could have separate passes for the vertical and horizontal scaling In between each SwsPass lies a fully allocated image buffer Graph passes may have different levels of e g we can have a single threaded error diffusion pass following a multi threaded scaling pass SwsGraph is internally recreated whenever the image format
Definition: swscale-v2.txt:14
FLAGS
#define FLAGS
Definition: vf_bilateral_cuda.c:342
AVFilterContext::inputs
AVFilterLink ** inputs
array of pointers to input links
Definition: avfilter.h:282
av_default_item_name
const char * av_default_item_name(void *ptr)
Return the context name.
Definition: log.c:241
options
Definition: swscale.c:43
CUDABilateralContext::cu_func_uv
CUfunction cu_func_uv
Definition: vf_bilateral_cuda.c:66
cuda_bilateral_process_internal
static int cuda_bilateral_process_internal(AVFilterContext *ctx, AVFrame *out, AVFrame *in)
Definition: vf_bilateral_cuda.c:241
CUDABilateralContext::window_size
int window_size
Definition: vf_bilateral_cuda.c:59
ff_filter_link
static FilterLink * ff_filter_link(AVFilterLink *link)
Definition: filters.h:198
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:207
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
CUDABilateralContext::cu_module
CUmodule cu_module
Definition: vf_bilateral_cuda.c:64
height
#define height
Definition: dsp.h:89
cuda_bilateral_outputs
static const AVFilterPad cuda_bilateral_outputs[]
Definition: vf_bilateral_cuda.c:365
cudabilateral_uninit
static av_cold void cudabilateral_uninit(AVFilterContext *ctx)
Definition: vf_bilateral_cuda.c:70
ff_vf_bilateral_cuda
const FFFilter ff_vf_bilateral_cuda
Definition: vf_bilateral_cuda.c:373
supported_formats
static enum AVPixelFormat supported_formats[]
Definition: vf_bilateral_cuda.c:38
cuda_bilateral_config_props
static av_cold int cuda_bilateral_config_props(AVFilterLink *outlink)
Definition: vf_bilateral_cuda.c:178
AV_OPT_TYPE_FLOAT
@ AV_OPT_TYPE_FLOAT
Underlying C type is float.
Definition: opt.h:271
i
#define i(width, name, range_min, range_max)
Definition: cbs_h2645.c:256
CUDABilateralContext::sigmaS
float sigmaS
Definition: vf_bilateral_cuda.c:60
internal.h
set_format_info
static av_cold void set_format_info(AVFilterContext *ctx, enum AVPixelFormat format)
Definition: vf_bilateral_cuda.c:95
common.h
CUDABilateralContext::sigmaR
float sigmaR
Definition: vf_bilateral_cuda.c:61
AVFilterPad::name
const char * name
Pad name.
Definition: filters.h:45
BLOCKX
#define BLOCKX
Definition: vf_bilateral_cuda.c:45
AVHWFramesContext
This struct describes a set or pool of "hardware" frames (i.e.
Definition: hwcontext.h:118
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:204
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:81
AVHWFramesContext::device_ctx
AVHWDeviceContext * device_ctx
The parent AVHWDeviceContext.
Definition: hwcontext.h:137
cuda_check.h
cuda_bilateral_class
static const AVClass cuda_bilateral_class
Definition: vf_bilateral_cuda.c:350
AVFrame::height
int height
Definition: frame.h:499
AV_OPT_TYPE_INT
@ AV_OPT_TYPE_INT
Underlying C type is int.
Definition: opt.h:259
avfilter.h
Windows::Graphics::DirectX::Direct3D11::p
IDirect3DDxgiInterfaceAccess _COM_Outptr_ void ** p
Definition: vsrc_gfxcapture_winrt.hpp:53
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:78
AVFilterContext
An instance of a filter.
Definition: avfilter.h:274
AVMEDIA_TYPE_VIDEO
@ AVMEDIA_TYPE_VIDEO
Definition: avutil.h:200
FFFilter::p
AVFilter p
The public AVFilter.
Definition: filters.h:270
OFFSET
#define OFFSET(x)
Definition: vf_bilateral_cuda.c:341
AVPixFmtDescriptor
Descriptor that unambiguously describes how the bits of a pixel are stored in the up to 4 data planes...
Definition: pixdesc.h:69
CUDABilateralContext::cu_func
CUfunction cu_func
Definition: vf_bilateral_cuda.c:65
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:472
av_log
#define av_log(a,...)
Definition: tableprint_vlc.h:27
width
#define width
Definition: dsp.h:89
CHECK_CU
#define CHECK_CU(x)
Definition: vf_bilateral_cuda.c:48
CUDABilateralContext::in_planes
int in_planes
Definition: vf_bilateral_cuda.c:56
FILTER_SINGLE_PIXFMT
#define FILTER_SINGLE_PIXFMT(pix_fmt_)
Definition: filters.h:253
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:3376
Generated on Tue Nov 18 2025 19:23:07 for FFmpeg by   doxygen 1.8.17

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