[PR] avfilter/vf_minterpolate_cuda: add CUDA-accelerated motion interpolation filter (PR #23615)
stevenliu via ffmpeg-devel <[email protected]>
| Newsgroups | gmane.comp.video.ffmpeg.devel |
|---|---|
| Message-ID | <178255315103.59.1810882831850808010@29965ddac10e> |
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 <[email protected]> # Summary of changes Briefly describe what this PR does and why. <!-- If this PR requires new FATE test samples, attach them to the PR and list their target paths below (relative to the fate-suite root). Attached filenames must match the sample's filename: ```fate-samples # e.g. vorbis/new-sample.ogg ``` --> From 6c619efadb15e86fee272b33064b6ab4b0269e66 Mon Sep 17 00:00:00 2001 From: Steven Liu <[email protected]> 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 <[email protected]> --- 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 <[email protected]> + * + * 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 _______________________________________________ ffmpeg-devel mailing list -- [email protected] To unsubscribe send an email to [email protected]