PR #23741 opened by Diego de Souza (ddesouza) URL: https://code.ffmpeg.org/FFmpeg/FFmpeg/pulls/23741 Patch URL: https://code.ffmpeg.org/FFmpeg/FFmpeg/pulls/23741.patch
This PR fixes single-axis scaling through the existing generic weight-LUT path and improves the quality of Lanczos scaling in `scale_cuda`. The generic path previously ran an empty horizontal intermediate pass when only one dimension changed, producing an all-zero frame. The intermediate pass is now allocated and executed only when temporary filter weights exist. Lanczos now uses the existing generic separable filter path by default. This: - Generates normalized Lanczos weights through libswscale. - Widens the filter when downscaling for proper anti-aliasing. - Supports a configurable radius from 1 to 10 through `param`. - Uses a floating-point intermediate for two-axis scaling to avoid quantization between the horizontal and vertical passes. - Saturates and rounds generic output to the input format's valid code range while preserving MSB alignment. The legacy four-tap Lanczos implementation remains available with `use_filters=0`. The `scale_cuda` documentation is updated accordingly. No public API or ABI changes are introduced. No new FATE samples are included because these paths require CUDA hardware. >From f42c983edfcd6cb3c9609fc816098cb16128e498 Mon Sep 17 00:00:00 2001 From: Diego de Souza <[email protected]> Date: Wed, 8 Jul 2026 21:22:24 +0000 Subject: [PATCH 1/2] avfilter/vf_scale_cuda: skip empty intermediate pass The generic filter path only generates FILTER_TMP weights when both dimensions change. It nevertheless always allocated and ran the intermediate horizontal pass, so single-axis scaling used a filter with no weights and produced an all-zero frame. Only allocate and run the intermediate pass when temporary filter weights exist. Signed-off-by: Diego de Souza <[email protected]> --- libavfilter/vf_scale_cuda.c | 10 ++++++---- 1 file changed, 6 insertions(+), 4 deletions(-) diff --git a/libavfilter/vf_scale_cuda.c b/libavfilter/vf_scale_cuda.c index 240e40d43b..e7b6f78509 100644 --- a/libavfilter/vf_scale_cuda.c +++ b/libavfilter/vf_scale_cuda.c @@ -666,9 +666,11 @@ static av_cold int cudascale_setup_filters(AVFilterContext *ctx) } } - ret = inter_buf_init(ctx, outlink->w, inlink->h); - if (ret < 0) - goto fail; + if (pass_x == FILTER_TMP) { + ret = inter_buf_init(ctx, outlink->w, inlink->h); + if (ret < 0) + goto fail; + } ret = 0; @@ -915,7 +917,7 @@ static int cudascale_scale(AVFilterContext *ctx, AVFrame *out, AVFrame *in) goto fail; const CUDATex *src = &in_tex; - if (s->use_filters) { + if (s->use_filters && s->filters[FILTER_TMP].weights) { /* Handle first pass separately */ s->inter_tex.color_range = in->color_range; ret = scalecuda_resize(ctx, FILTER_TMP, &s->inter_tex, src); -- 2.52.0 >From 343665f22b290a52b82f9733e1050c3a3fa55f8e Mon Sep 17 00:00:00 2001 From: Diego de Souza <[email protected]> Date: Sat, 27 Jun 2026 20:54:14 +0000 Subject: [PATCH 2/2] avfilter/vf_scale_cuda: use high-quality separable Lanczos The fixed-function Lanczos implementation uses the four-tap bicubic framework, limiting its filter window and providing no proper anti-aliasing when downscaling. Use the generic separable filter path for Lanczos by default. Generate normalized Lanczos weights through libswscale, widen the filter when downscaling, and honor param as the configurable filter radius. For two-axis scaling, store the horizontal-pass result in a floating- point intermediate texture to avoid quantization before the vertical pass. Also saturate and round generic filter output to the input format's valid code range while preserving MSB alignment. The legacy four-tap implementation remains available with use_filters=0. Signed-off-by: Diego de Souza <[email protected]> --- doc/filters.texi | 12 ++-- libavfilter/vf_scale_cuda.c | 54 +++++++++++----- libavfilter/vf_scale_cuda.cu | 117 ++++++++++++++++++++++++++++++++--- libavfilter/vf_scale_cuda.h | 2 + 4 files changed, 158 insertions(+), 27 deletions(-) diff --git a/doc/filters.texi b/doc/filters.texi index 5497421f34..895985fcfa 100644 --- a/doc/filters.texi +++ b/doc/filters.texi @@ -27467,7 +27467,8 @@ Bicubic This is the default. @item lanczos -Lanczos +High-quality separable Lanczos (windowed sinc). The filter radius can be +set with @option{param}; the default is 3. @end table @@ -27489,12 +27490,13 @@ parameters. This is the default behaviour. If set to 1, filter with a generic weight LUT instead of using fixed-function shader kernels. May be faster or slower depending on the hardware. A value of @code{auto} (the default) enables this automatically when required for -correct anti-aliasing when downscaling. +correct anti-aliasing when downscaling, and for high-quality Lanczos scaling. +Set this to 0 to use the legacy fixed 4-tap Lanczos kernel. @item param -Algorithm-Specific parameter. - -Affects the curves of the bicubic algorithm. +Algorithm-specific parameter. For @var{bicubic}, this affects the filter +curve. For @var{lanczos}, this sets the filter radius from 1 to 10 (default 3) +when using the generic filter path; it has no effect with @option{use_filters=0}. @item force_original_aspect_ratio @item force_divisible_by diff --git a/libavfilter/vf_scale_cuda.c b/libavfilter/vf_scale_cuda.c index e7b6f78509..3c62043ed7 100644 --- a/libavfilter/vf_scale_cuda.c +++ b/libavfilter/vf_scale_cuda.c @@ -263,7 +263,8 @@ fail: return ret; } -static av_cold int inter_buf_init(AVFilterContext *ctx, int out_width, int in_height) +static av_cold int inter_buf_init(AVFilterContext *ctx, int out_width, int in_height, + int use_float) { CUDAScaleContext *s = ctx->priv; CudaFunctions *cu = s->hwctx->internal->cuda_dl; @@ -285,7 +286,8 @@ static av_cold int inter_buf_init(AVFilterContext *ctx, int out_width, int in_he const int sub_y = is_chroma ? s->inter_tex.log2_chroma_h : 0; const int plane_w = AV_CEIL_RSHIFT(out_width, sub_x); const int plane_h = AV_CEIL_RSHIFT(in_height, sub_y); - const int sizeof_pixel = (s->in_plane_depths[i] <= 8 ? 1 : 2) * + const int sizeof_pixel = (use_float ? sizeof(float) : + s->in_plane_depths[i] <= 8 ? 1 : 2) * s->in_plane_channels[i]; size_t pitch; @@ -303,7 +305,8 @@ static av_cold int inter_buf_init(AVFilterContext *ctx, int out_width, int in_he CUDA_RESOURCE_DESC res_desc = { .resType = CU_RESOURCE_TYPE_PITCH2D, - .res.pitch2D.format = s->in_plane_depths[i] <= 8 ? + .res.pitch2D.format = use_float ? CU_AD_FORMAT_FLOAT : + s->in_plane_depths[i] <= 8 ? CU_AD_FORMAT_UNSIGNED_INT8 : CU_AD_FORMAT_UNSIGNED_INT16, .res.pitch2D.numChannels = s->in_plane_channels[i], @@ -418,12 +421,18 @@ static av_cold int init_processing_chain(AVFilterContext *ctx, int in_width, int in_format == out_format && s->interp_algo == INTERP_ALGO_DEFAULT) s->interp_algo = INTERP_ALGO_NEAREST; - if (s->interp_algo == INTERP_ALGO_NEAREST) { + if (in_width == out_width && in_height == out_height) { + /* No weights are generated when there is nothing to scale. */ s->use_filters = 0; - } else if (s->use_filters < 0 && (out_width < in_width || out_height < in_height)) - s->use_filters = 1; /* downscaling; needed for anti-aliasing */ - else if (s->use_filters < 0) + } else if (s->interp_algo == INTERP_ALGO_NEAREST) { s->use_filters = 0; + } else if (s->use_filters < 0) { + /* Lanczos needs the generic path for its full windowed-sinc + * kernel. Other algorithms need it when downscaling for correct + * anti-aliasing. */ + s->use_filters = s->interp_algo == INTERP_ALGO_LANCZOS || + out_width < in_width || out_height < in_height; + } } outl->hw_frames_ctx = av_buffer_ref(s->frames_ctx); @@ -511,15 +520,24 @@ static av_cold int cudascale_load_functions(AVFilterContext *ctx) if (s->use_filters) { /* Intermediate pass is always horizontal */ - snprintf(buf, sizeof(buf), "Subsample_Generic_h_%s_%s", in_fmt_name, in_fmt_name); + const char *tmp_infix = s->interp_algo == INTERP_ALGO_LANCZOS && + s->filters[FILTER_TMP].weights ? + "Generic_float_h" : "Generic_h"; + + snprintf(buf, sizeof(buf), "Subsample_%s_%s_%s", tmp_infix, + in_fmt_name, in_fmt_name); ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_func[FILTER_TMP], s->cu_module, buf)); if (ret < 0) goto fail; - snprintf(buf, sizeof(buf), "Subsample_Generic_h_%s_%s_uv", in_fmt_name, in_fmt_name); - ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_uv[FILTER_TMP], s->cu_module, buf)); - if (ret < 0) - goto fail; + if (s->in_planes > 1) { + snprintf(buf, sizeof(buf), "Subsample_%s_%s_%s_uv", tmp_infix, + in_fmt_name, in_fmt_name); + ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_uv[FILTER_TMP], + s->cu_module, buf)); + if (ret < 0) + goto fail; + } } fail: @@ -546,7 +564,11 @@ static av_cold int cudascale_filter_init(AVFilterContext *ctx, switch (s->interp_algo) { case INTERP_ALGO_NEAREST: return 0; /* no weights needed */ case INTERP_ALGO_BILINEAR: params.scaler = SWS_SCALE_BILINEAR; break; - case INTERP_ALGO_LANCZOS: params.scaler = SWS_SCALE_LANCZOS; break; + case INTERP_ALGO_LANCZOS: + params.scaler = SWS_SCALE_LANCZOS; + if (s->param != SCALE_CUDA_PARAM_DEFAULT) + params.scaler_params[0] = s->param; + break; case INTERP_ALGO_DEFAULT: case INTERP_ALGO_BICUBIC: params.scaler = SWS_SCALE_BICUBIC; @@ -667,7 +689,8 @@ static av_cold int cudascale_setup_filters(AVFilterContext *ctx) } if (pass_x == FILTER_TMP) { - ret = inter_buf_init(ctx, outlink->w, inlink->h); + ret = inter_buf_init(ctx, outlink->w, inlink->h, + s->interp_algo == INTERP_ALGO_LANCZOS); if (ret < 0) goto fail; } @@ -845,6 +868,9 @@ static int call_resize_kernel(AVFilterContext *ctx, CUfunction func, .src_height = src_height, .param = s->param, .mpeg_range = mpeg_range, + .src_depth = s->in_desc->comp[0].depth, + .src_max_value = ((1U << s->in_desc->comp[0].depth) - 1) << + s->in_desc->comp[0].shift, }; if (filter) { diff --git a/libavfilter/vf_scale_cuda.cu b/libavfilter/vf_scale_cuda.cu index 01398c3db8..2e957c8a2a 100644 --- a/libavfilter/vf_scale_cuda.cu +++ b/libavfilter/vf_scale_cuda.cu @@ -28,7 +28,8 @@ using subsample_function_t = T (*)(cudaTextureObject_t tex, int xo, int yo, int dst_width, int dst_height, int src_left, int src_top, int src_width, int src_height, - int bit_depth, float param, + int bit_depth, int src_depth, int max_value, + float param, const float *weights, const int *offsets, int filter_size); @@ -92,7 +93,7 @@ static inline __device__ ushort conv_16to10pl(ushort in) __device__ static inline void N(cudaTextureObject_t src_tex[4], T *dst[4], int xo, int yo, \ int dst_width, int dst_height, int dst_pitch, \ int src_left, int src_top, int src_width, int src_height, \ - float param, int mpeg_range, \ + int src_depth, int max_value, float param, int mpeg_range, \ const float *weights, const int *offsets, int filter_size) #define SUB_F(m, plane) \ @@ -100,7 +101,7 @@ static inline __device__ ushort conv_16to10pl(ushort in) dst_width, dst_height, \ src_left, src_top, \ src_width, src_height, \ - in_bit_depth, param, \ + in_bit_depth, src_depth, max_value, param, \ weights, offsets, filter_size) // FFmpeg passes pitch in bytes, CUDA uses potentially larger types @@ -1099,7 +1100,8 @@ __device__ static inline T Subsample_Nearest(cudaTextureObject_t tex, int dst_width, int dst_height, int src_left, int src_top, int src_width, int src_height, - int bit_depth, float param, + int bit_depth, int src_depth, int max_value, + float param, const float *weights, const int *offsets, int filter_size) { @@ -1117,7 +1119,8 @@ __device__ static inline T Subsample_Bilinear(cudaTextureObject_t tex, int dst_width, int dst_height, int src_left, int src_top, int src_width, int src_height, - int bit_depth, float param, + int bit_depth, int src_depth, int max_value, + float param, const float *weights, const int *offsets, int filter_size) { @@ -1151,7 +1154,8 @@ __device__ static inline T Subsample_Bicubic(cudaTextureObject_t tex, int dst_width, int dst_height, int src_left, int src_top, int src_width, int src_height, - int bit_depth, float param, + int bit_depth, int src_depth, int max_value, + float param, const float *weights, const int *offsets, int filter_size) { @@ -1197,11 +1201,14 @@ __device__ static inline T Subsample_Generic(cudaTextureObject_t tex, int dst_width, int dst_height, int src_left, int src_top, int src_width, int src_height, - int bit_depth, float param, + int bit_depth, int src_depth, int max_value, + float param, const float *weights, const int *offsets, int filter_size) { const float factor = bit_depth > 8 ? 0xFFFF : 0xFF; + const float code_max = (1U << src_depth) - 1; + const float sample_step = max_value / code_max; floatT sum; vec_set_scalar(sum, 0.0f); @@ -1220,7 +1227,29 @@ __device__ static inline T Subsample_Generic(cudaTextureObject_t tex, sum += tex2D<floatT>(tex, x, y + i) * col[i]; } - return from_floatN<T, floatT>(sum * factor); + return from_floatN<T, floatT>( + saturate_rintf(sum * (factor / max_value), code_max) * sample_step + ); +} + +template<typename T> +__device__ static inline floatT Subsample_GenericFloat(cudaTextureObject_t tex, + int xo, int yo, + int src_left, int src_top, + const float *weights, + const int *offsets, + int filter_size) +{ + const float *row = &weights[xo * filter_size]; + const float x = 0.5f + src_left + offsets[xo]; + const float y = 0.5f + src_top + yo; + floatT sum; + + vec_set_scalar(sum, 0.0f); + for (int i = 0; i < filter_size; i++) + sum += tex2D<floatT>(tex, x + i, y) * row[i]; + + return sum; } /// --- FUNCTION EXPORTS --- @@ -1244,6 +1273,7 @@ __device__ static inline T Subsample_Generic(cudaTextureObject_t tex, params.dst_width, params.dst_height, params.dst_pitch, \ params.src_left, params.src_top, \ params.src_width, params.src_height, \ + params.src_depth, params.src_max_value, \ params.param, params.mpeg_range, \ (const float*) params.weights, \ (const int*) params.offsets, \ @@ -1452,4 +1482,75 @@ GENERIC_KERNELS_RGB(rgb0) GENERIC_KERNELS_RGB(bgr0) GENERIC_KERNELS_RGB(rgba) GENERIC_KERNELS_RGB(bgra) + +#define GENERIC_FLOAT_KERNEL_Y(C, T, F) \ + __global__ void Subsample_Generic_float_h_##C##_##C( \ + CUDAScaleKernelParams params) \ + { \ + int xo = blockIdx.x * blockDim.x + threadIdx.x; \ + int yo = blockIdx.y * blockDim.y + threadIdx.y; \ + if (yo >= params.dst_height || xo >= params.dst_width) return; \ + F *dst = (F *)((char *)params.dst[0] + \ + yo * params.dst_pitch); \ + dst[xo] = Subsample_GenericFloat<T>( \ + params.src_tex[0], xo, yo, params.src_left, params.src_top, \ + (const float *)params.weights, (const int *)params.offsets, \ + params.filter_size); \ + } + +#define GENERIC_FLOAT_KERNEL_UV_PLANAR(C, T) \ + __global__ void Subsample_Generic_float_h_##C##_##C##_uv( \ + CUDAScaleKernelParams params) \ + { \ + int xo = blockIdx.x * blockDim.x + threadIdx.x; \ + int yo = blockIdx.y * blockDim.y + threadIdx.y; \ + if (yo >= params.dst_height || xo >= params.dst_width) return; \ + float *dst_u = (float *)((char *)params.dst[1] + \ + yo * params.dst_pitch); \ + float *dst_v = (float *)((char *)params.dst[2] + \ + yo * params.dst_pitch); \ + const float *weights = (const float *)params.weights; \ + const int *offsets = (const int *)params.offsets; \ + dst_u[xo] = Subsample_GenericFloat<T>(params.src_tex[1], xo, yo, \ + params.src_left, params.src_top, weights, offsets, \ + params.filter_size); \ + dst_v[xo] = Subsample_GenericFloat<T>(params.src_tex[2], xo, yo, \ + params.src_left, params.src_top, weights, offsets, \ + params.filter_size); \ + } + +#define GENERIC_FLOAT_KERNEL_UV_SEMI(C, T, F) \ + __global__ void Subsample_Generic_float_h_##C##_##C##_uv( \ + CUDAScaleKernelParams params) \ + { \ + int xo = blockIdx.x * blockDim.x + threadIdx.x; \ + int yo = blockIdx.y * blockDim.y + threadIdx.y; \ + if (yo >= params.dst_height || xo >= params.dst_width) return; \ + F *dst = (F *)((char *)params.dst[1] + \ + yo * params.dst_pitch); \ + dst[xo] = Subsample_GenericFloat<T>( \ + params.src_tex[1], xo, yo, params.src_left, params.src_top, \ + (const float *)params.weights, (const int *)params.offsets, \ + params.filter_size); \ + } + +#define GENERIC_FLOAT_KERNEL_PLANAR(C, T) \ + GENERIC_FLOAT_KERNEL_Y(C, T, float) \ + GENERIC_FLOAT_KERNEL_UV_PLANAR(C, T) + +#define GENERIC_FLOAT_KERNEL_SEMI(C, T, TUV) \ + GENERIC_FLOAT_KERNEL_Y(C, T, float) \ + GENERIC_FLOAT_KERNEL_UV_SEMI(C, TUV, float2) + +GENERIC_FLOAT_KERNEL_PLANAR(planar8, uchar) +GENERIC_FLOAT_KERNEL_PLANAR(planar10, ushort) +GENERIC_FLOAT_KERNEL_PLANAR(planar16, ushort) +GENERIC_FLOAT_KERNEL_SEMI(semiplanar8, uchar, uchar2) +GENERIC_FLOAT_KERNEL_SEMI(semiplanar10, ushort, ushort2) +GENERIC_FLOAT_KERNEL_SEMI(semiplanar16, ushort, ushort2) +GENERIC_FLOAT_KERNEL_Y(rgb0, uchar4, float4) +GENERIC_FLOAT_KERNEL_Y(bgr0, uchar4, float4) +GENERIC_FLOAT_KERNEL_Y(rgba, uchar4, float4) +GENERIC_FLOAT_KERNEL_Y(bgra, uchar4, float4) + } diff --git a/libavfilter/vf_scale_cuda.h b/libavfilter/vf_scale_cuda.h index 391631699e..dd15102e8c 100644 --- a/libavfilter/vf_scale_cuda.h +++ b/libavfilter/vf_scale_cuda.h @@ -44,6 +44,8 @@ typedef struct { int src_height; float param; int mpeg_range; + int src_depth; + int src_max_value; /* Weights for the generic filter kernel */ CUdeviceptr weights; -- 2.52.0 _______________________________________________ ffmpeg-devel mailing list -- [email protected] To unsubscribe send an email to [email protected]
