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,
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,
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;
366 ret =
CHECK_CU(cu->cuCtxPushCurrent(cuda_ctx));
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.
BYTE int const BYTE int src_pitch
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)
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 i(width, name, range_min, range_max)
#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.
const char AVS_Value args
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
static const AVFilterPad inputs[]
HW acceleration through CUDA.
static const AVFilterPad outputs[]
#define FF_ARRAY_ELEMS(a)
static av_cold void uninit(AVFilterContext *ctx)
static int thumbnail(AVFilterContext *ctx, int *histogram, AVFrame *in)
#define AV_LOG_INFO
Standard information.
int linesize[AV_NUM_DATA_POINTERS]
For video, size in bytes of each picture line.
uint8_t * data
The data buffer.
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)\", 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.
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.
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)
enum AVPixelFormat sw_format
The pixel format identifying the actual data layout of the hardware frames.
AVPixelFormat
Pixel format.
CUfunction cu_func_uchar2