From 2c86da9d65f3fff86d5e110110b13a49d7bd415c Mon Sep 17 00:00:00 2001 From: nyanmisaka Date: Tue, 30 Jun 2026 00:58:29 +0800 Subject: [PATCH] Fix typos in some shader based deint filters Signed-off-by: nyanmisaka --- .../0056-add-bwdif-videotoolbox-filter.patch | 54 +++++-- ...d-yadif-and-bwdif-opencl-filter-impl.patch | 28 ++-- ...ckport-fixes-for-cuda-from-upstream.patch} | 135 +++++++++++++++++- debian/patches/series | 2 +- 4 files changed, 190 insertions(+), 29 deletions(-) rename debian/patches/{0083-backport-fixes-for-nvenc-and-cuvid-from-upstream.patch => 0083-backport-fixes-for-cuda-from-upstream.patch} (55%) diff --git a/debian/patches/0056-add-bwdif-videotoolbox-filter.patch b/debian/patches/0056-add-bwdif-videotoolbox-filter.patch index 6c3e4eb50..53eb01578 100644 --- a/debian/patches/0056-add-bwdif-videotoolbox-filter.patch +++ b/debian/patches/0056-add-bwdif-videotoolbox-filter.patch @@ -2,7 +2,7 @@ Index: FFmpeg/configure =================================================================== --- FFmpeg.orig/configure +++ FFmpeg/configure -@@ -4120,6 +4120,7 @@ boxblur_opencl_filter_deps="opencl gpl" +@@ -4123,6 +4123,7 @@ boxblur_opencl_filter_deps="opencl gpl" bs2b_filter_deps="libbs2b" bwdif_cuda_filter_deps="ffnvcodec" bwdif_cuda_filter_deps_any="cuda_nvcc cuda_llvm" @@ -155,7 +155,7 @@ Index: FFmpeg/libavfilter/metal/vf_bwdif_videotoolbox.metal + 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)) / 4) ++ + coef_hf[2] * (prev2_mrefs4 + next2_mrefs4 + prev2_prefs4 + next2_prefs4)) / 4) + + coef_lf[0] * (c + e) - coef_lf[1] * (cur_mrefs3 + cur_prefs3)) / (1 << 13); + } else { + interpol = (coef_sp[0] * (c + e) - coef_sp[1] * (cur_mrefs3 + cur_prefs3)) / (1 << 13); @@ -196,10 +196,10 @@ Index: FFmpeg/libavfilter/metal/vf_bwdif_videotoolbox.metal + } + + // Calculate temporal prediction -+ texture2d prev2 = prev; -+ texture2d prev1 = is_second_field ? cur : prev; -+ texture2d next1 = is_second_field ? next : cur; -+ texture2d next2 = next; ++ texture2d prev2 = is_second_field ? cur : prev; ++ texture2d prev1 = prev; ++ texture2d next1 = next; ++ texture2d next2 = is_second_field ? next : cur; + + T prev2_prefs4 = tex2D(prev2, pos.x, pos.y+ 4); + T prev2_prefs2 = tex2D(prev2, pos.x, pos.y + 2); @@ -249,10 +249,10 @@ Index: FFmpeg/libavfilter/metal/vf_bwdif_videotoolbox.metal + } + + // Calculate temporal prediction -+ texture2d prev2 = prev; -+ texture2d prev1 = is_second_field ? cur : prev; -+ texture2d next1 = is_second_field ? next : cur; -+ texture2d next2 = next; ++ texture2d prev2 = is_second_field ? cur : prev; ++ texture2d prev1 = prev; ++ texture2d next1 = next; ++ texture2d next2 = is_second_field ? next : cur; + + T prev2_prefs4 = tex2D(prev2, pos.x, pos.y+ 4); + T prev2_prefs2 = tex2D(prev2, pos.x, pos.y + 2); @@ -312,6 +312,40 @@ Index: FFmpeg/libavfilter/metal/vf_bwdif_videotoolbox.metal + } + dst.write(pred.xyyy, pos); +} +Index: FFmpeg/libavfilter/metal/vf_yadif_videotoolbox.metal +=================================================================== +--- FFmpeg.orig/libavfilter/metal/vf_yadif_videotoolbox.metal ++++ FFmpeg/libavfilter/metal/vf_yadif_videotoolbox.metal +@@ -88,7 +88,7 @@ T spatial_predictor(T a, T b, T c, T d, + T h, T i, T j, T k, T l, T m, T n) + { + T spatial_pred = (d + k)/2; +- T spatial_score = abs(c - j) + abs(d - k) + abs(e - l); ++ T spatial_score = abs(c - j) + abs(d - k) + abs(e - l) - (1.0f / 255.0f); + + T score = abs(b - k) + abs(c - l) + abs(d - m); + if (score < spatial_score) { +@@ -128,7 +128,7 @@ T temporal_predictor(T A, T B, T C, T D, + T tdiff1 = (abs(A - F) + abs(B - G)) / 2; + T tdiff2 = (abs(K - F) + abs(G - L)) / 2; + +- T diff = max3(tdiff0, tdiff1, tdiff2); ++ T diff = max3(tdiff0 / 2, tdiff1, tdiff2); + + if (!skip_check) { + T maxi = max3(p2 - p3, p2 - p1, min(p0 - p1, p4 - p3)); +@@ -239,9 +239,9 @@ T yadif( + T spatial_pred = yadif_compute_spatial(cur, pos); + + if (params.is_second_field) { +- return yadif_compute_temporal(cur, prev, cur, next, next, spatial_pred, params.skip_spatial_check, pos); ++ return yadif_compute_temporal(cur, cur, prev, next, next, spatial_pred, params.skip_spatial_check, pos); + } else { +- return yadif_compute_temporal(cur, prev, prev, cur, next, spatial_pred, params.skip_spatial_check, pos); ++ return yadif_compute_temporal(cur, prev, prev, next, cur, spatial_pred, params.skip_spatial_check, pos); + } + } + Index: FFmpeg/libavfilter/vf_bwdif_videotoolbox.m =================================================================== --- /dev/null diff --git a/debian/patches/0069-add-yadif-and-bwdif-opencl-filter-impl.patch b/debian/patches/0069-add-yadif-and-bwdif-opencl-filter-impl.patch index d5c1b13f6..c623c5858 100644 --- a/debian/patches/0069-add-yadif-and-bwdif-opencl-filter-impl.patch +++ b/debian/patches/0069-add-yadif-and-bwdif-opencl-filter-impl.patch @@ -2,7 +2,7 @@ Index: FFmpeg/configure =================================================================== --- FFmpeg.orig/configure +++ FFmpeg/configure -@@ -4121,6 +4121,7 @@ boxblur_opencl_filter_deps="opencl gpl" +@@ -4124,6 +4124,7 @@ boxblur_opencl_filter_deps="opencl gpl" bs2b_filter_deps="libbs2b" bwdif_cuda_filter_deps="ffnvcodec" bwdif_cuda_filter_deps_any="cuda_nvcc cuda_llvm" @@ -10,7 +10,7 @@ Index: FFmpeg/configure bwdif_videotoolbox_filter_deps="metal corevideo videotoolbox" bwdif_vulkan_filter_deps="vulkan spirv_compiler" chromaber_vulkan_filter_deps="vulkan spirv_library" -@@ -4280,6 +4281,7 @@ xfade_opencl_filter_deps="opencl" +@@ -4283,6 +4284,7 @@ xfade_opencl_filter_deps="opencl" xfade_vulkan_filter_deps="vulkan spirv_library" yadif_cuda_filter_deps="ffnvcodec" yadif_cuda_filter_deps_any="cuda_nvcc cuda_llvm" @@ -52,7 +52,7 @@ Index: FFmpeg/libavfilter/allfilters.c extern const FFFilter ff_vf_bwdif_videotoolbox; extern const FFFilter ff_vf_bwdif_vulkan; extern const FFFilter ff_vf_cas; -@@ -562,6 +563,7 @@ extern const FFFilter ff_vf_xpsnr; +@@ -563,6 +564,7 @@ extern const FFFilter ff_vf_xpsnr; extern const FFFilter ff_vf_xstack; extern const FFFilter ff_vf_yadif; extern const FFFilter ff_vf_yadif_cuda; @@ -228,7 +228,7 @@ Index: FFmpeg/libavfilter/opencl/bwdif.cl + if (fabs(c - e) > temporal_diff0) { + interpol = native_divide((((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)) * 0.25f) ++ + coef_hf[2] * (prev2_mrefs4 + next2_mrefs4 + prev2_prefs4 + next2_prefs4)) * 0.25f) + + coef_lf[0] * (c + e) - coef_lf[1] * (cur_mrefs3 + cur_prefs3)), (float)(1 << 13)); + } else { + interpol = native_divide((coef_sp[0] * (c + e) - coef_sp[1] * (cur_mrefs3 + cur_prefs3)), (float)(1 << 13)); @@ -322,17 +322,17 @@ Index: FFmpeg/libavfilter/opencl/bwdif.cl + + if (channels == 1) { + float pred = is_second_field -+ ? bwdif_compute_float(dst, cur, prev, cur, next, next, ++ ? bwdif_compute_float(dst, cur, cur, prev, next, next, + parity, is_field_end, pos) -+ : bwdif_compute_float(dst, cur, prev, prev, cur, next, ++ : bwdif_compute_float(dst, cur, prev, prev, next, cur, + parity, is_field_end, pos); + + write_imagef(dst, pos, (float4)(pred, 0.0f, 0.0f, 1.0f)); + } else if (channels == 2) { + float2 pred = is_second_field -+ ? bwdif_compute_float2(dst, cur, prev, cur, next, next, ++ ? bwdif_compute_float2(dst, cur, cur, prev, next, next, + parity, is_field_end, pos) -+ : bwdif_compute_float2(dst, cur, prev, prev, cur, next, ++ : bwdif_compute_float2(dst, cur, prev, prev, next, cur, + parity, is_field_end, pos); + + write_imagef(dst, pos, (float4)(pred.x, pred.y, 0.0f, 1.0f)); @@ -375,7 +375,7 @@ Index: FFmpeg/libavfilter/opencl/yadif.cl + float h, float i, float j, float k, float l, float m, float n) +{ + float spatial_pred = (d + k) * 0.5f; -+ float spatial_score = fabs(c - j) + fabs(d - k) + fabs(e - l); ++ float spatial_score = fabs(c - j) + fabs(d - k) + fabs(e - l) - (1.0f / 255.0f); + + float score = fabs(b - k) + fabs(c - l) + fabs(d - m); + if (score < spatial_score) { @@ -424,7 +424,7 @@ Index: FFmpeg/libavfilter/opencl/yadif.cl + float tdiff1 = (fabs(A - F) + fabs(B - G)) * 0.5f; + float tdiff2 = (fabs(K - F) + fabs(G - L)) * 0.5f; + -+ float diff = max3(tdiff0, tdiff1, tdiff2); ++ float diff = max3(tdiff0 * 0.5f, tdiff1, tdiff2); + + if (!skip_check) { + float maxi = max3(p2 - p3, p2 - p1, min(p0 - p1, p4 - p3)); @@ -523,18 +523,18 @@ Index: FFmpeg/libavfilter/opencl/yadif.cl + if (channels == 1) { + float spatial_pred = yadif_compute_spatial_float(cur, pos); + float pred = is_second_field -+ ? yadif_compute_temporal_float(cur, prev, cur, next, next, ++ ? yadif_compute_temporal_float(cur, cur, prev, next, next, + spatial_pred, skip_spatial_check, pos) -+ : yadif_compute_temporal_float(cur, prev, prev, cur, next, ++ : yadif_compute_temporal_float(cur, prev, prev, next, cur, + spatial_pred, skip_spatial_check, pos); + + write_imagef(dst, pos, (float4)(pred, 0.0f, 0.0f, 1.0f)); + } else if (channels == 2) { + float2 spatial_pred = yadif_compute_spatial_float2(cur, pos); + float2 pred = is_second_field -+ ? yadif_compute_temporal_float2(cur, prev, cur, next, next, ++ ? yadif_compute_temporal_float2(cur, cur, prev, next, next, + spatial_pred, skip_spatial_check, pos) -+ : yadif_compute_temporal_float2(cur, prev, prev, cur, next, ++ : yadif_compute_temporal_float2(cur, prev, prev, next, cur, + spatial_pred, skip_spatial_check, pos); + + write_imagef(dst, pos, (float4)(pred.x, pred.y, 0.0f, 1.0f)); diff --git a/debian/patches/0083-backport-fixes-for-nvenc-and-cuvid-from-upstream.patch b/debian/patches/0083-backport-fixes-for-cuda-from-upstream.patch similarity index 55% rename from debian/patches/0083-backport-fixes-for-nvenc-and-cuvid-from-upstream.patch rename to debian/patches/0083-backport-fixes-for-cuda-from-upstream.patch index d40e3cf21..1599dd6d1 100644 --- a/debian/patches/0083-backport-fixes-for-nvenc-and-cuvid-from-upstream.patch +++ b/debian/patches/0083-backport-fixes-for-cuda-from-upstream.patch @@ -83,7 +83,7 @@ Index: FFmpeg/libavcodec/nvenc.c =================================================================== --- FFmpeg.orig/libavcodec/nvenc.c +++ FFmpeg/libavcodec/nvenc.c -@@ -592,12 +594,24 @@ static int nvenc_check_capabilities(AVCo +@@ -594,12 +594,24 @@ static int nvenc_check_capabilities(AVCo #ifdef NVENC_HAVE_BFRAME_REF_MODE tmp = (ctx->b_ref_mode >= 0) ? ctx->b_ref_mode : NV_ENC_BFRAME_REF_MODE_DISABLED; ret = nvenc_check_cap(avctx, NV_ENC_CAPS_SUPPORT_BFRAME_REF_MODE); @@ -114,7 +114,7 @@ Index: FFmpeg/libavcodec/nvenc.c } #else tmp = (ctx->b_ref_mode >= 0) ? ctx->b_ref_mode : 0; -@@ -683,7 +697,9 @@ static int nvenc_check_capabilities(AVCo +@@ -685,7 +697,9 @@ static int nvenc_check_capabilities(AVCo #ifdef NVENC_HAVE_MVHEVC ctx->multiview_supported = nvenc_check_cap(avctx, NV_ENC_CAPS_SUPPORT_MVHEVC_ENCODE) > 0; @@ -125,7 +125,7 @@ Index: FFmpeg/libavcodec/nvenc.c av_log(avctx, AV_LOG_WARNING, "Multiview not supported by the device\n"); return AVERROR(ENOSYS); } -@@ -1163,6 +1179,13 @@ static av_cold int nvenc_setup_rate_cont +@@ -1165,6 +1179,13 @@ static av_cold int nvenc_setup_rate_cont #ifdef NVENC_HAVE_QP_CHROMA_OFFSETS ctx->encode_config.rcParams.cbQPIndexOffset = ctx->qp_cb_offset; ctx->encode_config.rcParams.crQPIndexOffset = ctx->qp_cr_offset; @@ -139,7 +139,7 @@ Index: FFmpeg/libavcodec/nvenc.c #else if (ctx->qp_cb_offset || ctx->qp_cr_offset) av_log(avctx, AV_LOG_WARNING, "Failed setting QP CB/CR offsets, SDK 11.1 or greater required at compile time.\n"); -@@ -1359,6 +1382,11 @@ static av_cold int nvenc_setup_h264_conf +@@ -1361,6 +1382,11 @@ static av_cold int nvenc_setup_h264_conf case NV_ENC_H264_PROFILE_BASELINE: cc->profileGUID = NV_ENC_H264_PROFILE_BASELINE_GUID; avctx->profile = AV_PROFILE_H264_BASELINE; @@ -151,6 +151,49 @@ Index: FFmpeg/libavcodec/nvenc.c break; case NV_ENC_H264_PROFILE_MAIN: cc->profileGUID = NV_ENC_H264_PROFILE_MAIN_GUID; +Index: FFmpeg/libavfilter/vf_bwdif_cuda.cu +=================================================================== +--- FFmpeg.orig/libavfilter/vf_bwdif_cuda.cu ++++ FFmpeg/libavfilter/vf_bwdif_cuda.cu +@@ -91,7 +91,7 @@ __inline__ __device__ T filter(T cur_pre + 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( + // 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(prev2, xo, yo + 4); + T prev2_prefs2 = tex2D(prev2, xo, yo + 2); +@@ -213,10 +213,10 @@ __inline__ __device__ void bwdif_double( + + 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(prev2, xo, yo + 4); + T prev2_prefs2 = tex2D(prev2, xo, yo + 2); Index: FFmpeg/libavfilter/vf_scale_cuda.h =================================================================== --- FFmpeg.orig/libavfilter/vf_scale_cuda.h @@ -166,3 +209,87 @@ Index: FFmpeg/libavfilter/vf_scale_cuda.h #else #include #endif +Index: FFmpeg/libavfilter/vf_yadif_cuda.cu +=================================================================== +--- FFmpeg.orig/libavfilter/vf_yadif_cuda.cu ++++ FFmpeg/libavfilter/vf_yadif_cuda.cu +@@ -23,7 +23,7 @@ __inline__ __device__ T spatial_predicto + T h, T i, T j, T k, T l, T m, T n) + { + int spatial_pred = (d + k)/2; +- int spatial_score = abs(c - j) + abs(d - k) + abs(e - l); ++ int spatial_score = abs(c - j) + abs(d - k) + abs(e - l) - 1; + + int score = abs(b - k) + abs(c - l) + abs(d - m); + if (score < spatial_score) { +@@ -75,7 +75,7 @@ __inline__ __device__ T temporal_predict + int tdiff1 = (abs(A - F) + abs(B - G)) / 2; + int tdiff2 = (abs(K - F) + abs(G - L)) / 2; + +- int diff = max3(tdiff0, tdiff1, tdiff2); ++ int diff = max3(tdiff0 >> 1, tdiff1, tdiff2); + + if (!skip_check) { + int maxi = max3(p2 - p3, p2 - p1, min(p0 - p1, p4 - p3)); +@@ -85,8 +85,7 @@ __inline__ __device__ T temporal_predict + + if (spatial_pred > p2 + diff) { + spatial_pred = p2 + diff; +- } +- if (spatial_pred < p2 - diff) { ++ } else if (spatial_pred < p2 - diff) { + spatial_pred = p2 - diff; + } + +@@ -139,10 +138,10 @@ __inline__ __device__ void yadif_single( + // 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 A = tex2D(prev2, xo, yo - 1); + T B = tex2D(prev2, xo, yo + 1); +@@ -210,10 +209,10 @@ __inline__ __device__ void yadif_double( + // 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 A = tex2D(prev2, xo, yo - 1); + T B = tex2D(prev2, xo, yo + 1); +Index: FFmpeg/libavutil/hwcontext_cuda.c +=================================================================== +--- FFmpeg.orig/libavutil/hwcontext_cuda.c ++++ FFmpeg/libavutil/hwcontext_cuda.c +@@ -207,7 +207,7 @@ static int cuda_get_buffer(AVHWFramesCon + if (ctx->sw_format == AV_PIX_FMT_YUV420P) { + frame->linesize[1] = frame->linesize[2] = frame->linesize[0] / 2; + frame->data[2] = frame->data[1]; +- frame->data[1] = frame->data[2] + frame->linesize[2] * (ctx->height / 2); ++ frame->data[1] = frame->data[2] + frame->linesize[2] * AV_CEIL_RSHIFT(ctx->height, 1); + } + + frame->format = AV_PIX_FMT_CUDA; +@@ -259,7 +259,7 @@ static int cuda_transfer_data(AVHWFrames + .srcPitch = src->linesize[i], + .dstPitch = dst->linesize[i], + .WidthInBytes = FFMIN(src->linesize[i], dst->linesize[i]), +- .Height = src->height >> ((i == 0 || i == 3) ? 0 : priv->shift_height), ++ .Height = AV_CEIL_RSHIFT(src->height, ((i == 0 || i == 3) ? 0 : priv->shift_height)), + }; + + if (src->hw_frames_ctx) { diff --git a/debian/patches/series b/debian/patches/series index 97f7763a9..448f6ffd4 100644 --- a/debian/patches/series +++ b/debian/patches/series @@ -80,7 +80,7 @@ 0080-prefer-vulkan-device-with-higher-api-version.patch 0081-signal-vulkan-semaphores-in-drm-prime-reverse-mapping.patch 0082-add-fixes-and-enhancements-for-vulkan-filters.patch -0083-backport-fixes-for-nvenc-and-cuvid-from-upstream.patch +0083-backport-fixes-for-cuda-from-upstream.patch 0084-backport-fixes-for-detecting-ac4-streams-in-mpegts.patch 0085-fix-hdr10-sidedata-not-removed-in-libplacebo-outlink.patch 0086-workaround-amd-hevc-vaapi-encoder-out-of-band-vtag-hvc1.patch