June 2026 - ffmpeg-devel - lists.ffmpeg.org

[PR] avfilter/vf_bwdif_cuda: fix typos in algorithm (PR #23619)
by philipl 27 Jun '26

27 Jun '26

PR #23619 opened by philipl URL: https://code.ffmpeg.org/FFmpeg/FFmpeg/pulls/23619 Patch URL: https://code.ffmpeg.org/FFmpeg/FFmpeg/pulls/23619.patch 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. # Summary of changes Briefly describe what this PR does and why.  From d33873dd5194f79b9287e32b9520431fac1a8311 Mon Sep 17 00:00:00 2001 From: Philip Langdale <philipl(a)overt.org> Date: Fri, 26 Jun 2026 22:56:34 -0700 Subject: [PATCH] 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. --- 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 3d4c29d8c3..c09b6f7f36 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.52.0

1 0

[PR] avformat/hls: reject Master Playlist tags in a Media Playlist (PR #23618)
by Kacper Michajłow 27 Jun '26

27 Jun '26

PR #23618 opened by Kacper Michajłow (kasper93) URL: https://code.ffmpeg.org/FFmpeg/FFmpeg/pulls/23618 Patch URL: https://code.ffmpeg.org/FFmpeg/FFmpeg/pulls/23618.patch Fixes: recursion.m3u8 Fixes: cEy2cxyyPaLH Fixes: cd223e0b4e (Add Apple HTTP Live Streaming demuxer) Found-by: BapToutatis From 83401897d7fd7bb5ea032ec74cadee8931e4a9e6 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Kacper=20Michaj=C5=82ow?= <kasper93(a)gmail.com> Date: Sat, 27 Jun 2026 17:43:33 +0200 Subject: [PATCH] avformat/hls: reject Master Playlist tags in a Media Playlist MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Fixes infinite loop on crafted HLS playlists. Fixes: recursion.m3u8 Fixes: cEy2cxyyPaLH Fixes: cd223e0b4e (Add Apple HTTP Live Streaming demuxer) Found-by: BapToutatis Signed-off-by: Kacper Michajłow <kasper93(a)gmail.com> --- libavformat/hls.c | 12 ++++++++++++ 1 file changed, 12 insertions(+) diff --git a/libavformat/hls.c b/libavformat/hls.c index b3b8f3a80a..e76e68ed51 100644 --- a/libavformat/hls.c +++ b/libavformat/hls.c @@ -890,6 +890,12 @@ static int parse_playlist(HLSContext *c, const char *url, while (!avio_feof(in)) { ff_get_chomp_line(in, line, sizeof(line)); if (av_strstart(line, "#EXT-X-STREAM-INF:", &ptr)) { + if (pls) { + av_log(c->ctx, AV_LOG_ERROR, + "Master Playlist tag found in a Media Playlist\n"); + ret = AVERROR_INVALIDDATA; + goto fail; + } is_variant = 1; memset(&variant_info, 0, sizeof(variant_info)); ff_parse_key_value(ptr, handle_variant_args, &variant_info); @@ -909,6 +915,12 @@ static int parse_playlist(HLSContext *c, const char *url, av_strlcpy(key, info.uri, sizeof(key)); } else if (av_strstart(line, "#EXT-X-MEDIA:", &ptr)) { struct rendition_info info = {{0}}; + if (pls) { + av_log(c->ctx, AV_LOG_ERROR, + "Master Playlist tag found in a Media Playlist\n"); + ret = AVERROR_INVALIDDATA; + goto fail; + } ff_parse_key_value(ptr, handle_rendition_args, &info); new_rendition(c, &info, url); } else if (av_strstart(line, "#EXT-X-TARGETDURATION:", &ptr)) { -- 2.52.0

1 0

[PR] avformat/movenc: unify remainder samples signaling (PR #23617)
by James Almer 27 Jun '26

27 Jun '26

PR #23617 opened by James Almer (jamrial) URL: https://code.ffmpeg.org/FFmpeg/FFmpeg/pulls/23617 Patch URL: https://code.ffmpeg.org/FFmpeg/FFmpeg/pulls/23617.patch When writing edit lists and the last packet has skip samples side data to report discard padding, for codecs that have a fixed frame size it's best to always use the full packet duration to calculate track duration (and thus bitrate in some cases) and let the edit list be the sole source of triming information. From a475017603705b7f2ebca4383213ef63ba28f91b Mon Sep 17 00:00:00 2001 From: James Almer <jamrial(a)gmail.com> Date: Sat, 27 Jun 2026 13:50:28 +0000 Subject: [PATCH] avformat/movenc: unify remainder samples signaling When writing edit lists and the last packet has skip samples side data to report discard padding, for codecs that have a fixed frame size it's best to always use the full packet duration to calculate track duration (and thus bitrate in some cases) and let the edit list be the sole source of triming information. Signed-off-by: James Almer <jamrial(a)gmail.com> --- libavformat/movenc.c | 2 ++ tests/ref/fate/autorotate | 4 ++-- tests/ref/fate/mov-mp4-iamf-5_1_4 | 16 ++++++++-------- tests/ref/fate/mov-mp4-iamf-7_1_4-video-first | 14 +++++++------- .../ref/fate/mov-mp4-iamf-7_1_4-video-first-2 | 2 +- .../ref/fate/mov-mp4-iamf-7_1_4-video-first-3 | 2 +- tests/ref/fate/mov-mp4-iamf-7_1_4-video-last | 18 +++++++++--------- tests/ref/fate/mov-mp4-iamf-ambisonic_1 | 12 ++++++------ tests/ref/fate/mov-mp4-iamf-stereo | 6 +++--- 9 files changed, 39 insertions(+), 37 deletions(-) diff --git a/libavformat/movenc.c b/libavformat/movenc.c index c14f05ecc8..2d3994e7c0 100644 --- a/libavformat/movenc.c +++ b/libavformat/movenc.c @@ -7310,6 +7310,8 @@ int ff_mov_write_packet(AVFormatContext *s, AVPacket *pkt) if (sd && sd->size >= 10 && trk->par->frame_size) { duration = FFMAX(av_rescale_q(trk->par->frame_size, (AVRational){ 1, trk->par->sample_rate }, trk->st->time_base), duration); + if (mov->use_editlist) + pkt->duration = duration; duration -= av_rescale_q(AV_RL32(sd->data + 4), (AVRational){ 1, trk->par->sample_rate }, trk->st->time_base); if (duration < 0) diff --git a/tests/ref/fate/autorotate b/tests/ref/fate/autorotate index 1166bb63fc..212035c52c 100644 --- a/tests/ref/fate/autorotate +++ b/tests/ref/fate/autorotate @@ -1,5 +1,5 @@ -bb03affe94dfbe39c8cad8eb69dc367c *tests/data/fate/autorotate.mov -197366 tests/data/fate/autorotate.mov +585a9e2d06f94a887a4bab9a7054a8c4 *tests/data/fate/autorotate.mov +197358 tests/data/fate/autorotate.mov #extradata 0: 34, 0x9d7d073f #tb 0: 1/15360 #media_type 0: video diff --git a/tests/ref/fate/mov-mp4-iamf-5_1_4 b/tests/ref/fate/mov-mp4-iamf-5_1_4 index 437663b099..e7bcce54d0 100644 --- a/tests/ref/fate/mov-mp4-iamf-5_1_4 +++ b/tests/ref/fate/mov-mp4-iamf-5_1_4 @@ -1,5 +1,5 @@ -3e12bd652a42de8329a4171c36b6d5a7 *tests/data/fate/mov-mp4-iamf-5_1_4.mp4 -86164 tests/data/fate/mov-mp4-iamf-5_1_4.mp4 +b8a3031c9ce9abe17368842c9744a18c *tests/data/fate/mov-mp4-iamf-5_1_4.mp4 +86156 tests/data/fate/mov-mp4-iamf-5_1_4.mp4 #extradata 0: 34, 0xa77c0cd6 #extradata 1: 34, 0xa77c0cd6 #extradata 2: 34, 0xa7500cd4 @@ -96,12 +96,12 @@ 3, 36864, 36864, 4096, 1251, 0x0b380fcb 4, 36864, 36864, 4096, 1254, 0xf6f01051 5, 36864, 36864, 4096, 1254, 0xf6f01051 -0, 40960, 40960, 3140, 1235, 0xcf181f77, S=1, Skip Samples, 10, 0x047700bf -1, 40960, 40960, 3140, 1235, 0xcf181f77, S=1, Skip Samples, 10, 0x047700bf -2, 40960, 40960, 3140, 1231, 0x98611f9f, S=1, Skip Samples, 10, 0x047700bf -3, 40960, 40960, 3140, 1231, 0x98611f9f, S=1, Skip Samples, 10, 0x047700bf -4, 40960, 40960, 3140, 1235, 0xcf181f77, S=1, Skip Samples, 10, 0x047700bf -5, 40960, 40960, 3140, 1235, 0xcf181f77, S=1, Skip Samples, 10, 0x047700bf +0, 40960, 40960, 4096, 1235, 0xcf181f77, S=1, Skip Samples, 10, 0x047700bf +1, 40960, 40960, 4096, 1235, 0xcf181f77, S=1, Skip Samples, 10, 0x047700bf +2, 40960, 40960, 4096, 1231, 0x98611f9f, S=1, Skip Samples, 10, 0x047700bf +3, 40960, 40960, 4096, 1231, 0x98611f9f, S=1, Skip Samples, 10, 0x047700bf +4, 40960, 40960, 4096, 1235, 0xcf181f77, S=1, Skip Samples, 10, 0x047700bf +5, 40960, 40960, 4096, 1235, 0xcf181f77, S=1, Skip Samples, 10, 0x047700bf [STREAM_GROUP] index=0 id=0x1 diff --git a/tests/ref/fate/mov-mp4-iamf-7_1_4-video-first b/tests/ref/fate/mov-mp4-iamf-7_1_4-video-first index af1620575e..8154e09e19 100644 --- a/tests/ref/fate/mov-mp4-iamf-7_1_4-video-first +++ b/tests/ref/fate/mov-mp4-iamf-7_1_4-video-first @@ -140,13 +140,13 @@ 0, 10752, 10752, 512, 1884, 0x2c567db6, F=0x0 0, 11264, 11264, 512, 1880, 0x72237c52, F=0x0 0, 11776, 11776, 512, 1970, 0x60708ff8, F=0x0 -1, 40960, 40960, 3140, 1235, 0xcf181f77, S=1, Skip Samples, 10, 0x047700bf -2, 40960, 40960, 3140, 1235, 0xcf181f77, S=1, Skip Samples, 10, 0x047700bf -3, 40960, 40960, 3140, 1231, 0x98611f9f, S=1, Skip Samples, 10, 0x047700bf -4, 40960, 40960, 3140, 1231, 0x98611f9f, S=1, Skip Samples, 10, 0x047700bf -5, 40960, 40960, 3140, 1235, 0xcf181f77, S=1, Skip Samples, 10, 0x047700bf -6, 40960, 40960, 3140, 1235, 0xcf181f77, S=1, Skip Samples, 10, 0x047700bf -7, 40960, 40960, 3140, 1235, 0xcf181f77, S=1, Skip Samples, 10, 0x047700bf +1, 40960, 40960, 4096, 1235, 0xcf181f77, S=1, Skip Samples, 10, 0x047700bf +2, 40960, 40960, 4096, 1235, 0xcf181f77, S=1, Skip Samples, 10, 0x047700bf +3, 40960, 40960, 4096, 1231, 0x98611f9f, S=1, Skip Samples, 10, 0x047700bf +4, 40960, 40960, 4096, 1231, 0x98611f9f, S=1, Skip Samples, 10, 0x047700bf +5, 40960, 40960, 4096, 1235, 0xcf181f77, S=1, Skip Samples, 10, 0x047700bf +6, 40960, 40960, 4096, 1235, 0xcf181f77, S=1, Skip Samples, 10, 0x047700bf +7, 40960, 40960, 4096, 1235, 0xcf181f77, S=1, Skip Samples, 10, 0x047700bf 0, 12288, 12288, 512, 11659, 0x6bcb830e [STREAM_GROUP] index=0 diff --git a/tests/ref/fate/mov-mp4-iamf-7_1_4-video-first-2 b/tests/ref/fate/mov-mp4-iamf-7_1_4-video-first-2 index a0d0a63244..2cca39ef85 100644 --- a/tests/ref/fate/mov-mp4-iamf-7_1_4-video-first-2 +++ b/tests/ref/fate/mov-mp4-iamf-7_1_4-video-first-2 @@ -14,4 +14,4 @@ 0, 28672, 28672, 4096, 1313, 0x2e79ffa2 0, 32768, 32768, 4096, 1323, 0xb9f70fce 0, 36864, 36864, 4096, 1254, 0xf6f01051 -0, 40960, 40960, 3140, 1235, 0xcf181f77, S=1, Skip Samples, 10, 0x047700bf +0, 40960, 40960, 4096, 1235, 0xcf181f77, S=1, Skip Samples, 10, 0x047700bf diff --git a/tests/ref/fate/mov-mp4-iamf-7_1_4-video-first-3 b/tests/ref/fate/mov-mp4-iamf-7_1_4-video-first-3 index 8097c9df1c..49130fe1e3 100644 --- a/tests/ref/fate/mov-mp4-iamf-7_1_4-video-first-3 +++ b/tests/ref/fate/mov-mp4-iamf-7_1_4-video-first-3 @@ -14,4 +14,4 @@ 0, 28672, 28672, 4096, 1310, 0x1f77ff91 0, 32768, 32768, 4096, 1320, 0xea180f1a 0, 36864, 36864, 4096, 1251, 0x0b380fcb -0, 40960, 40960, 3140, 1231, 0x98611f9f, S=1, Skip Samples, 10, 0x047700bf +0, 40960, 40960, 4096, 1231, 0x98611f9f, S=1, Skip Samples, 10, 0x047700bf diff --git a/tests/ref/fate/mov-mp4-iamf-7_1_4-video-last b/tests/ref/fate/mov-mp4-iamf-7_1_4-video-last index 7f858d8cf4..99d89c7c90 100644 --- a/tests/ref/fate/mov-mp4-iamf-7_1_4-video-last +++ b/tests/ref/fate/mov-mp4-iamf-7_1_4-video-last @@ -1,5 +1,5 @@ -b41682705ad3331c1e8a1c1700d9ddd7 *tests/data/fate/mov-mp4-iamf-7_1_4-video-last.mp4 -416432 tests/data/fate/mov-mp4-iamf-7_1_4-video-last.mp4 +1e81db49be5b9631d52ef2bc9ca42913 *tests/data/fate/mov-mp4-iamf-7_1_4-video-last.mp4 +416424 tests/data/fate/mov-mp4-iamf-7_1_4-video-last.mp4 #extradata 0: 34, 0xa77c0cd6 #extradata 1: 34, 0xa77c0cd6 #extradata 2: 34, 0xa7500cd4 @@ -142,13 +142,13 @@ b41682705ad3331c1e8a1c1700d9ddd7 *tests/data/fate/mov-mp4-iamf-7_1_4-video-last. 7, 10752, 10752, 512, 1884, 0x2c567db6, F=0x0 7, 11264, 11264, 512, 1880, 0x72237c52, F=0x0 7, 11776, 11776, 512, 1970, 0x60708ff8, F=0x0 -0, 40960, 40960, 3140, 1235, 0xcf181f77, S=1, Skip Samples, 10, 0x047700bf -1, 40960, 40960, 3140, 1235, 0xcf181f77, S=1, Skip Samples, 10, 0x047700bf -2, 40960, 40960, 3140, 1231, 0x98611f9f, S=1, Skip Samples, 10, 0x047700bf -3, 40960, 40960, 3140, 1231, 0x98611f9f, S=1, Skip Samples, 10, 0x047700bf -4, 40960, 40960, 3140, 1235, 0xcf181f77, S=1, Skip Samples, 10, 0x047700bf -5, 40960, 40960, 3140, 1235, 0xcf181f77, S=1, Skip Samples, 10, 0x047700bf -6, 40960, 40960, 3140, 1235, 0xcf181f77, S=1, Skip Samples, 10, 0x047700bf +0, 40960, 40960, 4096, 1235, 0xcf181f77, S=1, Skip Samples, 10, 0x047700bf +1, 40960, 40960, 4096, 1235, 0xcf181f77, S=1, Skip Samples, 10, 0x047700bf +2, 40960, 40960, 4096, 1231, 0x98611f9f, S=1, Skip Samples, 10, 0x047700bf +3, 40960, 40960, 4096, 1231, 0x98611f9f, S=1, Skip Samples, 10, 0x047700bf +4, 40960, 40960, 4096, 1235, 0xcf181f77, S=1, Skip Samples, 10, 0x047700bf +5, 40960, 40960, 4096, 1235, 0xcf181f77, S=1, Skip Samples, 10, 0x047700bf +6, 40960, 40960, 4096, 1235, 0xcf181f77, S=1, Skip Samples, 10, 0x047700bf 7, 12288, 12288, 512, 11659, 0x6bcb830e [STREAM_GROUP] index=0 diff --git a/tests/ref/fate/mov-mp4-iamf-ambisonic_1 b/tests/ref/fate/mov-mp4-iamf-ambisonic_1 index c789b9b57f..d3504a6bc7 100644 --- a/tests/ref/fate/mov-mp4-iamf-ambisonic_1 +++ b/tests/ref/fate/mov-mp4-iamf-ambisonic_1 @@ -1,5 +1,5 @@ -a7107b8604e26720cea8b2b55dd517b9 *tests/data/fate/mov-mp4-iamf-ambisonic_1.mp4 -57619 tests/data/fate/mov-mp4-iamf-ambisonic_1.mp4 +5b419591f207d4a2636ef93a41446244 *tests/data/fate/mov-mp4-iamf-ambisonic_1.mp4 +57611 tests/data/fate/mov-mp4-iamf-ambisonic_1.mp4 #extradata 0: 34, 0x9f9f0c6b #extradata 1: 34, 0x9f9f0c6b #extradata 2: 34, 0x9f9f0c6b @@ -64,10 +64,10 @@ a7107b8604e26720cea8b2b55dd517b9 *tests/data/fate/mov-mp4-iamf-ambisonic_1.mp4 1, 36864, 36864, 4096, 1251, 0x0b380fcb 2, 36864, 36864, 4096, 1251, 0x0b380fcb 3, 36864, 36864, 4096, 1251, 0x0b380fcb -0, 40960, 40960, 3140, 1231, 0x98611f9f, S=1, Skip Samples, 10, 0x047700bf -1, 40960, 40960, 3140, 1231, 0x98611f9f, S=1, Skip Samples, 10, 0x047700bf -2, 40960, 40960, 3140, 1231, 0x98611f9f, S=1, Skip Samples, 10, 0x047700bf -3, 40960, 40960, 3140, 1231, 0x98611f9f, S=1, Skip Samples, 10, 0x047700bf +0, 40960, 40960, 4096, 1231, 0x98611f9f, S=1, Skip Samples, 10, 0x047700bf +1, 40960, 40960, 4096, 1231, 0x98611f9f, S=1, Skip Samples, 10, 0x047700bf +2, 40960, 40960, 4096, 1231, 0x98611f9f, S=1, Skip Samples, 10, 0x047700bf +3, 40960, 40960, 4096, 1231, 0x98611f9f, S=1, Skip Samples, 10, 0x047700bf [STREAM_GROUP] index=0 id=0x1 diff --git a/tests/ref/fate/mov-mp4-iamf-stereo b/tests/ref/fate/mov-mp4-iamf-stereo index 11b3aca65a..ce00ebc72e 100644 --- a/tests/ref/fate/mov-mp4-iamf-stereo +++ b/tests/ref/fate/mov-mp4-iamf-stereo @@ -1,5 +1,5 @@ -cd36f114d163e2cd93e6e80a68726170 *tests/data/fate/mov-mp4-iamf-stereo.mp4 -15138 tests/data/fate/mov-mp4-iamf-stereo.mp4 +e3da79a8ab08184a51dd0bc2c752b3b3 *tests/data/fate/mov-mp4-iamf-stereo.mp4 +15130 tests/data/fate/mov-mp4-iamf-stereo.mp4 #extradata 0: 34, 0xa77c0cd6 #tb 0: 1/44100 #media_type 0: audio @@ -16,7 +16,7 @@ cd36f114d163e2cd93e6e80a68726170 *tests/data/fate/mov-mp4-iamf-stereo.mp4 0, 28672, 28672, 4096, 1313, 0x2e79ffa2 0, 32768, 32768, 4096, 1323, 0xb9f70fce 0, 36864, 36864, 4096, 1254, 0xf6f01051 -0, 40960, 40960, 3140, 1235, 0xcf181f77, S=1, Skip Samples, 10, 0x047700bf +0, 40960, 40960, 4096, 1235, 0xcf181f77, S=1, Skip Samples, 10, 0x047700bf [STREAM_GROUP] index=0 id=0x1 -- 2.52.0

1 0

[PR] hevc-interlace (PR #23616)
by Torbjörn Einarsson 27 Jun '26

27 Jun '26

PR #23616 opened by Torbjörn Einarsson (tobbee) URL: https://code.ffmpeg.org/FFmpeg/FFmpeg/pulls/23616 Patch URL: https://code.ffmpeg.org/FFmpeg/FFmpeg/pulls/23616.patch # Summary of changes Unlike H.264, where field coding (PAFF/MBAFF) is part of the bitstream and the decoder already returns a complementary field pair as one full-height frame, HEVC has no field coding at all: interlaced content is just a sequence of independent half-height field pictures with distinct POCs, flagged only via the PTL source-scan flags, VUI field_seq_flag and the picture timing SEI. A plain decoder therefore emits half-height pictures at field rate. Currently, FFmpeg does not properly handle interlaced HEVC. This PR makes the decoder weave the two coded fields of a pair into one full-height frame, bringing field-coded HEVC to parity with H.264 PAFF: a 1080i50 stream decodes to 1920x1080 @ 25fps, flagged INTERLACED with TOP_FIELD_FIRST set for top-field-first content and cleared for bottom-field-first. A previous attempt (Jose Santiago, V1..V7, Oct-Nov 2024) was not merged as too complex for "just weave inside a decoder"; this is a deliberately minimal take: - No copy: the leader (first field in decode order) allocates one full-height buffer; both fields are per-plane views into it with doubled linesize (top->even lines, bottom->odd). Reconstruction is unchanged; inter-field prediction works through the normal POC-based DPB. - No new subsystem: no construction context, no extra FIFO, no mutex. Pairing is a small leader/follower state finalized at the second field's frame_start before ff_thread_finish_setup(), so it is correct under frame threading via the existing thread-context propagation. Output is byte-identical single- vs multi-threaded. - Contained scope: SW decoding only; hwaccel is untouched (opaque surfaces can't be woven via strides), re-checked after get_format(). Field order for bare pic_struct 1/2 (no pairing hint) is anchored on IRAP pictures; the explicit 9..12 hints are handled directly. Testing: full FATE passes after update; fate-hevc-paired-fields (10-bit 4:2:2, explicit hints) is updated, and a new fate-hevc-paired-fields-420 (8-bit 4:2:0, bare pic_struct 1/2) exercises the IRAP-anchored path. The new test needs the attached sample. ```fate-samples hevc/paired_fields_420.hevc ``` From 684914b7b1565fe96ab0c7dc40725b9a19b041c0 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Torbjo=CC=88rn=20Einarsson?= <torbjorn(a)einarssons.se> Date: Tue, 9 Jun 2026 15:20:50 +0200 Subject: [PATCH 1/4] avcodec/hevc: store raw pic_struct and source_scan_type from pic timing SEI MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit When frame_field_info_present_flag is set, the picture timing SEI carries pic_struct, source_scan_type and duplicate_flag (H.265 D.3.3), but only the derived picture_struct was kept and the latter two were left unread. Store the raw pic_struct and source_scan_type, flag the timing info as present, and read the previously skipped source_scan_type and duplicate_flag. A following commit uses these to report the field order. Signed-off-by: Torbjörn Einarsson <torbjorn(a)einarssons.se> --- libavcodec/hevc/sei.c | 4 ++++ libavcodec/hevc/sei.h | 3 +++ 2 files changed, 7 insertions(+) diff --git a/libavcodec/hevc/sei.c b/libavcodec/hevc/sei.c index 83c726a217..70b0a6b17f 100644 --- a/libavcodec/hevc/sei.c +++ b/libavcodec/hevc/sei.c @@ -60,6 +60,8 @@ static int decode_nal_sei_pic_timing(HEVCSEI *s, GetBitContext *gb, if (sps->vui.frame_field_info_present_flag) { int pic_struct = get_bits(gb, 4); + h->present = 1; + h->pic_struct = pic_struct; h->picture_struct = AV_PICTURE_STRUCTURE_UNKNOWN; if (pic_struct == 2 || pic_struct == 10 || pic_struct == 12) { av_log(logctx, AV_LOG_DEBUG, "BOTTOM Field\n"); @@ -74,6 +76,8 @@ static int decode_nal_sei_pic_timing(HEVCSEI *s, GetBitContext *gb, av_log(logctx, AV_LOG_DEBUG, "Frame/Field Tripling\n"); h->picture_struct = HEVC_SEI_PIC_STRUCT_FRAME_TRIPLING; } + h->source_scan_type = get_bits(gb, 2); + skip_bits1(gb); // duplicate_flag } return 0; diff --git a/libavcodec/hevc/sei.h b/libavcodec/hevc/sei.h index 59bd2b45f8..55f03e0ede 100644 --- a/libavcodec/hevc/sei.h +++ b/libavcodec/hevc/sei.h @@ -52,6 +52,9 @@ typedef struct HEVCSEIFramePacking { typedef struct HEVCSEIPictureTiming { int picture_struct; + int pic_struct; ///< raw pic_struct (H.265 Table D.2), valid if present + int source_scan_type; ///< 0: interlaced, 1: progressive, 2: unknown + int present; ///< a pic_timing SEI with frame_field_info was parsed } HEVCSEIPictureTiming; typedef struct HEVCSEIAlternativeTransfer { -- 2.52.0 From 426a5c6af0b6b2748172759c46df8ce8f1f9c844 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Torbjo=CC=88rn=20Einarsson?= <torbjorn(a)einarssons.se> Date: Tue, 9 Jun 2026 15:20:50 +0200 Subject: [PATCH 2/4] avcodec/hevc: set field_order from picture timing SEI in the parser MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit hevc_parse_slice_header assigned the SEI picture_struct, an AV_PICTURE_STRUCTURE_* value, directly to AVCodecParserContext.field_order which is an enum AVFieldOrder. The enums are unrelated, so a top field was reported as "progressive" and a coded frame as "bb". Derive the field order from the picture timing SEI pic_struct, falling back to the profile_tier_level source-scan flags, mirroring the H.264 parser. ffprobe now reports interlaced HEVC as tt/bb instead of progressive. Signed-off-by: Torbjörn Einarsson <torbjorn(a)einarssons.se> --- libavcodec/hevc/parser.c | 37 ++++++++++++++++++++++++++++++++++++- 1 file changed, 36 insertions(+), 1 deletion(-) diff --git a/libavcodec/hevc/parser.c b/libavcodec/hevc/parser.c index 47a1ac70d0..508d093807 100644 --- a/libavcodec/hevc/parser.c +++ b/libavcodec/hevc/parser.c @@ -53,6 +53,40 @@ typedef struct HEVCParserContext { int pocTid0; } HEVCParserContext; +/* + * Derive the stream field order. HEVC has no field coding, so interlace is + * carried as metadata: primarily by the pic_struct of the picture timing SEI + * (H.265 Table D.2), with the profile_tier_level source-scan flags as a + * fallback. This mirrors the H.264 parser (h264_parser.c). Top-field-first + * pic_struct values are {1,3,5,9,11}, bottom-field-first {2,4,6,10,12}. + */ +static enum AVFieldOrder hevc_field_order(const HEVCSPS *sps, + const HEVCSEIPictureTiming *pt) +{ + const PTLCommon *ptl = &sps->ptl.general_ptl; + + if (pt->present) { + switch (pt->pic_struct) { + case 1: case 3: case 5: case 9: case 11: + return AV_FIELD_TT; + case 2: case 4: case 6: case 10: case 12: + return AV_FIELD_BB; + case 0: case 7: case 8: + return AV_FIELD_PROGRESSIVE; + } + if (pt->source_scan_type == 1) + return AV_FIELD_PROGRESSIVE; + } + + /* No usable pic_struct: fall back to the profile_tier_level flags. */ + if (ptl->progressive_source_flag && !ptl->interlaced_source_flag) + return AV_FIELD_PROGRESSIVE; + if (ptl->interlaced_source_flag && !ptl->progressive_source_flag) + return AV_FIELD_TT; /* interlaced; order not signalled, assume top first */ + + return AV_FIELD_UNKNOWN; +} + static int hevc_parse_slice_header(AVCodecParserContext *s, H2645NAL *nal, AVCodecContext *avctx) { @@ -70,7 +104,6 @@ static int hevc_parse_slice_header(AVCodecParserContext *s, H2645NAL *nal, first_slice_in_pic_flag = get_bits1(gb); s->picture_structure = sei->picture_timing.picture_struct; - s->field_order = sei->picture_timing.picture_struct; if (IS_IRAP_NAL(nal)) { s->key_frame = 1; @@ -85,6 +118,8 @@ static int hevc_parse_slice_header(AVCodecParserContext *s, H2645NAL *nal, pps = ps->pps_list[pps_id]; sps = pps->sps; + s->field_order = hevc_field_order(sps, &sei->picture_timing); + ow = &sps->output_window; s->coded_width = sps->width; -- 2.52.0 From 7faa4fb6d5df9ec6483d55617bf619b27c44d76a Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Torbjo=CC=88rn=20Einarsson?= <torbjorn(a)einarssons.se> Date: Fri, 12 Jun 2026 00:12:55 +0200 Subject: [PATCH 3/4] avcodec/hevc: combine interlaced field pairs into full frames MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit HEVC has no field coding; broadcast-style interlaced content is coded as a sequence of half-height field pictures, signalled only through the profile_tier_level source-scan flags, VUI field_seq_flag and the picture timing SEI. Decoders that present those pictures as-is yield half-height output at field rate, unlike H.264 PAFF where the two fields of a pair are returned as one full-height frame. Decode the two fields of a complementary pair into one shared full-height buffer and output it once as a woven frame carrying the INTERLACED and TOP_FIELD_FIRST flags, so field-coded HEVC is presented like H.264 PAFF: a 1080i50 stream decodes to 1920x1080 frames at 25 fps. There is no copying: the first field of a pair (the leader, in decode order) allocates the full-height buffer and both fields' frames become per-plane views into it with doubled linesize, the top field on even and the bottom field on odd lines, independently of which one leads, so both TFF and BFF streams work. The reconstruction code takes strides as parameters throughout and needs no changes; inter-field prediction works through the regular POC-based DPB since the fields keep distinct POCs. Pairing follows the pic_struct pairing hints (Table D.2): 11/12 lead, 9/10 follow; bare 1/2, which carry no pairing hint, take their role from the field order locked at the first IRAP field of the sequence (an IRAP starts a frame, so it leads its pair and its parity gives the field order), falling back to decode order before the first IRAP. Combining engages only for software decoding of streams that signal interlaced_source_flag and not progressive_source_flag, set VUI field_seq_flag, and carry picture timing SEI marking field pictures; everything else, including hwaccel decoding, is unchanged (the geometry is re-exported after get_format() so it follows the hwaccel decision). When combining, the exported coded/display height is doubled and the frame rate halved so callers see the woven geometry. The pair is finalized at the second field's frame_start: the leader gets its OUTPUT flag, and the woven frame its flags, doubled crop and duration, and the leader's side data, all before ff_thread_finish_setup(). This makes the scheme work under frame threading: the DPB state every later worker snapshots is consistent, so completed pairs are emitted at sequence boundaries and at EOF. The woven frame cannot reach the caller before both fields are decoded, because frames are returned by the decode call that bumps them from the DPB and the caller receives a worker's output only after that worker and all earlier ones (including the leader's) have finished; the pending leader is carried across worker contexts in update_thread_context(). Output is byte-identical to single-threaded decoding and deterministic over repeated runs, verified for TFF and BFF streams in 8-bit 4:2:0 and 10-bit 4:2:2. An unpaired field is dropped with a warning. Film grain synthesis is not applied to woven frames. verify_md5() uses the frame's own (field) dimensions instead of the doubled avctx coded size. fate-hevc-paired-fields now outputs two full-height frames instead of four fields; extend it to also check the frame dimensions. Add fate-hevc-paired-fields-420, an 8-bit 4:2:0 stream signalling bare pic_struct 1/2, which exercises field-order inference from the IRAP anchor; it needs the new paired_fields_420.hevc sample. Signed-off-by: Torbjörn Einarsson <torbjorn(a)einarssons.se> --- libavcodec/hevc/hevcdec.c | 274 +++++++++++++++++++++++++- libavcodec/hevc/hevcdec.h | 36 ++++ libavcodec/hevc/refs.c | 75 ++++++- tests/fate/hevc.mak | 5 +- tests/ref/fate/hevc-paired-fields | 12 +- tests/ref/fate/hevc-paired-fields-420 | 14 ++ 6 files changed, 397 insertions(+), 19 deletions(-) create mode 100644 tests/ref/fate/hevc-paired-fields-420 diff --git a/libavcodec/hevc/hevcdec.c b/libavcodec/hevc/hevcdec.c index b4c2d82e8d..a0215dec65 100644 --- a/libavcodec/hevc/hevcdec.c +++ b/libavcodec/hevc/hevcdec.c @@ -329,18 +329,40 @@ static int decode_lt_rps(const HEVCSPS *sps, LongTermRPS *rps, return 0; } +/* + * combine_fields: whether the woven full-height geometry should be advertised + * for this SPS. Fields are only combined in software decoding of an + * interlaced-source stream whose coded pictures are individual fields + * (VUI field_seq_flag, D.3.3); a frame-coded interlaced stream must not have + * its geometry doubled. Advertising doubled dimensions while emitting + * per-field frames would corrupt buffer allocation, so the predicate must + * match the runtime combine gate (combine_decide_role() additionally requires + * per-picture timing SEI). + */ +static int combine_geometry_active(const HEVCContext *s, const HEVCSPS *sps) +{ + const PTLCommon *ptl = &sps->ptl.general_ptl; + + return !s->avctx->hwaccel && sps->vui.field_seq_flag && + ptl->interlaced_source_flag && !ptl->progressive_source_flag; +} + static void export_stream_params(HEVCContext *s, const HEVCSPS *sps) { AVCodecContext *avctx = s->avctx; const HEVCVPS *vps = sps->vps; const HEVCWindow *ow = &sps->output_window; unsigned int num = 0, den = 0; + /* combine_fields weaves two field pictures into one frame of twice the + * height at half the field rate; advertise that geometry to the caller. */ + const int combine = combine_geometry_active(s, sps); + const int vmul = combine ? 2 : 1; avctx->pix_fmt = sps->pix_fmt; avctx->coded_width = sps->width; - avctx->coded_height = sps->height; + avctx->coded_height = sps->height * vmul; avctx->width = sps->width - ow->left_offset - ow->right_offset; - avctx->height = sps->height - ow->top_offset - ow->bottom_offset; + avctx->height = (sps->height - ow->top_offset - ow->bottom_offset) * vmul; avctx->has_b_frames = sps->temporal_layer[sps->max_sub_layers - 1].num_reorder_pics; avctx->profile = sps->ptl.general_ptl.profile_idc; avctx->level = sps->ptl.general_ptl.level_idc; @@ -381,8 +403,11 @@ static void export_stream_params(HEVCContext *s, const HEVCSPS *sps) } if (num > 0 && den > 0) + /* num is num_units_in_tick (the framerate denominator), den is + * time_scale (the numerator); combining two fields into one frame + * halves the rate, i.e. doubles the tick count per output frame. */ av_reduce(&avctx->framerate.den, &avctx->framerate.num, - num, den, 1 << 30); + num * vmul, den, 1 << 30); } static int export_stream_params_from_sei(HEVCContext *s) @@ -3194,9 +3219,159 @@ static int find_finish_setup_nal(const HEVCContext *s) return nal_idx; } +/* + * combine_fields: decide the pairing role of the field picture about to be + * started. Returns HEVC_COMBINE_LEADER for the first field of a pair (which + * allocates the full-height buffer), HEVC_COMBINE_FOLLOWER for the second field + * (which completes a pending pair), or HEVC_COMBINE_NONE otherwise (combining + * disabled, progressive content, hwaccel, non-base layer, or no usable picture + * timing SEI). On a non-NONE return *bottom is set to the field's spatial + * parity: 1 if it occupies the bottom (odd) lines of the woven frame, 0 for the + * top (even) lines. The leader/follower role follows decode order and is + * independent of parity, so both top-field-first (leader = top) and + * bottom-field-first (leader = bottom) streams are handled. + * + * The pairing direction comes from the picture timing pic_struct (H.265 Table + * D.2): values 11/12 are a top/bottom field paired with the *next* field in + * output order (the pair leader), 9/10 a top/bottom field paired with the + * *previous* field (the follower). Bare top/bottom fields (1/2) carry no + * pairing hint; their role is taken from the field order locked at the first + * IRAP field of the sequence (an IRAP starts a frame, hence leads its pair, so + * its parity reveals whether the stream is top- or bottom-field-first). Before + * the first IRAP the order is provisionally seeded from decode order. + */ +static int combine_decide_role(HEVCContext *s, const HEVCLayerContext *l, + int *bottom) +{ + const HEVCSEIPictureTiming *pt = &s->sei.picture_timing; + + *bottom = 0; + + if (l != &s->layers[0] || !combine_geometry_active(s, l->sps)) + return HEVC_COMBINE_NONE; + + if (!pt->present || pt->source_scan_type == 1 /* progressive */) + return HEVC_COMBINE_NONE; + + switch (pt->pic_struct) { + case 11: /* top field, paired with next bottom field -> leader, top */ + *bottom = 0; + return HEVC_COMBINE_LEADER; + case 12: /* bottom field, paired with next top field -> leader, bottom */ + *bottom = 1; + return HEVC_COMBINE_LEADER; + case 9: /* top field, paired with previous bottom field -> follower, top */ + *bottom = 0; + return s->combine_leader ? HEVC_COMBINE_FOLLOWER : HEVC_COMBINE_NONE; + case 10: /* bottom field, paired with previous top field -> follower, bottom */ + *bottom = 1; + return s->combine_leader ? HEVC_COMBINE_FOLLOWER : HEVC_COMBINE_NONE; + case 1: /* bare top field -> role from the locked field order */ + case 2: /* bare bottom field -> role from the locked field order */ + { + int parity = pt->pic_struct == 2; /* spatial parity: 0 top, 1 bottom */ + *bottom = parity; + + /* An IRAP field necessarily starts a frame, so it is a pair leader, + * and its parity fixes the field order for the sequence (top -> TFF, + * bottom -> BFF). Real interlaced streams place the IRAP on the first + * field only, so this seeds and re-anchors the pairing phase, also + * recovering from any earlier dropped field. */ + if (IS_IRAP(s)) { + s->combine_lead_bottom = parity; + return HEVC_COMBINE_LEADER; + } + + /* Field order known: a field matching the leading parity opens a pair, + * the opposite parity closes the pending one. */ + if (s->combine_lead_bottom >= 0) { + if (parity == s->combine_lead_bottom) + return HEVC_COMBINE_LEADER; + return s->combine_leader ? HEVC_COMBINE_FOLLOWER : HEVC_COMBINE_NONE; + } + + /* Field order not yet established (entry before the first IRAP): with + * no leader pending this field opens the pair and provisionally seeds + * the order; the next IRAP corrects it if the guess was wrong. */ + if (!s->combine_leader) { + s->combine_lead_bottom = parity; + return HEVC_COMBINE_LEADER; + } + return HEVC_COMBINE_FOLLOWER; + } + } + + return HEVC_COMBINE_NONE; +} + +/* + * combine_fields: the second field of the pair has started decoding into the + * shared buffer, so finalize the full-height "combined" frame the leader owns + * and flag the leader for output. The combined frame already carries pts/ + * duration/SAR/colorimetry from its own get_buffer() at leader frame start; + * here we set the per-frame properties the decoder normally fills and fix up + * the geometry that differs from a single field. + */ +static int combine_finalize(HEVCFrame *leader, int tff) +{ + AVFrame *combined = leader->combined; + + combined->pict_type = leader->f->pict_type; + combined->flags |= AV_FRAME_FLAG_INTERLACED; + if (tff) + combined->flags |= AV_FRAME_FLAG_TOP_FIELD_FIRST; + else + combined->flags &= ~AV_FRAME_FLAG_TOP_FIELD_FIRST; + if (leader->f->flags & AV_FRAME_FLAG_KEY) + combined->flags |= AV_FRAME_FLAG_KEY; + + /* one field line crops to two frame lines; columns are unaffected */ + combined->crop_top = leader->f->crop_top * 2; + combined->crop_bottom = leader->f->crop_bottom * 2; + combined->crop_left = leader->f->crop_left; + combined->crop_right = leader->f->crop_right; + + if (combined->duration > 0) + combined->duration *= 2; + + /* per-frame SEI side data (A53 CC, HDR metadata, ...) was attached to the + * leader's field view in set_side_data(); the woven frame only has the + * packet-level properties from get_buffer() */ + for (int i = 0; i < leader->f->nb_side_data; i++) { + int ret = av_frame_side_data_clone(&combined->side_data, + &combined->nb_side_data, + leader->f->side_data[i], + AV_FRAME_SIDE_DATA_FLAG_UNIQUE); + if (ret < 0) + return ret; + } + + leader->flags |= HEVC_FRAME_FLAG_OUTPUT; + + return 0; +} + +/* + * combine_fields: drop a pending pair leader whose second field can no longer + * arrive (sequence boundary or another leader showing up). The half-filled + * woven buffer is never output; the field itself stays in the DPB as a + * reference until it ages out. + */ +static void combine_drop_leader(HEVCContext *s) +{ + if (s->combine_leader) { + av_log(s->avctx, AV_LOG_WARNING, + "combine_fields: dropping unpaired field with POC %d\n", + s->combine_leader->poc); + s->combine_leader = NULL; + } +} + static int hevc_frame_start(HEVCContext *s, HEVCLayerContext *l, unsigned nal_idx) { + int combine_role = HEVC_COMBINE_NONE; + int combine_bottom = 0; const HEVCPPS *const pps = s->ps.pps_list[s->sh.pps_id]; const HEVCSPS *const sps = pps->sps; int pic_size_in_ctb = ((sps->width >> sps->log2_min_cb_size) + 1) * @@ -3248,6 +3423,7 @@ static int hevc_frame_start(HEVCContext *s, HEVCLayerContext *l, } } + combine_drop_leader(s); ff_hevc_clear_refs(l); ret = set_sps(s, l, sps); @@ -3263,6 +3439,16 @@ static int hevc_frame_start(HEVCContext *s, HEVCLayerContext *l, return ret; } + /* get_format() may have enabled a hwaccel, which disables field + * combining; re-export so the advertised geometry matches the + * per-field frames that will be produced, keeping the negotiated + * pixel format */ + if (s->avctx->hwaccel) { + enum AVPixelFormat pix_fmt = s->avctx->pix_fmt; + export_stream_params(s, sps); + s->avctx->pix_fmt = pix_fmt; + } + new_sequence = 1; } } @@ -3273,8 +3459,10 @@ static int hevc_frame_start(HEVCContext *s, HEVCLayerContext *l, memset(l->is_pcm, 0, (sps->min_pu_width + 1) * (sps->min_pu_height + 1)); memset(l->tab_slice_address, -1, pic_size_in_ctb * sizeof(*l->tab_slice_address)); - if (IS_IDR(s)) + if (IS_IDR(s)) { + combine_drop_leader(s); ff_hevc_clear_refs(l); + } s->slice_idx = 0; s->first_nal_type = s->nal_unit_type; @@ -3319,10 +3507,41 @@ static int hevc_frame_start(HEVCContext *s, HEVCLayerContext *l, if (ret < 0) return ret; + combine_role = combine_decide_role(s, l, &combine_bottom); + s->combine_field_role = combine_role; + s->combine_field_bottom = combine_bottom; ret = ff_hevc_set_new_ref(s, l, s->poc); + s->combine_field_role = HEVC_COMBINE_NONE; + s->combine_field_bottom = 0; if (ret < 0) goto fail; + if (combine_role == HEVC_COMBINE_LEADER) { + /* a still-pending leader means the previous pair never received its + * second field */ + combine_drop_leader(s); + /* hold the leader's output until the matching second field arrives */ + s->combine_leader = s->cur_frame; + s->combine_leader_bottom = combine_bottom; + s->cur_frame->flags &= ~HEVC_FRAME_FLAG_OUTPUT; + } else if (combine_role == HEVC_COMBINE_FOLLOWER) { + /* The second field completes the pair: release the woven frame for + * output, ordered by the leader's POC. Its samples are only fully + * decoded once this field finishes, but it cannot reach the caller + * earlier: output frames are returned by the decode call that bumps + * them from the DPB, and under frame threading the caller receives a + * worker's frames only after that worker (and all earlier ones, + * including the leader's) fully finished. Finalizing here, before + * ff_thread_finish_setup(), keeps the DPB state that later workers + * snapshot consistent, so the pair is also emitted when a sequence + * boundary or EOF follows immediately. */ + s->cur_frame->flags &= ~HEVC_FRAME_FLAG_OUTPUT; + ret = combine_finalize(s->combine_leader, !s->combine_leader_bottom); + s->combine_leader = NULL; + if (ret < 0) + goto fail; + } + ret = ff_hevc_frame_rps(s, l); if (ret < 0) { av_log(s->avctx, AV_LOG_ERROR, "Error constructing the frame RPS.\n"); @@ -3340,6 +3559,11 @@ static int hevc_frame_start(HEVCContext *s, HEVCLayerContext *l, !(s->avctx->export_side_data & AV_CODEC_EXPORT_DATA_FILM_GRAIN) && !s->avctx->hwaccel; + /* film grain would have to be applied to the woven buffer, not the field + * views; not supported when combining fields */ + if (combine_role != HEVC_COMBINE_NONE) + s->cur_frame->needs_fg = 0; + ret = set_side_data(s); if (ret < 0) goto fail; @@ -3395,6 +3619,8 @@ static int hevc_frame_start(HEVCContext *s, HEVCLayerContext *l, return 0; fail: + if (s->combine_leader == l->cur_frame) + s->combine_leader = NULL; if (l->cur_frame) ff_hevc_unref_frame(l->cur_frame, ~0); l->cur_frame = NULL; @@ -3429,8 +3655,10 @@ static int verify_md5(HEVCContext *s, AVFrame *frame) msg_buf[0] = '\0'; for (i = 0; frame->data[i]; i++) { - int width = s->avctx->coded_width; - int height = s->avctx->coded_height; + /* the frame's own dimensions, not avctx->coded_*: when combining + * fields the checksum covers the half-height field view */ + int width = frame->width; + int height = frame->height; int w = (i == 1 || i == 2) ? (width >> desc->log2_chroma_w) : width; int h = (i == 1 || i == 2) ? (height >> desc->log2_chroma_h) : height; uint8_t md5[16]; @@ -3899,6 +4127,24 @@ static int hevc_ref_frame(HEVCFrame *dst, const HEVCFrame *src) dst->needs_fg = 1; } + /* combine_fields: the leader's full-height woven buffer must travel with the + * frame across thread contexts (the follower writes into it and it is the + * frame finally output). It is a plain refcounted AVFrame, so give each DPB + * entry its own reference rather than sharing the pointer (which would + * double-free in ff_hevc_unref_frame). */ + if (src->combined) { + dst->combined = av_frame_alloc(); + if (!dst->combined) { + ff_hevc_unref_frame(dst, ~0); + return AVERROR(ENOMEM); + } + ret = av_frame_ref(dst->combined, src->combined); + if (ret < 0) { + ff_hevc_unref_frame(dst, ~0); + return ret; + } + } + dst->pps = av_refstruct_ref_c(src->pps); dst->tab_mvf = av_refstruct_ref(src->tab_mvf); dst->rpl_tab = av_refstruct_ref(src->rpl_tab); @@ -3999,6 +4245,7 @@ static av_cold int hevc_init_context(AVCodecContext *avctx) s->dovi_ctx.logctx = avctx; s->eos = 0; + s->combine_lead_bottom = -1; ff_hevc_reset_sei(&s->sei); @@ -4048,6 +4295,19 @@ static int hevc_update_thread_context(AVCodecContext *dst, s->eos = s0->eos; s->no_rasl_output_flag = s0->no_rasl_output_flag; + /* combine_fields: carry the pending pair leader to the worker that will + * decode the follower. The DPB was copied index-for-index above, so the + * leader lives at the same slot; re-point into this context's base-layer + * DPB. */ + s->combine_leader_bottom = s0->combine_leader_bottom; + s->combine_lead_bottom = s0->combine_lead_bottom; + if (s0->combine_leader) { + size_t idx = s0->combine_leader - s0->layers[0].DPB; + s->combine_leader = &s->layers[0].DPB[idx]; + } else { + s->combine_leader = NULL; + } + s->is_nalff = s0->is_nalff; s->nal_length_size = s0->nal_length_size; s->layers_active_decode = s0->layers_active_decode; @@ -4187,6 +4447,8 @@ static av_cold int hevc_decode_init(AVCodecContext *avctx) static av_cold void hevc_decode_flush(AVCodecContext *avctx) { HEVCContext *s = avctx->priv_data; + s->combine_leader = NULL; + s->combine_lead_bottom = -1; ff_hevc_flush_dpb(s); ff_hevc_reset_sei(&s->sei); ff_dovi_ctx_flush(&s->dovi_ctx); diff --git a/libavcodec/hevc/hevcdec.h b/libavcodec/hevc/hevcdec.h index 8394740c4b..21460e19b3 100644 --- a/libavcodec/hevc/hevcdec.h +++ b/libavcodec/hevc/hevcdec.h @@ -357,6 +357,17 @@ typedef struct DBParams { #define HEVC_FRAME_FLAG_UNAVAILABLE (1 << 3) #define HEVC_FRAME_FLAG_CORRUPT (1 << 4) +/* combine_fields: pairing role of the field picture currently being started, + * used to decide how its decode buffer is allocated (see alloc_frame()). The + * role reflects decode order (leader is the first field of the pair), which is + * independent of spatial parity: for top-field-first content the leader is the + * top field, for bottom-field-first content it is the bottom field. The spatial + * parity (which lines of the woven frame a field writes) is carried separately + * by combine_field_bottom. */ +#define HEVC_COMBINE_NONE 0 +#define HEVC_COMBINE_LEADER 1 ///< first field of the pair: owns the full-height buffer +#define HEVC_COMBINE_FOLLOWER 2 ///< second field of the pair: a view into the leader's buffer + typedef struct HEVCFrame { union { struct { @@ -376,6 +387,15 @@ typedef struct HEVCFrame { RefPicListTab *rpl; ///< RefStruct reference int nb_rpl_elems; + /** + * combine_fields: for the leader (the first field, in decode order) of a + * combined field pair, this is the full-height AVFrame that both fields are + * decoded into and which is output once as a single full-height frame. f is + * a doubled-stride view into this buffer. NULL for followers and for + * normally-decoded frames. + */ + AVFrame *combined; + void *hwaccel_picture_private; ///< RefStruct reference // for secondary-layer frames, this is the DPB index of the base-layer frame @@ -561,6 +581,22 @@ typedef struct HEVCContext { ///< as a format defined in 14496-15 int apply_defdispwin; + /// while mid-pair, the leader (first) field frame the follower combines into + HEVCFrame *combine_leader; + /// spatial parity of the held leader: 1 if it writes the bottom (odd) lines, + /// i.e. the pair is bottom-field-first; used to set TOP_FIELD_FIRST + int combine_leader_bottom; + /// transient HEVC_COMBINE_* role for the frame currently being allocated + int combine_field_role; + /// transient spatial parity (1 = bottom/odd lines) of the frame being allocated + int combine_field_bottom; + /// combine_fields: locked field order for bare top/bottom fields (pic_struct + /// 1/2, which carry no pairing hint). Parity of the field that leads each + /// pair, matching combine_leader_bottom: -1 unknown, 0 top-field-first, + /// 1 bottom-field-first. Set from the first IRAP field of the sequence (an + /// IRAP starts a frame, hence leads its pair) and re-anchored at every IRAP. + int combine_lead_bottom; + // multi-layer AVOptions int *view_ids; unsigned nb_view_ids; diff --git a/libavcodec/hevc/refs.c b/libavcodec/hevc/refs.c index 2acffd72d9..17b9810ef8 100644 --- a/libavcodec/hevc/refs.c +++ b/libavcodec/hevc/refs.c @@ -39,6 +39,7 @@ void ff_hevc_unref_frame(HEVCFrame *frame, int flags) frame->flags = 0; if (!frame->flags) { ff_progress_frame_unref(&frame->tf); + av_frame_free(&frame->combined); av_frame_unref(frame->frame_grain); frame->needs_fg = 0; @@ -107,6 +108,63 @@ static int replace_alpha_plane(AVFrame *alpha, AVFrame *base) return AVERROR_BUG; } +/* + * combine_fields: set up frame->f as a doubled-stride field view into a + * full-height "combined" buffer that holds both fields of an interlaced pair. + * + * For the pair leader (the first field in decode order) a fresh full-height + * (2*sps->height) AVFrame is allocated and stored in frame->combined; for the + * follower the leader's combined buffer is shared. Independently of which field + * leads, the spatial parity selects the lines written: a top field maps the + * even lines, a bottom field the odd lines (data offset by one line). In both + * cases f keeps the field's logical dimensions (sps->width x sps->height); only + * data[] and linesize[] describe the interleaved view, so reconstruction, SAO, + * deblocking and motion compensation all operate in field geometry and + * naturally write every other line of the shared frame. + */ +static int combine_setup_field_view(HEVCContext *s, HEVCLayerContext *l, + HEVCFrame *frame, int is_leader, int bottom) +{ + AVFrame *f = frame->f; + AVFrame *combined; + int ret, i; + + if (is_leader) { + combined = av_frame_alloc(); + if (!combined) + return AVERROR(ENOMEM); + combined->format = s->avctx->pix_fmt; + combined->width = l->sps->width; + combined->height = l->sps->height * 2; + ret = ff_thread_get_buffer(s->avctx, combined, AV_GET_BUFFER_FLAG_REF); + if (ret < 0) { + av_frame_free(&combined); + return ret; + } + frame->combined = combined; + } else { + if (!s->combine_leader || !s->combine_leader->combined) + return AVERROR_BUG; + combined = s->combine_leader->combined; + } + + f->format = combined->format; + f->width = l->sps->width; + f->height = l->sps->height; + for (i = 0; i < FF_ARRAY_ELEMS(f->buf) && combined->buf[i]; i++) { + f->buf[i] = av_buffer_ref(combined->buf[i]); + if (!f->buf[i]) + return AVERROR(ENOMEM); + } + for (i = 0; i < FF_ARRAY_ELEMS(f->data) && combined->data[i]; i++) { + f->data[i] = combined->data[i] + (bottom ? combined->linesize[i] : 0); + f->linesize[i] = combined->linesize[i] * 2; + } + f->extended_data = f->data; + + return 0; +} + static HEVCFrame *alloc_frame(HEVCContext *s, HEVCLayerContext *l) { const HEVCVPS *vps = l->sps->vps; @@ -157,9 +215,17 @@ static HEVCFrame *alloc_frame(HEVCContext *s, HEVCLayerContext *l) } } - ret = ff_thread_get_buffer(s->avctx, frame->f, AV_GET_BUFFER_FLAG_REF); - if (ret < 0) - goto fail; + if (s->combine_field_role != HEVC_COMBINE_NONE) { + ret = combine_setup_field_view(s, l, frame, + s->combine_field_role == HEVC_COMBINE_LEADER, + s->combine_field_bottom); + if (ret < 0) + goto fail; + } else { + ret = ff_thread_get_buffer(s->avctx, frame->f, AV_GET_BUFFER_FLAG_REF); + if (ret < 0) + goto fail; + } size_t rpl_bytes; if (av_size_mult(s->pkt.nb_nals, sizeof(*frame->rpl), &rpl_bytes) < 0) @@ -303,7 +369,8 @@ int ff_hevc_output_frames(HEVCContext *s, (nb_output && (nb_dpb[0] > max_dpb || nb_dpb[1] > max_dpb))) { HEVCFrame *frame = &s->layers[min_layer].DPB[min_idx]; - AVFrame *f = frame->needs_fg ? frame->frame_grain : frame->f; + AVFrame *f = frame->needs_fg ? frame->frame_grain : + frame->combined ? frame->combined : frame->f; int output = !discard && (layers_active_output & (1 << min_layer)); if (output) { diff --git a/tests/fate/hevc.mak b/tests/fate/hevc.mak index 7ce9ff403b..8045c99c8e 100644 --- a/tests/fate/hevc.mak +++ b/tests/fate/hevc.mak @@ -266,9 +266,12 @@ FATE_HEVC-$(call FRAMEMD5, HEVC, HEVC, HEVC_PARSER) += fate-hevc-skiploopfilter FATE_HEVC-$(call FRAMEMD5, MOV, HEVC, SCALE_FILTER) += fate-hevc-extradata-reload fate-hevc-extradata-reload: CMD = framemd5 -i $(TARGET_SAMPLES)/hevc/extradata-reload-multi-stsd.mov -sws_flags bitexact -fate-hevc-paired-fields: CMD = probeframes -show_entries frame=interlaced_frame,top_field_first $(TARGET_SAMPLES)/hevc/paired_fields.hevc +fate-hevc-paired-fields: CMD = probeframes -show_entries frame=width,height,interlaced_frame,top_field_first $(TARGET_SAMPLES)/hevc/paired_fields.hevc FATE_HEVC_FFPROBE-$(call DEMDEC, HEVC, HEVC) += fate-hevc-paired-fields +fate-hevc-paired-fields-420: CMD = probeframes -show_entries frame=width,height,pix_fmt,interlaced_frame,top_field_first $(TARGET_SAMPLES)/hevc/paired_fields_420.hevc +FATE_HEVC_FFPROBE-$(call DEMDEC, HEVC, HEVC) += fate-hevc-paired-fields-420 + fate-hevc-monochrome-crop: CMD = probeframes -show_entries frame=width,height:stream=width,height $(TARGET_SAMPLES)/hevc/hevc-monochrome.hevc FATE_HEVC_FFPROBE-$(call PARSERDEMDEC, HEVC, HEVC, HEVC) += fate-hevc-monochrome-crop diff --git a/tests/ref/fate/hevc-paired-fields b/tests/ref/fate/hevc-paired-fields index f2223e770b..ace522cc9b 100644 --- a/tests/ref/fate/hevc-paired-fields +++ b/tests/ref/fate/hevc-paired-fields @@ -1,16 +1,12 @@ [FRAME] +width=1920 +height=1080 interlaced_frame=1 top_field_first=1 [/FRAME] [FRAME] -interlaced_frame=1 -top_field_first=0 -[/FRAME] -[FRAME] +width=1920 +height=1080 interlaced_frame=1 top_field_first=1 [/FRAME] -[FRAME] -interlaced_frame=1 -top_field_first=0 -[/FRAME] diff --git a/tests/ref/fate/hevc-paired-fields-420 b/tests/ref/fate/hevc-paired-fields-420 new file mode 100644 index 0000000000..d667e26079 --- /dev/null +++ b/tests/ref/fate/hevc-paired-fields-420 @@ -0,0 +1,14 @@ +[FRAME] +width=1920 +height=1080 +pix_fmt=yuv420p +interlaced_frame=1 +top_field_first=1 +[/FRAME] +[FRAME] +width=1920 +height=1080 +pix_fmt=yuv420p +interlaced_frame=1 +top_field_first=1 +[/FRAME] -- 2.52.0 From 370460d5a79be5632f45eee51af28e5c4e4ca474 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Torbjo=CC=88rn=20Einarsson?= <torbjorn(a)einarssons.se> Date: Fri, 12 Jun 2026 00:04:59 +0200 Subject: [PATCH 4/4] avcodec/hevc: mark HEVC as field-based and report field-unit repeat_pict MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Add AV_CODEC_PROP_FIELDS to the HEVC codec descriptor and make the parser report repeat_pict in field periods minus one (0 field, 1 frame, 2 frame with a repeated field, 3 frame doubling, 5 frame tripling) and avctx->framerate as the frame rate: the HEVC VUI tick is per coded picture, which for field_seq streams is the field rate, so it is halved for field pictures. This mirrors the H.264 parser. libavformat then derives correct packet durations for field-coded streams and estimates e.g. avg_frame_rate 25/1 with r_frame_rate 50/1 for a 1080i50 stream in MPEG-TS, the same presentation as for interlaced H.264, instead of reporting the field rate as the frame rate. Signed-off-by: Torbjörn Einarsson <torbjorn(a)einarssons.se> --- libavcodec/codec_desc.c | 3 ++- libavcodec/hevc/parser.c | 41 ++++++++++++++++++++++++++++++++++++---- 2 files changed, 39 insertions(+), 5 deletions(-) diff --git a/libavcodec/codec_desc.c b/libavcodec/codec_desc.c index 81c095bea7..46c01ca14c 100644 --- a/libavcodec/codec_desc.c +++ b/libavcodec/codec_desc.c @@ -1268,7 +1268,8 @@ static const AVCodecDescriptor codec_descriptors[] = { .type = AVMEDIA_TYPE_VIDEO, .name = "hevc", .long_name = NULL_IF_CONFIG_SMALL("H.265 / HEVC (High Efficiency Video Coding)"), - .props = AV_CODEC_PROP_LOSSY | AV_CODEC_PROP_LOSSLESS | AV_CODEC_PROP_REORDER, + .props = AV_CODEC_PROP_LOSSY | AV_CODEC_PROP_LOSSLESS | + AV_CODEC_PROP_REORDER | AV_CODEC_PROP_FIELDS, .profiles = NULL_IF_CONFIG_SMALL(ff_hevc_profiles), }, { diff --git a/libavcodec/hevc/parser.c b/libavcodec/hevc/parser.c index 508d093807..51c605cb4b 100644 --- a/libavcodec/hevc/parser.c +++ b/libavcodec/hevc/parser.c @@ -87,6 +87,18 @@ static enum AVFieldOrder hevc_field_order(const HEVCSPS *sps, return AV_FIELD_UNKNOWN; } +/* whether the current picture is an individual field (H.265 Table D.2) */ +static int hevc_pic_struct_is_field(const HEVCSEIPictureTiming *pt) +{ + if (!pt->present) + return 0; + switch (pt->pic_struct) { + case 1: case 2: case 9: case 10: case 11: case 12: + return 1; + } + return 0; +} + static int hevc_parse_slice_header(AVCodecParserContext *s, H2645NAL *nal, AVCodecContext *avctx) { @@ -139,8 +151,12 @@ static int hevc_parse_slice_header(AVCodecParserContext *s, H2645NAL *nal, } if (num > 0 && den > 0) + /* Report the frame rate: the VUI tick is per coded picture, so for + * field pictures the tick rate is the field rate and one frame + * spans two ticks (AV_CODEC_PROP_FIELDS semantics, as for H.264). */ av_reduce(&avctx->framerate.den, &avctx->framerate.num, - num, den, 1 << 30); + num * (int64_t)(hevc_pic_struct_is_field(&sei->picture_timing) ? 2 : 1), + den, 1 << 30); if (!first_slice_in_pic_flag) { unsigned int slice_segment_addr; @@ -271,10 +287,27 @@ static int parse_nal_units(AVCodecParserContext *s, const uint8_t *buf, case HEVC_NAL_RADL_R: case HEVC_NAL_RASL_N: case HEVC_NAL_RASL_R: - if (ctx->sei.picture_timing.picture_struct == HEVC_SEI_PIC_STRUCT_FRAME_DOUBLING) { + /* repeat_pict is the picture's duration in field periods minus + * one (AV_CODEC_PROP_FIELDS semantics, as in the H.264 parser) */ + if (ctx->sei.picture_timing.present) { + switch (ctx->sei.picture_timing.pic_struct) { + case 1: case 2: case 9: case 10: case 11: case 12: + s->repeat_pict = 0; /* individual field */ + break; + case 5: case 6: + s->repeat_pict = 2; /* field pair with one field repeated */ + break; + case 7: + s->repeat_pict = 3; /* frame doubling */ + break; + case 8: + s->repeat_pict = 5; /* frame tripling */ + break; + default: + s->repeat_pict = 1; /* frame */ + } + } else { s->repeat_pict = 1; - } else if (ctx->sei.picture_timing.picture_struct == HEVC_SEI_PIC_STRUCT_FRAME_TRIPLING) { - s->repeat_pict = 2; } ret = hevc_parse_slice_header(s, nal, avctx); if (ret) -- 2.52.0

1 0

[PR] avfilter/vf_minterpolate_cuda: add CUDA-accelerated motion interpolation filter (PR #23615)
by stevenliu 27 Jun '26

27 Jun '26

PR #23615 opened by stevenliu URL: https://code.ffmpeg.org/FFmpeg/FFmpeg/pulls/23615 Patch URL: https://code.ffmpeg.org/FFmpeg/FFmpeg/pulls/23615.patch This adds a new GPU-accelerated minterpolate filter using CUDA, supporting frame rate conversion via motion-compensated interpolation (MCI), frame blending, and frame duplication modes. Architecture: - 11 motion estimation kernels: ESA (exhaustive), TSS (three-step), TDLS (2D log), NTSS (new three-step), FSS (four-step), DS (diamond), HEXBS (hexagon-based), EPZS (enhanced predictive zonal, 2-pass), UMH (uneven multi-hexagon, 2-pass) - Bilateral and bidirectional ME modes with overlapped block motion compensation (OBMC) and adaptive OBMC weighting - Variable-size block motion compensation (VSBM) with CPU-side iterative MV clustering and GPU sub-block refinement - Scene change detection via SAD-based frame difference on GPU - 4-frame ring buffer for interpolation between consecutive input pairs - Supports YUV420P, YUV422P, YUV444P, NV12, and GRAY8 pixel formats The ME cost function uses a 2xMB overlapped-block bilateral SAD matching the CPU get_sbad_ob implementation. The bidirectional OBMC uses a window-shift approach matching the CPU bidirectional_obmc, interpolating between frames using two independent motion vector fields with proper forward MV propagation across iterations. Performance (NVIDIA RTX 2060, 1280x720, 24fps to 60fps, EPZS): - CPU minterpolate: 0.16x realtime (12.0 seconds for 2 seconds of video) - CUDA minterpolate: 0.98x realtime (1.9 seconds for 2 seconds of video) - Speedup: approximately 6.3x faster on GPU Test results (raw YUV input, 640x360, 4fps to 12fps, 6 output frames): - BLEND mode: 100% bit-exact identical to CPU (6/6 frames, PASS) - DUP mode: 100% bit-exact identical to CPU (6/6 frames, PASS) - MCI BILAT ESA: 35.7 dB Y-PSNR, 2/6 frames bit-exact - MCI BILAT EPZS: 35.5 dB Y-PSNR, 2/6 frames bit-exact - MCI BILAT HEXBS: 35.8 dB Y-PSNR, 2/6 frames bit-exact - MCI BILAT UMH: 34.9 dB Y-PSNR, 2/6 frames bit-exact - MCI BIDIR EPZS: 33.4 dB Y-PSNR, 2/6 frames bit-exact - MCI AOBMC EPZS: 35.4 dB Y-PSNR, 2/6 frames bit-exact - ESA global motion: 100% bit-exact identical to CPU (3/3 frames, PASS) - hwupload/hwdownload roundtrip: 100% bit-exact (4/4 frames, PASS) Signed-off-by: Steven Liu <lq(a)chinaffmpeg.org> # Summary of changes Briefly describe what this PR does and why.  From 6c619efadb15e86fee272b33064b6ab4b0269e66 Mon Sep 17 00:00:00 2001 From: Steven Liu <lq(a)chinaffmpeg.org> Date: Sat, 27 Jun 2026 17:22:45 +0800 Subject: [PATCH] avfilter/vf_minterpolate_cuda: add CUDA-accelerated motion interpolation filter This adds a new GPU-accelerated minterpolate filter using CUDA, supporting frame rate conversion via motion-compensated interpolation (MCI), frame blending, and frame duplication modes. Architecture: - 11 motion estimation kernels: ESA (exhaustive), TSS (three-step), TDLS (2D log), NTSS (new three-step), FSS (four-step), DS (diamond), HEXBS (hexagon-based), EPZS (enhanced predictive zonal, 2-pass), UMH (uneven multi-hexagon, 2-pass) - Bilateral and bidirectional ME modes with overlapped block motion compensation (OBMC) and adaptive OBMC weighting - Variable-size block motion compensation (VSBM) with CPU-side iterative MV clustering and GPU sub-block refinement - Scene change detection via SAD-based frame difference on GPU - 4-frame ring buffer for interpolation between consecutive input pairs - Supports YUV420P, YUV422P, YUV444P, NV12, and GRAY8 pixel formats The ME cost function uses a 2xMB overlapped-block bilateral SAD matching the CPU get_sbad_ob implementation. The bidirectional OBMC uses a window-shift approach matching the CPU bidirectional_obmc, interpolating between frames using two independent motion vector fields with proper forward MV propagation across iterations. Performance (NVIDIA RTX 2060, 1280x720, 24fps to 60fps, EPZS): - CPU minterpolate: 0.16x realtime (12.0 seconds for 2 seconds of video) - CUDA minterpolate: 0.98x realtime (1.9 seconds for 2 seconds of video) - Speedup: approximately 6.3x faster on GPU Test results (raw YUV input, 640x360, 4fps to 12fps, 6 output frames): - BLEND mode: 100% bit-exact identical to CPU (6/6 frames, PASS) - DUP mode: 100% bit-exact identical to CPU (6/6 frames, PASS) - MCI BILAT ESA: 35.7 dB Y-PSNR, 2/6 frames bit-exact - MCI BILAT EPZS: 35.5 dB Y-PSNR, 2/6 frames bit-exact - MCI BILAT HEXBS: 35.8 dB Y-PSNR, 2/6 frames bit-exact - MCI BILAT UMH: 34.9 dB Y-PSNR, 2/6 frames bit-exact - MCI BIDIR EPZS: 33.4 dB Y-PSNR, 2/6 frames bit-exact - MCI AOBMC EPZS: 35.4 dB Y-PSNR, 2/6 frames bit-exact - ESA global motion: 100% bit-exact identical to CPU (3/3 frames, PASS) - hwupload/hwdownload roundtrip: 100% bit-exact (4/4 frames, PASS) Signed-off-by: Steven Liu <lq(a)chinaffmpeg.org> --- Changelog | 1 + configure | 2 + libavfilter/Makefile | 1 + libavfilter/allfilters.c | 1 + libavfilter/vf_minterpolate_cuda.c | 1313 ++++++++++++++++ libavfilter/vf_minterpolate_cuda.cu | 2168 +++++++++++++++++++++++++++ 6 files changed, 3486 insertions(+) create mode 100644 libavfilter/vf_minterpolate_cuda.c create mode 100644 libavfilter/vf_minterpolate_cuda.cu diff --git a/Changelog b/Changelog index 3268ab2cca..5055fef262 100644 --- a/Changelog +++ b/Changelog @@ -22,6 +22,7 @@ version 9.0: - Add AMF hardware memory mapping support. - ONNX Runtime DNN backend with GPU execution provider support - Remove deprecated NVENC options and support for pre-11.1 SDK versions +- minterpolate_cuda filter version 8.1: diff --git a/configure b/configure index 8941ced1f4..5d31096c86 100755 --- a/configure +++ b/configure @@ -4221,6 +4221,8 @@ movie_filter_deps="avcodec avformat" mpdecimate_filter_deps="gpl" mpdecimate_filter_select="pixelutils" minterpolate_filter_select="scene_sad" +minterpolate_cuda_filter_deps="ffnvcodec" +minterpolate_cuda_filter_deps_any="cuda_nvcc cuda_llvm" mptestsrc_filter_deps="gpl" msad_filter_select="scene_sad" negate_filter_deps="lut_filter" diff --git a/libavfilter/Makefile b/libavfilter/Makefile index cbae5f4ffd..d356a1a0e8 100644 --- a/libavfilter/Makefile +++ b/libavfilter/Makefile @@ -400,6 +400,7 @@ OBJS-$(CONFIG_MESTIMATE_D3D12_FILTER) += vf_mestimate_d3d12.o OBJS-$(CONFIG_METADATA_FILTER) += f_metadata.o OBJS-$(CONFIG_MIDEQUALIZER_FILTER) += vf_midequalizer.o framesync.o OBJS-$(CONFIG_MINTERPOLATE_FILTER) += vf_minterpolate.o motion_estimation.o +OBJS-$(CONFIG_MINTERPOLATE_CUDA_FILTER) += vf_minterpolate_cuda.o vf_minterpolate_cuda.ptx.o cuda/load_helper.o OBJS-$(CONFIG_MIX_FILTER) += vf_mix.o framesync.o OBJS-$(CONFIG_MONOCHROME_FILTER) += vf_monochrome.o OBJS-$(CONFIG_MORPHO_FILTER) += vf_morpho.o framesync.o diff --git a/libavfilter/allfilters.c b/libavfilter/allfilters.c index 402b843649..9d002b9fd3 100644 --- a/libavfilter/allfilters.c +++ b/libavfilter/allfilters.c @@ -374,6 +374,7 @@ extern const FFFilter ff_vf_mestimate_d3d12; extern const FFFilter ff_vf_metadata; extern const FFFilter ff_vf_midequalizer; extern const FFFilter ff_vf_minterpolate; +extern const FFFilter ff_vf_minterpolate_cuda; extern const FFFilter ff_vf_mix; extern const FFFilter ff_vf_monochrome; extern const FFFilter ff_vf_morpho; diff --git a/libavfilter/vf_minterpolate_cuda.c b/libavfilter/vf_minterpolate_cuda.c new file mode 100644 index 0000000000..b55cc71e76 --- /dev/null +++ b/libavfilter/vf_minterpolate_cuda.c @@ -0,0 +1,1313 @@ +/* + * CUDA-accelerated Motion Interpolation Filter + * Copyright (c) 2026 Steven Liu <lq(a)chinaffmpeg.org> + * + * This file is part of FFmpeg. + * + * FFmpeg is free software; you can redistribute it and/or + * modify it under the terms of the GNU Lesser General Public + * License as published by the Free Software Foundation; either + * version 2.1 of the License, or (at your option) any later version. + * + * FFmpeg is distributed in the hope that it will be useful, + * but WITHOUT ANY WARRANTY; without even the implied warranty of + * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU + * Lesser General Public License for more details. + * + * You should have received a copy of the GNU Lesser General Public + * License along with FFmpeg; if not, write to the Free Software + * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA + */ + +#include <stdint.h> +#include <stdio.h> +#include <math.h> + +#include "libavutil/avassert.h" +#include "libavutil/common.h" +#include "libavutil/hwcontext.h" +#include "libavutil/hwcontext_cuda_internal.h" +#include "libavutil/cuda_check.h" +#include "libavutil/internal.h" +#include "libavutil/mem.h" +#include "libavutil/opt.h" +#include "libavutil/pixdesc.h" + +#include "avfilter.h" +#include "filters.h" +#include "video.h" + +#include "cuda/load_helper.h" +#include "motion_estimation.h" /* AV_ME_METHOD_* enums */ + +#define NB_FRAMES 4 +#define ALPHA_MAX 1024 +#define BLOCKX 32 +#define BLOCKY 16 +#define ME_THREADS_PER_BLOCK 256 + +#define DIV_UP(a, b) (((a) + (b) - 1) / (b)) +#define CHECK_CU(x) FF_CUDA_CHECK_DL(ctx, s->hwctx->internal->cuda_dl, x) + +static const enum AVPixelFormat supported_formats[] = { + AV_PIX_FMT_YUV420P, + AV_PIX_FMT_YUV422P, + AV_PIX_FMT_YUV444P, + AV_PIX_FMT_NV12, + AV_PIX_FMT_GRAY8, + AV_PIX_FMT_NONE +}; + +enum MIMode { + MI_MODE_DUP = 0, + MI_MODE_BLEND = 1, + MI_MODE_MCI = 2, +}; + +enum MEMode { + ME_MODE_BIDIR = 0, + ME_MODE_BILAT = 1, +}; + +enum MCMode { + MC_MODE_OBMC = 0, + MC_MODE_AOBMC = 1, +}; + +typedef struct FrameEntry { + AVFrame *avf; + int tex_valid; + CUtexObject tex[4]; /* cached texture per plane */ +} FrameEntry; + +typedef struct CUDAMInterpolateContext { + const AVClass *class; + + AVCUDADeviceContext *hwctx; + AVBufferRef *device_ref; + AVBufferRef *input_frames_ref; + AVHWFramesContext *input_frames; + + const AVPixFmtDescriptor *desc; + int nb_planes; + int log2_chroma_w; + int log2_chroma_h; + int plane_channels[4]; + + AVRational frame_rate; + int mi_mode; + int me_mode; /* ME_MODE_BIDIR / ME_MODE_BILAT */ + int mc_mode; /* MC_MODE_OBMC / MC_MODE_AOBMC */ + int me_method; /* AV_ME_METHOD_ESA .. AV_ME_METHOD_UMH */ + int mb_size; + int search_param; + int vsbmc; /* 0/1 variable-size block MC (stub) */ + int scd_method; + double scd_threshold; + double prev_mafd; + + /* CUDA state */ + CUmodule cu_module; + + /* ME kernels */ + CUfunction cu_func_me_esa; + CUfunction cu_func_me_tss; + CUfunction cu_func_me_tdls; + CUfunction cu_func_me_ntss; + CUfunction cu_func_me_fss; + CUfunction cu_func_me_ds; + CUfunction cu_func_me_hexbs; + CUfunction cu_func_me_epzs_pass1; + CUfunction cu_func_me_epzs_pass2; + CUfunction cu_func_me_umh_pass1; + CUfunction cu_func_me_umh_pass2; + + /* OBMC / blend kernels */ + CUfunction cu_func_obmc; + CUfunction cu_func_obmc_uchar2; + CUfunction cu_func_obmc_bidir; + CUfunction cu_func_obmc_uchar2_bidir; + CUfunction cu_func_obmc_aobmc; + CUfunction cu_func_obmc_uchar2_aobmc; + CUfunction cu_func_blend; + CUfunction cu_func_blend_uchar2; + + /* AOBMC helper */ + CUfunction cu_func_sbad; + + /* VSBM */ + CUfunction cu_func_vsbm_sub_me; + + /* Scene detection */ + CUfunction cu_func_scene_sad; + CUfunction cu_func_scene_sad_reduce; + + CUstream cu_stream; + + /* Motion estimation parameters */ + int log2_mb_size; + int b_width; + int b_height; + int b_count; + + CUdeviceptr mv_buffer; /* b_count*2*int16_t — primary MV output */ + CUdeviceptr mv_buffer_pass1; /* b_count*2*int16_t — EPZS/UMH pass1 temp */ + CUdeviceptr mv_buffer_dir0; /* b_count*2*int16_t — BIDIR dir0 */ + CUdeviceptr mv_buffer_dir1; /* b_count*2*int16_t — BIDIR dir1 */ + CUdeviceptr mv_buffer_fwd; /* b_count*2*int16_t — BIDIR prev forward MV */ + CUdeviceptr mv_table_prev1; /* b_count*2*int16_t — EPZS temporal t-1 */ + CUdeviceptr mv_table_prev2; /* b_count*2*int16_t — EPZS temporal t-2 */ + + /* CUDA device memory — AOBMC */ + CUdeviceptr sbad_buffer; /* b_count*uint64_t */ + + /* CUDA device memory — VSBM */ + CUdeviceptr vsbm_cid; /* b_count*int — cluster IDs */ + CUdeviceptr sub_mv_buffer; /* b_count*4*2*int16_t — sub-block MVs */ + CUdeviceptr sub_valid; /* b_count*int — 1 if split */ + + /* CUDA device memory — scene detection */ + CUdeviceptr sad_buffer; /* per-block SAD partials */ + CUdeviceptr sad_total; /* single uint64_t for reduced total */ + + /* Frame ring buffer (4 frames) */ + FrameEntry frames[NB_FRAMES]; + + /* Output tracking */ + int64_t out_pts; + int scene_changed; +} CUDAMInterpolateContext; + +#define OFFSET(x) offsetof(CUDAMInterpolateContext, x) +#define FLAGS (AV_OPT_FLAG_VIDEO_PARAM | AV_OPT_FLAG_FILTERING_PARAM) +#define CONST(name, help, val, u) { name, help, 0, AV_OPT_TYPE_CONST, {.i64 = val}, 0, 0, FLAGS, .unit = u } + +static const AVOption minterpolate_cuda_options[] = { + { "fps", "output frame rate", OFFSET(frame_rate), AV_OPT_TYPE_VIDEO_RATE, {.str = "60"}, 0, INT_MAX, FLAGS }, + { "mi_mode", "motion interpolation mode", OFFSET(mi_mode), AV_OPT_TYPE_INT, {.i64 = MI_MODE_MCI}, MI_MODE_DUP, MI_MODE_MCI, FLAGS, .unit = "mi_mode" }, + CONST("dup", "duplicate frames", MI_MODE_DUP, "mi_mode"), + CONST("blend", "blend frames", MI_MODE_BLEND, "mi_mode"), + CONST("mci", "motion compensated interp.", MI_MODE_MCI, "mi_mode"), + { "mc_mode", "motion compensation mode", OFFSET(mc_mode), AV_OPT_TYPE_INT, {.i64 = MC_MODE_OBMC}, MC_MODE_OBMC, MC_MODE_AOBMC, FLAGS, .unit = "mc_mode" }, + CONST("obmc", "overlapped block MC", MC_MODE_OBMC, "mc_mode"), + CONST("aobmc", "adaptive overlapped block MC", MC_MODE_AOBMC, "mc_mode"), + { "me_mode", "motion estimation mode", OFFSET(me_mode), AV_OPT_TYPE_INT, {.i64 = ME_MODE_BILAT}, ME_MODE_BIDIR, ME_MODE_BILAT, FLAGS, .unit = "me_mode" }, + CONST("bidir", "bidirectional ME", ME_MODE_BIDIR, "me_mode"), + CONST("bilat", "bilateral ME", ME_MODE_BILAT, "me_mode"), + { "me", "motion estimation method", OFFSET(me_method), AV_OPT_TYPE_INT, {.i64 = AV_ME_METHOD_EPZS}, AV_ME_METHOD_ESA, AV_ME_METHOD_UMH, FLAGS, .unit = "me" }, + CONST("esa", "exhaustive search", AV_ME_METHOD_ESA, "me"), + CONST("tss", "three step search", AV_ME_METHOD_TSS, "me"), + CONST("tdls", "two dimensional log search", AV_ME_METHOD_TDLS, "me"), + CONST("ntss", "new three step search", AV_ME_METHOD_NTSS, "me"), + CONST("fss", "four step search", AV_ME_METHOD_FSS, "me"), + CONST("ds", "diamond search", AV_ME_METHOD_DS, "me"), + CONST("hexbs", "hexagon-based search", AV_ME_METHOD_HEXBS, "me"), + CONST("epzs", "enhanced predictive zonal", AV_ME_METHOD_EPZS, "me"), + CONST("umh", "uneven multi-hexagon", AV_ME_METHOD_UMH, "me"), + { "mb_size", "macroblock size", OFFSET(mb_size), AV_OPT_TYPE_INT, {.i64 = 16}, 4, 16, FLAGS }, + { "search_param", "search parameter", OFFSET(search_param), AV_OPT_TYPE_INT, {.i64 = 32}, 4, INT_MAX, FLAGS }, + { "vsbmc", "variable-size block MC", OFFSET(vsbmc), AV_OPT_TYPE_INT, {.i64 = 0}, 0, 1, FLAGS }, + { "scd", "scene change detection", OFFSET(scd_method), AV_OPT_TYPE_INT, {.i64 = 1}, 0, 1, FLAGS, .unit = "scene" }, + CONST("none", "disable detection", 0, "scene"), + CONST("fdiff", "frame difference", 1, "scene"), + { "scd_threshold", "scene change threshold", OFFSET(scd_threshold),AV_OPT_TYPE_DOUBLE, {.dbl = 10.0}, 0, 100.0, FLAGS }, + { NULL } +}; + +AVFILTER_DEFINE_CLASS(minterpolate_cuda); + + +static void destroy_frame_textures(AVFilterContext *ctx, + CUDAMInterpolateContext *s, FrameEntry *fe) +{ + CudaFunctions *cu = s->hwctx->internal->cuda_dl; + if (!fe->tex_valid) + return; + for (int i = 0; i < s->nb_planes; i++) { + if (fe->tex[i]) + CHECK_CU(cu->cuTexObjectDestroy(fe->tex[i])); + fe->tex[i] = 0; + } + fe->tex_valid = 0; +} + +static int create_frame_textures(AVFilterContext *ctx, FrameEntry *fe) +{ + CUDAMInterpolateContext *s = ctx->priv; + CudaFunctions *cu = s->hwctx->internal->cuda_dl; + int ret; + + for (int plane = 0; plane < s->nb_planes; plane++) { + int pw = !plane ? fe->avf->width + : AV_CEIL_RSHIFT(fe->avf->width, s->log2_chroma_w); + int ph = !plane ? fe->avf->height + : AV_CEIL_RSHIFT(fe->avf->height, s->log2_chroma_h); + int channels = s->plane_channels[plane]; + + CUDA_TEXTURE_DESC tex_desc = { + .filterMode = CU_TR_FILTER_MODE_POINT, + .flags = CU_TRSF_READ_AS_INTEGER, + }; + CUDA_RESOURCE_DESC res_desc = { + .resType = CU_RESOURCE_TYPE_PITCH2D, + .res.pitch2D.format = CU_AD_FORMAT_UNSIGNED_INT8, + .res.pitch2D.numChannels = channels, + .res.pitch2D.width = pw, + .res.pitch2D.height = ph, + .res.pitch2D.pitchInBytes = fe->avf->linesize[plane], + .res.pitch2D.devPtr = (CUdeviceptr)fe->avf->data[plane], + }; + + ret = CHECK_CU(cu->cuTexObjectCreate(&fe->tex[plane], + &res_desc, &tex_desc, NULL)); + if (ret < 0) { + for (int j = 0; j < plane; j++) { + CHECK_CU(cu->cuTexObjectDestroy(fe->tex[j])); + fe->tex[j] = 0; + } + return ret; + } + } + fe->tex_valid = 1; + return 0; +} + +static void free_buf(AVFilterContext *ctx, CUDAMInterpolateContext *s, + CUdeviceptr *ptr) +{ + CudaFunctions *cu = s->hwctx->internal->cuda_dl; + if (*ptr) { + CHECK_CU(cu->cuMemFree(*ptr)); + *ptr = 0; + } +} + +static av_cold void minterpolate_cuda_uninit(AVFilterContext *ctx) +{ + CUDAMInterpolateContext *s = ctx->priv; + + if (s->hwctx) { + CudaFunctions *cu = s->hwctx->internal->cuda_dl; + CHECK_CU(cu->cuCtxPushCurrent(s->hwctx->cuda_ctx)); + + for (int i = 0; i < NB_FRAMES; i++) + destroy_frame_textures(ctx, s, &s->frames[i]); + + if (s->cu_module) { + CHECK_CU(cu->cuModuleUnload(s->cu_module)); + s->cu_module = NULL; + } + + free_buf(ctx, s, &s->mv_buffer); + free_buf(ctx, s, &s->mv_buffer_pass1); + free_buf(ctx, s, &s->mv_buffer_dir0); + free_buf(ctx, s, &s->mv_buffer_dir1); + free_buf(ctx, s, &s->mv_buffer_fwd); + free_buf(ctx, s, &s->mv_table_prev1); + free_buf(ctx, s, &s->mv_table_prev2); + free_buf(ctx, s, &s->sbad_buffer); + free_buf(ctx, s, &s->vsbm_cid); + free_buf(ctx, s, &s->sub_mv_buffer); + free_buf(ctx, s, &s->sub_valid); + free_buf(ctx, s, &s->sad_buffer); + free_buf(ctx, s, &s->sad_total); + + { CUcontext pctx; CHECK_CU(cu->cuCtxPopCurrent(&pctx)); } + } + + for (int i = 0; i < NB_FRAMES; i++) + av_frame_free(&s->frames[i].avf); + + av_buffer_unref(&s->device_ref); + s->hwctx = NULL; + av_buffer_unref(&s->input_frames_ref); + s->input_frames = NULL; +} + +static int format_is_supported(enum AVPixelFormat fmt) +{ + for (int i = 0; supported_formats[i] != AV_PIX_FMT_NONE; i++) + if (supported_formats[i] == fmt) + return 1; + return 0; +} + +static av_cold int load_functions(AVFilterContext *ctx) +{ + CUDAMInterpolateContext *s = ctx->priv; + CUcontext dummy, cuda_ctx = s->hwctx->cuda_ctx; + CudaFunctions *cu = s->hwctx->internal->cuda_dl; + int ret; +#define LOAD(name) do { \ + ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_##name, s->cu_module, \ + "minterpolate_" #name)); \ + if (ret < 0) { \ + av_log(ctx, AV_LOG_ERROR, "Failed loading " #name "\n"); \ + goto fail; \ + } \ +} while (0) + + extern const unsigned char ff_vf_minterpolate_cuda_ptx_data[]; + extern const unsigned int ff_vf_minterpolate_cuda_ptx_len; + + ret = CHECK_CU(cu->cuCtxPushCurrent(cuda_ctx)); + if (ret < 0) + return ret; + + ret = ff_cuda_load_module(ctx, s->hwctx, &s->cu_module, + ff_vf_minterpolate_cuda_ptx_data, + ff_vf_minterpolate_cuda_ptx_len); + if (ret < 0) + goto fail; + + LOAD(me_esa); + LOAD(me_tss); + LOAD(me_tdls); + LOAD(me_ntss); + LOAD(me_fss); + LOAD(me_ds); + LOAD(me_hexbs); + LOAD(me_epzs_pass1); + LOAD(me_epzs_pass2); + LOAD(me_umh_pass1); + LOAD(me_umh_pass2); + LOAD(obmc); + LOAD(obmc_uchar2); + LOAD(obmc_bidir); + LOAD(obmc_uchar2_bidir); + LOAD(obmc_aobmc); + LOAD(obmc_uchar2_aobmc); + LOAD(blend); + LOAD(blend_uchar2); + LOAD(sbad); + LOAD(vsbm_sub_me); + LOAD(scene_sad); + LOAD(scene_sad_reduce); + +fail: + CHECK_CU(cu->cuCtxPopCurrent(&dummy)); + return ret; +#undef LOAD +} + +static int alloc_buf(AVFilterContext *ctx, CUDAMInterpolateContext *s, + CUdeviceptr *ptr, size_t size, const char *name) +{ + CudaFunctions *cu = s->hwctx->internal->cuda_dl; + int ret = CHECK_CU(cu->cuMemAlloc(ptr, size)); + if (ret < 0) + av_log(ctx, AV_LOG_ERROR, "Failed to allocate %s\n", name); + return ret; +} + +static int config_input(AVFilterLink *inlink) +{ + FilterLink *l = ff_filter_link(inlink); + AVFilterContext *ctx = inlink->dst; + CUDAMInterpolateContext *s = ctx->priv; + + if (!l->hw_frames_ctx) { + av_log(ctx, AV_LOG_ERROR, "A hardware frames reference is required.\n"); + return AVERROR(EINVAL); + } + + s->input_frames_ref = av_buffer_ref(l->hw_frames_ctx); + if (!s->input_frames_ref) + return AVERROR(ENOMEM); + + s->input_frames = (AVHWFramesContext *)s->input_frames_ref->data; + + return 0; +} + +static int config_output(AVFilterLink *outlink) +{ + FilterLink *l = ff_filter_link(outlink); + AVHWFramesContext *output_frames; + AVFilterContext *ctx = outlink->src; + AVFilterLink *inlink = ctx->inputs[0]; + CUDAMInterpolateContext *s = ctx->priv; + CudaFunctions *cu; + CUcontext dummy; + int ret = 0; + + av_assert0(s->input_frames); + s->device_ref = av_buffer_ref(s->input_frames->device_ref); + if (!s->device_ref) + return AVERROR(ENOMEM); + + s->hwctx = ((AVHWDeviceContext *)s->device_ref->data)->hwctx; + cu = s->hwctx->internal->cuda_dl; + s->cu_stream = s->hwctx->stream; + + if (!format_is_supported(s->input_frames->sw_format)) { + av_log(ctx, AV_LOG_ERROR, "Unsupported format: %s\n", + av_get_pix_fmt_name(s->input_frames->sw_format)); + return AVERROR(ENOSYS); + } + + s->desc = av_pix_fmt_desc_get(s->input_frames->sw_format); + s->nb_planes = av_pix_fmt_count_planes(s->input_frames->sw_format); + s->log2_chroma_w = s->desc->log2_chroma_w; + s->log2_chroma_h = s->desc->log2_chroma_h; + + memset(s->plane_channels, 0, sizeof(s->plane_channels)); + for (int i = 0; i < s->desc->nb_components; i++) { + int d = (s->desc->comp[i].depth + 7) / 8; + int p = s->desc->comp[i].plane; + s->plane_channels[p] = FFMAX(s->plane_channels[p], + s->desc->comp[i].step / d); + } + + s->log2_mb_size = av_ceil_log2(s->mb_size); + s->mb_size = 1 << s->log2_mb_size; + s->b_width = inlink->w >> s->log2_mb_size; + s->b_height = inlink->h >> s->log2_mb_size; + s->b_count = s->b_width * s->b_height; + + if (s->mi_mode == MI_MODE_MCI) { + if (s->b_width < 2 || s->b_height < 2) { + av_log(ctx, AV_LOG_ERROR, "Height or width < %d\n", 2 * s->mb_size); + return AVERROR(EINVAL); + } + } + + l->frame_rate = s->frame_rate; + outlink->time_base = av_inv_q(s->frame_rate); + + l->hw_frames_ctx = av_hwframe_ctx_alloc(s->device_ref); + if (!l->hw_frames_ctx) + return AVERROR(ENOMEM); + + output_frames = (AVHWFramesContext *)l->hw_frames_ctx->data; + output_frames->format = AV_PIX_FMT_CUDA; + output_frames->sw_format = s->input_frames->sw_format; + output_frames->width = inlink->w; + output_frames->height = inlink->h; + output_frames->initial_pool_size = 4; + + ret = ff_filter_init_hw_frames(ctx, outlink, 10); + if (ret < 0) + goto exit; + + ret = av_hwframe_ctx_init(l->hw_frames_ctx); + if (ret < 0) { + av_log(ctx, AV_LOG_ERROR, "hwframe init: %d\n", ret); + goto exit; + } + + ret = CHECK_CU(cu->cuCtxPushCurrent(s->hwctx->cuda_ctx)); + if (ret < 0) + goto exit; + + /* Allocate all GPU buffers */ + size_t mv_size = s->b_count * 2 * sizeof(int16_t); + if ((ret = alloc_buf(ctx, s, &s->mv_buffer, mv_size, "mv_buffer")) < 0) goto gpu_fail; + if ((ret = alloc_buf(ctx, s, &s->mv_buffer_pass1, mv_size, "mv_buffer_pass1")) < 0) goto gpu_fail; + if ((ret = alloc_buf(ctx, s, &s->mv_buffer_dir0, mv_size, "mv_buffer_dir0")) < 0) goto gpu_fail; + if ((ret = alloc_buf(ctx, s, &s->mv_buffer_dir1, mv_size, "mv_buffer_dir1")) < 0) goto gpu_fail; + if ((ret = alloc_buf(ctx, s, &s->mv_buffer_fwd, mv_size, "mv_buffer_fwd")) < 0) goto gpu_fail; + /* Initialize EPZS MV history to zero */ + if ((ret = alloc_buf(ctx, s, &s->mv_table_prev1, mv_size, "mv_table_prev1")) < 0) goto gpu_fail; + if ((ret = alloc_buf(ctx, s, &s->mv_table_prev2, mv_size, "mv_table_prev2")) < 0) goto gpu_fail; + /* Zero-initialize MV history */ + size_t zero_size = s->b_count * 2 * sizeof(int16_t); + void *zeros = av_mallocz(zero_size); + if (!zeros) { + ret = AVERROR(ENOMEM); + goto gpu_fail; + } + CHECK_CU(cu->cuMemcpyHtoD(s->mv_table_prev1, zeros, zero_size)); + CHECK_CU(cu->cuMemcpyHtoD(s->mv_table_prev2, zeros, zero_size)); + CHECK_CU(cu->cuMemcpyHtoD(s->mv_buffer_fwd, zeros, zero_size)); + av_free(zeros); + + if ((ret = alloc_buf(ctx, s, &s->sbad_buffer, s->b_count * sizeof(uint64_t), + "sbad_buffer")) < 0) goto gpu_fail; + + /* VSBM buffers */ + if ((ret = alloc_buf(ctx, s, &s->vsbm_cid, s->b_count * sizeof(int), + "vsbm_cid")) < 0) goto gpu_fail; + if ((ret = alloc_buf(ctx, s, &s->sub_mv_buffer, + s->b_count * 4 * 2 * sizeof(int16_t), + "sub_mv_buffer")) < 0) goto gpu_fail; + if ((ret = alloc_buf(ctx, s, &s->sub_valid, s->b_count * sizeof(int), + "sub_valid")) < 0) goto gpu_fail; + /* Initialize sub_valid to zero */ + int *z = av_mallocz(s->b_count * sizeof(int)); + if (!z) { + ret = AVERROR(ENOMEM); + goto gpu_fail; + } + CHECK_CU(cu->cuMemcpyHtoD(s->sub_valid, z, s->b_count * sizeof(int))); + av_free(z); + + /* Scene detection buffers */ + int sad_blocks_x = DIV_UP(inlink->w, BLOCKX); + int sad_blocks_y = DIV_UP(inlink->h, BLOCKY); + int sad_blocks = sad_blocks_x * sad_blocks_y; + if ((ret = alloc_buf(ctx, s, &s->sad_buffer, + sad_blocks * sizeof(uint64_t), "sad_buffer")) < 0) + goto gpu_fail; + if ((ret = alloc_buf(ctx, s, &s->sad_total, + sizeof(uint64_t), "sad_total")) < 0) + goto gpu_fail; + + ret = load_functions(ctx); + +gpu_fail: + CHECK_CU(cu->cuCtxPopCurrent(&dummy)); + if (ret < 0) { + /* Cleanup on error */ + free_buf(ctx, s, &s->mv_buffer); + free_buf(ctx, s, &s->mv_buffer_pass1); + free_buf(ctx, s, &s->mv_buffer_dir0); + free_buf(ctx, s, &s->mv_buffer_dir1); + free_buf(ctx, s, &s->mv_buffer_fwd); + free_buf(ctx, s, &s->mv_table_prev1); + free_buf(ctx, s, &s->mv_table_prev2); + free_buf(ctx, s, &s->sbad_buffer); + free_buf(ctx, s, &s->vsbm_cid); + free_buf(ctx, s, &s->sub_mv_buffer); + free_buf(ctx, s, &s->sub_valid); + free_buf(ctx, s, &s->sad_buffer); + free_buf(ctx, s, &s->sad_total); + } +exit: + if (ret < 0) { + av_buffer_unref(&l->hw_frames_ctx); + av_buffer_unref(&s->device_ref); + s->hwctx = NULL; + av_buffer_unref(&s->input_frames_ref); + s->input_frames = NULL; + } + return ret; +} + +static int launch_me_kernel(AVFilterContext *ctx, + FrameEntry *cur, FrameEntry *ref, + CUdeviceptr mv_out, + int pass) +{ + CUDAMInterpolateContext *s = ctx->priv; + CudaFunctions *cu = s->hwctx->internal->cuda_dl; + CUfunction func = NULL; + void *args[13]; + + if (!cur->tex_valid || !ref->tex_valid) + return AVERROR(EINVAL); + + int n = 0; + args[n++] = &cur->tex[0]; + args[n++] = &ref->tex[0]; + args[n++] = &mv_out; + + switch (s->me_method) { + case AV_ME_METHOD_ESA: + func = s->cu_func_me_esa; + break; + case AV_ME_METHOD_TSS: + func = s->cu_func_me_tss; + break; + case AV_ME_METHOD_TDLS: + func = s->cu_func_me_tdls; + break; + case AV_ME_METHOD_NTSS: + func = s->cu_func_me_ntss; + break; + case AV_ME_METHOD_FSS: + func = s->cu_func_me_fss; + break; + case AV_ME_METHOD_DS: + func = s->cu_func_me_ds; + break; + case AV_ME_METHOD_HEXBS: + func = s->cu_func_me_hexbs; + break; + + case AV_ME_METHOD_EPZS: + if (pass == 1) { + func = s->cu_func_me_epzs_pass1; + } else { + func = s->cu_func_me_epzs_pass2; + args[n++] = &s->mv_buffer_pass1; /* spatial source */ + } + args[n++] = &s->mv_table_prev1; + args[n++] = &s->mv_table_prev2; + break; + + case AV_ME_METHOD_UMH: + if (pass == 1) { + func = s->cu_func_me_umh_pass1; + } else { + func = s->cu_func_me_umh_pass2; + args[n++] = &s->mv_buffer_pass1; /* spatial source */ + } + args[n++] = &s->mv_table_prev1; + + break; + + default: + av_log(ctx, AV_LOG_ERROR, "Unknown ME method %d\n", s->me_method); + return AVERROR(EINVAL); + } + + args[n++] = &s->b_width; + args[n++] = &s->b_height; + args[n++] = &s->mb_size; + args[n++] = &s->search_param; + args[n++] = &cur->avf->width; + args[n++] = &cur->avf->height; + + return CHECK_CU(cu->cuLaunchKernel( + func, + s->b_count, 1, 1, + s->me_method == AV_ME_METHOD_ESA ? ME_THREADS_PER_BLOCK : 1, + 1, 1, + 0, s->cu_stream, args, NULL)); +} + +static int run_me_cycle(AVFilterContext *ctx, + FrameEntry *cur, FrameEntry *ref, + CUdeviceptr mv_out) +{ + CUDAMInterpolateContext *s = ctx->priv; + CudaFunctions *cu = s->hwctx->internal->cuda_dl; + int ret; + + if (s->me_method == AV_ME_METHOD_EPZS || + s->me_method == AV_ME_METHOD_UMH) { + /* 2-pass ME: pass1 → pass2 */ + ret = launch_me_kernel(ctx, cur, ref, s->mv_buffer_pass1, 1); + if (ret < 0) + return ret; + + ret = launch_me_kernel(ctx, cur, ref, mv_out, 2); + if (ret < 0) + return ret; + + /* Rotate MV history: prev2 ← prev1 ← current */ + CHECK_CU(cu->cuMemcpyDtoD(s->mv_table_prev2, s->mv_table_prev1, + s->b_count * 2 * sizeof(int16_t))); + CHECK_CU(cu->cuMemcpyDtoD(s->mv_table_prev1, mv_out, + s->b_count * 2 * sizeof(int16_t))); + } else { + /* Single-pass ME */ + ret = launch_me_kernel(ctx, cur, ref, mv_out, 0); + } + return ret; +} + +static int launch_obmc_plane(AVFilterContext *ctx, + FrameEntry *prev, FrameEntry *next, + AVFrame *dst, int plane, int alpha) +{ + CUDAMInterpolateContext *s = ctx->priv; + CudaFunctions *cu = s->hwctx->internal->cuda_dl; + int pw = !plane ? dst->width + : AV_CEIL_RSHIFT(dst->width, s->log2_chroma_w); + int ph = !plane ? dst->height + : AV_CEIL_RSHIFT(dst->height, s->log2_chroma_h); + int csw = !plane ? 0 : s->log2_chroma_w; + int csh = !plane ? 0 : s->log2_chroma_h; + int channels = s->plane_channels[plane]; + int dpitch = dst->linesize[plane] / ((channels > 1) ? 2 : 1); + CUdeviceptr dst_dev = (CUdeviceptr)dst->data[plane]; + + CUfunction func; + if (channels > 1) { + func = (s->mc_mode == MC_MODE_AOBMC) ? s->cu_func_obmc_uchar2_aobmc + : s->cu_func_obmc_uchar2; + } else { + func = (s->mc_mode == MC_MODE_AOBMC) ? s->cu_func_obmc_aobmc + : s->cu_func_obmc; + } + + void *args[17]; + int n = 0; + + args[n++] = &prev->tex[plane]; + args[n++] = &next->tex[plane]; + args[n++] = &s->mv_buffer; + args[n++] = &s->sub_mv_buffer; /* VSBM sub-block MVs */ + args[n++] = &s->sub_valid; /* VSBM split flags */ + + if (s->mc_mode == MC_MODE_AOBMC) + args[n++] = &s->sbad_buffer; + + args[n++] = &dst_dev; + args[n++] = &pw; + args[n++] = &ph; + args[n++] = &dpitch; + args[n++] = &s->b_width; + args[n++] = &s->b_height; + args[n++] = &s->mb_size; + args[n++] = &s->log2_mb_size; + args[n++] = α + args[n++] = &csw; + args[n++] = &csh; + + return CHECK_CU(cu->cuLaunchKernel( + func, + DIV_UP(pw, BLOCKX), DIV_UP(ph, BLOCKY), 1, + BLOCKX, BLOCKY, 1, + 0, s->cu_stream, args, NULL)); +} + +/* BIDIR OBMC — matches CPU bidirectional_obmc: window-shift, 2 frames */ +static int launch_obmc_bidir_plane(AVFilterContext *ctx, + FrameEntry *prev, /* frame[1] */ + FrameEntry *cur, /* frame[2] */ + AVFrame *dst, int plane, int alpha) +{ + CUDAMInterpolateContext *s = ctx->priv; + CudaFunctions *cu = s->hwctx->internal->cuda_dl; + int pw = !plane ? dst->width + : AV_CEIL_RSHIFT(dst->width, s->log2_chroma_w); + int ph = !plane ? dst->height + : AV_CEIL_RSHIFT(dst->height, s->log2_chroma_h); + int csw = !plane ? 0 : s->log2_chroma_w; + int csh = !plane ? 0 : s->log2_chroma_h; + int channels = s->plane_channels[plane]; + int dpitch = dst->linesize[plane] / ((channels > 1) ? 2 : 1); + CUdeviceptr dst_dev = (CUdeviceptr)dst->data[plane]; + + CUfunction func = (channels > 1) ? s->cu_func_obmc_uchar2_bidir + : s->cu_func_obmc_bidir; + + void *args[] = { + &prev->tex[plane], &cur->tex[plane], + &s->mv_buffer_dir0, &s->mv_buffer_fwd, + &dst_dev, + &pw, &ph, &dpitch, + &s->b_width, &s->b_height, + &s->mb_size, &s->log2_mb_size, + &alpha, + &csw, &csh, + }; + + return CHECK_CU(cu->cuLaunchKernel( + func, + DIV_UP(pw, BLOCKX), DIV_UP(ph, BLOCKY), 1, + BLOCKX, BLOCKY, 1, + 0, s->cu_stream, args, NULL)); +} + +/* Blend launch helper */ +static int launch_blend_plane(AVFilterContext *ctx, + FrameEntry *prev, FrameEntry *next, + AVFrame *dst, int plane, int alpha) +{ + CUDAMInterpolateContext *s = ctx->priv; + CudaFunctions *cu = s->hwctx->internal->cuda_dl; + int pw = !plane ? dst->width + : AV_CEIL_RSHIFT(dst->width, s->log2_chroma_w); + int ph = !plane ? dst->height + : AV_CEIL_RSHIFT(dst->height, s->log2_chroma_h); + int channels = s->plane_channels[plane]; + int dpitch = dst->linesize[plane] / ((channels > 1) ? 2 : 1); + CUdeviceptr dst_dev = (CUdeviceptr)dst->data[plane]; + CUfunction func = (channels > 1) ? s->cu_func_blend_uchar2 + : s->cu_func_blend; + + void *args[] = { + &prev->tex[plane], &next->tex[plane], + &dst_dev, + &pw, &ph, &dpitch, + &alpha, + }; + + return CHECK_CU(cu->cuLaunchKernel( + func, + DIV_UP(pw, BLOCKX), DIV_UP(ph, BLOCKY), 1, + BLOCKX, BLOCKY, 1, + 0, s->cu_stream, args, NULL)); +} + +static int launch_sbad_kernel(AVFilterContext *ctx, + FrameEntry *cur, FrameEntry *ref) +{ + CUDAMInterpolateContext *s = ctx->priv; + CudaFunctions *cu = s->hwctx->internal->cuda_dl; + + void *args[] = { + &cur->tex[0], &ref->tex[0], + &s->mv_buffer, + &s->sbad_buffer, + &s->b_width, &s->b_height, + &s->mb_size, &s->log2_mb_size, + &cur->avf->width, &cur->avf->height, + }; + + /* Launch with 256 threads per block for cooperative SAD */ + return CHECK_CU(cu->cuLaunchKernel( + s->cu_func_sbad, + s->b_count, 1, 1, + 256, 1, 1, + 0, s->cu_stream, args, NULL)); +} + +#define NB_CLUSTERS 128 +#define CLUSTER_THRESHOLD 4 + +typedef struct { + int64_t sum[2]; + int nb; +} VSBMCluster; + +static int launch_vsbm(AVFilterContext *ctx) +{ + CUDAMInterpolateContext *s = ctx->priv; + CudaFunctions *cu = s->hwctx->internal->cuda_dl; + int ret = 0; + + if (!s->vsbmc || s->me_mode != ME_MODE_BILAT) + return 0; + + /* 1. Copy MVs from GPU to CPU for clustering */ + int16_t *mv_host = av_malloc(s->b_count * 2 * sizeof(int16_t)); + int *cid_host = av_malloc(s->b_count * sizeof(int)); + if (!mv_host || !cid_host) { + av_free(mv_host); av_free(cid_host); + return AVERROR(ENOMEM); + } + + ret = CHECK_CU(cu->cuMemcpyDtoH(mv_host, s->mv_buffer, + s->b_count * 2 * sizeof(int16_t))); + if (ret < 0) { + av_free(mv_host); + av_free(cid_host); + return ret; + } + + /* 2. CPU-side iterative MV clustering */ + VSBMCluster clusters[NB_CLUSTERS]; + memset(clusters, 0, sizeof(clusters)); + + /* Initialize: all blocks in cluster 0 */ + for (int i = 0; i < s->b_count; i++) { + clusters[0].sum[0] += mv_host[i * 2]; + clusters[0].sum[1] += mv_host[i * 2 + 1]; + cid_host[i] = 0; + } + clusters[0].nb = s->b_count; + int c_max = 0; + + /* Iterate until stable */ + int changed; + do { + changed = 0; + for (int mb_y = 0; mb_y < s->b_height; mb_y++) { + for (int mb_x = 0; mb_x < s->b_width; mb_x++) { + int mb_i = mb_x + mb_y * s->b_width; + int c = cid_host[mb_i]; + VSBMCluster *cl = &clusters[c]; + int mv_x = mv_host[mb_i * 2]; + int mv_y = mv_host[mb_i * 2 + 1]; + + if (cl->nb < 2) + continue; + + int avg_x = (int)(cl->sum[0] / cl->nb); + int avg_y = (int)(cl->sum[1] / cl->nb); + int dx = avg_x - mv_x; + int dy = avg_y - mv_y; + + if (abs(dx) > CLUSTER_THRESHOLD || + abs(dy) > CLUSTER_THRESHOLD) { + /* Find nearest neighboring cluster */ + int new_c = c; + for (int d = 1; d < 5; d++) { + int y0 = FFMAX(mb_y - d, 0); + int y1 = FFMIN(mb_y + d + 1, s->b_height); + int x0 = FFMAX(mb_x - d, 0); + int x1 = FFMIN(mb_x + d + 1, s->b_width); + for (int y = y0; y < y1; y++) { + for (int x = x0; x < x1; x++) { + int nb_c = cid_host[x + y * s->b_width]; + if (nb_c > c && + (nb_c < new_c || new_c == c)) + new_c = nb_c; + } + } + } + + if (new_c == c) + new_c = c_max + 1; + if (new_c >= NB_CLUSTERS) + continue; + + VSBMCluster *nc = &clusters[new_c]; + nc->sum[0] += mv_x; + nc->sum[1] += mv_y; + nc->nb++; + cl->sum[0] -= mv_x; + cl->sum[1] -= mv_y; + cl->nb--; + + c_max = FFMAX(c_max, new_c); + cid_host[mb_i] = new_c; + changed = 1; + } + } + } + } while (changed); + + /* 3. Upload cluster IDs to GPU */ + ret = CHECK_CU(cu->cuMemcpyHtoD(s->vsbm_cid, cid_host, + s->b_count * sizeof(int))); + if (ret < 0) { + av_free(mv_host); + av_free(cid_host); + return ret; + } + + /* 4. Launch sub-block ME kernel for boundary blocks */ + { + void *args[] = { + &s->frames[1].tex[0], &s->frames[2].tex[0], + &s->mv_buffer, + &s->vsbm_cid, + &s->sub_mv_buffer, + &s->sub_valid, + &s->b_width, &s->b_height, + &s->mb_size, &s->log2_mb_size, + &s->frames[1].avf->width, &s->frames[1].avf->height, + }; + + ret = CHECK_CU(cu->cuLaunchKernel( + s->cu_func_vsbm_sub_me, + s->b_count, 1, 1, + 1, 1, 1, + 0, s->cu_stream, args, NULL)); + } + + av_free(mv_host); + av_free(cid_host); + return ret; +} + +static int detect_scene_change_gpu(AVFilterContext *ctx) +{ + CUDAMInterpolateContext *s = ctx->priv; + AVFilterLink *inlink = ctx->inputs[0]; + CudaFunctions *cu = s->hwctx->internal->cuda_dl; + int sad_blocks_x = DIV_UP(inlink->w, BLOCKX); + int sad_blocks_y = DIV_UP(inlink->h, BLOCKY); + int sad_blocks = sad_blocks_x * sad_blocks_y; + uint64_t total_sad = 0; + double mafd, diff; + int ret; + CUcontext dummy; + + if (!s->scd_method) + return 0; + if (!s->frames[1].tex_valid || !s->frames[2].tex_valid) + return 0; + + ret = CHECK_CU(cu->cuCtxPushCurrent(s->hwctx->cuda_ctx)); + if (ret < 0) + return 0; + + /* Stage 1 */ + { + void *args[] = { + &s->frames[1].tex[0], &s->frames[2].tex[0], + &inlink->w, &inlink->h, + &s->sad_buffer, + }; + ret = CHECK_CU(cu->cuLaunchKernel(s->cu_func_scene_sad, + sad_blocks_x, sad_blocks_y, 1, BLOCKX, BLOCKY, 1, + 0, s->cu_stream, args, NULL)); + if (ret < 0) + goto exit; + } + + /* Stage 2 */ + { + void *args[] = { &s->sad_buffer, &sad_blocks, &s->sad_total }; + ret = CHECK_CU(cu->cuLaunchKernel(s->cu_func_scene_sad_reduce, + 1, 1, 1, 256, 1, 1, 0, s->cu_stream, args, NULL)); + if (ret < 0) + goto exit; + } + + ret = CHECK_CU(cu->cuMemcpyDtoH(&total_sad, s->sad_total, sizeof(uint64_t))); + if (ret < 0) + goto exit; + + mafd = (double)total_sad * 100.0 / (inlink->h * inlink->w) / 255.0; + diff = fabs(mafd - s->prev_mafd); + s->prev_mafd = mafd; + + CHECK_CU(cu->cuCtxPopCurrent(&dummy)); + return FFMIN(mafd, diff) >= s->scd_threshold ? 1 : 0; + +exit: + CHECK_CU(cu->cuCtxPopCurrent(&dummy)); + return 0; +} + +static int inject_frame(AVFilterLink *inlink, AVFrame *avf_in) +{ + AVFilterContext *ctx = inlink->dst; + CUDAMInterpolateContext *s = ctx->priv; + CudaFunctions *cu = s->hwctx->internal->cuda_dl; + FrameEntry frame_tmp; + int ret; + + ret = CHECK_CU(cu->cuCtxPushCurrent(s->hwctx->cuda_ctx)); + if (ret < 0) + return ret; + + destroy_frame_textures(ctx, s, &s->frames[0]); + + av_frame_free(&s->frames[0].avf); + frame_tmp = s->frames[0]; + memmove(&s->frames[0], &s->frames[1], + sizeof(s->frames[0]) * (NB_FRAMES - 1)); + s->frames[NB_FRAMES - 1] = frame_tmp; + s->frames[NB_FRAMES - 1].avf = avf_in; + s->frames[NB_FRAMES - 1].tex_valid = 0; + memset(s->frames[NB_FRAMES - 1].tex, 0, sizeof(s->frames[NB_FRAMES - 1].tex)); + + ret = create_frame_textures(ctx, &s->frames[NB_FRAMES - 1]); + if (ret < 0) { + /* Don't own the frame; let the caller free it */ + s->frames[NB_FRAMES - 1].avf = NULL; + goto exit_inject; + } + + /* Run ME if we have enough frames */ + if (s->mi_mode == MI_MODE_MCI && + s->frames[1].tex_valid && s->frames[2].tex_valid) { + + if (s->me_mode == ME_MODE_BIDIR && s->frames[3].tex_valid) { + /* BIDIR: two ME searches from frame[2] */ + /* dir0: frame[2]->frame[1] */ + ret = run_me_cycle(ctx, &s->frames[2], &s->frames[1], s->mv_buffer_dir0); + /* dir1: frame[2]->frame[3] */ + if (ret >= 0) + ret = run_me_cycle(ctx, &s->frames[2], &s->frames[3], + s->mv_buffer_dir1); + } else { + /* BILAT: from frame[1] to frame[2] */ + ret = run_me_cycle(ctx, &s->frames[1], &s->frames[2], s->mv_buffer); + /* VSBM: MV clustering + sub-block refinement */ + if (ret >= 0 && s->vsbmc) + ret = launch_vsbm(ctx); + } + } + +exit_inject: + CUcontext pctx; + CHECK_CU(cu->cuCtxPopCurrent(&pctx)); + + return ret; +} + +static int interpolate(AVFilterLink *inlink, AVFrame *avf_out) +{ + AVFilterContext *ctx = inlink->dst; + AVFilterLink *outlink = ctx->outputs[0]; + CUDAMInterpolateContext *s = ctx->priv; + CudaFunctions *cu = s->hwctx->internal->cuda_dl; + int plane, ret; + int64_t pts; + int alpha = 0; + + pts = av_rescale(avf_out->pts, + (int64_t)ALPHA_MAX * outlink->time_base.num * inlink->time_base.den, + (int64_t)outlink->time_base.den * inlink->time_base.num); + + if (s->frames[2].avf->pts > s->frames[1].avf->pts) { + alpha = (int)((pts - s->frames[1].avf->pts * ALPHA_MAX) / + (s->frames[2].avf->pts - s->frames[1].avf->pts)); + alpha = av_clip(alpha, 0, ALPHA_MAX); + } + + if (!alpha || alpha == ALPHA_MAX) { + av_frame_copy(avf_out, alpha ? s->frames[2].avf : s->frames[1].avf); + return 0; + } + + if (s->scene_changed) { + av_frame_copy(avf_out, + alpha > ALPHA_MAX / 2 ? s->frames[2].avf : s->frames[1].avf); + return 0; + } + + ret = CHECK_CU(cu->cuCtxPushCurrent(s->hwctx->cuda_ctx)); + if (ret < 0) + return ret; + + switch (s->mi_mode) { + case MI_MODE_DUP: + av_frame_copy(avf_out, + alpha > ALPHA_MAX / 2 ? s->frames[2].avf : s->frames[1].avf); + break; + + case MI_MODE_BLEND: + for (plane = 0; plane < s->nb_planes; plane++) { + ret = launch_blend_plane(ctx, &s->frames[1], &s->frames[2], + avf_out, plane, alpha); + if (ret < 0) + goto exit; + } + break; + + case MI_MODE_MCI: + if (s->me_mode == ME_MODE_BIDIR) { + /* BIDIR OBMC: 2 frames, window-shift interpolation */ + for (plane = 0; plane < s->nb_planes; plane++) { + ret = launch_obmc_bidir_plane(ctx, + &s->frames[1], /* prev */ + &s->frames[2], /* cur */ + avf_out, plane, alpha); + if (ret < 0) + goto exit; + } + } else { + /* BILAT OBMC */ + if (s->mc_mode == MC_MODE_AOBMC) { + /* Compute sbad first (once, not per plane) */ + ret = launch_sbad_kernel(ctx, &s->frames[1], &s->frames[2]); + if (ret < 0) + goto exit; + } + for (plane = 0; plane < s->nb_planes; plane++) { + ret = launch_obmc_plane(ctx, &s->frames[1], &s->frames[2], + avf_out, plane, alpha); + if (ret < 0) + goto exit; + } + } + break; + } + +exit: + CUcontext pctx; + CHECK_CU(cu->cuCtxPopCurrent(&pctx)); + + return ret; +} + +static int filter_frame(AVFilterLink *inlink, AVFrame *avf_in) +{ + AVFilterContext *ctx = inlink->dst; + AVFilterLink *outlink = ctx->outputs[0]; + CUDAMInterpolateContext *s = ctx->priv; + int ret; + + if (avf_in->pts == AV_NOPTS_VALUE) + return ff_filter_frame(outlink, avf_in); + + if (!s->frames[NB_FRAMES - 1].avf || + avf_in->pts < s->frames[NB_FRAMES - 1].avf->pts) { + av_log(ctx, AV_LOG_VERBOSE, + "Initializing out pts from input pts %"PRId64"\n", avf_in->pts); + s->out_pts = av_rescale_q(avf_in->pts, inlink->time_base, + outlink->time_base); + } + + if (!s->frames[NB_FRAMES - 1].avf) { + AVFrame *clone = av_frame_clone(avf_in); + if (!clone) { + ret = AVERROR(ENOMEM); + goto fail; + } + ret = inject_frame(inlink, clone); + if (ret < 0) { + av_frame_free(&clone); + goto fail; + } + } + + ret = inject_frame(inlink, avf_in); + if (ret < 0) + goto fail; + avf_in = NULL; /* ring buffer now owns the frame */ + + if (!s->frames[0].avf) + return 0; + + s->scene_changed = detect_scene_change_gpu(ctx); + + for (;;) { + AVFrame *avf_out; + + if (av_compare_ts(s->out_pts, outlink->time_base, + s->frames[2].avf->pts, inlink->time_base) > 0) + break; + + avf_out = ff_get_video_buffer(outlink, inlink->w, inlink->h); + if (!avf_out) { + ret = AVERROR(ENOMEM); + goto fail; + } + + ret = av_frame_copy_props(avf_out, s->frames[NB_FRAMES - 1].avf); + if (ret < 0) { + av_frame_free(&avf_out); + goto fail; + } + + avf_out->pts = s->out_pts++; + avf_out->duration = 1; + + ret = interpolate(inlink, avf_out); + if (ret < 0) { + av_frame_free(&avf_out); + goto fail; + } + + ret = ff_filter_frame(outlink, avf_out); + if (ret < 0) + goto fail; + } + + /* BIDIR: save dir1 as fwd MV for next iteration (after ring shift, + * mv_buffer_dir1 becomes frame[1]→frame[2]) */ + if (s->me_mode == ME_MODE_BIDIR && s->frames[1].tex_valid && + s->frames[2].tex_valid && s->frames[3].tex_valid) { + CudaFunctions *cuf = s->hwctx->internal->cuda_dl; + CHECK_CU(cuf->cuMemcpyDtoD(s->mv_buffer_fwd, + s->mv_buffer_dir1, + s->b_count * 2 * sizeof(int16_t))); + } + + return 0; + +fail: + av_frame_free(&avf_in); + return ret; +} + +static const AVFilterPad minterpolate_cuda_inputs[] = { + { + .name = "default", + .type = AVMEDIA_TYPE_VIDEO, + .filter_frame = filter_frame, + .config_props = config_input, + }, +}; + +static const AVFilterPad minterpolate_cuda_outputs[] = { + { + .name = "default", + .type = AVMEDIA_TYPE_VIDEO, + .config_props = config_output, + }, +}; + +const FFFilter ff_vf_minterpolate_cuda = { + .p.name = "minterpolate_cuda", + .p.description = NULL_IF_CONFIG_SMALL("GPU-accelerated motion interpolation."), + .p.priv_class = &minterpolate_cuda_class, + .priv_size = sizeof(CUDAMInterpolateContext), + .uninit = minterpolate_cuda_uninit, + FILTER_INPUTS(minterpolate_cuda_inputs), + FILTER_OUTPUTS(minterpolate_cuda_outputs), + FILTER_SINGLE_PIXFMT(AV_PIX_FMT_CUDA), + .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE, +}; diff --git a/libavfilter/vf_minterpolate_cuda.cu b/libavfilter/vf_minterpolate_cuda.cu new file mode 100644 index 0000000000..6f914d2ac2 --- /dev/null +++ b/libavfilter/vf_minterpolate_cuda.cu @@ -0,0 +1,2168 @@ +/* + * GPU kernels for CUDA-accelerated Motion Interpolation Filter + * Copyright (c) 2026 + * + * This file is part of FFmpeg. + * + * FFmpeg is free software; you can redistribute it and/or + * modify it under the terms of the GNU Lesser General Public + * License as published by the Free Software Foundation; either + * version 2.1 of the License, or (at your option) any later version. + * + * FFmpeg is distributed in the hope that it will be useful, + * but WITHOUT ANY WARRANTY; without even the implied warranty of + * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU + * Lesser General Public License for more details. + */ + +typedef unsigned char uchar; +typedef unsigned long long uint64_t; +typedef unsigned short ushort; + +#define ALPHA_MAX 1024 +#define COST_PRED_SCALE 64 +#define MAX_MB_SIZE 16 + +/* =================================================================== + * OBMC weight tables — shared with CPU vf_minterpolate.c + * =================================================================== */ +__constant__ uchar d_obmc_linear32[1024] = { + 0, 0, 0, 0, 4, 4, 4, 4, 4, 4, 4, 4, 8, 8, 8, 8, 8, 8, 8, 8, 4, 4, 4, 4, 4, 4, 4, 4, 0, 0, 0, 0, + 0, 4, 4, 4, 8, 8, 8, 12, 12, 16, 16, 16, 20, 20, 20, 24, 24, 20, 20, 20, 16, 16, 16, 12, 12, 8, 8, 8, 4, 4, 4, 0, + 0, 4, 8, 8, 12, 12, 16, 20, 20, 24, 28, 28, 32, 32, 36, 40, 40, 36, 32, 32, 28, 28, 24, 20, 20, 16, 12, 12, 8, 8, 4, 0, + 0, 4, 8, 12, 16, 20, 24, 28, 28, 32, 36, 40, 44, 48, 52, 56, 56, 52, 48, 44, 40, 36, 32, 28, 28, 24, 20, 16, 12, 8, 4, 0, + 4, 8, 12, 16, 20, 24, 28, 32, 40, 44, 48, 52, 56, 60, 64, 68, 68, 64, 60, 56, 52, 48, 44, 40, 32, 28, 24, 20, 16, 12, 8, 4, + 4, 8, 12, 20, 24, 32, 36, 40, 48, 52, 56, 64, 68, 76, 80, 84, 84, 80, 76, 68, 64, 56, 52, 48, 40, 36, 32, 24, 20, 12, 8, 4, + 4, 8, 16, 24, 28, 36, 44, 48, 56, 60, 68, 76, 80, 88, 96,100,100, 96, 88, 80, 76, 68, 60, 56, 48, 44, 36, 28, 24, 16, 8, 4, + 4, 12, 20, 28, 32, 40, 48, 56, 64, 72, 80, 88, 92,100,108,116,116,108,100, 92, 88, 80, 72, 64, 56, 48, 40, 32, 28, 20, 12, 4, + 4, 12, 20, 28, 40, 48, 56, 64, 72, 80, 88, 96,108,116,124,132,132,124,116,108, 96, 88, 80, 72, 64, 56, 48, 40, 28, 20, 12, 4, + 4, 16, 24, 32, 44, 52, 60, 72, 80, 92,100,108,120,128,136,148,148,136,128,120,108,100, 92, 80, 72, 60, 52, 44, 32, 24, 16, 4, + 4, 16, 28, 36, 48, 56, 68, 80, 88,100,112,120,132,140,152,164,164,152,140,132,120,112,100, 88, 80, 68, 56, 48, 36, 28, 16, 4, + 4, 16, 28, 40, 52, 64, 76, 88, 96,108,120,132,144,156,168,180,180,168,156,144,132,120,108, 96, 88, 76, 64, 52, 40, 28, 16, 4, + 8, 20, 32, 44, 56, 68, 80, 92,108,120,132,144,156,168,180,192,192,180,168,156,144,132,120,108, 92, 80, 68, 56, 44, 32, 20, 8, + 8, 20, 32, 48, 60, 76, 88,100,116,128,140,156,168,184,196,208,208,196,184,168,156,140,128,116,100, 88, 76, 60, 48, 32, 20, 8, + 8, 20, 36, 52, 64, 80, 96,108,124,136,152,168,180,196,212,224,224,212,196,180,168,152,136,124,108, 96, 80, 64, 52, 36, 20, 8, + 8, 24, 40, 56, 68, 84,100,116,132,148,164,180,192,208,224,240,240,224,208,192,180,164,148,132,116,100, 84, 68, 56, 40, 24, 8, + 8, 24, 40, 56, 68, 84,100,116,132,148,164,180,192,208,224,240,240,224,208,192,180,164,148,132,116,100, 84, 68, 56, 40, 24, 8, + 8, 20, 36, 52, 64, 80, 96,108,124,136,152,168,180,196,212,224,224,212,196,180,168,152,136,124,108, 96, 80, 64, 52, 36, 20, 8, + 8, 20, 32, 48, 60, 76, 88,100,116,128,140,156,168,184,196,208,208,196,184,168,156,140,128,116,100, 88, 76, 60, 48, 32, 20, 8, + 8, 20, 32, 44, 56, 68, 80, 92,108,120,132,144,156,168,180,192,192,180,168,156,144,132,120,108, 92, 80, 68, 56, 44, 32, 20, 8, + 4, 16, 28, 40, 52, 64, 76, 88, 96,108,120,132,144,156,168,180,180,168,156,144,132,120,108, 96, 88, 76, 64, 52, 40, 28, 16, 4, + 4, 16, 28, 36, 48, 56, 68, 80, 88,100,112,120,132,140,152,164,164,152,140,132,120,112,100, 88, 80, 68, 56, 48, 36, 28, 16, 4, + 4, 16, 24, 32, 44, 52, 60, 72, 80, 92,100,108,120,128,136,148,148,136,128,120,108,100, 92, 80, 72, 60, 52, 44, 32, 24, 16, 4, + 4, 12, 20, 28, 40, 48, 56, 64, 72, 80, 88, 96,108,116,124,132,132,124,116,108, 96, 88, 80, 72, 64, 56, 48, 40, 28, 20, 12, 4, + 4, 12, 20, 28, 32, 40, 48, 56, 64, 72, 80, 88, 92,100,108,116,116,108,100, 92, 88, 80, 72, 64, 56, 48, 40, 32, 28, 20, 12, 4, + 4, 8, 16, 24, 28, 36, 44, 48, 56, 60, 68, 76, 80, 88, 96,100,100, 96, 88, 80, 76, 68, 60, 56, 48, 44, 36, 28, 24, 16, 8, 4, + 4, 8, 12, 20, 24, 32, 36, 40, 48, 52, 56, 64, 68, 76, 80, 84, 84, 80, 76, 68, 64, 56, 52, 48, 40, 36, 32, 24, 20, 12, 8, 4, + 4, 8, 12, 16, 20, 24, 28, 32, 40, 44, 48, 52, 56, 60, 64, 68, 68, 64, 60, 56, 52, 48, 44, 40, 32, 28, 24, 20, 16, 12, 8, 4, + 0, 4, 8, 12, 16, 20, 24, 28, 28, 32, 36, 40, 44, 48, 52, 56, 56, 52, 48, 44, 40, 36, 32, 28, 28, 24, 20, 16, 12, 8, 4, 0, + 0, 4, 8, 8, 12, 12, 16, 20, 20, 24, 28, 28, 32, 32, 36, 40, 40, 36, 32, 32, 28, 28, 24, 20, 20, 16, 12, 12, 8, 8, 4, 0, + 0, 4, 4, 4, 8, 8, 8, 12, 12, 16, 16, 16, 20, 20, 20, 24, 24, 20, 20, 20, 16, 16, 16, 12, 12, 8, 8, 8, 4, 4, 4, 0, + 0, 0, 0, 0, 4, 4, 4, 4, 4, 4, 4, 4, 8, 8, 8, 8, 8, 8, 8, 8, 4, 4, 4, 4, 4, 4, 4, 4, 0, 0, 0, 0, +}; + +__constant__ uchar d_obmc_linear16[256] = { + 0, 4, 4, 8, 8, 12, 12, 16, 16, 12, 12, 8, 8, 4, 4, 0, + 4, 8, 16, 20, 28, 32, 40, 44, 44, 40, 32, 28, 20, 16, 8, 4, + 4, 16, 24, 36, 44, 56, 64, 76, 76, 64, 56, 44, 36, 24, 16, 4, + 8, 20, 36, 48, 64, 76, 92,104,104, 92, 76, 64, 48, 36, 20, 8, + 8, 28, 44, 64, 80,100,116,136,136,116,100, 80, 64, 44, 28, 8, + 12, 32, 56, 76,100,120,144,164,164,144,120,100, 76, 56, 32, 12, + 12, 40, 64, 92,116,144,168,196,196,168,144,116, 92, 64, 40, 12, + 16, 44, 76,104,136,164,196,224,224,196,164,136,104, 76, 44, 16, + 16, 44, 76,104,136,164,196,224,224,196,164,136,104, 76, 44, 16, + 12, 40, 64, 92,116,144,168,196,196,168,144,116, 92, 64, 40, 12, + 12, 32, 56, 76,100,120,144,164,164,144,120,100, 76, 56, 32, 12, + 8, 28, 44, 64, 80,100,116,136,136,116,100, 80, 64, 44, 28, 8, + 8, 20, 36, 48, 64, 76, 92,104,104, 92, 76, 64, 48, 36, 20, 8, + 4, 16, 24, 36, 44, 56, 64, 76, 76, 64, 56, 44, 36, 24, 16, 4, + 4, 8, 16, 20, 28, 32, 40, 44, 44, 40, 32, 28, 20, 16, 8, 4, + 0, 4, 4, 8, 8, 12, 12, 16, 16, 12, 12, 8, 8, 4, 4, 0, +}; + +__constant__ uchar d_obmc_linear8[64] = { + 4, 12, 20, 28, 28, 20, 12, 4, + 12, 36, 60, 84, 84, 60, 36, 12, + 20, 60,100,140,140,100, 60, 20, + 28, 84,140,196,196,140, 84, 28, + 28, 84,140,196,196,140, 84, 28, + 20, 60,100,140,140,100, 60, 20, + 12, 36, 60, 84, 84, 60, 36, 12, + 4, 12, 20, 28, 28, 20, 12, 4, +}; + +__constant__ uchar d_obmc_linear4[16] = { + 16, 48, 48, 16, + 48,144,144, 48, + 48,144,144, 48, + 16, 48, 48, 16, +}; + +/* =================================================================== + * ME search patterns — matching CPU motion_estimation.c + * =================================================================== */ +__device__ const int d_sqr1[8][2] = { + { 0,-1},{ 0, 1},{-1, 0},{ 1, 0},{-1,-1},{ 1,-1},{-1, 1},{ 1, 1}}; +__device__ const int d_dia1[4][2] = { + { 0,-1},{ 0, 1},{-1, 0},{ 1, 0}}; +__device__ const int d_dia2[8][2] = { + {-2, 0},{-1,-1},{ 0,-2},{ 1,-1},{ 2, 0},{ 1, 1},{ 0, 2},{-1, 1}}; +__device__ const int d_hex2[6][2] = { + {-2, 0},{-1,-2},{-1, 2},{ 1,-2},{ 1, 2},{ 2, 0}}; +__device__ const int d_hex4_umh[16][2] = { + { 4, 0},{ 2, 3},{-2, 3},{-4, 0},{-2,-3},{ 2,-3},{ 6,-2},{ 6, 2}, + {-6, 2},{-6,-2},{ 2, 5},{-2, 5},{ 4,-5},{-4,-5},{ 0,-7},{ 0, 7}}; + +/* =================================================================== + * Device helper functions + * =================================================================== */ + +/* Device-side OBMC weight lookup (force inline) */ +static __device__ __forceinline__ uchar obmc_weight_lut(int rx, int ry, + int log2_mb_size) +{ + int stride = 2 << log2_mb_size; + int idx = rx + ry * stride; + switch (1 << log2_mb_size) { + case 16: return d_obmc_linear32[idx]; + case 8: return d_obmc_linear16[idx]; + case 4: return d_obmc_linear8[idx]; + default: return d_obmc_linear4[idx]; + } +} + +/* Three-value median — matches CPU mid_pred() */ +static __device__ __forceinline__ int mid_pred(int a, int b, int c) +{ + if (a > b) { int t = a; a = b; b = t; } + if (b > c) { b = c; } + if (a > b) { b = a; } + return b; +} + +/* Compute SAD over mb_size×mb_size block comparing s_cur vs tex_ref */ +static __device__ __forceinline__ uint64_t mb_sad( + cudaTextureObject_t tex_ref, const uchar *s_cur, + int x_mv, int y_mv, int mb_size) +{ + uint64_t sad = 0; + for (int j = 0; j < mb_size; j++) { + for (int i = 0; i < mb_size; i++) { + uchar cur_val = s_cur[i + j * mb_size]; + uchar ref_val = tex2D<uchar>(tex_ref, x_mv + i, y_mv + j); + sad += abs((int)cur_val - (int)ref_val); + } + } + return sad; +} + +/* Compute bilateral SAD (2xMB window) — matches CPU get_sbad_ob clip-before-loop */ +static __device__ __forceinline__ uint64_t mb_sbad( + cudaTextureObject_t tex_cur, cudaTextureObject_t tex_ref, + int x_mb, int y_mb, int mv_x, int mv_y, int mb_size, + int width, int height) +{ + int half = mb_size / 2; + int x_min = half; + int x_max = width - mb_size * 3 / 2; + int y_min = half; + int y_max = height - mb_size * 3 / 2; + + /* Clip MB position (matches CPU get_sbad_ob) */ + int x = max(x_min, min(x_max, x_mb)); + int y = max(y_min, min(y_max, y_mb)); + + /* Clip MV to keep 2xMB window fully in bounds */ + int cmv_x = mv_x + (x_mb - x); + cmv_x = max(-(x - x_min), min(x_max - x, cmv_x)); + int cmv_y = mv_y + (y_mb - y); + cmv_y = max(-(y - y_min), min(y_max - y, cmv_y)); + + uint64_t sbad = 0; + for (int j = -half; j < mb_size * 3 / 2; j++) { + int yc = y + cmv_y + j; + int yr = y - cmv_y + j; + for (int i = -half; i < mb_size * 3 / 2; i++) { + int xc = x + cmv_x + i; + int xr = x - cmv_x + i; + int cur_val = tex2D<uchar>(tex_cur, xc, yc); + int ref_val = tex2D<uchar>(tex_ref, xr, yr); + sbad += abs(cur_val - ref_val); + } + } + return sbad; +} + +/* Shared-memory block-level min reduction */ +static __device__ __forceinline__ void block_reduce_min( + uint64_t *s_cost, int16_t *s_mv_x, int16_t *s_mv_y, int tid, int n) +{ + for (int stride = n / 2; stride > 0; stride >>= 1) { + if (tid < stride) { + if (s_cost[tid + stride] < s_cost[tid]) { + s_cost[tid] = s_cost[tid + stride]; + s_mv_x[tid] = s_mv_x[tid + stride]; + s_mv_y[tid] = s_mv_y[tid + stride]; + } + } + __syncthreads(); + } +} + +extern "C" +{ + +/* =================================================================== + * ESA — Exhaustive Search Motion Estimation (optimized, Phase 0) + * Grid: (b_count, 1, 1) Block: (256, 1, 1) + * Shared memory: cur block prefetch + reduction arrays + * =================================================================== */ +__global__ void minterpolate_me_esa( + cudaTextureObject_t tex_cur, + cudaTextureObject_t tex_ref, + int16_t *__restrict__ mv_buffer, + int b_width, int b_height, + int mb_size, int search_param, + int width, int height) +{ + __shared__ uint64_t s_best_cost[256]; + __shared__ int16_t s_best_mv_x[256]; + __shared__ int16_t s_best_mv_y[256]; + + int mb_idx = blockIdx.x; + int tid = threadIdx.x; + + if (mb_idx >= b_width * b_height) + return; + + int mb_x = mb_idx % b_width; + int mb_y = mb_idx / b_width; + int x_mb = mb_x * mb_size; + int y_mb = mb_y * mb_size; + + /* Search bounds */ + int sy_min = max(mb_size/2, y_mb - search_param); + int sy_max = min(height - mb_size*3/2, y_mb + search_param); + int sx_min = max(mb_size/2, x_mb - search_param); + int sx_max = min(width - mb_size*3/2, x_mb + search_param); + int n_cand_x = sx_max - sx_min + 1; + int n_cand_y = sy_max - sy_min + 1; + int total_candidates = n_cand_x * n_cand_y; + + uint64_t local_best_cost = 0xFFFFFFFFFFFFFFFFULL; + int16_t local_best_mv_x = 0, local_best_mv_y = 0; + + for (int base = 0; base < total_candidates; base += blockDim.x) { + int cand_idx = base + tid; + uint64_t sad = 0xFFFFFFFFFFFFFFFFULL; + int16_t mv_x = 0, mv_y = 0; + + if (cand_idx < total_candidates) { + int cy = cand_idx / n_cand_x; + int cx = cand_idx % n_cand_x; + mv_x = (int16_t)((sx_min + cx) - x_mb); + mv_y = (int16_t)((sy_min + cy) - y_mb); + int x_mv = x_mb + mv_x; + int y_mv = y_mb + mv_y; + + sad = mb_sbad(tex_cur, tex_ref, x_mb, y_mb, x_mv - x_mb, y_mv - y_mb, mb_size, width, height); + sad += (abs((int)mv_x) + abs((int)mv_y)) * COST_PRED_SCALE; + } + + if (sad < local_best_cost) { + local_best_cost = sad; + local_best_mv_x = mv_x; + local_best_mv_y = mv_y; + } + } + + s_best_cost[tid] = local_best_cost; + s_best_mv_x[tid] = local_best_mv_x; + s_best_mv_y[tid] = local_best_mv_y; + __syncthreads(); + + block_reduce_min(s_best_cost, s_best_mv_x, s_best_mv_y, tid, 256); + + if (tid == 0) { + mv_buffer[mb_idx * 2] = s_best_mv_x[0]; + mv_buffer[mb_idx * 2 + 1] = s_best_mv_y[0]; + } +} + +/* =================================================================== + * TSS — Three Step Search + * Grid: (b_count, 1, 1) Block: (1, 1, 1) + * =================================================================== */ +__global__ void minterpolate_me_tss( + cudaTextureObject_t tex_cur, + cudaTextureObject_t tex_ref, + int16_t *__restrict__ mv_buffer, + int b_width, int b_height, + int mb_size, int search_param, + int width, int height) +{ + int mb_idx = blockIdx.x; + if (mb_idx >= b_width * b_height) return; + + int mb_x = mb_idx % b_width; + int mb_y = mb_idx / b_width; + int x_mb = mb_x * mb_size; + int y_mb = mb_y * mb_size; + + + int mv_x = 0, mv_y = 0; + uint64_t best_cost = mb_sbad(tex_cur, tex_ref, x_mb, y_mb, 0, 0, mb_size, width, height); + + int step = (search_param + 1) / 2; + while (step > 0) { + for (int k = 0; k < 8; k++) { + int cx = mv_x + d_sqr1[k][0] * step; + int cy = mv_y + d_sqr1[k][1] * step; + int x_mv = x_mb + cx, y_mv = y_mb + cy; + if (x_mv < mb_size/2 || x_mv + mb_size*3/2 > width || + y_mv < mb_size/2 || y_mv + mb_size*3/2 > height) + continue; + uint64_t cost = mb_sbad(tex_cur, tex_ref, x_mb, y_mb, x_mv - x_mb, y_mv - y_mb, mb_size, width, height) + + (abs(cx) + abs(cy)) * COST_PRED_SCALE; + if (cost < best_cost) { + best_cost = cost; + mv_x = cx; mv_y = cy; + } + } + step /= 2; + } + + mv_buffer[mb_idx * 2] = (int16_t)mv_x; + mv_buffer[mb_idx * 2 + 1] = (int16_t)mv_y; +} + +/* =================================================================== + * TDLS — Two Dimensional Logarithmic Search + * Grid: (b_count, 1, 1) Block: (1, 1, 1) + * =================================================================== */ +__global__ void minterpolate_me_tdls( + cudaTextureObject_t tex_cur, + cudaTextureObject_t tex_ref, + int16_t *__restrict__ mv_buffer, + int b_width, int b_height, + int mb_size, int search_param, + int width, int height) +{ + int mb_idx = blockIdx.x; + if (mb_idx >= b_width * b_height) return; + + int mb_x = mb_idx % b_width; + int mb_y = mb_idx / b_width; + int x_mb = mb_x * mb_size; + int y_mb = mb_y * mb_size; + + + int mv_x = 0, mv_y = 0; + uint64_t best_cost = mb_sbad(tex_cur, tex_ref, x_mb, y_mb, 0, 0, mb_size, width, height); + int step = search_param / 2; + + while (step > 0) { + int improved = 0; + for (int k = 0; k < 4; k++) { + int cx = mv_x + d_dia1[k][0] * step; + int cy = mv_y + d_dia1[k][1] * step; + int x_mv = x_mb + cx, y_mv = y_mb + cy; + if (x_mv < mb_size/2 || x_mv + mb_size*3/2 > width || + y_mv < mb_size/2 || y_mv + mb_size*3/2 > height) + continue; + uint64_t cost = mb_sbad(tex_cur, tex_ref, x_mb, y_mb, x_mv - x_mb, y_mv - y_mb, mb_size, width, height) + + (abs(cx) + abs(cy)) * COST_PRED_SCALE; + if (cost < best_cost) { + best_cost = cost; + mv_x = cx; mv_y = cy; + improved = 1; + } + } + if (!improved) + step /= 2; + } + + mv_buffer[mb_idx * 2] = (int16_t)mv_x; + mv_buffer[mb_idx * 2 + 1] = (int16_t)mv_y; +} + +/* =================================================================== + * NTSS — New Three Step Search + * Grid: (b_count, 1, 1) Block: (1, 1, 1) + * =================================================================== */ +__global__ void minterpolate_me_ntss( + cudaTextureObject_t tex_cur, + cudaTextureObject_t tex_ref, + int16_t *__restrict__ mv_buffer, + int b_width, int b_height, + int mb_size, int search_param, + int width, int height) +{ + int mb_idx = blockIdx.x; + if (mb_idx >= b_width * b_height) return; + + int mb_x = mb_idx % b_width; + int mb_y = mb_idx / b_width; + int x_mb = mb_x * mb_size; + int y_mb = mb_y * mb_size; + + + int mv_x = 0, mv_y = 0; + uint64_t best_cost = mb_sbad(tex_cur, tex_ref, x_mb, y_mb, 0, 0, mb_size, width, height); + + int step = max(4, (search_param + 1) / 2); + + /* First step: check neighbors at step=1 (8 extra points) */ + if (step >= 4) { + for (int k = 0; k < 8; k++) { + int cx = d_sqr1[k][0], cy = d_sqr1[k][1]; + int x_mv = x_mb + cx, y_mv = y_mb + cy; + if (x_mv < mb_size/2 || x_mv + mb_size*3/2 > width || + y_mv < mb_size/2 || y_mv + mb_size*3/2 > height) + continue; + uint64_t cost = mb_sbad(tex_cur, tex_ref, x_mb, y_mb, x_mv - x_mb, y_mv - y_mb, mb_size, width, height) + + (abs(cx) + abs(cy)) * COST_PRED_SCALE; + if (cost < best_cost) { + best_cost = cost; + mv_x = cx; mv_y = cy; + } + } + /* Early termination: if center still best, stop */ + if (mv_x == 0 && mv_y == 0) + goto done; + } + + while (step > 0) { + int improved = 0; + for (int k = 0; k < 8; k++) { + int cx = mv_x + d_sqr1[k][0] * step; + int cy = mv_y + d_sqr1[k][1] * step; + int x_mv = x_mb + cx, y_mv = y_mb + cy; + if (x_mv < mb_size/2 || x_mv + mb_size*3/2 > width || + y_mv < mb_size/2 || y_mv + mb_size*3/2 > height) + continue; + uint64_t cost = mb_sbad(tex_cur, tex_ref, x_mb, y_mb, x_mv - x_mb, y_mv - y_mb, mb_size, width, height) + + (abs(cx) + abs(cy)) * COST_PRED_SCALE; + if (cost < best_cost) { + best_cost = cost; + mv_x = cx; mv_y = cy; + improved = 1; + } + } + if (!improved) step /= 2; + /* If best is at ±1, refine with all 8 neighbors and stop */ + if (abs(mv_x) <= 1 && abs(mv_y) <= 1 && step >= 2) { + for (int k = 0; k < 8; k++) { + int cx = mv_x + d_sqr1[k][0]; + int cy = mv_y + d_sqr1[k][1]; + int x_mv = x_mb + cx, y_mv = y_mb + cy; + if (x_mv < mb_size/2 || x_mv + mb_size*3/2 > width || + y_mv < mb_size/2 || y_mv + mb_size*3/2 > height) + continue; + uint64_t cost = mb_sbad(tex_cur, tex_ref, x_mb, y_mb, x_mv - x_mb, y_mv - y_mb, mb_size, width, height) + + (abs(cx) + abs(cy)) * COST_PRED_SCALE; + if (cost < best_cost) { + best_cost = cost; + mv_x = cx; mv_y = cy; + } + } + goto done; + } + } + +done: + mv_buffer[mb_idx * 2] = (int16_t)mv_x; + mv_buffer[mb_idx * 2 + 1] = (int16_t)mv_y; +} + +/* =================================================================== + * FSS — Four Step Search + * Grid: (b_count, 1, 1) Block: (1, 1, 1) + * =================================================================== */ +__global__ void minterpolate_me_fss( + cudaTextureObject_t tex_cur, + cudaTextureObject_t tex_ref, + int16_t *__restrict__ mv_buffer, + int b_width, int b_height, + int mb_size, int search_param, + int width, int height) +{ + int mb_idx = blockIdx.x; + if (mb_idx >= b_width * b_height) return; + + int mb_x = mb_idx % b_width; + int mb_y = mb_idx / b_width; + int x_mb = mb_x * mb_size; + int y_mb = mb_y * mb_size; + + + int mv_x = 0, mv_y = 0; + uint64_t best_cost = mb_sbad(tex_cur, tex_ref, x_mb, y_mb, 0, 0, mb_size, width, height); + int step = 2; + + while (step > 0) { + int improved = 0; + for (int k = 0; k < 8; k++) { + int cx = mv_x + d_sqr1[k][0] * step; + int cy = mv_y + d_sqr1[k][1] * step; + int x_mv = x_mb + cx, y_mv = y_mb + cy; + if (x_mv < mb_size/2 || x_mv + mb_size*3/2 > width || + y_mv < mb_size/2 || y_mv + mb_size*3/2 > height) + continue; + uint64_t cost = mb_sbad(tex_cur, tex_ref, x_mb, y_mb, x_mv - x_mb, y_mv - y_mb, mb_size, width, height) + + (abs(cx) + abs(cy)) * COST_PRED_SCALE; + if (cost < best_cost) { + best_cost = cost; + mv_x = cx; mv_y = cy; + improved = 1; + } + } + if (!improved) step /= 2; + } + + mv_buffer[mb_idx * 2] = (int16_t)mv_x; + mv_buffer[mb_idx * 2 + 1] = (int16_t)mv_y; +} + +/* =================================================================== + * DS — Diamond Search + * Grid: (b_count, 1, 1) Block: (1, 1, 1) + * =================================================================== */ +__global__ void minterpolate_me_ds( + cudaTextureObject_t tex_cur, + cudaTextureObject_t tex_ref, + int16_t *__restrict__ mv_buffer, + int b_width, int b_height, + int mb_size, int search_param, + int width, int height) +{ + int mb_idx = blockIdx.x; + if (mb_idx >= b_width * b_height) return; + + int mb_x = mb_idx % b_width; + int mb_y = mb_idx / b_width; + int x_mb = mb_x * mb_size; + int y_mb = mb_y * mb_size; + + + int mv_x = 0, mv_y = 0; + uint64_t best_cost = mb_sbad(tex_cur, tex_ref, x_mb, y_mb, 0, 0, mb_size, width, height); + + /* Stage 1: Large diamond (dia2) — iterate until center is best */ + for (;;) { + int improved = 0; + for (int k = 0; k < 8; k++) { + int cx = mv_x + d_dia2[k][0]; + int cy = mv_y + d_dia2[k][1]; + int x_mv = x_mb + cx, y_mv = y_mb + cy; + if (x_mv < mb_size/2 || x_mv + mb_size*3/2 > width || + y_mv < mb_size/2 || y_mv + mb_size*3/2 > height) + continue; + uint64_t cost = mb_sbad(tex_cur, tex_ref, x_mb, y_mb, x_mv - x_mb, y_mv - y_mb, mb_size, width, height) + + (abs(cx) + abs(cy)) * COST_PRED_SCALE; + if (cost < best_cost) { + best_cost = cost; + mv_x = cx; mv_y = cy; + improved = 1; + } + } + if (!improved) break; + } + + /* Stage 2: Small diamond (dia1) refinement */ + { + int improved; + do { + improved = 0; + for (int k = 0; k < 4; k++) { + int cx = mv_x + d_dia1[k][0]; + int cy = mv_y + d_dia1[k][1]; + int x_mv = x_mb + cx, y_mv = y_mb + cy; + if (x_mv < mb_size/2 || x_mv + mb_size*3/2 > width || + y_mv < mb_size/2 || y_mv + mb_size*3/2 > height) + continue; + uint64_t cost = mb_sbad(tex_cur, tex_ref, x_mb, y_mb, x_mv - x_mb, y_mv - y_mb, mb_size, width, height) + + (abs(cx) + abs(cy)) * COST_PRED_SCALE; + if (cost < best_cost) { + best_cost = cost; + mv_x = cx; mv_y = cy; + improved = 1; + } + } + } while (improved); + } + + mv_buffer[mb_idx * 2] = (int16_t)mv_x; + mv_buffer[mb_idx * 2 + 1] = (int16_t)mv_y; +} + +/* =================================================================== + * HEXBS — Hexagon-Based Search + * Grid: (b_count, 1, 1) Block: (1, 1, 1) + * =================================================================== */ +__global__ void minterpolate_me_hexbs( + cudaTextureObject_t tex_cur, + cudaTextureObject_t tex_ref, + int16_t *__restrict__ mv_buffer, + int b_width, int b_height, + int mb_size, int search_param, + int width, int height) +{ + int mb_idx = blockIdx.x; + if (mb_idx >= b_width * b_height) return; + + int mb_x = mb_idx % b_width; + int mb_y = mb_idx / b_width; + int x_mb = mb_x * mb_size; + int y_mb = mb_y * mb_size; + + + int mv_x = 0, mv_y = 0; + uint64_t best_cost = mb_sbad(tex_cur, tex_ref, x_mb, y_mb, 0, 0, mb_size, width, height); + + /* Stage 1: Large hexagon — iterate until center is best */ + for (;;) { + int improved = 0; + for (int k = 0; k < 6; k++) { + int cx = mv_x + d_hex2[k][0]; + int cy = mv_y + d_hex2[k][1]; + int x_mv = x_mb + cx, y_mv = y_mb + cy; + if (x_mv < mb_size/2 || x_mv + mb_size*3/2 > width || + y_mv < mb_size/2 || y_mv + mb_size*3/2 > height) + continue; + uint64_t cost = mb_sbad(tex_cur, tex_ref, x_mb, y_mb, x_mv - x_mb, y_mv - y_mb, mb_size, width, height) + + (abs(cx) + abs(cy)) * COST_PRED_SCALE; + if (cost < best_cost) { + best_cost = cost; + mv_x = cx; mv_y = cy; + improved = 1; + } + } + if (!improved) break; + } + + /* Stage 2: Small diamond refinement */ + { + int improved; + do { + improved = 0; + for (int k = 0; k < 4; k++) { + int cx = mv_x + d_dia1[k][0]; + int cy = mv_y + d_dia1[k][1]; + int x_mv = x_mb + cx, y_mv = y_mb + cy; + if (x_mv < mb_size/2 || x_mv + mb_size*3/2 > width || + y_mv < mb_size/2 || y_mv + mb_size*3/2 > height) + continue; + uint64_t cost = mb_sbad(tex_cur, tex_ref, x_mb, y_mb, x_mv - x_mb, y_mv - y_mb, mb_size, width, height) + + (abs(cx) + abs(cy)) * COST_PRED_SCALE; + if (cost < best_cost) { + best_cost = cost; + mv_x = cx; mv_y = cy; + improved = 1; + } + } + } while (improved); + } + + mv_buffer[mb_idx * 2] = (int16_t)mv_x; + mv_buffer[mb_idx * 2 + 1] = (int16_t)mv_y; +} + +/* =================================================================== + * EPZS Pass 1 — temporal predictors only (fully parallel) + * Grid: (b_count, 1, 1) Block: (1, 1, 1) + * + * Uses mv_table_prev1 as source of temporal predictors. + * No spatial dependency — all MBs processable in parallel. + * =================================================================== */ +__global__ void minterpolate_me_epzs_pass1( + cudaTextureObject_t tex_cur, + cudaTextureObject_t tex_ref, + int16_t *__restrict__ mv_buffer_out, + const int16_t *__restrict__ mv_table_prev1, + const int16_t *__restrict__ mv_table_prev2, + int b_width, int b_height, + int mb_size, int search_param, + int width, int height) +{ + int mb_idx = blockIdx.x; + if (mb_idx >= b_width * b_height) return; + + int mb_x = mb_idx % b_width; + int mb_y = mb_idx / b_width; + int x_mb = mb_x * mb_size; + int y_mb = mb_y * mb_size; + + + int16_t pred_mvs[10][2]; + int nb_pred = 0; + + /* (0,0) always */ + pred_mvs[nb_pred][0] = 0; + pred_mvs[nb_pred][1] = 0; nb_pred++; + + /* Collocated MB from prev frame */ + { + int px = mv_table_prev1[mb_idx * 2]; + int py = mv_table_prev1[mb_idx * 2 + 1]; + pred_mvs[nb_pred][0] = (int16_t)px; + pred_mvs[nb_pred][1] = (int16_t)py; nb_pred++; + } + + /* Accelerator MV */ + { + int p1x = mv_table_prev1[mb_idx * 2]; + int p1y = mv_table_prev1[mb_idx * 2 + 1]; + int p2x = mv_table_prev2[mb_idx * 2]; + int p2y = mv_table_prev2[mb_idx * 2 + 1]; + int ax = p1x + (p1x - p2x); + int ay = p1y + (p1y - p2y); + pred_mvs[nb_pred][0] = (int16_t)ax; + pred_mvs[nb_pred][1] = (int16_t)ay; nb_pred++; + } + + /* Neighbors from prev frame */ + if (mb_x > 0) { + int ni = mb_idx - 1; + pred_mvs[nb_pred][0] = mv_table_prev1[ni * 2]; + pred_mvs[nb_pred][1] = mv_table_prev1[ni * 2 + 1]; nb_pred++; + } + if (mb_y > 0) { + int ni = mb_idx - b_width; + pred_mvs[nb_pred][0] = mv_table_prev1[ni * 2]; + pred_mvs[nb_pred][1] = mv_table_prev1[ni * 2 + 1]; nb_pred++; + } + if (mb_x + 1 < b_width) { + int ni = mb_idx + 1; + pred_mvs[nb_pred][0] = mv_table_prev1[ni * 2]; + pred_mvs[nb_pred][1] = mv_table_prev1[ni * 2 + 1]; nb_pred++; + } + if (mb_y + 1 < b_height) { + int ni = mb_idx + b_width; + pred_mvs[nb_pred][0] = mv_table_prev1[ni * 2]; + pred_mvs[nb_pred][1] = mv_table_prev1[ni * 2 + 1]; nb_pred++; + } + + /* Evaluate all predictors, pick best */ + int16_t best_mv_x = 0, best_mv_y = 0; + int16_t pred_x = 0, pred_y = 0; + uint64_t best_cost = 0xFFFFFFFFFFFFFFFFULL; + + for (int p = 0; p < nb_pred; p++) { + int cx = pred_mvs[p][0], cy = pred_mvs[p][1]; + int x_mv = x_mb + cx, y_mv = y_mb + cy; + if (x_mv < mb_size/2 || x_mv + mb_size*3/2 > width || + y_mv < mb_size/2 || y_mv + mb_size*3/2 > height) + continue; + uint64_t cost = mb_sbad(tex_cur, tex_ref, x_mb, y_mb, x_mv - x_mb, y_mv - y_mb, mb_size, width, height) + + (abs(cx) + abs(cy)) * COST_PRED_SCALE; + if (cost < best_cost) { + best_cost = cost; + best_mv_x = (int16_t)cx; + best_mv_y = (int16_t)cy; + } + } + pred_x = best_mv_x; pred_y = best_mv_y; + + /* Diamond refinement around best predictor */ + { + int16_t cur_x = best_mv_x, cur_y = best_mv_y; + int improved; + do { + improved = 0; + for (int k = 0; k < 4; k++) { + int cx = cur_x + d_dia1[k][0]; + int cy = cur_y + d_dia1[k][1]; + int x_mv = x_mb + cx, y_mv = y_mb + cy; + if (x_mv < mb_size/2 || x_mv + mb_size*3/2 > width || + y_mv < mb_size/2 || y_mv + mb_size*3/2 > height) + continue; + uint64_t cost = mb_sbad(tex_cur, tex_ref, x_mb, y_mb, x_mv - x_mb, y_mv - y_mb, mb_size, width, height) + + (abs(cx - pred_x) + abs(cy - pred_y)) * COST_PRED_SCALE; + if (cost < best_cost) { + best_cost = cost; + cur_x = (int16_t)cx; + cur_y = (int16_t)cy; + improved = 1; + } + } + } while (improved); + best_mv_x = cur_x; best_mv_y = cur_y; + } + + mv_buffer_out[mb_idx * 2] = best_mv_x; + mv_buffer_out[mb_idx * 2 + 1] = best_mv_y; +} + +/* =================================================================== + * EPZS Pass 2 — spatial predictors from Pass1 + temporal from prev frame + * Grid: (b_count, 1, 1) Block: (1, 1, 1) + * + * Reads left/top/top-right spatial neighbor MVs from mv_buffer_pass1. + * Computes median predictor, adds more predictors, does diamond refinement. + * =================================================================== */ +__global__ void minterpolate_me_epzs_pass2( + cudaTextureObject_t tex_cur, + cudaTextureObject_t tex_ref, + int16_t *__restrict__ mv_buffer_final, + const int16_t *__restrict__ mv_buffer_pass1, + const int16_t *__restrict__ mv_table_prev1, + const int16_t *__restrict__ mv_table_prev2, + int b_width, int b_height, + int mb_size, int search_param, + int width, int height) +{ + int mb_idx = blockIdx.x; + if (mb_idx >= b_width * b_height) return; + + int mb_x = mb_idx % b_width; + int mb_y = mb_idx / b_width; + int x_mb = mb_x * mb_size; + int y_mb = mb_y * mb_size; + + + int16_t preds0[10][2], preds1[10][2]; + int nb0 = 0, nb1 = 0; + + /* --- Predictor set 0 (spatial + collocated) --- */ + /* (0,0) */ + preds0[nb0][0] = 0; preds0[nb0][1] = 0; nb0++; + + /* Left neighbor from pass1 */ + if (mb_x > 0) { + int ni = mb_idx - 1; + preds0[nb0][0] = mv_buffer_pass1[ni * 2]; + preds0[nb0][1] = mv_buffer_pass1[ni * 2 + 1]; nb0++; + } + /* Top neighbor from pass1 */ + if (mb_y > 0) { + int ni = mb_idx - b_width; + preds0[nb0][0] = mv_buffer_pass1[ni * 2]; + preds0[nb0][1] = mv_buffer_pass1[ni * 2 + 1]; nb0++; + } + /* Top-right neighbor from pass1 */ + if (mb_y > 0 && mb_x + 1 < b_width) { + int ni = mb_idx - b_width + 1; + preds0[nb0][0] = mv_buffer_pass1[ni * 2]; + preds0[nb0][1] = mv_buffer_pass1[ni * 2 + 1]; nb0++; + } + + /* Median predictor */ + int med_x = 0, med_y = 0; + if (nb0 == 5) { /* (0,0) + left + top + top-r = 4 spatial preds */ + med_x = mid_pred(preds0[1][0], preds0[2][0], preds0[3][0]); + med_y = mid_pred(preds0[1][1], preds0[2][1], preds0[3][1]); + } else if (nb0 == 4) { + med_x = mid_pred(0, preds0[1][0], preds0[2][0]); + med_y = mid_pred(0, preds0[1][1], preds0[2][1]); + } else if (nb0 == 3) { + med_x = preds0[1][0]; + med_y = preds0[1][1]; + } + /* Add median as predictor */ + if (med_x != 0 || med_y != 0) { + preds0[nb0][0] = (int16_t)med_x; + preds0[nb0][1] = (int16_t)med_y; nb0++; + } + + /* Collocated MB from prev frame */ + { + int px = mv_table_prev1[mb_idx * 2]; + int py = mv_table_prev1[mb_idx * 2 + 1]; + preds0[nb0][0] = (int16_t)px; + preds0[nb0][1] = (int16_t)py; nb0++; + } + + /* --- Predictor set 1 (temporal neighbors) --- */ + /* Accelerator MV */ + { + int p1x = mv_table_prev1[mb_idx * 2]; + int p1y = mv_table_prev1[mb_idx * 2 + 1]; + int p2x = mv_table_prev2[mb_idx * 2]; + int p2y = mv_table_prev2[mb_idx * 2 + 1]; + preds1[nb1][0] = (int16_t)(p1x + (p1x - p2x)); + preds1[nb1][1] = (int16_t)(p1y + (p1y - p2y)); nb1++; + } + /* Prev frame spatial neighbors */ + if (mb_x > 0) { + int ni = mb_idx - 1; + preds1[nb1][0] = mv_table_prev1[ni * 2]; + preds1[nb1][1] = mv_table_prev1[ni * 2 + 1]; nb1++; + } + if (mb_y > 0) { + int ni = mb_idx - b_width; + preds1[nb1][0] = mv_table_prev1[ni * 2]; + preds1[nb1][1] = mv_table_prev1[ni * 2 + 1]; nb1++; + } + if (mb_x + 1 < b_width) { + int ni = mb_idx + 1; + preds1[nb1][0] = mv_table_prev1[ni * 2]; + preds1[nb1][1] = mv_table_prev1[ni * 2 + 1]; nb1++; + } + if (mb_y + 1 < b_height) { + int ni = mb_idx + b_width; + preds1[nb1][0] = mv_table_prev1[ni * 2]; + preds1[nb1][1] = mv_table_prev1[ni * 2 + 1]; nb1++; + } + + /* Evaluate all predictors */ + int16_t best_mv_x = 0, best_mv_y = 0; + int16_t pred_x = 0, pred_y = 0; + uint64_t best_cost = 0xFFFFFFFFFFFFFFFFULL; + + #define EVAL_PRED(px, py) do { \ + int cx = (px), cy = (py); \ + int x_mv = x_mb + cx, y_mv = y_mb + cy; \ + if (x_mv >= mb_size/2 && x_mv + mb_size*3/2 <= width && \ + y_mv >= mb_size/2 && y_mv + mb_size*3/2 <= height) { \ + uint64_t cost = mb_sbad(tex_cur, tex_ref, x_mb, y_mb, x_mv - x_mb, y_mv - y_mb, mb_size, width, height) \ + + (abs(cx) + abs(cy)) * COST_PRED_SCALE; \ + if (cost < best_cost) { \ + best_cost = cost; best_mv_x = (int16_t)cx; best_mv_y = (int16_t)cy; \ + } \ + } \ + } while(0) + + /* Evaluate preds0 (including median) */ + for (int p = 0; p < nb0; p++) + EVAL_PRED(preds0[p][0], preds0[p][1]); + + /* Evaluate preds1 */ + for (int p = 0; p < nb1; p++) + EVAL_PRED(preds1[p][0], preds1[p][1]); + + #undef EVAL_PRED + + pred_x = best_mv_x; pred_y = best_mv_y; + + /* Diamond refinement */ + { + int16_t cur_x = best_mv_x, cur_y = best_mv_y; + int improved; + do { + improved = 0; + for (int k = 0; k < 4; k++) { + int cx = cur_x + d_dia1[k][0]; + int cy = cur_y + d_dia1[k][1]; + int x_mv = x_mb + cx, y_mv = y_mb + cy; + if (x_mv < mb_size/2 || x_mv + mb_size*3/2 > width || + y_mv < mb_size/2 || y_mv + mb_size*3/2 > height) + continue; + uint64_t cost = mb_sbad(tex_cur, tex_ref, x_mb, y_mb, x_mv - x_mb, y_mv - y_mb, mb_size, width, height) + + (abs(cx) + abs(cy)) * COST_PRED_SCALE; + if (cost < best_cost) { + best_cost = cost; + cur_x = (int16_t)cx; + cur_y = (int16_t)cy; + improved = 1; + } + } + } while (improved); + best_mv_x = cur_x; best_mv_y = cur_y; + } + + mv_buffer_final[mb_idx * 2] = best_mv_x; + mv_buffer_final[mb_idx * 2 + 1] = best_mv_y; +} + +/* =================================================================== + * UMH Pass 1 — temporal predictors + UMH search pattern + * Grid: (b_count, 1, 1) Block: (1, 1, 1) + * =================================================================== */ +__global__ void minterpolate_me_umh_pass1( + cudaTextureObject_t tex_cur, + cudaTextureObject_t tex_ref, + int16_t *__restrict__ mv_buffer_out, + const int16_t *__restrict__ mv_table_prev1, + int b_width, int b_height, + int mb_size, int search_param, + int width, int height) +{ + int mb_idx = blockIdx.x; + if (mb_idx >= b_width * b_height) return; + + int mb_x = mb_idx % b_width; + int mb_y = mb_idx / b_width; + int x_mb = mb_x * mb_size; + int y_mb = mb_y * mb_size; + + + int16_t preds[5][2]; + int nb_pred = 0; + + /* (0,0), left, top, top-right/left from prev frame (simplified for pass1) */ + preds[nb_pred][0] = 0; preds[nb_pred][1] = 0; nb_pred++; + if (mb_x > 0) { + int ni = mb_idx - 1; + preds[nb_pred][0] = mv_table_prev1[ni * 2]; + preds[nb_pred][1] = mv_table_prev1[ni * 2 + 1]; nb_pred++; + } + if (mb_y > 0) { + int ni = mb_idx - b_width; + preds[nb_pred][0] = mv_table_prev1[ni * 2]; + preds[nb_pred][1] = mv_table_prev1[ni * 2 + 1]; nb_pred++; + if (mb_x + 1 < b_width) { + ni = mb_idx - b_width + 1; + preds[nb_pred][0] = mv_table_prev1[ni * 2]; + preds[nb_pred][1] = mv_table_prev1[ni * 2 + 1]; nb_pred++; + } else if (mb_x > 0) { + ni = mb_idx - b_width - 1; + preds[nb_pred][0] = mv_table_prev1[ni * 2]; + preds[nb_pred][1] = mv_table_prev1[ni * 2 + 1]; nb_pred++; + } + } + + /* Median predictor (add to list if valid) */ + if (nb_pred >= 3) { + int med_x, med_y; + if (nb_pred >= 4) { + med_x = mid_pred(preds[1][0], preds[2][0], preds[3][0]); + med_y = mid_pred(preds[1][1], preds[2][1], preds[3][1]); + } else { + med_x = preds[1][0]; + med_y = preds[1][1]; + } + if (med_x != 0 || med_y != 0) { + preds[nb_pred][0] = (int16_t)med_x; + preds[nb_pred][1] = (int16_t)med_y; + nb_pred++; + } + } + + /* Eval best predictor */ + int16_t best_mv_x = 0, best_mv_y = 0; + uint64_t best_cost = 0xFFFFFFFFFFFFFFFFULL; + + for (int p = 0; p < nb_pred; p++) { + int cx = preds[p][0], cy = preds[p][1]; + int x_mv = x_mb + cx, y_mv = y_mb + cy; + if (x_mv < mb_size/2 || x_mv + mb_size*3/2 > width || + y_mv < mb_size/2 || y_mv + mb_size*3/2 > height) + continue; + uint64_t cost = mb_sbad(tex_cur, tex_ref, x_mb, y_mb, x_mv - x_mb, y_mv - y_mb, mb_size, width, height) + + (abs(cx) + abs(cy)) * COST_PRED_SCALE; + if (cost < best_cost) { best_cost = cost; best_mv_x = (int16_t)cx; best_mv_y = (int16_t)cy; } + } + + /* UMH: unsymmetrical cross search */ + for (int d = 1; d <= search_param; d += 2) { + int cx = best_mv_x + d, cy = best_mv_y; + if (cx >= mb_size/2 - x_mb && cx + mb_size*3/2 <= width - x_mb && cy >= mb_size/2 - y_mb && cy + mb_size*3/2 <= height - y_mb) { + uint64_t cost = mb_sbad(tex_cur, tex_ref, x_mb, y_mb, cx, cy, mb_size, width, height) + (abs(cx) + abs(cy)) * COST_PRED_SCALE; + if (cost < best_cost) { best_cost = cost; best_mv_x = (int16_t)cx; } + } + cx = best_mv_x - d; + if (cx >= mb_size/2 - x_mb && cx + mb_size*3/2 <= width - x_mb && cy >= mb_size/2 - y_mb && cy + mb_size*3/2 <= height - y_mb) { + uint64_t cost = mb_sbad(tex_cur, tex_ref, x_mb, y_mb, cx, cy, mb_size, width, height) + (abs(cx) + abs(cy)) * COST_PRED_SCALE; + if (cost < best_cost) { best_cost = cost; best_mv_x = (int16_t)cx; } + } + } + for (int d = 1; d <= search_param / 2; d += 2) { + int cx = best_mv_x, cy = best_mv_y + d; + if (cx >= mb_size/2 - x_mb && cx + mb_size*3/2 <= width - x_mb && cy >= mb_size/2 - y_mb && cy + mb_size*3/2 <= height - y_mb) { + uint64_t cost = mb_sbad(tex_cur, tex_ref, x_mb, y_mb, cx, cy, mb_size, width, height) + (abs(cx) + abs(cy)) * COST_PRED_SCALE; + if (cost < best_cost) { best_cost = cost; best_mv_y = (int16_t)cy; } + } + cy = best_mv_y - d; + if (cx >= mb_size/2 - x_mb && cx + mb_size*3/2 <= width - x_mb && cy >= mb_size/2 - y_mb && cy + mb_size*3/2 <= height - y_mb) { + uint64_t cost = mb_sbad(tex_cur, tex_ref, x_mb, y_mb, cx, cy, mb_size, width, height) + (abs(cx) + abs(cy)) * COST_PRED_SCALE; + if (cost < best_cost) { best_cost = cost; best_mv_y = (int16_t)cy; } + } + } + + /* 5x5 full search refinement around cross winner */ + for (int dy = -2; dy <= 2; dy++) { + for (int dx = -2; dx <= 2; dx++) { + int cx = best_mv_x + dx, cy = best_mv_y + dy; + int x_mv = x_mb + cx, y_mv = y_mb + cy; + if (x_mv < mb_size/2 || x_mv + mb_size*3/2 > width || + y_mv < mb_size/2 || y_mv + mb_size*3/2 > height) + continue; + uint64_t cost = mb_sbad(tex_cur, tex_ref, x_mb, y_mb, x_mv - x_mb, y_mv - y_mb, mb_size, width, height) + + (abs(cx) + abs(cy)) * COST_PRED_SCALE; + if (cost < best_cost) { best_cost = cost; best_mv_x = (int16_t)cx; best_mv_y = (int16_t)cy; } + } + } + + /* Uneven multi-hexagon grid */ + for (int d = 1; d <= search_param / 4; d++) { + for (int k = 1; k < 16; k++) { /* skip k=0 (duplicate of diamond) */ + int cx = best_mv_x + d_hex4_umh[k][0] * d; + int cy = best_mv_y + d_hex4_umh[k][1] * d; + int x_mv = x_mb + cx, y_mv = y_mb + cy; + if (x_mv < mb_size/2 || x_mv + mb_size*3/2 > width || + y_mv < mb_size/2 || y_mv + mb_size*3/2 > height) + continue; + uint64_t cost = mb_sbad(tex_cur, tex_ref, x_mb, y_mb, x_mv - x_mb, y_mv - y_mb, mb_size, width, height) + + (abs(cx) + abs(cy)) * COST_PRED_SCALE; + if (cost < best_cost) { best_cost = cost; best_mv_x = (int16_t)cx; best_mv_y = (int16_t)cy; } + } + } + + /* Extended hexagon search */ + for (;;) { + int improved = 0; + for (int k = 0; k < 6; k++) { + int cx = best_mv_x + d_hex2[k][0]; + int cy = best_mv_y + d_hex2[k][1]; + int x_mv = x_mb + cx, y_mv = y_mb + cy; + if (x_mv < mb_size/2 || x_mv + mb_size*3/2 > width || + y_mv < mb_size/2 || y_mv + mb_size*3/2 > height) + continue; + uint64_t cost = mb_sbad(tex_cur, tex_ref, x_mb, y_mb, x_mv - x_mb, y_mv - y_mb, mb_size, width, height) + + (abs(cx) + abs(cy)) * COST_PRED_SCALE; + if (cost < best_cost) { best_cost = cost; best_mv_x = (int16_t)cx; best_mv_y = (int16_t)cy; improved = 1; } + } + if (!improved) break; + } + + /* Diamond refinement */ + { + int improved; + do { + improved = 0; + for (int k = 0; k < 4; k++) { + int cx = best_mv_x + d_dia1[k][0]; + int cy = best_mv_y + d_dia1[k][1]; + int x_mv = x_mb + cx, y_mv = y_mb + cy; + if (x_mv < mb_size/2 || x_mv + mb_size*3/2 > width || + y_mv < mb_size/2 || y_mv + mb_size*3/2 > height) + continue; + uint64_t cost = mb_sbad(tex_cur, tex_ref, x_mb, y_mb, x_mv - x_mb, y_mv - y_mb, mb_size, width, height) + + (abs(cx) + abs(cy)) * COST_PRED_SCALE; + if (cost < best_cost) { best_cost = cost; best_mv_x = (int16_t)cx; best_mv_y = (int16_t)cy; improved = 1; } + } + } while (improved); + } + + mv_buffer_out[mb_idx * 2] = best_mv_x; + mv_buffer_out[mb_idx * 2 + 1] = best_mv_y; +} + +/* =================================================================== + * UMH Pass 2 — spatial predictors from Pass1 + full UMH search + * Grid: (b_count, 1, 1) Block: (1, 1, 1) + * =================================================================== */ +__global__ void minterpolate_me_umh_pass2( + cudaTextureObject_t tex_cur, + cudaTextureObject_t tex_ref, + int16_t *__restrict__ mv_buffer_final, + const int16_t *__restrict__ mv_buffer_pass1, + const int16_t *__restrict__ mv_table_prev1, + int b_width, int b_height, + int mb_size, int search_param, + int width, int height) +{ + int mb_idx = blockIdx.x; + if (mb_idx >= b_width * b_height) return; + + int mb_x = mb_idx % b_width; + int mb_y = mb_idx / b_width; + int x_mb = mb_x * mb_size; + int y_mb = mb_y * mb_size; + + + /* Spatial predictors from pass1: (0,0), left, top, top-r, median */ + int16_t preds[6][2]; + int nb = 0; + preds[nb][0] = 0; preds[nb][1] = 0; nb++; + if (mb_x > 0) { + int ni = mb_idx - 1; + preds[nb][0] = mv_buffer_pass1[ni * 2]; + preds[nb][1] = mv_buffer_pass1[ni * 2 + 1]; nb++; + } + if (mb_y > 0) { + int ni = mb_idx - b_width; + preds[nb][0] = mv_buffer_pass1[ni * 2]; + preds[nb][1] = mv_buffer_pass1[ni * 2 + 1]; nb++; + if (mb_x + 1 < b_width) { + ni = mb_idx - b_width + 1; + preds[nb][0] = mv_buffer_pass1[ni * 2]; + preds[nb][1] = mv_buffer_pass1[ni * 2 + 1]; nb++; + } else if (mb_x > 0) { + ni = mb_idx - b_width - 1; + preds[nb][0] = mv_buffer_pass1[ni * 2]; + preds[nb][1] = mv_buffer_pass1[ni * 2 + 1]; nb++; + } + } + + /* Median */ + int med_x = 0, med_y = 0; + if (nb == 5) { + med_x = mid_pred(preds[1][0], preds[2][0], preds[3][0]); + med_y = mid_pred(preds[1][1], preds[2][1], preds[3][1]); + } else if (nb == 4) { + med_x = mid_pred(0, preds[1][0], preds[2][0]); + med_y = mid_pred(0, preds[1][1], preds[2][1]); + } else if (nb == 3) { + med_x = preds[1][0]; med_y = preds[1][1]; + } + if (med_x != 0 || med_y != 0) { + preds[nb][0] = (int16_t)med_x; preds[nb][1] = (int16_t)med_y; nb++; + } + + /* Eval best predictor */ + int16_t best_mv_x = 0, best_mv_y = 0; + uint64_t best_cost = 0xFFFFFFFFFFFFFFFFULL; + for (int p = 0; p < nb; p++) { + int cx = preds[p][0], cy = preds[p][1]; + int x_mv = x_mb + cx, y_mv = y_mb + cy; + if (x_mv < mb_size/2 || x_mv + mb_size*3/2 > width || + y_mv < mb_size/2 || y_mv + mb_size*3/2 > height) + continue; + uint64_t cost = mb_sbad(tex_cur, tex_ref, x_mb, y_mb, x_mv - x_mb, y_mv - y_mb, mb_size, width, height) + + (abs(cx) + abs(cy)) * COST_PRED_SCALE; + if (cost < best_cost) { best_cost = cost; best_mv_x = (int16_t)cx; best_mv_y = (int16_t)cy; } + } + + /* Same UMH search as pass1 */ + #define UMH_SEARCH() do { \ + for (int d = 1; d <= search_param; d += 2) { \ + for (int sign = -1; sign <= 1; sign += 2) { \ + int cx = best_mv_x + sign * d, cy = best_mv_y; \ + if (x_mb + cx >= mb_size/2 && x_mb + cx + mb_size*3/2 <= width && \ + y_mb + cy >= mb_size/2 && y_mb + cy + mb_size*3/2 <= height) { \ + uint64_t cost = mb_sbad(tex_cur, tex_ref, x_mb, y_mb, cx, cy, mb_size, width, height) \ + + (abs(cx) + abs(cy)) * COST_PRED_SCALE; \ + if (cost < best_cost) { best_cost = cost; best_mv_x = (int16_t)cx; best_mv_y = (int16_t)cy; } \ + } \ + } \ + } \ + for (int d = 1; d <= search_param / 2; d += 2) { \ + for (int sign = -1; sign <= 1; sign += 2) { \ + int cx = best_mv_x, cy = best_mv_y + sign * d; \ + if (x_mb + cx >= mb_size/2 && x_mb + cx + mb_size*3/2 <= width && \ + y_mb + cy >= mb_size/2 && y_mb + cy + mb_size*3/2 <= height) { \ + uint64_t cost = mb_sbad(tex_cur, tex_ref, x_mb, y_mb, cx, cy, mb_size, width, height) \ + + (abs(cx) + abs(cy)) * COST_PRED_SCALE; \ + if (cost < best_cost) { best_cost = cost; best_mv_y = (int16_t)cy; } \ + } \ + } \ + } \ + for (int dy = -2; dy <= 2; dy++) \ + for (int dx = -2; dx <= 2; dx++) { \ + int cx = best_mv_x + dx, cy = best_mv_y + dy; \ + int x_mv = x_mb + cx, y_mv = y_mb + cy; \ + if (x_mv < mb_size/2 || x_mv + mb_size*3/2 > width || y_mv < mb_size/2 || y_mv + mb_size*3/2 > height) continue; \ + uint64_t cost = mb_sbad(tex_cur, tex_ref, x_mb, y_mb, x_mv - x_mb, y_mv - y_mb, mb_size, width, height) \ + + (abs(cx) + abs(cy)) * COST_PRED_SCALE; \ + if (cost < best_cost) { best_cost = cost; best_mv_x = (int16_t)cx; best_mv_y = (int16_t)cy; } \ + } \ + for (int d = 1; d <= search_param / 4; d++) \ + for (int k = 1; k < 16; k++) { \ + int cx = best_mv_x + d_hex4_umh[k][0] * d; \ + int cy = best_mv_y + d_hex4_umh[k][1] * d; \ + int x_mv = x_mb + cx, y_mv = y_mb + cy; \ + if (x_mv < mb_size/2 || x_mv + mb_size*3/2 > width || y_mv < mb_size/2 || y_mv + mb_size*3/2 > height) continue; \ + uint64_t cost = mb_sbad(tex_cur, tex_ref, x_mb, y_mb, x_mv - x_mb, y_mv - y_mb, mb_size, width, height) \ + + (abs(cx) + abs(cy)) * COST_PRED_SCALE; \ + if (cost < best_cost) { best_cost = cost; best_mv_x = (int16_t)cx; best_mv_y = (int16_t)cy; } \ + } \ + for (;;) { \ + int improved = 0; \ + for (int k = 0; k < 6; k++) { \ + int cx = best_mv_x + d_hex2[k][0], cy = best_mv_y + d_hex2[k][1]; \ + int x_mv = x_mb + cx, y_mv = y_mb + cy; \ + if (x_mv < mb_size/2 || x_mv + mb_size*3/2 > width || y_mv < mb_size/2 || y_mv + mb_size*3/2 > height) continue; \ + uint64_t cost = mb_sbad(tex_cur, tex_ref, x_mb, y_mb, x_mv - x_mb, y_mv - y_mb, mb_size, width, height) \ + + (abs(cx) + abs(cy)) * COST_PRED_SCALE; \ + if (cost < best_cost) { best_cost = cost; best_mv_x = (int16_t)cx; best_mv_y = (int16_t)cy; improved = 1; } \ + } \ + if (!improved) break; \ + } \ + { \ + int improved; \ + do { improved = 0; \ + for (int k = 0; k < 4; k++) { \ + int cx = best_mv_x + d_dia1[k][0], cy = best_mv_y + d_dia1[k][1]; \ + int x_mv = x_mb + cx, y_mv = y_mb + cy; \ + if (x_mv < mb_size/2 || x_mv + mb_size*3/2 > width || y_mv < mb_size/2 || y_mv + mb_size*3/2 > height) continue; \ + uint64_t cost = mb_sbad(tex_cur, tex_ref, x_mb, y_mb, x_mv - x_mb, y_mv - y_mb, mb_size, width, height) \ + + (abs(cx) + abs(cy)) * COST_PRED_SCALE; \ + if (cost < best_cost) { best_cost = cost; best_mv_x = (int16_t)cx; best_mv_y = (int16_t)cy; improved = 1; } \ + } \ + } while (improved); \ + } \ + } while(0) + + UMH_SEARCH(); + + mv_buffer_final[mb_idx * 2] = best_mv_x; + mv_buffer_final[mb_idx * 2 + 1] = best_mv_y; +} + +/* =================================================================== + * OBMC — standard bilat OBMC (existing, kept for bilat mode) + * =================================================================== */ +__global__ void minterpolate_obmc( + cudaTextureObject_t tex_prev, + cudaTextureObject_t tex_next, + const int16_t *__restrict__ mv_buffer, + const int16_t *__restrict__ sub_mv_buffer, + const int *__restrict__ sub_valid, + uchar *__restrict__ dst, + int width, int height, int dst_pitch, + int b_width, int b_height, + int mb_size, int log2_mb_size, + int alpha, + int chroma_shift_w, int chroma_shift_h) +{ + int x = blockIdx.x * blockDim.x + threadIdx.x; + int y = blockIdx.y * blockDim.y + threadIdx.y; + if (x >= width || y >= height) return; + + int weight_sum = 0, val = 0; + int half_mb = mb_size / 2; + int luma_x = x << chroma_shift_w; + int luma_y = y << chroma_shift_h; + int mb_x0 = luma_x / mb_size; + int mb_y0 = luma_y / mb_size; + + for (int dy = -1; dy <= 0; dy++) { + int mb_y = mb_y0 + dy; + if (mb_y < 0 || mb_y >= b_height) continue; + for (int dx = -1; dx <= 0; dx++) { + int mb_x = mb_x0 + dx; + if (mb_x < 0 || mb_x >= b_width) continue; + + int start_lx = (mb_x << log2_mb_size) - half_mb; + int start_ly = (mb_y << log2_mb_size) - half_mb; + int rx = luma_x - start_lx, ry = luma_y - start_ly; + if (rx < 0 || rx >= 2 * mb_size || ry < 0 || ry >= 2 * mb_size) continue; + + uchar obmc_weight = obmc_weight_lut(rx, ry, log2_mb_size); + if (obmc_weight == 0) continue; + + int mb_idx = mb_x + mb_y * b_width; + int mv_lx, mv_ly; + if (sub_valid && sub_valid[mb_idx]) { + int sub_x = (luma_x - (mb_x << log2_mb_size)) >> (log2_mb_size - 1); + int sub_y = (luma_y - (mb_y << log2_mb_size)) >> (log2_mb_size - 1); + int si = sub_x + sub_y * 2; + mv_lx = sub_mv_buffer[(mb_idx * 4 + si) * 2]; + mv_ly = sub_mv_buffer[(mb_idx * 4 + si) * 2 + 1]; + } else { + mv_lx = mv_buffer[mb_idx * 2]; + mv_ly = mv_buffer[mb_idx * 2 + 1]; + } + int mv_x = mv_lx >> chroma_shift_w; + int mv_y = mv_ly >> chroma_shift_h; + + /* Frame[1] */ + { + int src_x = x + (mv_x * alpha) / ALPHA_MAX; + int src_y = y + (mv_y * alpha) / ALPHA_MAX; + src_x = max(0, min(width - 1, src_x)); + src_y = max(0, min(height - 1, src_y)); + int w = obmc_weight * (ALPHA_MAX - alpha); + val += w * tex2D<uchar>(tex_prev, src_x, src_y); + weight_sum += w; + } + /* Frame[2] */ + { + int src_x = x - (mv_x * (ALPHA_MAX - alpha)) / ALPHA_MAX; + int src_y = y - (mv_y * (ALPHA_MAX - alpha)) / ALPHA_MAX; + src_x = max(0, min(width - 1, src_x)); + src_y = max(0, min(height - 1, src_y)); + int w = obmc_weight * alpha; + val += w * tex2D<uchar>(tex_next, src_x, src_y); + weight_sum += w; + } + } + } + if (weight_sum > 0) + val = (val + weight_sum / 2) / weight_sum; + dst[y * dst_pitch + x] = (uchar)val; +} + +/* =================================================================== + * OBMC BIDIR — bidirectional OBMC: 2 MV sets, 8 contributions per pixel + * Grid: 2D (BLOCKX, BLOCKY) + * =================================================================== */ +/* =================================================================== + * OBMC BIDIR — matches CPU bidirectional_obmc: window-shift, 2 frames + * Grid: 2D (BLOCKX, BLOCKY) + * =================================================================== */ +__global__ void minterpolate_obmc_bidir( + cudaTextureObject_t tex_prev, /* frame[1] */ + cudaTextureObject_t tex_cur, /* frame[2] */ + const int16_t *__restrict__ mv_buffer_dir0, /* frame[2]→frame[1] */ + const int16_t *__restrict__ mv_buffer_fwd, /* frame[1]→frame[2] prev */ + uchar *__restrict__ dst, + int width, int height, int dst_pitch, + int b_width, int b_height, + int mb_size, int log2_mb_size, + int alpha, + int chroma_shift_w, int chroma_shift_h) +{ + int x = blockIdx.x * blockDim.x + threadIdx.x; + int y = blockIdx.y * blockDim.y + threadIdx.y; + if (x >= width || y >= height) return; + + int weight_sum = 0, val = 0; + int half_mb = mb_size / 2; + int luma_x = x << chroma_shift_w; + int luma_y = y << chroma_shift_h; + int mb_x0 = luma_x / mb_size; + int mb_y0 = luma_y / mb_size; + + for (int dy = -1; dy <= 0; dy++) { + int mb_y = mb_y0 + dy; + if (mb_y < 0 || mb_y >= b_height) continue; + for (int dx = -1; dx <= 0; dx++) { + int mb_x = mb_x0 + dx; + if (mb_x < 0 || mb_x >= b_width) continue; + + int mb_idx = mb_x + mb_y * b_width; + int mb_ox = mb_x << log2_mb_size; + int mb_oy = mb_y << log2_mb_size; + + /* --- Dir 0: frame[2]→frame[1], window shift=(ALPHA_MAX-alpha) --- */ + { + int mv_lx = mv_buffer_dir0[mb_idx * 2]; + int mv_ly = mv_buffer_dir0[mb_idx * 2 + 1]; + int mv_x = mv_lx >> chroma_shift_w; + int mv_y = mv_ly >> chroma_shift_h; + + int ws_x = mb_ox - half_mb + mv_x * (ALPHA_MAX - alpha) / ALPHA_MAX; + int ws_y = mb_oy - half_mb + mv_y * (ALPHA_MAX - alpha) / ALPHA_MAX; + int rx = luma_x - ws_x; + int ry = luma_y - ws_y; + + if (rx >= 0 && rx < 2 * mb_size && ry >= 0 && ry < 2 * mb_size) { + uchar w = obmc_weight_lut(rx, ry, log2_mb_size); + if (w > 0) { + int sx = max(0, min(width - 1, x + mv_x)); + int sy = max(0, min(height - 1, y + mv_y)); + val += w * tex2D<uchar>(tex_cur, sx, sy); + weight_sum += w; + } + } + } + + /* --- Dir 1: frame[1]→frame[2] (prev fwd), window shift=alpha --- */ + { + int mv_lx = mv_buffer_fwd[mb_idx * 2]; + int mv_ly = mv_buffer_fwd[mb_idx * 2 + 1]; + int mv_x = mv_lx >> chroma_shift_w; + int mv_y = mv_ly >> chroma_shift_h; + + int ws_x = mb_ox - half_mb + mv_x * alpha / ALPHA_MAX; + int ws_y = mb_oy - half_mb + mv_y * alpha / ALPHA_MAX; + int rx = luma_x - ws_x; + int ry = luma_y - ws_y; + + if (rx >= 0 && rx < 2 * mb_size && ry >= 0 && ry < 2 * mb_size) { + uchar w = obmc_weight_lut(rx, ry, log2_mb_size); + if (w > 0) { + /* Negated MV for dir=1 (matching CPU) */ + int sx = max(0, min(width - 1, x - mv_x)); + int sy = max(0, min(height - 1, y - mv_y)); + val += w * tex2D<uchar>(tex_prev, sx, sy); + weight_sum += w; + } + } + } + } + } + if (weight_sum > 0) + val = (val + weight_sum / 2) / weight_sum; + dst[y * dst_pitch + x] = (uchar)val; +} + +/* =================================================================== + * Blend kernel + * =================================================================== */ +__global__ void minterpolate_blend( + cudaTextureObject_t tex_prev, + cudaTextureObject_t tex_next, + uchar *__restrict__ dst, + int width, int height, int dst_pitch, + int alpha) +{ + int x = blockIdx.x * blockDim.x + threadIdx.x; + int y = blockIdx.y * blockDim.y + threadIdx.y; + if (x >= width || y >= height) return; + + int v_prev = tex2D<uchar>(tex_prev, x, y); + int v_next = tex2D<uchar>(tex_next, x, y); + int val = ((ALPHA_MAX - alpha) * v_prev + alpha * v_next + 512) >> 10; + dst[y * dst_pitch + x] = (uchar)max(0, min(255, val)); +} + +/* =================================================================== + * Scene SAD — Stage 1: per-block reduction + * =================================================================== */ +__global__ void minterpolate_scene_sad( + cudaTextureObject_t tex_a, + cudaTextureObject_t tex_b, + int width, int height, + uint64_t *__restrict__ block_sad) +{ + __shared__ uint64_t shared[512]; + int x = blockIdx.x * blockDim.x + threadIdx.x; + int y = blockIdx.y * blockDim.y + threadIdx.y; + int tid = threadIdx.y * blockDim.x + threadIdx.x; + + uint64_t sad = 0; + if (x < width && y < height) { + int a = tex2D<uchar>(tex_a, x, y); + int b = tex2D<uchar>(tex_b, x, y); + sad = (uint64_t)abs(a - b); + } + shared[tid] = sad; + __syncthreads(); + for (int s = 256; s > 0; s >>= 1) { + if (tid < s) shared[tid] += shared[tid + s]; + __syncthreads(); + } + if (tid == 0) { + int bid = blockIdx.y * gridDim.x + blockIdx.x; + block_sad[bid] = shared[0]; + } +} + +/* =================================================================== + * Scene SAD — Stage 2: final reduction to single value + * =================================================================== */ +__global__ void minterpolate_scene_sad_reduce( + uint64_t *__restrict__ block_sad, + int num_blocks, + uint64_t *__restrict__ total_sad) +{ + __shared__ uint64_t shared[256]; + int tid = threadIdx.x; + uint64_t sum = 0; + for (int i = tid; i < num_blocks; i += blockDim.x) + sum += block_sad[i]; + shared[tid] = sum; + __syncthreads(); + for (int s = 128; s > 0; s >>= 1) { + if (tid < s) shared[tid] += shared[tid + s]; + __syncthreads(); + } + if (tid == 0) *total_sad = shared[0]; +} + +/* =================================================================== + * SBAD — Bilateral block matching cost for AOBMC + * Grid: (b_count, 1, 1) Block: (256, 1, 1) + * =================================================================== */ +__global__ void minterpolate_sbad( + cudaTextureObject_t tex_cur, /* frame[1] */ + cudaTextureObject_t tex_ref, /* frame[2] */ + const int16_t *__restrict__ mv_buffer, + uint64_t *__restrict__ sbad_buffer, + int b_width, int b_height, + int mb_size, int log2_mb_size, + int width, int height) +{ + int mb_idx = blockIdx.x; + if (mb_idx >= b_width * b_height) return; + + int mb_x = mb_idx % b_width; + int mb_y = mb_idx / b_width; + int x_mb = mb_x * mb_size; + int y_mb = mb_y * mb_size; + + int mv_x = mv_buffer[mb_idx * 2]; + int mv_y = mv_buffer[mb_idx * 2 + 1]; + + /* Use cooperative SAD if multiple threads, else single-thread it */ + uint64_t sbad = 0; + int half = mb_size / 2; + if (blockDim.x > 1) { + /* Cooperative: each thread handles a subset of rows */ + int total_rows = mb_size * 2; + for (int row = threadIdx.x; row < total_rows; row += blockDim.x) { + int j = row - half; + int yc = y_mb + mv_y + j; + int yr = y_mb - mv_y + j; + yc = max(0, min(height - 1, yc)); + yr = max(0, min(height - 1, yr)); + for (int i = -half; i < mb_size * 3 / 2; i++) { + int xc = x_mb + mv_x + i; + int xr = x_mb - mv_x + i; + xc = max(0, min(width - 1, xc)); + xr = max(0, min(width - 1, xr)); + sbad += abs((int)tex2D<uchar>(tex_cur, xc, yc) - + (int)tex2D<uchar>(tex_ref, xr, yr)); + } + } + /* Warp reduction */ + for (int s = 16; s > 0; s >>= 1) { + sbad += __shfl_xor_sync(0xFFFFFFFF, sbad, s); + } + if (threadIdx.x == 0) + sbad_buffer[mb_idx] = sbad; + } else { + /* Single-thread fallback */ + for (int j = -half; j < mb_size * 3 / 2; j++) { + int yc = y_mb + mv_y + j; + int yr = y_mb - mv_y + j; + yc = max(0, min(height - 1, yc)); + yr = max(0, min(height - 1, yr)); + for (int i = -half; i < mb_size * 3 / 2; i++) { + int xc = x_mb + mv_x + i; + int xr = x_mb - mv_x + i; + xc = max(0, min(width - 1, xc)); + xr = max(0, min(width - 1, xr)); + sbad += abs((int)tex2D<uchar>(tex_cur, xc, yc) - + (int)tex2D<uchar>(tex_ref, xr, yr)); + } + } + sbad_buffer[mb_idx] = sbad; + } +} + +/* =================================================================== + * OBMC AOBMC — adaptive OBMC with sbad-based weight adjustment + * Grid: 2D (BLOCKX, BLOCKY) + * =================================================================== */ +__global__ void minterpolate_obmc_aobmc( + cudaTextureObject_t tex_prev, + cudaTextureObject_t tex_next, + const int16_t *__restrict__ mv_buffer, + const int16_t *__restrict__ sub_mv_buffer, + const int *__restrict__ sub_valid, + const uint64_t *__restrict__ sbad_buffer, + uchar *__restrict__ dst, + int width, int height, int dst_pitch, + int b_width, int b_height, + int mb_size, int log2_mb_size, + int alpha, + int chroma_shift_w, int chroma_shift_h) +{ + int x = blockIdx.x * blockDim.x + threadIdx.x; + int y = blockIdx.y * blockDim.y + threadIdx.y; + if (x >= width || y >= height) return; + + int weight_sum = 0, val = 0; + int half_mb = mb_size / 2; + int luma_x = x << chroma_shift_w; + int luma_y = y << chroma_shift_h; + int mb_x0 = luma_x / mb_size; + int mb_y0 = luma_y / mb_size; + + /* Determine which "sub-quadrant" of the OBMC window this pixel is in */ + int base_start_lx = (mb_x0 << log2_mb_size) - half_mb; + int base_start_ly = (mb_y0 << log2_mb_size) - half_mb; + int base_rx = luma_x - base_start_lx; + int base_ry = luma_y - base_start_ly; + + /* Map to neighbor offsets (-1, 0) — matches CPU AOBMC nb_x/nb_y logic */ + int nb_x = (((base_rx >> (log2_mb_size - 1)) * 2 - 3) / 2); + int nb_y = (((base_ry >> (log2_mb_size - 1)) * 2 - 3) / 2); + /* Clamp to valid range */ + if (mb_x0 + nb_x < 0 || mb_x0 + nb_x >= b_width) nb_x = 0; + if (mb_y0 + nb_y < 0 || mb_y0 + nb_y >= b_height) nb_y = 0; + + int nb_mb_idx = (mb_x0 + nb_x) + (mb_y0 + nb_y) * b_width; + uint64_t nb_sbad = sbad_buffer[nb_mb_idx]; + uint64_t cur_sbad = sbad_buffer[mb_x0 + mb_y0 * b_width]; + + for (int dy = -1; dy <= 0; dy++) { + int mb_y = mb_y0 + dy; + if (mb_y < 0 || mb_y >= b_height) continue; + for (int dx = -1; dx <= 0; dx++) { + int mb_x = mb_x0 + dx; + if (mb_x < 0 || mb_x >= b_width) continue; + + int start_lx = (mb_x << log2_mb_size) - half_mb; + int start_ly = (mb_y << log2_mb_size) - half_mb; + int rx = luma_x - start_lx, ry = luma_y - start_ly; + if (rx < 0 || rx >= 2 * mb_size || ry < 0 || ry >= 2 * mb_size) continue; + + uchar obmc_weight = obmc_weight_lut(rx, ry, log2_mb_size); + if (obmc_weight == 0) continue; + + /* AOBMC weight adjustment */ + if ((dx != 0 || dy != 0) && /* this is a neighbor block */ + (nb_x != 0 || nb_y != 0) && /* pixel is in overlapping region */ + nb_sbad > 0 && cur_sbad > 0) { + int phi = min(ALPHA_MAX, (int)(ALPHA_MAX * nb_sbad / cur_sbad)); + obmc_weight = (uchar)((int)obmc_weight * phi / ALPHA_MAX); + } + + int mb_idx = mb_x + mb_y * b_width; + int mv_lx, mv_ly; + if (sub_valid && sub_valid[mb_idx]) { + int sub_x = (luma_x - (mb_x << log2_mb_size)) >> (log2_mb_size - 1); + int sub_y = (luma_y - (mb_y << log2_mb_size)) >> (log2_mb_size - 1); + int si = sub_x + sub_y * 2; + mv_lx = sub_mv_buffer[(mb_idx * 4 + si) * 2]; + mv_ly = sub_mv_buffer[(mb_idx * 4 + si) * 2 + 1]; + } else { + mv_lx = mv_buffer[mb_idx * 2]; + mv_ly = mv_buffer[mb_idx * 2 + 1]; + } + int mv_x = mv_lx >> chroma_shift_w; + int mv_y = mv_ly >> chroma_shift_h; + + /* Frame[1] */ + { + int src_x = x + (mv_x * alpha) / ALPHA_MAX; + int src_y = y + (mv_y * alpha) / ALPHA_MAX; + src_x = max(0, min(width - 1, src_x)); + src_y = max(0, min(height - 1, src_y)); + int w = obmc_weight * (ALPHA_MAX - alpha); + val += w * tex2D<uchar>(tex_prev, src_x, src_y); + weight_sum += w; + } + /* Frame[2] */ + { + int src_x = x - (mv_x * (ALPHA_MAX - alpha)) / ALPHA_MAX; + int src_y = y - (mv_y * (ALPHA_MAX - alpha)) / ALPHA_MAX; + src_x = max(0, min(width - 1, src_x)); + src_y = max(0, min(height - 1, src_y)); + int w = obmc_weight * alpha; + val += w * tex2D<uchar>(tex_next, src_x, src_y); + weight_sum += w; + } + } + } + if (weight_sum > 0) + val = (val + weight_sum / 2) / weight_sum; + dst[y * dst_pitch + x] = (uchar)val; +} + +/* =================================================================== + * uchar2 kernel variants (NV12 chroma) — OBMC, blend, AOBMC, BIDIR + * =================================================================== */ + +__global__ void minterpolate_obmc_uchar2( + cudaTextureObject_t tex_prev, + cudaTextureObject_t tex_next, + const int16_t *__restrict__ mv_buffer, + const int16_t *__restrict__ sub_mv_buffer, + const int *__restrict__ sub_valid, + uchar2 *__restrict__ dst, + int width, int height, int dst_pitch, + int b_width, int b_height, + int mb_size, int log2_mb_size, + int alpha, + int chroma_shift_w, int chroma_shift_h) +{ + int x = blockIdx.x * blockDim.x + threadIdx.x; + int y = blockIdx.y * blockDim.y + threadIdx.y; + if (x >= width || y >= height) return; + + int weight_sum = 0, val_u = 0, val_v = 0; + int half_mb = mb_size / 2; + int luma_x = x << chroma_shift_w; + int luma_y = y << chroma_shift_h; + int mb_x0 = luma_x / mb_size; + int mb_y0 = luma_y / mb_size; + + for (int dy = -1; dy <= 0; dy++) { + int mb_y = mb_y0 + dy; + if (mb_y < 0 || mb_y >= b_height) continue; + for (int dx = -1; dx <= 0; dx++) { + int mb_x = mb_x0 + dx; + if (mb_x < 0 || mb_x >= b_width) continue; + int start_lx = (mb_x << log2_mb_size) - half_mb; + int start_ly = (mb_y << log2_mb_size) - half_mb; + int rx = luma_x - start_lx, ry = luma_y - start_ly; + if (rx < 0 || rx >= 2 * mb_size || ry < 0 || ry >= 2 * mb_size) continue; + uchar obmc_weight = obmc_weight_lut(rx, ry, log2_mb_size); + if (obmc_weight == 0) continue; + + int mb_idx = mb_x + mb_y * b_width; + int mv_lx, mv_ly; + if (sub_valid && sub_valid[mb_idx]) { + int sub_x = (luma_x - (mb_x << log2_mb_size)) >> (log2_mb_size - 1); + int sub_y = (luma_y - (mb_y << log2_mb_size)) >> (log2_mb_size - 1); + int si = sub_x + sub_y * 2; + mv_lx = sub_mv_buffer[(mb_idx * 4 + si) * 2]; + mv_ly = sub_mv_buffer[(mb_idx * 4 + si) * 2 + 1]; + } else { + mv_lx = mv_buffer[mb_idx * 2]; + mv_ly = mv_buffer[mb_idx * 2 + 1]; + } + int mv_x = mv_lx >> chroma_shift_w; + int mv_y = mv_ly >> chroma_shift_h; + + /* Frame[1] */ + { + int src_x = x + (mv_x * alpha) / ALPHA_MAX; + int src_y = y + (mv_y * alpha) / ALPHA_MAX; + src_x = max(0, min(width - 1, src_x)); + src_y = max(0, min(height - 1, src_y)); + uchar2 p = tex2D<uchar2>(tex_prev, src_x, src_y); + int w = obmc_weight * (ALPHA_MAX - alpha); + val_u += w * p.x; val_v += w * p.y; + weight_sum += w; + } + /* Frame[2] */ + { + int src_x = x - (mv_x * (ALPHA_MAX - alpha)) / ALPHA_MAX; + int src_y = y - (mv_y * (ALPHA_MAX - alpha)) / ALPHA_MAX; + src_x = max(0, min(width - 1, src_x)); + src_y = max(0, min(height - 1, src_y)); + uchar2 p = tex2D<uchar2>(tex_next, src_x, src_y); + int w = obmc_weight * alpha; + val_u += w * p.x; val_v += w * p.y; + weight_sum += w; + } + } + } + if (weight_sum > 0) { + val_u = (val_u + weight_sum / 2) / weight_sum; + val_v = (val_v + weight_sum / 2) / weight_sum; + } + dst[y * dst_pitch + x] = make_uchar2((uchar)val_u, (uchar)val_v); +} + +__global__ void minterpolate_blend_uchar2( + cudaTextureObject_t tex_prev, + cudaTextureObject_t tex_next, + uchar2 *__restrict__ dst, + int width, int height, int dst_pitch, + int alpha) +{ + int x = blockIdx.x * blockDim.x + threadIdx.x; + int y = blockIdx.y * blockDim.y + threadIdx.y; + if (x >= width || y >= height) return; + uchar2 p_prev = tex2D<uchar2>(tex_prev, x, y); + uchar2 p_next = tex2D<uchar2>(tex_next, x, y); + int vu = ((ALPHA_MAX - alpha) * p_prev.x + alpha * p_next.x + 512) >> 10; + int vv = ((ALPHA_MAX - alpha) * p_prev.y + alpha * p_next.y + 512) >> 10; + dst[y * dst_pitch + x] = make_uchar2((uchar)max(0, min(255, vu)), + (uchar)max(0, min(255, vv))); +} + +/* BIDIR uchar2 variant — matches CPU bidirectional_obmc: window-shift */ +__global__ void minterpolate_obmc_uchar2_bidir( + cudaTextureObject_t tex_prev, /* frame[1] */ + cudaTextureObject_t tex_cur, /* frame[2] */ + const int16_t *__restrict__ mv_buffer_dir0, /* frame[2]→frame[1] */ + const int16_t *__restrict__ mv_buffer_fwd, /* frame[1]→frame[2] prev */ + uchar2 *__restrict__ dst, + int width, int height, int dst_pitch, + int b_width, int b_height, + int mb_size, int log2_mb_size, + int alpha, + int chroma_shift_w, int chroma_shift_h) +{ + int x = blockIdx.x * blockDim.x + threadIdx.x; + int y = blockIdx.y * blockDim.y + threadIdx.y; + if (x >= width || y >= height) return; + + int weight_sum = 0, val_u = 0, val_v = 0; + int half_mb = mb_size / 2; + int luma_x = x << chroma_shift_w; + int luma_y = y << chroma_shift_h; + int mb_x0 = luma_x / mb_size; + int mb_y0 = luma_y / mb_size; + + for (int dy = -1; dy <= 0; dy++) { + int mb_y = mb_y0 + dy; + if (mb_y < 0 || mb_y >= b_height) continue; + for (int dx = -1; dx <= 0; dx++) { + int mb_x = mb_x0 + dx; + if (mb_x < 0 || mb_x >= b_width) continue; + + int mb_idx = mb_x + mb_y * b_width; + int mb_ox = mb_x << log2_mb_size; + int mb_oy = mb_y << log2_mb_size; + + /* Dir 0: frame[2]→frame[1], window shift=(ALPHA_MAX-alpha) */ + { + int mv_lx = mv_buffer_dir0[mb_idx * 2]; + int mv_ly = mv_buffer_dir0[mb_idx * 2 + 1]; + int mv_x = mv_lx >> chroma_shift_w; + int mv_y = mv_ly >> chroma_shift_h; + + int ws_x = mb_ox - half_mb + mv_x * (ALPHA_MAX - alpha) / ALPHA_MAX; + int ws_y = mb_oy - half_mb + mv_y * (ALPHA_MAX - alpha) / ALPHA_MAX; + int rx = luma_x - ws_x; + int ry = luma_y - ws_y; + + if (rx >= 0 && rx < 2 * mb_size && ry >= 0 && ry < 2 * mb_size) { + uchar w = obmc_weight_lut(rx, ry, log2_mb_size); + if (w > 0) { + int sx = max(0, min(width - 1, x + mv_x)); + int sy = max(0, min(height - 1, y + mv_y)); + uchar2 p = tex2D<uchar2>(tex_cur, sx, sy); + val_u += w * p.x; val_v += w * p.y; weight_sum += w; + } + } + } + + /* Dir 1: frame[1]→frame[2] (prev fwd), window shift=alpha, negated */ + { + int mv_lx = mv_buffer_fwd[mb_idx * 2]; + int mv_ly = mv_buffer_fwd[mb_idx * 2 + 1]; + int mv_x = mv_lx >> chroma_shift_w; + int mv_y = mv_ly >> chroma_shift_h; + + int ws_x = mb_ox - half_mb + mv_x * alpha / ALPHA_MAX; + int ws_y = mb_oy - half_mb + mv_y * alpha / ALPHA_MAX; + int rx = luma_x - ws_x; + int ry = luma_y - ws_y; + + if (rx >= 0 && rx < 2 * mb_size && ry >= 0 && ry < 2 * mb_size) { + uchar w = obmc_weight_lut(rx, ry, log2_mb_size); + if (w > 0) { + int sx = max(0, min(width - 1, x - mv_x)); + int sy = max(0, min(height - 1, y - mv_y)); + uchar2 p = tex2D<uchar2>(tex_prev, sx, sy); + val_u += w * p.x; val_v += w * p.y; weight_sum += w; + } + } + } + } + } + if (weight_sum > 0) { + val_u = (val_u + weight_sum / 2) / weight_sum; + val_v = (val_v + weight_sum / 2) / weight_sum; + } + dst[y * dst_pitch + x] = make_uchar2((uchar)val_u, (uchar)val_v); +} + +__global__ void minterpolate_obmc_uchar2_aobmc( + cudaTextureObject_t tex_prev, + cudaTextureObject_t tex_next, + const int16_t *__restrict__ mv_buffer, + const int16_t *__restrict__ sub_mv_buffer, + const int *__restrict__ sub_valid, + const uint64_t *__restrict__ sbad_buffer, + uchar2 *__restrict__ dst, + int width, int height, int dst_pitch, + int b_width, int b_height, + int mb_size, int log2_mb_size, + int alpha, + int chroma_shift_w, int chroma_shift_h) +{ + int x = blockIdx.x * blockDim.x + threadIdx.x; + int y = blockIdx.y * blockDim.y + threadIdx.y; + if (x >= width || y >= height) return; + + int weight_sum = 0, val_u = 0, val_v = 0; + int half_mb = mb_size / 2; + int luma_x = x << chroma_shift_w; + int luma_y = y << chroma_shift_h; + int mb_x0 = luma_x / mb_size; + int mb_y0 = luma_y / mb_size; + + /* Map to neighbor offsets for AOBMC */ + int base_start_lx = (mb_x0 << log2_mb_size) - half_mb; + int base_start_ly = (mb_y0 << log2_mb_size) - half_mb; + int base_rx = luma_x - base_start_lx; + int base_ry = luma_y - base_start_ly; + int nb_x = (((base_rx >> (log2_mb_size - 1)) * 2 - 3) / 2); + int nb_y = (((base_ry >> (log2_mb_size - 1)) * 2 - 3) / 2); + if (mb_x0 + nb_x < 0 || mb_x0 + nb_x >= b_width) nb_x = 0; + if (mb_y0 + nb_y < 0 || mb_y0 + nb_y >= b_height) nb_y = 0; + int nb_mb_idx = (mb_x0 + nb_x) + (mb_y0 + nb_y) * b_width; + uint64_t nb_sbad = sbad_buffer[nb_mb_idx]; + uint64_t cur_sbad = sbad_buffer[mb_x0 + mb_y0 * b_width]; + + for (int dy = -1; dy <= 0; dy++) { + int mb_y = mb_y0 + dy; + if (mb_y < 0 || mb_y >= b_height) continue; + for (int dx = -1; dx <= 0; dx++) { + int mb_x = mb_x0 + dx; + if (mb_x < 0 || mb_x >= b_width) continue; + int start_lx = (mb_x << log2_mb_size) - half_mb; + int start_ly = (mb_y << log2_mb_size) - half_mb; + int rx = luma_x - start_lx, ry = luma_y - start_ly; + if (rx < 0 || rx >= 2 * mb_size || ry < 0 || ry >= 2 * mb_size) continue; + uchar obmc_weight = obmc_weight_lut(rx, ry, log2_mb_size); + if (obmc_weight == 0) continue; + + if ((dx != 0 || dy != 0) && (nb_x != 0 || nb_y != 0) && + nb_sbad > 0 && cur_sbad > 0) { + int phi = min(ALPHA_MAX, (int)(ALPHA_MAX * nb_sbad / cur_sbad)); + obmc_weight = (uchar)((int)obmc_weight * phi / ALPHA_MAX); + } + + int mb_idx = mb_x + mb_y * b_width; + int mv_lx, mv_ly; + if (sub_valid && sub_valid[mb_idx]) { + int sub_x = (luma_x - (mb_x << log2_mb_size)) >> (log2_mb_size - 1); + int sub_y = (luma_y - (mb_y << log2_mb_size)) >> (log2_mb_size - 1); + int si = sub_x + sub_y * 2; + mv_lx = sub_mv_buffer[(mb_idx * 4 + si) * 2]; + mv_ly = sub_mv_buffer[(mb_idx * 4 + si) * 2 + 1]; + } else { + mv_lx = mv_buffer[mb_idx * 2]; + mv_ly = mv_buffer[mb_idx * 2 + 1]; + } + int mv_x = mv_lx >> chroma_shift_w; + int mv_y = mv_ly >> chroma_shift_h; + + { + int sx = x + (mv_x * alpha) / ALPHA_MAX; + int sy = y + (mv_y * alpha) / ALPHA_MAX; + sx = max(0, min(width - 1, sx)); + sy = max(0, min(height - 1, sy)); + uchar2 p = tex2D<uchar2>(tex_prev, sx, sy); + int w = obmc_weight * (ALPHA_MAX - alpha); + val_u += w * p.x; val_v += w * p.y; weight_sum += w; + } + { + int sx = x - (mv_x * (ALPHA_MAX - alpha)) / ALPHA_MAX; + int sy = y - (mv_y * (ALPHA_MAX - alpha)) / ALPHA_MAX; + sx = max(0, min(width - 1, sx)); + sy = max(0, min(height - 1, sy)); + uchar2 p = tex2D<uchar2>(tex_next, sx, sy); + int w = obmc_weight * alpha; + val_u += w * p.x; val_v += w * p.y; weight_sum += w; + } + } + } + if (weight_sum > 0) { + val_u = (val_u + weight_sum / 2) / weight_sum; + val_v = (val_v + weight_sum / 2) / weight_sum; + } + dst[y * dst_pitch + x] = make_uchar2((uchar)val_u, (uchar)val_v); +} + +/* =================================================================== + * VSBM — Variable-Size Block ME: sub-block diamond search at boundaries + * Grid: (b_count, 1, 1) Block: (1, 1, 1) + * + * For blocks with neighbors in different MV clusters, split into 4 + * sub-blocks. Each sub-block runs a diamond search starting from the + * parent MV. If all 4 sub-blocks have lower per-pixel cost than the + * parent, the split is accepted. + * =================================================================== */ +__global__ void minterpolate_vsbm_sub_me( + cudaTextureObject_t tex_cur, /* frame[1] */ + cudaTextureObject_t tex_ref, /* frame[2] */ + const int16_t *__restrict__ mv_buffer, /* parent MVs */ + const int *__restrict__ cid_buffer, /* cluster IDs */ + int16_t *__restrict__ sub_mv_buffer, /* [b_count][4][2] */ + int *__restrict__ sub_valid, /* [b_count] 1=split */ + int b_width, int b_height, + int mb_size, int log2_mb_size, + int width, int height) +{ + int mb_idx = blockIdx.x; + if (mb_idx >= b_width * b_height) return; + + int mb_x = mb_idx % b_width; + int mb_y = mb_idx / b_width; + int x_mb = mb_x * mb_size; + int y_mb = mb_y * mb_size; + + sub_valid[mb_idx] = 0; + + /* Skip frame edges (like CPU: !mb_x || !mb_y || edge check) */ + if (mb_x == 0 || mb_y == 0 || + mb_x == b_width - 1 || mb_y == b_height - 1) + return; + + int my_cid = cid_buffer[mb_idx]; + int parent_mv_x = mv_buffer[mb_idx * 2]; + int parent_mv_y = mv_buffer[mb_idx * 2 + 1]; + + /* Check if any 4-connected neighbor belongs to a different cluster */ + int is_boundary = 0; + /* up */ + if (cid_buffer[mb_idx - b_width] != my_cid) is_boundary = 1; + /* down */ + if (cid_buffer[mb_idx + b_width] != my_cid) is_boundary = 1; + /* left */ + if (cid_buffer[mb_idx - 1] != my_cid) is_boundary = 1; + /* right */ + if (cid_buffer[mb_idx + 1] != my_cid) is_boundary = 1; + + if (!is_boundary) + return; + + /* Compute parent cost using bilateral SAD (2xMB window) */ + uint64_t cost_old = 0; + int half = mb_size / 2; + for (int j = -half; j < mb_size * 3 / 2; j++) { + int yc = y_mb + parent_mv_y + j; + int yr = y_mb - parent_mv_y + j; + yc = max(0, min(height - 1, yc)); + yr = max(0, min(height - 1, yr)); + for (int i = -half; i < mb_size * 3 / 2; i++) { + int xc = x_mb + parent_mv_x + i; + int xr = x_mb - parent_mv_x + i; + xc = max(0, min(width - 1, xc)); + xr = max(0, min(width - 1, xr)); + cost_old += abs((int)tex2D<uchar>(tex_cur, xc, yc) - + (int)tex2D<uchar>(tex_ref, xr, yr)); + } + } + if (cost_old == 0) return; + + int sub_mb = mb_size / 2; + int16_t sub_mvs[4][2]; + int all_pass = 1; + + for (int sy = 0; sy < 2 && all_pass; sy++) { + for (int sx = 0; sx < 2 && all_pass; sx++) { + int sx_mb = x_mb + sx * sub_mb; + int sy_mb = y_mb + sy * sub_mb; + int sub_idx = sx + sy * 2; + + /* Start from parent MV target */ + int mv_x = parent_mv_x, mv_y = parent_mv_y; + /* Diamond search with search_param=2, sub-block size */ + int best_dx = 0, best_dy = 0; + uint64_t best_sad = 0; + /* Compute initial cost at parent MV position for this sub-block */ + for (int j = 0; j < sub_mb; j++) { + for (int i = 0; i < sub_mb; i++) { + int cur_val = tex2D<uchar>(tex_cur, sx_mb + mv_x + i, sy_mb + mv_y + j); + int ref_val = tex2D<uchar>(tex_ref, sx_mb - mv_x + i, sy_mb - mv_y + j); + best_sad += abs(cur_val - ref_val); + } + } + + /* Small diamond search (search_param=2) */ + int improved; + do { + improved = 0; + for (int k = 0; k < 4; k++) { + int cx = best_dx + d_dia1[k][0]; + int cy = best_dy + d_dia1[k][1]; + int x_mv = sx_mb + mv_x + cx; + int y_mv = sy_mb + mv_y + cy; + int x_ref = sx_mb - mv_x - cx; + int y_ref = sy_mb - mv_y - cy; + if (x_mv < 0 || x_mv + sub_mb > width || + y_mv < 0 || y_mv + sub_mb > height || + x_ref < 0 || x_ref + sub_mb > width || + y_ref < 0 || y_ref + sub_mb > height) + continue; + uint64_t sad = 0; + for (int j = 0; j < sub_mb; j++) { + for (int i = 0; i < sub_mb; i++) { + sad += abs((int)tex2D<uchar>(tex_cur, x_mv + i, y_mv + j) - + (int)tex2D<uchar>(tex_ref, x_ref + i, y_ref + j)); + } + } + if (sad < best_sad) { + best_sad = sad; + best_dx = cx; best_dy = cy; + improved = 1; + } + } + } while (improved); + + sub_mvs[sub_idx][0] = (int16_t)(mv_x + best_dx); + sub_mvs[sub_idx][1] = (int16_t)(mv_y + best_dy); + + /* Check if sub-block cost < parent cost / 4 */ + if (best_sad >= cost_old / 4) { + all_pass = 0; + } + } + } + + if (all_pass) { + sub_valid[mb_idx] = 1; + for (int si = 0; si < 4; si++) { + sub_mv_buffer[(mb_idx * 4 + si) * 2] = sub_mvs[si][0]; + sub_mv_buffer[(mb_idx * 4 + si) * 2 + 1] = sub_mvs[si][1]; + } + } +} + +} /* extern "C" */ \ No newline at end of file -- 2.52.0

1 0

[PR] fix/utils-review (PR #23614)
by Jun Zhao 27 Jun '26

27 Jun '26

PR #23614 opened by Jun Zhao (mypopydev) URL: https://code.ffmpeg.org/FFmpeg/FFmpeg/pulls/23614 Patch URL: https://code.ffmpeg.org/FFmpeg/FFmpeg/pulls/23614.patch Four fixes in libavformat: a typo in internal.h, a swallowed error in ff_mkdir_p(), orphan OpenSSL declarations in tls.h, and a broken error check in avformat_network_init(). From 9039f07614eb1396ad5d37afdfaabfabfea353f7 Mon Sep 17 00:00:00 2001 From: Jun Zhao <barryjzhao(a)tencent.com> Date: Sat, 27 Jun 2026 14:06:24 +0800 Subject: [PATCH 1/4] avformat/internal: fix typo allow_unkown -> allow_unknown The declaration of ff_parse_opts_from_query_string() misspelled "unknown" as "unkown". The implementation in utils.c already uses the correct spelling. Signed-off-by: Jun Zhao <barryjzhao(a)tencent.com> --- libavformat/internal.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/libavformat/internal.h b/libavformat/internal.h index a04343989b..cf42389062 100644 --- a/libavformat/internal.h +++ b/libavformat/internal.h @@ -680,7 +680,7 @@ int ff_dict_set_timestamp(AVDictionary **dict, const char *key, int64_t timestam * nested protocols are used. * @return <0 on error */ -int ff_parse_opts_from_query_string(void *obj, const char *str, int allow_unkown); +int ff_parse_opts_from_query_string(void *obj, const char *str, int allow_unknown); /** * Make a RFC 4281/6381 like string describing a codec. -- 2.52.0 From 386845a681f245c5ecef9e361d2150588d50f714 Mon Sep 17 00:00:00 2001 From: Jun Zhao <barryjzhao(a)tencent.com> Date: Sat, 27 Jun 2026 14:06:48 +0800 Subject: [PATCH 2/4] avformat/utils: fix ff_mkdir_p() swallowing intermediate mkdir errors When creating nested directories (e.g. /a/b/c), if mkdir() fails for an intermediate component, the function continued creating deeper directories instead of stopping. This caused the original error code to be overwritten by subsequent failures, making debugging harder. Return immediately when an intermediate mkdir() fails, preserving the errno for the caller. Signed-off-by: Jun Zhao <barryjzhao(a)tencent.com> --- libavformat/utils.c | 4 ++++ 1 file changed, 4 insertions(+) diff --git a/libavformat/utils.c b/libavformat/utils.c index c64ff27418..5ee41bfdbe 100644 --- a/libavformat/utils.c +++ b/libavformat/utils.c @@ -443,6 +443,10 @@ int ff_mkdir_p(const char *path) tmp_ch = *pos; *pos = '\0'; ret = mkdir(temp, 0755); + if (ret < 0) { + av_free(temp); + return ret; + } *pos = tmp_ch; } } -- 2.52.0 From ed451a86004df272153d8c9d6f15b56e02911b29 Mon Sep 17 00:00:00 2001 From: Jun Zhao <barryjzhao(a)tencent.com> Date: Sat, 27 Jun 2026 15:10:37 +0800 Subject: [PATCH 3/4] avformat/tls: remove orphan ff_openssl_init/ff_openssl_deinit declarations These declarations were left behind by f0913b3f16 which removed support for OpenSSL < 1.1.0. Since OpenSSL >= 1.1.0 auto-initializes, neither function has an implementation or caller. Signed-off-by: Jun Zhao <barryjzhao(a)tencent.com> --- libavformat/tls.h | 3 --- 1 file changed, 3 deletions(-) diff --git a/libavformat/tls.h b/libavformat/tls.h index f2f4f8991f..570c245ff4 100644 --- a/libavformat/tls.h +++ b/libavformat/tls.h @@ -125,9 +125,6 @@ int ff_ssl_gen_key_cert(char *key_buf, size_t key_sz, char *cert_buf, size_t cer void ff_gnutls_init(void); void ff_gnutls_deinit(void); -int ff_openssl_init(void); -void ff_openssl_deinit(void); - /** * Whether the packet is a DTLS packet, as defined by RFC 5764 Section 5.1.2. */ -- 2.52.0 From e84b414e9b589e878adbc069a5c92178d6cd1f9f Mon Sep 17 00:00:00 2001 From: Jun Zhao <barryjzhao(a)tencent.com> Date: Sat, 27 Jun 2026 15:22:10 +0800 Subject: [PATCH 4/4] avformat/utils: fix avformat_network_init() error check for ff_network_init() ff_network_init() returns 0 on failure and 1 on success, never a negative AVERROR code. All other callers (sapdec.c, avio.c, rtspdec.c, rtsp.c, sapenc.c) use !ff_network_init() to check for failure, but avformat_network_init() used a < 0 comparison which never matched. Change the check to match the other callers and return AVERROR(EIO) on failure. Signed-off-by: Jun Zhao <barryjzhao(a)tencent.com> --- libavformat/network.c | 7 +++++++ libavformat/utils.c | 4 ++-- 2 files changed, 9 insertions(+), 2 deletions(-) diff --git a/libavformat/network.c b/libavformat/network.c index 2b7e665bc4..764b62a35b 100644 --- a/libavformat/network.c +++ b/libavformat/network.c @@ -55,6 +55,13 @@ void ff_tls_deinit(void) #endif } +/** + * Initialize the network subsystem. On Windows, this calls WSAStartup(). + * + * @return 1 on success, 0 on failure. Note: this does NOT follow the + * AVERROR convention (negative on error); callers should use + * !ff_network_init() to check for failure. + */ int ff_network_init(void) { #if HAVE_WINSOCK2_H diff --git a/libavformat/utils.c b/libavformat/utils.c index 5ee41bfdbe..64dab53197 100644 --- a/libavformat/utils.c +++ b/libavformat/utils.c @@ -566,8 +566,8 @@ int avformat_network_init(void) { #if CONFIG_NETWORK int ret; - if ((ret = ff_network_init()) < 0) - return ret; + if (!ff_network_init()) + return AVERROR(EIO); if ((ret = ff_tls_init()) < 0) return ret; #endif -- 2.52.0

1 0

[PR] avdevice/avfoundation: add device selection by USB serial and unique ID (PR #23613)
by Gabriel Balaich 27 Jun '26

27 Jun '26

PR #23613 opened by Gabriel Balaich (ninbura) URL: https://code.ffmpeg.org/FFmpeg/FFmpeg/pulls/23613 Patch URL: https://code.ffmpeg.org/FFmpeg/FFmpeg/pulls/23613.patch # Summary of changes This PR adds the ability for macOS users to list & select avfoundation video/audio inputs via unique ID & serial, where-as previously only name & index were the available options. # Why selecting video/audio inputs via name/index is unreliable. 1. When multiple video/audio inputs of the same model are attached to the same system, you cannot uniquely select said inputs via name. 2. Video input indexes can change after hot plugs and system restarts. 3. Audio input indexes can change when connecting/disconnecting audio inputs. 4. When multiple video inputs of the same model are attached to the same system, the source selected via index can randomly change from run-to-run (see example below). See attached `index-frame-test.zsh` & `capture.log` for full command & output. ```Zsh === frame 1 @ 2026-06-26_20-05-42.877 (index 0, 3840x2160@60) === Input #0, avfoundation, from '': Duration: N/A, start: 4412.272300, bitrate: N/A Stream #0:0: Video: rawvideo (NV12 / 0x3231564E), nv12, 3840x2160, 60 fps, 60 tbr, 1000k tbn, start 4412.272300 Stream mapping: Stream #0:0 -> #0:0 (rawvideo (native) -> mjpeg (native)) Output #0, image2, to 'index-frame-test/run_2026-06-26_20-05-42_idx0/frame-01_2026-06-26_20-05-42.877.jpg': Metadata: encoder : Lavf62.19.101 Stream #0:0: Video: mjpeg, yuv420p(pc, progressive), 3840x2160, q=2-31, 200 kb/s, 60 fps, 60 tbn Metadata: encoder : Lavc62.36.101 mjpeg Side data: CPB properties: bitrate max/min/avg: 0/0/200000 buffer size: 0 vbv_delay: N/A [image2 @ 0xbe1018280] The specified filename 'index-frame-test/run_2026-06-26_20-05-42_idx0/frame-01_2026-06-26_20-05-42.877.jpg' does not contain an image sequence pattern or a pattern is invalid. [image2 @ 0xbe1018280] Use a pattern such as %03d for an image sequence or use the -update option (with -frames:v 1 if needed) to write a single image. [out#0/image2 @ 0xbe12900c0] video:378KiB audio:0KiB subtitle:0KiB other streams:0KiB global headers:0KiB muxing overhead: unknown frame= 1 fps=0.0 q=12.7 Lsize=N/A time=00:00:00.01 bitrate=N/A speed=0.657x elapsed=0:00:00.02 # Repeated 10 times... ``` Output changes despite looping the same FFmpeg command 10x times against the same input. As far as I can tell this is due to 1 or many inputs sharing the same name (same model). <img src="/attachments/ed927d62-8258-4797-82eb-49ee014de73c" alt="Screenshot 2026-06-26 at 23.06.49" width="1200"> # Why unique ID is desirable, and when it's undesirable. ### Desirable - Unique ID allows you to select an audio device with an identifier that *never changes*. - Persistent across restarts, hot plugs, and when increasing or decreasing the amount of audio I/O attached to you system (index can change in these scenarios). - Allows for selecting 1 of many audio devices with the same name. - ```Zsh [AVFoundation indev @ 0xb0cc18180] [4] same-name-test uid=com.rogueamoeba.Loopback::F43848CE-D7CE-4C9C-B5F6-04F55190C157 serial=N/A [AVFoundation indev @ 0xb0cc18180] [6] same-name-test uid=com.rogueamoeba.Loopback::18AAEF28-5258-490C-8DC8-85FB09624E83 serial=N/A ``` ### Undesirable - A video source's unique ID *can change* after hotplugs & system restarts. - ```Zsh # Baseline before hot plug [1] Elgato 4K X uid=0x232000000fd9009b serial=A7SNB504219J0R [3] Elgato 4K X uid=0x222000000fd9009b serial=A7SNB50423R73R [6] Elgato 4K X uid=0x252000000fd9009c serial=A7SNB50424UBQI # After hot plug (index & uid can change) [1] Elgato 4K X uid=0x232000000fd9009b serial=A7SNB504219J0R [3] Elgato 4K X uid=0x222000000fd9009b serial=A7SNB50423R73R [4] Elgato 4K X uid=0x252000000fd9009b serial=A7SNB50424UBQI # After restart (index & uid can change) [0] Elgato 4K X uid=0x232000000fd9009b serial=A7SNB50424UBQI [2] Elgato 4K X uid=0x252000000fd9009b serial=A7SNB504219J0R [3] Elgato 4K X uid=0x262000000fd9009b serial=A7SNB50423R73R ``` # Why serial is desirable, and its limitations. ### Desirable - Serial allows you to select video/audio inputs with an identifier that *never changes* (tested with hot plugs & restarts). - Does not suffer from the same behavior as index where feed can randomly change run-to-run (tested with loops). - Unique across multiple units of the same device/model & name. - Provides a common identifier, like name, for selecting/correlating a pair of video/audio inputs from the same device. ### Limitations - Serial is only available for USB devices, but nearly all consumer-grade capture cards for macOS are connected via USB. From e4cd3786b7b88e4f36f2278de8e30f94d1ffaa35 Mon Sep 17 00:00:00 2001 From: Gabriel Balaich <ffmpeg(a)ninbura.com> Date: Fri, 26 Jun 2026 21:08:55 -0600 Subject: [PATCH] avdevice/avfoundation: add device selection by USB serial and unique ID A capture device could previously be selected only by index or by name. Both are unreliable when multiple audio/video devices share a name or across reboots. AVFoundation reorders the device indices, and the USB video uniqueID embeds the macOS locationID, which is reassigned on reboot or replug and can then resolve to a different physical device. Add four options: -video_device_serial / -audio_device_serial -video_device_uid / -audio_device_uid The USB serial number is the only identifier that stays pegged to a given physical unit. For video it is resolved to the device's current locationID via IOKit. For audio it is matched against the uniqueID, which already embeds it. The unique ID covers devices that have no serial, such as virtual camera/audio devices. -list_devices additionally prints each device's uniqueID and USB serial so the values can be discovered. IOKit is detected in configure and used only when available. Signed-off-by: Gabriel Balaich <ffmpeg(a)ninbura.com> --- Changelog | 1 + configure | 4 +- doc/indevs.texi | 21 ++- libavdevice/avfoundation.m | 263 ++++++++++++++++++++++++++++++++++++- 4 files changed, 281 insertions(+), 8 deletions(-) diff --git a/Changelog b/Changelog index 2ad3ee255f..714a45322c 100644 --- a/Changelog +++ b/Changelog @@ -18,6 +18,7 @@ version <next>: - Remove ogg/celt parsing - Bitstream filter to split Dolby Vision multi-layer HEVC - Add AMF hardware memory mapping support. +- AVFoundation input device selection by unique ID and USB serial number version 8.1: diff --git a/configure b/configure index a6bbb86807..f594e7c881 100755 --- a/configure +++ b/configure @@ -2723,6 +2723,7 @@ HAVE_LIST=" $TYPES_LIST gzip ioctl_posix + iokit libdrm_getfb2 makeinfo_html opencl_d3d11 @@ -4039,7 +4040,7 @@ android_camera_indev_deps="android camera2ndk mediandk pthreads" alsa_indev_deps="alsa" alsa_outdev_deps="alsa" avfoundation_indev_deps="avfoundation corevideo coremedia pthreads AVCaptureSession" -avfoundation_indev_suggest="coregraphics applicationservices" +avfoundation_indev_suggest="coregraphics applicationservices iokit" avfoundation_indev_extralibs="-framework Foundation" audiotoolbox_outdev_deps="audiotoolbox pthreads AudioObjectPropertyAddress" audiotoolbox_outdev_extralibs="-framework AudioToolbox -framework CoreAudio" @@ -7145,6 +7146,7 @@ enabled avfoundation && { disable coregraphics applicationservices check_lib coregraphics CoreGraphics/CoreGraphics.h CGGetActiveDisplayList "-framework CoreGraphics" || check_lib applicationservices ApplicationServices/ApplicationServices.h CGGetActiveDisplayList "-framework ApplicationServices" + check_lib iokit IOKit/IOKitLib.h IOServiceGetMatchingServices "-framework IOKit" check_objc_class AVFoundation/AVFoundation.h AVCaptureSession } diff --git a/doc/indevs.texi b/doc/indevs.texi index 8822e070fe..6554203509 100644 --- a/doc/indevs.texi +++ b/doc/indevs.texi @@ -148,8 +148,8 @@ AVFoundation supports the following options: @table @option @item -list_devices <TRUE|FALSE> -If set to true, a list of all available input devices is given showing all -device names and indices. +If set to true, a list of all available input devices is given showing their +index, name, unique ID, and USB serial number (when available). @item -video_device_index <INDEX> Specify the video device by its index. Overrides anything given in the input filename. @@ -157,6 +157,23 @@ Specify the video device by its index. Overrides anything given in the input fil @item -audio_device_index <INDEX> Specify the audio device by its index. Overrides anything given in the input filename. +@item -video_device_serial <SERIAL> +Specify the video device by its USB serial number (USB devices only). Unlike the +index, this is stable across reboots and device reordering. Overrides anything +given in the input filename. + +@item -audio_device_serial <SERIAL> +Specify the audio device by its USB serial number, matched against the device's +unique ID on a best-effort basis. See @option{-video_device_serial}. + +@item -video_device_uid <UID> +Specify the video device by its AVFoundation unique ID (as shown by +@option{-list_devices}). Unique-ID stability is device-dependent and is not +reliable for USB capture devices; prefer @option{-video_device_serial} for those. + +@item -audio_device_uid <UID> +Specify the audio device by its AVFoundation unique ID. See @option{-video_device_uid}. + @item -pixel_format <FORMAT> Request the video device to use a specific pixel format. If the specified format is not supported, a list of available formats is given diff --git a/libavdevice/avfoundation.m b/libavdevice/avfoundation.m index ebec1ac4f2..2f5988b318 100644 --- a/libavdevice/avfoundation.m +++ b/libavdevice/avfoundation.m @@ -25,7 +25,19 @@ * @author Thilo Borgmann <thilo.borgmann(a)mail.de> */ +#include "config.h" + #import <AVFoundation/AVFoundation.h> +#if HAVE_IOKIT +# import <IOKit/IOKitLib.h> + /* kIOMainPortDefault replaced kIOMasterPortDefault in the macOS 12 SDK. */ +# if defined(__MAC_OS_X_VERSION_MAX_ALLOWED) && __MAC_OS_X_VERSION_MAX_ALLOWED >= 120000 +# define AVF_IO_MAIN_PORT_DEFAULT kIOMainPortDefault +# else +# define AVF_IO_MAIN_PORT_DEFAULT kIOMasterPortDefault +# endif +#endif + #include <pthread.h> #include "libavutil/channel_layout.h" @@ -111,6 +123,10 @@ typedef struct char *url; char *video_filename; char *audio_filename; + char *video_device_serial; + char *audio_device_serial; + char *video_device_uid; + char *audio_device_uid; int num_video_devices; @@ -821,6 +837,230 @@ static NSArray* getDevicesWithMediaType(AVMediaType mediaType) { #endif } +#if HAVE_IOKIT +/* Newer macOS exposes USB devices as IOUSBHostDevice; pre-10.11 systems and some + * devices still use the legacy IOUSBDevice class. Try the modern class first. */ +static const char *const avf_usb_match_classes[] = { "IOUSBHostDevice", "IOUSBDevice" }; + +static int avf_io_get_string(io_service_t service, CFStringRef key, char *buf, size_t size) +{ + CFTypeRef ref = IORegistryEntryCreateCFProperty(service, key, kCFAllocatorDefault, 0); + int ok = ref && CFGetTypeID(ref) == CFStringGetTypeID() && CFStringGetCString(ref, buf, size, kCFStringEncodingUTF8); + if (ref) + CFRelease(ref); + return ok; +} + +static int avf_io_get_uint32(io_service_t service, CFStringRef key, uint32_t *out) +{ + CFTypeRef ref = IORegistryEntryCreateCFProperty(service, key, kCFAllocatorDefault, 0); + int ok = ref && CFGetTypeID(ref) == CFNumberGetTypeID() && CFNumberGetValue(ref, kCFNumberSInt32Type, out); + if (ref) + CFRelease(ref); + return ok; +} +#endif + +static int64_t avf_usb_location_for_serial(const char *serial) +{ +#if HAVE_IOKIT + int64_t location = -1; + + for (size_t i = 0; i < FF_ARRAY_ELEMS(avf_usb_match_classes) && location < 0; i++) { + io_iterator_t iterator = 0; + io_service_t service; + + if (IOServiceGetMatchingServices(AVF_IO_MAIN_PORT_DEFAULT, + IOServiceMatching(avf_usb_match_classes[i]), &iterator) != KERN_SUCCESS) + continue; + + while (location < 0 && (service = IOIteratorNext(iterator))) { + char found[256]; + uint32_t loc; + if (avf_io_get_string(service, CFSTR("USB Serial Number"), found, sizeof(found)) && + !strcmp(found, serial) && + avf_io_get_uint32(service, CFSTR("locationID"), &loc)) + location = loc; + IOObjectRelease(service); + } + IOObjectRelease(iterator); + } + + return location; +#else + return -1; +#endif +} + +static NSString *avf_usb_serial_for_location(uint32_t location) +{ +#if HAVE_IOKIT + NSString *serial = nil; + + for (size_t i = 0; i < FF_ARRAY_ELEMS(avf_usb_match_classes) && !serial; i++) { + io_iterator_t iterator = 0; + io_service_t service; + + if (IOServiceGetMatchingServices(AVF_IO_MAIN_PORT_DEFAULT, + IOServiceMatching(avf_usb_match_classes[i]), &iterator) != KERN_SUCCESS) + continue; + + while (!serial && (service = IOIteratorNext(iterator))) { + char found[256]; + uint32_t loc; + if (avf_io_get_uint32(service, CFSTR("locationID"), &loc) && loc == location && + avf_io_get_string(service, CFSTR("USB Serial Number"), found, sizeof(found))) + serial = [NSString stringWithUTF8String:found]; + IOObjectRelease(service); + } + IOObjectRelease(iterator); + } + + return serial; +#else + return nil; +#endif +} + +// USB video uniqueID = locationID<<32 | VID<<16 | PID; match on the locationID. +static AVCaptureDevice *avf_video_device_with_serial(const char *serial, + NSArray *devices, NSArray *devices_muxed, int *is_muxed) +{ + int64_t location = avf_usb_location_for_serial(serial); + NSArray *lists[2] = { devices, devices_muxed }; + + if (location < 0) + return nil; + + for (int i = 0; i < 2; i++) { + for (AVCaptureDevice *device in lists[i]) { + NSString *uid = [device uniqueID]; + if ([uid hasPrefix:@"0x"] && + (uint32_t)(strtoull([uid UTF8String], NULL, 16) >> 32) == (uint32_t)location) { + *is_muxed = (i == 1); + return device; + } + } + } + + return nil; +} + +static AVCaptureDevice *avf_video_device_with_uid(const char *uid, + NSArray *devices, NSArray *devices_muxed, int *is_muxed) +{ + NSString *want = [NSString stringWithUTF8String:uid]; + NSArray *lists[2] = { devices, devices_muxed }; + + for (int i = 0; i < 2; i++) { + for (AVCaptureDevice *device in lists[i]) { + if ([[device uniqueID] isEqualToString:want]) { + *is_muxed = (i == 1); + return device; + } + } + } + + return nil; +} + +// USB audio uniqueID embeds the serial (...:SERIAL:N); match it directly. +static AVCaptureDevice *avf_audio_device_with_serial(const char *serial, NSArray *devices) +{ + NSString *needle = [NSString stringWithFormat:@":%s:", serial]; + + for (AVCaptureDevice *device in devices) { + if ([[device uniqueID] rangeOfString:needle].location != NSNotFound) + return device; + } + + return nil; +} + +static AVCaptureDevice *avf_audio_device_with_uid(const char *uid, NSArray *devices) +{ + NSString *want = [NSString stringWithUTF8String:uid]; + + for (AVCaptureDevice *device in devices) { + if ([[device uniqueID] isEqualToString:want]) + return device; + } + + return nil; +} + +// Best-effort serial string for -list_devices, or "N/A". +static const char *avf_device_listing_serial(AVCaptureDevice *device) +{ + NSString *uid = [device uniqueID]; + + if ([uid hasPrefix:@"0x"]) { + unsigned long long value = strtoull([uid UTF8String], NULL, 16); + NSString *serial = avf_usb_serial_for_location((uint32_t)(value >> 32)); + return serial.length ? [serial UTF8String] : "N/A"; + } else { + // Trust the serial field only when the trailing field is the numeric engine index. + NSArray<NSString *> *fields = [uid componentsSeparatedByString:@":"]; + if (fields.count >= 2) { + NSString *serial = fields[fields.count - 2]; + NSString *index = fields[fields.count - 1]; + if (serial.length && index.length && + [index rangeOfCharacterFromSet: + [[NSCharacterSet decimalDigitCharacterSet] invertedSet]].location == NSNotFound) + return [serial UTF8String]; + } + return "N/A"; + } +} + +// Returns 1 if a serial/uid selector was set (*device = match, or nil after logging on miss), else 0. +static int avf_video_device_from_string(AVFContext *ctx, NSArray *devices, + NSArray *devices_muxed, AVCaptureDevice **device) +{ + if (ctx->video_device_serial && *ctx->video_device_serial) { + *device = avf_video_device_with_serial(ctx->video_device_serial, + devices, devices_muxed, &ctx->video_is_muxed); + if (!*device) + av_log(ctx, AV_LOG_ERROR, + "Video capture device with serial number '%s' not found\n", + ctx->video_device_serial); + return 1; + } + if (ctx->video_device_uid && *ctx->video_device_uid) { + *device = avf_video_device_with_uid(ctx->video_device_uid, + devices, devices_muxed, &ctx->video_is_muxed); + if (!*device) + av_log(ctx, AV_LOG_ERROR, + "Video capture device with unique ID '%s' not found\n", + ctx->video_device_uid); + return 1; + } + return 0; +} + +static int avf_audio_device_from_string(AVFContext *ctx, AVCaptureDevice **device) +{ + if (ctx->audio_device_serial && *ctx->audio_device_serial) { + NSArray *devices = getDevicesWithMediaType(AVMediaTypeAudio); + *device = avf_audio_device_with_serial(ctx->audio_device_serial, devices); + if (!*device) + av_log(ctx, AV_LOG_ERROR, + "Audio capture device with serial number '%s' not found\n", + ctx->audio_device_serial); + return 1; + } + if (ctx->audio_device_uid && *ctx->audio_device_uid) { + NSArray *devices = getDevicesWithMediaType(AVMediaTypeAudio); + *device = avf_audio_device_with_uid(ctx->audio_device_uid, devices); + if (!*device) + av_log(ctx, AV_LOG_ERROR, + "Audio capture device with unique ID '%s' not found\n", + ctx->audio_device_uid); + return 1; + } + return 0; +} + static int avf_read_header(AVFormatContext *s) { int ret = 0; @@ -848,12 +1088,14 @@ static int avf_read_header(AVFormatContext *s) for (AVCaptureDevice *device in devices) { const char *name = [[device localizedName] UTF8String]; index = [devices indexOfObject:device]; - av_log(ctx, AV_LOG_INFO, "[%d] %s\n", index, name); + av_log(ctx, AV_LOG_INFO, "[%d] %s uid=%s serial=%s\n", index, name, + [[device uniqueID] UTF8String], avf_device_listing_serial(device)); } for (AVCaptureDevice *device in devices_muxed) { const char *name = [[device localizedName] UTF8String]; index = [devices count] + [devices_muxed indexOfObject:device]; - av_log(ctx, AV_LOG_INFO, "[%d] %s\n", index, name); + av_log(ctx, AV_LOG_INFO, "[%d] %s uid=%s serial=%s\n", index, name, + [[device uniqueID] UTF8String], avf_device_listing_serial(device)); } #if !TARGET_OS_IPHONE && __MAC_OS_X_VERSION_MIN_REQUIRED >= 1070 if (num_screens > 0) { @@ -870,7 +1112,8 @@ static int avf_read_header(AVFormatContext *s) for (AVCaptureDevice *device in devices) { const char *name = [[device localizedName] UTF8String]; int index = [devices indexOfObject:device]; - av_log(ctx, AV_LOG_INFO, "[%d] %s\n", index, name); + av_log(ctx, AV_LOG_INFO, "[%d] %s uid=%s serial=%s\n", index, name, + [[device uniqueID] UTF8String], avf_device_listing_serial(device)); } goto fail; } @@ -888,7 +1131,10 @@ static int avf_read_header(AVFormatContext *s) sscanf(ctx->audio_filename, "%d", &ctx->audio_device_index); } - if (ctx->video_device_index >= 0) { + if (avf_video_device_from_string(ctx, devices, devices_muxed, &video_device)) { + if (!video_device) + goto fail; + } else if (ctx->video_device_index >= 0) { if (ctx->video_device_index < ctx->num_video_devices) { if (ctx->video_device_index < [devices count]) { video_device = [devices objectAtIndex:ctx->video_device_index]; @@ -989,7 +1235,10 @@ static int avf_read_header(AVFormatContext *s) } // get audio device - if (ctx->audio_device_index >= 0) { + if (avf_audio_device_from_string(ctx, &audio_device)) { + if (!audio_device) + goto fail; + } else if (ctx->audio_device_index >= 0) { NSArray *devices = getDevicesWithMediaType(AVMediaTypeAudio); if (ctx->audio_device_index >= [devices count]) { @@ -1278,6 +1527,10 @@ static const AVOption options[] = { { "list_devices", "list available devices", offsetof(AVFContext, list_devices), AV_OPT_TYPE_BOOL, {.i64=0}, 0, 1, AV_OPT_FLAG_DECODING_PARAM }, { "video_device_index", "select video device by index for devices with same name (starts at 0)", offsetof(AVFContext, video_device_index), AV_OPT_TYPE_INT, {.i64 = -1}, -1, INT_MAX, AV_OPT_FLAG_DECODING_PARAM }, { "audio_device_index", "select audio device by index for devices with same name (starts at 0)", offsetof(AVFContext, audio_device_index), AV_OPT_TYPE_INT, {.i64 = -1}, -1, INT_MAX, AV_OPT_FLAG_DECODING_PARAM }, + { "video_device_serial", "select video device by USB serial number (stable across reboots and reordering)", offsetof(AVFContext, video_device_serial), AV_OPT_TYPE_STRING, {.str = NULL}, 0, 0, AV_OPT_FLAG_DECODING_PARAM }, + { "audio_device_serial", "select audio device by USB serial number (best-effort match within its unique ID)", offsetof(AVFContext, audio_device_serial), AV_OPT_TYPE_STRING, {.str = NULL}, 0, 0, AV_OPT_FLAG_DECODING_PARAM }, + { "video_device_uid", "select video device by its AVFoundation unique ID", offsetof(AVFContext, video_device_uid), AV_OPT_TYPE_STRING, {.str = NULL}, 0, 0, AV_OPT_FLAG_DECODING_PARAM }, + { "audio_device_uid", "select audio device by its AVFoundation unique ID", offsetof(AVFContext, audio_device_uid), AV_OPT_TYPE_STRING, {.str = NULL}, 0, 0, AV_OPT_FLAG_DECODING_PARAM }, { "pixel_format", "set pixel format", offsetof(AVFContext, pixel_format), AV_OPT_TYPE_PIXEL_FMT, {.i64 = AV_PIX_FMT_YUV420P}, 0, INT_MAX, AV_OPT_FLAG_DECODING_PARAM}, { "framerate", "set frame rate", offsetof(AVFContext, framerate), AV_OPT_TYPE_VIDEO_RATE, {.str = "ntsc"}, 0, INT_MAX, AV_OPT_FLAG_DECODING_PARAM }, { "video_size", "set video size", offsetof(AVFContext, width), AV_OPT_TYPE_IMAGE_SIZE, {.str = NULL}, 0, 0, AV_OPT_FLAG_DECODING_PARAM }, -- 2.52.0

1 0

[PR] hwcontext_vulkan: probe host image transfer support before enabling (PR #23612)
by Kacper Michajłow 27 Jun '26

27 Jun '26

PR #23612 opened by Kacper Michajłow (kasper93) URL: https://code.ffmpeg.org/FFmpeg/FFmpeg/pulls/23612 Patch URL: https://code.ffmpeg.org/FFmpeg/FFmpeg/pulls/23612.patch From 1b0ad6620972a60b0b30cdad1588a5f6cd1ad616 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Kacper=20Michaj=C5=82ow?= <kasper93(a)gmail.com> Date: Sat, 27 Jun 2026 07:20:46 +0200 Subject: [PATCH 1/2] hwcontext_vulkan: probe host image transfer support before enabling MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Notably video decode/encode images are often not possible to allocate from host visible memory type. Fixes: #21704 Signed-off-by: Kacper Michajłow <kasper93(a)gmail.com> --- libavutil/hwcontext_vulkan.c | 70 ++++++++++++++++++++++++++++++++++++ libavutil/vulkan_functions.h | 1 + 2 files changed, 71 insertions(+) diff --git a/libavutil/hwcontext_vulkan.c b/libavutil/hwcontext_vulkan.c index 20b6ed46f8..40a943667e 100644 --- a/libavutil/hwcontext_vulkan.c +++ b/libavutil/hwcontext_vulkan.c @@ -2934,6 +2934,69 @@ static void vulkan_frames_uninit(AVHWFramesContext *hwfc) av_buffer_pool_uninit(&fp->tmp); } +/* Probes whether the host transfer usage bit is actually usable in combination + * with the rest of the usage flags. Notably video decode/encode images may not + * be possible to allocate from the host visible memory type. */ +static int vulkan_host_transfer_usable(AVHWFramesContext *hwfc) +{ + VulkanFramesPriv *fp = hwfc->hwctx; + AVVulkanFramesContext *hwctx = &fp->p; + VulkanDevicePriv *p = hwfc->device_ctx->hwctx; + AVVulkanDeviceContext *dev_hwctx = &p->p; + FFVulkanFunctions *vk = &p->vkctx.vkfn; + + /* DRM modifier tiling can't be probed without the actual modifier, which the + * driver only picks at creation time. */ + if (hwctx->tiling == VK_IMAGE_TILING_DRM_FORMAT_MODIFIER_EXT) + return 0; + + VkImageFormatProperties2 props = { + .sType = VK_STRUCTURE_TYPE_IMAGE_FORMAT_PROPERTIES_2, + }; + VkPhysicalDeviceImageFormatInfo2 pinfo = { + .sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_IMAGE_FORMAT_INFO_2, + .pNext = hwctx->create_pnext, + .format = hwctx->format[0], + .type = VK_IMAGE_TYPE_2D, + .tiling = hwctx->tiling, + .usage = hwctx->usage, + .flags = hwctx->img_flags, + }; + + if (vk->GetPhysicalDeviceImageFormatProperties2(dev_hwctx->phys_dev, + &pinfo, &props) != VK_SUCCESS) + return 0; + + if (!p->vkctx.host_image_props.identicalMemoryTypeRequirements) { + VkImageCreateInfo create_info = { + .sType = VK_STRUCTURE_TYPE_IMAGE_CREATE_INFO, + .pNext = hwctx->create_pnext, + .imageType = VK_IMAGE_TYPE_2D, + .format = hwctx->format[0], + .extent = { hwfc->width, hwfc->height, 1 }, + .mipLevels = 1, + .arrayLayers = hwctx->nb_layers, + .flags = hwctx->img_flags, + .tiling = hwctx->tiling, + .usage = hwctx->usage, + .samples = VK_SAMPLE_COUNT_1_BIT, + }; + VkDeviceImageMemoryRequirements req_info = { + .sType = VK_STRUCTURE_TYPE_DEVICE_IMAGE_MEMORY_REQUIREMENTS, + .pCreateInfo = &create_info, + }; + VkMemoryRequirements2 req = { + .sType = VK_STRUCTURE_TYPE_MEMORY_REQUIREMENTS_2, + }; + + vk->GetDeviceImageMemoryRequirements(dev_hwctx->act_dev, &req_info, &req); + if (!req.memoryRequirements.memoryTypeBits) + return 0; + } + + return 1; +} + static int vulkan_frames_init(AVHWFramesContext *hwfc) { int err; @@ -3066,6 +3129,13 @@ static int vulkan_frames_init(AVHWFramesContext *hwfc) } } + /* Drop the host transfer if it isn't usable for this image configuration. */ + if ((hwctx->usage & VK_IMAGE_USAGE_HOST_TRANSFER_BIT_EXT) && + !vulkan_host_transfer_usable(hwfc)) + { + hwctx->usage &= ~VK_IMAGE_USAGE_HOST_TRANSFER_BIT_EXT; + } + if (!hwctx->lock_frame) hwctx->lock_frame = lock_frame; diff --git a/libavutil/vulkan_functions.h b/libavutil/vulkan_functions.h index 86dff046b2..185a8e290d 100644 --- a/libavutil/vulkan_functions.h +++ b/libavutil/vulkan_functions.h @@ -172,6 +172,7 @@ typedef uint64_t FFVulkanExtensions; /* Image */ \ MACRO(1, 1, FF_VK_EXT_DRM_MODIFIER_FLAGS, GetImageDrmFormatModifierPropertiesEXT) \ MACRO(1, 1, FF_VK_EXT_NO_FLAG, GetImageMemoryRequirements2) \ + MACRO(1, 1, FF_VK_EXT_NO_FLAG, GetDeviceImageMemoryRequirements) \ MACRO(1, 1, FF_VK_EXT_NO_FLAG, CreateImage) \ MACRO(1, 1, FF_VK_EXT_NO_FLAG, BindImageMemory2) \ MACRO(1, 1, FF_VK_EXT_NO_FLAG, GetImageSubresourceLayout) \ -- 2.52.0 From bd9d9e4fc319644f01cbce537d268f03ff769f8d Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Kacper=20Michaj=C5=82ow?= <kasper93(a)gmail.com> Date: Sat, 27 Jun 2026 07:31:38 +0200 Subject: [PATCH 2/2] hwcontext_vulkan: drop NVIDIA/MoltenVK host transfer blocklist MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Now that we correctly test if images are valid for host transfer, we can drop the hardcoded blocklist. Signed-off-by: Kacper Michajłow <kasper93(a)gmail.com> --- libavutil/hwcontext_vulkan.c | 7 ++----- 1 file changed, 2 insertions(+), 5 deletions(-) diff --git a/libavutil/hwcontext_vulkan.c b/libavutil/hwcontext_vulkan.c index 40a943667e..dc7b21936b 100644 --- a/libavutil/hwcontext_vulkan.c +++ b/libavutil/hwcontext_vulkan.c @@ -3080,9 +3080,7 @@ static int vulkan_frames_init(AVHWFramesContext *hwfc) VK_IMAGE_USAGE_STORAGE_BIT | VK_IMAGE_USAGE_SAMPLED_BIT); - if (p->vkctx.extensions & FF_VK_EXT_HOST_IMAGE_COPY && - !(p->dprops.driverID == VK_DRIVER_ID_NVIDIA_PROPRIETARY) && - !(p->dprops.driverID == VK_DRIVER_ID_MOLTENVK)) + if (p->vkctx.extensions & FF_VK_EXT_HOST_IMAGE_COPY) hwctx->usage |= supported_usage & VK_IMAGE_USAGE_HOST_TRANSFER_BIT_EXT; /* Enables encoding of images, if supported by format and extensions */ @@ -4756,8 +4754,7 @@ static int vulkan_transfer_frame(AVHWFramesContext *hwfc, if (swf->width > hwfc->width || swf->height > hwfc->height) return AVERROR(EINVAL); - if (hwctx->usage & VK_IMAGE_USAGE_HOST_TRANSFER_BIT_EXT && - !(p->dprops.driverID == VK_DRIVER_ID_NVIDIA_PROPRIETARY)) + if (hwctx->usage & VK_IMAGE_USAGE_HOST_TRANSFER_BIT_EXT) return vulkan_transfer_host(hwfc, hwf, swf, upload); for (int i = 0; i < av_pix_fmt_count_planes(swf->format); i++) { -- 2.52.0

1 0

[PR] forgejo/workflows: change to targeting 9.0 release branch (PR #23611)
by Timo Rothenpieler 27 Jun '26

27 Jun '26

PR #23611 opened by Timo Rothenpieler (BtbN) URL: https://code.ffmpeg.org/FFmpeg/FFmpeg/pulls/23611 Patch URL: https://code.ffmpeg.org/FFmpeg/FFmpeg/pulls/23611.patch Just general housekeeping to update the 9.0 branches CI. From 16e59dfabf92a95a3cc99af38de54644fe995bb0 Mon Sep 17 00:00:00 2001 From: Timo Rothenpieler <timo(a)rothenpieler.org> Date: Sat, 27 Jun 2026 00:03:34 +0200 Subject: [PATCH] forgejo/workflows: change to targeting 9.0 release branch --- .forgejo/labeler/labeler.js | 74 -------------------------------- .forgejo/labeler/labeler.yml | 35 --------------- .forgejo/workflows/autolabel.yml | 32 -------------- .forgejo/workflows/lint.yml | 2 +- .forgejo/workflows/test.yml | 4 +- 5 files changed, 3 insertions(+), 144 deletions(-) delete mode 100644 .forgejo/labeler/labeler.js delete mode 100644 .forgejo/labeler/labeler.yml delete mode 100644 .forgejo/workflows/autolabel.yml diff --git a/.forgejo/labeler/labeler.js b/.forgejo/labeler/labeler.js deleted file mode 100644 index 173d98080c..0000000000 --- a/.forgejo/labeler/labeler.js +++ /dev/null @@ -1,74 +0,0 @@ -module.exports = async ({github, context}) => { - const title = (context.payload.pull_request?.title || context.payload.issue?.title || '').toLowerCase(); - const labels = []; - const issueNumber = context.payload.pull_request?.number || context.payload.issue?.number; - - const kwmap = { - 'avcodec': 'avcodec', - 'avdevice': 'avdevice', - 'avfilter': 'avfilter', - 'avformat': 'avformat', - 'avutil': 'avutil', - 'swresample': 'swresample', - 'swscale': 'swscale', - 'fftools': 'CLI', - 'vulkan': 'vulkan' - }; - - async function isOrgMember(username) { - try { - const response = await github.rest.orgs.checkMembershipForUser({ - org: context.repo.owner, - username: username - }); - return response.status === 204; - } catch (error) { - return false; - } - } - - if (context.payload.action === 'closed' || - (context.payload.action !== 'opened' && ( - context.payload.action === 'assigned' || - context.payload.action === 'label_updated' || - context.payload.action === 'labeled' || - context.payload.comment) && - await isOrgMember(context.payload.sender.login)) - ) { - try { - await github.rest.issues.removeLabel({ - owner: context.repo.owner, - repo: context.repo.repo, - issue_number: issueNumber, - // this should say 'new', but forgejo deviates from GitHub API here and expects the ID - name: '41' - }); - console.log('Removed "new" label'); - } catch (error) { - if (error.status !== 404 && error.status !== 410) { - console.log('Could not remove "new" label'); - } - } - } else if (context.payload.action === 'opened') { - labels.push('new'); - console.log('Detected label: new'); - } - - if ((context.payload.action === 'opened' || context.payload.action === 'edited') && context.eventName !== 'issue_comment') { - for (const [kw, label] of Object.entries(kwmap)) { - if (title.includes(kw)) { - labels.push(label); - console.log('Detected label: ' + label); - } - } - } - - if (labels.length > 0) { - await github.rest.issues.addLabels({ - owner: context.repo.owner, - repo: context.repo.repo, - issue_number: issueNumber, - labels: labels, - }); - } -} diff --git a/.forgejo/labeler/labeler.yml b/.forgejo/labeler/labeler.yml deleted file mode 100644 index 446a675316..0000000000 --- a/.forgejo/labeler/labeler.yml +++ /dev/null @@ -1,35 +0,0 @@ -avcodec: - - changed-files: - - any-glob-to-any-file: 'libavcodec/**' - -avdevice: - - changed-files: - - any-glob-to-any-file: 'libavdevice/**' - -avfilter: - - changed-files: - - any-glob-to-any-file: 'libavfilter/**' - -avformat: - - changed-files: - - any-glob-to-any-file: 'libavformat/**' - -avutil: - - changed-files: - - any-glob-to-any-file: 'libavutil/**' - -swresample: - - changed-files: - - any-glob-to-any-file: 'libswresample/**' - -swscale: - - changed-files: - - any-glob-to-any-file: 'libswscale/**' - -CLI: - - changed-files: - - any-glob-to-any-file: 'fftools/**' - -vulkan: - - changed-files: - - any-glob-to-any-file: '**/*vulkan*' diff --git a/.forgejo/workflows/autolabel.yml b/.forgejo/workflows/autolabel.yml deleted file mode 100644 index 0235bdd7d4..0000000000 --- a/.forgejo/workflows/autolabel.yml +++ /dev/null @@ -1,32 +0,0 @@ -name: Autolabel - -on: - pull_request_target: - types: [opened, edited, synchronize, closed, assigned, labeled, unlabeled] - issues: - types: [opened, edited, closed, assigned, labeled, unlabeled] - issue_comment: - types: [created] - -jobs: - pr_labeler: - name: Labeler - runs-on: utilities - if: ${{ github.event.sender.login != 'ffmpeg-devel' }} - steps: - - name: Checkout - uses: actions/checkout@v6 - - name: Label by file-changes - uses: actions/labeler@v6 - if: ${{ forge.event_name == 'pull_request_target' }} - with: - configuration-path: .forgejo/labeler/labeler.yml - repo-token: ${{ secrets.AUTOLABELER_TOKEN }} - sync-labels: true - - name: Label by title-match - uses: actions/github-script@v8 - with: - script: | - const script = require('.forgejo/labeler/labeler.js') - await script({github, context}) - github-token: ${{ secrets.AUTOLABELER_TOKEN }} diff --git a/.forgejo/workflows/lint.yml b/.forgejo/workflows/lint.yml index 32237eecc7..d4c025104e 100644 --- a/.forgejo/workflows/lint.yml +++ b/.forgejo/workflows/lint.yml @@ -3,7 +3,7 @@ name: Lint on: push: branches: - - master + - release/9.0 pull_request: concurrency: diff --git a/.forgejo/workflows/test.yml b/.forgejo/workflows/test.yml index 3af1522b88..35e1064712 100644 --- a/.forgejo/workflows/test.yml +++ b/.forgejo/workflows/test.yml @@ -3,7 +3,7 @@ name: Test on: push: branches: - - master + - release/9.0 pull_request: concurrency: @@ -77,7 +77,7 @@ jobs: strategy: fail-fast: false matrix: - image: ['ghcr.io/btbn/ffmpeg-builds/win64-gpl:latest'] + image: ['ghcr.io/btbn/ffmpeg-builds/win64-gpl-9.0:latest'] target_exec: ['wine'] runs-on: linux-amd64 container: ${{ matrix.image }} -- 2.52.0

1 0

[PR] avcodec: odd dimensions bayer fixes (PR #23610)
by michaelni 26 Jun '26

26 Jun '26

PR #23610 opened by michaelni URL: https://code.ffmpeg.org/FFmpeg/FFmpeg/pulls/23610 Patch URL: https://code.ffmpeg.org/FFmpeg/FFmpeg/pulls/23610.patch From bbad4e19ff4508ce52f9f416f8355dde9f20a6b6 Mon Sep 17 00:00:00 2001 From: Michael Niedermayer <michael(a)niedermayer.cc> Date: Fri, 26 Jun 2026 20:58:49 +0200 Subject: [PATCH 1/3] avcodec: Align bayer pixel formats to 2x2 Fixes: out of array access with odd dimensioned bayer Fixes: fa6F4c0xA8el Signed-off-by: Michael Niedermayer <michael(a)niedermayer.cc> --- libavcodec/utils.c | 15 +++++++++++++++ 1 file changed, 15 insertions(+) diff --git a/libavcodec/utils.c b/libavcodec/utils.c index 550d818572..256fb4ed39 100644 --- a/libavcodec/utils.c +++ b/libavcodec/utils.c @@ -315,6 +315,21 @@ void avcodec_align_dimensions2(AVCodecContext *s, int *width, int *height, h_align = 8; } break; + case AV_PIX_FMT_BAYER_BGGR8: + case AV_PIX_FMT_BAYER_RGGB8: + case AV_PIX_FMT_BAYER_GBRG8: + case AV_PIX_FMT_BAYER_GRBG8: + case AV_PIX_FMT_BAYER_BGGR16LE: + case AV_PIX_FMT_BAYER_BGGR16BE: + case AV_PIX_FMT_BAYER_RGGB16LE: + case AV_PIX_FMT_BAYER_RGGB16BE: + case AV_PIX_FMT_BAYER_GBRG16LE: + case AV_PIX_FMT_BAYER_GBRG16BE: + case AV_PIX_FMT_BAYER_GRBG16LE: + case AV_PIX_FMT_BAYER_GRBG16BE: + w_align = FFMAX(w_align, 2); + h_align = FFMAX(h_align, 2); + break; default: break; } -- 2.52.0 From a23ad40bce779f2b3b799dc0868b9e3dc734a707 Mon Sep 17 00:00:00 2001 From: Michael Niedermayer <michael(a)niedermayer.cc> Date: Fri, 26 Jun 2026 20:58:49 +0200 Subject: [PATCH 2/3] avcodec/ffv1dec: Reject bayer with unaligned slice coordinates Bayer was introduced (4f509c9e43) with combined_version 0x4000A, so no real bayer stream uses the old unaligned coordinates; reject it. Odd picture dimensions remain valid and are handled by aligning the bayer allocation to 2x2. Fixes: out of array write Fixes: fa6F4c0xA8el Fixes: 4f509c9e43 (ffv1dec: implement Bayer pixel format encoding) Found-by: Anonymous Signed-off-by: Michael Niedermayer <michael(a)niedermayer.cc> --- libavcodec/ffv1dec.c | 6 ++++++ 1 file changed, 6 insertions(+) diff --git a/libavcodec/ffv1dec.c b/libavcodec/ffv1dec.c index 6677e5c8b4..0c0689c81f 100644 --- a/libavcodec/ffv1dec.c +++ b/libavcodec/ffv1dec.c @@ -583,6 +583,12 @@ static int read_header(FFV1Context *f, RangeCoder *c) if (ret < 0) return ret; + if (f->bayer && f->combined_version <= 0x40002) { + av_log(f->avctx, AV_LOG_ERROR, + "Bayer requires aligned slice coordinates (combined_version > 0x40002)\n"); + return AVERROR_INVALIDDATA; + } + if (f->configured_pix_fmt != f->pix_fmt || f->configured_width != f->width || f->configured_height != f->height || -- 2.52.0 From deafebff037d0edb559275796ac16134c1a5b4fe Mon Sep 17 00:00:00 2001 From: Michael Niedermayer <michael(a)niedermayer.cc> Date: Fri, 26 Jun 2026 20:58:49 +0200 Subject: [PATCH 3/3] avcodec/ffv1enc: Reject odd dimensions for bayer Signed-off-by: Michael Niedermayer <michael(a)niedermayer.cc> --- libavcodec/ffv1enc.c | 5 +++++ 1 file changed, 5 insertions(+) diff --git a/libavcodec/ffv1enc.c b/libavcodec/ffv1enc.c index eb55d3ba32..e6b7506557 100644 --- a/libavcodec/ffv1enc.c +++ b/libavcodec/ffv1enc.c @@ -1016,6 +1016,11 @@ static av_cold int encode_init_internal(AVCodecContext *avctx) if (ret < 0) return ret; + if (s->bayer && (avctx->width & 1 || avctx->height & 1)) { + av_log(avctx, AV_LOG_ERROR, "bayer requires even dimensions\n"); + return AVERROR(EINVAL); + } + if (s->bits_per_raw_sample > (s->version > 3 ? 16 : 8) && !s->remap_mode) { if (s->ac == AC_GOLOMB_RICE) { av_log(avctx, AV_LOG_INFO, -- 2.52.0

1 0