jellyfin/jellyfin-ffmpeg
1
0
Fork 0
mirror of https://github.com/jellyfin/jellyfin-ffmpeg.git synced 2026-09-03 05:09:58 +03:00
Code Issues Packages Projects Releases Wiki Activity

Fix typos in some shader based deint filters

Browse Source 
Signed-off-by: nyanmisaka <nst799610810@gmail.com>
This commit is contained in:
nyanmisaka
2026-06-30 01:00:05 +08:00
parent 5cb8e04fae
commit 2c86da9d65
4 changed files with 190 additions and 29 deletions
Download Patch File Download Diff File Expand all files Collapse all files
debian/patches/0056-add-bwdif-videotoolbox-filter.patch
Vendored
+44 -10
View File
@@ -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<float, accesstype> prev2 = prev;
+ texture2d<float, accesstype> prev1 = is_second_field ? cur : prev;
+ texture2d<float, accesstype> next1 = is_second_field ? next : cur;
+ texture2d<float, accesstype> next2 = next;
+ texture2d<float, accesstype> prev2 = is_second_field ? cur : prev;
+ texture2d<float, accesstype> prev1 = prev;
+ texture2d<float, accesstype> next1 = next;
+ texture2d<float, accesstype> next2 = is_second_field ? next : cur;
+
+ T prev2_prefs4 = tex2D<T>(prev2, pos.x, pos.y+ 4);
+ T prev2_prefs2 = tex2D<T>(prev2, pos.x, pos.y + 2);
@@ -249,10 +249,10 @@ Index: FFmpeg/libavfilter/metal/vf_bwdif_videotoolbox.metal
+ }
+
+ // Calculate temporal prediction
+ texture2d<float, accesstype> prev2 = prev;
+ texture2d<float, accesstype> prev1 = is_second_field ? cur : prev;
+ texture2d<float, accesstype> next1 = is_second_field ? next : cur;
+ texture2d<float, accesstype> next2 = next;
+ texture2d<float, accesstype> prev2 = is_second_field ? cur : prev;
+ texture2d<float, accesstype> prev1 = prev;
+ texture2d<float, accesstype> next1 = next;
+ texture2d<float, accesstype> next2 = is_second_field ? next : cur;
+
+ T prev2_prefs4 = tex2D<T>(prev2, pos.x, pos.y+ 4);
+ T prev2_prefs2 = tex2D<T>(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<T>(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
debian/patches/0069-add-yadif-and-bwdif-opencl-filter-impl.patch
Vendored
+14 -14
View File
@@ -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));
debian/patches/0083-backport-fixes-for-nvenc-and-cuvid-from-upstream.patch → debian/patches/0083-backport-fixes-for-cuda-from-upstream.patch
Vendored
+131 -4
View File
@@ -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<T>(prev2, xo, yo + 4);
T prev2_prefs2 = tex2D<T>(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<T>(prev2, xo, yo + 4);
T prev2_prefs2 = tex2D<T>(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 <ffnvcodec/dynlink_cuda.h>
#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<T>(prev2, xo, yo - 1);
T B = tex2D<T>(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<T>(prev2, xo, yo - 1);
T B = tex2D<T>(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) {
debian/patches/series
Vendored
+1 -1
View File
@@ -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