FFmpeg
vf_nlmeans_opencl.c
Go to the documentation of this file.
1 /*
2  * This file is part of FFmpeg.
3  *
4  * FFmpeg is free software; you can redistribute it and/or
5  * modify it under the terms of the GNU Lesser General Public
6  * License as published by the Free Software Foundation; either
7  * version 2.1 of the License, or (at your option) any later version.
8  *
9  * FFmpeg is distributed in the hope that it will be useful,
10  * but WITHOUT ANY WARRANTY; without even the implied warranty of
11  * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU
12  * Lesser General Public License for more details.
13  *
14  * You should have received a copy of the GNU Lesser General Public
15  * License along with FFmpeg; if not, write to the Free Software
16  * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA
17  */
18 #include <float.h>
19 
20 #include "libavutil/avassert.h"
21 #include "libavutil/common.h"
22 #include "libavutil/imgutils.h"
23 #include "libavutil/mem.h"
24 #include "libavutil/opt.h"
25 #include "libavutil/pixdesc.h"
26 
27 #include "avfilter.h"
28 #include "internal.h"
29 #include "opencl.h"
30 #include "opencl_source.h"
31 #include "video.h"
32 
33 // TODO:
34 // the integral image may overflow 32bit, consider using 64bit
35 
36 static const enum AVPixelFormat supported_formats[] = {
37  AV_PIX_FMT_YUV420P,
38  AV_PIX_FMT_YUV444P,
39  AV_PIX_FMT_GBRP,
40 };
41 
42 static int is_format_supported(enum AVPixelFormat fmt)
43 {
44  int i;
45 
46  for (i = 0; i < FF_ARRAY_ELEMS(supported_formats); i++)
47  if (supported_formats[i] == fmt)
48  return 1;
49  return 0;
50 }
51 
52 typedef struct NLMeansOpenCLContext {
53  OpenCLFilterContext ocf;
54  int initialised;
55  cl_kernel vert_kernel;
56  cl_kernel horiz_kernel;
57  cl_kernel accum_kernel;
58  cl_kernel average_kernel;
59  cl_mem integral_img;
60  cl_mem weight;
61  cl_mem sum;
62  cl_mem overflow; // overflow in integral image?
63  double sigma;
64  float h;
65  int chroma_w;
66  int chroma_h;
67  int patch_size;
68  int patch_size_uv;
69  int research_size;
70  int research_size_uv;
71  cl_command_queue command_queue;
72 } NLMeansOpenCLContext;
73 
74 static int nlmeans_opencl_init(AVFilterContext *avctx, int width, int height)
75 {
76  NLMeansOpenCLContext *ctx = avctx->priv;
77  cl_int cle;
78  int err;
79  int weight_buf_size = width * height * sizeof(float);
80 
81  ctx->h = ctx->sigma * 10;
82  if (!(ctx->research_size & 1)) {
83  ctx->research_size |= 1;
84  av_log(avctx, AV_LOG_WARNING,
85  "research_size should be odd, set to %d",
86  ctx->research_size);
87  }
88 
89  if (!(ctx->patch_size & 1)) {
90  ctx->patch_size |= 1;
91  av_log(avctx, AV_LOG_WARNING,
92  "patch_size should be odd, set to %d",
93  ctx->patch_size);
94  }
95 
96  if (!ctx->research_size_uv)
97  ctx->research_size_uv = ctx->research_size;
98  if (!ctx->patch_size_uv)
99  ctx->patch_size_uv = ctx->patch_size;
100 
101  err = ff_opencl_filter_load_program(avctx, &ff_source_nlmeans_cl, 1);
102  if (err < 0)
103  goto fail;
104 
105  ctx->command_queue = clCreateCommandQueue(ctx->ocf.hwctx->context,
106  ctx->ocf.hwctx->device_id,
107  0, &cle);
108  CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create OpenCL "
109  "command queue %d.\n", cle);
110 
111  ctx->vert_kernel = clCreateKernel(ctx->ocf.program,
112  "vert_sum", &cle);
113  CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create "
114  "vert_sum kernel %d.\n", cle);
115 
116  ctx->horiz_kernel = clCreateKernel(ctx->ocf.program,
117  "horiz_sum", &cle);
118  CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create "
119  "horiz_sum kernel %d.\n", cle);
120 
121  ctx->accum_kernel = clCreateKernel(ctx->ocf.program,
122  "weight_accum", &cle);
123  CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create "
124  "accum kernel %d.\n", cle);
125 
126  ctx->average_kernel = clCreateKernel(ctx->ocf.program,
127  "average", &cle);
128  CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create "
129  "average kernel %d.\n", cle);
130 
131  ctx->integral_img = clCreateBuffer(ctx->ocf.hwctx->context, 0,
132  4 * width * height * sizeof(cl_int),
133  NULL, &cle);
134  CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create "
135  "integral image %d.\n", cle);
136 
137  ctx->weight = clCreateBuffer(ctx->ocf.hwctx->context, 0,
138  weight_buf_size, NULL, &cle);
139  CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create "
140  "weight buffer %d.\n", cle);
141 
142  ctx->sum = clCreateBuffer(ctx->ocf.hwctx->context, 0,
143  weight_buf_size, NULL, &cle);
144  CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create "
145  "sum buffer %d.\n", cle);
146 
147  ctx->overflow = clCreateBuffer(ctx->ocf.hwctx->context, 0,
148  sizeof(cl_int), NULL, &cle);
149  CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create "
150  "overflow buffer %d.\n", cle);
151 
152  ctx->initialised = 1;
153  return 0;
154 
155 fail:
156  CL_RELEASE_KERNEL(ctx->vert_kernel);
157  CL_RELEASE_KERNEL(ctx->horiz_kernel);
158  CL_RELEASE_KERNEL(ctx->accum_kernel);
159  CL_RELEASE_KERNEL(ctx->average_kernel);
160 
161  CL_RELEASE_MEMORY(ctx->integral_img);
162  CL_RELEASE_MEMORY(ctx->weight);
163  CL_RELEASE_MEMORY(ctx->sum);
164  CL_RELEASE_MEMORY(ctx->overflow);
165 
166  CL_RELEASE_QUEUE(ctx->command_queue);
167  return err;
168 }
169 
170 static int nlmeans_plane(AVFilterContext *avctx, cl_mem dst, cl_mem src,
171  cl_int width, cl_int height, cl_int p, cl_int r)
172 {
173  NLMeansOpenCLContext *ctx = avctx->priv;
174  const float zero = 0.0f;
175  const size_t worksize1[] = {height};
176  const size_t worksize2[] = {width};
177  const size_t worksize3[2] = {width, height};
178  int i, dx, dy, err = 0, weight_buf_size;
179  cl_int cle;
180  int nb_pixel, *tmp = NULL, idx = 0;
181  cl_int *dxdy = NULL;
182 
183  weight_buf_size = width * height * sizeof(float);
184  cle = clEnqueueFillBuffer(ctx->command_queue, ctx->weight,
185  &zero, sizeof(float), 0, weight_buf_size,
186  0, NULL, NULL);
187  CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to fill weight buffer: %d.\n",
188  cle);
189  cle = clEnqueueFillBuffer(ctx->command_queue, ctx->sum,
190  &zero, sizeof(float), 0, weight_buf_size,
191  0, NULL, NULL);
192  CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to fill sum buffer: %d.\n",
193  cle);
194 
195  nb_pixel = (2 * r + 1) * (2 * r + 1) - 1;
196  dxdy = av_malloc(nb_pixel * 2 * sizeof(cl_int));
197  tmp = av_malloc(nb_pixel * 2 * sizeof(int));
198 
199  if (!dxdy || !tmp)
200  goto fail;
201 
202  for (dx = -r; dx <= r; dx++) {
203  for (dy = -r; dy <= r; dy++) {
204  if (dx || dy) {
205  tmp[idx++] = dx;
206  tmp[idx++] = dy;
207  }
208  }
209  }
210  // repack dx/dy seperately, as we want to do four pairs of dx/dy in a batch
211  for (i = 0; i < nb_pixel / 4; i++) {
212  dxdy[i * 8] = tmp[i * 8]; // dx0
213  dxdy[i * 8 + 1] = tmp[i * 8 + 2]; // dx1
214  dxdy[i * 8 + 2] = tmp[i * 8 + 4]; // dx2
215  dxdy[i * 8 + 3] = tmp[i * 8 + 6]; // dx3
216  dxdy[i * 8 + 4] = tmp[i * 8 + 1]; // dy0
217  dxdy[i * 8 + 5] = tmp[i * 8 + 3]; // dy1
218  dxdy[i * 8 + 6] = tmp[i * 8 + 5]; // dy2
219  dxdy[i * 8 + 7] = tmp[i * 8 + 7]; // dy3
220  }
221  av_freep(&tmp);
222 
223  for (i = 0; i < nb_pixel / 4; i++) {
224  cl_int *dx_cur = dxdy + 8 * i;
225  cl_int *dy_cur = dxdy + 8 * i + 4;
226 
227  // horizontal pass
228  // integral(x,y) = sum([u(v,y) - u(v+dx,y+dy)]^2) for v in [0, x]
229  CL_SET_KERNEL_ARG(ctx->horiz_kernel, 0, cl_mem, &ctx->integral_img);
230  CL_SET_KERNEL_ARG(ctx->horiz_kernel, 1, cl_mem, &src);
231  CL_SET_KERNEL_ARG(ctx->horiz_kernel, 2, cl_int, &width);
232  CL_SET_KERNEL_ARG(ctx->horiz_kernel, 3, cl_int, &height);
233  CL_SET_KERNEL_ARG(ctx->horiz_kernel, 4, cl_int4, dx_cur);
234  CL_SET_KERNEL_ARG(ctx->horiz_kernel, 5, cl_int4, dy_cur);
235  cle = clEnqueueNDRangeKernel(ctx->command_queue, ctx->horiz_kernel, 1,
236  NULL, worksize1, NULL, 0, NULL, NULL);
237  CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to enqueue horiz_kernel: %d.\n",
238  cle);
239  // vertical pass
240  // integral(x, y) = sum(integral(x, v)) for v in [0, y]
241  CL_SET_KERNEL_ARG(ctx->vert_kernel, 0, cl_mem, &ctx->integral_img);
242  CL_SET_KERNEL_ARG(ctx->vert_kernel, 1, cl_mem, &ctx->overflow);
243  CL_SET_KERNEL_ARG(ctx->vert_kernel, 2, cl_int, &width);
244  CL_SET_KERNEL_ARG(ctx->vert_kernel, 3, cl_int, &height);
245  cle = clEnqueueNDRangeKernel(ctx->command_queue, ctx->vert_kernel,
246  1, NULL, worksize2, NULL, 0, NULL, NULL);
247  CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to enqueue vert_kernel: %d.\n",
248  cle);
249 
250  // accumulate weights
251  CL_SET_KERNEL_ARG(ctx->accum_kernel, 0, cl_mem, &ctx->sum);
252  CL_SET_KERNEL_ARG(ctx->accum_kernel, 1, cl_mem, &ctx->weight);
253  CL_SET_KERNEL_ARG(ctx->accum_kernel, 2, cl_mem, &ctx->integral_img);
254  CL_SET_KERNEL_ARG(ctx->accum_kernel, 3, cl_mem, &src);
255  CL_SET_KERNEL_ARG(ctx->accum_kernel, 4, cl_int, &width);
256  CL_SET_KERNEL_ARG(ctx->accum_kernel, 5, cl_int, &height);
257  CL_SET_KERNEL_ARG(ctx->accum_kernel, 6, cl_int, &p);
258  CL_SET_KERNEL_ARG(ctx->accum_kernel, 7, cl_float, &ctx->h);
259  CL_SET_KERNEL_ARG(ctx->accum_kernel, 8, cl_int4, dx_cur);
260  CL_SET_KERNEL_ARG(ctx->accum_kernel, 9, cl_int4, dy_cur);
261  cle = clEnqueueNDRangeKernel(ctx->command_queue, ctx->accum_kernel,
262  2, NULL, worksize3, NULL, 0, NULL, NULL);
263  CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to enqueue kernel: %d.\n", cle);
264  }
265  av_freep(&dxdy);
266 
267  // average
268  CL_SET_KERNEL_ARG(ctx->average_kernel, 0, cl_mem, &dst);
269  CL_SET_KERNEL_ARG(ctx->average_kernel, 1, cl_mem, &src);
270  CL_SET_KERNEL_ARG(ctx->average_kernel, 2, cl_mem, &ctx->sum);
271  CL_SET_KERNEL_ARG(ctx->average_kernel, 3, cl_mem, &ctx->weight);
272  cle = clEnqueueNDRangeKernel(ctx->command_queue, ctx->average_kernel, 2,
273  NULL, worksize3, NULL, 0, NULL, NULL);
274  CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to enqueue average kernel: %d.\n",
275  cle);
276  cle = clFlush(ctx->command_queue);
277  CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to flush command queue: %d.\n", cle);
278 fail:
279  if (tmp)
280  av_freep(&tmp);
281  if (dxdy)
282  av_freep(&dxdy);
283  return err;
284 }
285 
286 static int nlmeans_opencl_filter_frame(AVFilterLink *inlink, AVFrame *input)
287 {
288  AVFilterContext *avctx = inlink->dst;
289  AVFilterLink *outlink = avctx->outputs[0];
290  NLMeansOpenCLContext *ctx = avctx->priv;
291  AVFrame *output = NULL;
292  AVHWFramesContext *input_frames_ctx;
293  const AVPixFmtDescriptor *desc;
294  enum AVPixelFormat in_format;
295  cl_mem src, dst;
296  const cl_int zero = 0;
297  int w, h, err, cle, overflow, p, patch, research;
298 
299  av_log(ctx, AV_LOG_DEBUG, "Filter input: %s, %ux%u (%"PRId64").\n",
300  av_get_pix_fmt_name(input->format),
301  input->width, input->height, input->pts);
302 
303  if (!input->hw_frames_ctx)
304  return AVERROR(EINVAL);
305  input_frames_ctx = (AVHWFramesContext*)input->hw_frames_ctx->data;
306  in_format = input_frames_ctx->sw_format;
307 
308  output = ff_get_video_buffer(outlink, outlink->w, outlink->h);
309  if (!output) {
310  err = AVERROR(ENOMEM);
311  goto fail;
312  }
313 
314  err = av_frame_copy_props(output, input);
315  if (err < 0)
316  goto fail;
317 
318  if (!ctx->initialised) {
319  desc = av_pix_fmt_desc_get(in_format);
320  if (!is_format_supported(in_format)) {
321  err = AVERROR(EINVAL);
322  av_log(avctx, AV_LOG_ERROR, "input format %s not supported\n",
323  av_get_pix_fmt_name(in_format));
324  goto fail;
325  }
326  ctx->chroma_w = AV_CEIL_RSHIFT(inlink->w, desc->log2_chroma_w);
327  ctx->chroma_h = AV_CEIL_RSHIFT(inlink->h, desc->log2_chroma_h);
328 
329  err = nlmeans_opencl_init(avctx, inlink->w, inlink->h);
330  if (err < 0)
331  goto fail;
332  }
333 
334  cle = clEnqueueWriteBuffer(ctx->command_queue, ctx->overflow, CL_FALSE,
335  0, sizeof(cl_int), &zero, 0, NULL, NULL);
336  CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to initialize overflow"
337  "detection buffer %d.\n", cle);
338 
339  for (p = 0; p < FF_ARRAY_ELEMS(output->data); p++) {
340  src = (cl_mem) input->data[p];
341  dst = (cl_mem) output->data[p];
342 
343  if (!dst)
344  break;
345  av_assert0(src);
346  w = p ? ctx->chroma_w : inlink->w;
347  h = p ? ctx->chroma_h : inlink->h;
348  patch = (p ? ctx->patch_size_uv : ctx->patch_size) / 2;
349  research = (p ? ctx->research_size_uv : ctx->research_size) / 2;
350  err = nlmeans_plane(avctx, dst, src, w, h, patch, research);
351  if (err < 0)
352  goto fail;
353  }
354  // overflow occurred?
355  cle = clEnqueueReadBuffer(ctx->command_queue, ctx->overflow, CL_FALSE,
356  0, sizeof(cl_int), &overflow, 0, NULL, NULL);
357  CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to read overflow: %d.\n", cle);
358 
359  cle = clFinish(ctx->command_queue);
360  CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to finish kernel: %d.\n", cle);
361 
362  if (overflow > 0)
363  av_log(avctx, AV_LOG_ERROR, "integral image overflow %d\n", overflow);
364 
365  av_frame_free(&input);
366 
367  av_log(ctx, AV_LOG_DEBUG, "Filter output: %s, %ux%u (%"PRId64").\n",
368  av_get_pix_fmt_name(output->format),
369  output->width, output->height, output->pts);
370 
371  return ff_filter_frame(outlink, output);
372 
373 fail:
374  clFinish(ctx->command_queue);
375  av_frame_free(&input);
376  av_frame_free(&output);
377  return err;
378 }
379 
380 static av_cold void nlmeans_opencl_uninit(AVFilterContext *avctx)
381 {
382  NLMeansOpenCLContext *ctx = avctx->priv;
383  cl_int cle;
384 
385  CL_RELEASE_KERNEL(ctx->vert_kernel);
386  CL_RELEASE_KERNEL(ctx->horiz_kernel);
387  CL_RELEASE_KERNEL(ctx->accum_kernel);
388  CL_RELEASE_KERNEL(ctx->average_kernel);
389 
390  CL_RELEASE_MEMORY(ctx->integral_img);
391  CL_RELEASE_MEMORY(ctx->weight);
392  CL_RELEASE_MEMORY(ctx->sum);
393  CL_RELEASE_MEMORY(ctx->overflow);
394 
395  CL_RELEASE_QUEUE(ctx->command_queue);
396 
397  ff_opencl_filter_uninit(avctx);
398 }
399 
400 #define OFFSET(x) offsetof(NLMeansOpenCLContext, x)
401 #define FLAGS (AV_OPT_FLAG_FILTERING_PARAM | AV_OPT_FLAG_VIDEO_PARAM)
402 static const AVOption nlmeans_opencl_options[] = {
403  { "s", "denoising strength", OFFSET(sigma), AV_OPT_TYPE_DOUBLE, { .dbl = 1.0 }, 1.0, 30.0, FLAGS },
404  { "p", "patch size", OFFSET(patch_size), AV_OPT_TYPE_INT, { .i64 = 2*3+1 }, 0, 99, FLAGS },
405  { "pc", "patch size for chroma planes", OFFSET(patch_size_uv), AV_OPT_TYPE_INT, { .i64 = 0 }, 0, 99, FLAGS },
406  { "r", "research window", OFFSET(research_size), AV_OPT_TYPE_INT, { .i64 = 7*2+1 }, 0, 99, FLAGS },
407  { "rc", "research window for chroma planes", OFFSET(research_size_uv), AV_OPT_TYPE_INT, { .i64 = 0 }, 0, 99, FLAGS },
408  { NULL }
409 };
410 
411 AVFILTER_DEFINE_CLASS(nlmeans_opencl);
412 
413 static const AVFilterPad nlmeans_opencl_inputs[] = {
414  {
415  .name = "default",
416  .type = AVMEDIA_TYPE_VIDEO,
417  .filter_frame = &nlmeans_opencl_filter_frame,
418  .config_props = &ff_opencl_filter_config_input,
419  },
420 };
421 
422 static const AVFilterPad nlmeans_opencl_outputs[] = {
423  {
424  .name = "default",
425  .type = AVMEDIA_TYPE_VIDEO,
426  .config_props = &ff_opencl_filter_config_output,
427  },
428 };
429 
430 const AVFilter ff_vf_nlmeans_opencl = {
431  .name = "nlmeans_opencl",
432  .description = NULL_IF_CONFIG_SMALL("Non-local means denoiser through OpenCL"),
433  .priv_size = sizeof(NLMeansOpenCLContext),
434  .priv_class = &nlmeans_opencl_class,
435  .init = &ff_opencl_filter_init,
436  .uninit = &nlmeans_opencl_uninit,
437  FILTER_INPUTS(nlmeans_opencl_inputs),
438  FILTER_OUTPUTS(nlmeans_opencl_outputs),
439  FILTER_SINGLE_PIXFMT(AV_PIX_FMT_OPENCL),
440  .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE,
441  .flags = AVFILTER_FLAG_HWDEVICE,
442 };
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:108
NLMeansOpenCLContext::horiz_kernel
cl_kernel horiz_kernel
Definition: vf_nlmeans_opencl.c:56
AV_LOG_WARNING
#define AV_LOG_WARNING
Something somehow does not look correct.
Definition: log.h:186
nlmeans_opencl_options
static const AVOption nlmeans_opencl_options[]
Definition: vf_nlmeans_opencl.c:402
AVPixelFormat
AVPixelFormat
Pixel format.
Definition: pixfmt.h:64
NLMeansOpenCLContext::research_size_uv
int research_size_uv
Definition: vf_nlmeans_opencl.c:70
r
const char * r
Definition: vf_curves.c:126
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
NLMeansOpenCLContext::chroma_h
int chroma_h
Definition: vf_nlmeans_opencl.c:66
nlmeans_opencl_init
static int nlmeans_opencl_init(AVFilterContext *avctx, int width, int height)
Definition: vf_nlmeans_opencl.c:74
CL_SET_KERNEL_ARG
#define CL_SET_KERNEL_ARG(kernel, arg_num, type, arg)
set argument to specific Kernel.
Definition: opencl.h:61
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:364
ff_filter_frame
int ff_filter_frame(AVFilterLink *link, AVFrame *frame)
Send a frame of data to the next filter.
Definition: avfilter.c:978
AVFILTER_DEFINE_CLASS
AVFILTER_DEFINE_CLASS(nlmeans_opencl)
NLMeansOpenCLContext
Definition: vf_nlmeans_opencl.c:52
av_pix_fmt_desc_get
const AVPixFmtDescriptor * av_pix_fmt_desc_get(enum AVPixelFormat pix_fmt)
Definition: pixdesc.c:2964
output
filter_frame For filters that do not use the this method is called when a frame is pushed to the filter s input It can be called at any time except in a reentrant way If the input frame is enough to produce output
Definition: filter_design.txt:225
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:100
test::height
int height
Definition: vc1dsp.c:39
nlmeans_opencl_outputs
static const AVFilterPad nlmeans_opencl_outputs[]
Definition: vf_nlmeans_opencl.c:422
AVFrame
This structure describes decoded (raw) audio or video data.
Definition: frame.h:340
tmp
static uint8_t tmp[11]
Definition: aes_ctr.c:28
pixdesc.h
w
uint8_t w
Definition: llviddspenc.c:38
opencl.h
AVOption
AVOption.
Definition: opt.h:251
float.h
ff_opencl_filter_load_program
int ff_opencl_filter_load_program(AVFilterContext *avctx, const char **program_source_array, int nb_strings)
Load a new OpenCL program from strings in memory.
Definition: opencl.c:156
is_format_supported
static int is_format_supported(enum AVPixelFormat fmt)
Definition: vf_nlmeans_opencl.c:42
AVFilter::name
const char * name
Filter name.
Definition: avfilter.h:170
video.h
ff_vf_nlmeans_opencl
const AVFilter ff_vf_nlmeans_opencl
Definition: vf_nlmeans_opencl.c:430
av_malloc
#define av_malloc(s)
Definition: tableprint_vlc.h:30
ff_source_nlmeans_cl
const char * ff_source_nlmeans_cl
NLMeansOpenCLContext::patch_size
int patch_size
Definition: vf_nlmeans_opencl.c:67
AVFilterContext::priv
void * priv
private data for use by the filter
Definition: avfilter.h:412
fail
#define fail()
Definition: checkasm.h:142
FLAGS
#define FLAGS
Definition: vf_nlmeans_opencl.c:401
NLMeansOpenCLContext::patch_size_uv
int patch_size_uv
Definition: vf_nlmeans_opencl.c:68
ff_opencl_filter_config_output
int ff_opencl_filter_config_output(AVFilterLink *outlink)
Create a suitable hardware frames context for the output.
Definition: opencl.c:81
AVFilterPad
A filter pad used for either input or output.
Definition: internal.h:47
avassert.h
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
float
float
Definition: af_crystalizer.c:121
width
#define width
AV_CEIL_RSHIFT
#define AV_CEIL_RSHIFT(a, b)
Definition: common.h:51
AV_OPT_TYPE_DOUBLE
@ AV_OPT_TYPE_DOUBLE
Definition: opt.h:227
NLMeansOpenCLContext::h
float h
Definition: vf_nlmeans_opencl.c:64
init
int(* init)(AVBSFContext *ctx)
Definition: dts2pts_bsf.c:365
av_assert0
#define av_assert0(cond)
assert() equivalent, that is always enabled.
Definition: avassert.h:40
AV_LOG_DEBUG
#define AV_LOG_DEBUG
Stuff which is only useful for libav* developers.
Definition: log.h:201
NLMeansOpenCLContext::integral_img
cl_mem integral_img
Definition: vf_nlmeans_opencl.c:59
ctx
AVFormatContext * ctx
Definition: movenc.c:48
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:192
NLMeansOpenCLContext::weight
cl_mem weight
Definition: vf_nlmeans_opencl.c:60
CL_RELEASE_KERNEL
#define CL_RELEASE_KERNEL(k)
release an OpenCL Kernel
Definition: opencl.h:101
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:736
AV_PIX_FMT_OPENCL
@ AV_PIX_FMT_OPENCL
Hardware surfaces for OpenCL.
Definition: pixfmt.h:355
NLMeansOpenCLContext::vert_kernel
cl_kernel vert_kernel
Definition: vf_nlmeans_opencl.c:55
NLMeansOpenCLContext::initialised
int initialised
Definition: vf_nlmeans_opencl.c:54
NLMeansOpenCLContext::sum
cl_mem sum
Definition: vf_nlmeans_opencl.c:61
test::width
int width
Definition: vc1dsp.c:38
NULL_IF_CONFIG_SMALL
#define NULL_IF_CONFIG_SMALL(x)
Return NULL if CONFIG_SMALL is true, otherwise the argument without modification.
Definition: internal.h:106
CL_RELEASE_MEMORY
#define CL_RELEASE_MEMORY(m)
release an OpenCL Memory Object
Definition: opencl.h:114
AVFILTER_FLAG_HWDEVICE
#define AVFILTER_FLAG_HWDEVICE
The filter can create hardware frames using AVFilterContext.hw_device_ctx.
Definition: avfilter.h:138
nlmeans_opencl_inputs
static const AVFilterPad nlmeans_opencl_inputs[]
Definition: vf_nlmeans_opencl.c:413
NLMeansOpenCLContext::accum_kernel
cl_kernel accum_kernel
Definition: vf_nlmeans_opencl.c:57
NLMeansOpenCLContext::average_kernel
cl_kernel average_kernel
Definition: vf_nlmeans_opencl.c:58
NLMeansOpenCLContext::overflow
cl_mem overflow
Definition: vf_nlmeans_opencl.c:62
NLMeansOpenCLContext::sigma
double sigma
Definition: vf_nlmeans_opencl.c:63
height
#define height
nlmeans_opencl_uninit
static av_cold void nlmeans_opencl_uninit(AVFilterContext *avctx)
Definition: vf_nlmeans_opencl.c:380
opencl_source.h
input
and forward the test the status of outputs and forward it to the corresponding return FFERROR_NOT_READY If the filters stores internally one or a few frame for some input
Definition: filter_design.txt:172
ff_opencl_filter_config_input
int ff_opencl_filter_config_input(AVFilterLink *inlink)
Check that the input link contains a suitable hardware frames context and extract the device from it.
Definition: opencl.c:45
internal.h
FILTER_SINGLE_PIXFMT
#define FILTER_SINGLE_PIXFMT(pix_fmt_)
Definition: internal.h:182
NLMeansOpenCLContext::command_queue
cl_command_queue command_queue
Definition: vf_nlmeans_opencl.c:71
i
#define i(width, name, range_min, range_max)
Definition: cbs_h2645.c:245
common.h
AVFilterPad::name
const char * name
Pad name.
Definition: internal.h:53
AVFilter
Filter definition.
Definition: avfilter.h:166
AVHWFramesContext
This struct describes a set or pool of "hardware" frames (i.e.
Definition: hwcontext.h:124
ff_opencl_filter_init
int ff_opencl_filter_init(AVFilterContext *avctx)
Initialise an OpenCL filter context.
Definition: opencl.c:132
NLMeansOpenCLContext::chroma_w
int chroma_w
Definition: vf_nlmeans_opencl.c:65
NLMeansOpenCLContext::ocf
OpenCLFilterContext ocf
Definition: vf_nlmeans_opencl.c:53
nlmeans_plane
static int nlmeans_plane(AVFilterContext *avctx, cl_mem dst, cl_mem src, cl_int width, cl_int height, cl_int p, cl_int r)
Definition: vf_nlmeans_opencl.c:170
AV_OPT_TYPE_INT
@ AV_OPT_TYPE_INT
Definition: opt.h:225
avfilter.h
nlmeans_opencl_filter_frame
static int nlmeans_opencl_filter_frame(AVFilterLink *inlink, AVFrame *input)
Definition: vf_nlmeans_opencl.c:286
OpenCLFilterContext
Definition: opencl.h:36
ff_opencl_filter_uninit
void ff_opencl_filter_uninit(AVFilterContext *avctx)
Uninitialise an OpenCL filter context.
Definition: opencl.c:141
OFFSET
#define OFFSET(x)
Definition: vf_nlmeans_opencl.c:400
CL_RELEASE_QUEUE
#define CL_RELEASE_QUEUE(q)
release an OpenCL Command Queue
Definition: opencl.h:127
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:397
NLMeansOpenCLContext::research_size
int research_size
Definition: vf_nlmeans_opencl.c:69
AV_PIX_FMT_GBRP
@ AV_PIX_FMT_GBRP
planar GBR 4:4:4 24bpp
Definition: pixfmt.h:158
desc
const char * desc
Definition: libsvtav1.c:83
AVMEDIA_TYPE_VIDEO
@ AVMEDIA_TYPE_VIDEO
Definition: avutil.h:201
zero
#define zero
Definition: regdef.h:64
mem.h
overflow
Undefined Behavior In the C some operations are like signed integer overflow
Definition: undefined.txt:3
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:193
av_freep
#define av_freep(p)
Definition: tableprint_vlc.h:34
src
INIT_CLIP pixel * src
Definition: h264pred_template.c:418
supported_formats
static enum AVPixelFormat supported_formats[]
Definition: vf_nlmeans_opencl.c:36
imgutils.h
av_log
#define av_log(a,...)
Definition: tableprint_vlc.h:27
CL_FAIL_ON_ERROR
#define CL_FAIL_ON_ERROR(errcode,...)
A helper macro to handle OpenCL errors.
Definition: opencl.h:74
uninit
static av_cold int uninit(AVCodecContext *avctx)
Definition: crystalhd.c:285
h
h
Definition: vp9dsp_template.c:2038
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:2884
AVFilterContext::outputs
AVFilterLink ** outputs
array of pointers to output links
Definition: avfilter.h:409
Generated on Mon Dec 4 2023 19:22:06 for FFmpeg by   doxygen 1.8.17