FFmpeg  4.3.8
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/mem.h"
21 #include "libavutil/opt.h"
22 #include "libavutil/pixdesc.h"
23 
24 #include "avfilter.h"
25 #include "filters.h"
26 #include "internal.h"
27 #include "opencl.h"
28 #include "opencl_source.h"
29 #include "video.h"
30 
43 };
44 
45 typedef struct XFadeOpenCLContext {
47 
49  const char *source_file;
50  const char *kernel_name;
53 
55  cl_kernel kernel;
56  cl_command_queue command_queue;
57 
58  int nb_planes;
59 
67  int eof[2];
68  AVFrame *xf[2];
70 
72  enum AVPixelFormat main_format,
73  enum AVPixelFormat xfade_format)
74 {
75  XFadeOpenCLContext *ctx = avctx->priv;
76  cl_int cle;
77  const AVPixFmtDescriptor *main_desc;
78  int err, main_planes;
79  const char *kernel_name;
80 
81  main_desc = av_pix_fmt_desc_get(main_format);
82  if (main_format != xfade_format) {
83  av_log(avctx, AV_LOG_ERROR, "Input formats are not same.\n");
84  return AVERROR(EINVAL);
85  }
86 
87  main_planes = 0;
88  for (int i = 0; i < main_desc->nb_components; i++)
89  main_planes = FFMAX(main_planes,
90  main_desc->comp[i].plane + 1);
91 
92  ctx->nb_planes = main_planes;
93 
94  if (ctx->transition == CUSTOM) {
96  } else {
98  }
99  if (err < 0)
100  return err;
101 
102  ctx->command_queue = clCreateCommandQueue(ctx->ocf.hwctx->context,
103  ctx->ocf.hwctx->device_id,
104  0, &cle);
105  CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create OpenCL "
106  "command queue %d.\n", cle);
107 
108  switch (ctx->transition) {
109  case CUSTOM: kernel_name = ctx->kernel_name; break;
110  case FADE: kernel_name = "fade"; break;
111  case WIPELEFT: kernel_name = "wipeleft"; break;
112  case WIPERIGHT: kernel_name = "wiperight"; break;
113  case WIPEUP: kernel_name = "wipeup"; break;
114  case WIPEDOWN: kernel_name = "wipedown"; break;
115  case SLIDELEFT: kernel_name = "slideleft"; break;
116  case SLIDERIGHT: kernel_name = "slideright"; break;
117  case SLIDEUP: kernel_name = "slideup"; break;
118  case SLIDEDOWN: kernel_name = "slidedown"; break;
119  default:
120  err = AVERROR_BUG;
121  goto fail;
122  }
123 
124  ctx->kernel = clCreateKernel(ctx->ocf.program, kernel_name, &cle);
125  CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create kernel %d.\n", cle);
126 
127  ctx->initialised = 1;
128 
129  return 0;
130 
131 fail:
132  if (ctx->command_queue)
133  clReleaseCommandQueue(ctx->command_queue);
134  if (ctx->kernel)
135  clReleaseKernel(ctx->kernel);
136  return err;
137 }
138 
140 {
141  AVFilterLink *outlink = avctx->outputs[0];
142  XFadeOpenCLContext *ctx = avctx->priv;
143  AVFrame *output;
144  cl_int cle;
145  cl_float progress = av_clipf(1.f - ((cl_float)(ctx->pts - ctx->first_pts - ctx->offset_pts) / ctx->duration_pts), 0.f, 1.f);
146  size_t global_work[2];
147  int kernel_arg = 0;
148  int err, plane;
149 
150  if (!ctx->initialised) {
151  AVHWFramesContext *main_fc =
153  AVHWFramesContext *xfade_fc =
155 
156  err = xfade_opencl_load(avctx, main_fc->sw_format,
157  xfade_fc->sw_format);
158  if (err < 0)
159  return err;
160  }
161 
162  output = ff_get_video_buffer(outlink, outlink->w, outlink->h);
163  if (!output) {
164  err = AVERROR(ENOMEM);
165  goto fail;
166  }
167 
168  for (plane = 0; plane < ctx->nb_planes; plane++) {
169  cl_mem mem;
170  kernel_arg = 0;
171 
172  mem = (cl_mem)output->data[plane];
173  CL_SET_KERNEL_ARG(ctx->kernel, kernel_arg, cl_mem, &mem);
174  kernel_arg++;
175 
176  mem = (cl_mem)ctx->xf[0]->data[plane];
177  CL_SET_KERNEL_ARG(ctx->kernel, kernel_arg, cl_mem, &mem);
178  kernel_arg++;
179 
180  mem = (cl_mem)ctx->xf[1]->data[plane];
181  CL_SET_KERNEL_ARG(ctx->kernel, kernel_arg, cl_mem, &mem);
182  kernel_arg++;
183 
184  CL_SET_KERNEL_ARG(ctx->kernel, kernel_arg, cl_float, &progress);
185  kernel_arg++;
186 
187  err = ff_opencl_filter_work_size_from_image(avctx, global_work,
188  output, plane, 0);
189  if (err < 0)
190  goto fail;
191 
192  cle = clEnqueueNDRangeKernel(ctx->command_queue, ctx->kernel, 2, NULL,
193  global_work, NULL, 0, NULL, NULL);
194  CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to enqueue xfade kernel "
195  "for plane %d: %d.\n", plane, cle);
196  }
197 
198  cle = clFinish(ctx->command_queue);
199  CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to finish command queue: %d.\n", cle);
200 
201  err = av_frame_copy_props(output, ctx->xf[0]);
202  if (err < 0)
203  goto fail;
204 
205  output->pts = ctx->pts;
206 
207  return ff_filter_frame(outlink, output);
208 
209 fail:
210  av_frame_free(&output);
211  return err;
212 }
213 
215 {
216  AVFilterContext *avctx = outlink->src;
217  XFadeOpenCLContext *ctx = avctx->priv;
218  AVFilterLink *inlink0 = avctx->inputs[0];
219  AVFilterLink *inlink1 = avctx->inputs[1];
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  outlink->frame_rate = inlink0->frame_rate;
250 
251  if (ctx->duration)
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 
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  ff_inlink_consume_frame(avctx->inputs[0], &in);
298  return ff_filter_frame(outlink, in);
299  }
300 
301  ctx->need_second = 1;
302  }
303  }
304 
305  if (ctx->xf[0] && ff_inlink_queued_frames(avctx->inputs[1]) > 0) {
306  ff_inlink_consume_frame(avctx->inputs[0], &ctx->xf[0]);
307  ff_inlink_consume_frame(avctx->inputs[1], &ctx->xf[1]);
308 
309  ctx->last_pts = ctx->xf[1]->pts;
310  ctx->pts = ctx->xf[0]->pts;
311  if (ctx->xf[0]->pts - (ctx->first_pts + ctx->offset_pts) > ctx->duration_pts)
312  ctx->xfade_is_over = 1;
313  ret = xfade_frame(avctx, ctx->xf[0], ctx->xf[1]);
314  av_frame_free(&ctx->xf[0]);
315  av_frame_free(&ctx->xf[1]);
316  return ret;
317  }
318 
319  if (ff_inlink_queued_frames(avctx->inputs[0]) > 0 &&
320  ff_inlink_queued_frames(avctx->inputs[1]) > 0) {
321  ff_filter_set_ready(avctx, 100);
322  return 0;
323  }
324 
325  if (ff_outlink_frame_wanted(outlink)) {
326  if (!ctx->eof[0] && ff_outlink_get_status(avctx->inputs[0])) {
327  ctx->eof[0] = 1;
328  ctx->xfade_is_over = 1;
329  }
330  if (!ctx->eof[1] && ff_outlink_get_status(avctx->inputs[1])) {
331  ctx->eof[1] = 1;
332  }
333  if (!ctx->eof[0] && !ctx->xf[0])
334  ff_inlink_request_frame(avctx->inputs[0]);
335  if (!ctx->eof[1] && (ctx->need_second || ctx->eof[0]))
336  ff_inlink_request_frame(avctx->inputs[1]);
337  if (ctx->eof[0] && ctx->eof[1] && (
338  ff_inlink_queued_frames(avctx->inputs[0]) <= 0 ||
339  ff_inlink_queued_frames(avctx->inputs[1]) <= 0))
341  return 0;
342  }
343 
344  return FFERROR_NOT_READY;
345 }
346 
348 {
349  XFadeOpenCLContext *ctx = avctx->priv;
350  cl_int cle;
351 
352  if (ctx->kernel) {
353  cle = clReleaseKernel(ctx->kernel);
354  if (cle != CL_SUCCESS)
355  av_log(avctx, AV_LOG_ERROR, "Failed to release "
356  "kernel: %d.\n", cle);
357  }
358 
359  if (ctx->command_queue) {
360  cle = clReleaseCommandQueue(ctx->command_queue);
361  if (cle != CL_SUCCESS)
362  av_log(avctx, AV_LOG_ERROR, "Failed to release "
363  "command queue: %d.\n", cle);
364  }
365 
367 }
368 
369 static AVFrame *get_video_buffer(AVFilterLink *inlink, int w, int h)
370 {
371  XFadeOpenCLContext *s = inlink->dst->priv;
372 
373  return s->xfade_is_over || !s->need_second ?
374  ff_null_get_video_buffer (inlink, w, h) :
375  ff_default_get_video_buffer(inlink, w, h);
376 }
377 
378 #define OFFSET(x) offsetof(XFadeOpenCLContext, x)
379 #define FLAGS (AV_OPT_FLAG_FILTERING_PARAM | AV_OPT_FLAG_VIDEO_PARAM)
380 
381 static const AVOption xfade_opencl_options[] = {
382  { "transition", "set cross fade transition", OFFSET(transition), AV_OPT_TYPE_INT, {.i64=1}, 0, NB_TRANSITIONS-1, FLAGS, "transition" },
383  { "custom", "custom transition", 0, AV_OPT_TYPE_CONST, {.i64=CUSTOM}, 0, 0, FLAGS, "transition" },
384  { "fade", "fade transition", 0, AV_OPT_TYPE_CONST, {.i64=FADE}, 0, 0, FLAGS, "transition" },
385  { "wipeleft", "wipe left transition", 0, AV_OPT_TYPE_CONST, {.i64=WIPELEFT}, 0, 0, FLAGS, "transition" },
386  { "wiperight", "wipe right transition", 0, AV_OPT_TYPE_CONST, {.i64=WIPERIGHT}, 0, 0, FLAGS, "transition" },
387  { "wipeup", "wipe up transition", 0, AV_OPT_TYPE_CONST, {.i64=WIPEUP}, 0, 0, FLAGS, "transition" },
388  { "wipedown", "wipe down transition", 0, AV_OPT_TYPE_CONST, {.i64=WIPEDOWN}, 0, 0, FLAGS, "transition" },
389  { "slideleft", "slide left transition", 0, AV_OPT_TYPE_CONST, {.i64=SLIDELEFT}, 0, 0, FLAGS, "transition" },
390  { "slideright", "slide right transition", 0, AV_OPT_TYPE_CONST, {.i64=SLIDERIGHT}, 0, 0, FLAGS, "transition" },
391  { "slideup", "slide up transition", 0, AV_OPT_TYPE_CONST, {.i64=SLIDEUP}, 0, 0, FLAGS, "transition" },
392  { "slidedown", "slide down transition", 0, AV_OPT_TYPE_CONST, {.i64=SLIDEDOWN}, 0, 0, FLAGS, "transition" },
393  { "source", "set OpenCL program source file for custom transition", OFFSET(source_file), AV_OPT_TYPE_STRING, {.str = NULL}, .flags = FLAGS },
394  { "kernel", "set kernel name in program file for custom transition", OFFSET(kernel_name), AV_OPT_TYPE_STRING, {.str = NULL}, .flags = FLAGS },
395  { "duration", "set cross fade duration", OFFSET(duration), AV_OPT_TYPE_DURATION, {.i64=1000000}, 0, 60000000, FLAGS },
396  { "offset", "set cross fade start relative to first input stream", OFFSET(offset), AV_OPT_TYPE_DURATION, {.i64=0}, INT64_MIN, INT64_MAX, FLAGS },
397  { NULL }
398 };
399 
400 AVFILTER_DEFINE_CLASS(xfade_opencl);
401 
403  {
404  .name = "main",
405  .type = AVMEDIA_TYPE_VIDEO,
406  .get_video_buffer = get_video_buffer,
407  .config_props = &ff_opencl_filter_config_input,
408  },
409  {
410  .name = "xfade",
411  .type = AVMEDIA_TYPE_VIDEO,
412  .get_video_buffer = get_video_buffer,
413  .config_props = &ff_opencl_filter_config_input,
414  },
415  { NULL }
416 };
417 
419  {
420  .name = "default",
421  .type = AVMEDIA_TYPE_VIDEO,
422  .config_props = &xfade_opencl_config_output,
423  },
424  { NULL }
425 };
426 
428  .name = "xfade_opencl",
429  .description = NULL_IF_CONFIG_SMALL("Cross fade one video with another video."),
430  .priv_size = sizeof(XFadeOpenCLContext),
431  .priv_class = &xfade_opencl_class,
436  .inputs = xfade_opencl_inputs,
437  .outputs = xfade_opencl_outputs,
438  .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE,
439 };
int plane
Which of the 4 planes contains the component.
Definition: pixdesc.h:35
int ff_inlink_consume_frame(AVFilterLink *link, AVFrame **rframe)
Take a frame from the link&#39;s FIFO and update the link&#39;s stats.
Definition: avfilter.c:1476
const char * kernel_name
#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:365
static const AVFilterPad xfade_opencl_outputs[]
const AVPixFmtDescriptor * av_pix_fmt_desc_get(enum AVPixelFormat pix_fmt)
Definition: pixdesc.c:2549
This structure describes decoded (raw) audio or video data.
Definition: frame.h:300
AVOption.
Definition: opt.h:246
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:278
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
int ff_opencl_filter_query_formats(AVFilterContext *avctx)
Return that all inputs and outputs support only AV_PIX_FMT_OPENCL.
Definition: opencl.c:28
#define FF_FILTER_FORWARD_STATUS_BACK_ALL(outlink, filter)
Forward the status on an output link to all input links.
Definition: filters.h:212
static av_cold void xfade_opencl_uninit(AVFilterContext *avctx)
Main libavfilter public API header.
Memory handling functions.
static av_cold int init(AVCodecContext *avctx)
Definition: avrndec.c:35
static int xfade_frame(AVFilterContext *avctx, AVFrame *a, AVFrame *b)
AVFILTER_DEFINE_CLASS(xfade_opencl)
int num
Numerator.
Definition: rational.h:59
const char * b
Definition: vf_curves.c:116
AVFrame * ff_null_get_video_buffer(AVFilterLink *link, int w, int h)
Definition: video.c:39
#define FFERROR_NOT_READY
Filters implementation helper functions.
Definition: filters.h:34
AVFrame * ff_get_video_buffer(AVFilterLink *link, int w, int h)
Request a picture buffer with a specific set of permissions.
Definition: video.c:104
int ff_opencl_filter_init(AVFilterContext *avctx)
Initialise an OpenCL filter context.
Definition: opencl.c:147
static int xfade_opencl_activate(AVFilterContext *avctx)
#define OFFSET(x)
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
void ff_inlink_request_frame(AVFilterLink *link)
Mark that a frame is wanted on the link.
Definition: avfilter.c:1602
static int ff_outlink_frame_wanted(AVFilterLink *link)
Test if a frame is wanted on an output link.
Definition: filters.h:172
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:639
const char * name
Pad name.
Definition: internal.h:60
AVFilterLink ** inputs
array of pointers to input links
Definition: avfilter.h:346
int ff_filter_frame(AVFilterLink *link, AVFrame *frame)
Send a frame of data to the next filter.
Definition: avfilter.c:1075
AVComponentDescriptor comp[4]
Parameters that describe how pixels are packed.
Definition: pixdesc.h:117
#define av_cold
Definition: attributes.h:88
static av_cold int uninit(AVCodecContext *avctx)
Definition: crystalhd.c:279
AVOptions.
#define FLAGS
#define f(width, name)
Definition: cbs_vp9.c:255
static int xfade_opencl_load(AVFilterContext *avctx, enum AVPixelFormat main_format, enum AVPixelFormat xfade_format)
cl_device_id device_id
The primary device ID of the device.
int64_t pts
Presentation timestamp in time_base units (time when frame should be shown to user).
Definition: frame.h:393
#define AVERROR_EOF
End of file.
Definition: error.h:55
int ff_opencl_filter_config_output(AVFilterLink *outlink)
Create a suitable hardware frames context for the output.
Definition: opencl.c:96
#define av_log(a,...)
A filter pad used for either input or output.
Definition: internal.h:54
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
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:1431
AVFilterPad * input_pads
array of input pads
Definition: avfilter.h:345
OpenCLFilterContext ocf
#define i(width, name, range_min, range_max)
Definition: cbs_h2645.c:269
#define AV_LOG_ERROR
Something went wrong and cannot losslessly be recovered.
Definition: log.h:176
AVFilter ff_vf_xfade_opencl
#define AVERROR(e)
Definition: error.h:43
const char * source_file
void av_frame_free(AVFrame **frame)
Free the frame and any dynamically allocated objects in it, e.g.
Definition: frame.c:203
#define NULL_IF_CONFIG_SMALL(x)
Return NULL if CONFIG_SMALL is true, otherwise the argument without modification. ...
Definition: internal.h:188
void * priv
private data for use by the filter
Definition: avfilter.h:353
#define FFMAX(a, b)
Definition: common.h:94
#define fail()
Definition: checkasm.h:123
cl_command_queue command_queue
static const AVOption xfade_opencl_options[]
uint8_t nb_components
The number of components each pixel has, (1-4)
Definition: pixdesc.h:83
XFadeTransitions
Definition: vf_xfade.c:31
uint8_t w
Definition: llviddspenc.c:38
AVFormatContext * ctx
Definition: movenc.c:48
static int activate(AVFilterContext *ctx)
Definition: af_adeclick.c:622
#define s(width, name)
Definition: cbs_vp9.c:257
int ff_opencl_filter_load_program_from_file(AVFilterContext *avctx, const char *filename)
Load a new OpenCL program from a file.
Definition: opencl.c:219
static const AVFilterPad inputs[]
Definition: af_acontrast.c:193
AVFrame * ff_inlink_peek_frame(AVFilterLink *link, size_t idx)
Access a frame in the link fifo without consuming it.
Definition: avfilter.c:1515
static const AVFilterPad outputs[]
Definition: af_acontrast.c:203
#define CL_SET_KERNEL_ARG(kernel, arg_num, type, arg)
set argument to specific Kernel.
Definition: opencl.h:61
#define AV_TIME_BASE_Q
Internal time base represented as fractional value.
Definition: avutil.h:260
Descriptor that unambiguously describes how the bits of a pixel are stored in the up to 4 data planes...
Definition: pixdesc.h:81
long long int64_t
Definition: coverity.c:34
uint8_t * data
The data buffer.
Definition: buffer.h:89
#define AVERROR_BUG
Internal bug, also see AVERROR_BUG2.
Definition: error.h:50
uint8_t pi<< 24) CONV_FUNC_GROUP(AV_SAMPLE_FMT_FLT, float, AV_SAMPLE_FMT_U8, uint8_t,(*(const uint8_t *) pi - 0x80) *(1.0f/(1<< 7))) CONV_FUNC_GROUP(AV_SAMPLE_FMT_DBL, double, AV_SAMPLE_FMT_U8, uint8_t,(*(const uint8_t *) pi - 0x80) *(1.0/(1<< 7))) CONV_FUNC_GROUP(AV_SAMPLE_FMT_U8, uint8_t, AV_SAMPLE_FMT_S16, int16_t,(*(const int16_t *) pi >> 8)+0x80) CONV_FUNC_GROUP(AV_SAMPLE_FMT_FLT, float, AV_SAMPLE_FMT_S16, int16_t, *(const int16_t *) pi *(1.0f/(1<< 15))) CONV_FUNC_GROUP(AV_SAMPLE_FMT_DBL, double, AV_SAMPLE_FMT_S16, int16_t, *(const int16_t *) pi *(1.0/(1<< 15))) CONV_FUNC_GROUP(AV_SAMPLE_FMT_U8, uint8_t, AV_SAMPLE_FMT_S32, int32_t,(*(const int32_t *) pi >> 24)+0x80) CONV_FUNC_GROUP(AV_SAMPLE_FMT_FLT, float, AV_SAMPLE_FMT_S32, int32_t, *(const int32_t *) pi *(1.0f/(1U<< 31))) CONV_FUNC_GROUP(AV_SAMPLE_FMT_DBL, double, AV_SAMPLE_FMT_S32, int32_t, *(const int32_t *) pi *(1.0/(1U<< 31))) CONV_FUNC_GROUP(AV_SAMPLE_FMT_U8, uint8_t, AV_SAMPLE_FMT_FLT, float, av_clip_uint8(lrintf(*(const float *) pi *(1<< 7))+0x80)) CONV_FUNC_GROUP(AV_SAMPLE_FMT_S16, int16_t, AV_SAMPLE_FMT_FLT, float, av_clip_int16(lrintf(*(const float *) pi *(1<< 15)))) CONV_FUNC_GROUP(AV_SAMPLE_FMT_S32, int32_t, AV_SAMPLE_FMT_FLT, float, av_clipl_int32(llrintf(*(const float *) pi *(1U<< 31)))) CONV_FUNC_GROUP(AV_SAMPLE_FMT_U8, uint8_t, AV_SAMPLE_FMT_DBL, double, av_clip_uint8(lrint(*(const double *) pi *(1<< 7))+0x80)) CONV_FUNC_GROUP(AV_SAMPLE_FMT_S16, int16_t, AV_SAMPLE_FMT_DBL, double, av_clip_int16(lrint(*(const double *) pi *(1<< 15)))) CONV_FUNC_GROUP(AV_SAMPLE_FMT_S32, int32_t, AV_SAMPLE_FMT_DBL, double, av_clipl_int32(llrint(*(const double *) pi *(1U<< 31)))) #define SET_CONV_FUNC_GROUP(ofmt, ifmt) static void set_generic_function(AudioConvert *ac) { } void ff_audio_convert_free(AudioConvert **ac) { if(! *ac) return;ff_dither_free(&(*ac) ->dc);av_freep(ac);} AudioConvert *ff_audio_convert_alloc(AVAudioResampleContext *avr, enum AVSampleFormat out_fmt, enum AVSampleFormat in_fmt, int channels, int sample_rate, int apply_map) { AudioConvert *ac;int in_planar, out_planar;ac=av_mallocz(sizeof(*ac));if(!ac) return NULL;ac->avr=avr;ac->out_fmt=out_fmt;ac->in_fmt=in_fmt;ac->channels=channels;ac->apply_map=apply_map;if(avr->dither_method !=AV_RESAMPLE_DITHER_NONE &&av_get_packed_sample_fmt(out_fmt)==AV_SAMPLE_FMT_S16 &&av_get_bytes_per_sample(in_fmt) > 2) { ac->dc=ff_dither_alloc(avr, out_fmt, in_fmt, channels, sample_rate, apply_map);if(!ac->dc) { av_free(ac);return NULL;} return ac;} in_planar=ff_sample_fmt_is_planar(in_fmt, channels);out_planar=ff_sample_fmt_is_planar(out_fmt, channels);if(in_planar==out_planar) { ac->func_type=CONV_FUNC_TYPE_FLAT;ac->planes=in_planar ? ac->channels :1;} else if(in_planar) ac->func_type=CONV_FUNC_TYPE_INTERLEAVE;else ac->func_type=CONV_FUNC_TYPE_DEINTERLEAVE;set_generic_function(ac);if(ARCH_AARCH64) ff_audio_convert_init_aarch64(ac);if(ARCH_ARM) ff_audio_convert_init_arm(ac);if(ARCH_X86) ff_audio_convert_init_x86(ac);return ac;} int ff_audio_convert(AudioConvert *ac, AudioData *out, AudioData *in) { int use_generic=1;int len=in->nb_samples;int p;if(ac->dc) { av_log(ac->avr, AV_LOG_TRACE, "%d samples - audio_convert: %s to %s (dithered)\", len, av_get_sample_fmt_name(ac->in_fmt), av_get_sample_fmt_name(ac->out_fmt));return ff_convert_dither(ac-> in
Filter definition.
Definition: avfilter.h:144
int ff_outlink_get_status(AVFilterLink *link)
Get the status on an output link.
Definition: avfilter.c:1625
This struct describes a set or pool of "hardware" frames (i.e.
Definition: hwcontext.h:124
const char * name
Filter name.
Definition: avfilter.h:148
size_t ff_inlink_queued_frames(AVFilterLink *link)
Get the number of frames available on the link.
Definition: avfilter.c:1446
AVFilterLink ** outputs
array of pointers to output links
Definition: avfilter.h:350
#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:314
AVFrame * ff_default_get_video_buffer(AVFilterLink *link, int w, int h)
Definition: video.c:99
const char * ff_opencl_source_xfade
Definition: xfade.c:2
static int xfade_opencl_config_output(AVFilterLink *outlink)
static int query_formats(AVFilterContext *ctx)
Definition: aeval.c:244
static const AVFilterPad xfade_opencl_inputs[]
void ff_opencl_filter_uninit(AVFilterContext *avctx)
Uninitialise an OpenCL filter context.
Definition: opencl.c:156
void ff_filter_set_ready(AVFilterContext *filter, unsigned priority)
Mark a filter ready and schedule it for activation.
Definition: avfilter.c:193
int den
Denominator.
Definition: rational.h:60
static AVFrame * get_video_buffer(AVFilterLink *inlink, int w, int h)
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
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
internal API functions
enum AVPixelFormat sw_format
The pixel format identifying the actual data layout of the hardware frames.
Definition: hwcontext.h:222
AVPixelFormat
Pixel format.
Definition: pixfmt.h:64
int av_frame_copy_props(AVFrame *dst, const AVFrame *src)
Copy only "metadata" fields from src to dst.
Definition: frame.c:659
#define AV_NOPTS_VALUE
Undefined timestamp value.
Definition: avutil.h:248