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