FFmpeg
 All Data Structures Namespaces Files Functions Variables Typedefs Enumerations Enumerator Macros Groups Pages
vf_program_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/avstring.h"
20 #include "libavutil/log.h"
21 #include "libavutil/mem.h"
22 #include "libavutil/opt.h"
23 #include "libavutil/pixdesc.h"
24 
25 #include "avfilter.h"
26 #include "framesync.h"
27 #include "internal.h"
28 #include "opencl.h"
29 #include "video.h"
30 
31 typedef struct ProgramOpenCLContext {
33 
34  int loaded;
35  cl_uint index;
36  cl_kernel kernel;
37  cl_command_queue command_queue;
38 
41 
42  const char *source_file;
43  const char *kernel_name;
44  int nb_inputs;
45  int width, height;
49 
51 {
52  ProgramOpenCLContext *ctx = avctx->priv;
53  cl_int cle;
54  int err;
55 
57  if (err < 0)
58  return err;
59 
60  ctx->command_queue = clCreateCommandQueue(ctx->ocf.hwctx->context,
61  ctx->ocf.hwctx->device_id,
62  0, &cle);
63  if (!ctx->command_queue) {
64  av_log(avctx, AV_LOG_ERROR, "Failed to create OpenCL "
65  "command queue: %d.\n", cle);
66  return AVERROR(EIO);
67  }
68 
69  ctx->kernel = clCreateKernel(ctx->ocf.program, ctx->kernel_name, &cle);
70  if (!ctx->kernel) {
71  if (cle == CL_INVALID_KERNEL_NAME) {
72  av_log(avctx, AV_LOG_ERROR, "Kernel function '%s' not found in "
73  "program.\n", ctx->kernel_name);
74  } else {
75  av_log(avctx, AV_LOG_ERROR, "Failed to create kernel: %d.\n", cle);
76  }
77  return AVERROR(EIO);
78  }
79 
80  ctx->loaded = 1;
81  return 0;
82 }
83 
85 {
86  AVFilterLink *outlink = avctx->outputs[0];
87  ProgramOpenCLContext *ctx = avctx->priv;
88  AVFrame *output = NULL;
89  cl_int cle;
90  size_t global_work[2];
91  cl_mem src, dst;
92  int err, input, plane;
93 
94  if (!ctx->loaded) {
95  err = program_opencl_load(avctx);
96  if (err < 0)
97  return err;
98  }
99 
100  output = ff_get_video_buffer(outlink, outlink->w, outlink->h);
101  if (!output) {
102  err = AVERROR(ENOMEM);
103  goto fail;
104  }
105 
106  for (plane = 0; plane < FF_ARRAY_ELEMS(output->data); plane++) {
107  dst = (cl_mem)output->data[plane];
108  if (!dst)
109  break;
110 
111  cle = clSetKernelArg(ctx->kernel, 0, sizeof(cl_mem), &dst);
112  if (cle != CL_SUCCESS) {
113  av_log(avctx, AV_LOG_ERROR, "Failed to set kernel "
114  "destination image argument: %d.\n", cle);
115  err = AVERROR_UNKNOWN;
116  goto fail;
117  }
118  cle = clSetKernelArg(ctx->kernel, 1, sizeof(cl_uint), &ctx->index);
119  if (cle != CL_SUCCESS) {
120  av_log(avctx, AV_LOG_ERROR, "Failed to set kernel "
121  "index argument: %d.\n", cle);
122  err = AVERROR_UNKNOWN;
123  goto fail;
124  }
125 
126  for (input = 0; input < ctx->nb_inputs; input++) {
127  av_assert0(ctx->frames[input]);
128 
129  src = (cl_mem)ctx->frames[input]->data[plane];
130  av_assert0(src);
131 
132  cle = clSetKernelArg(ctx->kernel, 2 + input, sizeof(cl_mem), &src);
133  if (cle != CL_SUCCESS) {
134  av_log(avctx, AV_LOG_ERROR, "Failed to set kernel "
135  "source image argument %d: %d.\n", input, cle);
136  err = AVERROR_UNKNOWN;
137  goto fail;
138  }
139  }
140 
141  err = ff_opencl_filter_work_size_from_image(avctx, global_work,
142  output, plane, 0);
143  if (err < 0)
144  goto fail;
145 
146  av_log(avctx, AV_LOG_DEBUG, "Run kernel on plane %d "
147  "(%zux%zu).\n", plane, global_work[0], global_work[1]);
148 
149  cle = clEnqueueNDRangeKernel(ctx->command_queue, ctx->kernel, 2, NULL,
150  global_work, NULL, 0, NULL, NULL);
151  CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to enqueue kernel: %d.\n", cle);
152  }
153 
154  cle = clFinish(ctx->command_queue);
155  CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to finish command queue: %d.\n", cle);
156 
157  if (ctx->nb_inputs > 0) {
158  err = av_frame_copy_props(output, ctx->frames[0]);
159  if (err < 0)
160  goto fail;
161  } else {
162  output->pts = ctx->index;
163  }
164  ++ctx->index;
165 
166  av_log(ctx, AV_LOG_DEBUG, "Filter output: %s, %ux%u (%"PRId64").\n",
167  av_get_pix_fmt_name(output->format),
168  output->width, output->height, output->pts);
169 
170  return ff_filter_frame(outlink, output);
171 
172 fail:
173  clFinish(ctx->command_queue);
174  av_frame_free(&output);
175  return err;
176 }
177 
179 {
180  AVFilterContext *avctx = outlink->src;
181 
182  return program_opencl_run(avctx);
183 }
184 
186 {
187  AVFilterContext *avctx = fs->parent;
188  ProgramOpenCLContext *ctx = avctx->priv;
189  int err, i;
190 
191  for (i = 0; i < ctx->nb_inputs; i++) {
192  err = ff_framesync_get_frame(&ctx->fs, i, &ctx->frames[i], 0);
193  if (err < 0)
194  return err;
195  }
196 
197  return program_opencl_run(avctx);
198 }
199 
201 {
202  ProgramOpenCLContext *ctx = avctx->priv;
203 
204  av_assert0(ctx->nb_inputs > 0);
205 
206  return ff_framesync_activate(&ctx->fs);
207 }
208 
210 {
211  AVFilterContext *avctx = outlink->src;
212  ProgramOpenCLContext *ctx = avctx->priv;
213  int err;
214 
215  err = ff_opencl_filter_config_output(outlink);
216  if (err < 0)
217  return err;
218 
219  if (ctx->nb_inputs > 0) {
220  FFFrameSyncIn *in;
221  int i;
222 
223  err = ff_framesync_init(&ctx->fs, avctx, ctx->nb_inputs);
224  if (err < 0)
225  return err;
226 
227  ctx->fs.opaque = ctx;
229 
230  in = ctx->fs.in;
231  for (i = 0; i < ctx->nb_inputs; i++) {
232  const AVFilterLink *inlink = avctx->inputs[i];
233 
234  in[i].time_base = inlink->time_base;
235  in[i].sync = 1;
236  in[i].before = EXT_STOP;
237  in[i].after = EXT_INFINITY;
238  }
239 
240  err = ff_framesync_configure(&ctx->fs);
241  if (err < 0)
242  return err;
243 
244  } else {
245  outlink->time_base = av_inv_q(ctx->source_rate);
246  }
247 
248  return 0;
249 }
250 
252 {
253  ProgramOpenCLContext *ctx = avctx->priv;
254  int err;
255 
256  ff_opencl_filter_init(avctx);
257 
258  ctx->ocf.output_width = ctx->width;
259  ctx->ocf.output_height = ctx->height;
260 
261  if (!strcmp(avctx->filter->name, "openclsrc")) {
262  if (!ctx->ocf.output_width || !ctx->ocf.output_height) {
263  av_log(avctx, AV_LOG_ERROR, "OpenCL source requires output "
264  "dimensions to be specified.\n");
265  return AVERROR(EINVAL);
266  }
267 
268  ctx->nb_inputs = 0;
269  ctx->ocf.output_format = ctx->source_format;
270  } else {
271  int i;
272 
273  ctx->frames = av_mallocz_array(ctx->nb_inputs,
274  sizeof(*ctx->frames));
275  if (!ctx->frames)
276  return AVERROR(ENOMEM);
277 
278  for (i = 0; i < ctx->nb_inputs; i++) {
279  AVFilterPad input;
280  memset(&input, 0, sizeof(input));
281 
282  input.type = AVMEDIA_TYPE_VIDEO;
283  input.name = av_asprintf("input%d", i);
284  if (!input.name)
285  return AVERROR(ENOMEM);
286 
288 
289  err = ff_insert_inpad(avctx, i, &input);
290  if (err < 0) {
291  av_freep(&input.name);
292  return err;
293  }
294  }
295  }
296 
297  return 0;
298 }
299 
301 {
302  ProgramOpenCLContext *ctx = avctx->priv;
303  cl_int cle;
304  int i;
305 
306  if (ctx->nb_inputs > 0) {
307  ff_framesync_uninit(&ctx->fs);
308 
309  av_freep(&ctx->frames);
310  for (i = 0; i < avctx->nb_inputs; i++)
311  av_freep(&avctx->input_pads[i].name);
312  }
313 
314  if (ctx->kernel) {
315  cle = clReleaseKernel(ctx->kernel);
316  if (cle != CL_SUCCESS)
317  av_log(avctx, AV_LOG_ERROR, "Failed to release "
318  "kernel: %d.\n", cle);
319  }
320 
321  if (ctx->command_queue) {
322  cle = clReleaseCommandQueue(ctx->command_queue);
323  if (cle != CL_SUCCESS)
324  av_log(avctx, AV_LOG_ERROR, "Failed to release "
325  "command queue: %d.\n", cle);
326  }
327 
329 }
330 
331 #define OFFSET(x) offsetof(ProgramOpenCLContext, x)
332 #define FLAGS (AV_OPT_FLAG_VIDEO_PARAM | AV_OPT_FLAG_FILTERING_PARAM)
333 
334 #if CONFIG_PROGRAM_OPENCL_FILTER
335 
336 static const AVOption program_opencl_options[] = {
337  { "source", "OpenCL program source file", OFFSET(source_file),
338  AV_OPT_TYPE_STRING, { .str = NULL }, .flags = FLAGS },
339  { "kernel", "Kernel name in program", OFFSET(kernel_name),
340  AV_OPT_TYPE_STRING, { .str = NULL }, .flags = FLAGS },
341 
342  { "inputs", "Number of inputs", OFFSET(nb_inputs),
343  AV_OPT_TYPE_INT, { .i64 = 1 }, 1, INT_MAX, FLAGS },
344 
345  { "size", "Video size", OFFSET(width),
346  AV_OPT_TYPE_IMAGE_SIZE, { .str = NULL }, 0, 0, FLAGS },
347  { "s", "Video size", OFFSET(width),
348  AV_OPT_TYPE_IMAGE_SIZE, { .str = NULL }, 0, 0, FLAGS },
349 
350  { NULL },
351 };
352 
354 
355 static const AVFilterPad program_opencl_outputs[] = {
356  {
357  .name = "default",
358  .type = AVMEDIA_TYPE_VIDEO,
359  .config_props = &program_opencl_config_output,
360  },
361  { NULL }
362 };
363 
365  .name = "program_opencl",
366  .description = NULL_IF_CONFIG_SMALL("Filter video using an OpenCL program"),
367  .priv_size = sizeof(ProgramOpenCLContext),
368  .priv_class = &program_opencl_class,
369  .preinit = &program_opencl_framesync_preinit,
374  .inputs = NULL,
375  .outputs = program_opencl_outputs,
376  .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE,
377 };
378 
379 #endif
380 
381 #if CONFIG_OPENCLSRC_FILTER
382 
383 static const AVOption openclsrc_options[] = {
384  { "source", "OpenCL program source file", OFFSET(source_file),
385  AV_OPT_TYPE_STRING, { .str = NULL }, .flags = FLAGS },
386  { "kernel", "Kernel name in program", OFFSET(kernel_name),
387  AV_OPT_TYPE_STRING, { .str = NULL }, .flags = FLAGS },
388 
389  { "size", "Video size", OFFSET(width),
390  AV_OPT_TYPE_IMAGE_SIZE, { .str = NULL }, 0, 0, FLAGS },
391  { "s", "Video size", OFFSET(width),
392  AV_OPT_TYPE_IMAGE_SIZE, { .str = NULL }, 0, 0, FLAGS },
393 
394  { "format", "Video format", OFFSET(source_format),
395  AV_OPT_TYPE_PIXEL_FMT, { .i64 = AV_PIX_FMT_NONE }, -1, INT_MAX, FLAGS },
396 
397  { "rate", "Video frame rate", OFFSET(source_rate),
398  AV_OPT_TYPE_VIDEO_RATE, { .str = "25" }, 0, INT_MAX, FLAGS },
399  { "r", "Video frame rate", OFFSET(source_rate),
400  AV_OPT_TYPE_VIDEO_RATE, { .str = "25" }, 0, INT_MAX, FLAGS },
401 
402  { NULL },
403 };
404 
405 AVFILTER_DEFINE_CLASS(openclsrc);
406 
407 static const AVFilterPad openclsrc_outputs[] = {
408  {
409  .name = "default",
410  .type = AVMEDIA_TYPE_VIDEO,
411  .config_props = &program_opencl_config_output,
412  .request_frame = &program_opencl_request_frame,
413  },
414  { NULL }
415 };
416 
418  .name = "openclsrc",
419  .description = NULL_IF_CONFIG_SMALL("Generate video using an OpenCL program"),
420  .priv_size = sizeof(ProgramOpenCLContext),
421  .priv_class = &openclsrc_class,
425  .inputs = NULL,
426  .outputs = openclsrc_outputs,
427  .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE,
428 };
429 
430 #endif
int plane
Definition: avisynth_c.h:422
#define NULL
Definition: coverity.c:32
#define FRAMESYNC_DEFINE_CLASS(name, context, field)
Definition: framesync.h:300
#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
static av_cold int program_opencl_init(AVFilterContext *avctx)
This structure describes decoded (raw) audio or video data.
Definition: frame.h:226
static int activate(AVFilterContext *ctx)
Definition: af_adelay.c:237
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
Main libavfilter public API header.
Memory handling functions.
static av_cold int init(AVCodecContext *avctx)
Definition: avrndec.c:35
int(* on_event)(struct FFFrameSync *fs)
Callback called when a frame event is ready.
Definition: framesync.h:172
OpenCLFilterContext ocf
#define FLAGS
enum AVMediaType type
AVFilterPad type.
Definition: internal.h:65
int ff_framesync_configure(FFFrameSync *fs)
Configure a frame sync structure.
Definition: framesync.c:117
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
enum FFFrameSyncExtMode before
Extrapolation mode for timestamps before the first frame.
Definition: framesync.h:86
AVOpenCLDeviceContext * hwctx
Definition: opencl.h:40
const char * name
Pad name.
Definition: internal.h:60
AVFilterContext * parent
Parent filter context.
Definition: framesync.h:152
AVFilterLink ** inputs
array of pointers to input links
Definition: avfilter.h:346
#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:1080
#define av_cold
Definition: attributes.h:82
static av_cold int uninit(AVCodecContext *avctx)
Definition: crystalhd.c:279
#define fs(width, name, subs,...)
Definition: cbs_vp9.c:259
AVOptions.
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:319
FFFrameSyncIn * in
Pointer to array of inputs.
Definition: framesync.h:203
static int program_opencl_run(AVFilterContext *avctx)
enum FFFrameSyncExtMode after
Extrapolation mode for timestamps after the last frame.
Definition: framesync.h:91
Input stream structure.
Definition: framesync.h:81
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
AVFilterPad * input_pads
array of input pads
Definition: avfilter.h:345
static int program_opencl_request_frame(AVFilterLink *outlink)
int width
Definition: frame.h:284
#define AV_LOG_ERROR
Something went wrong and cannot losslessly be recovered.
Definition: log.h:176
static int program_opencl_activate(AVFilterContext *avctx)
cl_command_queue command_queue
void ff_framesync_uninit(FFFrameSync *fs)
Free all memory currently allocated.
Definition: framesync.c:293
Frame sync structure.
Definition: framesync.h:146
#define AVERROR(e)
Definition: error.h:43
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
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
static av_cold void program_opencl_uninit(AVFilterContext *avctx)
AVRational time_base
Time base for the incoming frames.
Definition: framesync.h:96
int ff_framesync_activate(FFFrameSync *fs)
Examine the frames in the filter's input and try to produce output.
Definition: framesync.c:344
#define fail()
Definition: checkasm.h:117
enum AVPixelFormat output_format
Definition: opencl.h:44
char * av_asprintf(const char *fmt,...)
Definition: avstring.c:113
unsigned nb_inputs
number of input pads
Definition: avfilter.h:347
AVFilter ff_vsrc_openclsrc
#define width
AVFormatContext * ctx
Definition: movenc.c:48
enum AVPixelFormat source_format
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
void * opaque
Opaque pointer, not used by the API.
Definition: framesync.h:177
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:299
Extend the frame to infinity.
Definition: framesync.h:75
int ff_framesync_init(FFFrameSync *fs, AVFilterContext *parent, unsigned nb_in)
Initialize a frame sync structure.
Definition: framesync.c:77
uint8_t pi<< 24) CONV_FUNC_GROUP(AV_SAMPLE_FMT_FLT, float, AV_SAMPLE_FMT_U8, uint8_t,(*(constuint8_t *) pi-0x80)*(1.0f/(1<< 7))) CONV_FUNC_GROUP(AV_SAMPLE_FMT_DBL, double, AV_SAMPLE_FMT_U8, uint8_t,(*(constuint8_t *) pi-0x80)*(1.0/(1<< 7))) CONV_FUNC_GROUP(AV_SAMPLE_FMT_U8, uint8_t, AV_SAMPLE_FMT_S16, int16_t,(*(constint16_t *) pi >>8)+0x80) CONV_FUNC_GROUP(AV_SAMPLE_FMT_FLT, float, AV_SAMPLE_FMT_S16, int16_t,*(constint16_t *) pi *(1.0f/(1<< 15))) CONV_FUNC_GROUP(AV_SAMPLE_FMT_DBL, double, AV_SAMPLE_FMT_S16, int16_t,*(constint16_t *) pi *(1.0/(1<< 15))) CONV_FUNC_GROUP(AV_SAMPLE_FMT_U8, uint8_t, AV_SAMPLE_FMT_S32, int32_t,(*(constint32_t *) pi >>24)+0x80) CONV_FUNC_GROUP(AV_SAMPLE_FMT_FLT, float, AV_SAMPLE_FMT_S32, int32_t,*(constint32_t *) pi *(1.0f/(1U<< 31))) CONV_FUNC_GROUP(AV_SAMPLE_FMT_DBL, double, AV_SAMPLE_FMT_S32, int32_t,*(constint32_t *) pi *(1.0/(1U<< 31))) CONV_FUNC_GROUP(AV_SAMPLE_FMT_U8, uint8_t, AV_SAMPLE_FMT_FLT, float, av_clip_uint8(lrintf(*(constfloat *) pi *(1<< 7))+0x80)) CONV_FUNC_GROUP(AV_SAMPLE_FMT_S16, int16_t, AV_SAMPLE_FMT_FLT, float, av_clip_int16(lrintf(*(constfloat *) pi *(1<< 15)))) CONV_FUNC_GROUP(AV_SAMPLE_FMT_S32, int32_t, AV_SAMPLE_FMT_FLT, float, av_clipl_int32(llrintf(*(constfloat *) pi *(1U<< 31)))) CONV_FUNC_GROUP(AV_SAMPLE_FMT_U8, uint8_t, AV_SAMPLE_FMT_DBL, double, av_clip_uint8(lrint(*(constdouble *) pi *(1<< 7))+0x80)) CONV_FUNC_GROUP(AV_SAMPLE_FMT_S16, int16_t, AV_SAMPLE_FMT_DBL, double, av_clip_int16(lrint(*(constdouble *) pi *(1<< 15)))) CONV_FUNC_GROUP(AV_SAMPLE_FMT_S32, int32_t, AV_SAMPLE_FMT_DBL, double, av_clipl_int32(llrint(*(constdouble *) pi *(1U<< 31))))#defineSET_CONV_FUNC_GROUP(ofmt, ifmt) staticvoidset_generic_function(AudioConvert *ac){}voidff_audio_convert_free(AudioConvert **ac){if(!*ac) return;ff_dither_free(&(*ac) ->dc);av_freep(ac);}AudioConvert *ff_audio_convert_alloc(AVAudioResampleContext *avr, enumAVSampleFormatout_fmt, enumAVSampleFormatin_fmt, intchannels, intsample_rate, intapply_map){AudioConvert *ac;intin_planar, out_planar;ac=av_mallocz(sizeof(*ac));if(!ac) returnNULL;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);returnNULL;}returnac;}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;}elseif(in_planar) ac->func_type=CONV_FUNC_TYPE_INTERLEAVE;elseac->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);returnac;}intff_audio_convert(AudioConvert *ac, AudioData *out, AudioData *in){intuse_generic=1;intlen=in->nb_samples;intp;if(ac->dc){av_log(ac->avr, AV_LOG_TRACE,"%dsamples-audio_convert:%sto%s(dithered)\n", len, av_get_sample_fmt_name(ac->in_fmt), av_get_sample_fmt_name(ac->out_fmt));returnff_convert_dither(ac-> in
unsigned sync
Synchronization level: frames on input at the highest sync level will generate output frame events...
Definition: framesync.h:139
Filter definition.
Definition: avfilter.h:144
Rational number (pair of numerator and denominator).
Definition: rational.h:58
offset must point to AVRational
Definition: opt.h:236
const char * name
Filter name.
Definition: avfilter.h:148
offset must point to two consecutive integers
Definition: opt.h:233
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:68
uint8_t * data[AV_NUM_DATA_POINTERS]
pointer to the picture/channel planes.
Definition: frame.h:240
static int program_opencl_load(AVFilterContext *avctx)
static av_always_inline AVRational av_inv_q(AVRational q)
Invert a rational.
Definition: rational.h:159
static int query_formats(AVFilterContext *ctx)
Definition: aeval.c:244
if(ret< 0)
Definition: vf_mcdeint.c:279
AVFilter ff_vf_program_opencl
void ff_opencl_filter_uninit(AVFilterContext *avctx)
Uninitialise an OpenCL filter context.
Definition: opencl.c:156
#define AVERROR_UNKNOWN
Unknown error, typically from an external library.
Definition: error.h:71
Completely stop all streams with this one.
Definition: framesync.h:65
#define AVFILTER_DEFINE_CLASS(fname)
Definition: internal.h:334
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
#define OFFSET(x)
int height
Definition: frame.h:284
#define av_freep(p)
int(* config_props)(AVFilterLink *link)
Link configuration callback.
Definition: internal.h:129
cl_program program
Definition: opencl.h:42
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:2362
internal API functions
int ff_framesync_get_frame(FFFrameSync *fs, unsigned in, AVFrame **rframe, unsigned get)
Get the current frame in an input.
Definition: framesync.c:256
static int program_opencl_filter(FFFrameSync *fs)
AVPixelFormat
Pixel format.
Definition: pixfmt.h:64
const AVFilter * filter
the AVFilter of which this is an instance
Definition: avfilter.h:341
static int program_opencl_config_output(AVFilterLink *outlink)
int av_frame_copy_props(AVFrame *dst, const AVFrame *src)
Copy only "metadata" fields from src to dst.
Definition: frame.c:654
void * av_mallocz_array(size_t nmemb, size_t size)
Definition: mem.c:191
static int ff_insert_inpad(AVFilterContext *f, unsigned index, AVFilterPad *p)
Insert a new input pad for the filter.
Definition: internal.h:277