PR #25323 opened by Forgejo_Fairy URL: https://code.ffmpeg.org/FFmpeg/FFmpeg/pulls/25323 Patch URL: https://code.ffmpeg.org/FFmpeg/FFmpeg/pulls/25323.patch
Fixes #25105. CUDA bwdif reads beyond the picture boundaries through clamped textures and applies interior temporal interpolation at the edges. On the reported 128×64 TFF clip, rows 61 and 63 contain 255 and 125 where CPU bwdif produces 235. Apply oittaa's attached CUDA patch, preserving authorship. Match CPU bwdif's missing-row substitution for intra processing and its boundary interpolation and spatial-check conditions for temporal processing, in both scalar and vector kernels. Add 16 GPU FATE cases covering 8- and 16-bit planar and semiplanar formats, both field orders and both output modes. References come from CPU bwdif; the tests require no external sample and are included only in fate-hw or the explicit fate-filter-bwdif_cuda target. Validation on Linux with an RTX 5090 and NVIDIA driver 595.84: - The report's H.264/NVDEC test matches CPU output byte for byte in every plane for both modes and field orders after the fix. - All 16 new FATE cases fail with the original kernel and pass with the patch. - Another 64 deterministic five-frame comparisons pass exactly across four formats, four frame heights, both modes and both field orders. - git diff --check passes. The separate release temporal errors and loader dependency are handled by existing public master commits, not by this change. From e9fddb30342cbb6ec1c339d994fd07699ed7a275 Mon Sep 17 00:00:00 2001 From: oittaa <[email protected]> Date: Thu, 15 Oct 2026 20:18:39 +0200 Subject: [PATCH 1/2] lavfi/vf_bwdif_cuda: filter the top and bottom edges like bwdif The kernel has no case for the rows near the top and bottom edges. It reads the rows past an edge from the texture, which clamps them to the outermost row, a row of the field being replaced. So the outermost rebuilt row comes out halfway between bwdif's value and the old row, and the rebuilt row next to it overshoots. bwdif replaces a row past an edge with the nearest row on the other side. It filters the 4 rows nearest each edge with filter_edge, which averages the rows next to it and uses the rows 2 away only when both are in the picture. Do the same here, in the intra path too. Tested on an RTX 4080 SUPER with interlaced H.264 clips decoded by NVDEC: white blocks on black touching both edges at 128x64 and 1920x1080, and the blocks over moving waves at 1920x1080. With top and bottom field first, and send_frame and send_field, the output now equals bwdif's in every plane. Before, a white (235) block touching the bottom edge gave 125 in row 63 and 255 in row 61, where bwdif gives 235. Co-Authored-By: Claude Opus 5.5 <[email protected]> Assisted-by: Fairy --- libavfilter/vf_bwdif_cuda.cu | 72 +++++++++++++++++++++++++----------- 1 file changed, 50 insertions(+), 22 deletions(-) diff --git a/libavfilter/vf_bwdif_cuda.cu b/libavfilter/vf_bwdif_cuda.cu index c09b6f7f36c..09aede27e2d 100644 --- a/libavfilter/vf_bwdif_cuda.cu +++ b/libavfilter/vf_bwdif_cuda.cu @@ -63,7 +63,7 @@ __inline__ __device__ T filter(T cur_prefs3, T cur_prefs, T cur_mrefs, T cur_mre T prev2_prefs4, T prev2_prefs2, T prev2_0, T prev2_mrefs2, T prev2_mrefs4, T prev_prefs, T prev_mrefs, T next_prefs, T next_mrefs, T next2_prefs4, T next2_prefs2, T next2_0, T next2_mrefs2, T next2_mrefs4, - int clip_max) + int edge, int spat, int clip_max) { T final; @@ -85,10 +85,14 @@ __inline__ __device__ T filter(T cur_prefs3, T cur_prefs, T cur_mrefs, T cur_mre int de = d - e; int mmax = max3(de, dc, min(b, f)); int mmin = min3(de, dc, max(b, f)); - diff = max3(diff, mmin, -mmax); + if (spat) { + diff = max3(diff, mmin, -mmax); + } int interpol; - if (abs(c - e) > temporal_diff0) { + if (edge) { + interpol = (c + e) >> 1; + } else 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_prefs4)) >> 2) @@ -132,10 +136,17 @@ __inline__ __device__ void bwdif_single(T *dst, return; } - T cur_prefs3 = tex2D<T>(cur, xo, yo + 3); - T cur_prefs = tex2D<T>(cur, xo, yo + 1); - T cur_mrefs = tex2D<T>(cur, xo, yo - 1); - T cur_mrefs3 = tex2D<T>(cur, xo, yo - 3); + // As in bwdif, a row past an edge is replaced by the nearest row on the + // other side + int prefs3 = yo + 3 < dst_height ? 3 : -1; + int prefs = yo + 1 < dst_height ? 1 : -1; + int mrefs = yo > 0 ? -1 : 1; + int mrefs3 = yo > 2 ? -3 : 1; + + T cur_prefs3 = tex2D<T>(cur, xo, yo + prefs3); + T cur_prefs = tex2D<T>(cur, xo, yo + prefs); + T cur_mrefs = tex2D<T>(cur, xo, yo + mrefs); + T cur_mrefs3 = tex2D<T>(cur, xo, yo + mrefs3); if (is_field_end) { dst[yo*dst_pitch+xo] = @@ -156,21 +167,26 @@ __inline__ __device__ void bwdif_single(T *dst, T prev2_0 = tex2D<T>(prev2, xo, yo + 0); T prev2_mrefs2 = tex2D<T>(prev2, xo, yo - 2); T prev2_mrefs4 = tex2D<T>(prev2, xo, yo - 4); - T prev_prefs = tex2D<T>(prev1, xo, yo + 1); - T prev_mrefs = tex2D<T>(prev1, xo, yo - 1); - T next_prefs = tex2D<T>(next1, xo, yo + 1); - T next_mrefs = tex2D<T>(next1, xo, yo - 1); + T prev_prefs = tex2D<T>(prev1, xo, yo + prefs); + T prev_mrefs = tex2D<T>(prev1, xo, yo + mrefs); + T next_prefs = tex2D<T>(next1, xo, yo + prefs); + T next_mrefs = tex2D<T>(next1, xo, yo + mrefs); T next2_prefs4 = tex2D<T>(next2, xo, yo + 4); T next2_prefs2 = tex2D<T>(next2, xo, yo + 2); T next2_0 = tex2D<T>(next2, xo, yo + 0); T next2_mrefs2 = tex2D<T>(next2, xo, yo - 2); T next2_mrefs4 = tex2D<T>(next2, xo, yo - 4); + // The 4 rows nearest each edge only interpolate their neighbours, and use + // the rows 2 away only when both are in the picture + int edge = yo < 4 || yo + 5 > dst_height; + int spat = !(yo < 2 || yo + 3 > dst_height); + dst[yo*dst_pitch+xo] = filter(cur_prefs3, cur_prefs, cur_mrefs, cur_mrefs3, prev2_prefs4, prev2_prefs2, prev2_0, prev2_mrefs2, prev2_mrefs4, prev_prefs, prev_mrefs, next_prefs, next_mrefs, next2_prefs4, next2_prefs2, next2_0, next2_mrefs2, next2_mrefs4, - clip_max); + edge, spat, clip_max); } template <typename T> @@ -196,10 +212,17 @@ __inline__ __device__ void bwdif_double(T *dst, return; } - T cur_prefs3 = tex2D<T>(cur, xo, yo + 3); - T cur_prefs = tex2D<T>(cur, xo, yo + 1); - T cur_mrefs = tex2D<T>(cur, xo, yo - 1); - T cur_mrefs3 = tex2D<T>(cur, xo, yo - 3); + // As in bwdif, a row past an edge is replaced by the nearest row on the + // other side + int prefs3 = yo + 3 < dst_height ? 3 : -1; + int prefs = yo + 1 < dst_height ? 1 : -1; + int mrefs = yo > 0 ? -1 : 1; + int mrefs3 = yo > 2 ? -3 : 1; + + T cur_prefs3 = tex2D<T>(cur, xo, yo + prefs3); + T cur_prefs = tex2D<T>(cur, xo, yo + prefs); + T cur_mrefs = tex2D<T>(cur, xo, yo + mrefs); + T cur_mrefs3 = tex2D<T>(cur, xo, yo + mrefs3); if (is_field_end) { T final; @@ -223,27 +246,32 @@ __inline__ __device__ void bwdif_double(T *dst, T prev2_0 = tex2D<T>(prev2, xo, yo + 0); T prev2_mrefs2 = tex2D<T>(prev2, xo, yo - 2); T prev2_mrefs4 = tex2D<T>(prev2, xo, yo - 4); - T prev_prefs = tex2D<T>(prev1, xo, yo + 1); - T prev_mrefs = tex2D<T>(prev1, xo, yo - 1); - T next_prefs = tex2D<T>(next1, xo, yo + 1); - T next_mrefs = tex2D<T>(next1, xo, yo - 1); + T prev_prefs = tex2D<T>(prev1, xo, yo + prefs); + T prev_mrefs = tex2D<T>(prev1, xo, yo + mrefs); + T next_prefs = tex2D<T>(next1, xo, yo + prefs); + T next_mrefs = tex2D<T>(next1, xo, yo + mrefs); T next2_prefs4 = tex2D<T>(next2, xo, yo + 4); T next2_prefs2 = tex2D<T>(next2, xo, yo + 2); T next2_0 = tex2D<T>(next2, xo, yo + 0); T next2_mrefs2 = tex2D<T>(next2, xo, yo - 2); T next2_mrefs4 = tex2D<T>(next2, xo, yo - 4); + // The 4 rows nearest each edge only interpolate their neighbours, and use + // the rows 2 away only when both are in the picture + int edge = yo < 4 || yo + 5 > dst_height; + int spat = !(yo < 2 || yo + 3 > dst_height); + T final; final.x = filter(cur_prefs3.x, cur_prefs.x, cur_mrefs.x, cur_mrefs3.x, prev2_prefs4.x, prev2_prefs2.x, prev2_0.x, prev2_mrefs2.x, prev2_mrefs4.x, prev_prefs.x, prev_mrefs.x, next_prefs.x, next_mrefs.x, next2_prefs4.x, next2_prefs2.x, next2_0.x, next2_mrefs2.x, next2_mrefs4.x, - clip_max); + edge, spat, clip_max); final.y = filter(cur_prefs3.y, cur_prefs.y, cur_mrefs.y, cur_mrefs3.y, prev2_prefs4.y, prev2_prefs2.y, prev2_0.y, prev2_mrefs2.y, prev2_mrefs4.y, prev_prefs.y, prev_mrefs.y, next_prefs.y, next_mrefs.y, next2_prefs4.y, next2_prefs2.y, next2_0.y, next2_mrefs2.y, next2_mrefs4.y, - clip_max); + edge, spat, clip_max); dst[yo*dst_pitch+xo] = final; } -- 2.54.0 From 69dcec3074432bcd95dbf1e831f2e0694eb856f8 Mon Sep 17 00:00:00 2001 From: Michael Niedermayer <[email protected]> Date: Wed, 21 Oct 2026 14:30:03 +0000 Subject: [PATCH 2/2] fate: cover bwdif_cuda picture edges Exercise 8- and 16-bit scalar and vector kernels, both field orders, and both output modes. Generate five frames with varying values in every plane to cover initial intra, temporal, and EOF processing. Use CPU bwdif output for the references. Keep these CUDA-device tests in fate-hw, with an explicit fate-filter-bwdif_cuda target. All 16 cases fail without the edge fix and pass with it. Assisted-by: Fairy --- tests/fate/filter-video.mak | 17 +++++++++++++++++ .../filter-bwdif_cuda-yuv420p-bff-send_field | 15 +++++++++++++++ .../filter-bwdif_cuda-yuv420p-bff-send_frame | 10 ++++++++++ .../filter-bwdif_cuda-yuv420p-tff-send_field | 15 +++++++++++++++ .../filter-bwdif_cuda-yuv420p-tff-send_frame | 10 ++++++++++ ...filter-bwdif_cuda-yuv420p16le-bff-send_field | 15 +++++++++++++++ ...filter-bwdif_cuda-yuv420p16le-bff-send_frame | 10 ++++++++++ ...filter-bwdif_cuda-yuv420p16le-tff-send_field | 15 +++++++++++++++ ...filter-bwdif_cuda-yuv420p16le-tff-send_frame | 10 ++++++++++ ...filter-bwdif_cuda-yuv444p16le-bff-send_field | 15 +++++++++++++++ ...filter-bwdif_cuda-yuv444p16le-bff-send_frame | 10 ++++++++++ ...filter-bwdif_cuda-yuv444p16le-tff-send_field | 15 +++++++++++++++ ...filter-bwdif_cuda-yuv444p16le-tff-send_frame | 10 ++++++++++ 13 files changed, 167 insertions(+) create mode 100644 tests/ref/fate/filter-bwdif_cuda-yuv420p-bff-send_field create mode 100644 tests/ref/fate/filter-bwdif_cuda-yuv420p-bff-send_frame create mode 100644 tests/ref/fate/filter-bwdif_cuda-yuv420p-tff-send_field create mode 100644 tests/ref/fate/filter-bwdif_cuda-yuv420p-tff-send_frame create mode 100644 tests/ref/fate/filter-bwdif_cuda-yuv420p16le-bff-send_field create mode 100644 tests/ref/fate/filter-bwdif_cuda-yuv420p16le-bff-send_frame create mode 100644 tests/ref/fate/filter-bwdif_cuda-yuv420p16le-tff-send_field create mode 100644 tests/ref/fate/filter-bwdif_cuda-yuv420p16le-tff-send_frame create mode 100644 tests/ref/fate/filter-bwdif_cuda-yuv444p16le-bff-send_field create mode 100644 tests/ref/fate/filter-bwdif_cuda-yuv444p16le-bff-send_frame create mode 100644 tests/ref/fate/filter-bwdif_cuda-yuv444p16le-tff-send_field create mode 100644 tests/ref/fate/filter-bwdif_cuda-yuv444p16le-tff-send_frame diff --git a/tests/fate/filter-video.mak b/tests/fate/filter-video.mak index 4a065ec0320..0af62b35b61 100644 --- a/tests/fate/filter-video.mak +++ b/tests/fate/filter-video.mak @@ -1016,6 +1016,23 @@ FATE_FFMPEG += $(FATE_FILTER-yes) FATE-$(CONFIG_AVFILTER) += $(FATE_LIBAVFILTER-yes) +# bwdif_cuda tests require a CUDA device and are only included in fate-hw. +# Both planar and semiplanar formats exercise the scalar and vector kernels. +define BWDIF_CUDA_TEST +FATE_BWDIF_CUDA-$(call FILTERFRAMECRC, COLOR GEQ FORMAT SETFIELD SCALE HWUPLOAD_CUDA BWDIF_CUDA HWDOWNLOAD, FFMPEG) += fate-filter-bwdif_cuda-$(1)-$(2)-$(3) +fate-filter-bwdif_cuda-$(1)-$(2)-$(3): CMD = framecrc -lavfi "color=s=32x20:r=25:d=0.2,format=$(4),geq=lum=mod(1973*X+3911*Y+7127*N\,$(5)):cb=mod(733*X+2053*Y+4101*N\,$(5)):cr=mod(1397*X+3029*Y+5067*N\,$(5)),scale,format=$(1),setfield=$(2),hwupload_cuda,bwdif_cuda=mode=$(3):parity=auto,hwdownload,format=$(1),scale,format=$(4)" +fate-filter-bwdif_cuda-$(1)-$(2)-$(3): REF = $(SRC_PATH)/tests/ref/fate/filter-bwdif_cuda-$(4)-$(2)-$(3) +endef + +$(foreach order,tff bff,$(foreach mode,send_frame send_field,$(eval $(call BWDIF_CUDA_TEST,yuv420p,$(order),$(mode),yuv420p,256)))) +$(foreach order,tff bff,$(foreach mode,send_frame send_field,$(eval $(call BWDIF_CUDA_TEST,nv12,$(order),$(mode),yuv420p,256)))) +$(foreach order,tff bff,$(foreach mode,send_frame send_field,$(eval $(call BWDIF_CUDA_TEST,yuv444p16le,$(order),$(mode),yuv444p16le,65536)))) +$(foreach order,tff bff,$(foreach mode,send_frame send_field,$(eval $(call BWDIF_CUDA_TEST,p016le,$(order),$(mode),yuv420p16le,65536)))) + +FATE_HW-yes += $(FATE_BWDIF_CUDA-yes) +$(FATE_BWDIF_CUDA-yes): $(FFMPEG) +fate-filter-bwdif_cuda: $(FATE_BWDIF_CUDA-yes) + # scale_cuda filter tests (YUV<->RGB color space conversion) # These require a CUDA capable GPU, so they are only run as part of fate-hw # or when requested directly with make fate-filter-scale_cuda. diff --git a/tests/ref/fate/filter-bwdif_cuda-yuv420p-bff-send_field b/tests/ref/fate/filter-bwdif_cuda-yuv420p-bff-send_field new file mode 100644 index 00000000000..2d97f47b486 --- /dev/null +++ b/tests/ref/fate/filter-bwdif_cuda-yuv420p-bff-send_field @@ -0,0 +1,15 @@ +#tb 0: 1/50 +#media_type 0: video +#codec_id 0: rawvideo +#dimensions 0: 32x20 +#sar 0: 1/1 +0, 0, 0, 1, 960, 0x65d9d98a +0, 1, 1, 1, 960, 0x3dc2d496 +0, 2, 2, 1, 960, 0x672ad580 +0, 3, 3, 1, 960, 0x63dfde0b +0, 4, 4, 1, 960, 0xf417d92e +0, 5, 5, 1, 960, 0xb9a4dd69 +0, 6, 6, 1, 960, 0xbc92d81d +0, 7, 7, 1, 960, 0x8f81dc5e +0, 8, 8, 1, 960, 0xdd8ddbf2 +0, 9, 9, 1, 960, 0xb33dd4a3 diff --git a/tests/ref/fate/filter-bwdif_cuda-yuv420p-bff-send_frame b/tests/ref/fate/filter-bwdif_cuda-yuv420p-bff-send_frame new file mode 100644 index 00000000000..5e446fc0c68 --- /dev/null +++ b/tests/ref/fate/filter-bwdif_cuda-yuv420p-bff-send_frame @@ -0,0 +1,10 @@ +#tb 0: 1/25 +#media_type 0: video +#codec_id 0: rawvideo +#dimensions 0: 32x20 +#sar 0: 1/1 +0, 0, 0, 1, 960, 0x65d9d98a +0, 1, 1, 1, 960, 0x672ad580 +0, 2, 2, 1, 960, 0xf417d92e +0, 3, 3, 1, 960, 0xbc92d81d +0, 4, 4, 1, 960, 0xdd8ddbf2 diff --git a/tests/ref/fate/filter-bwdif_cuda-yuv420p-tff-send_field b/tests/ref/fate/filter-bwdif_cuda-yuv420p-tff-send_field new file mode 100644 index 00000000000..539488627ea --- /dev/null +++ b/tests/ref/fate/filter-bwdif_cuda-yuv420p-tff-send_field @@ -0,0 +1,15 @@ +#tb 0: 1/50 +#media_type 0: video +#codec_id 0: rawvideo +#dimensions 0: 32x20 +#sar 0: 1/1 +0, 0, 0, 1, 960, 0x1b90d9c7 +0, 1, 1, 1, 960, 0x28ead634 +0, 2, 2, 1, 960, 0x55ded917 +0, 3, 3, 1, 960, 0x6591da2f +0, 4, 4, 1, 960, 0x50efd921 +0, 5, 5, 1, 960, 0x7148dd0b +0, 6, 6, 1, 960, 0x94fcd899 +0, 7, 7, 1, 960, 0x0576db79 +0, 8, 8, 1, 960, 0xce36d92c +0, 9, 9, 1, 960, 0xc722d7b3 diff --git a/tests/ref/fate/filter-bwdif_cuda-yuv420p-tff-send_frame b/tests/ref/fate/filter-bwdif_cuda-yuv420p-tff-send_frame new file mode 100644 index 00000000000..c6b71651e48 --- /dev/null +++ b/tests/ref/fate/filter-bwdif_cuda-yuv420p-tff-send_frame @@ -0,0 +1,10 @@ +#tb 0: 1/25 +#media_type 0: video +#codec_id 0: rawvideo +#dimensions 0: 32x20 +#sar 0: 1/1 +0, 0, 0, 1, 960, 0x1b90d9c7 +0, 1, 1, 1, 960, 0x55ded917 +0, 2, 2, 1, 960, 0x50efd921 +0, 3, 3, 1, 960, 0x94fcd899 +0, 4, 4, 1, 960, 0xce36d92c diff --git a/tests/ref/fate/filter-bwdif_cuda-yuv420p16le-bff-send_field b/tests/ref/fate/filter-bwdif_cuda-yuv420p16le-bff-send_field new file mode 100644 index 00000000000..8dba160e2e1 --- /dev/null +++ b/tests/ref/fate/filter-bwdif_cuda-yuv420p16le-bff-send_field @@ -0,0 +1,15 @@ +#tb 0: 1/50 +#media_type 0: video +#codec_id 0: rawvideo +#dimensions 0: 32x20 +#sar 0: 1/1 +0, 0, 0, 1, 1920, 0xd2fc7cca +0, 1, 1, 1, 1920, 0x4e5b73e0 +0, 2, 2, 1, 1920, 0xf7148b3e +0, 3, 3, 1, 1920, 0xd67f8a4b +0, 4, 4, 1, 1920, 0x3cd5ae46 +0, 5, 5, 1, 1920, 0x5614a907 +0, 6, 6, 1, 1920, 0x3164bdbe +0, 7, 7, 1, 1920, 0x99bcb8e3 +0, 8, 8, 1, 1920, 0x6471c84b +0, 9, 9, 1, 1920, 0x455fcb56 diff --git a/tests/ref/fate/filter-bwdif_cuda-yuv420p16le-bff-send_frame b/tests/ref/fate/filter-bwdif_cuda-yuv420p16le-bff-send_frame new file mode 100644 index 00000000000..342a3e30ad9 --- /dev/null +++ b/tests/ref/fate/filter-bwdif_cuda-yuv420p16le-bff-send_frame @@ -0,0 +1,10 @@ +#tb 0: 1/25 +#media_type 0: video +#codec_id 0: rawvideo +#dimensions 0: 32x20 +#sar 0: 1/1 +0, 0, 0, 1, 1920, 0xd2fc7cca +0, 1, 1, 1, 1920, 0xf7148b3e +0, 2, 2, 1, 1920, 0x3cd5ae46 +0, 3, 3, 1, 1920, 0x3164bdbe +0, 4, 4, 1, 1920, 0x6471c84b diff --git a/tests/ref/fate/filter-bwdif_cuda-yuv420p16le-tff-send_field b/tests/ref/fate/filter-bwdif_cuda-yuv420p16le-tff-send_field new file mode 100644 index 00000000000..3d61a559aae --- /dev/null +++ b/tests/ref/fate/filter-bwdif_cuda-yuv420p16le-tff-send_field @@ -0,0 +1,15 @@ +#tb 0: 1/50 +#media_type 0: video +#codec_id 0: rawvideo +#dimensions 0: 32x20 +#sar 0: 1/1 +0, 0, 0, 1, 1920, 0xd82578df +0, 1, 1, 1, 1920, 0x5e6f7afc +0, 2, 2, 1, 1920, 0xaa7b8de1 +0, 3, 3, 1, 1920, 0xd3ba8711 +0, 4, 4, 1, 1920, 0xd665a77d +0, 5, 5, 1, 1920, 0x2fe8a9c6 +0, 6, 6, 1, 1920, 0x3b8ab69d +0, 7, 7, 1, 1920, 0x7190b481 +0, 8, 8, 1, 1920, 0xc546c4ba +0, 9, 9, 1, 1920, 0xc062ccb0 diff --git a/tests/ref/fate/filter-bwdif_cuda-yuv420p16le-tff-send_frame b/tests/ref/fate/filter-bwdif_cuda-yuv420p16le-tff-send_frame new file mode 100644 index 00000000000..ebc82b55ff9 --- /dev/null +++ b/tests/ref/fate/filter-bwdif_cuda-yuv420p16le-tff-send_frame @@ -0,0 +1,10 @@ +#tb 0: 1/25 +#media_type 0: video +#codec_id 0: rawvideo +#dimensions 0: 32x20 +#sar 0: 1/1 +0, 0, 0, 1, 1920, 0xd82578df +0, 1, 1, 1, 1920, 0xaa7b8de1 +0, 2, 2, 1, 1920, 0xd665a77d +0, 3, 3, 1, 1920, 0x3b8ab69d +0, 4, 4, 1, 1920, 0xc546c4ba diff --git a/tests/ref/fate/filter-bwdif_cuda-yuv444p16le-bff-send_field b/tests/ref/fate/filter-bwdif_cuda-yuv444p16le-bff-send_field new file mode 100644 index 00000000000..9308d16884f --- /dev/null +++ b/tests/ref/fate/filter-bwdif_cuda-yuv444p16le-bff-send_field @@ -0,0 +1,15 @@ +#tb 0: 1/50 +#media_type 0: video +#codec_id 0: rawvideo +#dimensions 0: 32x20 +#sar 0: 1/1 +0, 0, 0, 1, 3840, 0x11fc70ed +0, 1, 1, 1, 3840, 0x358a655e +0, 2, 2, 1, 3840, 0xb64b92f6 +0, 3, 3, 1, 3840, 0x6ca08bef +0, 4, 4, 1, 3840, 0xd171bc91 +0, 5, 5, 1, 3840, 0xc914ad4d +0, 6, 6, 1, 3840, 0x7d39bdd2 +0, 7, 7, 1, 3840, 0x8ccab70b +0, 8, 8, 1, 3840, 0x1fbeb4d5 +0, 9, 9, 1, 3840, 0x8b19bad6 diff --git a/tests/ref/fate/filter-bwdif_cuda-yuv444p16le-bff-send_frame b/tests/ref/fate/filter-bwdif_cuda-yuv444p16le-bff-send_frame new file mode 100644 index 00000000000..98327e4fefc --- /dev/null +++ b/tests/ref/fate/filter-bwdif_cuda-yuv444p16le-bff-send_frame @@ -0,0 +1,10 @@ +#tb 0: 1/25 +#media_type 0: video +#codec_id 0: rawvideo +#dimensions 0: 32x20 +#sar 0: 1/1 +0, 0, 0, 1, 3840, 0x11fc70ed +0, 1, 1, 1, 3840, 0xb64b92f6 +0, 2, 2, 1, 3840, 0xd171bc91 +0, 3, 3, 1, 3840, 0x7d39bdd2 +0, 4, 4, 1, 3840, 0x1fbeb4d5 diff --git a/tests/ref/fate/filter-bwdif_cuda-yuv444p16le-tff-send_field b/tests/ref/fate/filter-bwdif_cuda-yuv444p16le-tff-send_field new file mode 100644 index 00000000000..72986056ed5 --- /dev/null +++ b/tests/ref/fate/filter-bwdif_cuda-yuv444p16le-tff-send_field @@ -0,0 +1,15 @@ +#tb 0: 1/50 +#media_type 0: video +#codec_id 0: rawvideo +#dimensions 0: 32x20 +#sar 0: 1/1 +0, 0, 0, 1, 3840, 0xb4a8679d +0, 1, 1, 1, 3840, 0x830a6c42 +0, 2, 2, 1, 3840, 0x040e945e +0, 3, 3, 1, 3840, 0x015e8be4 +0, 4, 4, 1, 3840, 0xf247b03e +0, 5, 5, 1, 3840, 0xee1dad8c +0, 6, 6, 1, 3840, 0x1dd0b949 +0, 7, 7, 1, 3840, 0x0ea9abc1 +0, 8, 8, 1, 3840, 0xef9daf2e +0, 9, 9, 1, 3840, 0x7d57bbad diff --git a/tests/ref/fate/filter-bwdif_cuda-yuv444p16le-tff-send_frame b/tests/ref/fate/filter-bwdif_cuda-yuv444p16le-tff-send_frame new file mode 100644 index 00000000000..98a6eebe0ad --- /dev/null +++ b/tests/ref/fate/filter-bwdif_cuda-yuv444p16le-tff-send_frame @@ -0,0 +1,10 @@ +#tb 0: 1/25 +#media_type 0: video +#codec_id 0: rawvideo +#dimensions 0: 32x20 +#sar 0: 1/1 +0, 0, 0, 1, 3840, 0xb4a8679d +0, 1, 1, 1, 3840, 0x040e945e +0, 2, 2, 1, 3840, 0xf247b03e +0, 3, 3, 1, 3840, 0x1dd0b949 +0, 4, 4, 1, 3840, 0xef9daf2e -- 2.54.0 _______________________________________________ ffmpeg-devel mailing list -- [email protected] To unsubscribe send an email to [email protected]
