FFmpeg  4.0
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  if (cle != CL_SUCCESS) {
152  av_log(avctx, AV_LOG_ERROR, "Failed to enqueue kernel: %d.\n",
153  cle);
154  err = AVERROR(EIO);
155  goto fail;
156  }
157  }
158 
159  cle = clFinish(ctx->command_queue);
160  if (cle != CL_SUCCESS) {
161  av_log(avctx, AV_LOG_ERROR, "Failed to finish command queue: %d.\n",
162  cle);
163  err = AVERROR(EIO);
164  goto fail;
165  }
166 
167  if (ctx->nb_inputs > 0) {
168  err = av_frame_copy_props(output, ctx->frames[0]);
169  if (err < 0)
170  goto fail;
171  } else {
172  output->pts = ctx->index;
173  }
174  ++ctx->index;
175 
176  av_log(ctx, AV_LOG_DEBUG, "Filter output: %s, %ux%u (%"PRId64").\n",
177  av_get_pix_fmt_name(output->format),
178  output->width, output->height, output->pts);
179 
180  return ff_filter_frame(outlink, output);
181 
182 fail:
183  clFinish(ctx->command_queue);
184  av_frame_free(&output);
185  return err;
186 }
187 
189 {
190  AVFilterContext *avctx = outlink->src;
191 
192  return program_opencl_run(avctx);
193 }
194 
196 {
197  AVFilterContext *avctx = fs->parent;
198  ProgramOpenCLContext *ctx = avctx->priv;
199  int err, i;
200 
201  for (i = 0; i < ctx->nb_inputs; i++) {
202  err = ff_framesync_get_frame(&ctx->fs, i, &ctx->frames[i], 0);
203  if (err < 0)
204  return err;
205  }
206 
207  return program_opencl_run(avctx);
208 }
209 
211 {
212  ProgramOpenCLContext *ctx = avctx->priv;
213 
214  av_assert0(ctx->nb_inputs > 0);
215 
216  return ff_framesync_activate(&ctx->fs);
217 }
218 
220 {
221  AVFilterContext *avctx = outlink->src;
222  ProgramOpenCLContext *ctx = avctx->priv;
223  int err;
224 
225  err = ff_opencl_filter_config_output(outlink);
226  if (err < 0)
227  return err;
228 
229  if (ctx->nb_inputs > 0) {
230  FFFrameSyncIn *in;
231  int i;
232 
233  err = ff_framesync_init(&ctx->fs, avctx, ctx->nb_inputs);
234  if (err < 0)
235  return err;
236 
237  ctx->fs.opaque = ctx;
239 
240  in = ctx->fs.in;
241  for (i = 0; i < ctx->nb_inputs; i++) {
242  const AVFilterLink *inlink = avctx->inputs[i];
243 
244  in[i].time_base = inlink->time_base;
245  in[i].sync = 1;
246  in[i].before = EXT_STOP;
247  in[i].after = EXT_INFINITY;
248  }
249 
250  err = ff_framesync_configure(&ctx->fs);
251  if (err < 0)
252  return err;
253 
254  } else {
255  outlink->time_base = av_inv_q(ctx->source_rate);
256  }
257 
258  return 0;
259 }
260 
262 {
263  ProgramOpenCLContext *ctx = avctx->priv;
264  int err;
265 
266  ff_opencl_filter_init(avctx);
267 
268  ctx->ocf.output_width = ctx->width;
269  ctx->ocf.output_height = ctx->height;
270 
271  if (!strcmp(avctx->filter->name, "openclsrc")) {
272  if (!ctx->ocf.output_width || !ctx->ocf.output_height) {
273  av_log(avctx, AV_LOG_ERROR, "OpenCL source requires output "
274  "dimensions to be specified.\n");
275  return AVERROR(EINVAL);
276  }
277 
278  ctx->nb_inputs = 0;
279  ctx->ocf.output_format = ctx->source_format;
280  } else {
281  int i;
282 
283  ctx->frames = av_mallocz_array(ctx->nb_inputs,
284  sizeof(*ctx->frames));
285  if (!ctx->frames)
286  return AVERROR(ENOMEM);
287 
288  for (i = 0; i < ctx->nb_inputs; i++) {
289  AVFilterPad input;
290  memset(&input, 0, sizeof(input));
291 
292  input.type = AVMEDIA_TYPE_VIDEO;
293  input.name = av_asprintf("input%d", i);
294  if (!input.name)
295  return AVERROR(ENOMEM);
296 
298 
299  err = ff_insert_inpad(avctx, i, &input);
300  if (err < 0) {
301  av_freep(&input.name);
302  return err;
303  }
304  }
305  }
306 
307  return 0;
308 }
309 
311 {
312  ProgramOpenCLContext *ctx = avctx->priv;
313  cl_int cle;
314  int i;
315 
316  if (ctx->nb_inputs > 0) {
317  ff_framesync_uninit(&ctx->fs);
318 
319  av_freep(&ctx->frames);
320  for (i = 0; i < avctx->nb_inputs; i++)
321  av_freep(&avctx->input_pads[i].name);
322  }
323 
324  if (ctx->kernel) {
325  cle = clReleaseKernel(ctx->kernel);
326  if (cle != CL_SUCCESS)
327  av_log(avctx, AV_LOG_ERROR, "Failed to release "
328  "kernel: %d.\n", cle);
329  }
330 
331  if (ctx->command_queue) {
332  cle = clReleaseCommandQueue(ctx->command_queue);
333  if (cle != CL_SUCCESS)
334  av_log(avctx, AV_LOG_ERROR, "Failed to release "
335  "command queue: %d.\n", cle);
336  }
337 
339 }
340 
341 #define OFFSET(x) offsetof(ProgramOpenCLContext, x)
342 #define FLAGS (AV_OPT_FLAG_VIDEO_PARAM | AV_OPT_FLAG_FILTERING_PARAM)
343 
344 #if CONFIG_PROGRAM_OPENCL_FILTER
345 
346 static const AVOption program_opencl_options[] = {
347  { "source", "OpenCL program source file", OFFSET(source_file),
348  AV_OPT_TYPE_STRING, { .str = NULL }, .flags = FLAGS },
349  { "kernel", "Kernel name in program", OFFSET(kernel_name),
350  AV_OPT_TYPE_STRING, { .str = NULL }, .flags = FLAGS },
351 
352  { "inputs", "Number of inputs", OFFSET(nb_inputs),
353  AV_OPT_TYPE_INT, { .i64 = 1 }, 1, INT_MAX, FLAGS },
354 
355  { "size", "Video size", OFFSET(width),
356  AV_OPT_TYPE_IMAGE_SIZE, { .str = NULL }, 0, 0, FLAGS },
357  { "s", "Video size", OFFSET(width),
358  AV_OPT_TYPE_IMAGE_SIZE, { .str = NULL }, 0, 0, FLAGS },
359 
360  { NULL },
361 };
362 
364 
365 static const AVFilterPad program_opencl_outputs[] = {
366  {
367  .name = "default",
368  .type = AVMEDIA_TYPE_VIDEO,
369  .config_props = &program_opencl_config_output,
370  },
371  { NULL }
372 };
373 
375  .name = "program_opencl",
376  .description = NULL_IF_CONFIG_SMALL("Filter video using an OpenCL program"),
377  .priv_size = sizeof(ProgramOpenCLContext),
378  .priv_class = &program_opencl_class,
379  .preinit = &program_opencl_framesync_preinit,
384  .inputs = NULL,
385  .outputs = program_opencl_outputs,
386  .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE,
387 };
388 
389 #endif
390 
391 #if CONFIG_OPENCLSRC_FILTER
392 
393 static const AVOption openclsrc_options[] = {
394  { "source", "OpenCL program source file", OFFSET(source_file),
395  AV_OPT_TYPE_STRING, { .str = NULL }, .flags = FLAGS },
396  { "kernel", "Kernel name in program", OFFSET(kernel_name),
397  AV_OPT_TYPE_STRING, { .str = NULL }, .flags = FLAGS },
398 
399  { "size", "Video size", OFFSET(width),
400  AV_OPT_TYPE_IMAGE_SIZE, { .str = NULL }, 0, 0, FLAGS },
401  { "s", "Video size", OFFSET(width),
402  AV_OPT_TYPE_IMAGE_SIZE, { .str = NULL }, 0, 0, FLAGS },
403 
404  { "format", "Video format", OFFSET(source_format),
405  AV_OPT_TYPE_PIXEL_FMT, { .i64 = AV_PIX_FMT_NONE }, -1, INT_MAX, FLAGS },
406 
407  { "rate", "Video frame rate", OFFSET(source_rate),
408  AV_OPT_TYPE_VIDEO_RATE, { .str = "25" }, 0, INT_MAX, FLAGS },
409  { "r", "Video frame rate", OFFSET(source_rate),
410  AV_OPT_TYPE_VIDEO_RATE, { .str = "25" }, 0, INT_MAX, FLAGS },
411 
412  { NULL },
413 };
414 
415 AVFILTER_DEFINE_CLASS(openclsrc);
416 
417 static const AVFilterPad openclsrc_outputs[] = {
418  {
419  .name = "default",
420  .type = AVMEDIA_TYPE_VIDEO,
421  .config_props = &program_opencl_config_output,
422  .request_frame = &program_opencl_request_frame,
423  },
424  { NULL }
425 };
426 
428  .name = "openclsrc",
429  .description = NULL_IF_CONFIG_SMALL("Generate video using an OpenCL program"),
430  .priv_size = sizeof(ProgramOpenCLContext),
431  .priv_class = &openclsrc_class,
435  .inputs = NULL,
436  .outputs = openclsrc_outputs,
437  .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE,
438 };
439 
440 #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:218
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:281
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:63
int ff_opencl_filter_query_formats(AVFilterContext *avctx)
Return that all inputs and outputs support only AV_PIX_FMT_OPENCL.
Definition: opencl.c:31
Main libavfilter public API header.
Memory handling functions.
static av_cold int init(AVCodecContext *avctx)
Definition: avrndec.c:35
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:150
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
static int activate(AVFilterContext *ctx)
Definition: af_amix.c:421
#define av_cold
Definition: attributes.h:82
static av_cold int uninit(AVCodecContext *avctx)
Definition: crystalhd.c:279
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:311
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:99
#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:276
#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&#39;s input and try to produce output.
Definition: framesync.c:344
int(* on_event)(struct FFFrameSync *fs)
Callback called when a frame event is ready.
Definition: framesync.h:172
#define fail()
Definition: checkasm.h:116
enum AVPixelFormat output_format
Definition: opencl.h:44
char * av_asprintf(const char *fmt,...)
Definition: avstring.c:113
int(* config_props)(AVFilterLink *link)
Link configuration callback.
Definition: internal.h:129
unsigned nb_inputs
number of input pads
Definition: avfilter.h:347
AVFilter ff_vsrc_openclsrc
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:222
static const AVFilterPad inputs[]
Definition: af_acontrast.c:193
void * opaque
Opaque pointer, not used by the API.
Definition: framesync.h:177
if(ret< 0)
Definition: vf_mcdeint.c:279
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:291
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,(*(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
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
uint8_t * data[AV_NUM_DATA_POINTERS]
pointer to the picture/channel planes.
Definition: frame.h:232
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
AVFilter ff_vf_program_opencl
void ff_opencl_filter_uninit(AVFilterContext *avctx)
Uninitialise an OpenCL filter context.
Definition: opencl.c:159
#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:276
#define av_freep(p)
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:2279
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:60
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:652
void * av_mallocz_array(size_t nmemb, size_t size)
Allocate a memory block for an array with av_mallocz().
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