32 #define CHECK_CU(x) FF_CUDA_CHECK_DL(ctx, s->hwctx->internal->cuda_dl, x) 34 #define HIST_SIZE (3*256) 35 #define DIV_UP(a, b) ( ((a) + (b) - 1) / (b) ) 75 #define OFFSET(x) offsetof(ThumbnailCudaContext, x) 76 #define FLAGS AV_OPT_FLAG_VIDEO_PARAM|AV_OPT_FLAG_FILTERING_PARAM 92 "Allocation failure, try to lower the number of frames\n");
108 double err, sum_sq_err = 0;
111 err = median[
i] - (double)hist[i];
112 sum_sq_err += err*err;
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;
127 for (i = 0; i < nb_frames; i++)
129 avg_hist[j] /= nb_frames;
133 for (i = 0; i < nb_frames; i++) {
135 if (i == 0 || sq_err < min_sq_err)
136 best_frame_idx =
i, min_sq_err = sq_err;
140 for (i = 0; i < nb_frames; i++) {
142 if (i != best_frame_idx)
150 "from a set of %d images\n", best_frame_idx,
158 int *
histogram,
uint8_t *src_dptr,
int src_width,
int src_height,
int src_pitch,
int pixel_size)
164 void *args[] = { &tex, &
histogram, &src_width, &src_height };
166 CUDA_TEXTURE_DESC tex_desc = {
167 .filterMode = CU_TR_FILTER_MODE_LINEAR,
168 .flags = CU_TRSF_READ_AS_INTEGER,
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,
183 ret =
CHECK_CU(cu->cuTexObjectCreate(&tex, &res_desc, &tex_desc,
NULL));
187 ret =
CHECK_CU(cu->cuLaunchKernel(func,
192 CHECK_CU(cu->cuTexObjectDestroy(tex));
256 CUDA_MEMCPY2D cpy = { 0 };
262 ret =
CHECK_CU(cu->cuCtxPushCurrent(
s->hwctx->cuda_ctx));
270 cpy.srcMemoryType = CU_MEMORYTYPE_DEVICE;
271 cpy.dstMemoryType = CU_MEMORYTYPE_HOST;
272 cpy.srcDevice =
s->data;
279 ret =
CHECK_CU(cu->cuMemcpy2DAsync(&cpy,
s->cu_stream));
288 hist[i] = 4 * hist[i];
297 if (
s->n <
s->n_frames)
357 CUcontext
dummy, cuda_ctx = device_hwctx->cuda_ctx;
358 CudaFunctions *cu = device_hwctx->internal->cuda_dl;
361 extern char vf_thumbnail_cuda_ptx[];
363 s->
hwctx = device_hwctx;
394 CHECK_CU(cu->cuCtxPopCurrent(&dummy));
444 .
name =
"thumbnail_cuda",
445 .description =
NULL_IF_CONFIG_SMALL(
"Select the most representative frame in a given sequence of consecutive frames."),
450 .
inputs = thumbnail_cuda_inputs,
451 .
outputs = thumbnail_cuda_outputs,
452 .priv_class = &thumbnail_cuda_class,
#define FF_FILTER_FLAG_HWFRAME_AWARE
The filter is aware of hardware frames, and any hardware frame context should not be automatically pr...
static const AVFilterPad thumbnail_cuda_inputs[]
AVFilter ff_vf_thumbnail_cuda
AVCUDADeviceContext * hwctx
This structure describes decoded (raw) audio or video data.
planar YUV 4:4:4, 24bpp, (1 Cr & Cb sample per 1x1 Y samples)
Main libavfilter public API header.
AVCUDADeviceContextInternal * internal
static enum AVPixelFormat supported_formats[]
static AVFrame * get_best_frame(AVFilterContext *ctx)
#define FF_ARRAY_ELEMS(a)
void * av_calloc(size_t nmemb, size_t size)
Non-inlined equivalent of av_mallocz_array().
AVBufferRef * hw_frames_ctx
For hwaccel-format frames, this should be a reference to the AVHWFramesContext describing the frame...
const char * name
Pad name.
AVFilterLink ** inputs
array of pointers to input links
int ff_filter_frame(AVFilterLink *link, AVFrame *frame)
Send a frame of data to the next filter.
static int filter_frame(AVFilterLink *inlink, AVFrame *frame)
AVFrame * buf
cached frame
int64_t pts
Presentation timestamp in time_base units (time when frame should be shown to user).
void * hwctx
The format-specific data, allocated and freed by libavutil along with this context.
static int config_props(AVFilterLink *inlink)
static double av_q2d(AVRational a)
Convert an AVRational to a double.
#define AVERROR_EOF
End of file.
#define AV_LOG_VERBOSE
Detailed information.
#define AV_PIX_FMT_YUV444P16
static int format_is_supported(enum AVPixelFormat fmt)
AVBufferRef * hw_frames_ctx
A filter pad used for either input or output.
A link between two filters.
#define AV_LOG_ERROR
Something went wrong and cannot losslessly be recovered.
void av_frame_free(AVFrame **frame)
Free the frame and any dynamically allocated objects in it, e.g.
#define NULL_IF_CONFIG_SMALL(x)
Return NULL if CONFIG_SMALL is true, otherwise the argument without modification. ...
void * priv
private data for use by the filter
planar YUV 4:2:0, 12bpp, 1 plane for Y and 1 plane for the UV components, which are interleaved (firs...
AVRational time_base
Define the time base used by the PTS of the frames/samples which will pass through this link...
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
AVFILTER_DEFINE_CLASS(thumbnail_cuda)
static int request_frame(AVFilterLink *link)
CUfunction cu_func_ushort2
static av_cold int init(AVFilterContext *ctx)
AVBufferRef * hw_frames_ctx
For hwaccel pixel formats, this should be a reference to the AVHWFramesContext describing the frames...
AVHWDeviceContext * device_ctx
The parent AVHWDeviceContext.
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
FFmpeg internal API for CUDA.
AVFilterContext * src
source filter
like NV12, with 10bpp per component, data in the high bits, zeros in the low bits, little-endian
HW acceleration through CUDA.
static const AVFilterPad outputs[]
static av_cold void uninit(AVFilterContext *ctx)
static int thumbnail(AVFilterContext *ctx, int *histogram, AVFrame *in)
#define AV_LOG_INFO
Standard information.
int linesize[AV_NUM_DATA_POINTERS]
For video, size in bytes of each picture line.
uint8_t * data
The data buffer.
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
AVRational tb
copy of the input timebase to ease access
#define AVERROR_BUG
Internal bug, also see AVERROR_BUG2.
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
This struct is allocated as AVHWDeviceContext.hwctx.
int(* func)(AVBPrint *dst, const char *in, const char *arg)
Describe the class of an AVClass context structure.
Rational number (pair of numerator and denominator).
This struct describes a set or pool of "hardware" frames (i.e.
const char * name
Filter name.
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
AVFilterLink ** outputs
array of pointers to output links
static enum AVPixelFormat pix_fmts[]
uint8_t * data[AV_NUM_DATA_POINTERS]
pointer to the picture/channel planes.
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
A reference to a data buffer.
planar YUV 4:2:0, 12bpp, (1 Cr & Cb sample per 2x2 Y samples)
static int query_formats(AVFilterContext *ctx)
AVBufferRef * av_buffer_ref(AVBufferRef *buf)
Create a new reference to an AVBuffer.
int n_frames
number of frames for analysis
CUfunction cu_func_ushort
AVFilterContext * dst
dest filter
static const AVFilterPad thumbnail_cuda_outputs[]
static const AVOption thumbnail_cuda_options[]
int histogram[HIST_SIZE]
RGB color distribution histogram of the frame.
int ff_request_frame(AVFilterLink *link)
Request an input frame from the filter at the other end of the link.
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.
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)
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
enum AVPixelFormat sw_format
The pixel format identifying the actual data layout of the hardware frames.
AVPixelFormat
Pixel format.
CUfunction cu_func_uchar2