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 "internal.h"
26 #include "opencl.h"
27 #include "opencl_source.h"
28 #include "video.h"
29 
42 };
43 
44 typedef struct XFadeOpenCLContext {
46 
48  const char *source_file;
49  const char *kernel_name;
50  int64_t duration;
51  int64_t offset;
52 
54  cl_kernel kernel;
55  cl_command_queue command_queue;
56 
57  int nb_planes;
58 
59  int64_t duration_pts;
60  int64_t offset_pts;
61  int64_t first_pts;
62  int64_t last_pts;
63  int64_t pts;
66  int eof[2];
67  AVFrame *xf[2];
69 
71  enum AVPixelFormat main_format,
72  enum AVPixelFormat xfade_format)
73 {
74  XFadeOpenCLContext *ctx = avctx->priv;
75  cl_int cle;
76  const AVPixFmtDescriptor *main_desc;
77  int err, main_planes;
78  const char *kernel_name;
79 
80  main_desc = av_pix_fmt_desc_get(main_format);
81  if (main_format != xfade_format) {
82  av_log(avctx, AV_LOG_ERROR, "Input formats are not same.\n");
83  return AVERROR(EINVAL);
84  }
85 
86  main_planes = 0;
87  for (int i = 0; i < main_desc->nb_components; i++)
88  main_planes = FFMAX(main_planes,
89  main_desc->comp[i].plane + 1);
90 
91  ctx->nb_planes = main_planes;
92 
93  if (ctx->transition == CUSTOM) {
94  err = ff_opencl_filter_load_program_from_file(avctx, ctx->source_file);
95  } else {
97  }
98  if (err < 0)
99  return err;
100 
101  ctx->command_queue = clCreateCommandQueue(ctx->ocf.hwctx->context,
102  ctx->ocf.hwctx->device_id,
103  0, &cle);
104  CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create OpenCL "
105  "command queue %d.\n", cle);
106 
107  switch (ctx->transition) {
108  case CUSTOM: kernel_name = ctx->kernel_name; break;
109  case FADE: kernel_name = "fade"; break;
110  case WIPELEFT: kernel_name = "wipeleft"; break;
111  case WIPERIGHT: kernel_name = "wiperight"; break;
112  case WIPEUP: kernel_name = "wipeup"; break;
113  case WIPEDOWN: kernel_name = "wipedown"; break;
114  case SLIDELEFT: kernel_name = "slideleft"; break;
115  case SLIDERIGHT: kernel_name = "slideright"; break;
116  case SLIDEUP: kernel_name = "slideup"; break;
117  case SLIDEDOWN: kernel_name = "slidedown"; break;
118  default:
119  err = AVERROR_BUG;
120  goto fail;
121  }
122 
123  ctx->kernel = clCreateKernel(ctx->ocf.program, kernel_name, &cle);
124  CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create kernel %d.\n", cle);
125 
126  ctx->initialised = 1;
127 
128  return 0;
129 
130 fail:
131  if (ctx->command_queue)
132  clReleaseCommandQueue(ctx->command_queue);
133  if (ctx->kernel)
134  clReleaseKernel(ctx->kernel);
135  return err;
136 }
137 
139 {
140  AVFilterLink *outlink = avctx->outputs[0];
141  XFadeOpenCLContext *ctx = avctx->priv;
142  AVFrame *output;
143  cl_int cle;
144  cl_float progress = av_clipf(1.f - ((cl_float)(ctx->pts - ctx->first_pts - ctx->offset_pts) / ctx->duration_pts), 0.f, 1.f);
145  size_t global_work[2];
146  int kernel_arg = 0;
147  int err, plane;
148 
149  if (!ctx->initialised) {
150  AVHWFramesContext *main_fc =
151  (AVHWFramesContext*)a->hw_frames_ctx->data;
152  AVHWFramesContext *xfade_fc =
153  (AVHWFramesContext*)b->hw_frames_ctx->data;
154 
155  err = xfade_opencl_load(avctx, main_fc->sw_format,
156  xfade_fc->sw_format);
157  if (err < 0)
158  return err;
159  }
160 
161  output = ff_get_video_buffer(outlink, outlink->w, outlink->h);
162  if (!output) {
163  err = AVERROR(ENOMEM);
164  goto fail;
165  }
166 
167  for (plane = 0; plane < ctx->nb_planes; plane++) {
168  cl_mem mem;
169  kernel_arg = 0;
170 
171  mem = (cl_mem)output->data[plane];
172  CL_SET_KERNEL_ARG(ctx->kernel, kernel_arg, cl_mem, &mem);
173  kernel_arg++;
174 
175  mem = (cl_mem)ctx->xf[0]->data[plane];
176  CL_SET_KERNEL_ARG(ctx->kernel, kernel_arg, cl_mem, &mem);
177  kernel_arg++;
178 
179  mem = (cl_mem)ctx->xf[1]->data[plane];
180  CL_SET_KERNEL_ARG(ctx->kernel, kernel_arg, cl_mem, &mem);
181  kernel_arg++;
182 
183  CL_SET_KERNEL_ARG(ctx->kernel, kernel_arg, cl_float, &progress);
184  kernel_arg++;
185 
186  err = ff_opencl_filter_work_size_from_image(avctx, global_work,
187  output, plane, 0);
188  if (err < 0)
189  goto fail;
190 
191  cle = clEnqueueNDRangeKernel(ctx->command_queue, ctx->kernel, 2, NULL,
192  global_work, NULL, 0, NULL, NULL);
193  CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to enqueue xfade kernel "
194  "for plane %d: %d.\n", plane, cle);
195  }
196 
197  cle = clFinish(ctx->command_queue);
198  CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to finish command queue: %d.\n", cle);
199 
200  err = av_frame_copy_props(output, ctx->xf[0]);
201  if (err < 0)
202  goto fail;
203 
204  output->pts = ctx->pts;
205 
206  return ff_filter_frame(outlink, output);
207 
208 fail:
210  return err;
211 }
212 
214 {
215  AVFilterContext *avctx = outlink->src;
216  XFadeOpenCLContext *ctx = avctx->priv;
217  AVFilterLink *inlink0 = avctx->inputs[0];
218  AVFilterLink *inlink1 = avctx->inputs[1];
219  int err;
220 
221  err = ff_opencl_filter_config_output(outlink);
222  if (err < 0)
223  return err;
224 
225  if (inlink0->w != inlink1->w || inlink0->h != inlink1->h) {
226  av_log(avctx, AV_LOG_ERROR, "First input link %s parameters "
227  "(size %dx%d) do not match the corresponding "
228  "second input link %s parameters (size %dx%d)\n",
229  avctx->input_pads[0].name, inlink0->w, inlink0->h,
230  avctx->input_pads[1].name, inlink1->w, inlink1->h);
231  return AVERROR(EINVAL);
232  }
233 
234  if (inlink0->time_base.num != inlink1->time_base.num ||
235  inlink0->time_base.den != inlink1->time_base.den) {
236  av_log(avctx, AV_LOG_ERROR, "First input link %s timebase "
237  "(%d/%d) do not match the corresponding "
238  "second input link %s timebase (%d/%d)\n",
239  avctx->input_pads[0].name, inlink0->time_base.num, inlink0->time_base.den,
240  avctx->input_pads[1].name, inlink1->time_base.num, inlink1->time_base.den);
241  return AVERROR(EINVAL);
242  }
243 
244  ctx->first_pts = ctx->last_pts = ctx->pts = AV_NOPTS_VALUE;
245 
246  outlink->time_base = inlink0->time_base;
247  outlink->sample_aspect_ratio = inlink0->sample_aspect_ratio;
248  outlink->frame_rate = inlink0->frame_rate;
249 
250  if (ctx->duration)
251  ctx->duration_pts = av_rescale_q(ctx->duration, AV_TIME_BASE_Q, outlink->time_base);
252  if (ctx->offset)
253  ctx->offset_pts = av_rescale_q(ctx->offset, AV_TIME_BASE_Q, outlink->time_base);
254 
255  return 0;
256 }
257 
259 {
260  XFadeOpenCLContext *ctx = avctx->priv;
261  AVFilterLink *outlink = avctx->outputs[0];
262  AVFrame *in = NULL;
263  int ret = 0, status;
264  int64_t pts;
265 
266  FF_FILTER_FORWARD_STATUS_BACK_ALL(outlink, avctx);
267 
268  if (ctx->xfade_is_over) {
269  ret = ff_inlink_consume_frame(avctx->inputs[1], &in);
270  if (ret < 0) {
271  return ret;
272  } else if (ret > 0) {
273  in->pts = (in->pts - ctx->last_pts) + ctx->pts;
274  return ff_filter_frame(outlink, in);
275  } else if (ff_inlink_acknowledge_status(avctx->inputs[1], &status, &pts)) {
276  ff_outlink_set_status(outlink, status, ctx->pts);
277  return 0;
278  } else if (!ret) {
279  if (ff_outlink_frame_wanted(outlink)) {
280  ff_inlink_request_frame(avctx->inputs[1]);
281  return 0;
282  }
283  }
284  }
285 
286  if (ff_inlink_queued_frames(avctx->inputs[0]) > 0) {
287  ctx->xf[0] = ff_inlink_peek_frame(avctx->inputs[0], 0);
288  if (ctx->xf[0]) {
289  if (ctx->first_pts == AV_NOPTS_VALUE) {
290  ctx->first_pts = ctx->xf[0]->pts;
291  }
292  ctx->pts = ctx->xf[0]->pts;
293  if (ctx->first_pts + ctx->offset_pts > ctx->xf[0]->pts) {
294  ctx->xf[0] = NULL;
295  ctx->need_second = 0;
296  ff_inlink_consume_frame(avctx->inputs[0], &in);
297  return ff_filter_frame(outlink, in);
298  }
299 
300  ctx->need_second = 1;
301  }
302  }
303 
304  if (ctx->xf[0] && ff_inlink_queued_frames(avctx->inputs[1]) > 0) {
305  ff_inlink_consume_frame(avctx->inputs[0], &ctx->xf[0]);
306  ff_inlink_consume_frame(avctx->inputs[1], &ctx->xf[1]);
307 
308  ctx->last_pts = ctx->xf[1]->pts;
309  ctx->pts = ctx->xf[0]->pts;
310  if (ctx->xf[0]->pts - (ctx->first_pts + ctx->offset_pts) > ctx->duration_pts)
311  ctx->xfade_is_over = 1;
312  ret = xfade_frame(avctx, ctx->xf[0], ctx->xf[1]);
313  av_frame_free(&ctx->xf[0]);
314  av_frame_free(&ctx->xf[1]);
315  return ret;
316  }
317 
318  if (ff_inlink_queued_frames(avctx->inputs[0]) > 0 &&
319  ff_inlink_queued_frames(avctx->inputs[1]) > 0) {
320  ff_filter_set_ready(avctx, 100);
321  return 0;
322  }
323 
324  if (ff_outlink_frame_wanted(outlink)) {
325  if (!ctx->eof[0] && ff_outlink_get_status(avctx->inputs[0])) {
326  ctx->eof[0] = 1;
327  ctx->xfade_is_over = 1;
328  }
329  if (!ctx->eof[1] && ff_outlink_get_status(avctx->inputs[1])) {
330  ctx->eof[1] = 1;
331  }
332  if (!ctx->eof[0] && !ctx->xf[0])
333  ff_inlink_request_frame(avctx->inputs[0]);
334  if (!ctx->eof[1] && (ctx->need_second || ctx->eof[0]))
335  ff_inlink_request_frame(avctx->inputs[1]);
336  if (ctx->eof[0] && ctx->eof[1] && (
337  ff_inlink_queued_frames(avctx->inputs[0]) <= 0 ||
338  ff_inlink_queued_frames(avctx->inputs[1]) <= 0))
340  return 0;
341  }
342 
343  return FFERROR_NOT_READY;
344 }
345 
347 {
348  XFadeOpenCLContext *ctx = avctx->priv;
349  cl_int cle;
350 
351  if (ctx->kernel) {
352  cle = clReleaseKernel(ctx->kernel);
353  if (cle != CL_SUCCESS)
354  av_log(avctx, AV_LOG_ERROR, "Failed to release "
355  "kernel: %d.\n", cle);
356  }
357 
358  if (ctx->command_queue) {
359  cle = clReleaseCommandQueue(ctx->command_queue);
360  if (cle != CL_SUCCESS)
361  av_log(avctx, AV_LOG_ERROR, "Failed to release "
362  "command queue: %d.\n", cle);
363  }
364 
366 }
367 
369 {
370  XFadeOpenCLContext *s = inlink->dst->priv;
371 
372  return s->xfade_is_over || !s->need_second ?
375 }
376 
377 #define OFFSET(x) offsetof(XFadeOpenCLContext, x)
378 #define FLAGS (AV_OPT_FLAG_FILTERING_PARAM | AV_OPT_FLAG_VIDEO_PARAM)
379 
380 static const AVOption xfade_opencl_options[] = {
381  { "transition", "set cross fade transition", OFFSET(transition), AV_OPT_TYPE_INT, {.i64=1}, 0, NB_TRANSITIONS-1, FLAGS, "transition" },
382  { "custom", "custom transition", 0, AV_OPT_TYPE_CONST, {.i64=CUSTOM}, 0, 0, FLAGS, "transition" },
383  { "fade", "fade transition", 0, AV_OPT_TYPE_CONST, {.i64=FADE}, 0, 0, FLAGS, "transition" },
384  { "wipeleft", "wipe left transition", 0, AV_OPT_TYPE_CONST, {.i64=WIPELEFT}, 0, 0, FLAGS, "transition" },
385  { "wiperight", "wipe right transition", 0, AV_OPT_TYPE_CONST, {.i64=WIPERIGHT}, 0, 0, FLAGS, "transition" },
386  { "wipeup", "wipe up transition", 0, AV_OPT_TYPE_CONST, {.i64=WIPEUP}, 0, 0, FLAGS, "transition" },
387  { "wipedown", "wipe down transition", 0, AV_OPT_TYPE_CONST, {.i64=WIPEDOWN}, 0, 0, FLAGS, "transition" },
388  { "slideleft", "slide left transition", 0, AV_OPT_TYPE_CONST, {.i64=SLIDELEFT}, 0, 0, FLAGS, "transition" },
389  { "slideright", "slide right transition", 0, AV_OPT_TYPE_CONST, {.i64=SLIDERIGHT}, 0, 0, FLAGS, "transition" },
390  { "slideup", "slide up transition", 0, AV_OPT_TYPE_CONST, {.i64=SLIDEUP}, 0, 0, FLAGS, "transition" },
391  { "slidedown", "slide down transition", 0, AV_OPT_TYPE_CONST, {.i64=SLIDEDOWN}, 0, 0, FLAGS, "transition" },
392  { "source", "set OpenCL program source file for custom transition", OFFSET(source_file), AV_OPT_TYPE_STRING, {.str = NULL}, .flags = FLAGS },
393  { "kernel", "set kernel name in program file for custom transition", OFFSET(kernel_name), AV_OPT_TYPE_STRING, {.str = NULL}, .flags = FLAGS },
394  { "duration", "set cross fade duration", OFFSET(duration), AV_OPT_TYPE_DURATION, {.i64=1000000}, 0, 60000000, FLAGS },
395  { "offset", "set cross fade start relative to first input stream", OFFSET(offset), AV_OPT_TYPE_DURATION, {.i64=0}, INT64_MIN, INT64_MAX, FLAGS },
396  { NULL }
397 };
398 
399 AVFILTER_DEFINE_CLASS(xfade_opencl);
400 
402  {
403  .name = "main",
404  .type = AVMEDIA_TYPE_VIDEO,
405  .get_buffer.video = get_video_buffer,
406  .config_props = &ff_opencl_filter_config_input,
407  },
408  {
409  .name = "xfade",
410  .type = AVMEDIA_TYPE_VIDEO,
411  .get_buffer.video = get_video_buffer,
412  .config_props = &ff_opencl_filter_config_input,
413  },
414 };
415 
417  {
418  .name = "default",
419  .type = AVMEDIA_TYPE_VIDEO,
420  .config_props = &xfade_opencl_config_output,
421  },
422 };
423 
425  .name = "xfade_opencl",
426  .description = NULL_IF_CONFIG_SMALL("Cross fade one video with another video."),
427  .priv_size = sizeof(XFadeOpenCLContext),
428  .priv_class = &xfade_opencl_class,
435  .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE,
436 };
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:98
AVPixelFormat
AVPixelFormat
Pixel format.
Definition: pixfmt.h:64
status
they must not be accessed directly The fifo field contains the frames that are queued in the input for processing by the filter The status_in and status_out fields contains the queued status(EOF or error) of the link
FLAGS
#define FLAGS
Definition: vf_xfade_opencl.c:378
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_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:371
ff_filter_frame
int ff_filter_frame(AVFilterLink *link, AVFrame *frame)
Send a frame of data to the next filter.
Definition: avfilter.c:1018
av_pix_fmt_desc_get
const AVPixFmtDescriptor * av_pix_fmt_desc_get(enum AVPixelFormat pix_fmt)
Definition: pixdesc.c:2660
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:260
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:109
AVFrame
This structure describes decoded (raw) audio or video data.
Definition: frame.h:317
pixdesc.h
AVFrame::pts
int64_t pts
Presentation timestamp in time_base units (time when frame should be shown to user).
Definition: frame.h:424
w
uint8_t w
Definition: llviddspenc.c:38
opencl.h
AVOption
AVOption.
Definition: opt.h:247
b
#define b
Definition: input.c:40
xfade_opencl_activate
static int xfade_opencl_activate(AVFilterContext *avctx)
Definition: vf_xfade_opencl.c:258
AV_OPT_TYPE_DURATION
@ AV_OPT_TYPE_DURATION
Definition: opt.h:238
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
SLIDEUP
@ SLIDEUP
Definition: vf_xfade_opencl.c:39
FFMAX
#define FFMAX(a, b)
Definition: macros.h:47
AVFilter::name
const char * name
Filter name.
Definition: avfilter.h:169
SLIDELEFT
@ SLIDELEFT
Definition: vf_xfade_opencl.c:37
video.h
XFadeOpenCLContext::first_pts
int64_t first_pts
Definition: vf_xfade_opencl.c:61
XFadeOpenCLContext::xfade_is_over
int xfade_is_over
Definition: vf_xfade_opencl.c:64
init
static int init
Definition: av_tx.c:47
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:1417
ff_default_get_video_buffer
AVFrame * ff_default_get_video_buffer(AVFilterLink *link, int w, int h)
Definition: video.c:43
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:263
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:212
XFadeOpenCLContext::xf
AVFrame * xf[2]
Definition: vf_xfade_opencl.c:67
AVFilterContext::priv
void * priv
private data for use by the filter
Definition: avfilter.h:417
XFadeOpenCLContext::offset
int64_t offset
Definition: vf_xfade_opencl.c:51
fail
#define fail()
Definition: checkasm.h:127
xfade_opencl_uninit
static av_cold void xfade_opencl_uninit(AVFilterContext *avctx)
Definition: vf_xfade_opencl.c:346
FADE
@ FADE
Definition: vf_xfade_opencl.c:32
XFadeOpenCLContext::need_second
int need_second
Definition: vf_xfade_opencl.c:65
pts
static int64_t pts
Definition: transcode_aac.c:653
XFadeOpenCLContext
Definition: vf_xfade_opencl.c:44
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:81
AVFilterPad
A filter pad used for either input or output.
Definition: internal.h:50
AVFilterContext::input_pads
AVFilterPad * input_pads
array of input pads
Definition: avfilter.h:409
AV_LOG_ERROR
#define AV_LOG_ERROR
Something went wrong and cannot losslessly be recovered.
Definition: log.h:180
av_cold
#define av_cold
Definition: attributes.h:90
SLIDEDOWN
@ SLIDEDOWN
Definition: vf_xfade_opencl.c:40
duration
int64_t duration
Definition: movenc.c:64
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:189
ff_inlink_request_frame
void ff_inlink_request_frame(AVFilterLink *link)
Mark that a frame is wanted on the link.
Definition: avfilter.c:1534
s
#define s(width, name)
Definition: cbs_vp9.c:257
filters.h
WIPERIGHT
@ WIPERIGHT
Definition: vf_xfade_opencl.c:34
ctx
AVFormatContext * ctx
Definition: movenc.c:48
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:141
CUSTOM
@ CUSTOM
Definition: vf_xfade_opencl.c:31
OFFSET
#define OFFSET(x)
Definition: vf_xfade_opencl.c:377
f
#define f(width, name)
Definition: cbs_vp9.c:255
FILTER_INPUTS
#define FILTER_INPUTS(array)
Definition: internal.h:191
XFadeOpenCLContext::kernel_name
const char * kernel_name
Definition: vf_xfade_opencl.c:49
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:1456
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:537
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:410
av_clipf
#define av_clipf
Definition: common.h:144
AVComponentDescriptor::plane
int plane
Which of the 4 planes contains the component.
Definition: pixdesc.h:34
SLIDERIGHT
@ SLIDERIGHT
Definition: vf_xfade_opencl.c:38
AV_PIX_FMT_OPENCL
@ AV_PIX_FMT_OPENCL
Hardware surfaces for OpenCL.
Definition: pixfmt.h:325
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:1371
ff_inlink_queued_frames
size_t ff_inlink_queued_frames(AVFilterLink *link)
Get the number of frames available on the link.
Definition: avfilter.c:1386
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:117
XFadeOpenCLContext::initialised
int initialised
Definition: vf_xfade_opencl.c:53
XFadeOpenCLContext::last_pts
int64_t last_pts
Definition: vf_xfade_opencl.c:62
NB_TRANSITIONS
@ NB_TRANSITIONS
Definition: vf_xfade_opencl.c:41
AV_NOPTS_VALUE
#define AV_NOPTS_VALUE
Undefined timestamp value.
Definition: avutil.h:248
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:416
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:45
ff_null_get_video_buffer
AVFrame * ff_null_get_video_buffer(AVFilterLink *link, int w, int h)
Definition: video.c:38
internal.h
FILTER_SINGLE_PIXFMT
#define FILTER_SINGLE_PIXFMT(pix_fmt_)
Definition: internal.h:181
xfade_opencl_load
static int xfade_opencl_load(AVFilterContext *avctx, enum AVPixelFormat main_format, enum AVPixelFormat xfade_format)
Definition: vf_xfade_opencl.c:70
log.h
i
#define i(width, name, range_min, range_max)
Definition: cbs_h2645.c:271
WIPEDOWN
@ WIPEDOWN
Definition: vf_xfade_opencl.c:36
xfade_opencl_options
static const AVOption xfade_opencl_options[]
Definition: vf_xfade_opencl.c:380
XFadeOpenCLContext::eof
int eof[2]
Definition: vf_xfade_opencl.c:66
XFadeOpenCLContext::transition
int transition
Definition: vf_xfade_opencl.c:47
AVFilterPad::name
const char * name
Pad name.
Definition: internal.h:56
AVFilter
Filter definition.
Definition: avfilter.h:165
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
ret
ret
Definition: filter_design.txt:187
XFadeOpenCLContext::pts
int64_t pts
Definition: vf_xfade_opencl.c:63
XFadeOpenCLContext::kernel
cl_kernel kernel
Definition: vf_xfade_opencl.c:54
XFadeOpenCLContext::ocf
OpenCLFilterContext ocf
Definition: vf_xfade_opencl.c:45
XFadeOpenCLContext::source_file
const char * source_file
Definition: vf_xfade_opencl.c:48
AVRational::den
int den
Denominator.
Definition: rational.h:60
AV_OPT_TYPE_INT
@ AV_OPT_TYPE_INT
Definition: opt.h:224
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:204
ff_vf_xfade_opencl
const AVFilter ff_vf_xfade_opencl
Definition: vf_xfade_opencl.c:424
XFadeOpenCLContext::duration
int64_t duration
Definition: vf_xfade_opencl.c:50
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:141
AVFormatContext::duration
int64_t duration
Duration of the stream, in AV_TIME_BASE fractional seconds.
Definition: avformat.h:1302
ff_outlink_get_status
int ff_outlink_get_status(AVFilterLink *link)
Get the status on an output link.
Definition: avfilter.c:1557
AVFilterContext
An instance of a filter.
Definition: avfilter.h:402
get_video_buffer
static AVFrame * get_video_buffer(AVFilterLink *inlink, int w, int h)
Definition: vf_xfade_opencl.c:368
AVMEDIA_TYPE_VIDEO
@ AVMEDIA_TYPE_VIDEO
Definition: avutil.h:201
XFadeOpenCLContext::command_queue
cl_command_queue command_queue
Definition: vf_xfade_opencl.c:55
XFadeOpenCLContext::duration_pts
int64_t duration_pts
Definition: vf_xfade_opencl.c:59
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:31
FILTER_OUTPUTS
#define FILTER_OUTPUTS(array)
Definition: internal.h:192
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:28
XFadeOpenCLContext::nb_planes
int nb_planes
Definition: vf_xfade_opencl.c:57
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:138
uninit
static av_cold int uninit(AVCodecContext *avctx)
Definition: crystalhd.c:282
h
h
Definition: vp9dsp_template.c:2038
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:60
AV_OPT_TYPE_STRING
@ AV_OPT_TYPE_STRING
Definition: opt.h:228
xfade_opencl_config_output
static int xfade_opencl_config_output(AVFilterLink *outlink)
Definition: vf_xfade_opencl.c:213
AVFILTER_DEFINE_CLASS
AVFILTER_DEFINE_CLASS(xfade_opencl)
ff_opencl_source_xfade
const char * ff_opencl_source_xfade
WIPELEFT
@ WIPELEFT
Definition: vf_xfade_opencl.c:33
AV_OPT_TYPE_CONST
@ AV_OPT_TYPE_CONST
Definition: opt.h:233
WIPEUP
@ WIPEUP
Definition: vf_xfade_opencl.c:35
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:211
xfade_opencl_inputs
static const AVFilterPad xfade_opencl_inputs[]
Definition: vf_xfade_opencl.c:401
AVFilterContext::outputs
AVFilterLink ** outputs
array of pointers to output links
Definition: avfilter.h:414