PR #25324 opened by Forgejo_Fairy URL: https://code.ffmpeg.org/FFmpeg/FFmpeg/pulls/25324 Patch URL: https://code.ffmpeg.org/FFmpeg/FFmpeg/pulls/25324.patch
Address the release/9.0 temporal and build defects reported in #25105 by backporting two public master commits: - ffb15c94c286: link cuda/load_helper.o into the three CUDA filters that call it but omit the dependency. - 587db45caf2a: correct bwdif_cuda's temporal frame selection and duplicated row tap, from #23619. Preserve the original authorship and record the source commits with cherry-pick trailers. Validation on release/9.0 c89d6ee1ca84 with an RTX 5090: the minimal bwdif_cuda build initially fails with an undefined reference to ff_cuda_load_module and links after the loader backport. With only that build fix applied, the supplied H.264/NVDEC example reproduces the report's output exactly. Adding the temporal backport removes the interior-row differences. Both field orders and both send_frame/send_field modes were exercised. The loader fix was additionally verified with identical minimal builds at ffb15c94c286 and its parent: the parent fails and the fixing commit succeeds. git diff --check passes. The picture-edge defect remains. Its new patch is proposed separately for master and is not included in this release backport. >From 8f488ab6b57405e4751909133810f0dcd55ed2f0 Mon Sep 17 00:00:00 2001 From: lusoris <[email protected]> Date: Sun, 18 Oct 2026 14:50:52 +0200 Subject: [PATCH 1/2] avfilter/Makefile: link cuda/load_helper.o into three CUDA filters bilateral_cuda, bwdif_cuda and chromakey_cuda call ff_cuda_load_module(), which lives in cuda/load_helper.c, but their OBJS lines do not pull that object in; the other CUDA filters that load a module list it. A build where one of these is the only CUDA filter fails to link, for example: vf_bwdif_cuda.c:312: undefined reference to `ff_cuda_load_module' Reproduced for each filter with --disable-everything --enable-avfilter --enable-filter=<name> --enable-cuda-llvm --enable-ffnvcodec; each links after this change. Signed-off-by: lusoris <[email protected]> (cherry picked from commit ffb15c94c286c9d1af911991b3ad08724427372a) Assisted-by: Fairy --- libavfilter/Makefile | 8 +++++--- 1 file changed, 5 insertions(+), 3 deletions(-) diff --git a/libavfilter/Makefile b/libavfilter/Makefile index ce40cf3e65b..1c8f9196718 100644 --- a/libavfilter/Makefile +++ b/libavfilter/Makefile @@ -207,7 +207,8 @@ OBJS-$(CONFIG_BACKGROUNDKEY_FILTER) += vf_backgroundkey.o OBJS-$(CONFIG_BBOX_FILTER) += bbox.o vf_bbox.o OBJS-$(CONFIG_BENCH_FILTER) += f_bench.o OBJS-$(CONFIG_BILATERAL_FILTER) += vf_bilateral.o -OBJS-$(CONFIG_BILATERAL_CUDA_FILTER) += vf_bilateral_cuda.o vf_bilateral_cuda.ptx.o +OBJS-$(CONFIG_BILATERAL_CUDA_FILTER) += vf_bilateral_cuda.o vf_bilateral_cuda.ptx.o \ + cuda/load_helper.o OBJS-$(CONFIG_BITPLANENOISE_FILTER) += vf_bitplanenoise.o OBJS-$(CONFIG_BLACKDETECT_FILTER) += vf_blackdetect.o OBJS-$(CONFIG_BLACKDETECT_VULKAN_FILTER) += vf_blackdetect_vulkan.o @@ -222,14 +223,15 @@ OBJS-$(CONFIG_BOXBLUR_OPENCL_FILTER) += vf_avgblur_opencl.o opencl.o \ opencl/avgblur.o boxblur.o OBJS-$(CONFIG_BWDIF_FILTER) += vf_bwdif.o bwdifdsp.o yadif_common.o OBJS-$(CONFIG_BWDIF_CUDA_FILTER) += vf_bwdif_cuda.o vf_bwdif_cuda.ptx.o \ - yadif_common.o + yadif_common.o cuda/load_helper.o OBJS-$(CONFIG_BWDIF_VULKAN_FILTER) += vf_bwdif_vulkan.o yadif_common.o vulkan.o vulkan_filter.o OBJS-$(CONFIG_CAS_FILTER) += vf_cas.o OBJS-$(CONFIG_CCREPACK_FILTER) += vf_ccrepack.o OBJS-$(CONFIG_CHROMABER_VULKAN_FILTER) += vf_chromaber_vulkan.o vulkan.o vulkan_filter.o OBJS-$(CONFIG_CHROMAHOLD_FILTER) += vf_chromakey.o OBJS-$(CONFIG_CHROMAKEY_FILTER) += vf_chromakey.o -OBJS-$(CONFIG_CHROMAKEY_CUDA_FILTER) += vf_chromakey_cuda.o vf_chromakey_cuda.ptx.o +OBJS-$(CONFIG_CHROMAKEY_CUDA_FILTER) += vf_chromakey_cuda.o vf_chromakey_cuda.ptx.o \ + cuda/load_helper.o OBJS-$(CONFIG_CHROMANR_FILTER) += vf_chromanr.o OBJS-$(CONFIG_CHROMASHIFT_FILTER) += vf_chromashift.o -- 2.54.0 >From 8c15cf46a72cf353d979c59b2eac5209953b4fe5 Mon Sep 17 00:00:00 2001 From: Philip Langdale <[email protected]> Date: Fri, 26 Jun 2026 22:56:34 -0700 Subject: [PATCH 2/2] avfilter/vf_bwdif_cuda: fix typos in algorithm It's been years, but took another look at the bwdif_cuda implementation and there are a couple of typos sitting in there. Found them when I was doing a comparison with the Vulkan implementation. This probably explains the small PSNR differences we've noted in the past. (cherry picked from commit 587db45caf2a5be4bfdc1704b89c662681ad1142) Assisted-by: Fairy --- libavfilter/vf_bwdif_cuda.cu | 18 +++++++++--------- 1 file changed, 9 insertions(+), 9 deletions(-) diff --git a/libavfilter/vf_bwdif_cuda.cu b/libavfilter/vf_bwdif_cuda.cu index 3d4c29d8c39..c09b6f7f36c 100644 --- a/libavfilter/vf_bwdif_cuda.cu +++ b/libavfilter/vf_bwdif_cuda.cu @@ -91,7 +91,7 @@ __inline__ __device__ T filter(T cur_prefs3, T cur_prefs, T cur_mrefs, T cur_mre if (abs(c - e) > temporal_diff0) { interpol = (((coef_hf[0] * (prev2_0 + next2_0) - coef_hf[1] * (prev2_mrefs2 + next2_mrefs2 + prev2_prefs2 + next2_prefs2) - + coef_hf[2] * (prev2_mrefs4 + next2_mrefs4 + prev2_prefs4 + next2_mrefs4)) >> 2) + + coef_hf[2] * (prev2_mrefs4 + next2_mrefs4 + prev2_prefs4 + next2_prefs4)) >> 2) + coef_lf[0] * (c + e) - coef_lf[1] * (cur_mrefs3 + cur_prefs3)) >> 13; } else { interpol = (coef_sp[0] * (c + e) - coef_sp[1] * (cur_mrefs3 + cur_prefs3)) >> 13; @@ -146,10 +146,10 @@ __inline__ __device__ void bwdif_single(T *dst, // Calculate temporal prediction int is_second_field = !(parity ^ tff); - cudaTextureObject_t prev2 = prev; - cudaTextureObject_t prev1 = is_second_field ? cur : prev; - cudaTextureObject_t next1 = is_second_field ? next : cur; - cudaTextureObject_t next2 = next; + cudaTextureObject_t prev2 = is_second_field ? cur : prev; + cudaTextureObject_t prev1 = prev; + cudaTextureObject_t next1 = next; + cudaTextureObject_t next2 = is_second_field ? next : cur; T prev2_prefs4 = tex2D<T>(prev2, xo, yo + 4); T prev2_prefs2 = tex2D<T>(prev2, xo, yo + 2); @@ -213,10 +213,10 @@ __inline__ __device__ void bwdif_double(T *dst, int is_second_field = !(parity ^ tff); - cudaTextureObject_t prev2 = prev; - cudaTextureObject_t prev1 = is_second_field ? cur : prev; - cudaTextureObject_t next1 = is_second_field ? next : cur; - cudaTextureObject_t next2 = next; + cudaTextureObject_t prev2 = is_second_field ? cur : prev; + cudaTextureObject_t prev1 = prev; + cudaTextureObject_t next1 = next; + cudaTextureObject_t next2 = is_second_field ? next : cur; T prev2_prefs4 = tex2D<T>(prev2, xo, yo + 4); T prev2_prefs2 = tex2D<T>(prev2, xo, yo + 2); -- 2.54.0 _______________________________________________ ffmpeg-devel mailing list -- [email protected] To unsubscribe send an email to [email protected]
