FFmpeg: libavfilter/vf_xfade_opencl.c Source File

FFmpeg
vf_xfade_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 
19 #include "libavutil/log.h"
20 #include "libavutil/opt.h"
21 #include "libavutil/pixdesc.h"
22 
23 #include "avfilter.h"
24 #include "filters.h"
25 #include "opencl.h"
26 #include "opencl_source.h"
27 #include "video.h"
28 
29  enum XFadeTransitions {
30   CUSTOM,
31   FADE,
32   WIPELEFT,
33   WIPERIGHT,
34   WIPEUP,
35   WIPEDOWN,
36   SLIDELEFT,
37   SLIDERIGHT,
38   SLIDEUP,
39   SLIDEDOWN,
40   NB_TRANSITIONS,
41 };
42 
43  typedef struct XFadeOpenCLContext {
44   OpenCLFilterContext ocf;
45 
46   int transition;
47   const char *source_file;
48   const char *kernel_name;
49   int64_t duration;
50   int64_t offset;
51 
52   int initialised;
53   cl_kernel kernel;
54   cl_command_queue command_queue;
55 
56   int nb_planes;
57 
58   int64_t duration_pts;
59   int64_t offset_pts;
60   int64_t first_pts;
61   int64_t last_pts;
62   int64_t pts;
63   int xfade_is_over;
64   int need_second;
65   int eof[2];
66   AVFrame *xf[2];
67 } XFadeOpenCLContext;
68 
69  static int xfade_opencl_load(AVFilterContext *avctx,
70  enum AVPixelFormat main_format,
71  enum AVPixelFormat xfade_format)
72 {
73  XFadeOpenCLContext *ctx = avctx->priv;
74  cl_int cle;
75  const AVPixFmtDescriptor *main_desc;
76  int err, main_planes;
77  const char *kernel_name;
78 
79  main_desc = av_pix_fmt_desc_get(main_format);
80  if (main_format != xfade_format) {
81  av_log(avctx, AV_LOG_ERROR, "Input formats are not same.\n");
82  return AVERROR(EINVAL);
83  }
84 
85  main_planes = 0;
86  for (int i = 0; i < main_desc->nb_components; i++)
87  main_planes = FFMAX(main_planes,
88  main_desc->comp[i].plane + 1);
89 
90  ctx->nb_planes = main_planes;
91 
92  if (ctx->transition == CUSTOM) {
93  err = ff_opencl_filter_load_program_from_file(avctx, ctx->source_file);
94  } else {
95  err = ff_opencl_filter_load_program(avctx, &ff_source_xfade_cl, 1);
96  }
97  if (err < 0)
98  return err;
99 
100  ctx->command_queue = clCreateCommandQueue(ctx->ocf.hwctx->context,
101  ctx->ocf.hwctx->device_id,
102  0, &cle);
103  CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create OpenCL "
104  "command queue %d.\n", cle);
105 
106  switch (ctx->transition) {
107  case CUSTOM: kernel_name = ctx->kernel_name; break;
108  case FADE: kernel_name = "fade"; break;
109  case WIPELEFT: kernel_name = "wipeleft"; break;
110  case WIPERIGHT: kernel_name = "wiperight"; break;
111  case WIPEUP: kernel_name = "wipeup"; break;
112  case WIPEDOWN: kernel_name = "wipedown"; break;
113  case SLIDELEFT: kernel_name = "slideleft"; break;
114  case SLIDERIGHT: kernel_name = "slideright"; break;
115  case SLIDEUP: kernel_name = "slideup"; break;
116  case SLIDEDOWN: kernel_name = "slidedown"; break;
117  default:
118  err = AVERROR_BUG;
119  goto fail;
120  }
121 
122  ctx->kernel = clCreateKernel(ctx->ocf.program, kernel_name, &cle);
123  CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create kernel %d.\n", cle);
124 
125  ctx->initialised = 1;
126 
127  return 0;
128 
129 fail:
130  if (ctx->command_queue)
131  clReleaseCommandQueue(ctx->command_queue);
132  if (ctx->kernel)
133  clReleaseKernel(ctx->kernel);
134  return err;
135 }
136 
137  static int xfade_frame(AVFilterContext *avctx, AVFrame *a, AVFrame *b)
138 {
139  AVFilterLink *outlink = avctx->outputs[0];
140  XFadeOpenCLContext *ctx = avctx->priv;
141  AVFrame *output;
142  cl_int cle;
143  cl_float progress = av_clipf(1.f - ((cl_float)(ctx->pts - ctx->first_pts - ctx->offset_pts) / ctx->duration_pts), 0.f, 1.f);
144  size_t global_work[2];
145  int kernel_arg = 0;
146  int err, plane;
147 
148  if (!ctx->initialised) {
149  AVHWFramesContext *main_fc =
150  (AVHWFramesContext*)a->hw_frames_ctx->data;
151  AVHWFramesContext *xfade_fc =
152  (AVHWFramesContext*)b->hw_frames_ctx->data;
153 
154  err = xfade_opencl_load(avctx, main_fc->sw_format,
155  xfade_fc->sw_format);
156  if (err < 0)
157  return err;
158  }
159 
160  output = ff_get_video_buffer(outlink, outlink->w, outlink->h);
161  if (!output) {
162  err = AVERROR(ENOMEM);
163  goto fail;
164  }
165 
166  for (plane = 0; plane < ctx->nb_planes; plane++) {
167  cl_mem mem;
168  kernel_arg = 0;
169 
170  mem = (cl_mem)output->data[plane];
171  CL_SET_KERNEL_ARG(ctx->kernel, kernel_arg, cl_mem, &mem);
172  kernel_arg++;
173 
174  mem = (cl_mem)ctx->xf[0]->data[plane];
175  CL_SET_KERNEL_ARG(ctx->kernel, kernel_arg, cl_mem, &mem);
176  kernel_arg++;
177 
178  mem = (cl_mem)ctx->xf[1]->data[plane];
179  CL_SET_KERNEL_ARG(ctx->kernel, kernel_arg, cl_mem, &mem);
180  kernel_arg++;
181 
182  CL_SET_KERNEL_ARG(ctx->kernel, kernel_arg, cl_float, &progress);
183  kernel_arg++;
184 
185  err = ff_opencl_filter_work_size_from_image(avctx, global_work,
186  output, plane, 0);
187  if (err < 0)
188  goto fail;
189 
190  cle = clEnqueueNDRangeKernel(ctx->command_queue, ctx->kernel, 2, NULL,
191  global_work, NULL, 0, NULL, NULL);
192  CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to enqueue xfade kernel "
193  "for plane %d: %d.\n", plane, cle);
194  }
195 
196  cle = clFinish(ctx->command_queue);
197  CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to finish command queue: %d.\n", cle);
198 
199  err = av_frame_copy_props(output, ctx->xf[0]);
200  if (err < 0)
201  goto fail;
202 
203  output->pts = ctx->pts;
204 
205  return ff_filter_frame(outlink, output);
206 
207 fail:
208  av_frame_free(&output);
209  return err;
210 }
211 
212  static int xfade_opencl_config_output(AVFilterLink *outlink)
213 {
214  AVFilterContext *avctx = outlink->src;
215  XFadeOpenCLContext *ctx = avctx->priv;
216  AVFilterLink *inlink0 = avctx->inputs[0];
217  AVFilterLink *inlink1 = avctx->inputs[1];
218  FilterLink *il = ff_filter_link(inlink0);
219  FilterLink *ol = ff_filter_link(outlink);
220  int err;
221 
222  err = ff_opencl_filter_config_output(outlink);
223  if (err < 0)
224  return err;
225 
226  if (inlink0->w != inlink1->w || inlink0->h != inlink1->h) {
227  av_log(avctx, AV_LOG_ERROR, "First input link %s parameters "
228  "(size %dx%d) do not match the corresponding "
229  "second input link %s parameters (size %dx%d)\n",
230  avctx->input_pads[0].name, inlink0->w, inlink0->h,
231  avctx->input_pads[1].name, inlink1->w, inlink1->h);
232  return AVERROR(EINVAL);
233  }
234 
235  if (inlink0->time_base.num != inlink1->time_base.num ||
236  inlink0->time_base.den != inlink1->time_base.den) {
237  av_log(avctx, AV_LOG_ERROR, "First input link %s timebase "
238  "(%d/%d) do not match the corresponding "
239  "second input link %s timebase (%d/%d)\n",
240  avctx->input_pads[0].name, inlink0->time_base.num, inlink0->time_base.den,
241  avctx->input_pads[1].name, inlink1->time_base.num, inlink1->time_base.den);
242  return AVERROR(EINVAL);
243  }
244 
245  ctx->first_pts = ctx->last_pts = ctx->pts = AV_NOPTS_VALUE;
246 
247  outlink->time_base = inlink0->time_base;
248  outlink->sample_aspect_ratio = inlink0->sample_aspect_ratio;
249  ol->frame_rate = il->frame_rate;
250 
251  if (ctx->duration)
252  ctx->duration_pts = av_rescale_q(ctx->duration, AV_TIME_BASE_Q, outlink->time_base);
253  if (ctx->offset)
254  ctx->offset_pts = av_rescale_q(ctx->offset, AV_TIME_BASE_Q, outlink->time_base);
255 
256  return 0;
257 }
258 
259  static int xfade_opencl_activate(AVFilterContext *avctx)
260 {
261  XFadeOpenCLContext *ctx = avctx->priv;
262  AVFilterLink *outlink = avctx->outputs[0];
263  AVFrame *in = NULL;
264  int ret = 0, status;
265  int64_t pts;
266 
267  FF_FILTER_FORWARD_STATUS_BACK_ALL(outlink, avctx);
268 
269  if (ctx->xfade_is_over) {
270  ret = ff_inlink_consume_frame(avctx->inputs[1], &in);
271  if (ret < 0) {
272  return ret;
273  } else if (ret > 0) {
274  in->pts = (in->pts - ctx->last_pts) + ctx->pts;
275  return ff_filter_frame(outlink, in);
276  } else if (ff_inlink_acknowledge_status(avctx->inputs[1], &status, &pts)) {
277  ff_outlink_set_status(outlink, status, ctx->pts);
278  return 0;
279  } else if (!ret) {
280  if (ff_outlink_frame_wanted(outlink)) {
281  ff_inlink_request_frame(avctx->inputs[1]);
282  return 0;
283  }
284  }
285  }
286 
287  if (ff_inlink_queued_frames(avctx->inputs[0]) > 0) {
288  ctx->xf[0] = ff_inlink_peek_frame(avctx->inputs[0], 0);
289  if (ctx->xf[0]) {
290  if (ctx->first_pts == AV_NOPTS_VALUE) {
291  ctx->first_pts = ctx->xf[0]->pts;
292  }
293  ctx->pts = ctx->xf[0]->pts;
294  if (ctx->first_pts + ctx->offset_pts > ctx->xf[0]->pts) {
295  ctx->xf[0] = NULL;
296  ctx->need_second = 0;
297  ret = ff_inlink_consume_frame(avctx->inputs[0], &in);
298  if (ret < 0)
299  return ret;
300  return ff_filter_frame(outlink, in);
301  }
302 
303  ctx->need_second = 1;
304  }
305  }
306 
307  if (ctx->xf[0] && ff_inlink_queued_frames(avctx->inputs[1]) > 0) {
308  ret = ff_inlink_consume_frame(avctx->inputs[0], &ctx->xf[0]);
309  if (ret < 0)
310  return ret;
311  ret = ff_inlink_consume_frame(avctx->inputs[1], &ctx->xf[1]);
312  if (ret < 0) {
313  av_frame_free(&ctx->xf[0]);
314  return ret;
315  }
316 
317  ctx->last_pts = ctx->xf[1]->pts;
318  ctx->pts = ctx->xf[0]->pts;
319  if (ctx->xf[0]->pts - (ctx->first_pts + ctx->offset_pts) > ctx->duration_pts)
320  ctx->xfade_is_over = 1;
321  ret = xfade_frame(avctx, ctx->xf[0], ctx->xf[1]);
322  av_frame_free(&ctx->xf[0]);
323  av_frame_free(&ctx->xf[1]);
324  return ret;
325  }
326 
327  if (ff_inlink_queued_frames(avctx->inputs[0]) > 0 &&
328  ff_inlink_queued_frames(avctx->inputs[1]) > 0) {
329  ff_filter_set_ready(avctx, 100);
330  return 0;
331  }
332 
333  if (ff_outlink_frame_wanted(outlink)) {
334  if (!ctx->eof[0] && ff_outlink_get_status(avctx->inputs[0])) {
335  ctx->eof[0] = 1;
336  ctx->xfade_is_over = 1;
337  }
338  if (!ctx->eof[1] && ff_outlink_get_status(avctx->inputs[1])) {
339  ctx->eof[1] = 1;
340  }
341  if (!ctx->eof[0] && !ctx->xf[0])
342  ff_inlink_request_frame(avctx->inputs[0]);
343  if (!ctx->eof[1] && (ctx->need_second || ctx->eof[0]))
344  ff_inlink_request_frame(avctx->inputs[1]);
345  if (ctx->eof[0] && ctx->eof[1] && (
346  ff_inlink_queued_frames(avctx->inputs[0]) <= 0 ||
347  ff_inlink_queued_frames(avctx->inputs[1]) <= 0))
348  ff_outlink_set_status(outlink, AVERROR_EOF, AV_NOPTS_VALUE);
349  return 0;
350  }
351 
352  return FFERROR_NOT_READY;
353 }
354 
355  static av_cold void xfade_opencl_uninit(AVFilterContext *avctx)
356 {
357  XFadeOpenCLContext *ctx = avctx->priv;
358  cl_int cle;
359 
360  if (ctx->kernel) {
361  cle = clReleaseKernel(ctx->kernel);
362  if (cle != CL_SUCCESS)
363  av_log(avctx, AV_LOG_ERROR, "Failed to release "
364  "kernel: %d.\n", cle);
365  }
366 
367  if (ctx->command_queue) {
368  cle = clReleaseCommandQueue(ctx->command_queue);
369  if (cle != CL_SUCCESS)
370  av_log(avctx, AV_LOG_ERROR, "Failed to release "
371  "command queue: %d.\n", cle);
372  }
373 
374  ff_opencl_filter_uninit(avctx);
375 }
376 
377  static AVFrame *get_video_buffer(AVFilterLink *inlink, int w, int h)
378 {
379  XFadeOpenCLContext *s = inlink->dst->priv;
380 
381  return s->xfade_is_over || !s->need_second ?
382  ff_null_get_video_buffer (inlink, w, h) :
383  ff_default_get_video_buffer(inlink, w, h);
384 }
385 
386  #define OFFSET(x) offsetof(XFadeOpenCLContext, x)
387  #define FLAGS (AV_OPT_FLAG_FILTERING_PARAM | AV_OPT_FLAG_VIDEO_PARAM)
388 
389  static const AVOption xfade_opencl_options[] = {
390  { "transition", "set cross fade transition", OFFSET(transition), AV_OPT_TYPE_INT, {.i64=1}, 0, NB_TRANSITIONS-1, FLAGS, .unit = "transition" },
391  { "custom", "custom transition", 0, AV_OPT_TYPE_CONST, {.i64=CUSTOM}, 0, 0, FLAGS, .unit = "transition" },
392  { "fade", "fade transition", 0, AV_OPT_TYPE_CONST, {.i64=FADE}, 0, 0, FLAGS, .unit = "transition" },
393  { "wipeleft", "wipe left transition", 0, AV_OPT_TYPE_CONST, {.i64=WIPELEFT}, 0, 0, FLAGS, .unit = "transition" },
394  { "wiperight", "wipe right transition", 0, AV_OPT_TYPE_CONST, {.i64=WIPERIGHT}, 0, 0, FLAGS, .unit = "transition" },
395  { "wipeup", "wipe up transition", 0, AV_OPT_TYPE_CONST, {.i64=WIPEUP}, 0, 0, FLAGS, .unit = "transition" },
396  { "wipedown", "wipe down transition", 0, AV_OPT_TYPE_CONST, {.i64=WIPEDOWN}, 0, 0, FLAGS, .unit = "transition" },
397  { "slideleft", "slide left transition", 0, AV_OPT_TYPE_CONST, {.i64=SLIDELEFT}, 0, 0, FLAGS, .unit = "transition" },
398  { "slideright", "slide right transition", 0, AV_OPT_TYPE_CONST, {.i64=SLIDERIGHT}, 0, 0, FLAGS, .unit = "transition" },
399  { "slideup", "slide up transition", 0, AV_OPT_TYPE_CONST, {.i64=SLIDEUP}, 0, 0, FLAGS, .unit = "transition" },
400  { "slidedown", "slide down transition", 0, AV_OPT_TYPE_CONST, {.i64=SLIDEDOWN}, 0, 0, FLAGS, .unit = "transition" },
401  { "source", "set OpenCL program source file for custom transition", OFFSET(source_file), AV_OPT_TYPE_STRING, {.str = NULL}, .flags = FLAGS },
402  { "kernel", "set kernel name in program file for custom transition", OFFSET(kernel_name), AV_OPT_TYPE_STRING, {.str = NULL}, .flags = FLAGS },
403  { "duration", "set cross fade duration", OFFSET(duration), AV_OPT_TYPE_DURATION, {.i64=1000000}, 0, 60000000, FLAGS },
404  { "offset", "set cross fade start relative to first input stream", OFFSET(offset), AV_OPT_TYPE_DURATION, {.i64=0}, INT64_MIN, INT64_MAX, FLAGS },
405  { NULL }
406 };
407 
408 AVFILTER_DEFINE_CLASS(xfade_opencl);
409 
410  static const AVFilterPad xfade_opencl_inputs[] = {
411  {
412  .name = "main",
413  .type = AVMEDIA_TYPE_VIDEO,
414  .get_buffer.video = get_video_buffer,
415  .config_props = &ff_opencl_filter_config_input,
416  },
417  {
418  .name = "xfade",
419  .type = AVMEDIA_TYPE_VIDEO,
420  .get_buffer.video = get_video_buffer,
421  .config_props = &ff_opencl_filter_config_input,
422  },
423 };
424 
425  static const AVFilterPad xfade_opencl_outputs[] = {
426  {
427  .name = "default",
428  .type = AVMEDIA_TYPE_VIDEO,
429  .config_props = &xfade_opencl_config_output,
430  },
431 };
432 
433  const FFFilter ff_vf_xfade_opencl = {
434  .p.name = "xfade_opencl",
435  .p.description = NULL_IF_CONFIG_SMALL("Cross fade one video with another video."),
436  .p.priv_class = &xfade_opencl_class,
437  .p.flags = AVFILTER_FLAG_HWDEVICE,
438  .priv_size = sizeof(XFadeOpenCLContext),
439  .init = &ff_opencl_filter_init,
440  .uninit = &xfade_opencl_uninit,
441  .activate = &xfade_opencl_activate,
442  FILTER_INPUTS(xfade_opencl_inputs),
443  FILTER_OUTPUTS(xfade_opencl_outputs),
444  FILTER_SINGLE_PIXFMT(AV_PIX_FMT_OPENCL),
445  .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE,
446 };
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
AVPixelFormat
AVPixelFormat
Pixel format.
Definition: pixfmt.h:71
FLAGS
#define FLAGS
Definition: vf_xfade_opencl.c:387
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
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_frame
int ff_filter_frame(AVFilterLink *link, AVFrame *frame)
Send a frame of data to the next filter.
Definition: avfilter.c:1067
av_pix_fmt_desc_get
const AVPixFmtDescriptor * av_pix_fmt_desc_get(enum AVPixelFormat pix_fmt)
Definition: pixdesc.c:3447
AVERROR_EOF
#define AVERROR_EOF
End of file.
Definition: error.h:57
FFERROR_NOT_READY
return FFERROR_NOT_READY
Definition: filter_design.txt:204
AV_TIME_BASE_Q
#define AV_TIME_BASE_Q
Internal time base represented as fractional value.
Definition: avutil.h:263
int64_t
long long int64_t
Definition: coverity.c:34
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:226
ff_vf_xfade_opencl
const FFFilter ff_vf_xfade_opencl
Definition: vf_xfade_opencl.c:433
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: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::pts
int64_t pts
Presentation timestamp in time_base units (time when frame should be shown to user).
Definition: frame.h:529
w
uint8_t w
Definition: llviddspenc.c:38
opencl.h
AVOption
AVOption.
Definition: opt.h:429
b
#define b
Definition: input.c:42
xfade_opencl_activate
static int xfade_opencl_activate(AVFilterContext *avctx)
Definition: vf_xfade_opencl.c:259
AV_OPT_TYPE_DURATION
@ AV_OPT_TYPE_DURATION
Underlying C type is int64_t.
Definition: opt.h:319
ff_source_xfade_cl
const char * ff_source_xfade_cl
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:159
SLIDEUP
@ SLIDEUP
Definition: vf_xfade_opencl.c:38
FFMAX
#define FFMAX(a, b)
Definition: macros.h:47
AVFilter::name
const char * name
Filter name.
Definition: avfilter.h:220
SLIDELEFT
@ SLIDELEFT
Definition: vf_xfade_opencl.c:36
video.h
XFadeOpenCLContext::first_pts
int64_t first_pts
Definition: vf_xfade_opencl.c:60
XFadeOpenCLContext::xfade_is_over
int xfade_is_over
Definition: vf_xfade_opencl.c:63
ff_inlink_consume_frame
int ff_inlink_consume_frame(AVFilterLink *link, AVFrame **rframe)
Take a frame from the link's FIFO and update the link's stats.
Definition: avfilter.c:1517
ff_default_get_video_buffer
AVFrame * ff_default_get_video_buffer(AVFilterLink *link, int w, int h)
Definition: video.c:112
ff_opencl_filter_work_size_from_image
int ff_opencl_filter_work_size_from_image(AVFilterContext *avctx, size_t *work_size, AVFrame *frame, int plane, int block_alignment)
Find the work size needed needed for a given plane of an image.
Definition: opencl.c:266
FF_FILTER_FORWARD_STATUS_BACK_ALL
#define FF_FILTER_FORWARD_STATUS_BACK_ALL(outlink, filter)
Forward the status on an output link to all input links.
Definition: filters.h:651
XFadeOpenCLContext::xf
AVFrame * xf[2]
Definition: vf_xfade_opencl.c:66
AVFilterContext::priv
void * priv
private data for use by the filter
Definition: avfilter.h:289
XFadeOpenCLContext::offset
int64_t offset
Definition: vf_xfade_opencl.c:50
fail
#define fail()
Definition: checkasm.h:206
xfade_opencl_uninit
static av_cold void xfade_opencl_uninit(AVFilterContext *avctx)
Definition: vf_xfade_opencl.c:355
FADE
@ FADE
Definition: vf_xfade_opencl.c:31
XFadeOpenCLContext::need_second
int need_second
Definition: vf_xfade_opencl.c:64
pts
static int64_t pts
Definition: transcode_aac.c:644
XFadeOpenCLContext
Definition: vf_xfade_opencl.c:43
AVRational::num
int num
Numerator.
Definition: rational.h:59
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:83
AVFilterPad
A filter pad used for either input or output.
Definition: filters.h:39
AVFilterContext::input_pads
AVFilterPad * input_pads
array of input pads
Definition: avfilter.h:281
AV_LOG_ERROR
#define AV_LOG_ERROR
Something went wrong and cannot losslessly be recovered.
Definition: log.h:210
av_cold
#define av_cold
Definition: attributes.h:106
SLIDEDOWN
@ SLIDEDOWN
Definition: vf_xfade_opencl.c:39
FFFilter
Definition: filters.h:266
duration
int64_t duration
Definition: movenc.c:65
ff_outlink_set_status
static void ff_outlink_set_status(AVFilterLink *link, int status, int64_t pts)
Set the status field of a link from the source filter.
Definition: filters.h:628
ff_inlink_request_frame
void ff_inlink_request_frame(AVFilterLink *link)
Mark that a frame is wanted on the link.
Definition: avfilter.c:1620
s
#define s(width, name)
Definition: cbs_vp9.c:198
filters.h
WIPERIGHT
@ WIPERIGHT
Definition: vf_xfade_opencl.c:33
ctx
AVFormatContext * ctx
Definition: movenc.c:49
av_rescale_q
int64_t av_rescale_q(int64_t a, AVRational bq, AVRational cq)
Rescale a 64-bit integer by 2 rational numbers.
Definition: mathematics.c:142
CUSTOM
@ CUSTOM
Definition: vf_xfade_opencl.c:30
FILTER_OUTPUTS
#define FILTER_OUTPUTS(array)
Definition: filters.h:264
OFFSET
#define OFFSET(x)
Definition: vf_xfade_opencl.c:386
XFadeOpenCLContext::kernel_name
const char * kernel_name
Definition: vf_xfade_opencl.c:48
ff_inlink_peek_frame
AVFrame * ff_inlink_peek_frame(AVFilterLink *link, size_t idx)
Access a frame in the link fifo without consuming it.
Definition: avfilter.c:1558
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
AVPixFmtDescriptor::nb_components
uint8_t nb_components
The number of components each pixel has, (1-4)
Definition: pixdesc.h:71
activate
filter_frame For filters that do not use the activate() callback
AVFilterContext::inputs
AVFilterLink ** inputs
array of pointers to input links
Definition: avfilter.h:282
AVComponentDescriptor::plane
int plane
Which of the 4 planes contains the component.
Definition: pixdesc.h:34
SLIDERIGHT
@ SLIDERIGHT
Definition: vf_xfade_opencl.c:37
AV_PIX_FMT_OPENCL
@ AV_PIX_FMT_OPENCL
Hardware surfaces for OpenCL.
Definition: pixfmt.h:358
av_clipf
av_clipf
Definition: af_crystalizer.c:122
ff_inlink_acknowledge_status
int ff_inlink_acknowledge_status(AVFilterLink *link, int *rstatus, int64_t *rpts)
Test and acknowledge the change of status on the link.
Definition: avfilter.c:1464
ff_inlink_queued_frames
size_t ff_inlink_queued_frames(AVFilterLink *link)
Get the number of frames available on the link.
Definition: avfilter.c:1480
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
f
f
Definition: af_crystalizer.c:122
init
int(* init)(AVBSFContext *ctx)
Definition: dts2pts.c:368
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
XFadeOpenCLContext::initialised
int initialised
Definition: vf_xfade_opencl.c:52
XFadeOpenCLContext::last_pts
int64_t last_pts
Definition: vf_xfade_opencl.c:61
AVFILTER_FLAG_HWDEVICE
#define AVFILTER_FLAG_HWDEVICE
The filter can create hardware frames using AVFilterContext.hw_device_ctx.
Definition: avfilter.h:188
NB_TRANSITIONS
@ NB_TRANSITIONS
Definition: vf_xfade_opencl.c:40
AV_NOPTS_VALUE
#define AV_NOPTS_VALUE
Undefined timestamp value.
Definition: avutil.h:247
a
The reader does not expect b to be semantically here and if the code is changed by maybe adding a a division or other the signedness will almost certainly be mistaken To avoid this confusion a new type was SUINT is the C unsigned type but it holds a signed int to use the same example SUINT a
Definition: undefined.txt:41
xfade_opencl_outputs
static const AVFilterPad xfade_opencl_outputs[]
Definition: vf_xfade_opencl.c:425
offset
it s the only field you need to keep assuming you have a context There is some magic you don t need to care about around this just let it vf offset
Definition: writing_filters.txt:86
opencl_source.h
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:46
ff_null_get_video_buffer
AVFrame * ff_null_get_video_buffer(AVFilterLink *link, int w, int h)
Definition: video.c:44
xfade_opencl_load
static int xfade_opencl_load(AVFilterContext *avctx, enum AVPixelFormat main_format, enum AVPixelFormat xfade_format)
Definition: vf_xfade_opencl.c:69
uninit
static void uninit(AVBSFContext *ctx)
Definition: pcm_rechunk.c:68
log.h
i
#define i(width, name, range_min, range_max)
Definition: cbs_h2645.c:256
WIPEDOWN
@ WIPEDOWN
Definition: vf_xfade_opencl.c:35
xfade_opencl_options
static const AVOption xfade_opencl_options[]
Definition: vf_xfade_opencl.c:389
XFadeOpenCLContext::eof
int eof[2]
Definition: vf_xfade_opencl.c:65
XFadeOpenCLContext::transition
int transition
Definition: vf_xfade_opencl.c:46
AVFilterPad::name
const char * name
Pad name.
Definition: filters.h:45
AVHWFramesContext
This struct describes a set or pool of "hardware" frames (i.e.
Definition: hwcontext.h:118
ff_opencl_filter_init
int ff_opencl_filter_init(AVFilterContext *avctx)
Initialise an OpenCL filter context.
Definition: opencl.c:135
ret
ret
Definition: filter_design.txt:187
XFadeOpenCLContext::pts
int64_t pts
Definition: vf_xfade_opencl.c:62
XFadeOpenCLContext::kernel
cl_kernel kernel
Definition: vf_xfade_opencl.c:53
XFadeOpenCLContext::ocf
OpenCLFilterContext ocf
Definition: vf_xfade_opencl.c:44
status
ov_status_e status
Definition: dnn_backend_openvino.c:100
XFadeOpenCLContext::source_file
const char * source_file
Definition: vf_xfade_opencl.c:47
AVRational::den
int den
Denominator.
Definition: rational.h:60
AV_OPT_TYPE_INT
@ AV_OPT_TYPE_INT
Underlying C type is int.
Definition: opt.h:259
avfilter.h
ff_opencl_filter_load_program_from_file
int ff_opencl_filter_load_program_from_file(AVFilterContext *avctx, const char *filename)
Load a new OpenCL program from a file.
Definition: opencl.c:207
XFadeOpenCLContext::duration
int64_t duration
Definition: vf_xfade_opencl.c:49
AVPixFmtDescriptor::comp
AVComponentDescriptor comp[4]
Parameters that describe how pixels are packed.
Definition: pixdesc.h:105
OpenCLFilterContext
Definition: opencl.h:36
ff_opencl_filter_uninit
void ff_opencl_filter_uninit(AVFilterContext *avctx)
Uninitialise an OpenCL filter context.
Definition: opencl.c:144
AVFormatContext::duration
int64_t duration
Duration of the stream, in AV_TIME_BASE fractional seconds.
Definition: avformat.h:1399
ff_outlink_get_status
int ff_outlink_get_status(AVFilterLink *link)
Get the status on an output link.
Definition: avfilter.c:1645
AVFilterContext
An instance of a filter.
Definition: avfilter.h:274
get_video_buffer
static AVFrame * get_video_buffer(AVFilterLink *inlink, int w, int h)
Definition: vf_xfade_opencl.c:377
AVMEDIA_TYPE_VIDEO
@ AVMEDIA_TYPE_VIDEO
Definition: avutil.h:200
FFFilter::p
AVFilter p
The public AVFilter.
Definition: filters.h:270
XFadeOpenCLContext::command_queue
cl_command_queue command_queue
Definition: vf_xfade_opencl.c:54
XFadeOpenCLContext::duration_pts
int64_t duration_pts
Definition: vf_xfade_opencl.c:58
AVPixFmtDescriptor
Descriptor that unambiguously describes how the bits of a pixel are stored in the up to 4 data planes...
Definition: pixdesc.h:69
XFadeTransitions
XFadeTransitions
Definition: vf_xfade.c:29
AVERROR_BUG
#define AVERROR_BUG
Internal bug, also see AVERROR_BUG2.
Definition: error.h:52
av_log
#define av_log(a,...)
Definition: tableprint_vlc.h:27
XFadeOpenCLContext::nb_planes
int nb_planes
Definition: vf_xfade_opencl.c:56
CL_FAIL_ON_ERROR
#define CL_FAIL_ON_ERROR(errcode,...)
A helper macro to handle OpenCL errors.
Definition: opencl.h:74
xfade_frame
static int xfade_frame(AVFilterContext *avctx, AVFrame *a, AVFrame *b)
Definition: vf_xfade_opencl.c:137
h
h
Definition: vp9dsp_template.c:2070
ff_outlink_frame_wanted
the definition of that something depends on the semantic of the filter The callback must examine the status of the filter s links and proceed accordingly The status of output links is stored in the status_in and status_out fields and tested by the ff_outlink_frame_wanted() function. If this function returns true
XFadeOpenCLContext::offset_pts
int64_t offset_pts
Definition: vf_xfade_opencl.c:59
AV_OPT_TYPE_STRING
@ AV_OPT_TYPE_STRING
Underlying C type is a uint8_t* that is either NULL or points to a C string allocated with the av_mal...
Definition: opt.h:276
xfade_opencl_config_output
static int xfade_opencl_config_output(AVFilterLink *outlink)
Definition: vf_xfade_opencl.c:212
AVFILTER_DEFINE_CLASS
AVFILTER_DEFINE_CLASS(xfade_opencl)
WIPELEFT
@ WIPELEFT
Definition: vf_xfade_opencl.c:32
AV_OPT_TYPE_CONST
@ AV_OPT_TYPE_CONST
Special option type for declaring named constants.
Definition: opt.h:299
FILTER_SINGLE_PIXFMT
#define FILTER_SINGLE_PIXFMT(pix_fmt_)
Definition: filters.h:253
WIPEUP
@ WIPEUP
Definition: vf_xfade_opencl.c:34
ff_filter_set_ready
void ff_filter_set_ready(AVFilterContext *filter, unsigned priority)
Mark a filter ready and schedule it for activation.
Definition: avfilter.c:229
xfade_opencl_inputs
static const AVFilterPad xfade_opencl_inputs[]
Definition: vf_xfade_opencl.c:410
AVFilterContext::outputs
AVFilterLink ** outputs
array of pointers to output links
Definition: avfilter.h:286
Generated on Sat Oct 18 2025 19:23:14 for FFmpeg by   doxygen 1.8.17

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