This is an automated email from the git hooks/post-receive script. Git pushed a commit to branch master in repository ffmpeg.
commit eb4c99d49c455bad5e3747135b3df95fcacaa09f Author: Diego de Souza <[email protected]> AuthorDate: Fri Jul 10 19:25:13 2026 +0200 Commit: Niklas Haas <[email protected]> CommitDate: Wed Aug 26 09:52:59 2026 +0000 avfilter/vf_scale_cuda: use floating-point Lanczos intermediates Store horizontal Lanczos results in floating-point textures so two-axis filtering quantizes only after the final pass. Signed-off-by: Diego de Souza <[email protected]> --- libavfilter/vf_scale_cuda.c | 12 ++++-- libavfilter/vf_scale_cuda.cu | 91 ++++++++++++++++++++++++++++++++++++++++++++ 2 files changed, 99 insertions(+), 4 deletions(-) diff --git a/libavfilter/vf_scale_cuda.c b/libavfilter/vf_scale_cuda.c index 6f924b42a6..ca7f375760 100644 --- a/libavfilter/vf_scale_cuda.c +++ b/libavfilter/vf_scale_cuda.c @@ -372,7 +372,7 @@ fail: static int inter_buf_init(AVFilterContext *ctx, CUDATex *tex, int out_width, int in_height, - unsigned int planes) + int use_float, unsigned int planes) { CUDAScaleContext *s = ctx->priv; CudaFunctions *cu = s->hwctx->internal->cuda_dl; @@ -395,7 +395,8 @@ static int inter_buf_init(AVFilterContext *ctx, CUDATex *tex, const int sub_y = is_chroma ? 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]; if (!(planes & plane)) @@ -416,7 +417,8 @@ static int inter_buf_init(AVFilterContext *ctx, CUDATex *tex, 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], @@ -617,7 +619,8 @@ static av_cold int cudascale_load_functions(AVFilterContext *ctx) if (s->use_filters) { static const char *const infix[] = { "Generic_h", "Generic_v" }; - const char *tmp_infix = "Generic_h"; + const char *tmp_infix = s->interp_algo == INTERP_ALGO_LANCZOS ? + "Generic_float_h" : "Generic_h"; s->interp_use_linear = 0; s->interp_as_integer = 0; @@ -862,6 +865,7 @@ static int cudascale_filter_set_init(AVFilterContext *ctx, if (pass_x == FILTER_TMP) { ret = inter_buf_init(ctx, &set->inter_tex, outlink->w, in_height, + s->interp_algo == INTERP_ALGO_LANCZOS, set->tmp_planes); if (ret < 0) goto fail; diff --git a/libavfilter/vf_scale_cuda.cu b/libavfilter/vf_scale_cuda.cu index 8c9ed0a41f..7243777651 100644 --- a/libavfilter/vf_scale_cuda.cu +++ b/libavfilter/vf_scale_cuda.cu @@ -1232,6 +1232,26 @@ __device__ static inline T Subsample_Generic(cudaTextureObject_t tex, ); } +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 --- #define KERNEL_ARGS(T) CUDAScaleKernelParams params @@ -1462,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) + } -- To stop receiving notification emails like this one, please contact [email protected]. _______________________________________________ ffmpeg-cvslog mailing list -- [email protected] To unsubscribe send an email to [email protected]
