PR #22697 opened by nyanmisaka URL: https://code.ffmpeg.org/FFmpeg/FFmpeg/pulls/22697 Patch URL: https://code.ffmpeg.org/FFmpeg/FFmpeg/pulls/22697.patch
Prior to this, the results were not saturated to the range [0.f, 1.f] before being written. The characteristics of the Lanczos filter exposed this issue. Example: ``` ffmpeg -init_hw_device cuda -f lavfi -i testsrc2=s=960x540,format=yuv420p \ -vf hwupload,scale_cuda=format=yuv420p:w=-2:h=720:interp_algo=lanczos \ -c:v h264_nvenc -qp:v 20 -t 1 <OUTPUT> ``` Fix #20784 Signed-off-by: nyanmisaka <[email protected]> Before: <img src="/attachments/d00cbb35-708f-433a-a42e-14e9808122e4" height="200"> After: <img src="/attachments/40244e14-5fb0-4653-8e61-248f52bf71bd" height="200"> >From c58604643760782b9e0562afccbfbefa1e4dd9ec Mon Sep 17 00:00:00 2001 From: nyanmisaka <[email protected]> Date: Fri, 3 Apr 2026 19:28:21 +0800 Subject: [PATCH] avfilter/scale_cuda: fix color bleeding in lanczos scaling Prior to this, the results were not saturated to the range [0.f, 1.f] before being written. The characteristics of the Lanczos filter exposed this issue. Example: ffmpeg -init_hw_device cuda -f lavfi -i testsrc2=s=960x540,format=yuv420p \ -vf hwupload,scale_cuda=format=yuv420p:w=-2:h=720:interp_algo=lanczos \ -c:v h264_nvenc -qp:v 20 -t 1 <OUTPUT> Fix #20784 Signed-off-by: nyanmisaka <[email protected]> --- libavfilter/cuda/vector_helpers.cuh | 23 +++++++++++++++++++++++ libavfilter/vf_scale_cuda.cu | 12 +++++++----- 2 files changed, 30 insertions(+), 5 deletions(-) diff --git a/libavfilter/cuda/vector_helpers.cuh b/libavfilter/cuda/vector_helpers.cuh index 67332ef030..01ed145ec6 100644 --- a/libavfilter/cuda/vector_helpers.cuh +++ b/libavfilter/cuda/vector_helpers.cuh @@ -109,4 +109,27 @@ inline __device__ float4 lerp_scalar<float4>(float4 v0, float4 v1, float t) { ); } +template<typename T> +inline __device__ T saturate(T a) { + return __saturatef(a); +} + +template<> +inline __device__ float2 saturate<float2>(float2 a) { + return make_float2( + saturate(a.x), + saturate(a.y) + ); +} + +template<> +inline __device__ float4 saturate<float4>(float4 a) { + return make_float4( + saturate(a.x), + saturate(a.y), + saturate(a.z), + saturate(a.w) + ); +} + #endif diff --git a/libavfilter/vf_scale_cuda.cu b/libavfilter/vf_scale_cuda.cu index d674c0885a..e71704a8ba 100644 --- a/libavfilter/vf_scale_cuda.cu +++ b/libavfilter/vf_scale_cuda.cu @@ -1162,11 +1162,13 @@ __device__ static inline T Subsample_Bicubic(cudaTextureObject_t tex, #define PIX(x, y) tex2D<floatT>(tex, (x), (y)) return from_floatN<T, floatT>( - apply_coeffs<floatT>(coeffsY, - apply_coeffs<floatT>(coeffsX, PIX(px - 1, py - 1), PIX(px, py - 1), PIX(px + 1, py - 1), PIX(px + 2, py - 1)), - apply_coeffs<floatT>(coeffsX, PIX(px - 1, py ), PIX(px, py ), PIX(px + 1, py ), PIX(px + 2, py )), - apply_coeffs<floatT>(coeffsX, PIX(px - 1, py + 1), PIX(px, py + 1), PIX(px + 1, py + 1), PIX(px + 2, py + 1)), - apply_coeffs<floatT>(coeffsX, PIX(px - 1, py + 2), PIX(px, py + 2), PIX(px + 1, py + 2), PIX(px + 2, py + 2)) + saturate( + apply_coeffs<floatT>(coeffsY, + apply_coeffs<floatT>(coeffsX, PIX(px - 1, py - 1), PIX(px, py - 1), PIX(px + 1, py - 1), PIX(px + 2, py - 1)), + apply_coeffs<floatT>(coeffsX, PIX(px - 1, py ), PIX(px, py ), PIX(px + 1, py ), PIX(px + 2, py )), + apply_coeffs<floatT>(coeffsX, PIX(px - 1, py + 1), PIX(px, py + 1), PIX(px + 1, py + 1), PIX(px + 2, py + 1)), + apply_coeffs<floatT>(coeffsX, PIX(px - 1, py + 2), PIX(px, py + 2), PIX(px + 1, py + 2), PIX(px + 2, py + 2)) + ) ) * factor ); -- 2.52.0 _______________________________________________ ffmpeg-devel mailing list -- [email protected] To unsubscribe send an email to [email protected]
