55 #define DIV_UP(a, b) ( ((a) + (b) - 1) / (b) ) 59 #define CHECK_CU(x) FF_CUDA_CHECK_DL(ctx, s->hwctx->internal->cuda_dl, x) 145 CHECK_CU(cu->cuCtxPopCurrent(&dummy));
213 int out_width,
int out_height)
246 if (s->
passthrough && in_width == out_width && in_height == out_height && in_format == out_format) {
272 CUcontext
dummy, cuda_ctx = device_hwctx->cuda_ctx;
273 CudaFunctions *cu = device_hwctx->internal->cuda_dl;
279 const char *function_infix =
"";
281 extern char vf_scale_cuda_ptx[];
282 extern char vf_scale_cuda_bicubic_ptx[];
286 scaler_ptx = vf_scale_cuda_ptx;
287 function_infix =
"_Nearest";
292 scaler_ptx = vf_scale_cuda_ptx;
293 function_infix =
"_Bilinear";
299 scaler_ptx = vf_scale_cuda_bicubic_ptx;
300 function_infix =
"_Bicubic";
305 scaler_ptx = vf_scale_cuda_bicubic_ptx;
306 function_infix =
"_Lanczos";
315 s->
hwctx = device_hwctx;
326 snprintf(buf,
sizeof(buf),
"Subsample%s_uchar", function_infix);
331 snprintf(buf,
sizeof(buf),
"Subsample%s_uchar2", function_infix);
336 snprintf(buf,
sizeof(buf),
"Subsample%s_uchar4", function_infix);
341 snprintf(buf,
sizeof(buf),
"Subsample%s_ushort", function_infix);
346 snprintf(buf,
sizeof(buf),
"Subsample%s_ushort2", function_infix);
351 snprintf(buf,
sizeof(buf),
"Subsample%s_ushort4", function_infix);
357 CHECK_CU(cu->cuCtxPopCurrent(&dummy));
368 if (((int64_t)
h * inlink->
w) > INT_MAX ||
369 ((int64_t)w * inlink->
h) > INT_MAX)
380 inlink->
w, inlink->
h, outlink->
w, outlink->
h, s->
passthrough ?
" (passthrough)" :
"");
384 outlink->
w*inlink->
h},
397 uint8_t *src_dptr,
int src_width,
int src_height,
int src_pitch,
398 uint8_t *dst_dptr,
int dst_width,
int dst_height,
int dst_pitch,
403 CUdeviceptr dst_devptr = (CUdeviceptr)dst_dptr;
405 void *args_uchar[] = { &tex, &dst_devptr, &dst_width, &dst_height, &dst_pitch,
409 CUDA_TEXTURE_DESC tex_desc = {
411 CU_TR_FILTER_MODE_LINEAR :
412 CU_TR_FILTER_MODE_POINT,
416 CUDA_RESOURCE_DESC res_desc = {
417 .resType = CU_RESOURCE_TYPE_PITCH2D,
418 .res.pitch2D.format = pixel_size == 1 ?
419 CU_AD_FORMAT_UNSIGNED_INT8 :
420 CU_AD_FORMAT_UNSIGNED_INT16,
421 .res.pitch2D.numChannels =
channels,
422 .res.pitch2D.width = src_width,
423 .res.pitch2D.height = src_height,
424 .res.pitch2D.pitchInBytes = src_pitch,
425 .res.pitch2D.devPtr = (CUdeviceptr)src_dptr,
430 dst_pitch /= channels * pixel_size;
432 ret =
CHECK_CU(cu->cuTexObjectCreate(&tex, &res_desc, &tex_desc,
NULL));
436 ret =
CHECK_CU(cu->cuLaunchKernel(func,
442 CHECK_CU(cu->cuTexObjectDestroy(tex));
595 CHECK_CU(cu->cuCtxPopCurrent(&dummy));
621 #define OFFSET(x) offsetof(CUDAScaleContext, x) 622 #define FLAGS (AV_OPT_FLAG_FILTERING_PARAM|AV_OPT_FLAG_VIDEO_PARAM) 637 {
"force_divisible_by",
"enforce that the output resolution is divisible by a defined integer when force_original_aspect_ratio is used",
OFFSET(
force_divisible_by),
AV_OPT_TYPE_INT, { .i64 = 1 }, 1, 256, FLAGS },
668 .
name =
"scale_cuda",
676 .priv_class = &cudascale_class,
678 .
inputs = cudascale_inputs,
int force_original_aspect_ratio
#define FF_FILTER_FLAG_HWFRAME_AWARE
The filter is aware of hardware frames, and any hardware frame context should not be automatically pr...
enum AVPixelFormat format
Output sw format.
void av_buffer_unref(AVBufferRef **buf)
Free a given reference and automatically free the buffer if there are no more references to it...
This structure describes decoded (raw) audio or video data.
planar YUV 4:4:4, 24bpp, (1 Cr & Cb sample per 1x1 Y samples)
#define LIBAVUTIL_VERSION_INT
Main libavfilter public API header.
int h
agreed upon image height
AVCUDADeviceContextInternal * internal
const char * av_default_item_name(void *ptr)
Return the context name.
int width
The allocated dimensions of the frames in this pool.
AVFrame * ff_null_get_video_buffer(AVFilterLink *link, int w, int h)
enum AVPixelFormat format
The pixel format identifying the underlying HW surface type.
void av_frame_move_ref(AVFrame *dst, AVFrame *src)
Move everything contained in src to dst and reset src.
char * w_expr
width expression string
AVCUDADeviceContext * hwctx
CUfunction cu_func_ushort
AVBufferRef * hw_frames_ctx
For hwaccel-format frames, this should be a reference to the AVHWFramesContext describing the frame...
const char * name
Pad name.
const char * class_name
The name of the class; usually it is the same name as the context structure type to which the AVClass...
AVFilterLink ** inputs
array of pointers to input links
static int scalecuda_resize(AVFilterContext *ctx, AVFrame *out, AVFrame *in)
int ff_filter_frame(AVFilterLink *link, AVFrame *frame)
Send a frame of data to the next filter.
AVFrame * av_frame_alloc(void)
Allocate an AVFrame and set its fields to default values.
static enum AVPixelFormat supported_formats[]
static const AVFilterPad cudascale_outputs[]
CUfunction cu_func_ushort4
void * hwctx
The format-specific data, allocated and freed by libavutil along with this context.
static int cudascale_scale(AVFilterContext *ctx, AVFrame *out, AVFrame *in)
#define AV_LOG_VERBOSE
Detailed information.
#define AV_PIX_FMT_YUV444P16
int av_reduce(int *dst_num, int *dst_den, int64_t num, int64_t den, int64_t max)
Reduce a fraction.
#define SCALE_CUDA_PARAM_DEFAULT
static av_cold int init_processing_chain(AVFilterContext *ctx, int in_width, int in_height, int out_width, int out_height)
A filter pad used for either input or output.
A link between two filters.
static av_cold int init_hwframe_ctx(CUDAScaleContext *s, AVBufferRef *device_ctx, int width, int height)
#define AV_LOG_ERROR
Something went wrong and cannot losslessly be recovered.
static int format_is_supported(enum AVPixelFormat fmt)
int ff_scale_eval_dimensions(void *log_ctx, const char *w_expr, const char *h_expr, AVFilterLink *inlink, AVFilterLink *outlink, int *ret_w, int *ret_h)
Parse and evaluate string expressions for width and height.
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
char * h_expr
height expression string
planar YUV 4:2:0, 12bpp, 1 plane for Y and 1 plane for the UV components, which are interleaved (firs...
#define AV_PIX_FMT_0BGR32
like NV12, with 16bpp per component, little-endian
int av_hwframe_ctx_init(AVBufferRef *ref)
Finalize the context before use.
int av_hwframe_get_buffer(AVBufferRef *hwframe_ref, AVFrame *frame, int flags)
Allocate a new frame attached to the given AVHWFramesContext.
static int cudascale_filter_frame(AVFilterLink *link, AVFrame *in)
int w
agreed upon image width
common internal API header
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.
static void bit_depth(AudioStatsContext *s, uint64_t mask, uint64_t imask, AVRational *depth)
static const AVClass cudascale_class
FFmpeg internal API for CUDA.
CUfunction cu_func_uchar2
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 int call_resize_kernel(AVFilterContext *ctx, CUfunction func, int channels, uint8_t *src_dptr, int src_width, int src_height, int src_pitch, uint8_t *dst_dptr, int dst_width, int dst_height, int dst_pitch, int pixel_size, int bit_depth)
static const AVOption options[]
static const AVFilterPad outputs[]
#define FF_ARRAY_ELEMS(a)
int linesize[AV_NUM_DATA_POINTERS]
For video, size in bytes of each picture line.
AVFilter ff_vf_scale_cuda
uint8_t * data
The data buffer.
AVRational sample_aspect_ratio
Sample aspect ratio for the video frame, 0/1 if unknown/unspecified.
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
#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
static av_cold void cudascale_uninit(AVFilterContext *ctx)
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.
static av_cold int cudascale_config_props(AVFilterLink *outlink)
const char * name
Filter name.
AVRational sample_aspect_ratio
agreed upon sample aspect ratio
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[]
int ff_scale_adjust_dimensions(AVFilterLink *inlink, int *ret_w, int *ret_h, int force_original_aspect_ratio, int force_divisible_by)
Transform evaluated width and height obtained from ff_scale_eval_dimensions into actual target width ...
static const AVFilterPad cudascale_inputs[]
void av_frame_unref(AVFrame *frame)
Unreference all the buffers referenced by frame and reset the frame fields.
static AVFrame * cudascale_get_video_buffer(AVFilterLink *inlink, int w, int h)
uint8_t * data[AV_NUM_DATA_POINTERS]
pointer to the picture/channel planes.
CUfunction cu_func_ushort2
AVBufferRef * device_ref
A reference to the parent AVHWDeviceContext.
AVFrame * ff_default_get_video_buffer(AVFilterLink *link, int w, int h)
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)
common internal and external API header
static av_cold int cudascale_init(AVFilterContext *ctx)
AVBufferRef * av_hwframe_ctx_alloc(AVBufferRef *device_ref_in)
Allocate an AVHWFramesContext tied to a given device context.
static int cudascale_query_formats(AVFilterContext *ctx)
AVBufferRef * av_buffer_ref(AVBufferRef *buf)
Create a new reference to an AVBuffer.
enum AVPixelFormat in_fmt
AVFilterContext * dst
dest filter
CUfunction cu_func_uchar4
AVRational av_mul_q(AVRational b, AVRational c)
Multiply two rationals.
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.
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.
enum AVPixelFormat out_fmt
int av_frame_copy_props(AVFrame *dst, const AVFrame *src)
Copy only "metadata" fields from src to dst.
#define AV_PIX_FMT_0RGB32