FFmpeg
 All Data Structures Namespaces Files Functions Variables Typedefs Enumerations Enumerator Macros Groups Pages
vf_thumbnail_cuda.c
Go to the documentation of this file.
1 /*
2 * Copyright (c) 2017, NVIDIA CORPORATION. All rights reserved.
3 *
4 * Permission is hereby granted, free of charge, to any person obtaining a
5 * copy of this software and associated documentation files (the "Software"),
6 * to deal in the Software without restriction, including without limitation
7 * the rights to use, copy, modify, merge, publish, distribute, sublicense,
8 * and/or sell copies of the Software, and to permit persons to whom the
9 * Software is furnished to do so, subject to the following conditions:
10 *
11 * The above copyright notice and this permission notice shall be included in
12 * all copies or substantial portions of the Software.
13 *
14 * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
15 * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
16 * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL
17 * THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
18 * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
19 * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER
20 * DEALINGS IN THE SOFTWARE.
21 */
22 
23 #include <cuda.h>
24 
25 #include "libavutil/hwcontext.h"
27 #include "libavutil/opt.h"
28 #include "libavutil/pixdesc.h"
29 
30 #include "avfilter.h"
31 #include "internal.h"
32 
33 #define HIST_SIZE (3*256)
34 #define DIV_UP(a, b) ( ((a) + (b) - 1) / (b) )
35 #define BLOCKX 32
36 #define BLOCKY 16
37 
38 static const enum AVPixelFormat supported_formats[] = {
45 };
46 
47 struct thumb_frame {
48  AVFrame *buf; ///< cached frame
49  int histogram[HIST_SIZE]; ///< RGB color distribution histogram of the frame
50 };
51 
52 typedef struct ThumbnailCudaContext {
53  const AVClass *class;
54  int n; ///< current frame
55  int n_frames; ///< number of frames for analysis
56  struct thumb_frame *frames; ///< the n_frames frames
57  AVRational tb; ///< copy of the input timebase to ease access
58 
60 
61  CUmodule cu_module;
62 
63  CUfunction cu_func_uchar;
64  CUfunction cu_func_uchar2;
65  CUfunction cu_func_ushort;
66  CUfunction cu_func_ushort2;
67  CUtexref cu_tex_uchar;
68  CUtexref cu_tex_uchar2;
69  CUtexref cu_tex_ushort;
70  CUtexref cu_tex_ushort2;
71 
72  CUdeviceptr data;
74 
75 #define OFFSET(x) offsetof(ThumbnailCudaContext, x)
76 #define FLAGS AV_OPT_FLAG_VIDEO_PARAM|AV_OPT_FLAG_FILTERING_PARAM
77 
78 static const AVOption thumbnail_cuda_options[] = {
79  { "n", "set the frames batch size", OFFSET(n_frames), AV_OPT_TYPE_INT, {.i64=100}, 2, INT_MAX, FLAGS },
80  { NULL }
81 };
82 
83 AVFILTER_DEFINE_CLASS(thumbnail_cuda);
84 
86 {
87  ThumbnailCudaContext *s = ctx->priv;
88 
89  s->frames = av_calloc(s->n_frames, sizeof(*s->frames));
90  if (!s->frames) {
91  av_log(ctx, AV_LOG_ERROR,
92  "Allocation failure, try to lower the number of frames\n");
93  return AVERROR(ENOMEM);
94  }
95  av_log(ctx, AV_LOG_VERBOSE, "batch size: %d frames\n", s->n_frames);
96  return 0;
97 }
98 
99 /**
100  * @brief Compute Sum-square deviation to estimate "closeness".
101  * @param hist color distribution histogram
102  * @param median average color distribution histogram
103  * @return sum of squared errors
104  */
105 static double frame_sum_square_err(const int *hist, const double *median)
106 {
107  int i;
108  double err, sum_sq_err = 0;
109 
110  for (i = 0; i < HIST_SIZE; i++) {
111  err = median[i] - (double)hist[i];
112  sum_sq_err += err*err;
113  }
114  return sum_sq_err;
115 }
116 
118 {
119  AVFrame *picref;
120  ThumbnailCudaContext *s = ctx->priv;
121  int i, j, best_frame_idx = 0;
122  int nb_frames = s->n;
123  double avg_hist[HIST_SIZE] = {0}, sq_err, min_sq_err = -1;
124 
125  // average histogram of the N frames
126  for (j = 0; j < FF_ARRAY_ELEMS(avg_hist); j++) {
127  for (i = 0; i < nb_frames; i++)
128  avg_hist[j] += (double)s->frames[i].histogram[j];
129  avg_hist[j] /= nb_frames;
130  }
131 
132  // find the frame closer to the average using the sum of squared errors
133  for (i = 0; i < nb_frames; i++) {
134  sq_err = frame_sum_square_err(s->frames[i].histogram, avg_hist);
135  if (i == 0 || sq_err < min_sq_err)
136  best_frame_idx = i, min_sq_err = sq_err;
137  }
138 
139  // free and reset everything (except the best frame buffer)
140  for (i = 0; i < nb_frames; i++) {
141  memset(s->frames[i].histogram, 0, sizeof(s->frames[i].histogram));
142  if (i != best_frame_idx)
143  av_frame_free(&s->frames[i].buf);
144  }
145  s->n = 0;
146 
147  // raise the chosen one
148  picref = s->frames[best_frame_idx].buf;
149  av_log(ctx, AV_LOG_INFO, "frame id #%d (pts_time=%f) selected "
150  "from a set of %d images\n", best_frame_idx,
151  picref->pts * av_q2d(s->tb), nb_frames);
152  s->frames[best_frame_idx].buf = NULL;
153 
154  return picref;
155 }
156 
157 static int thumbnail_kernel(ThumbnailCudaContext *s, CUfunction func, CUtexref tex, int channels,
158  int *histogram, uint8_t *src_dptr, int src_width, int src_height, int src_pitch, int pixel_size)
159 {
160  CUdeviceptr src_devptr = (CUdeviceptr)src_dptr;
161  void *args[] = { &histogram, &src_width, &src_height };
162  CUDA_ARRAY_DESCRIPTOR desc;
163 
164  desc.Width = src_width;
165  desc.Height = src_height;
166  desc.NumChannels = channels;
167  if (pixel_size == 1) {
168  desc.Format = CU_AD_FORMAT_UNSIGNED_INT8;
169  }
170  else {
171  desc.Format = CU_AD_FORMAT_UNSIGNED_INT16;
172  }
173 
174  cuTexRefSetAddress2D_v3(tex, &desc, src_devptr, src_pitch);
175  cuLaunchKernel(func, DIV_UP(src_width, BLOCKX), DIV_UP(src_height, BLOCKY), 1, BLOCKX, BLOCKY, 1, 0, 0, args, NULL);
176 
177  return 0;
178 }
179 
181 {
182  AVHWFramesContext *in_frames_ctx = (AVHWFramesContext*)in->hw_frames_ctx->data;
183  ThumbnailCudaContext *s = ctx->priv;
184 
185  switch (in_frames_ctx->sw_format) {
186  case AV_PIX_FMT_NV12:
187  thumbnail_kernel(s, s->cu_func_uchar, s->cu_tex_uchar, 1,
188  histogram, in->data[0], in->width, in->height, in->linesize[0], 1);
189  thumbnail_kernel(s, s->cu_func_uchar2, s->cu_tex_uchar2, 2,
190  histogram + 256, in->data[1], in->width / 2, in->height / 2, in->linesize[1], 1);
191  break;
192  case AV_PIX_FMT_YUV420P:
193  thumbnail_kernel(s, s->cu_func_uchar, s->cu_tex_uchar, 1,
194  histogram, in->data[0], in->width, in->height, in->linesize[0], 1);
195  thumbnail_kernel(s, s->cu_func_uchar, s->cu_tex_uchar, 1,
196  histogram + 256, in->data[1], in->width / 2, in->height / 2, in->linesize[1], 1);
197  thumbnail_kernel(s, s->cu_func_uchar, s->cu_tex_uchar, 1,
198  histogram + 512, in->data[2], in->width / 2, in->height / 2, in->linesize[2], 1);
199  break;
200  case AV_PIX_FMT_YUV444P:
201  thumbnail_kernel(s, s->cu_func_uchar, s->cu_tex_uchar, 1,
202  histogram, in->data[0], in->width, in->height, in->linesize[0], 1);
203  thumbnail_kernel(s, s->cu_func_uchar, s->cu_tex_uchar, 1,
204  histogram + 256, in->data[1], in->width, in->height, in->linesize[1], 1);
205  thumbnail_kernel(s, s->cu_func_uchar, s->cu_tex_uchar, 1,
206  histogram + 512, in->data[2], in->width, in->height, in->linesize[2], 1);
207  break;
208  case AV_PIX_FMT_P010LE:
209  case AV_PIX_FMT_P016LE:
210  thumbnail_kernel(s, s->cu_func_ushort, s->cu_tex_ushort, 1,
211  histogram, in->data[0], in->width, in->height, in->linesize[0], 2);
212  thumbnail_kernel(s, s->cu_func_ushort2, s->cu_tex_ushort2, 2,
213  histogram + 256, in->data[1], in->width / 2, in->height / 2, in->linesize[1], 2);
214  break;
216  thumbnail_kernel(s, s->cu_func_ushort2, s->cu_tex_uchar, 1,
217  histogram, in->data[0], in->width, in->height, in->linesize[0], 2);
218  thumbnail_kernel(s, s->cu_func_ushort2, s->cu_tex_uchar, 1,
219  histogram + 256, in->data[1], in->width, in->height, in->linesize[1], 2);
220  thumbnail_kernel(s, s->cu_func_ushort2, s->cu_tex_uchar, 1,
221  histogram + 512, in->data[2], in->width, in->height, in->linesize[2], 2);
222  break;
223  default:
224  return AVERROR_BUG;
225  }
226 
227  return 0;
228 }
229 
230 static int filter_frame(AVFilterLink *inlink, AVFrame *frame)
231 {
232  AVFilterContext *ctx = inlink->dst;
233  ThumbnailCudaContext *s = ctx->priv;
234  AVFilterLink *outlink = ctx->outputs[0];
235  int *hist = s->frames[s->n].histogram;
236  AVHWFramesContext *hw_frames_ctx = (AVHWFramesContext*)s->hw_frames_ctx->data;
237  AVCUDADeviceContext *device_hwctx = hw_frames_ctx->device_ctx->hwctx;
238  CUresult err;
239  CUcontext dummy;
240  CUDA_MEMCPY2D cpy = { 0 };
241  int ret = 0;
242 
243  // keep a reference of each frame
244  s->frames[s->n].buf = frame;
245 
246  err = cuCtxPushCurrent(device_hwctx->cuda_ctx);
247  if (err != CUDA_SUCCESS)
248  return AVERROR_UNKNOWN;
249 
250  cuMemsetD8(s->data, 0, HIST_SIZE * sizeof(int));
251 
252  thumbnail(ctx, (int*)s->data, frame);
253 
254  cpy.srcMemoryType = CU_MEMORYTYPE_DEVICE;
255  cpy.dstMemoryType = CU_MEMORYTYPE_HOST;
256  cpy.srcDevice = s->data;
257  cpy.dstHost = hist;
258  cpy.srcPitch = HIST_SIZE * sizeof(int);
259  cpy.dstPitch = HIST_SIZE * sizeof(int);
260  cpy.WidthInBytes = HIST_SIZE * sizeof(int);
261  cpy.Height = 1;
262 
263  err = cuMemcpy2D(&cpy);
264  if (err != CUDA_SUCCESS) {
265  av_log(ctx, AV_LOG_ERROR, "Error transferring the data from the CUDA frame\n");
266  return AVERROR_UNKNOWN;
267  }
268 
269  if (hw_frames_ctx->sw_format == AV_PIX_FMT_NV12 || hw_frames_ctx->sw_format == AV_PIX_FMT_YUV420P ||
270  hw_frames_ctx->sw_format == AV_PIX_FMT_P010LE || hw_frames_ctx->sw_format == AV_PIX_FMT_P016LE)
271  {
272  int i;
273  for (i = 256; i < HIST_SIZE; i++)
274  hist[i] = 4 * hist[i];
275  }
276 
277  cuCtxPopCurrent(&dummy);
278  if (ret < 0)
279  return ret;
280 
281  // no selection until the buffer of N frames is filled up
282  s->n++;
283  if (s->n < s->n_frames)
284  return 0;
285 
286  return ff_filter_frame(outlink, get_best_frame(ctx));
287 }
288 
290 {
291  int i;
292  ThumbnailCudaContext *s = ctx->priv;
293 
294  if (s->data) {
295  cuMemFree(s->data);
296  s->data = 0;
297  }
298 
299  if (s->cu_module) {
300  cuModuleUnload(s->cu_module);
301  s->cu_module = NULL;
302  }
303 
304  for (i = 0; i < s->n_frames && s->frames[i].buf; i++)
305  av_frame_free(&s->frames[i].buf);
306  av_freep(&s->frames);
307 }
308 
309 static int request_frame(AVFilterLink *link)
310 {
311  AVFilterContext *ctx = link->src;
312  ThumbnailCudaContext *s = ctx->priv;
313  int ret = ff_request_frame(ctx->inputs[0]);
314 
315  if (ret == AVERROR_EOF && s->n) {
316  ret = ff_filter_frame(link, get_best_frame(ctx));
317  if (ret < 0)
318  return ret;
319  ret = AVERROR_EOF;
320  }
321  if (ret < 0)
322  return ret;
323  return 0;
324 }
325 
327 {
328  int i;
329 
330  for (i = 0; i < FF_ARRAY_ELEMS(supported_formats); i++)
331  if (supported_formats[i] == fmt)
332  return 1;
333  return 0;
334 }
335 
336 static int config_props(AVFilterLink *inlink)
337 {
338  AVFilterContext *ctx = inlink->dst;
339  ThumbnailCudaContext *s = ctx->priv;
340  AVHWFramesContext *hw_frames_ctx = (AVHWFramesContext*)inlink->hw_frames_ctx->data;
341  AVCUDADeviceContext *device_hwctx = hw_frames_ctx->device_ctx->hwctx;
342  CUcontext dummy, cuda_ctx = device_hwctx->cuda_ctx;
343  CUresult err;
344 
345  extern char vf_thumbnail_cuda_ptx[];
346 
347  err = cuCtxPushCurrent(cuda_ctx);
348  if (err != CUDA_SUCCESS) {
349  av_log(ctx, AV_LOG_ERROR, "Error pushing cuda context\n");
350  return AVERROR_UNKNOWN;
351  }
352 
353  err = cuModuleLoadData(&s->cu_module, vf_thumbnail_cuda_ptx);
354  if (err != CUDA_SUCCESS) {
355  av_log(ctx, AV_LOG_ERROR, "Error loading module data\n");
356  return AVERROR_UNKNOWN;
357  }
358 
359  cuModuleGetFunction(&s->cu_func_uchar, s->cu_module, "Thumbnail_uchar");
360  cuModuleGetFunction(&s->cu_func_uchar2, s->cu_module, "Thumbnail_uchar2");
361  cuModuleGetFunction(&s->cu_func_ushort, s->cu_module, "Thumbnail_ushort");
362  cuModuleGetFunction(&s->cu_func_ushort2, s->cu_module, "Thumbnail_ushort2");
363 
364  cuModuleGetTexRef(&s->cu_tex_uchar, s->cu_module, "uchar_tex");
365  cuModuleGetTexRef(&s->cu_tex_uchar2, s->cu_module, "uchar2_tex");
366  cuModuleGetTexRef(&s->cu_tex_ushort, s->cu_module, "ushort_tex");
367  cuModuleGetTexRef(&s->cu_tex_ushort2, s->cu_module, "ushort2_tex");
368 
369  cuTexRefSetFlags(s->cu_tex_uchar, CU_TRSF_READ_AS_INTEGER);
370  cuTexRefSetFlags(s->cu_tex_uchar2, CU_TRSF_READ_AS_INTEGER);
371  cuTexRefSetFlags(s->cu_tex_ushort, CU_TRSF_READ_AS_INTEGER);
372  cuTexRefSetFlags(s->cu_tex_ushort2, CU_TRSF_READ_AS_INTEGER);
373 
374  cuTexRefSetFilterMode(s->cu_tex_uchar, CU_TR_FILTER_MODE_LINEAR);
375  cuTexRefSetFilterMode(s->cu_tex_uchar2, CU_TR_FILTER_MODE_LINEAR);
376  cuTexRefSetFilterMode(s->cu_tex_ushort, CU_TR_FILTER_MODE_LINEAR);
377  cuTexRefSetFilterMode(s->cu_tex_ushort2, CU_TR_FILTER_MODE_LINEAR);
378 
379  err = cuMemAlloc(&s->data, HIST_SIZE * sizeof(int));
380  if (err != CUDA_SUCCESS) {
381  av_log(ctx, AV_LOG_ERROR, "Error allocating cuda memory\n");
382  return AVERROR_UNKNOWN;
383  }
384 
385  cuCtxPopCurrent(&dummy);
386 
387  s->hw_frames_ctx = ctx->inputs[0]->hw_frames_ctx;
388 
390  if (!ctx->outputs[0]->hw_frames_ctx)
391  return AVERROR(ENOMEM);
392 
393  s->tb = inlink->time_base;
394 
395  if (!format_is_supported(hw_frames_ctx->sw_format)) {
396  av_log(ctx, AV_LOG_ERROR, "Unsupported input format: %s\n", av_get_pix_fmt_name(hw_frames_ctx->sw_format));
397  return AVERROR(ENOSYS);
398  }
399 
400  return 0;
401 }
402 
404 {
405  static const enum AVPixelFormat pix_fmts[] = {
408  };
409  AVFilterFormats *fmts_list = ff_make_format_list(pix_fmts);
410  if (!fmts_list)
411  return AVERROR(ENOMEM);
412  return ff_set_common_formats(ctx, fmts_list);
413 }
414 
416  {
417  .name = "default",
418  .type = AVMEDIA_TYPE_VIDEO,
419  .config_props = config_props,
420  .filter_frame = filter_frame,
421  },
422  { NULL }
423 };
424 
426  {
427  .name = "default",
428  .type = AVMEDIA_TYPE_VIDEO,
429  .request_frame = request_frame,
430  },
431  { NULL }
432 };
433 
435  .name = "thumbnail_cuda",
436  .description = NULL_IF_CONFIG_SMALL("Select the most representative frame in a given sequence of consecutive frames."),
437  .priv_size = sizeof(ThumbnailCudaContext),
438  .init = init,
439  .uninit = uninit,
441  .inputs = thumbnail_cuda_inputs,
442  .outputs = thumbnail_cuda_outputs,
443  .priv_class = &thumbnail_cuda_class,
444  .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE,
445 };
#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:385
const char * s
Definition: avisynth_c.h:768
static const AVFilterPad thumbnail_cuda_inputs[]
AVFilter ff_vf_thumbnail_cuda
This structure describes decoded (raw) audio or video data.
Definition: frame.h:218
BYTE int const BYTE int src_pitch
Definition: avisynth_c.h:813
AVOption.
Definition: opt.h:246
const char * fmt
Definition: avisynth_c.h:769
planar YUV 4:4:4, 24bpp, (1 Cr & Cb sample per 1x1 Y samples)
Definition: pixfmt.h:67
Main libavfilter public API header.
const char * desc
Definition: nvenc.c:65
channels
Definition: aptx.c:30
#define OFFSET(x)
#define DIV_UP(a, b)
static enum AVPixelFormat supported_formats[]
static AVFrame * get_best_frame(AVFilterContext *ctx)
void * av_calloc(size_t nmemb, size_t size)
Non-inlined equivalent of av_mallocz_array().
Definition: mem.c:244
AVFilterFormats * ff_make_format_list(const int *fmts)
Create a list of supported formats.
Definition: formats.c:283
AVBufferRef * hw_frames_ctx
For hwaccel-format frames, this should be a reference to the AVHWFramesContext describing the frame...
Definition: frame.h:556
#define AV_PIX_FMT_P016
Definition: pixfmt.h:414
const char * name
Pad name.
Definition: internal.h:60
#define AV_PIX_FMT_P010
Definition: pixfmt.h:413
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:1080
static int filter_frame(AVFilterLink *inlink, AVFrame *frame)
AVFrame * buf
cached frame
Definition: vf_thumbnail.c:37
uint8_t
#define av_cold
Definition: attributes.h:82
AVOptions.
int64_t pts
Presentation timestamp in time_base units (time when frame should be shown to user).
Definition: frame.h:311
static AVFrame * frame
void * hwctx
The format-specific data, allocated and freed by libavutil along with this context.
Definition: hwcontext.h:91
static int config_props(AVFilterLink *inlink)
static double av_q2d(AVRational a)
Convert an AVRational to a double.
Definition: rational.h:104
#define AVERROR_EOF
End of file.
Definition: error.h:55
#define AV_LOG_VERBOSE
Detailed information.
Definition: log.h:192
#define AV_PIX_FMT_YUV444P16
Definition: pixfmt.h:381
static int format_is_supported(enum AVPixelFormat fmt)
AVBufferRef * hw_frames_ctx
#define av_log(a,...)
A filter pad used for either input or output.
Definition: internal.h:54
#define BLOCKY
int width
Definition: frame.h:276
#define AV_LOG_ERROR
Something went wrong and cannot losslessly be recovered.
Definition: log.h:176
int ff_set_common_formats(AVFilterContext *ctx, AVFilterFormats *formats)
A helper for query_formats() which sets all links to the same list of formats.
Definition: formats.c:568
#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
planar YUV 4:2:0, 12bpp, 1 plane for Y and 1 plane for the UV components, which are interleaved (firs...
Definition: pixfmt.h:85
struct thumb_frame * frames
the n_frames frames
static double frame_sum_square_err(const int *hist, const double *median)
Compute Sum-square deviation to estimate "closeness".
like NV12, with 16bpp per component, little-endian
Definition: pixfmt.h:296
#define BLOCKX
AVFILTER_DEFINE_CLASS(thumbnail_cuda)
static int request_frame(AVFilterLink *link)
static int thumbnail_kernel(ThumbnailCudaContext *s, CUfunction func, CUtexref tex, int channels, int *histogram, uint8_t *src_dptr, int src_width, int src_height, int src_pitch, int pixel_size)
static av_cold int init(AVFilterContext *ctx)
#define FLAGS
AVHWDeviceContext * device_ctx
The parent AVHWDeviceContext.
Definition: hwcontext.h:148
AVFormatContext * ctx
Definition: movenc.c:48
FFmpeg internal API for CUDA.
int dummy
Definition: motion.c:64
#define HIST_SIZE
like NV12, with 10bpp per component, data in the high bits, zeros in the low bits, little-endian
Definition: pixfmt.h:280
static const AVFilterPad inputs[]
Definition: af_acontrast.c:193
HW acceleration through CUDA.
Definition: pixfmt.h:231
static const AVFilterPad outputs[]
Definition: af_acontrast.c:203
#define FF_ARRAY_ELEMS(a)
static av_cold void uninit(AVFilterContext *ctx)
static int thumbnail(AVFilterContext *ctx, int *histogram, AVFrame *in)
#define AV_LOG_INFO
Standard information.
Definition: log.h:187
int linesize[AV_NUM_DATA_POINTERS]
For video, size in bytes of each picture line.
Definition: frame.h:249
uint8_t * data
The data buffer.
Definition: buffer.h:89
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
AVRational tb
copy of the input timebase to ease access
#define AVERROR_BUG
Internal bug, also see AVERROR_BUG2.
Definition: error.h:50
This struct is allocated as AVHWDeviceContext.hwctx.
Describe the class of an AVClass context structure.
Definition: log.h:67
Filter definition.
Definition: avfilter.h:144
Rational number (pair of numerator and denominator).
Definition: rational.h:58
This struct describes a set or pool of "hardware" frames (i.e.
Definition: hwcontext.h:123
int(* func)(AVBPrint *dst, const char *in, const char *arg)
Definition: jacosubdec.c:67
int n
current frame
const char * name
Filter name.
Definition: avfilter.h:148
AVFilterLink ** outputs
array of pointers to output links
Definition: avfilter.h:350
static enum AVPixelFormat pix_fmts[]
Definition: libkvazaar.c:266
uint8_t * data[AV_NUM_DATA_POINTERS]
pointer to the picture/channel planes.
Definition: frame.h:232
A reference to a data buffer.
Definition: buffer.h:81
int
planar YUV 4:2:0, 12bpp, (1 Cr & Cb sample per 2x2 Y samples)
Definition: pixfmt.h:62
static int query_formats(AVFilterContext *ctx)
AVBufferRef * av_buffer_ref(AVBufferRef *buf)
Create a new reference to an AVBuffer.
Definition: buffer.c:93
#define AVERROR_UNKNOWN
Unknown error, typically from an external library.
Definition: error.h:71
int n_frames
number of frames for analysis
static const AVFilterPad thumbnail_cuda_outputs[]
static const AVOption thumbnail_cuda_options[]
A list of supported formats for one end of a filter link.
Definition: formats.h:64
An instance of a filter.
Definition: avfilter.h:338
int height
Definition: frame.h:276
#define av_freep(p)
int histogram[HIST_SIZE]
RGB color distribution histogram of the frame.
Definition: vf_thumbnail.c:38
int ff_request_frame(AVFilterLink *link)
Request an input frame from the filter at the other end of the link.
Definition: avfilter.c:407
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
enum AVPixelFormat sw_format
The pixel format identifying the actual data layout of the hardware frames.
Definition: hwcontext.h:221
AVPixelFormat
Pixel format.
Definition: pixfmt.h:60