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/mem.h"
27 #include "libavutil/opt.h"
28 #include "libavutil/pixdesc.h"
29 
30 #include "avfilter.h"
31 #include "filters.h"
32 
33 #include "cuda/load_helper.h"
34 
35 #define CHECK_CU(x) FF_CUDA_CHECK_DL(ctx, s->hwctx->internal->cuda_dl, x)
36 
37 #define HIST_SIZE (3*256)
38 #define DIV_UP(a, b) ( ((a) + (b) - 1) / (b) )
39 #define BLOCKX 32
40 #define BLOCKY 16
41 
42 static const enum AVPixelFormat supported_formats[] = {
49 };
50 
51 struct thumb_frame {
52  AVFrame *buf; ///< cached frame
53  int histogram[HIST_SIZE]; ///< RGB color distribution histogram of the frame
54 };
55 
56 typedef struct ThumbnailCudaContext {
57  const AVClass *class;
58  int n; ///< current frame
59  int n_frames; ///< number of frames for analysis
60  struct thumb_frame *frames; ///< the n_frames frames
61  AVRational tb; ///< copy of the input timebase to ease access
62 
65 
66  CUmodule cu_module;
67 
68  CUfunction cu_func_uchar;
69  CUfunction cu_func_uchar2;
70  CUfunction cu_func_ushort;
71  CUfunction cu_func_ushort2;
72  CUstream cu_stream;
73 
74  CUdeviceptr data;
75 
77 
78 #define OFFSET(x) offsetof(ThumbnailCudaContext, x)
79 #define FLAGS AV_OPT_FLAG_VIDEO_PARAM|AV_OPT_FLAG_FILTERING_PARAM
80 
81 static const AVOption thumbnail_cuda_options[] = {
82  { "n", "set the frames batch size", OFFSET(n_frames), AV_OPT_TYPE_INT, {.i64=100}, 2, INT_MAX, FLAGS },
83  { NULL }
84 };
85 
86 AVFILTER_DEFINE_CLASS(thumbnail_cuda);
87 
89 {
90  ThumbnailCudaContext *s = ctx->priv;
91 
92  s->frames = av_calloc(s->n_frames, sizeof(*s->frames));
93  if (!s->frames) {
95  "Allocation failure, try to lower the number of frames\n");
96  return AVERROR(ENOMEM);
97  }
98  av_log(ctx, AV_LOG_VERBOSE, "batch size: %d frames\n", s->n_frames);
99  return 0;
100 }
101 
102 /**
103  * @brief Compute Sum-square deviation to estimate "closeness".
104  * @param hist color distribution histogram
105  * @param median average color distribution histogram
106  * @return sum of squared errors
107  */
108 static double frame_sum_square_err(const int *hist, const double *median)
109 {
110  int i;
111  double err, sum_sq_err = 0;
112 
113  for (i = 0; i < HIST_SIZE; i++) {
114  err = median[i] - (double)hist[i];
115  sum_sq_err += err*err;
116  }
117  return sum_sq_err;
118 }
119 
121 {
122  AVFrame *picref;
123  ThumbnailCudaContext *s = ctx->priv;
124  int i, j, best_frame_idx = 0;
125  int nb_frames = s->n;
126  double avg_hist[HIST_SIZE] = {0}, sq_err, min_sq_err = -1;
127 
128  // average histogram of the N frames
129  for (j = 0; j < FF_ARRAY_ELEMS(avg_hist); j++) {
130  for (i = 0; i < nb_frames; i++)
131  avg_hist[j] += (double)s->frames[i].histogram[j];
132  avg_hist[j] /= nb_frames;
133  }
134 
135  // find the frame closer to the average using the sum of squared errors
136  for (i = 0; i < nb_frames; i++) {
137  sq_err = frame_sum_square_err(s->frames[i].histogram, avg_hist);
138  if (i == 0 || sq_err < min_sq_err)
139  best_frame_idx = i, min_sq_err = sq_err;
140  }
141 
142  // free and reset everything (except the best frame buffer)
143  for (i = 0; i < nb_frames; i++) {
144  memset(s->frames[i].histogram, 0, sizeof(s->frames[i].histogram));
145  if (i != best_frame_idx)
146  av_frame_free(&s->frames[i].buf);
147  }
148  s->n = 0;
149 
150  // raise the chosen one
151  picref = s->frames[best_frame_idx].buf;
152  av_log(ctx, AV_LOG_INFO, "frame id #%d (pts_time=%f) selected "
153  "from a set of %d images\n", best_frame_idx,
154  picref->pts * av_q2d(s->tb), nb_frames);
155  s->frames[best_frame_idx].buf = NULL;
156 
157  return picref;
158 }
159 
160 static int thumbnail_kernel(AVFilterContext *ctx, CUfunction func, int channels,
161  int *histogram, uint8_t *src_dptr, int src_width, int src_height, int src_pitch, int pixel_size)
162 {
163  int ret;
164  ThumbnailCudaContext *s = ctx->priv;
165  CudaFunctions *cu = s->hwctx->internal->cuda_dl;
166  CUtexObject tex = 0;
167  void *args[] = { &tex, &histogram, &src_width, &src_height };
168 
169  CUDA_TEXTURE_DESC tex_desc = {
170  .filterMode = CU_TR_FILTER_MODE_LINEAR,
171  .flags = CU_TRSF_READ_AS_INTEGER,
172  };
173 
174  CUDA_RESOURCE_DESC res_desc = {
175  .resType = CU_RESOURCE_TYPE_PITCH2D,
176  .res.pitch2D.format = pixel_size == 1 ?
177  CU_AD_FORMAT_UNSIGNED_INT8 :
178  CU_AD_FORMAT_UNSIGNED_INT16,
179  .res.pitch2D.numChannels = channels,
180  .res.pitch2D.width = src_width,
181  .res.pitch2D.height = src_height,
182  .res.pitch2D.pitchInBytes = src_pitch,
183  .res.pitch2D.devPtr = (CUdeviceptr)src_dptr,
184  };
185 
186  ret = CHECK_CU(cu->cuTexObjectCreate(&tex, &res_desc, &tex_desc, NULL));
187  if (ret < 0)
188  goto exit;
189 
190  ret = CHECK_CU(cu->cuLaunchKernel(func,
191  DIV_UP(src_width, BLOCKX), DIV_UP(src_height, BLOCKY), 1,
192  BLOCKX, BLOCKY, 1, 0, s->cu_stream, args, NULL));
193 exit:
194  if (tex)
195  CHECK_CU(cu->cuTexObjectDestroy(tex));
196 
197  return ret;
198 }
199 
201 {
202  AVHWFramesContext *in_frames_ctx = (AVHWFramesContext*)in->hw_frames_ctx->data;
203  ThumbnailCudaContext *s = ctx->priv;
204 
205  switch (in_frames_ctx->sw_format) {
206  case AV_PIX_FMT_NV12:
207  thumbnail_kernel(ctx, s->cu_func_uchar, 1,
208  histogram, in->data[0], in->width, in->height, in->linesize[0], 1);
209  thumbnail_kernel(ctx, s->cu_func_uchar2, 2,
210  histogram + 256, in->data[1], in->width / 2, in->height / 2, in->linesize[1], 1);
211  break;
212  case AV_PIX_FMT_YUV420P:
213  thumbnail_kernel(ctx, s->cu_func_uchar, 1,
214  histogram, in->data[0], in->width, in->height, in->linesize[0], 1);
215  thumbnail_kernel(ctx, s->cu_func_uchar, 1,
216  histogram + 256, in->data[1], in->width / 2, in->height / 2, in->linesize[1], 1);
217  thumbnail_kernel(ctx, s->cu_func_uchar, 1,
218  histogram + 512, in->data[2], in->width / 2, in->height / 2, in->linesize[2], 1);
219  break;
220  case AV_PIX_FMT_YUV444P:
221  thumbnail_kernel(ctx, s->cu_func_uchar, 1,
222  histogram, in->data[0], in->width, in->height, in->linesize[0], 1);
223  thumbnail_kernel(ctx, s->cu_func_uchar, 1,
224  histogram + 256, in->data[1], in->width, in->height, in->linesize[1], 1);
225  thumbnail_kernel(ctx, s->cu_func_uchar, 1,
226  histogram + 512, in->data[2], in->width, in->height, in->linesize[2], 1);
227  break;
228  case AV_PIX_FMT_P010LE:
229  case AV_PIX_FMT_P016LE:
230  thumbnail_kernel(ctx, s->cu_func_ushort, 1,
231  histogram, in->data[0], in->width, in->height, in->linesize[0], 2);
232  thumbnail_kernel(ctx, s->cu_func_ushort2, 2,
233  histogram + 256, in->data[1], in->width / 2, in->height / 2, in->linesize[1], 2);
234  break;
236  thumbnail_kernel(ctx, s->cu_func_ushort2, 1,
237  histogram, in->data[0], in->width, in->height, in->linesize[0], 2);
238  thumbnail_kernel(ctx, s->cu_func_ushort2, 1,
239  histogram + 256, in->data[1], in->width, in->height, in->linesize[1], 2);
240  thumbnail_kernel(ctx, s->cu_func_ushort2, 1,
241  histogram + 512, in->data[2], in->width, in->height, in->linesize[2], 2);
242  break;
243  default:
244  return AVERROR_BUG;
245  }
246 
247  return 0;
248 }
249 
251 {
252  AVFilterContext *ctx = inlink->dst;
253  ThumbnailCudaContext *s = ctx->priv;
254  CudaFunctions *cu = s->hwctx->internal->cuda_dl;
255  AVFilterLink *outlink = ctx->outputs[0];
256  int *hist = s->frames[s->n].histogram;
257  AVHWFramesContext *hw_frames_ctx = (AVHWFramesContext*)s->hw_frames_ctx->data;
258  CUcontext dummy;
259  CUDA_MEMCPY2D cpy = { 0 };
260  int ret = 0;
261 
262  // keep a reference of each frame
263  s->frames[s->n].buf = frame;
264 
265  ret = CHECK_CU(cu->cuCtxPushCurrent(s->hwctx->cuda_ctx));
266  if (ret < 0)
267  return ret;
268 
269  CHECK_CU(cu->cuMemsetD8Async(s->data, 0, HIST_SIZE * sizeof(int), s->cu_stream));
270 
271  thumbnail(ctx, (int*)s->data, frame);
272 
273  cpy.srcMemoryType = CU_MEMORYTYPE_DEVICE;
274  cpy.dstMemoryType = CU_MEMORYTYPE_HOST;
275  cpy.srcDevice = s->data;
276  cpy.dstHost = hist;
277  cpy.srcPitch = HIST_SIZE * sizeof(int);
278  cpy.dstPitch = HIST_SIZE * sizeof(int);
279  cpy.WidthInBytes = HIST_SIZE * sizeof(int);
280  cpy.Height = 1;
281 
282  ret = CHECK_CU(cu->cuMemcpy2DAsync(&cpy, s->cu_stream));
283  if (ret < 0)
284  return ret;
285 
286  if (hw_frames_ctx->sw_format == AV_PIX_FMT_NV12 || hw_frames_ctx->sw_format == AV_PIX_FMT_YUV420P ||
287  hw_frames_ctx->sw_format == AV_PIX_FMT_P010LE || hw_frames_ctx->sw_format == AV_PIX_FMT_P016LE)
288  {
289  int i;
290  for (i = 256; i < HIST_SIZE; i++)
291  hist[i] = 4 * hist[i];
292  }
293 
294  ret = CHECK_CU(cu->cuCtxPopCurrent(&dummy));
295  if (ret < 0)
296  return ret;
297 
298  // no selection until the buffer of N frames is filled up
299  s->n++;
300  if (s->n < s->n_frames)
301  return 0;
302 
303  return ff_filter_frame(outlink, get_best_frame(ctx));
304 }
305 
307 {
308  ThumbnailCudaContext *s = ctx->priv;
309 
310  if (s->hwctx) {
311  CudaFunctions *cu = s->hwctx->internal->cuda_dl;
312 
313  if (s->data) {
314  CHECK_CU(cu->cuMemFree(s->data));
315  s->data = 0;
316  }
317 
318  if (s->cu_module) {
319  CHECK_CU(cu->cuModuleUnload(s->cu_module));
320  s->cu_module = NULL;
321  }
322  }
323 
324  if (s->frames) {
325  for (int i = 0; i < s->n_frames && s->frames[i].buf; i++)
326  av_frame_free(&s->frames[i].buf);
327  av_freep(&s->frames);
328  }
329 }
330 
332 {
333  AVFilterContext *ctx = link->src;
334  ThumbnailCudaContext *s = ctx->priv;
335  int ret = ff_request_frame(ctx->inputs[0]);
336 
337  if (ret == AVERROR_EOF && s->n) {
339  if (ret < 0)
340  return ret;
341  ret = AVERROR_EOF;
342  }
343  if (ret < 0)
344  return ret;
345  return 0;
346 }
347 
348 static int format_is_supported(enum AVPixelFormat fmt)
349 {
350  int i;
351 
352  for (i = 0; i < FF_ARRAY_ELEMS(supported_formats); i++)
353  if (supported_formats[i] == fmt)
354  return 1;
355  return 0;
356 }
357 
359 {
360  AVFilterContext *ctx = inlink->dst;
362  FilterLink *outl = ff_filter_link(ctx->outputs[0]);
363  ThumbnailCudaContext *s = ctx->priv;
364  AVHWFramesContext *hw_frames_ctx = (AVHWFramesContext*)inl->hw_frames_ctx->data;
365  AVCUDADeviceContext *device_hwctx = hw_frames_ctx->device_ctx->hwctx;
366  CUcontext dummy, cuda_ctx = device_hwctx->cuda_ctx;
367  CudaFunctions *cu = device_hwctx->internal->cuda_dl;
368  int ret;
369 
370  extern const unsigned char ff_vf_thumbnail_cuda_ptx_data[];
371  extern const unsigned int ff_vf_thumbnail_cuda_ptx_len;
372 
373  s->hwctx = device_hwctx;
374  s->cu_stream = s->hwctx->stream;
375 
376  ret = CHECK_CU(cu->cuCtxPushCurrent(cuda_ctx));
377  if (ret < 0)
378  return ret;
379 
380  ret = ff_cuda_load_module(ctx, device_hwctx, &s->cu_module, ff_vf_thumbnail_cuda_ptx_data, ff_vf_thumbnail_cuda_ptx_len);
381  if (ret < 0)
382  return ret;
383 
384  ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_uchar, s->cu_module, "Thumbnail_uchar"));
385  if (ret < 0)
386  return ret;
387 
388  ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_uchar2, s->cu_module, "Thumbnail_uchar2"));
389  if (ret < 0)
390  return ret;
391 
392  ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_ushort, s->cu_module, "Thumbnail_ushort"));
393  if (ret < 0)
394  return ret;
395 
396  ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_ushort2, s->cu_module, "Thumbnail_ushort2"));
397  if (ret < 0)
398  return ret;
399 
400  ret = CHECK_CU(cu->cuMemAlloc(&s->data, HIST_SIZE * sizeof(int)));
401  if (ret < 0)
402  return ret;
403 
404  CHECK_CU(cu->cuCtxPopCurrent(&dummy));
405 
406  s->hw_frames_ctx = inl->hw_frames_ctx;
407 
408  outl->hw_frames_ctx = av_buffer_ref(s->hw_frames_ctx);
409  if (!outl->hw_frames_ctx)
410  return AVERROR(ENOMEM);
411 
412  s->tb = inlink->time_base;
413 
414  if (!format_is_supported(hw_frames_ctx->sw_format)) {
415  av_log(ctx, AV_LOG_ERROR, "Unsupported input format: %s\n", av_get_pix_fmt_name(hw_frames_ctx->sw_format));
416  return AVERROR(ENOSYS);
417  }
418 
419  return 0;
420 }
421 
423  {
424  .name = "default",
425  .type = AVMEDIA_TYPE_VIDEO,
426  .config_props = config_props,
427  .filter_frame = filter_frame,
428  },
429 };
430 
432  {
433  .name = "default",
434  .type = AVMEDIA_TYPE_VIDEO,
435  .request_frame = request_frame,
436  },
437 };
438 
440  .name = "thumbnail_cuda",
441  .description = NULL_IF_CONFIG_SMALL("Select the most representative frame in a given sequence of consecutive frames."),
442  .priv_size = sizeof(ThumbnailCudaContext),
443  .init = init,
444  .uninit = uninit,
448  .priv_class = &thumbnail_cuda_class,
449  .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE,
450 };
func
int(* func)(AVBPrint *dst, const char *in, const char *arg)
Definition: jacosubdec.c:68
AVHWDeviceContext::hwctx
void * hwctx
The format-specific data, allocated and freed by libavutil along with this context.
Definition: hwcontext.h:85
AV_PIX_FMT_CUDA
@ AV_PIX_FMT_CUDA
HW acceleration through CUDA.
Definition: pixfmt.h:260
AVPixelFormat
AVPixelFormat
Pixel format.
Definition: pixfmt.h:71
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:81
ThumbnailCudaContext::cu_module
CUmodule cu_module
Definition: vf_thumbnail_cuda.c:66
hwcontext_cuda_internal.h
ff_vf_thumbnail_cuda
const AVFilter ff_vf_thumbnail_cuda
Definition: vf_thumbnail_cuda.c:439
ff_filter_frame
int ff_filter_frame(AVFilterLink *link, AVFrame *frame)
Send a frame of data to the next filter.
Definition: avfilter.c:1061
filter_frame
static int filter_frame(AVFilterLink *inlink, AVFrame *frame)
Definition: vf_thumbnail_cuda.c:250
AVERROR_EOF
#define AVERROR_EOF
End of file.
Definition: error.h:57
AVBufferRef::data
uint8_t * data
The data buffer.
Definition: buffer.h:90
thumbnail_cuda_inputs
static const AVFilterPad thumbnail_cuda_inputs[]
Definition: vf_thumbnail_cuda.c:422
thumb_frame::histogram
int histogram[HIST_SIZE]
RGB color distribution histogram of the frame.
Definition: vf_thumbnail.c:40
get_best_frame
static AVFrame * get_best_frame(AVFilterContext *ctx)
Definition: vf_thumbnail_cuda.c:120
thumb_frame::buf
AVFrame * buf
cached frame
Definition: vf_thumbnail.c:39
ff_cuda_load_module
int ff_cuda_load_module(void *avctx, AVCUDADeviceContext *hwctx, CUmodule *cu_module, const unsigned char *data, const unsigned int length)
Loads a CUDA module and applies any decompression, if necessary.
Definition: load_helper.c:35
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:162
FILTER_INPUTS
#define FILTER_INPUTS(array)
Definition: filters.h:262
AVFrame
This structure describes decoded (raw) audio or video data.
Definition: frame.h:389
pixdesc.h
AVFrame::pts
int64_t pts
Presentation timestamp in time_base units (time when frame should be shown to user).
Definition: frame.h:501
AVFrame::width
int width
Definition: frame.h:461
AVOption
AVOption.
Definition: opt.h:429
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:475
AV_LOG_VERBOSE
#define AV_LOG_VERBOSE
Detailed information.
Definition: log.h:225
av_buffer_ref
AVBufferRef * av_buffer_ref(const AVBufferRef *buf)
Create a new reference to an AVBuffer.
Definition: buffer.c:103
AVFilter::name
const char * name
Filter name.
Definition: avfilter.h:205
ThumbnailCudaContext::n
int n
current frame
Definition: vf_thumbnail_cuda.c:58
AVFrame::data
uint8_t * data[AV_NUM_DATA_POINTERS]
pointer to the picture/channel planes.
Definition: frame.h:410
thumb_frame
Definition: vf_thumbnail.c:38
config_props
static int config_props(AVFilterLink *inlink)
Definition: vf_thumbnail_cuda.c:358
dummy
int dummy
Definition: motion.c:66
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:108
supported_formats
static enum AVPixelFormat supported_formats[]
Definition: vf_thumbnail_cuda.c:42
AVFilterPad
A filter pad used for either input or output.
Definition: filters.h:38
request_frame
static int request_frame(AVFilterLink *link)
Definition: vf_thumbnail_cuda.c:331
AV_LOG_ERROR
#define AV_LOG_ERROR
Something went wrong and cannot losslessly be recovered.
Definition: log.h:209
FF_ARRAY_ELEMS
#define FF_ARRAY_ELEMS(a)
Definition: sinewin_tablegen.c:29
av_cold
#define av_cold
Definition: attributes.h:90
s
#define s(width, name)
Definition: cbs_vp9.c:198
ThumbnailCudaContext::cu_func_ushort
CUfunction cu_func_ushort
Definition: vf_thumbnail_cuda.c:70
FLAGS
#define FLAGS
Definition: vf_thumbnail_cuda.c:79
AV_PIX_FMT_YUV444P16
#define AV_PIX_FMT_YUV444P16
Definition: pixfmt.h:515
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:348
filters.h
ctx
AVFormatContext * ctx
Definition: movenc.c:49
channels
channels
Definition: aptx.h:31
load_helper.h
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:73
FILTER_OUTPUTS
#define FILTER_OUTPUTS(array)
Definition: filters.h:263
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:37
ThumbnailCudaContext
Definition: vf_thumbnail_cuda.c:56
ThumbnailCudaContext::frames
struct thumb_frame * frames
the n_frames frames
Definition: vf_thumbnail_cuda.c:60
init
static av_cold int init(AVFilterContext *ctx)
Definition: vf_thumbnail_cuda.c:88
AVClass
Describe the class of an AVClass context structure.
Definition: log.h:75
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:210
AVRational
Rational number (pair of numerator and denominator).
Definition: rational.h:58
double
double
Definition: af_crystalizer.c:132
AVFILTER_DEFINE_CLASS
AVFILTER_DEFINE_CLASS(thumbnail_cuda)
BLOCKY
#define BLOCKY
Definition: vf_thumbnail_cuda.c:40
ThumbnailCudaContext::hwctx
AVCUDADeviceContext * hwctx
Definition: vf_thumbnail_cuda.c:64
ff_filter_link
static FilterLink * ff_filter_link(AVFilterLink *link)
Definition: filters.h:197
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: filters.h:206
uninit
static av_cold void uninit(AVFilterContext *ctx)
Definition: vf_thumbnail_cuda.c:306
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:94
BLOCKX
#define BLOCKX
Definition: vf_thumbnail_cuda.c:39
ThumbnailCudaContext::cu_func_uchar
CUfunction cu_func_uchar
Definition: vf_thumbnail_cuda.c:68
ThumbnailCudaContext::tb
AVRational tb
copy of the input timebase to ease access
Definition: vf_thumbnail_cuda.c:61
AV_LOG_INFO
#define AV_LOG_INFO
Standard information.
Definition: log.h:220
thumbnail
static int thumbnail(AVFilterContext *ctx, int *histogram, AVFrame *in)
Definition: vf_thumbnail_cuda.c:200
i
#define i(width, name, range_min, range_max)
Definition: cbs_h2645.c:256
CHECK_CU
#define CHECK_CU(x)
Definition: vf_thumbnail_cuda.c:35
AVFilterPad::name
const char * name
Pad name.
Definition: filters.h:44
av_calloc
void * av_calloc(size_t nmemb, size_t size)
Definition: mem.c:264
AV_PIX_FMT_P016
#define AV_PIX_FMT_P016
Definition: pixfmt.h:554
AVFilter
Filter definition.
Definition: avfilter.h:201
AVHWFramesContext
This struct describes a set or pool of "hardware" frames (i.e.
Definition: hwcontext.h:115
AVCUDADeviceContext
This struct is allocated as AVHWDeviceContext.hwctx.
Definition: hwcontext_cuda.h:42
ret
ret
Definition: filter_design.txt:187
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:96
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:134
cuda_check.h
AV_PIX_FMT_P016LE
@ AV_PIX_FMT_P016LE
like NV12, with 16bpp per component, little-endian
Definition: pixfmt.h:323
ThumbnailCudaContext::data
CUdeviceptr data
Definition: vf_thumbnail_cuda.c:74
AVFrame::hw_frames_ctx
AVBufferRef * hw_frames_ctx
For hwaccel-format frames, this should be a reference to the AVHWFramesContext describing the frame.
Definition: frame.h:740
AVFrame::height
int height
Definition: frame.h:461
AV_OPT_TYPE_INT
@ AV_OPT_TYPE_INT
Underlying C type is int.
Definition: opt.h:259
avfilter.h
thumbnail_cuda_outputs
static const AVFilterPad thumbnail_cuda_outputs[]
Definition: vf_thumbnail_cuda.c:431
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:78
AVFilterContext
An instance of a filter.
Definition: avfilter.h:457
ThumbnailCudaContext::cu_func_ushort2
CUfunction cu_func_ushort2
Definition: vf_thumbnail_cuda.c:71
AV_PIX_FMT_P010
#define AV_PIX_FMT_P010
Definition: pixfmt.h:552
AVMEDIA_TYPE_VIDEO
@ AVMEDIA_TYPE_VIDEO
Definition: avutil.h:201
mem.h
AVBufferRef
A reference to a data buffer.
Definition: buffer.h:82
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:307
ThumbnailCudaContext::cu_stream
CUstream cu_stream
Definition: vf_thumbnail_cuda.c:72
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:160
av_freep
#define av_freep(p)
Definition: tableprint_vlc.h:34
ThumbnailCudaContext::n_frames
int n_frames
number of frames for analysis
Definition: vf_thumbnail_cuda.c:59
ThumbnailCudaContext::hw_frames_ctx
AVBufferRef * hw_frames_ctx
Definition: vf_thumbnail_cuda.c:63
hwcontext.h
AVERROR_BUG
#define AVERROR_BUG
Internal bug, also see AVERROR_BUG2.
Definition: error.h:52
AVFrame::linesize
int linesize[AV_NUM_DATA_POINTERS]
For video, a positive or negative value, which is typically indicating the size in bytes of each pict...
Definition: frame.h:434
av_log
#define av_log(a,...)
Definition: tableprint_vlc.h:27
ThumbnailCudaContext::cu_func_uchar2
CUfunction cu_func_uchar2
Definition: vf_thumbnail_cuda.c:69
FILTER_SINGLE_PIXFMT
#define FILTER_SINGLE_PIXFMT(pix_fmt_)
Definition: filters.h:252
OFFSET
#define OFFSET(x)
Definition: vf_thumbnail_cuda.c:78
DIV_UP
#define DIV_UP(a, b)
Definition: vf_thumbnail_cuda.c:38
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:3090