FFmpeg
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 "libavutil/hwcontext.h"
25 #include "libavutil/cuda_check.h"
26 #include "libavutil/opt.h"
27 #include "libavutil/pixdesc.h"
28 
29 #include "avfilter.h"
30 #include "internal.h"
31 
32 #define CHECK_CU(x) FF_CUDA_CHECK_DL(ctx, s->hwctx->internal->cuda_dl, x)
33 
34 #define HIST_SIZE (3*256)
35 #define DIV_UP(a, b) ( ((a) + (b) - 1) / (b) )
36 #define BLOCKX 32
37 #define BLOCKY 16
38 
39 static const enum AVPixelFormat supported_formats[] = {
46 };
47 
48 struct thumb_frame {
49  AVFrame *buf; ///< cached frame
50  int histogram[HIST_SIZE]; ///< RGB color distribution histogram of the frame
51 };
52 
53 typedef struct ThumbnailCudaContext {
54  const AVClass *class;
55  int n; ///< current frame
56  int n_frames; ///< number of frames for analysis
57  struct thumb_frame *frames; ///< the n_frames frames
58  AVRational tb; ///< copy of the input timebase to ease access
59 
62 
63  CUmodule cu_module;
64 
65  CUfunction cu_func_uchar;
66  CUfunction cu_func_uchar2;
67  CUfunction cu_func_ushort;
68  CUfunction cu_func_ushort2;
69  CUstream cu_stream;
70 
71  CUdeviceptr data;
72 
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) {
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(AVFilterContext *ctx, CUfunction func, int channels,
158  int *histogram, uint8_t *src_dptr, int src_width, int src_height, int src_pitch, int pixel_size)
159 {
160  int ret;
161  ThumbnailCudaContext *s = ctx->priv;
162  CudaFunctions *cu = s->hwctx->internal->cuda_dl;
163  CUtexObject tex = 0;
164  void *args[] = { &tex, &histogram, &src_width, &src_height };
165 
166  CUDA_TEXTURE_DESC tex_desc = {
167  .filterMode = CU_TR_FILTER_MODE_LINEAR,
168  .flags = CU_TRSF_READ_AS_INTEGER,
169  };
170 
171  CUDA_RESOURCE_DESC res_desc = {
172  .resType = CU_RESOURCE_TYPE_PITCH2D,
173  .res.pitch2D.format = pixel_size == 1 ?
174  CU_AD_FORMAT_UNSIGNED_INT8 :
175  CU_AD_FORMAT_UNSIGNED_INT16,
176  .res.pitch2D.numChannels = channels,
177  .res.pitch2D.width = src_width,
178  .res.pitch2D.height = src_height,
179  .res.pitch2D.pitchInBytes = src_pitch,
180  .res.pitch2D.devPtr = (CUdeviceptr)src_dptr,
181  };
182 
183  ret = CHECK_CU(cu->cuTexObjectCreate(&tex, &res_desc, &tex_desc, NULL));
184  if (ret < 0)
185  goto exit;
186 
187  ret = CHECK_CU(cu->cuLaunchKernel(func,
188  DIV_UP(src_width, BLOCKX), DIV_UP(src_height, BLOCKY), 1,
189  BLOCKX, BLOCKY, 1, 0, s->cu_stream, args, NULL));
190 exit:
191  if (tex)
192  CHECK_CU(cu->cuTexObjectDestroy(tex));
193 
194  return ret;
195 }
196 
198 {
199  AVHWFramesContext *in_frames_ctx = (AVHWFramesContext*)in->hw_frames_ctx->data;
200  ThumbnailCudaContext *s = ctx->priv;
201 
202  switch (in_frames_ctx->sw_format) {
203  case AV_PIX_FMT_NV12:
204  thumbnail_kernel(ctx, s->cu_func_uchar, 1,
205  histogram, in->data[0], in->width, in->height, in->linesize[0], 1);
206  thumbnail_kernel(ctx, s->cu_func_uchar2, 2,
207  histogram + 256, in->data[1], in->width / 2, in->height / 2, in->linesize[1], 1);
208  break;
209  case AV_PIX_FMT_YUV420P:
210  thumbnail_kernel(ctx, s->cu_func_uchar, 1,
211  histogram, in->data[0], in->width, in->height, in->linesize[0], 1);
212  thumbnail_kernel(ctx, s->cu_func_uchar, 1,
213  histogram + 256, in->data[1], in->width / 2, in->height / 2, in->linesize[1], 1);
214  thumbnail_kernel(ctx, s->cu_func_uchar, 1,
215  histogram + 512, in->data[2], in->width / 2, in->height / 2, in->linesize[2], 1);
216  break;
217  case AV_PIX_FMT_YUV444P:
218  thumbnail_kernel(ctx, s->cu_func_uchar, 1,
219  histogram, in->data[0], in->width, in->height, in->linesize[0], 1);
220  thumbnail_kernel(ctx, s->cu_func_uchar, 1,
221  histogram + 256, in->data[1], in->width, in->height, in->linesize[1], 1);
222  thumbnail_kernel(ctx, s->cu_func_uchar, 1,
223  histogram + 512, in->data[2], in->width, in->height, in->linesize[2], 1);
224  break;
225  case AV_PIX_FMT_P010LE:
226  case AV_PIX_FMT_P016LE:
227  thumbnail_kernel(ctx, s->cu_func_ushort, 1,
228  histogram, in->data[0], in->width, in->height, in->linesize[0], 2);
229  thumbnail_kernel(ctx, s->cu_func_ushort2, 2,
230  histogram + 256, in->data[1], in->width / 2, in->height / 2, in->linesize[1], 2);
231  break;
233  thumbnail_kernel(ctx, s->cu_func_ushort2, 1,
234  histogram, in->data[0], in->width, in->height, in->linesize[0], 2);
235  thumbnail_kernel(ctx, s->cu_func_ushort2, 1,
236  histogram + 256, in->data[1], in->width, in->height, in->linesize[1], 2);
237  thumbnail_kernel(ctx, s->cu_func_ushort2, 1,
238  histogram + 512, in->data[2], in->width, in->height, in->linesize[2], 2);
239  break;
240  default:
241  return AVERROR_BUG;
242  }
243 
244  return 0;
245 }
246 
248 {
249  AVFilterContext *ctx = inlink->dst;
250  ThumbnailCudaContext *s = ctx->priv;
251  CudaFunctions *cu = s->hwctx->internal->cuda_dl;
252  AVFilterLink *outlink = ctx->outputs[0];
253  int *hist = s->frames[s->n].histogram;
254  AVHWFramesContext *hw_frames_ctx = (AVHWFramesContext*)s->hw_frames_ctx->data;
255  CUcontext dummy;
256  CUDA_MEMCPY2D cpy = { 0 };
257  int ret = 0;
258 
259  // keep a reference of each frame
260  s->frames[s->n].buf = frame;
261 
262  ret = CHECK_CU(cu->cuCtxPushCurrent(s->hwctx->cuda_ctx));
263  if (ret < 0)
264  return ret;
265 
266  CHECK_CU(cu->cuMemsetD8Async(s->data, 0, HIST_SIZE * sizeof(int), s->cu_stream));
267 
268  thumbnail(ctx, (int*)s->data, frame);
269 
270  cpy.srcMemoryType = CU_MEMORYTYPE_DEVICE;
271  cpy.dstMemoryType = CU_MEMORYTYPE_HOST;
272  cpy.srcDevice = s->data;
273  cpy.dstHost = hist;
274  cpy.srcPitch = HIST_SIZE * sizeof(int);
275  cpy.dstPitch = HIST_SIZE * sizeof(int);
276  cpy.WidthInBytes = HIST_SIZE * sizeof(int);
277  cpy.Height = 1;
278 
279  ret = CHECK_CU(cu->cuMemcpy2DAsync(&cpy, s->cu_stream));
280  if (ret < 0)
281  return ret;
282 
283  if (hw_frames_ctx->sw_format == AV_PIX_FMT_NV12 || hw_frames_ctx->sw_format == AV_PIX_FMT_YUV420P ||
284  hw_frames_ctx->sw_format == AV_PIX_FMT_P010LE || hw_frames_ctx->sw_format == AV_PIX_FMT_P016LE)
285  {
286  int i;
287  for (i = 256; i < HIST_SIZE; i++)
288  hist[i] = 4 * hist[i];
289  }
290 
291  CHECK_CU(cu->cuCtxPopCurrent(&dummy));
292  if (ret < 0)
293  return ret;
294 
295  // no selection until the buffer of N frames is filled up
296  s->n++;
297  if (s->n < s->n_frames)
298  return 0;
299 
300  return ff_filter_frame(outlink, get_best_frame(ctx));
301 }
302 
304 {
305  int i;
306  ThumbnailCudaContext *s = ctx->priv;
307  CudaFunctions *cu = s->hwctx->internal->cuda_dl;
308 
309  if (s->data) {
310  CHECK_CU(cu->cuMemFree(s->data));
311  s->data = 0;
312  }
313 
314  if (s->cu_module) {
315  CHECK_CU(cu->cuModuleUnload(s->cu_module));
316  s->cu_module = NULL;
317  }
318 
319  for (i = 0; i < s->n_frames && s->frames[i].buf; i++)
320  av_frame_free(&s->frames[i].buf);
321  av_freep(&s->frames);
322 }
323 
325 {
326  AVFilterContext *ctx = link->src;
327  ThumbnailCudaContext *s = ctx->priv;
328  int ret = ff_request_frame(ctx->inputs[0]);
329 
330  if (ret == AVERROR_EOF && s->n) {
332  if (ret < 0)
333  return ret;
334  ret = AVERROR_EOF;
335  }
336  if (ret < 0)
337  return ret;
338  return 0;
339 }
340 
341 static int format_is_supported(enum AVPixelFormat fmt)
342 {
343  int i;
344 
345  for (i = 0; i < FF_ARRAY_ELEMS(supported_formats); i++)
346  if (supported_formats[i] == fmt)
347  return 1;
348  return 0;
349 }
350 
352 {
353  AVFilterContext *ctx = inlink->dst;
354  ThumbnailCudaContext *s = ctx->priv;
355  AVHWFramesContext *hw_frames_ctx = (AVHWFramesContext*)inlink->hw_frames_ctx->data;
356  AVCUDADeviceContext *device_hwctx = hw_frames_ctx->device_ctx->hwctx;
357  CUcontext dummy, cuda_ctx = device_hwctx->cuda_ctx;
358  CudaFunctions *cu = device_hwctx->internal->cuda_dl;
359  int ret;
360 
361  extern char vf_thumbnail_cuda_ptx[];
362 
363  s->hwctx = device_hwctx;
364  s->cu_stream = s->hwctx->stream;
365 
366  ret = CHECK_CU(cu->cuCtxPushCurrent(cuda_ctx));
367  if (ret < 0)
368  return ret;
369 
370  ret = CHECK_CU(cu->cuModuleLoadData(&s->cu_module, vf_thumbnail_cuda_ptx));
371  if (ret < 0)
372  return ret;
373 
374  ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_uchar, s->cu_module, "Thumbnail_uchar"));
375  if (ret < 0)
376  return ret;
377 
378  ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_uchar2, s->cu_module, "Thumbnail_uchar2"));
379  if (ret < 0)
380  return ret;
381 
382  ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_ushort, s->cu_module, "Thumbnail_ushort"));
383  if (ret < 0)
384  return ret;
385 
386  ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_ushort2, s->cu_module, "Thumbnail_ushort2"));
387  if (ret < 0)
388  return ret;
389 
390  ret = CHECK_CU(cu->cuMemAlloc(&s->data, HIST_SIZE * sizeof(int)));
391  if (ret < 0)
392  return ret;
393 
394  CHECK_CU(cu->cuCtxPopCurrent(&dummy));
395 
396  s->hw_frames_ctx = ctx->inputs[0]->hw_frames_ctx;
397 
398  ctx->outputs[0]->hw_frames_ctx = av_buffer_ref(s->hw_frames_ctx);
399  if (!ctx->outputs[0]->hw_frames_ctx)
400  return AVERROR(ENOMEM);
401 
402  s->tb = inlink->time_base;
403 
404  if (!format_is_supported(hw_frames_ctx->sw_format)) {
405  av_log(ctx, AV_LOG_ERROR, "Unsupported input format: %s\n", av_get_pix_fmt_name(hw_frames_ctx->sw_format));
406  return AVERROR(ENOSYS);
407  }
408 
409  return 0;
410 }
411 
413 {
414  static const enum AVPixelFormat pix_fmts[] = {
417  };
419  if (!fmts_list)
420  return AVERROR(ENOMEM);
421  return ff_set_common_formats(ctx, fmts_list);
422 }
423 
425  {
426  .name = "default",
427  .type = AVMEDIA_TYPE_VIDEO,
428  .config_props = config_props,
429  .filter_frame = filter_frame,
430  },
431  { NULL }
432 };
433 
435  {
436  .name = "default",
437  .type = AVMEDIA_TYPE_VIDEO,
438  .request_frame = request_frame,
439  },
440  { NULL }
441 };
442 
444  .name = "thumbnail_cuda",
445  .description = NULL_IF_CONFIG_SMALL("Select the most representative frame in a given sequence of consecutive frames."),
446  .priv_size = sizeof(ThumbnailCudaContext),
447  .init = init,
448  .uninit = uninit,
452  .priv_class = &thumbnail_cuda_class,
453  .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE,
454 };
func
int(* func)(AVBPrint *dst, const char *in, const char *arg)
Definition: jacosubdec.c:67
AVHWDeviceContext::hwctx
void * hwctx
The format-specific data, allocated and freed by libavutil along with this context.
Definition: hwcontext.h:92
AV_PIX_FMT_CUDA
@ AV_PIX_FMT_CUDA
HW acceleration through CUDA.
Definition: pixfmt.h:235
AVPixelFormat
AVPixelFormat
Pixel format.
Definition: pixfmt.h:64
AVERROR
Filter the word “frame” indicates either a video frame or a group of audio as stored in an AVFrame structure Format for each input and each output the list of supported formats For video that means pixel format For audio that means channel sample they are references to shared objects When the negotiation mechanism computes the intersection of the formats supported at each end of a all references to both lists are replaced with a reference to the intersection And when a single format is eventually chosen for a link amongst the remaining all references to the list are updated That means that if a filter requires that its input and output have the same format amongst a supported all it has to do is use a reference to the same list of formats query_formats can leave some formats unset and return AVERROR(EAGAIN) to cause the negotiation mechanism toagain later. That can be used by filters with complex requirements to use the format negotiated on one link to set the formats supported on another. Frame references ownership and permissions
opt.h
thumbnail_cuda_options
static const AVOption thumbnail_cuda_options[]
Definition: vf_thumbnail_cuda.c:78
ff_make_format_list
AVFilterFormats * ff_make_format_list(const int *fmts)
Create a list of supported formats.
Definition: formats.c:286
ThumbnailCudaContext::cu_module
CUmodule cu_module
Definition: vf_thumbnail_cuda.c:63
hwcontext_cuda_internal.h
FF_FILTER_FLAG_HWFRAME_AWARE
#define FF_FILTER_FLAG_HWFRAME_AWARE
The filter is aware of hardware frames, and any hardware frame context should not be automatically pr...
Definition: internal.h:339
ff_filter_frame
int ff_filter_frame(AVFilterLink *link, AVFrame *frame)
Send a frame of data to the next filter.
Definition: avfilter.c:1096
filter_frame
static int filter_frame(AVFilterLink *inlink, AVFrame *frame)
Definition: vf_thumbnail_cuda.c:247
AVERROR_EOF
#define AVERROR_EOF
End of file.
Definition: error.h:55
thumbnail_cuda_inputs
static const AVFilterPad thumbnail_cuda_inputs[]
Definition: vf_thumbnail_cuda.c:424
thumb_frame::histogram
int histogram[HIST_SIZE]
RGB color distribution histogram of the frame.
Definition: vf_thumbnail.c:39
get_best_frame
static AVFrame * get_best_frame(AVFilterContext *ctx)
Definition: vf_thumbnail_cuda.c:117
ff_vf_thumbnail_cuda
AVFilter ff_vf_thumbnail_cuda
Definition: vf_thumbnail_cuda.c:443
thumb_frame::buf
AVFrame * buf
cached frame
Definition: vf_thumbnail.c:38
inlink
The exact code depends on how similar the blocks are and how related they are to the and needs to apply these operations to the correct inlink or outlink if there are several Macros are available to factor that when no extra processing is inlink
Definition: filter_design.txt:212
av_frame_free
void av_frame_free(AVFrame **frame)
Free the frame and any dynamically allocated objects in it, e.g.
Definition: frame.c:203
AVFrame
This structure describes decoded (raw) audio or video data.
Definition: frame.h:318
pixdesc.h
AVFrame::pts
int64_t pts
Presentation timestamp in time_base units (time when frame should be shown to user).
Definition: frame.h:411
AVOption
AVOption.
Definition: opt.h:248
ff_request_frame
int ff_request_frame(AVFilterLink *link)
Request an input frame from the filter at the other end of the link.
Definition: avfilter.c:408
AV_LOG_VERBOSE
#define AV_LOG_VERBOSE
Detailed information.
Definition: log.h:210
AVFilter::name
const char * name
Filter name.
Definition: avfilter.h:149
ThumbnailCudaContext::n
int n
current frame
Definition: vf_thumbnail_cuda.c:55
AVFilterFormats
A list of supported formats for one end of a filter link.
Definition: formats.h:65
thumb_frame
Definition: vf_thumbnail.c:37
config_props
static int config_props(AVFilterLink *inlink)
Definition: vf_thumbnail_cuda.c:351
frame_sum_square_err
static double frame_sum_square_err(const int *hist, const double *median)
Compute Sum-square deviation to estimate "closeness".
Definition: vf_thumbnail_cuda.c:105
supported_formats
static enum AVPixelFormat supported_formats[]
Definition: vf_thumbnail_cuda.c:39
AVFilterPad
A filter pad used for either input or output.
Definition: internal.h:54
request_frame
static int request_frame(AVFilterLink *link)
Definition: vf_thumbnail_cuda.c:324
AV_LOG_ERROR
#define AV_LOG_ERROR
Something went wrong and cannot losslessly be recovered.
Definition: log.h:194
FF_ARRAY_ELEMS
#define FF_ARRAY_ELEMS(a)
Definition: sinewin_tablegen.c:29
av_cold
#define av_cold
Definition: attributes.h:90
ff_set_common_formats
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:587
s
#define s(width, name)
Definition: cbs_vp9.c:257
ThumbnailCudaContext::cu_func_ushort
CUfunction cu_func_ushort
Definition: vf_thumbnail_cuda.c:67
FLAGS
#define FLAGS
Definition: vf_thumbnail_cuda.c:76
AV_PIX_FMT_YUV444P16
#define AV_PIX_FMT_YUV444P16
Definition: pixfmt.h:412
av_q2d
static double av_q2d(AVRational a)
Convert an AVRational to a double.
Definition: rational.h:104
format_is_supported
static int format_is_supported(enum AVPixelFormat fmt)
Definition: vf_thumbnail_cuda.c:341
outputs
static const AVFilterPad outputs[]
Definition: af_acontrast.c:203
pix_fmts
static enum AVPixelFormat pix_fmts[]
Definition: libkvazaar.c:309
ctx
AVFormatContext * ctx
Definition: movenc.c:48
channels
channels
Definition: aptx.h:33
AV_PIX_FMT_YUV420P
@ AV_PIX_FMT_YUV420P
planar YUV 4:2:0, 12bpp, (1 Cr & Cb sample per 2x2 Y samples)
Definition: pixfmt.h:66
link
Filter the word “frame” indicates either a video frame or a group of audio as stored in an AVFrame structure Format for each input and each output the list of supported formats For video that means pixel format For audio that means channel sample they are references to shared objects When the negotiation mechanism computes the intersection of the formats supported at each end of a link
Definition: filter_design.txt:23
HIST_SIZE
#define HIST_SIZE
Definition: vf_thumbnail_cuda.c:34
ThumbnailCudaContext
Definition: vf_thumbnail_cuda.c:53
ThumbnailCudaContext::frames
struct thumb_frame * frames
the n_frames frames
Definition: vf_thumbnail_cuda.c:57
init
static av_cold int init(AVFilterContext *ctx)
Definition: vf_thumbnail_cuda.c:85
AVClass
Describe the class of an AVClass context structure.
Definition: log.h:67
NULL
#define NULL
Definition: coverity.c:32
AVHWFramesContext::sw_format
enum AVPixelFormat sw_format
The pixel format identifying the actual data layout of the hardware frames.
Definition: hwcontext.h:222
AVRational
Rational number (pair of numerator and denominator).
Definition: rational.h:58
AVFILTER_DEFINE_CLASS
AVFILTER_DEFINE_CLASS(thumbnail_cuda)
inputs
these buffered frames must be flushed immediately if a new input produces new the filter must not call request_frame to get more It must just process the frame or queue it The task of requesting more frames is left to the filter s request_frame method or the application If a filter has several inputs
Definition: filter_design.txt:243
BLOCKY
#define BLOCKY
Definition: vf_thumbnail_cuda.c:37
ThumbnailCudaContext::hwctx
AVCUDADeviceContext * hwctx
Definition: vf_thumbnail_cuda.c:61
uninit
static av_cold void uninit(AVFilterContext *ctx)
Definition: vf_thumbnail_cuda.c:303
NULL_IF_CONFIG_SMALL
#define NULL_IF_CONFIG_SMALL(x)
Return NULL if CONFIG_SMALL is true, otherwise the argument without modification.
Definition: internal.h:117
BLOCKX
#define BLOCKX
Definition: vf_thumbnail_cuda.c:36
ThumbnailCudaContext::cu_func_uchar
CUfunction cu_func_uchar
Definition: vf_thumbnail_cuda.c:65
ThumbnailCudaContext::tb
AVRational tb
copy of the input timebase to ease access
Definition: vf_thumbnail_cuda.c:58
AV_LOG_INFO
#define AV_LOG_INFO
Standard information.
Definition: log.h:205
internal.h
in
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)\n", len, av_get_sample_fmt_name(ac->in_fmt), av_get_sample_fmt_name(ac->out_fmt));return ff_convert_dither(ac-> in
Definition: audio_convert.c:326
thumbnail
static int thumbnail(AVFilterContext *ctx, int *histogram, AVFrame *in)
Definition: vf_thumbnail_cuda.c:197
i
int i
Definition: input.c:407
CHECK_CU
#define CHECK_CU(x)
Definition: vf_thumbnail_cuda.c:32
uint8_t
uint8_t
Definition: audio_convert.c:194
AVFilterPad::name
const char * name
Pad name.
Definition: internal.h:60
AV_PIX_FMT_P016
#define AV_PIX_FMT_P016
Definition: pixfmt.h:449
AVFilter
Filter definition.
Definition: avfilter.h:145
AVHWFramesContext
This struct describes a set or pool of "hardware" frames (i.e.
Definition: hwcontext.h:124
AVCUDADeviceContext
This struct is allocated as AVHWDeviceContext.hwctx.
Definition: hwcontext_cuda.h:42
ret
ret
Definition: filter_design.txt:187
query_formats
static int query_formats(AVFilterContext *ctx)
Definition: vf_thumbnail_cuda.c:412
AV_PIX_FMT_NV12
@ AV_PIX_FMT_NV12
planar YUV 4:2:0, 12bpp, 1 plane for Y and 1 plane for the UV components, which are interleaved (firs...
Definition: pixfmt.h:89
frame
these buffered frames must be flushed immediately if a new input produces new the filter must not call request_frame to get more It must just process the frame or queue it The task of requesting more frames is left to the filter s request_frame method or the application If a filter has several the filter must be ready for frames arriving randomly on any input any filter with several inputs will most likely require some kind of queuing mechanism It is perfectly acceptable to have a limited queue and to drop frames when the inputs are too unbalanced request_frame For filters that do not use the this method is called when a frame is wanted on an output For a it should directly call filter_frame on the corresponding output For a if there are queued frames already one of these frames should be pushed If the filter should request a frame on one of its repeatedly until at least one frame has been pushed Return or at least make progress towards producing a frame
Definition: filter_design.txt:264
AVHWFramesContext::device_ctx
AVHWDeviceContext * device_ctx
The parent AVHWDeviceContext.
Definition: hwcontext.h:149
cuda_check.h
AV_PIX_FMT_P016LE
@ AV_PIX_FMT_P016LE
like NV12, with 16bpp per component, little-endian
Definition: pixfmt.h:300
ThumbnailCudaContext::data
CUdeviceptr data
Definition: vf_thumbnail_cuda.c:71
av_calloc
void * av_calloc(size_t nmemb, size_t size)
Non-inlined equivalent of av_mallocz_array().
Definition: mem.c:245
AV_PIX_FMT_NONE
@ AV_PIX_FMT_NONE
Definition: pixfmt.h:65
dummy
int dummy
Definition: motion.c:64
AV_OPT_TYPE_INT
@ AV_OPT_TYPE_INT
Definition: opt.h:225
avfilter.h
thumbnail_cuda_outputs
static const AVFilterPad thumbnail_cuda_outputs[]
Definition: vf_thumbnail_cuda.c:434
av_buffer_ref
AVBufferRef * av_buffer_ref(AVBufferRef *buf)
Create a new reference to an AVBuffer.
Definition: buffer.c:93
AV_PIX_FMT_YUV444P
@ AV_PIX_FMT_YUV444P
planar YUV 4:4:4, 24bpp, (1 Cr & Cb sample per 1x1 Y samples)
Definition: pixfmt.h:71
AVFilterContext
An instance of a filter.
Definition: avfilter.h:341
ThumbnailCudaContext::cu_func_ushort2
CUfunction cu_func_ushort2
Definition: vf_thumbnail_cuda.c:68
AV_PIX_FMT_P010
#define AV_PIX_FMT_P010
Definition: pixfmt.h:448
AVMEDIA_TYPE_VIDEO
@ AVMEDIA_TYPE_VIDEO
Definition: avutil.h:201
AVBufferRef
A reference to a data buffer.
Definition: buffer.h:84
AV_PIX_FMT_P010LE
@ AV_PIX_FMT_P010LE
like NV12, with 10bpp per component, data in the high bits, zeros in the low bits,...
Definition: pixfmt.h:284
ThumbnailCudaContext::cu_stream
CUstream cu_stream
Definition: vf_thumbnail_cuda.c:69
thumbnail_kernel
static int thumbnail_kernel(AVFilterContext *ctx, CUfunction func, int channels, int *histogram, uint8_t *src_dptr, int src_width, int src_height, int src_pitch, int pixel_size)
Definition: vf_thumbnail_cuda.c:157
av_freep
#define av_freep(p)
Definition: tableprint_vlc.h:35
ThumbnailCudaContext::n_frames
int n_frames
number of frames for analysis
Definition: vf_thumbnail_cuda.c:56
ThumbnailCudaContext::hw_frames_ctx
AVBufferRef * hw_frames_ctx
Definition: vf_thumbnail_cuda.c:60
hwcontext.h
AVERROR_BUG
#define AVERROR_BUG
Internal bug, also see AVERROR_BUG2.
Definition: error.h:50
av_log
#define av_log(a,...)
Definition: tableprint_vlc.h:28
ThumbnailCudaContext::cu_func_uchar2
CUfunction cu_func_uchar2
Definition: vf_thumbnail_cuda.c:66
int
int
Definition: ffmpeg_filter.c:170
OFFSET
#define OFFSET(x)
Definition: vf_thumbnail_cuda.c:75
DIV_UP
#define DIV_UP(a, b)
Definition: vf_thumbnail_cuda.c:35
av_get_pix_fmt_name
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:2489