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[] = {
40 };
41 
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 {
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;
71  cl_command_queue command_queue;
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 
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:
160 
163  CL_RELEASE_MEMORY(ctx->sum);
165 
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 
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 
381 {
382  NLMeansOpenCLContext *ctx = avctx->priv;
383  cl_int cle;
384 
389 
392  CL_RELEASE_MEMORY(ctx->sum);
394 
396 
398 }
399 
400 #define OFFSET(x) offsetof(NLMeansOpenCLContext, x)
401 #define FLAGS (AV_OPT_FLAG_FILTERING_PARAM | AV_OPT_FLAG_VIDEO_PARAM)
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 
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  { NULL }
421 };
422 
424  {
425  .name = "default",
426  .type = AVMEDIA_TYPE_VIDEO,
427  .config_props = &ff_opencl_filter_config_output,
428  },
429  { NULL }
430 };
431 
433  .name = "nlmeans_opencl",
434  .description = NULL_IF_CONFIG_SMALL("Non-local means denoiser through OpenCL"),
435  .priv_size = sizeof(NLMeansOpenCLContext),
436  .priv_class = &nlmeans_opencl_class,
440  .inputs = nlmeans_opencl_inputs,
441  .outputs = nlmeans_opencl_outputs,
442  .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE,
443 };
#define NULL
Definition: coverity.c:32
#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:385
const AVPixFmtDescriptor * av_pix_fmt_desc_get(enum AVPixelFormat pix_fmt)
Definition: pixdesc.c:2522
This structure describes decoded (raw) audio or video data.
Definition: frame.h:295
cl_command_queue command_queue
AVFILTER_DEFINE_CLASS(nlmeans_opencl)
AVOption.
Definition: opt.h:246
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:60
const char * fmt
Definition: avisynth_c.h:861
int ff_opencl_filter_query_formats(AVFilterContext *avctx)
Return that all inputs and outputs support only AV_PIX_FMT_OPENCL.
Definition: opencl.c:28
planar YUV 4:4:4, 24bpp, (1 Cr & Cb sample per 1x1 Y samples)
Definition: pixfmt.h:71
misc image utilities
#define AV_LOG_WARNING
Something somehow does not look correct.
Definition: log.h:182
Main libavfilter public API header.
Memory handling functions.
static const AVFilterPad nlmeans_opencl_outputs[]
const char * desc
Definition: nvenc.c:68
static av_cold int init(AVCodecContext *avctx)
Definition: avrndec.c:35
#define CL_RELEASE_MEMORY(m)
release an OpenCL Memory Object
Definition: opencl.h:114
OpenCLFilterContext ocf
planar GBR 4:4:4 24bpp
Definition: pixfmt.h:168
AVFrame * ff_get_video_buffer(AVFilterLink *link, int w, int h)
Request a picture buffer with a specific set of permissions.
Definition: video.c:99
#define src
Definition: vp8dsp.c:254
int ff_opencl_filter_init(AVFilterContext *avctx)
Initialise an OpenCL filter context.
Definition: opencl.c:147
uint8_t log2_chroma_w
Amount to shift the luma width right to find the chroma width.
Definition: pixdesc.h:92
static int is_format_supported(enum AVPixelFormat fmt)
AVOpenCLDeviceContext * hwctx
Definition: opencl.h:41
AVBufferRef * hw_frames_ctx
For hwaccel-format frames, this should be a reference to the AVHWFramesContext describing the frame...
Definition: frame.h:634
const char * name
Pad name.
Definition: internal.h:60
#define av_assert0(cond)
assert() equivalent, that is always enabled.
Definition: avassert.h:37
int ff_filter_frame(AVFilterLink *link, AVFrame *frame)
Send a frame of data to the next filter.
Definition: avfilter.c:1093
#define av_cold
Definition: attributes.h:82
#define av_malloc(s)
static av_cold int uninit(AVCodecContext *avctx)
Definition: crystalhd.c:279
AVOptions.
cl_device_id device_id
The primary device ID of the device.
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
int64_t pts
Presentation timestamp in time_base units (time when frame should be shown to user).
Definition: frame.h:388
#define height
int ff_opencl_filter_config_output(AVFilterLink *outlink)
Create a suitable hardware frames context for the output.
Definition: opencl.c:96
#define FLAGS
#define av_log(a,...)
A filter pad used for either input or output.
Definition: internal.h:54
#define CL_RELEASE_QUEUE(q)
release an OpenCL Command Queue
Definition: opencl.h:127
#define i(width, name, range_min, range_max)
Definition: cbs_h2645.c:259
int width
Definition: frame.h:353
#define AV_LOG_ERROR
Something went wrong and cannot losslessly be recovered.
Definition: log.h:176
uint8_t log2_chroma_h
Amount to shift the luma height right to find the chroma height.
Definition: pixdesc.h:101
void av_frame_free(AVFrame **frame)
Free the frame and any dynamically allocated objects in it, e.g.
Definition: frame.c:202
#define NULL_IF_CONFIG_SMALL(x)
Return NULL if CONFIG_SMALL is true, otherwise the argument without modification. ...
Definition: internal.h:186
const char * r
Definition: vf_curves.c:114
void * priv
private data for use by the filter
Definition: avfilter.h:353
#define AV_LOG_DEBUG
Stuff which is only useful for libav* developers.
Definition: log.h:197
#define zero
Definition: regdef.h:64
simple assert() macros that are a bit more flexible than ISO C assert().
#define fail()
Definition: checkasm.h:122
static const AVFilterPad nlmeans_opencl_inputs[]
#define width
uint8_t w
Definition: llviddspenc.c:38
AVFormatContext * ctx
Definition: movenc.c:48
AVFilter ff_vf_nlmeans_opencl
static const AVFilterPad outputs[]
Definition: af_acontrast.c:203
#define FF_ARRAY_ELEMS(a)
int format
format of the frame, -1 if unknown or unset Values correspond to enum AVPixelFormat for video frames...
Definition: frame.h:368
#define CL_SET_KERNEL_ARG(kernel, arg_num, type, arg)
set argument to specific Kernel.
Definition: opencl.h:61
Descriptor that unambiguously describes how the bits of a pixel are stored in the up to 4 data planes...
Definition: pixdesc.h:81
uint8_t * data
The data buffer.
Definition: buffer.h:89
static av_cold void nlmeans_opencl_uninit(AVFilterContext *avctx)
these buffered frames must be flushed immediately if a new input produces new the filter must not call request_frame to get more It must just process the frame or queue it The task of requesting more frames is left to the filter s request_frame method or the application If a filter has several inputs
Filter definition.
Definition: avfilter.h:144
This struct describes a set or pool of "hardware" frames (i.e.
Definition: hwcontext.h:123
const char * name
Filter name.
Definition: avfilter.h:148
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
static int nlmeans_plane(AVFilterContext *avctx, cl_mem dst, cl_mem src, cl_int width, cl_int height, cl_int p, cl_int r)
AVFilterLink ** outputs
array of pointers to output links
Definition: avfilter.h:350
static int nlmeans_opencl_filter_frame(AVFilterLink *inlink, AVFrame *input)
#define CL_FAIL_ON_ERROR(errcode,...)
A helper macro to handle OpenCL errors.
Definition: opencl.h:74
uint8_t * data[AV_NUM_DATA_POINTERS]
pointer to the picture/channel planes.
Definition: frame.h:309
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
const char * ff_opencl_source_nlmeans
planar YUV 4:2:0, 12bpp, (1 Cr & Cb sample per 2x2 Y samples)
Definition: pixfmt.h:66
static int query_formats(AVFilterContext *ctx)
Definition: aeval.c:244
common internal and external API header
void ff_opencl_filter_uninit(AVFilterContext *avctx)
Uninitialise an OpenCL filter context.
Definition: opencl.c:156
#define CL_RELEASE_KERNEL(k)
release an OpenCL Kernel
Definition: opencl.h:101
#define OFFSET(x)
static enum AVPixelFormat supported_formats[]
cl_context context
The OpenCL context which will contain all operations and frames on this device.
An instance of a filter.
Definition: avfilter.h:338
int height
Definition: frame.h:353
#define av_freep(p)
cl_program program
Definition: opencl.h:43
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:171
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:2438
internal API functions
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
enum AVPixelFormat sw_format
The pixel format identifying the actual data layout of the hardware frames.
Definition: hwcontext.h:221
AVPixelFormat
Pixel format.
Definition: pixfmt.h:64
static int nlmeans_opencl_init(AVFilterContext *avctx, int width, int height)
int av_frame_copy_props(AVFrame *dst, const AVFrame *src)
Copy only "metadata" fields from src to dst.
Definition: frame.c:654
#define AV_CEIL_RSHIFT(a, b)
Definition: common.h:58
static const AVOption nlmeans_opencl_options[]
static uint8_t tmp[11]
Definition: aes_ctr.c:26