diff --git a/.dockerignore b/.dockerignore index 0b219701..a575e575 100644 --- a/.dockerignore +++ b/.dockerignore @@ -8,6 +8,8 @@ tmp/ .vscode/ .rsync-backup-*/ ffmpeg/ +!deps/ffmpeg/ +!deps/ffmpeg/** *.log *.csv *.ts diff --git a/avpmixer/janus.py b/avpmixer/janus.py index 4395c1f2..d755ba0d 100644 --- a/avpmixer/janus.py +++ b/avpmixer/janus.py @@ -5,6 +5,7 @@ from dataclasses import dataclass RTP_PACKET_SIZE = 1_200 +DEFAULT_KEYFRAME_MIN_INTERVAL_MS = 150 @dataclass(frozen=True) @@ -16,6 +17,7 @@ class JanusVideoConfig: bitrate_kbps: int = 4_500 rtcp_bind: str = "0.0.0.0" rtcp_port: int = 0 + keyframe_min_interval_ms: int = DEFAULT_KEYFRAME_MIN_INTERVAL_MS def __post_init__(self) -> None: if not self.host: @@ -30,6 +32,9 @@ def __post_init__(self) -> None: raise ValueError("Janus bitrate must be positive") if not 0 <= self.rtcp_port <= 65535: raise ValueError("RTCP port must be between 0 and 65535") + if (type(self.keyframe_min_interval_ms) is not int + or not 0 <= self.keyframe_min_interval_ms <= 2_147_483_647): + raise ValueError("keyframe_min_interval_ms must be a non-negative integer") @property def rtcp_port_remote(self) -> int: @@ -58,6 +63,7 @@ def build_janus_output(avp, api, src_edge: str, janus: JanusVideoConfig, *, fps: avp.addNode(api.ForceKeyFrame({ "name": JANUS_KEYFRAME_NODE, "src": "janus_fps", "dst": "janus_keyframed", "interval_sec": "1/1", "auto_restart": "panic", "group": group, + "min_interval_ms": janus.keyframe_min_interval_ms, })) avp.addNode(api.AssumeVideoFormat({ "name": "janus_format", "src": "janus_keyframed", "dst": "janus_video", diff --git a/demos/cuda-overlay/Dockerfile b/demos/cuda-overlay/Dockerfile index a9ca27c8..8c3aa0ea 100644 --- a/demos/cuda-overlay/Dockerfile +++ b/demos/cuda-overlay/Dockerfile @@ -2,7 +2,7 @@ FROM nvidia/cuda:11.7.1-devel-ubuntu22.04 ARG DEBIAN_FRONTEND=noninteractive ARG AVPLUMBER_REVISION=workspace -ARG FFMPEG_TAG=n7.1.5 +ARG FFMPEG_TAG=n8.1 ARG NV_CODEC_HEADERS_TAG=n12.1.14.0 RUN apt-get update \ @@ -35,13 +35,13 @@ RUN git clone --quiet --branch "${NV_CODEC_HEADERS_TAG}" --depth 1 \ && make -C /tmp/nv-codec-headers install PREFIX=/usr/local \ && rm -rf /tmp/nv-codec-headers -COPY deps/ffmpeg-patches /build/deps/ffmpeg-patches +COPY deps/ffmpeg /build/deps/ffmpeg RUN git clone --quiet --branch "${FFMPEG_TAG}" --depth 1 \ https://github.com/FFmpeg/FFmpeg.git /tmp/ffmpeg \ && git -C /tmp/ffmpeg config user.name "cuda-overlay-demo builder" \ && git -C /tmp/ffmpeg config user.email "cuda-overlay-demo@local" \ - && git -C /tmp/ffmpeg am /build/deps/ffmpeg-patches/*.patch + && bash /build/deps/ffmpeg/apply.sh /tmp/ffmpeg RUN cd /tmp/ffmpeg \ && ./configure \ @@ -117,6 +117,7 @@ COPY demos/cuda-overlay /build/demos/cuda-overlay ENV LD_LIBRARY_PATH=/usr/local/lib:/usr/local/cuda/lib64 ENV AVPLUMBER_REVISION=${AVPLUMBER_REVISION} +ENV FFMPEG_TAG=${FFMPEG_TAG} ENV NVIDIA_DRIVER_CAPABILITIES=compute,utility,video ENV NVIDIA_VISIBLE_DEVICES=all ENV PYPLUMBER_PATH=/build diff --git a/demos/cuda-overlay/compose.yaml b/demos/cuda-overlay/compose.yaml index 535a292b..518cb3dc 100644 --- a/demos/cuda-overlay/compose.yaml +++ b/demos/cuda-overlay/compose.yaml @@ -5,7 +5,7 @@ services: dockerfile: demos/cuda-overlay/Dockerfile args: AVPLUMBER_REVISION: ${AVPLUMBER_REVISION:-workspace} - FFMPEG_TAG: n7.1.5 + FFMPEG_TAG: ${FFMPEG_TAG:-n8.1} image: avplumber-cuda-overlay-demo:local gpus: all environment: diff --git a/demos/cuda-overlay/docs/guide.md b/demos/cuda-overlay/docs/guide.md index 8dde2203..e38c450d 100644 --- a/demos/cuda-overlay/docs/guide.md +++ b/demos/cuda-overlay/docs/guide.md @@ -5,7 +5,8 @@ One 1080p base and 15 transparent, labeled overlays composed on the GPU by `overlay_many_cuda`. -Build the repository's FFmpeg patch stack on public FFmpeg `n7.1.5`, run the +Build the repository's shared FFmpeg 8.x patch stack on public FFmpeg `n8.1` +(or select `n8.0` with `FFMPEG_TAG`), run the patched `overlay_many_cuda` through a purpose-built PyPlumber graph, and compare every output plane against an independent CPU reference. diff --git a/demos/cuda-overlay/tests/test_run_matrix.py b/demos/cuda-overlay/tests/test_run_matrix.py index 51097c4a..a9f50b64 100644 --- a/demos/cuda-overlay/tests/test_run_matrix.py +++ b/demos/cuda-overlay/tests/test_run_matrix.py @@ -17,13 +17,43 @@ class RunMatrixTest(unittest.TestCase): def test_patch_identity_reads_the_repository_patch_series(self) -> None: expected_paths = sorted( - (REPOSITORY_DIR / "deps" / "ffmpeg-patches").glob("*.patch") + (REPOSITORY_DIR / "deps" / "ffmpeg" / "8").glob("*.patch") ) self.assertEqual(REPO_DIR, REPOSITORY_DIR) - self.assertEqual(len(expected_paths), 7) + self.assertGreater(len(expected_paths), 0) self.assertEqual(list(_patch_identity()), [path.name for path in expected_paths]) + def test_patch_identity_shares_series_between_ffmpeg80_and_81(self) -> None: + self.assertEqual(_patch_identity(), _patch_identity("n8.1")) + self.assertEqual(_patch_identity("n8.0"), _patch_identity("n8.1")) + + def test_patch_identity_rejects_unknown_tags(self) -> None: + with self.assertRaises(ValueError): + _patch_identity("n7.1.5") + + def test_cuda_images_keep_the_default_and_select_matching_patches(self) -> None: + for relative_path in ( + "demos/cuda-overlay/Dockerfile", + "demos/mixer/Dockerfile", + "demos/dmabuf-browser/consumer/Dockerfile.cuda", + ): + with self.subTest(dockerfile=relative_path): + dockerfile = (REPOSITORY_DIR / relative_path).read_text() + self.assertIn("ARG FFMPEG_TAG=n8.1", dockerfile) + self.assertIn("COPY deps/ffmpeg /build/deps/ffmpeg", dockerfile) + self.assertIn("/build/deps/ffmpeg/apply.sh /tmp/ffmpeg", dockerfile) + + def test_overlay_image_reports_the_selected_version(self) -> None: + dockerfile = (REPOSITORY_DIR / "demos/cuda-overlay/Dockerfile").read_text() + self.assertIn("ENV FFMPEG_TAG=${FFMPEG_TAG}", dockerfile) + + def test_mixer_checks_runtime_options_without_removed_command_marker(self) -> None: + dockerfile = (REPOSITORY_DIR / "demos/mixer/Dockerfile").read_text() + self.assertIn("-h filter=transition_cuda", dockerfile) + self.assertIn('$1 == "mode" && $3 ~ /T/', dockerfile) + self.assertNotIn('$1 ~ /C/', dockerfile) + if __name__ == "__main__": unittest.main() diff --git a/demos/cuda-overlay/tools/run_matrix.py b/demos/cuda-overlay/tools/run_matrix.py index 30b50f54..a4118bf3 100755 --- a/demos/cuda-overlay/tools/run_matrix.py +++ b/demos/cuda-overlay/tools/run_matrix.py @@ -33,9 +33,11 @@ def _command_text(command: list[str]) -> str: return output if output else f"exit {result.returncode}" -def _patch_identity() -> dict[str, str]: +def _patch_identity(ffmpeg_tag: str = "n8.1") -> dict[str, str]: + if ffmpeg_tag not in ("n8.0", "n8.1"): + raise ValueError(f"unsupported FFmpeg tag: {ffmpeg_tag}") identities: dict[str, str] = {} - for path in sorted((REPO_DIR / "deps" / "ffmpeg-patches").glob("*.patch")): + for path in sorted((REPO_DIR / "deps" / "ffmpeg" / "8").glob("*.patch")): identities[path.name] = hashlib.sha256(path.read_bytes()).hexdigest() return identities @@ -78,6 +80,8 @@ def main() -> int: parser.add_argument("--height", type=_positive_dimension, default=HEIGHT) args = parser.parse_args() + ffmpeg_tag = os.environ.get("FFMPEG_TAG", "n8.1") + started_at = dt.datetime.now(dt.timezone.utc) run_id = started_at.strftime("run-%Y%m%dT%H%M%SZ") run_dir = args.artifacts / run_id @@ -93,7 +97,7 @@ def main() -> int: "run_id": run_id, "started_at": started_at.isoformat(), "status": "running", - "ffmpeg_tag": "n7.1.5", + "ffmpeg_tag": ffmpeg_tag, "cuda_toolkit_minimum": "11.7", "dimensions": {"width": args.width, "height": args.height}, "requested_overlay_counts": args.counts, @@ -110,7 +114,7 @@ def main() -> int: ] ), "repository_commit": os.environ.get("AVPLUMBER_REVISION", "workspace"), - "patches_sha256": _patch_identity(), + "patches_sha256": _patch_identity(ffmpeg_tag), }, "cases": [], } diff --git a/demos/dmabuf-browser/consumer/Dockerfile.cuda b/demos/dmabuf-browser/consumer/Dockerfile.cuda index 010a443b..0d0036ae 100644 --- a/demos/dmabuf-browser/consumer/Dockerfile.cuda +++ b/demos/dmabuf-browser/consumer/Dockerfile.cuda @@ -16,7 +16,7 @@ ARG UBUNTU_VERSION=22.04 FROM nvidia/cuda:${CUDA_IMAGE_VERSION}-devel-ubuntu${UBUNTU_VERSION} AS builder ARG DEBIAN_FRONTEND=noninteractive -ARG FFMPEG_TAG=n7.1.5 +ARG FFMPEG_TAG=n8.1 ARG NV_CODEC_HEADERS_TAG=n12.1.14.0 ARG CUDA_NVCC_FLAGS="-gencode arch=compute_70,code=compute_70 -O2" @@ -50,13 +50,13 @@ RUN git clone --quiet --branch "${NV_CODEC_HEADERS_TAG}" --depth 1 \ && make -C /tmp/nv-codec-headers install PREFIX=/usr/local \ && rm -rf /tmp/nv-codec-headers -COPY deps/ffmpeg-patches /build/deps/ffmpeg-patches +COPY deps/ffmpeg /build/deps/ffmpeg RUN git clone --quiet --branch "${FFMPEG_TAG}" --depth 1 \ https://github.com/FFmpeg/FFmpeg.git /tmp/ffmpeg \ && git -C /tmp/ffmpeg config user.name "dmabuf-browser-demo builder" \ && git -C /tmp/ffmpeg config user.email "dmabuf-browser-demo@local" \ - && git -C /tmp/ffmpeg am /build/deps/ffmpeg-patches/*.patch + && bash /build/deps/ffmpeg/apply.sh /tmp/ffmpeg RUN cd /tmp/ffmpeg \ && ./configure \ diff --git a/demos/mixer/Dockerfile b/demos/mixer/Dockerfile index 01697f3b..69bf8341 100644 --- a/demos/mixer/Dockerfile +++ b/demos/mixer/Dockerfile @@ -2,7 +2,7 @@ FROM nvidia/cuda:11.7.1-devel-ubuntu22.04 ARG DEBIAN_FRONTEND=noninteractive ARG AVPLUMBER_REVISION=workspace -ARG FFMPEG_TAG=n7.1.5 +ARG FFMPEG_TAG=n8.1 ARG NV_CODEC_HEADERS_TAG=n12.1.14.0 RUN apt-get update \ @@ -36,13 +36,13 @@ RUN git clone --quiet --branch "${NV_CODEC_HEADERS_TAG}" --depth 1 \ && make -C /tmp/nv-codec-headers install PREFIX=/usr/local \ && rm -rf /tmp/nv-codec-headers -COPY deps/ffmpeg-patches /build/deps/ffmpeg-patches +COPY deps/ffmpeg /build/deps/ffmpeg RUN git clone --quiet --branch "${FFMPEG_TAG}" --depth 1 \ https://github.com/FFmpeg/FFmpeg.git /tmp/ffmpeg \ && git -C /tmp/ffmpeg config user.name "mixer-demo builder" \ && git -C /tmp/ffmpeg config user.email "mixer-demo@local" \ - && git -C /tmp/ffmpeg am /build/deps/ffmpeg-patches/*.patch + && bash /build/deps/ffmpeg/apply.sh /tmp/ffmpeg RUN cd /tmp/ffmpeg \ && ./configure \ @@ -77,8 +77,8 @@ RUN cd /tmp/ffmpeg \ && rm -rf /tmp/ffmpeg RUN /usr/local/bin/ffmpeg -hide_banner -filters | grep -q ' overlay_many_cuda ' \ - && /usr/local/bin/ffmpeg -hide_banner -filters \ - | awk '$2 == "transition_cuda" && $1 ~ /C/ { found = 1 } END { exit !found }' + && /usr/local/bin/ffmpeg -hide_banner -h filter=transition_cuda \ + | awk '$1 == "mode" && $3 ~ /T/ { found = 1 } END { exit !found }' WORKDIR /build diff --git a/demos/mixer/README.md b/demos/mixer/README.md index 503c7ad9..21eba91b 100644 --- a/demos/mixer/README.md +++ b/demos/mixer/README.md @@ -149,6 +149,12 @@ python3 demos/mixer/tests/frame_codes.py media --sources 16 --width 1920 --heigh ## Under the hood +Janus output limits forced keyframes to one per 150 ms by default (9 frames at +60 fps). Override with `--keyframe-min-interval-ms 200`; `0` disables the limit. +The option also applies to Janus renditions loaded with `--config`. Cuts and +ordinary frames are not delayed: pending cut/RTCP requests coalesce until the +next eligible frame. Periodic keyframes share the same limit. + Grouped mixer graph: inputs, two compositor slots, transitions, media wipe and output. Click for the full ungrouped graph. Two compositor slots draw every scene; a transition filter blends them and the diff --git a/demos/mixer/docs/guide.md b/demos/mixer/docs/guide.md index c8bb8cdd..ad89fa60 100644 --- a/demos/mixer/docs/guide.md +++ b/demos/mixer/docs/guide.md @@ -195,7 +195,7 @@ groups without reordering sources. Follow the [shared NVIDIA setup guide](../../README.md) first. It also provides a local Janus preview if you want WebRTC output. -The demo image builds FFmpeg 7.1 with `deps/ffmpeg-patches`, verifies the +The demo image defaults to FFmpeg 8.1 with the shared `deps/ffmpeg/8` series, verifies the patched CUDA overlay and transition filters, and builds the CUDA-enabled AVPlumber Python module against that FFmpeg installation: diff --git a/demos/mixer/mixer.py b/demos/mixer/mixer.py index be059d75..b566a9e6 100644 --- a/demos/mixer/mixer.py +++ b/demos/mixer/mixer.py @@ -18,7 +18,8 @@ from avpmixer.dmabuf_inputs import (dmabuf_cuda_input_nodes, is_dmabuf_url, open_browser_windows, open_windows, refresh_windows, wait_for_sockets, window_id) from avpmixer.inputs import build_input -from avpmixer.janus import JANUS_KEYFRAME_NODE, JanusVideoConfig, build_janus_output +from avpmixer.janus import (DEFAULT_KEYFRAME_MIN_INTERVAL_MS, JANUS_KEYFRAME_NODE, + JanusVideoConfig, build_janus_output) try: from .layouts import ( @@ -69,6 +70,7 @@ class GraphOptions: janus_video_pt: int = JANUS_DEFAULT_VIDEO_PT janus_video_ssrc: int = JANUS_DEFAULT_VIDEO_SSRC janus_video_bitrate_kbps: int = JANUS_DEFAULT_VIDEO_BITRATE_KBPS + keyframe_min_interval_ms: int = DEFAULT_KEYFRAME_MIN_INTERVAL_MS janus_rtcp_bind: str = "0.0.0.0" janus_rtcp_port: int = 0 preheat_timeout_sec: float = 60.0 @@ -107,6 +109,9 @@ def validate(self) -> None: raise ValueError("janus_rtcp_port must be between 0 and 65535") if self.preheat_timeout_sec <= 0: raise ValueError("preheat_timeout_sec must be positive") + if (type(self.keyframe_min_interval_ms) is not int + or not 0 <= self.keyframe_min_interval_ms <= 2_147_483_647): + raise ValueError("keyframe_min_interval_ms must be a non-negative integer") if any(v <= 0 for v in self.dmabuf_size): raise ValueError("--dmabuf-size must be WxH with positive numbers") ids = self.dmabuf_inputs @@ -502,6 +507,7 @@ def _build_outputs(avp, api, options: GraphOptions, mixer_edge: str, *, host=options.janus_host, video_port=options.janus_video_port, payload_type=options.janus_video_pt, ssrc=options.janus_video_ssrc, bitrate_kbps=options.janus_video_bitrate_kbps, + keyframe_min_interval_ms=options.keyframe_min_interval_ms, rtcp_bind=options.janus_rtcp_bind, rtcp_port=options.janus_rtcp_port, ), fps=options.fps, fps_den=FPS_DEN, width=width, height=height, @@ -609,6 +615,7 @@ def _build_renditions(avp, api, options: GraphOptions, cfg, mixer_edge: str): video_port=rendition.port or options.janus_video_port, payload_type=options.janus_video_pt, ssrc=options.janus_video_ssrc, bitrate_kbps=rendition.bitrate_kbps, + keyframe_min_interval_ms=options.keyframe_min_interval_ms, rtcp_bind=options.janus_rtcp_bind, rtcp_port=options.janus_rtcp_port, ), fps=rendition.fps, fps_den=FPS_DEN, width=rendition.width, height=rendition.height, @@ -755,6 +762,10 @@ def parse_args(argv: list[str] | None = None) -> GraphOptions: default=JANUS_DEFAULT_VIDEO_BITRATE_KBPS, ) parser.add_argument("--janus-rtcp-bind", default="0.0.0.0") + parser.add_argument("--keyframe-min-interval-ms", type=int, + default=DEFAULT_KEYFRAME_MIN_INTERVAL_MS, + help="Minimum forced-keyframe spacing for Janus output in media time " + "(default: 150 ms; 0 disables rate limiting)") parser.add_argument("--janus-rtcp-port", type=int, default=0) parser.add_argument("--preheat-timeout", type=float, default=60.0) parser.add_argument("--wipe-file", help="Alpha wipe clip to warm the media-wipe chain up with at start " @@ -787,6 +798,7 @@ def parse_args(argv: list[str] | None = None) -> GraphOptions: janus_video_pt=args.janus_video_pt, janus_video_ssrc=args.janus_video_ssrc, janus_video_bitrate_kbps=args.janus_video_bitrate_kbps, + keyframe_min_interval_ms=args.keyframe_min_interval_ms, janus_rtcp_bind=args.janus_rtcp_bind, janus_rtcp_port=args.janus_rtcp_port, preheat_timeout_sec=args.preheat_timeout, diff --git a/demos/mixer/tests/test_graph.py b/demos/mixer/tests/test_graph.py index 9d38aff9..3c424788 100644 --- a/demos/mixer/tests/test_graph.py +++ b/demos/mixer/tests/test_graph.py @@ -568,14 +568,52 @@ def test_cli_requires_inputs_or_config(): def test_transitions_trigger_a_keyframe_only_when_streaming(): FakeMixer.instances.clear() - build_application(GraphOptions(inputs=("a.mp4",), janus_output=True), api=fake_api()) + application = build_application(GraphOptions(inputs=("a.mp4",), janus_output=True), api=fake_api()) assert FakeMixer.instances[-1].parameters["keyframe_node"] == "janus_force_keyframe" + node = next(n for n in application.avp.nodes if n.parameters.get("name") == "janus_force_keyframe") + assert node.parameters["min_interval_ms"] == 150 + assert node.parameters["interval_sec"] == "1/1" FakeMixer.instances.clear() build_application(GraphOptions(inputs=("a.mp4",), output="p.mp4"), api=fake_api()) assert FakeMixer.instances[-1].parameters["keyframe_node"] is None +@pytest.mark.parametrize("minimum", [0, 100, 150, 200, 500]) +def test_janus_keyframe_limit_is_configurable(minimum): + from avpmixer.janus import JanusVideoConfig, build_janus_output + avp = FakeAvp() + build_janus_output(avp, fake_api(), "program", JanusVideoConfig(keyframe_min_interval_ms=minimum), + fps=60, width=1920, height=1080) + node = next(n for n in avp.nodes if n.parameters.get("name") == "janus_force_keyframe") + assert node.parameters["min_interval_ms"] == minimum + + +@pytest.mark.parametrize("minimum", [-1, 0.2, True, "200", 2**31]) +def test_janus_rejects_invalid_keyframe_limit(minimum): + from avpmixer.janus import JanusVideoConfig + with pytest.raises(ValueError, match="keyframe_min_interval_ms"): + JanusVideoConfig(keyframe_min_interval_ms=minimum) + + +@pytest.mark.parametrize("configured", [False, True]) +def test_mixer_keyframe_option_reaches_each_output_path(tmp_path, configured): + if configured: + path = tmp_path / "mixer.json" + path.write_text(json.dumps({ + **CONFIG, "sources": CONFIG["sources"][:1], "scenes": CONFIG["scenes"][:1], + "initial_scene": "full", "wipes": [], + "renditions": [{"id": "program", "target": "janus"}], + })) + options = GraphOptions(config=str(path), janus_output=True, keyframe_min_interval_ms=200) + else: + options = parse_args(["--input", "a.mp4", "--janus-output", "--keyframe-min-interval-ms", "200"]) + application = build_application(options, api=fake_api()) + node = next(n for n in application.avp.nodes if n.parameters.get("name") == "janus_force_keyframe") + assert node.parameters["min_interval_ms"] == 200 + assert parse_args(["--input", "a.mp4", "--janus-output"]).keyframe_min_interval_ms == 150 + + def test_wipe_dir_scans_a_library_and_explicit_entries_win(tmp_path): from avpmixer import config as mc for name in ("b_swoosh.mov", "a_dip.webm", "notes.txt", "c_star.mp4"): diff --git a/deps/avcpp b/deps/avcpp index 31de3f4f..dac00a23 160000 --- a/deps/avcpp +++ b/deps/avcpp @@ -1 +1 @@ -Subproject commit 31de3f4f937ed3bb30d083275e5e76192dfc9cb3 +Subproject commit dac00a2390000d6b16628cf0d5d0a4a3b0b5bbbb diff --git a/deps/ffmpeg-patches/0004-avcodec-nvdec-intra.patch b/deps/ffmpeg-patches/0004-avcodec-nvdec-intra.patch deleted file mode 100644 index e70a242a..00000000 --- a/deps/ffmpeg-patches/0004-avcodec-nvdec-intra.patch +++ /dev/null @@ -1,30 +0,0 @@ -From 100f0aa5a67eb5500a1c38576e4003db59255f4b Mon Sep 17 00:00:00 2001 -From: Jan Pietek -Date: Wed, 12 Aug 2026 15:54:13 +0000 -Subject: [PATCH 4/7] avcodec/nvdec: handle intra-only streams - -Preserve the legacy CUVID handling required for intra-only decoded streams. - -Original-commit: 4bb0479af64eab3994fecbbda2911609db892f22 ---- - libavcodec/cuviddec.c | 7 +++++++ - 1 file changed, 7 insertions(+) - -diff --git a/libavcodec/cuviddec.c b/libavcodec/cuviddec.c -index 3fae9c1..7ad59f1 100644 ---- a/libavcodec/cuviddec.c -+++ b/libavcodec/cuviddec.c -@@ -342,6 +342,13 @@ static int CUDAAPI cuvid_handle_video_sequence(void *opaque, CUVIDEOFORMAT* form - cuinfo.bitDepthMinus8 = format->bit_depth_luma_minus8; - cuinfo.DeinterlaceMode = ctx->deint_mode_current; - -+ if(avctx->gop_size == 100000) { -+ av_log(avctx, AV_LOG_INFO, "init decoder intra only\n"); -+ cuinfo.ulIntraDecodeOnly = 1; -+ } -+ -+ av_log(avctx, AV_LOG_INFO, "init decoder %d %d\n", ctx->nb_surfaces, avctx->refs); -+ - if (ctx->deint_mode_current != cudaVideoDeinterlaceMode_Weave && !ctx->drop_second_field) - avctx->framerate = av_mul_q(avctx->framerate, (AVRational){2, 1}); - diff --git a/deps/ffmpeg-patches/README.md b/deps/ffmpeg-patches/README.md deleted file mode 100644 index b4633f24..00000000 --- a/deps/ffmpeg-patches/README.md +++ /dev/null @@ -1,56 +0,0 @@ -# FFmpeg Patch Stack - -This directory contains the public FFmpeg patch stack used by avplumber's CUDA -composition and media-input workflows. - -## Base - -- Upstream repository: `https://github.com/FFmpeg/FFmpeg` -- Upstream tag: `n7.1.5` -- Upstream commit: `3a0867c2bfda4a4d4309ca1a8cbdc6175e67f587` -- Expected patched tree: `52361f7251069ef74fbb41460e6e1b65d6f9947c` - -## Series - -The seven patches are ordered by filename and grouped by feature rather than by -the chronology of incomplete ports and follow-up fixes: - -1. `0001-swscale-aarch64-argb-yuva420p.patch` — AArch64 fast color conversion. -2. `0002-avfilter-cuda-composition-suite.patch` — CUDA pad, convert, crop, - overlay, overlay-many, scale edge handling, transitions, and procedural - wipes. -3. `0003-avfilter-npp-cuda13-compat.patch` — CUDA 13 NPP compatibility. -4. `0004-avcodec-nvdec-intra.patch` — NVDEC intra-only stream handling. -5. `0005-avformat-rtp-rfc4175.patch` — RFC 4175 4:2:0 and incomplete-frame - handling. -6. `0006-avdevice-v4l2-compat.patch` — V4L2 timestamp compatibility. -7. `0007-avdevice-ndi-v5.patch` — NDI v5 device registration and documentation. - -Each patch message lists the original exported FFmpeg commits it replaces. - -The old FFmpeg `af_whisper` port is intentionally absent. Speech-to-text belongs -in an AVPlumber node and is not part of this FFmpeg variant. - -## Apply - -```bash -git clone --branch n7.1.5 --depth 1 \ - https://github.com/FFmpeg/FFmpeg clean-ffmpeg -git -C clean-ffmpeg config user.name "patch application" -git -C clean-ffmpeg config user.email "patch-application@local" -git -C clean-ffmpeg am /path/to/avplumber/deps/ffmpeg-patches/*.patch -``` - -## Verify - -Run the verifier with any FFmpeg Git checkout that contains the documented base -commit. It creates and removes an isolated temporary worktree; it does not alter -the checkout's active branch: - -```bash -deps/ffmpeg-patches/verify.sh /path/to/FFmpeg -``` - -Verification succeeds only when all seven patches apply and produce the exact -expected Git tree. Runtime validation is provided by `demos/cuda-overlay` and -`demos/mixer`, whose Dockerfiles build this series against FFmpeg `n7.1.5`. diff --git a/deps/ffmpeg-patches/verify.sh b/deps/ffmpeg-patches/verify.sh deleted file mode 100755 index f0913bb3..00000000 --- a/deps/ffmpeg-patches/verify.sh +++ /dev/null @@ -1,52 +0,0 @@ -#!/usr/bin/env bash -set -euo pipefail - -base_commit=3a0867c2bfda4a4d4309ca1a8cbdc6175e67f587 -expected_tree=52361f7251069ef74fbb41460e6e1b65d6f9947c -expected_patch_count=7 - -if [[ $# -ne 1 ]]; then - echo "usage: $0 /path/to/FFmpeg" >&2 - exit 2 -fi - -source_repo=$(git -C "$1" rev-parse --show-toplevel) -if ! git -C "$source_repo" cat-file -e "${base_commit}^{commit}"; then - echo "FFmpeg checkout does not contain base commit ${base_commit}" >&2 - exit 2 -fi - -script_dir=$(CDPATH= cd -- "$(dirname -- "$0")" && pwd) -shopt -s nullglob -patches=("$script_dir"/*.patch) -if [[ ${#patches[@]} -ne $expected_patch_count ]]; then - echo "expected ${expected_patch_count} patches, found ${#patches[@]}" >&2 - exit 1 -fi - -audit_root=$(mktemp -d "${TMPDIR:-/tmp}/avplumber-ffmpeg-verify.XXXXXX") -audit_worktree="$audit_root/ffmpeg" - -cleanup() { - if [[ -d "$audit_worktree" ]]; then - git -C "$source_repo" worktree remove --force "$audit_worktree" \ - >/dev/null 2>&1 || true - fi - rmdir "$audit_root" >/dev/null 2>&1 || true -} -trap cleanup EXIT - -git -C "$source_repo" worktree add --detach "$audit_worktree" "$base_commit" \ - >/dev/null -git -C "$audit_worktree" config user.name "avplumber patch verifier" -git -C "$audit_worktree" config user.email "patch-verifier@local" -git -C "$audit_worktree" am --whitespace=nowarn "${patches[@]}" >/dev/null - -actual_tree=$(git -C "$audit_worktree" rev-parse 'HEAD^{tree}') -if [[ "$actual_tree" != "$expected_tree" ]]; then - echo "unexpected patched tree: ${actual_tree}" >&2 - echo "expected patched tree: ${expected_tree}" >&2 - exit 1 -fi - -echo "FFmpeg patch stack verified: ${actual_tree}" diff --git a/deps/ffmpeg-patches/0001-swscale-aarch64-argb-yuva420p.patch b/deps/ffmpeg/8/0001-swscale-aarch64-argb-yuva420p.patch similarity index 87% rename from deps/ffmpeg-patches/0001-swscale-aarch64-argb-yuva420p.patch rename to deps/ffmpeg/8/0001-swscale-aarch64-argb-yuva420p.patch index d5a9aa69..c6e06fb6 100644 --- a/deps/ffmpeg-patches/0001-swscale-aarch64-argb-yuva420p.patch +++ b/deps/ffmpeg/8/0001-swscale-aarch64-argb-yuva420p.patch @@ -1,35 +1,27 @@ -From 127f4c991739e9af8cdc5ac1315511c5e1bd85ef Mon Sep 17 00:00:00 2001 -From: Teodor Wozniak -Date: Wed, 12 Aug 2026 15:53:23 +0000 -Subject: [PATCH 1/7] swscale/aarch64: add fast ARGB to YUVA420P conversion +From e65731e98a015ad5b8888367a7929f91fbcba4c9 Mon Sep 17 00:00:00 2001 +From: Jan Pietek +Date: Tue, 15 Sep 2026 09:35:57 +0200 +Subject: [PATCH 1/7] swscale/aarch64: port ARGB to YUVA420P fast paths to 8.1 -Add an unscaled NEON conversion path for ARGB input to YUVA420P output. - -Original-commit: 3708bfab707a6eeb151829fd5c92807470c0f6d8 --- - libswscale/aarch64/input.S | 36 ++++ + libswscale/aarch64/input.S | 32 ++++ libswscale/aarch64/swscale.c | 6 + libswscale/aarch64/swscale_unscaled.c | 125 +++++++++++++ libswscale/aarch64/swscale_unscaled_neon.S | 200 +++++++++++++++++++++ - 4 files changed, 367 insertions(+) + 4 files changed, 363 insertions(+) diff --git a/libswscale/aarch64/input.S b/libswscale/aarch64/input.S -index 5cb1871..bab2700 100644 +index c1c0adffc8..a671a2ed60 100644 --- a/libswscale/aarch64/input.S +++ b/libswscale/aarch64/input.S -@@ -313,3 +313,39 @@ rgbToUV_neon bgr24, rgb24, element=3 - rgbToUV_neon bgra32, rgba32, element=4 +@@ -314,6 +314,38 @@ rgbToUV_neon bgra32, rgba32, element=4 rgbToUV_neon abgr32, argb32, element=4, alpha_first=1 -+ -+// void ff_argb32ToA_neon(uint8_t *dst, const uint8_t *src, const uint8_t *unused1, -+// const uint8_t *unused2, int width, uint32_t *unused, void *opq) -+// dst stores one int16_t per pixel with the same scaling as abgrToA_c: -+// dst = (A << 6) | (A >> 2) + ++// dst = (A << 6) | (A >> 2), as in abgrToA_c. +function ff_argb32ToA_neon, export=1 + cmp w4, #0 + b.le 3f -+ + cmp w4, #16 + b.lt 2f +1: @@ -57,20 +49,24 @@ index 5cb1871..bab2700 100644 +3: + ret +endfunc ++ + #if HAVE_DOTPROD + ENABLE_DOTPROD + diff --git a/libswscale/aarch64/swscale.c b/libswscale/aarch64/swscale.c -index eb90728..a2abfe5 100644 +index 4f86364f4b..3844db23bd 100644 --- a/libswscale/aarch64/swscale.c +++ b/libswscale/aarch64/swscale.c -@@ -217,6 +217,8 @@ NEON_INPUT(bgr24); +@@ -291,6 +291,8 @@ NEON_INPUT(bgr24); NEON_INPUT(bgra32); NEON_INPUT(rgb24); NEON_INPUT(rgba32); +void ff_argb32ToA_neon(uint8_t *dst, const uint8_t *src, const uint8_t *, -+ const uint8_t *, int w, uint32_t *coeffs, void *opq); ++ const uint8_t *, int w, uint32_t *coeffs, void *opq); + NEON_INPUT_DOTPROD(bgra32); + NEON_INPUT_DOTPROD(rgba32); - void ff_lumRangeFromJpeg_neon(int16_t *dst, int width); - void ff_chrRangeFromJpeg_neon(int16_t *dstU, int16_t *dstV, int width); -@@ -256,6 +258,8 @@ av_cold void ff_sws_init_swscale_aarch64(SwsContext *c) +@@ -382,6 +384,8 @@ av_cold void ff_sws_init_swscale_aarch64(SwsInternal *c) c->chrToYV12 = ff_abgr32ToUV_half_neon; else c->chrToYV12 = ff_abgr32ToUV_neon; @@ -79,7 +75,7 @@ index eb90728..a2abfe5 100644 break; case AV_PIX_FMT_ARGB: -@@ -264,6 +268,8 @@ av_cold void ff_sws_init_swscale_aarch64(SwsContext *c) +@@ -390,6 +394,8 @@ av_cold void ff_sws_init_swscale_aarch64(SwsInternal *c) c->chrToYV12 = ff_argb32ToUV_half_neon; else c->chrToYV12 = ff_argb32ToUV_neon; @@ -89,7 +85,7 @@ index eb90728..a2abfe5 100644 case AV_PIX_FMT_BGR24: c->lumToYV12 = ff_bgr24ToY_neon; diff --git a/libswscale/aarch64/swscale_unscaled.c b/libswscale/aarch64/swscale_unscaled.c -index 9dfccc0..9626dd3 100644 +index ba24775210..75c3991ec0 100644 --- a/libswscale/aarch64/swscale_unscaled.c +++ b/libswscale/aarch64/swscale_unscaled.c @@ -21,6 +21,13 @@ @@ -106,7 +102,7 @@ index 9dfccc0..9626dd3 100644 #define YUV_TO_RGB_TABLE \ c->yuv2rgb_v2r_coeff, \ c->yuv2rgb_u2g_coeff, \ -@@ -164,6 +171,115 @@ static int nv24_to_yuv420p_neon_wrapper(SwsContext *c, const uint8_t *src[], +@@ -205,6 +212,115 @@ static int nv24_to_yuv420p_neon_wrapper(SwsInternal *c, const uint8_t *const src return srcSliceH; } @@ -188,11 +184,11 @@ index 9dfccc0..9626dd3 100644 + } +} + -+static int argb_to_yuva420p_neon_wrapper(SwsContext *c, const uint8_t *src[], -+ int srcStride[], int srcSliceY, int srcSliceH, -+ uint8_t *dst[], int dstStride[]) ++static int argb_to_yuva420p_neon_wrapper(SwsInternal *c, const uint8_t *const src[], ++ const int srcStride[], int srcSliceY, int srcSliceH, ++ uint8_t *const dst[], const int dstStride[]) +{ -+ const int width_aligned = c->srcW & ~15; ++ const int width_aligned = c->opts.src_w & ~15; + const uint8_t *src_ptr = src[0]; + uint8_t *ydst = dst[0] + srcSliceY * dstStride[0]; + uint8_t *udst = dst[1] + (srcSliceY >> 1) * dstStride[1]; @@ -201,9 +197,9 @@ index 9dfccc0..9626dd3 100644 + + /* Keep strict gating in wrapper too, so odd slice boundaries or non-16 widths + * never hit the assembly fast path. */ -+ if ((srcSliceY & 1) || (srcSliceH & 1) || (c->srcW & 15)) { ++ if ((srcSliceY & 1) || (srcSliceH & 1) || (c->opts.src_w & 15)) { + argb_to_yuva420p_c(src_ptr, ydst, udst, vdst, adst, -+ c->srcW, srcSliceH, dstStride[0], dstStride[1], ++ c->opts.src_w, srcSliceH, dstStride[0], dstStride[1], + srcStride[0], dstStride[3], c->input_rgb2yuv_table); + return srcSliceH; + } @@ -222,24 +218,24 @@ index 9dfccc0..9626dd3 100644 #define DECLARE_FF_NVX_TO_ALL_RGBX_FUNCS(nvx) \ DECLARE_FF_NVX_TO_RGBX_FUNCS(nvx, argb) \ DECLARE_FF_NVX_TO_RGBX_FUNCS(nvx, rgba) \ -@@ -207,6 +323,15 @@ static void get_unscaled_swscale_neon(SwsContext *c) { - (c->srcFormat == AV_PIX_FMT_NV24 || c->srcFormat == AV_PIX_FMT_NV42) && - !(c->srcH & 1) && !(c->srcW & 15) && !accurate_rnd) +@@ -259,6 +375,15 @@ static void get_unscaled_swscale_neon(SwsInternal *c) { + (c->opts.src_format == AV_PIX_FMT_NV24 || c->opts.src_format == AV_PIX_FMT_NV42) && + !(c->opts.src_h & 1) && !(c->opts.src_w & 15) && !accurate_rnd) c->convert_unscaled = nv24_to_yuv420p_neon_wrapper; + -+ if (c->srcFormat == AV_PIX_FMT_ARGB && -+ c->dstFormat == AV_PIX_FMT_YUVA420P && -+ !(c->srcH & 1) && -+ !(c->srcW & 15) && -+ !(c->srcW & 1) && !accurate_rnd) { ++ if (c->opts.src_format == AV_PIX_FMT_ARGB && ++ c->opts.dst_format == AV_PIX_FMT_YUVA420P && ++ !(c->opts.src_h & 1) && ++ !(c->opts.src_w & 15) && ++ !(c->opts.src_w & 1) && !accurate_rnd) { + c->convert_unscaled = argb_to_yuva420p_neon_wrapper; + c->dst_slice_align = 2; + } } - void ff_get_unscaled_swscale_aarch64(SwsContext *c) + void ff_get_unscaled_swscale_aarch64(SwsInternal *c) diff --git a/libswscale/aarch64/swscale_unscaled_neon.S b/libswscale/aarch64/swscale_unscaled_neon.S -index a074b18..0c961ad 100644 +index a074b187e5..0c961ad4af 100644 --- a/libswscale/aarch64/swscale_unscaled_neon.S +++ b/libswscale/aarch64/swscale_unscaled_neon.S @@ -20,6 +20,206 @@ diff --git a/deps/ffmpeg-patches/0002-avfilter-cuda-composition-suite.patch b/deps/ffmpeg/8/0002-avfilter-cuda-composition-suite.patch similarity index 81% rename from deps/ffmpeg-patches/0002-avfilter-cuda-composition-suite.patch rename to deps/ffmpeg/8/0002-avfilter-cuda-composition-suite.patch index ef437c10..f1e51adf 100644 --- a/deps/ffmpeg-patches/0002-avfilter-cuda-composition-suite.patch +++ b/deps/ffmpeg/8/0002-avfilter-cuda-composition-suite.patch @@ -1,51 +1,45 @@ -From 7bc571c2fdfcc90943f7d4c98ed99822fe181c71 Mon Sep 17 00:00:00 2001 -From: Jan Pietek -Date: Wed, 12 Aug 2026 15:53:55 +0000 -Subject: [PATCH 2/7] avfilter: add CUDA composition filter suite +From d9643102a8b37addf0905f9d18b199d6fb2c0321 Mon Sep 17 00:00:00 2001 +From: Jan Pietek +Date: Tue, 15 Sep 2026 09:37:19 +0200 +Subject: [PATCH 2/7] avfilter: port CUDA composition suite to FFmpeg 8.1 -Add and register the CUDA pad, convert, crop, overlay-many, and transition filters used by avplumber composition graphs. Include runtime overlay descriptors, procedural transitions, scale edge fixes, overlay context propagation, and CUDA architecture compatibility. - -This squashes the incomplete registration, missing-file, compatibility, and follow-up patches into the complete feature they implement. - -Original-commits: 3f0a41055cfc542662ce7ad68b131dbda2ed0bc4 1a54abb0cad9fd5992166ae23649bda9e9b14753 aa3aff11d4a032bda6ba2dffd6418969e66df953 534042a4c5b297fe7baca2adf9c231acb171ee12 de361ec1e6d9f8169f4ed8ebf0e9cbc469c57027 8a109d1d41400391ea3c3ac9ee6faeb925c402be 58a00475094535b234aa2a5b360fd8460b72fb86 ae1c26b083fbeb885b9850e1c54ec97ee9873d2a 2105ce91bf8f27d423d2ed55f49be0b4cf3557b4 0a384ecc48ba46ff8e5f15d960dd9bb652ea7b52 c3474bdec62f807fd4283669d9d5405f46da1390 5d0ac856ccf9738f6e84853172ea2656082f176c 1bbd5e268f460893dd900c2d7ba541f44446a350 b3c7f6584796fdd2d0dc439ba8cf8dc4746fb27a +Retain the custom pad implementation and existing processing behavior. Adapt FFFilter registration, preserve upstream scale_cuda format support, and use the upstream CUDA architecture fallback. Co-authored-by: Teodor Wozniak --- - configure | 5 +- - doc/filters.texi | 47 +++ - libavfilter/Makefile | 9 + - libavfilter/allfilters.c | 5 + - libavfilter/vf_convert_cuda.c | 511 +++++++++++++++++++++++ - libavfilter/vf_convert_cuda.cu | 273 ++++++++++++ - libavfilter/vf_crop_cuda.c | 561 +++++++++++++++++++++++++ + configure | 7 + + doc/filters.texi | 47 ++ + libavfilter/Makefile | 8 + + libavfilter/allfilters.c | 4 + + libavfilter/vf_convert_cuda.c | 511 +++++++++++++++ + libavfilter/vf_convert_cuda.cu | 273 ++++++++ + libavfilter/vf_crop_cuda.c | 561 +++++++++++++++++ libavfilter/vf_overlay_cuda.c | 1 + - libavfilter/vf_overlay_many_cuda.c | 559 +++++++++++++++++++++++++ - libavfilter/vf_overlay_many_cuda.cu | 252 +++++++++++ - libavfilter/vf_overlay_many_cuda.h | 68 +++ - libavfilter/vf_pad_cuda.c | 538 ++++++++++++++++++++++++ - libavfilter/vf_pad_cuda.cu | 37 ++ + libavfilter/vf_overlay_many_cuda.c | 559 +++++++++++++++++ + libavfilter/vf_overlay_many_cuda.cu | 252 ++++++++ + libavfilter/vf_overlay_many_cuda.h | 68 ++ + libavfilter/vf_pad_cuda.c | 922 +++++++++++++--------------- + libavfilter/vf_pad_cuda.cu | 84 +-- libavfilter/vf_scale_cuda.c | 5 + - libavfilter/vf_scale_cuda.cu | 104 ++++- - libavfilter/vf_transition_cuda.c | 621 ++++++++++++++++++++++++++++ - libavfilter/vf_transition_cuda.cu | 56 +++ + libavfilter/vf_scale_cuda.cu | 104 +++- + libavfilter/vf_transition_cuda.c | 621 +++++++++++++++++++ + libavfilter/vf_transition_cuda.cu | 56 ++ libavutil/hwcontext_cuda.c | 1 + - 18 files changed, 3636 insertions(+), 17 deletions(-) + 18 files changed, 3504 insertions(+), 580 deletions(-) create mode 100644 libavfilter/vf_convert_cuda.c create mode 100644 libavfilter/vf_convert_cuda.cu create mode 100644 libavfilter/vf_crop_cuda.c create mode 100644 libavfilter/vf_overlay_many_cuda.c create mode 100644 libavfilter/vf_overlay_many_cuda.cu create mode 100644 libavfilter/vf_overlay_many_cuda.h - create mode 100644 libavfilter/vf_pad_cuda.c - create mode 100644 libavfilter/vf_pad_cuda.cu create mode 100644 libavfilter/vf_transition_cuda.c create mode 100644 libavfilter/vf_transition_cuda.cu diff --git a/configure b/configure -index 1652b0d..2cd7884 100755 +index 1759694274..584e1df313 100755 --- a/configure +++ b/configure -@@ -3309,6 +3309,8 @@ chromakey_cuda_filter_deps="ffnvcodec" +@@ -3515,6 +3515,8 @@ chromakey_cuda_filter_deps="ffnvcodec" chromakey_cuda_filter_deps_any="cuda_nvcc cuda_llvm" colorspace_cuda_filter_deps="ffnvcodec" colorspace_cuda_filter_deps_any="cuda_nvcc cuda_llvm" @@ -54,30 +48,25 @@ index 1652b0d..2cd7884 100755 hwupload_cuda_filter_deps="ffnvcodec" scale_npp_filter_deps="ffnvcodec libnpp" scale2ref_npp_filter_deps="ffnvcodec libnpp" -@@ -3319,6 +3321,7 @@ thumbnail_cuda_filter_deps_any="cuda_nvcc cuda_llvm" +@@ -3525,6 +3527,11 @@ thumbnail_cuda_filter_deps_any="cuda_nvcc cuda_llvm" transpose_npp_filter_deps="ffnvcodec libnpp" overlay_cuda_filter_deps="ffnvcodec" overlay_cuda_filter_deps_any="cuda_nvcc cuda_llvm" +overlay_many_cuda_filter_deps="ffnvcodec cuda_nvcc" ++convert_cuda_filter_deps="ffnvcodec" ++convert_cuda_filter_deps_any="cuda_nvcc cuda_llvm" ++transition_cuda_filter_deps="ffnvcodec" ++transition_cuda_filter_deps_any="cuda_nvcc cuda_llvm" + pad_cuda_filter_deps="ffnvcodec" + pad_cuda_filter_deps_any="cuda_nvcc cuda_llvm" sharpen_npp_filter_deps="ffnvcodec libnpp" - - ddagrab_filter_deps="d3d11va IDXGIOutput1 DXGI_OUTDUPL_FRAME_INFO" -@@ -4705,7 +4708,7 @@ set_default nvcc - - if enabled cuda_nvcc; then - if $nvcc $nvccflags_default 2>&1 | grep -qi unsupported; then -- nvccflags_default="-gencode arch=compute_60,code=sm_60 -O2" -+ nvccflags_default="-gencode arch=compute_75,code=sm_75 -O2" - fi - fi - diff --git a/doc/filters.texi b/doc/filters.texi -index 428986a..a4ae86c 100644 +index 5d222c6b96..5ae2fc5080 100644 --- a/doc/filters.texi +++ b/doc/filters.texi -@@ -19050,6 +19050,53 @@ See @ref{framesync}. +@@ -19412,6 +19412,53 @@ testsrc=s=100x100, split=4 [in0][in1][in2][in3]; - This filter also supports the @ref{framesync} options. + @end itemize +@anchor{overlay_many_cuda} +@section overlay_many_cuda @@ -130,10 +119,10 @@ index 428986a..a4ae86c 100644 Apply Overcomplete Wavelet denoiser. diff --git a/libavfilter/Makefile b/libavfilter/Makefile -index 91487af..c80bece 100644 +index a530cfae29..2fae9853a6 100644 --- a/libavfilter/Makefile +++ b/libavfilter/Makefile -@@ -247,6 +247,8 @@ OBJS-$(CONFIG_COLORSPACE_CUDA_FILTER) += vf_colorspace_cuda.o \ +@@ -253,6 +253,8 @@ OBJS-$(CONFIG_COLORSPACE_CUDA_FILTER) += vf_colorspace_cuda.o \ vf_colorspace_cuda.ptx.o \ cuda/load_helper.o OBJS-$(CONFIG_COLORTEMPERATURE_FILTER) += vf_colortemperature.o @@ -142,7 +131,7 @@ index 91487af..c80bece 100644 OBJS-$(CONFIG_CONVOLUTION_FILTER) += vf_convolution.o OBJS-$(CONFIG_CONVOLUTION_OPENCL_FILTER) += vf_convolution_opencl.o opencl.o \ opencl/convolution.o -@@ -256,6 +258,7 @@ OBJS-$(CONFIG_COREIMAGE_FILTER) += vf_coreimage.o +@@ -262,6 +264,7 @@ OBJS-$(CONFIG_COREIMAGE_FILTER) += vf_coreimage.o OBJS-$(CONFIG_CORR_FILTER) += vf_corr.o framesync.o OBJS-$(CONFIG_COVER_RECT_FILTER) += vf_cover_rect.o lavfutils.o OBJS-$(CONFIG_CROP_FILTER) += vf_crop.o @@ -150,7 +139,7 @@ index 91487af..c80bece 100644 OBJS-$(CONFIG_CROPDETECT_FILTER) += vf_cropdetect.o edge_common.o OBJS-$(CONFIG_CUE_FILTER) += f_cue.o OBJS-$(CONFIG_CURVES_FILTER) += vf_curves.o -@@ -410,6 +413,9 @@ OBJS-$(CONFIG_OSCILLOSCOPE_FILTER) += vf_datascope.o +@@ -422,6 +425,9 @@ OBJS-$(CONFIG_OSCILLOSCOPE_FILTER) += vf_datascope.o OBJS-$(CONFIG_OVERLAY_FILTER) += vf_overlay.o framesync.o OBJS-$(CONFIG_OVERLAY_CUDA_FILTER) += vf_overlay_cuda.o framesync.o vf_overlay_cuda.ptx.o \ cuda/load_helper.o @@ -160,15 +149,7 @@ index 91487af..c80bece 100644 OBJS-$(CONFIG_OVERLAY_OPENCL_FILTER) += vf_overlay_opencl.o opencl.o \ opencl/overlay.o framesync.o OBJS-$(CONFIG_OVERLAY_QSV_FILTER) += vf_overlay_qsv.o framesync.o -@@ -417,6 +423,7 @@ OBJS-$(CONFIG_OVERLAY_VAAPI_FILTER) += vf_overlay_vaapi.o framesync.o v - OBJS-$(CONFIG_OVERLAY_VULKAN_FILTER) += vf_overlay_vulkan.o vulkan.o vulkan_filter.o - OBJS-$(CONFIG_OWDENOISE_FILTER) += vf_owdenoise.o - OBJS-$(CONFIG_PAD_FILTER) += vf_pad.o -+OBJS-$(CONFIG_PAD_CUDA_FILTER) += vf_pad_cuda.o vf_pad_cuda.ptx.o cuda/load_helper.o - OBJS-$(CONFIG_PAD_OPENCL_FILTER) += vf_pad_opencl.o opencl.o opencl/pad.o - OBJS-$(CONFIG_PALETTEGEN_FILTER) += vf_palettegen.o palette.o - OBJS-$(CONFIG_PALETTEUSE_FILTER) += vf_paletteuse.o framesync.o palette.o -@@ -528,6 +535,8 @@ OBJS-$(CONFIG_TONEMAP_OPENCL_FILTER) += vf_tonemap_opencl.o opencl.o \ +@@ -545,6 +551,8 @@ OBJS-$(CONFIG_TONEMAP_OPENCL_FILTER) += vf_tonemap_opencl.o opencl.o \ opencl/tonemap.o opencl/colorspace_common.o OBJS-$(CONFIG_TONEMAP_VAAPI_FILTER) += vf_tonemap_vaapi.o vaapi_vpp.o OBJS-$(CONFIG_TPAD_FILTER) += vf_tpad.o @@ -178,47 +159,44 @@ index 91487af..c80bece 100644 OBJS-$(CONFIG_TRANSPOSE_NPP_FILTER) += vf_transpose_npp.o OBJS-$(CONFIG_TRANSPOSE_OPENCL_FILTER) += vf_transpose_opencl.o opencl.o opencl/transpose.o diff --git a/libavfilter/allfilters.c b/libavfilter/allfilters.c -index 9819f0f..640b3f8 100644 +index e26859e159..01c75b2971 100644 --- a/libavfilter/allfilters.c +++ b/libavfilter/allfilters.c -@@ -226,6 +226,7 @@ extern const AVFilter ff_vf_colormatrix; - extern const AVFilter ff_vf_colorspace; - extern const AVFilter ff_vf_colorspace_cuda; - extern const AVFilter ff_vf_colortemperature; -+extern const AVFilter ff_vf_convert_cuda; - extern const AVFilter ff_vf_convolution; - extern const AVFilter ff_vf_convolution_opencl; - extern const AVFilter ff_vf_convolve; -@@ -234,6 +235,7 @@ extern const AVFilter ff_vf_coreimage; - extern const AVFilter ff_vf_corr; - extern const AVFilter ff_vf_cover_rect; - extern const AVFilter ff_vf_crop; -+extern const AVFilter ff_vf_crop_cuda; - extern const AVFilter ff_vf_cropdetect; - extern const AVFilter ff_vf_cue; - extern const AVFilter ff_vf_curves; -@@ -390,8 +392,10 @@ extern const AVFilter ff_vf_overlay_qsv; - extern const AVFilter ff_vf_overlay_vaapi; - extern const AVFilter ff_vf_overlay_vulkan; - extern const AVFilter ff_vf_overlay_cuda; -+extern const AVFilter ff_vf_overlay_many_cuda; - extern const AVFilter ff_vf_owdenoise; - extern const AVFilter ff_vf_pad; -+extern const AVFilter ff_vf_pad_cuda; - extern const AVFilter ff_vf_pad_opencl; - extern const AVFilter ff_vf_palettegen; - extern const AVFilter ff_vf_paletteuse; -@@ -497,6 +501,7 @@ extern const AVFilter ff_vf_tonemap; - extern const AVFilter ff_vf_tonemap_opencl; - extern const AVFilter ff_vf_tonemap_vaapi; - extern const AVFilter ff_vf_tpad; -+extern const AVFilter ff_vf_transition_cuda; - extern const AVFilter ff_vf_transpose; - extern const AVFilter ff_vf_transpose_npp; - extern const AVFilter ff_vf_transpose_opencl; +@@ -230,6 +230,7 @@ extern const FFFilter ff_vf_colormatrix; + extern const FFFilter ff_vf_colorspace; + extern const FFFilter ff_vf_colorspace_cuda; + extern const FFFilter ff_vf_colortemperature; ++extern const FFFilter ff_vf_convert_cuda; + extern const FFFilter ff_vf_convolution; + extern const FFFilter ff_vf_convolution_opencl; + extern const FFFilter ff_vf_convolve; +@@ -238,6 +239,7 @@ extern const FFFilter ff_vf_coreimage; + extern const FFFilter ff_vf_corr; + extern const FFFilter ff_vf_cover_rect; + extern const FFFilter ff_vf_crop; ++extern const FFFilter ff_vf_crop_cuda; + extern const FFFilter ff_vf_cropdetect; + extern const FFFilter ff_vf_cue; + extern const FFFilter ff_vf_curves; +@@ -399,6 +401,7 @@ extern const FFFilter ff_vf_overlay_qsv; + extern const FFFilter ff_vf_overlay_vaapi; + extern const FFFilter ff_vf_overlay_vulkan; + extern const FFFilter ff_vf_overlay_cuda; ++extern const FFFilter ff_vf_overlay_many_cuda; + extern const FFFilter ff_vf_owdenoise; + extern const FFFilter ff_vf_pad; + extern const FFFilter ff_vf_pad_cuda; +@@ -512,6 +515,7 @@ extern const FFFilter ff_vf_tonemap; + extern const FFFilter ff_vf_tonemap_opencl; + extern const FFFilter ff_vf_tonemap_vaapi; + extern const FFFilter ff_vf_tpad; ++extern const FFFilter ff_vf_transition_cuda; + extern const FFFilter ff_vf_transpose; + extern const FFFilter ff_vf_transpose_npp; + extern const FFFilter ff_vf_transpose_opencl; diff --git a/libavfilter/vf_convert_cuda.c b/libavfilter/vf_convert_cuda.c new file mode 100644 -index 0000000..12f9299 +index 0000000000..5f49d2126e --- /dev/null +++ b/libavfilter/vf_convert_cuda.c @@ -0,0 +1,511 @@ @@ -721,11 +699,11 @@ index 0000000..12f9299 + }, +}; + -+const AVFilter ff_vf_convert_cuda = { -+ .name = "convert_cuda", -+ .description = NULL_IF_CONFIG_SMALL("CUDA accelerated pixel format conversion"), ++const FFFilter ff_vf_convert_cuda = { ++ .p.name = "convert_cuda", ++ .p.description = NULL_IF_CONFIG_SMALL("CUDA accelerated pixel format conversion"), + .priv_size = sizeof(ConvertCUDAContext), -+ .priv_class = &convert_cuda_class, ++ .p.priv_class = &convert_cuda_class, + .init = init, + .uninit = uninit, + FILTER_INPUTS(inputs), @@ -735,7 +713,7 @@ index 0000000..12f9299 +}; diff --git a/libavfilter/vf_convert_cuda.cu b/libavfilter/vf_convert_cuda.cu new file mode 100644 -index 0000000..beabc0c +index 0000000000..beabc0cd91 --- /dev/null +++ b/libavfilter/vf_convert_cuda.cu @@ -0,0 +1,273 @@ @@ -1014,7 +992,7 @@ index 0000000..beabc0c +} diff --git a/libavfilter/vf_crop_cuda.c b/libavfilter/vf_crop_cuda.c new file mode 100644 -index 0000000..d08dc38 +index 0000000000..a8586793d7 --- /dev/null +++ b/libavfilter/vf_crop_cuda.c @@ -0,0 +1,561 @@ @@ -1566,11 +1544,11 @@ index 0000000..d08dc38 + }, +}; + -+const AVFilter ff_vf_crop_cuda = { -+ .name = "crop_cuda", -+ .description = NULL_IF_CONFIG_SMALL("Crop the input CUDA video."), ++const FFFilter ff_vf_crop_cuda = { ++ .p.name = "crop_cuda", ++ .p.description = NULL_IF_CONFIG_SMALL("Crop the input CUDA video."), + .priv_size = sizeof(CropCUDAContext), -+ .priv_class = &crop_cuda_class, ++ .p.priv_class = &crop_cuda_class, + .init = init, + .uninit = uninit, + FILTER_INPUTS(crop_cuda_inputs), @@ -1580,10 +1558,10 @@ index 0000000..d08dc38 + .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE, +}; diff --git a/libavfilter/vf_overlay_cuda.c b/libavfilter/vf_overlay_cuda.c -index a35f6ed..ac5a789 100644 +index 3d39458f3f..dcf768cab5 100644 --- a/libavfilter/vf_overlay_cuda.c +++ b/libavfilter/vf_overlay_cuda.c -@@ -495,6 +495,7 @@ static int overlay_cuda_config_output(AVFilterLink *outlink) +@@ -477,6 +477,7 @@ static int overlay_cuda_config_output(AVFilterLink *outlink) ctx->hwctx = ((AVHWDeviceContext*)ctx->hw_device_ctx->data)->hwctx; cuda_ctx = ctx->hwctx->cuda_ctx; @@ -1593,7 +1571,7 @@ index a35f6ed..ac5a789 100644 ctx->cu_stream = ctx->hwctx->stream; diff --git a/libavfilter/vf_overlay_many_cuda.c b/libavfilter/vf_overlay_many_cuda.c new file mode 100644 -index 0000000..da5b915 +index 0000000000..20e5e9c9e6 --- /dev/null +++ b/libavfilter/vf_overlay_many_cuda.c @@ -0,0 +1,559 @@ @@ -2142,12 +2120,12 @@ index 0000000..da5b915 + }, +}; + -+const AVFilter ff_vf_overlay_many_cuda = { -+ .name = "overlay_many_cuda", -+ .description = NULL_IF_CONFIG_SMALL( ++const FFFilter ff_vf_overlay_many_cuda = { ++ .p.name = "overlay_many_cuda", ++ .p.description = NULL_IF_CONFIG_SMALL( + "Overlay multiple videos on top of each other using CUDA"), + .priv_size = sizeof(OverlayManyCUDAContext), -+ .priv_class = &overlay_many_cuda_class, ++ .p.priv_class = &overlay_many_cuda_class, + .init = overlay_many_cuda_init, + .uninit = overlay_many_cuda_uninit, + .activate = overlay_many_cuda_activate, @@ -2158,7 +2136,7 @@ index 0000000..da5b915 +}; diff --git a/libavfilter/vf_overlay_many_cuda.cu b/libavfilter/vf_overlay_many_cuda.cu new file mode 100644 -index 0000000..46d45df +index 0000000000..46d45df61c --- /dev/null +++ b/libavfilter/vf_overlay_many_cuda.cu @@ -0,0 +1,252 @@ @@ -2416,7 +2394,7 @@ index 0000000..46d45df +} /* extern "C" */ diff --git a/libavfilter/vf_overlay_many_cuda.h b/libavfilter/vf_overlay_many_cuda.h new file mode 100644 -index 0000000..4f42b88 +index 0000000000..4f42b883cd --- /dev/null +++ b/libavfilter/vf_overlay_many_cuda.h @@ -0,0 +1,68 @@ @@ -2489,29 +2467,68 @@ index 0000000..4f42b88 + +#endif /* AVFILTER_VF_OVERLAY_MANY_CUDA_H */ diff --git a/libavfilter/vf_pad_cuda.c b/libavfilter/vf_pad_cuda.c -new file mode 100644 -index 0000000..7daa227 ---- /dev/null +index 425d6d2a94..7003c1a8df 100644 +--- a/libavfilter/vf_pad_cuda.c +++ b/libavfilter/vf_pad_cuda.c -@@ -0,0 +1,538 @@ -+#include "libavutil/hwcontext.h" -+#include "libavutil/hwcontext_cuda_internal.h" +@@ -1,636 +1,538 @@ +-/* +- * 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 +- */ +- +-/** +- * @file +- * CUDA video padding filter +- */ +- +-#include +- +-#include "filters.h" +-#include "libavutil/avstring.h" +-#include "libavutil/common.h" +-#include "libavutil/cuda_check.h" +-#include "libavutil/eval.h" + #include "libavutil/hwcontext.h" + #include "libavutil/hwcontext_cuda_internal.h" +-#include "libavutil/imgutils.h" +-#include "libavutil/internal.h" +#include "libavutil/cuda_check.h" -+#include "libavutil/opt.h" -+#include "libavutil/pixdesc.h" + #include "libavutil/opt.h" + #include "libavutil/pixdesc.h" +#include "libavutil/eval.h" -+#include "libavutil/colorspace.h" -+ + #include "libavutil/colorspace.h" + +#include "avfilter.h" +#include "formats.h" +#include "avfilter_internal.h" +#include "video.h" + -+#include "cuda/load_helper.h" -+ -+#define BLOCK_X 32 -+#define BLOCK_Y 16 -+ + #include "cuda/load_helper.h" + +-#define CHECK_CU(x) FF_CUDA_CHECK_DL(ctx, device_hwctx->internal->cuda_dl, x) +-#define DIV_UP(a, b) ( ((a) + (b) - 1) / (b) ) + #define BLOCK_X 32 + #define BLOCK_Y 16 + +-static const enum AVPixelFormat supported_formats[] = { +- AV_PIX_FMT_YUV420P, +- AV_PIX_FMT_YUV444P, +- AV_PIX_FMT_YUVA420P, +- AV_PIX_FMT_YUVA444P, +- AV_PIX_FMT_NV12, +enum var_name { + VAR_IN_W, VAR_IW, + VAR_IN_H, VAR_IH, @@ -2523,33 +2540,174 @@ index 0000000..7daa227 + VAR_SAR, + VAR_DAR, + VARS_NB -+}; -+ -+static const char *const var_names[] = { + }; + +-typedef struct CUDAPadContext { +- const AVClass *class; +- +- AVBufferRef *frames_ctx; +- +- int w, h; ///< output dimensions, a value of 0 will result in the input size +- int x, y; ///< offsets of the input area with respect to the padded area +- int in_w, in_h; ///< width and height for the padded input video +- +- char *w_expr; ///< width expression +- char *h_expr; ///< height expression +- char *x_expr; ///< x offset expression +- char *y_expr; ///< y offset expression +- +- uint8_t rgba_color[4]; ///< color for the padding area +- uint8_t parsed_color[4]; +- AVRational aspect; +- +- int eval_mode; +- +- int last_out_w, last_out_h; ///< used to evaluate the prior output width and height with the incoming frame +- +- AVCUDADeviceContext *hwctx; +- CUmodule cu_module; +- CUfunction cu_func_uchar; +- CUfunction cu_func_uchar2; +-} CUDAPadContext; +- + static const char *const var_names[] = { +- "in_w", "iw", +- "in_h", "ih", +- "out_w", "ow", +- "out_h", "oh", + "in_w", "iw", + "in_h", "ih", + "out_w", "ow", + "out_h", "oh", -+ "x", -+ "y", -+ "a", -+ "sar", -+ "dar", -+ NULL -+}; -+ + "x", + "y", + "a", + "sar", + "dar", +- "hsub", +- "vsub", + NULL + }; + +-enum { +- VAR_IN_W, +- VAR_IW, +- VAR_IN_H, +- VAR_IH, +- VAR_OUT_W, +- VAR_OW, +- VAR_OUT_H, +- VAR_OH, +- VAR_X, +- VAR_Y, +- VAR_A, +- VAR_SAR, +- VAR_DAR, +- VAR_HSUB, +- VAR_VSUB, +- VARS_NB +-}; +- +-enum EvalMode { +- EVAL_MODE_INIT, +- EVAL_MODE_FRAME, +- EVAL_MODE_NB +static enum AVPixelFormat supported_formats[] = { + AV_PIX_FMT_NV12, + AV_PIX_FMT_YUV420P, AV_PIX_FMT_YUVA420P, + AV_PIX_FMT_YUV444P, AV_PIX_FMT_YUVA444P, -+}; -+ + }; + +-static int eval_expr(AVFilterContext *ctx) +-{ +- CUDAPadContext *s = ctx->priv; +- AVFilterLink *inlink = ctx->inputs[0]; +- const AVPixFmtDescriptor *desc = av_pix_fmt_desc_get(inlink->format); +- +- double var_values[VARS_NB], res; +- char *expr; +- int ret; +- +- var_values[VAR_IN_W] = var_values[VAR_IW] = s->in_w; +- var_values[VAR_IN_H] = var_values[VAR_IH] = s->in_h; +- var_values[VAR_OUT_W] = var_values[VAR_OW] = NAN; +- var_values[VAR_OUT_H] = var_values[VAR_OH] = NAN; +- var_values[VAR_A] = (double)s->in_w / s->in_h; +- var_values[VAR_SAR] = inlink->sample_aspect_ratio.num ? +- (double)inlink->sample_aspect_ratio.num / +- inlink->sample_aspect_ratio.den : 1; +- var_values[VAR_DAR] = var_values[VAR_A] * var_values[VAR_SAR]; +- var_values[VAR_HSUB] = 1 << desc->log2_chroma_w; +- var_values[VAR_VSUB] = 1 << desc->log2_chroma_h; +- +- expr = s->w_expr; +- ret = av_expr_parse_and_eval(&res, expr, var_names, var_values, NULL, NULL, NULL, NULL, NULL, 0, ctx); +- if (ret < 0) +- goto fail; +- +- s->w = res; +- if (s->w < 0) { +- av_log(ctx, AV_LOG_ERROR, "Width expression is negative.\n"); +- ret = AVERROR(EINVAL); +- goto fail; +- } +- +- var_values[VAR_OUT_W] = var_values[VAR_OW] = s->w; +- +- expr = s->h_expr; +- ret = av_expr_parse_and_eval(&res, expr, var_names, var_values, NULL, NULL, NULL, NULL, NULL, 0, ctx); +- if (ret < 0) +- goto fail; +- +- s->h = res; +- if (s->h < 0) { +- av_log(ctx, AV_LOG_ERROR, "Height expression is negative.\n"); +- ret = AVERROR(EINVAL); +- goto fail; +- } +- var_values[VAR_OUT_H] = var_values[VAR_OH] = s->h; +- +- if (!s->h) +- s->h = s->in_h; +- +- var_values[VAR_OUT_H] = var_values[VAR_OH] = s->h; +- +- +- expr = s->w_expr; +- ret = av_expr_parse_and_eval(&res, expr, var_names, var_values, NULL, NULL, NULL, NULL, NULL, 0, ctx); +- if (ret < 0) +- goto fail; +- +- s->w = res; +- if (s->w < 0) { +- av_log(ctx, AV_LOG_ERROR, "Width expression is negative.\n"); +- ret = AVERROR(EINVAL); +- goto fail; +- } +- if (!s->w) +- s->w = s->in_w; +- +- var_values[VAR_OUT_W] = var_values[VAR_OW] = s->w; +- +- +- expr = s->x_expr; +- ret = av_expr_parse_and_eval(&res, expr, var_names, var_values, NULL, NULL, NULL, NULL, NULL, 0, ctx); +- if (ret < 0) +- goto fail; +- +- s->x = res; +#define DIV_UP(a, b) ( ((a) + (b) - 1) / (b) ) +#define BLOCKX 32 +#define BLOCKY 16 -+ + +#define CHECK_CU(x) FF_CUDA_CHECK_DL(ctx, cu, x) -+ + +- expr = s->y_expr; +- ret = av_expr_parse_and_eval(&res, expr, var_names, var_values, NULL, NULL, NULL, NULL, NULL, 0, ctx); +- if (ret < 0) +- goto fail; +- +- s->y = res; +typedef struct PadCUDAContext { + const AVClass *class; + AVCUDADeviceContext *hwctx; @@ -2558,74 +2716,194 @@ index 0000000..7daa227 + AVBufferRef *frames_ctx; + AVFrame *frame; + AVFrame *tmp_frame; -+ + +- if (s->x < 0 || s->x + s->in_w > s->w) { +- s->x = (s->w - s->in_w) / 2; +- av_log(ctx, AV_LOG_VERBOSE, "centering X offset.\n"); +- } + CUmodule cu_module; + CUfunction cu_func; + CUstream cu_stream; -+ + +- if (s->y < 0 || s->y + s->in_h > s->h) { +- s->y = (s->h - s->in_h) / 2; +- av_log(ctx, AV_LOG_VERBOSE, "centering Y offset.\n"); +- } + enum AVPixelFormat sw_format; -+ + +- s->w = av_clip(s->w, 1, INT_MAX); +- s->h = av_clip(s->h, 1, INT_MAX); + char *w_expr, *h_expr, *x_expr, *y_expr; + AVRational aspect; + int w, h, x, y; + uint8_t pad_rgba[4]; + uint8_t pad_color[4]; +} PadCUDAContext; -+ -+ -+ + +- if (s->w < s->in_w || s->h < s->in_h) { +- av_log(ctx, AV_LOG_ERROR, "Padded size < input size.\n"); +- return AVERROR(EINVAL); +- } + +- av_log(ctx, AV_LOG_DEBUG, +- "w:%d h:%d -> w:%d h:%d x:%d y:%d color:0x%02X%02X%02X%02X\n", +- inlink->w, inlink->h, s->w, s->h, s->x, s->y, s->rgba_color[0], +- s->rgba_color[1], s->rgba_color[2], s->rgba_color[3]); + +static int format_is_supported(enum AVPixelFormat fmt) +{ + for (int i = 0; i < FF_ARRAY_ELEMS(supported_formats); i++) + if (supported_formats[i] == fmt) + return 1; -+ return 0; -+} -+ + return 0; +- +-fail: +- av_log(ctx, AV_LOG_ERROR, "Error evaluating '%s'\n", expr); +- return ret; + } + +-static int cuda_pad_alloc_out_frames_ctx(AVFilterContext *ctx, AVBufferRef **out_frames_ctx, const int width, const int height) +static av_cold int pad_cuda_init(AVFilterContext *avctx) -+{ + { +- AVFilterLink *inlink = ctx->inputs[0]; +- FilterLink *inl = ff_filter_link(inlink); +- AVHWFramesContext *in_frames_ctx = (AVHWFramesContext *)inl->hw_frames_ctx->data; +- int ret; + PadCUDAContext *ctx = avctx->priv; -+ + +- *out_frames_ctx = av_hwframe_ctx_alloc(in_frames_ctx->device_ref); +- if (!*out_frames_ctx) { + ctx->frame = av_frame_alloc(); + if (!ctx->frame) -+ return AVERROR(ENOMEM); -+ + return AVERROR(ENOMEM); +- } + +- AVHWFramesContext *out_fc = (AVHWFramesContext *)(*out_frames_ctx)->data; +- out_fc->format = AV_PIX_FMT_CUDA; +- out_fc->sw_format = in_frames_ctx->sw_format; +- +- out_fc->width = FFALIGN(width, 32); +- out_fc->height = FFALIGN(height, 32); +- +- ret = av_hwframe_ctx_init(*out_frames_ctx); +- if (ret < 0) { +- av_log(ctx, AV_LOG_ERROR, "Failed to init output ctx\n"); +- av_buffer_unref(out_frames_ctx); +- return ret; +- } +- +- return 0; +-} +- +-static av_cold int cuda_pad_init(AVFilterContext *ctx) +-{ +- CUDAPadContext *s = ctx->priv; +- +- s->last_out_w = -1; +- s->last_out_h = -1; + ctx->tmp_frame = av_frame_alloc(); + if (!ctx->tmp_frame) + return AVERROR(ENOMEM); -+ -+ return 0; -+} -+ + + return 0; + } + +-static av_cold void cuda_pad_uninit(AVFilterContext *ctx) +static av_cold void pad_cuda_uninit(AVFilterContext *avctx) -+{ + { +- CUDAPadContext *s = ctx->priv; +- CUcontext dummy; + PadCUDAContext *ctx = avctx->priv; -+ + +- av_buffer_unref(&s->frames_ctx); + if (ctx->hwctx && ctx->cu_module) { + CudaFunctions *cu = ctx->hwctx->internal->cuda_dl; + CUcontext dummy; -+ + +- if (s->hwctx && s->cu_module) { +- CudaFunctions *cu = s->hwctx->internal->cuda_dl; +- AVCUDADeviceContext *device_hwctx = s->hwctx; +- CHECK_CU(cu->cuCtxPushCurrent(s->hwctx->cuda_ctx)); +- CHECK_CU(cu->cuModuleUnload(s->cu_module)); + CHECK_CU(cu->cuCtxPushCurrent(ctx->hwctx->cuda_ctx)); + CHECK_CU(cu->cuModuleUnload(ctx->cu_module)); -+ CHECK_CU(cu->cuCtxPopCurrent(&dummy)); -+ } -+ + CHECK_CU(cu->cuCtxPopCurrent(&dummy)); + } + +- s->cu_module = NULL; +- s->hwctx = NULL; + av_frame_free(&ctx->frame); + av_buffer_unref(&ctx->hw_device_ctx); + av_buffer_unref(&ctx->frames_ctx); + av_frame_free(&ctx->tmp_frame); + ctx->hwctx = NULL; -+} -+ + } + +-static av_cold int cuda_pad_load_functions(AVFilterContext *ctx) +static av_cold int pad_cuda_output_config_props(AVFilterLink *outlink) -+{ -+ extern const unsigned char ff_vf_pad_cuda_ptx_data[]; -+ extern const unsigned int ff_vf_pad_cuda_ptx_len; -+ + { +- CUDAPadContext *s = ctx->priv; +- CudaFunctions *cu = s->hwctx->internal->cuda_dl; +- CUcontext dummy_cu_ctx; +- int ret; +- +- AVCUDADeviceContext *device_hwctx = s->hwctx; +- + extern const unsigned char ff_vf_pad_cuda_ptx_data[]; + extern const unsigned int ff_vf_pad_cuda_ptx_len; + +- ret = CHECK_CU(cu->cuCtxPushCurrent(device_hwctx->cuda_ctx)); +- if (ret < 0) +- return ret; +- +- ret = ff_cuda_load_module(ctx, device_hwctx, &s->cu_module, +- ff_vf_pad_cuda_ptx_data, ff_vf_pad_cuda_ptx_len); +- if (ret < 0) { +- av_log(ctx, AV_LOG_ERROR, "Failed to load CUDA module\n"); +- goto end; +- } +- +- ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_uchar, s->cu_module, "pad_uchar")); +- if (ret < 0) { +- av_log(ctx, AV_LOG_ERROR, "Failed to load pad_planar_cuda\n"); +- goto end; +- } +- +- ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_uchar2, s->cu_module, "pad_uchar2")); +- if (ret < 0) +- av_log(ctx, AV_LOG_ERROR, "Failed to load pad_uv_cuda\n"); +- +-end: +- CHECK_CU(cu->cuCtxPopCurrent(&dummy_cu_ctx)); +- +- return ret; +-} +- +-static int cuda_pad_config_props(AVFilterLink *outlink) +-{ +- AVFilterContext *ctx = outlink->src; +- CUDAPadContext *s = ctx->priv; +- +- AVFilterLink *inlink = ctx->inputs[0]; + FilterLink *outl = ff_filter_link(outlink); + AVFilterContext *avctx = outlink->src; + AVFilterLink *inlink = outlink->src->inputs[0]; -+ FilterLink *inl = ff_filter_link(inlink); + FilterLink *inl = ff_filter_link(inlink); +- +- FilterLink *ol = ff_filter_link(outlink); +- +- AVHWFramesContext *in_frames_ctx; +- int format_supported = 0; +- int ret; +- +- s->in_w = inlink->w; +- s->in_h = inlink->h; +- ret = eval_expr(ctx); +- if (ret < 0) +- return ret; +- +- if (!inl->hw_frames_ctx) { +- av_log(ctx, AV_LOG_ERROR, "No hw context provided on input\n"); + PadCUDAContext *ctx = avctx->priv; + AVHWFramesContext *frames_ctx = (AVHWFramesContext*)inl->hw_frames_ctx->data; + @@ -2635,20 +2913,43 @@ index 0000000..7daa227 + + if (!frames_ctx) { + av_log(ctx, AV_LOG_ERROR, "No hw context provided on main input\n"); -+ return AVERROR(EINVAL); -+ } -+ + return AVERROR(EINVAL); + } + +- in_frames_ctx = (AVHWFramesContext *)inl->hw_frames_ctx->data; +- s->hwctx = in_frames_ctx->device_ctx->hwctx; +- +- for (int i = 0; i < FF_ARRAY_ELEMS(supported_formats); i++) { +- if (in_frames_ctx->sw_format == supported_formats[i]) { +- format_supported = 1; +- break; +- } + ctx->sw_format = frames_ctx->sw_format; + if (!format_is_supported(ctx->sw_format)) { + av_log(ctx, AV_LOG_ERROR, "Unsupported input format: %s\n", + av_get_pix_fmt_name(ctx->sw_format)); + return AVERROR(ENOSYS); -+ } -+ + } +- if (!format_supported) { +- av_log(ctx, AV_LOG_ERROR, "Unsupported input format.\n"); +- return AVERROR(EINVAL); +- } +- +- s->parsed_color[0] = RGB_TO_Y_BT709(s->rgba_color[0], s->rgba_color[1], s->rgba_color[2]); +- s->parsed_color[1] = RGB_TO_U_BT709(s->rgba_color[0], s->rgba_color[1], s->rgba_color[2], 0); +- s->parsed_color[2] = RGB_TO_V_BT709(s->rgba_color[0], s->rgba_color[1], s->rgba_color[2], 0); +- s->parsed_color[3] = s->rgba_color[3]; + +- ret = cuda_pad_alloc_out_frames_ctx(ctx, &s->frames_ctx, s->w, s->h); +- if (ret < 0) +- return ret; +- +- ol->hw_frames_ctx = av_buffer_ref(s->frames_ctx); +- if (!ol->hw_frames_ctx) + // initialize + ctx->hw_device_ctx = av_buffer_ref(frames_ctx->device_ref); + if (!ctx->hw_device_ctx) -+ return AVERROR(ENOMEM); + return AVERROR(ENOMEM); + ctx->hwctx = ((AVHWDeviceContext*)frames_ctx->device_ref->data)->hwctx; + ctx->cu_stream = ctx->hwctx->stream; + @@ -2662,22 +2963,58 @@ index 0000000..7daa227 + if (err < 0) { + return err; + } -+ + +- outlink->w = s->w; +- outlink->h = s->h; +- outlink->time_base = inlink->time_base; +- outlink->format = AV_PIX_FMT_CUDA; +- +- s->last_out_w = s->w; +- s->last_out_h = s->h; +- +- ret = cuda_pad_load_functions(ctx); +- if (ret < 0) +- return ret; +- +- return 0; +-} +- +-static int cuda_pad_pad(AVFilterContext *ctx, AVFrame *out, const AVFrame *in) +-{ +- CUDAPadContext *s = ctx->priv; +- FilterLink *inl = ff_filter_link(ctx->inputs[0]); +- +- AVHWFramesContext *in_frames_ctx = (AVHWFramesContext *)inl->hw_frames_ctx->data; +- const AVPixFmtDescriptor *pixdesc = av_pix_fmt_desc_get(in_frames_ctx->sw_format); +- +- CudaFunctions *cu = s->hwctx->internal->cuda_dl; +- AVCUDADeviceContext *device_hwctx = s->hwctx; +- int ret; +- + err = ff_cuda_load_module(ctx, ctx->hwctx, &ctx->cu_module, ff_vf_pad_cuda_ptx_data, ff_vf_pad_cuda_ptx_len); + if (err < 0) { + CHECK_CU(cu->cuCtxPopCurrent(&dummy)); + return err; + } -+ + +- const int nb_planes = av_pix_fmt_count_planes(in_frames_ctx->sw_format); +- for (int plane = 0; plane < nb_planes; plane++) { +- const AVComponentDescriptor *cur_comp = &pixdesc->comp[0]; +- for (int comp = 1; comp < pixdesc->nb_components && cur_comp->plane != plane; comp++) +- cur_comp = &pixdesc->comp[comp]; + err = CHECK_CU(cu->cuModuleGetFunction(&ctx->cu_func, ctx->cu_module, "Pad_Cuda")); + if (err < 0) { + CHECK_CU(cu->cuCtxPopCurrent(&dummy)); + return err; + } -+ + +- int hsub = (plane == 1 || plane == 2) ? pixdesc->log2_chroma_w : 0; +- int vsub = (plane == 1 || plane == 2) ? pixdesc->log2_chroma_h : 0; + CHECK_CU(cu->cuCtxPopCurrent(&dummy)); + } -+ + +- int src_w = AV_CEIL_RSHIFT(s->in_w, hsub); +- int src_h = AV_CEIL_RSHIFT(s->in_h, vsub); + // process filter parameters + { + const AVPixFmtDescriptor *desc = av_pix_fmt_desc_get(ctx->sw_format); @@ -2720,7 +3057,9 @@ index 0000000..7daa227 + ctx->w = var_values[VAR_OUT_W] = var_values[VAR_OW] = av_rescale(ctx->h, adjusted_aspect.num, adjusted_aspect.den); + } + } -+ + +- int dst_w = AV_CEIL_RSHIFT(s->w, hsub); +- int dst_h = AV_CEIL_RSHIFT(s->h, vsub); + /* evaluate x and y */ + av_expr_parse_and_eval(&res, ctx->x_expr, + var_names, var_values, @@ -2737,18 +3076,27 @@ index 0000000..7daa227 + NULL, NULL, NULL, NULL, NULL, 0, ctx)) < 0) + return ret; + ctx->x = var_values[VAR_X] = res; -+ + +- int y_plane_offset = AV_CEIL_RSHIFT(s->y, vsub); +- int x_plane_offset = AV_CEIL_RSHIFT(s->x, hsub); + if (ctx->x < 0 || ctx->x + inlink->w > ctx->w) + ctx->x = var_values[VAR_X] = (ctx->w - inlink->w) / 2; + if (ctx->y < 0 || ctx->y + inlink->h > ctx->h) + ctx->y = var_values[VAR_Y] = (ctx->h - inlink->h) / 2; -+ + +- if (x_plane_offset + src_w > dst_w || y_plane_offset + src_h > dst_h) { +- av_log(ctx, AV_LOG_ERROR, +- "ROI out of bounds in plane %d: offset=(%d,%d) in=(%dx%d) " +- "out=(%dx%d)\n", +- plane, x_plane_offset, y_plane_offset, src_w, src_h, dst_w, dst_h); + /* sanity check params */ + if (ctx->w < inlink->w || ctx->h < inlink->h) { + av_log(ctx, AV_LOG_ERROR, "Padded dimensions cannot be smaller than input dimensions.\n"); -+ return AVERROR(EINVAL); -+ } -+ + return AVERROR(EINVAL); + } + +- int dst_linesize = out->linesize[plane] / cur_comp->step; +- int src_linesize = in->linesize[plane] / cur_comp->step; + // Align x, y offsets between planes + ctx->x &= ~((1 << desc->log2_chroma_w) - 1); + ctx->y &= ~((1 << desc->log2_chroma_h) - 1); @@ -2759,18 +3107,32 @@ index 0000000..7daa227 + ctx->pad_color[3] = ctx->pad_rgba[3]; + } + -+ + +- CUdeviceptr d_dst = (CUdeviceptr)out->data[plane]; +- CUdeviceptr d_src = (CUdeviceptr)in->data[plane]; + outlink->w = ctx->w; + outlink->h = ctx->h; -+ -+ + +- CUfunction cuda_func; + +- if (cur_comp->step == 1 && cur_comp->depth == 8) +- cuda_func = s->cu_func_uchar; +- else if(cur_comp->step == 2 && cur_comp->depth == 8) +- cuda_func = s->cu_func_uchar2; +- else +- return AVERROR_BUG; + // prepare output buffer + { + AVHWFramesContext *out_ctx; + AVBufferRef *out_ref = av_hwframe_ctx_alloc(ctx->hw_device_ctx); + if (!out_ref) + return AVERROR(ENOMEM); -+ + +- void *kernel_args[] = { +- &d_dst, &dst_linesize, &dst_w, &dst_h, +- &d_src, &src_linesize, &src_w, &src_h, +- &x_plane_offset, &y_plane_offset, &s->parsed_color[plane] +- }; + out_ctx = (AVHWFramesContext*)out_ref->data; + out_ctx->format = AV_PIX_FMT_CUDA; + out_ctx->sw_format = ctx->sw_format; @@ -2780,7 +3142,9 @@ index 0000000..7daa227 + ret = av_hwframe_ctx_init(out_ref); + if (ret < 0) + goto output_buffer_fail; -+ + +- unsigned int grid_x = DIV_UP(dst_w, BLOCK_X); +- unsigned int grid_y = DIV_UP(dst_h, BLOCK_Y); + av_frame_unref(ctx->frame); + ret = av_hwframe_get_buffer(out_ref, ctx->frame, 0); + if (ret < 0) @@ -2788,24 +3152,38 @@ index 0000000..7daa227 + + ctx->frame->width = ctx->w; + ctx->frame->height = ctx->h; -+ + +- ret = CHECK_CU(cu->cuLaunchKernel(cuda_func, grid_x, grid_y, 1, +- BLOCK_X, BLOCK_Y, 1, +- 0, s->hwctx->stream, kernel_args, NULL)); + ctx->frames_ctx = out_ref; -+ -+ if (ret < 0) { + + if (ret < 0) { +- av_log(ctx, AV_LOG_ERROR, "Failed to launch kernel for plane %d\n", plane); +output_buffer_fail: + av_buffer_unref(&out_ref); -+ return ret; -+ } + return ret; + } + + outl->hw_frames_ctx = av_buffer_ref(inl->hw_frames_ctx); + if (!outl->hw_frames_ctx) { + return AVERROR(ENOMEM); + } -+ } -+ -+ return 0; -+} -+ + } + + return 0; + } + +-static int cuda_pad_filter_frame(AVFilterLink *inlink, AVFrame *in) +-{ +- AVFilterContext *ctx = inlink->dst; +- CUDAPadContext *s = ctx->priv; +- AVFilterLink *outlink = ctx->outputs[0]; +- +- FilterLink *outl = ff_filter_link(outlink); +- +- AVHWFramesContext *out_frames_ctx = (AVHWFramesContext *)outl->hw_frames_ctx->data; +- AVCUDADeviceContext *device_hwctx = out_frames_ctx->device_ctx->hwctx; +static int pad_cuda_call_kernel( + PadCUDAContext *ctx, + uint8_t *main_data, int main_linesize, @@ -2831,23 +3209,32 @@ index 0000000..7daa227 + BLOCK_X, BLOCK_Y, 1, + 0, ctx->cu_stream, kernel_args, NULL)); +} -+ + +static int pad_cuda_fill_buffers(AVFilterContext *avctx, AVFrame *out, AVFrame *in) +{ + PadCUDAContext *ctx = avctx->priv; + CudaFunctions *cu = ctx->hwctx->internal->cuda_dl; + CUcontext cuda_ctx = ctx->hwctx->cuda_ctx; + CUcontext dummy; -+ int ret; + int ret; + uint8_t color_y = ctx->pad_color[0]; + uint8_t color_u = ctx->pad_color[1]; + uint8_t color_v = ctx->pad_color[2]; + uint8_t color_a = ctx->pad_color[3]; -+ + +- if (s->eval_mode == EVAL_MODE_FRAME) { +- s->in_w = in->width; +- s->in_h = in->height; +- s->aspect = in->sample_aspect_ratio; + ret = CHECK_CU(cu->cuCtxPushCurrent(cuda_ctx)); + if (ret < 0) + return ret; -+ + +- ret = eval_expr(ctx); +- if (ret < 0) { +- av_frame_free(&in); +- return ret; +- } + // overlay first plane + pad_cuda_call_kernel(ctx, + out->data[0], out->linesize[0], @@ -2910,8 +3297,15 @@ index 0000000..7daa227 + av_frame_free(&out); + CHECK_CU(cu->cuCtxPopCurrent(&dummy)); + return AVERROR_BUG; -+ } -+ + } + +- +- if (s->x == 0 && s->y == 0 && +- s->w == in->width && s->h == in->height) { +- av_log(ctx, AV_LOG_DEBUG, "No border. Passing the frame unmodified.\n"); +- s->last_out_w = s->w; +- s->last_out_h = s->h; +- return ff_filter_frame(outlink, in); + if (in->data[3]) { + pad_cuda_call_kernel(ctx, + out->data[3], out->linesize[3], @@ -2920,29 +3314,32 @@ index 0000000..7daa227 + in->width, in->height, + ctx->x, ctx->y, + color_a, color_a); -+ } + } + + CHECK_CU(cu->cuCtxPopCurrent(&dummy)); + return 0; +} -+ + +static int pad_cuda_filter_frame(AVFilterLink *link, AVFrame *in) +{ + AVFilterContext *avctx = link->dst; + PadCUDAContext *ctx = avctx->priv; + AVFilterLink *outlink = avctx->outputs[0]; + CudaFunctions *cu = ctx->hwctx->internal->cuda_dl; -+ + +- if (s->w != s->last_out_w || s->h != s->last_out_h) { + AVFrame *out = NULL; + CUcontext dummy; + int ret = 0; -+ + +- av_buffer_unref(&s->frames_ctx); + out = av_frame_alloc(); + if (!out) { + ret = AVERROR(ENOMEM); + goto fail; + } -+ + +- ret = cuda_pad_alloc_out_frames_ctx(ctx, &s->frames_ctx, s->w, s->h); + ret = CHECK_CU(cu->cuCtxPushCurrent(ctx->hwctx->cuda_ctx)); + if (ret < 0) + goto fail; @@ -2950,41 +3347,105 @@ index 0000000..7daa227 + // fill and prepare the output frame + { + ret = pad_cuda_fill_buffers(avctx, ctx->frame, in); -+ if (ret < 0) -+ return ret; -+ + if (ret < 0) + return ret; + +- av_buffer_unref(&outl->hw_frames_ctx); +- outl->hw_frames_ctx = av_buffer_ref(s->frames_ctx); +- if (!outl->hw_frames_ctx) { +- av_frame_free(&in); +- av_log(ctx, AV_LOG_ERROR, "Failed to allocate output frame context.\n"); +- return AVERROR(ENOMEM); +- } +- outlink->w = s->w; +- outlink->h = s->h; +- +- s->last_out_w = s->w; +- s->last_out_h = s->h; +- } +- +- AVFrame *out = av_frame_alloc(); +- if (!out) { +- av_frame_free(&in); +- av_log(ctx, AV_LOG_ERROR, "Failed to allocate output AVFrame.\n"); +- return AVERROR(ENOMEM); +- } +- ret = av_hwframe_get_buffer(outl->hw_frames_ctx, out, 0); +- if (ret < 0) { +- av_log(ctx, AV_LOG_ERROR, "Unable to get output buffer: %s\n", +- av_err2str(ret)); +- av_frame_free(&out); +- av_frame_free(&in); +- return ret; +- } +- +- CUcontext dummy; +- ret = CHECK_CU(device_hwctx->internal->cuda_dl->cuCtxPushCurrent( +- device_hwctx->cuda_ctx)); +- if (ret < 0) { +- av_frame_free(&out); +- av_frame_free(&in); +- return ret; +- } + ret = av_hwframe_get_buffer(ctx->frame->hw_frames_ctx, ctx->tmp_frame, 0); + if (ret < 0) + return ret; -+ + +- ret = cuda_pad_pad(ctx, out, in); + av_frame_move_ref(out, ctx->frame); + av_frame_move_ref(ctx->frame, ctx->tmp_frame); -+ + +- CHECK_CU(device_hwctx->internal->cuda_dl->cuCtxPopCurrent(&dummy)); + ctx->frame->width = outlink->w; + ctx->frame->height = outlink->h; -+ + +- if (ret < 0) { +- av_frame_free(&out); +- av_frame_free(&in); +- return ret; + ret = av_frame_copy_props(out, in); + if (ret < 0) + return ret; -+ } -+ + } + +- av_frame_copy_props(out, in); +- out->width = s->w; +- out->height = s->h; +- + CHECK_CU(cu->cuCtxPopCurrent(&dummy)); + if (ret < 0) + goto fail; -+ -+ av_reduce(&out->sample_aspect_ratio.num, &out->sample_aspect_ratio.den, + + av_reduce(&out->sample_aspect_ratio.num, &out->sample_aspect_ratio.den, +- (int64_t)in->sample_aspect_ratio.num * out->height * in->width, +- (int64_t)in->sample_aspect_ratio.den * out->width * in->height, + (int64_t)in->sample_aspect_ratio.num * outlink->h * link->w, + (int64_t)in->sample_aspect_ratio.den * outlink->w * link->h, -+ INT_MAX); -+ -+ av_frame_free(&in); -+ return ff_filter_frame(outlink, out); + INT_MAX); + + av_frame_free(&in); + return ff_filter_frame(outlink, out); +fail: + av_frame_free(&in); + av_frame_free(&out); + return ret; -+} -+ + } + +-#define OFFSET(x) offsetof(CUDAPadContext, x) +-#define FLAGS (AV_OPT_FLAG_FILTERING_PARAM | AV_OPT_FLAG_VIDEO_PARAM) +- +-static const AVOption cuda_pad_options[] = { +- { "width", "set the pad area width expression", OFFSET(w_expr), AV_OPT_TYPE_STRING, {.str = "iw"}, 0, 0, FLAGS }, +- { "w", "set the pad area width expression", OFFSET(w_expr), AV_OPT_TYPE_STRING, {.str = "iw"}, 0, 0, FLAGS }, +- { "height", "set the pad area height expression", OFFSET(h_expr), AV_OPT_TYPE_STRING, {.str = "ih"}, 0, 0, FLAGS }, +- { "h", "set the pad area height expression", OFFSET(h_expr), AV_OPT_TYPE_STRING, {.str = "ih"}, 0, 0, FLAGS }, +- { "x", "set the x offset expression for the input image position", OFFSET(x_expr), AV_OPT_TYPE_STRING, {.str = "0"}, 0, 0, FLAGS }, +- { "y", "set the y offset expression for the input image position", OFFSET(y_expr), AV_OPT_TYPE_STRING, {.str = "0"}, 0, 0, FLAGS }, +- { "color", "set the color of the padded area border", OFFSET(rgba_color), AV_OPT_TYPE_COLOR, {.str = "black"}, .flags = FLAGS }, +- { "eval", "specify when to evaluate expressions", OFFSET(eval_mode), AV_OPT_TYPE_INT, {.i64 = EVAL_MODE_INIT}, 0, EVAL_MODE_NB-1, FLAGS, .unit = "eval" }, +- { "init", "eval expressions once during initialization", 0, AV_OPT_TYPE_CONST, {.i64=EVAL_MODE_INIT}, .flags = FLAGS, .unit = "eval" }, +- { "frame", "eval expressions during initialization and per-frame", 0, AV_OPT_TYPE_CONST, {.i64=EVAL_MODE_FRAME}, .flags = FLAGS, .unit = "eval" }, +- { "aspect", "pad to fit an aspect instead of a resolution", OFFSET(aspect), AV_OPT_TYPE_RATIONAL, {.dbl = 0}, 0, DBL_MAX, FLAGS }, + + +#define OFFSET(x) offsetof(PadCUDAContext, x) @@ -2999,11 +3460,22 @@ index 0000000..7daa227 + { "y", "set the y offset for the input image position", OFFSET(y_expr), AV_OPT_TYPE_STRING, {.str = "0"}, 0, INT16_MAX, FLAGS }, + { "color", "set the color of the padded area border", OFFSET(pad_rgba), AV_OPT_TYPE_COLOR, { .str = "black" }, 0, 0, FLAGS }, + { "aspect", "pad to fit an aspect instead of a resolution", OFFSET(aspect), AV_OPT_TYPE_RATIONAL, {.dbl = 0}, 0, INT16_MAX, FLAGS }, -+ { NULL } -+}; -+ + { NULL } + }; + +-static const AVClass cuda_pad_class = { +- .class_name = "pad_cuda", +- .item_name = av_default_item_name, +- .option = cuda_pad_options, +- .version = LIBAVUTIL_VERSION_INT, +-}; +AVFILTER_DEFINE_CLASS(pad_cuda); -+ + +-static const AVFilterPad cuda_pad_inputs[] = {{ +- .name = "default", +- .type = AVMEDIA_TYPE_VIDEO, +- .filter_frame = cuda_pad_filter_frame +-}}; +static const AVFilterPad pad_cuda_inputs[] = { + { + .name = "default", @@ -3011,42 +3483,82 @@ index 0000000..7daa227 + .filter_frame = pad_cuda_filter_frame, + }, +}; -+ + +-static const AVFilterPad cuda_pad_outputs[] = {{ +static const AVFilterPad pad_cuda_outputs[] = { + { -+ .name = "default", -+ .type = AVMEDIA_TYPE_VIDEO, + .name = "default", + .type = AVMEDIA_TYPE_VIDEO, +- .config_props = cuda_pad_config_props, +-}}; + .config_props = pad_cuda_output_config_props, + }, +}; -+ -+const AVFilter ff_vf_pad_cuda = { -+ .name = "pad_cuda", -+ .description = NULL_IF_CONFIG_SMALL("Pad CUDA accelerated video using solid color"), + + const FFFilter ff_vf_pad_cuda = { +- .p.name = "pad_cuda", +- .p.description = NULL_IF_CONFIG_SMALL("CUDA-based GPU padding filter"), +- .init = cuda_pad_init, +- .uninit = cuda_pad_uninit, +- +- .p.priv_class = &cuda_pad_class, +- +- FILTER_INPUTS(cuda_pad_inputs), +- FILTER_OUTPUTS(cuda_pad_outputs), +- ++ .p.name = "pad_cuda", ++ .p.description = NULL_IF_CONFIG_SMALL("Pad CUDA accelerated video using solid color"), + .priv_size = sizeof(PadCUDAContext), -+ .priv_class = &pad_cuda_class, ++ .p.priv_class = &pad_cuda_class, + .init = pad_cuda_init, + .uninit = pad_cuda_uninit, + FILTER_INPUTS(pad_cuda_inputs), + FILTER_OUTPUTS(pad_cuda_outputs), -+ FILTER_SINGLE_PIXFMT(AV_PIX_FMT_CUDA), -+ .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE, -+}; + FILTER_SINGLE_PIXFMT(AV_PIX_FMT_CUDA), +- +- .priv_size = sizeof(CUDAPadContext), + .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE, + }; diff --git a/libavfilter/vf_pad_cuda.cu b/libavfilter/vf_pad_cuda.cu -new file mode 100644 -index 0000000..62f7988 ---- /dev/null +index f1323d122f..62f7988de5 100644 +--- a/libavfilter/vf_pad_cuda.cu +++ b/libavfilter/vf_pad_cuda.cu -@@ -0,0 +1,37 @@ +@@ -1,61 +1,37 @@ +-/* +- * 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 +- */ +extern "C" { -+ + +-template +-__device__ void pad_impl(T* dst, int dst_pitch, int dst_w, int dst_h, +- const T* src, int src_pitch, int src_w, int src_h, +- int roi_x, int roi_y, T fill_val) +__global__ void Pad_Cuda( + unsigned char* main, int main_linesize, + unsigned char* overlay, int overlay_linesize, + int overlay_w, int overlay_h, + int offset_x, int offset_y, + unsigned char fill_color1, unsigned char fill_color2) -+{ + { +- const int x = blockIdx.x * blockDim.x + threadIdx.x; +- const int y = blockIdx.y * blockDim.y + threadIdx.y; +- +- if (x >= dst_w || y >= dst_h) { +- return; + // + // We pass two colors to handle NV12 plane with interleaved UV. + // So fill_color1 would correspond to U and fill_color2 to V. @@ -3065,21 +3577,49 @@ index 0000000..62f7988 + { + // Picking colors[0] for even value of x; and colors[1] for odd value of x; + *m = color[x & 1]; -+ } + } +- +- if (x >= roi_x && x < (roi_x + src_w) && y >= roi_y && y < (roi_y + src_h)) { +- const int src_x = x - roi_x; +- const int src_y = y - roi_y; +- dst[y * dst_pitch + x] = src[src_y * src_pitch + src_x]; +- } else { +- dst[y * dst_pitch + x] = fill_val; + else + { + int overlay_x = x - offset_x; + int overlay_y = y - offset_y; + *m = overlay[overlay_x + overlay_y*overlay_linesize]; -+ } -+} -+ -+} + } + } + +- +-extern "C" { +- +-__global__ void pad_uchar(unsigned char* dst, int dst_pitch, int dst_w, int dst_h, +- const unsigned char* src, int src_pitch, int src_w, int src_h, +- int roi_x, int roi_y, unsigned char fill_val) +-{ +- pad_impl(dst, dst_pitch, dst_w, dst_h, +- src, src_pitch, src_w, src_h, +- roi_x, roi_y, fill_val); +-} +- +-__global__ void pad_uchar2(uchar2* dst, int dst_pitch, int dst_w, int dst_h, +- const uchar2* src, int src_pitch, int src_w, int src_h, +- int roi_x, int roi_y, uchar2 fill_val) +-{ +- pad_impl(dst, dst_pitch, dst_w, dst_h, +- src, src_pitch, src_w, src_h, +- roi_x, roi_y, fill_val); +-} +- + } diff --git a/libavfilter/vf_scale_cuda.c b/libavfilter/vf_scale_cuda.c -index 54a3409..2db84dd 100644 +index 5fd757161b..56c4d6350d 100644 --- a/libavfilter/vf_scale_cuda.c +++ b/libavfilter/vf_scale_cuda.c -@@ -446,6 +446,11 @@ static int scalecuda_resize(AVFilterContext *ctx, +@@ -487,6 +487,11 @@ static int scalecuda_resize(AVFilterContext *ctx, for (i = 0; i < s->in_planes; i++) { CUDA_TEXTURE_DESC tex_desc = { @@ -3092,10 +3632,10 @@ index 54a3409..2db84dd 100644 CU_TR_FILTER_MODE_LINEAR : CU_TR_FILTER_MODE_POINT, diff --git a/libavfilter/vf_scale_cuda.cu b/libavfilter/vf_scale_cuda.cu -index de06ba9..414649e 100644 +index d674c0885a..eec20db1a1 100644 --- a/libavfilter/vf_scale_cuda.cu +++ b/libavfilter/vf_scale_cuda.cu -@@ -1013,24 +1013,29 @@ struct Convert_rgb0_rgba +@@ -1043,24 +1043,29 @@ struct Convert_rgb0_rgba typedef float4 (*coeffs_function_t)(float, float); @@ -3139,7 +3679,7 @@ index de06ba9..414649e 100644 return res / (res.x + res.y + res.z + res.w); } -@@ -1059,6 +1064,71 @@ __device__ static inline V apply_coeffs(float4 coeffs, V c0, V c1, V c2, V c3) +@@ -1089,6 +1094,71 @@ __device__ static inline V apply_coeffs(float4 coeffs, V c0, V c1, V c2, V c3) return res; } @@ -3211,7 +3751,7 @@ index de06ba9..414649e 100644 template __device__ static inline T Subsample_Nearest(cudaTextureObject_t tex, int xo, int yo, -@@ -1126,9 +1196,11 @@ __device__ static inline T Subsample_Bicubic(cudaTextureObject_t tex, +@@ -1159,9 +1229,11 @@ __device__ static inline T Subsample_Bicubic(cudaTextureObject_t tex, float4 coeffsX = coeffs_function(fx, param); float4 coeffsY = coeffs_function(fy, param); @@ -3227,7 +3767,7 @@ index de06ba9..414649e 100644 apply_coeffs(coeffsX, PIX(px - 1, py ), PIX(px, py ), PIX(px + 1, py ), PIX(px + 2, py )), diff --git a/libavfilter/vf_transition_cuda.c b/libavfilter/vf_transition_cuda.c new file mode 100644 -index 0000000..44a198f +index 0000000000..0c055a71e0 --- /dev/null +++ b/libavfilter/vf_transition_cuda.c @@ -0,0 +1,621 @@ @@ -3501,11 +4041,11 @@ index 0000000..44a198f + } + + if (ctx->eval_mode == EVAL_MODE_FRAME) { -+ int pos = input_main->pkt_pos; + ctx->var_values[VAR_N] = inl->frame_count_out; + ctx->var_values[VAR_T] = input_main->pts == AV_NOPTS_VALUE ? + NAN : input_main->pts * av_q2d(inlink->time_base); -+ ctx->var_values[VAR_POS] = pos == -1 ? NAN : pos; ++ // FFmpeg 8 no longer exposes the packet byte position on AVFrame. ++ ctx->var_values[VAR_POS] = NAN; + ctx->var_values[VAR_OVERLAY_W] = ctx->var_values[VAR_OW] = input_overlay->width; + ctx->var_values[VAR_OVERLAY_H] = ctx->var_values[VAR_OH] = input_overlay->height; + ctx->var_values[VAR_MAIN_W ] = ctx->var_values[VAR_MW] = input_main->width; @@ -3837,11 +4377,11 @@ index 0000000..44a198f + }, +}; + -+const AVFilter ff_vf_transition_cuda = { -+ .name = "transition_cuda", -+ .description = NULL_IF_CONFIG_SMALL("Transition between videos using CUDA"), ++const FFFilter ff_vf_transition_cuda = { ++ .p.name = "transition_cuda", ++ .p.description = NULL_IF_CONFIG_SMALL("Transition between videos using CUDA"), + .priv_size = sizeof(TransitionCUDAContext), -+ .priv_class = &transition_cuda_class, ++ .p.priv_class = &transition_cuda_class, + .init = &transition_cuda_init, + .uninit = &transition_cuda_uninit, + .activate = &transition_cuda_activate, @@ -3854,7 +4394,7 @@ index 0000000..44a198f +}; diff --git a/libavfilter/vf_transition_cuda.cu b/libavfilter/vf_transition_cuda.cu new file mode 100644 -index 0000000..c688e66 +index 0000000000..c688e66cca --- /dev/null +++ b/libavfilter/vf_transition_cuda.cu @@ -0,0 +1,56 @@ @@ -3915,14 +4455,14 @@ index 0000000..c688e66 + +} diff --git a/libavutil/hwcontext_cuda.c b/libavutil/hwcontext_cuda.c -index 3de3847..f8ff9a4 100644 +index b0b65b2446..372c4722b5 100644 --- a/libavutil/hwcontext_cuda.c +++ b/libavutil/hwcontext_cuda.c -@@ -44,6 +44,7 @@ static const enum AVPixelFormat supported_formats[] = { - AV_PIX_FMT_NV12, +@@ -46,6 +46,7 @@ static const enum AVPixelFormat supported_formats[] = { AV_PIX_FMT_YUV420P, AV_PIX_FMT_YUVA420P, -+ AV_PIX_FMT_YUVA444P, AV_PIX_FMT_YUV444P, ++ AV_PIX_FMT_YUVA444P, AV_PIX_FMT_P010, AV_PIX_FMT_P016, + AV_PIX_FMT_P210, diff --git a/deps/ffmpeg-patches/0003-avfilter-npp-cuda13-compat.patch b/deps/ffmpeg/8/0003-avfilter-npp-cuda13-compat.patch similarity index 72% rename from deps/ffmpeg-patches/0003-avfilter-npp-cuda13-compat.patch rename to deps/ffmpeg/8/0003-avfilter-npp-cuda13-compat.patch index f745d185..d82d6839 100644 --- a/deps/ffmpeg-patches/0003-avfilter-npp-cuda13-compat.patch +++ b/deps/ffmpeg/8/0003-avfilter-npp-cuda13-compat.patch @@ -1,11 +1,8 @@ -From e5ba36a6cfabc10f5c745fba9e8a91fef92cf4a7 Mon Sep 17 00:00:00 2001 -From: Jan Pietek -Date: Wed, 12 Aug 2026 15:54:13 +0000 -Subject: [PATCH 3/7] avfilter/npp: support CUDA 13 stream context APIs +From 585909cb4234d7ed8af27158a5fb53bf0c2a2c87 Mon Sep 17 00:00:00 2001 +From: x +Date: Wed, 16 Sep 2026 20:00:13 +0200 +Subject: [PATCH] avfilter/npp: CUDA 13 stream-context compatibility -Add an NPP compatibility layer and use explicit CUDA stream contexts in scale, sharpen, and transpose filters. - -Original-commit: 30b851f13d97bad9b0b28e7310122d914ac6fc1d --- libavfilter/cuda/npp_compat.h | 65 ++++++++++++++++++++++++++++++++++ libavfilter/vf_scale_npp.c | 56 ++++++++++++++++++++--------- @@ -86,20 +83,14 @@ index 0000000..bf5d73e + +#endif /* AVFILTER_CUDA_NPP_COMPAT_H */ diff --git a/libavfilter/vf_scale_npp.c b/libavfilter/vf_scale_npp.c -index 0c38987..f90b291 100644 +index 8e91135..de179ed 100644 --- a/libavfilter/vf_scale_npp.c +++ b/libavfilter/vf_scale_npp.c -@@ -36,6 +36,7 @@ - #include "libavutil/pixdesc.h" - +@@ -38,2 +38,3 @@ #include "avfilter.h" +#include "cuda/npp_compat.h" #include "filters.h" - #include "formats.h" - #include "scale_eval.h" -@@ -708,14 +709,22 @@ static int nppscale_deinterleave(AVFilterContext *ctx, NPPScaleStageContext *sta - AVFrame *out, AVFrame *in) - { +@@ -698,3 +699,10 @@ static int nppscale_deinterleave(AVFilterContext *ctx, NPPScaleStageContext *sta AVHWFramesContext *in_frames_ctx = (AVHWFramesContext*)in->hw_frames_ctx->data; + AVCUDADeviceContext *device_hwctx = in_frames_ctx->device_ctx->hwctx; + NppStreamContext npp_ctx; @@ -110,7 +101,7 @@ index 0c38987..f90b291 100644 + if (ret < 0) + return ret; - switch (in_frames_ctx->sw_format) { +@@ -702,6 +710,7 @@ static int nppscale_deinterleave(AVFilterContext *ctx, NPPScaleStageContext *sta case AV_PIX_FMT_NV12: - err = nppiYCbCr420_8u_P2P3R(in->data[0], in->linesize[0], - in->data[1], in->linesize[1], @@ -122,11 +113,7 @@ index 0c38987..f90b291 100644 + (NppiSize){ in->width, in->height }, + npp_ctx); break; - default: - return AVERROR_BUG; -@@ -731,9 +740,16 @@ static int nppscale_deinterleave(AVFilterContext *ctx, NPPScaleStageContext *sta - static int nppscale_resize(AVFilterContext *ctx, NPPScaleStageContext *stage, - AVFrame *out, AVFrame *in) +@@ -721,5 +730,12 @@ static int nppscale_resize(AVFilterContext *ctx, NPPScaleStageContext *stage, { + AVHWFramesContext *in_frames_ctx = (AVHWFramesContext*)in->hw_frames_ctx->data; + AVCUDADeviceContext *device_hwctx = in_frames_ctx->device_ctx->hwctx; @@ -140,11 +127,7 @@ index 0c38987..f90b291 100644 + if (ret < 0) + return ret; - for (i = 0; i < FF_ARRAY_ELEMS(stage->planes_in) && i < FF_ARRAY_ELEMS(in->data) && in->data[i]; i++) { - int iw = stage->planes_in[i].width; -@@ -741,12 +757,12 @@ static int nppscale_resize(AVFilterContext *ctx, NPPScaleStageContext *stage, - int ow = stage->planes_out[i].width; - int oh = stage->planes_out[i].height; +@@ -731,8 +747,8 @@ static int nppscale_resize(AVFilterContext *ctx, NPPScaleStageContext *stage, - err = nppiResizeSqrPixel_8u_C1R(in->data[i], (NppiSize){ iw, ih }, - in->linesize[i], (NppiRect){ 0, 0, iw, ih }, @@ -159,11 +142,7 @@ index 0c38987..f90b291 100644 + (double)ow / iw, (double)oh / ih, + 0.0, 0.0, s->interp_algo, npp_ctx); if (err != NPP_SUCCESS) { - av_log(ctx, AV_LOG_ERROR, "NPP resize error: %d\n", err); - return AVERROR_UNKNOWN; -@@ -760,15 +776,23 @@ static int nppscale_interleave(AVFilterContext *ctx, NPPScaleStageContext *stage - AVFrame *out, AVFrame *in) - { +@@ -750,3 +766,10 @@ static int nppscale_interleave(AVFilterContext *ctx, NPPScaleStageContext *stage AVHWFramesContext *out_frames_ctx = (AVHWFramesContext*)out->hw_frames_ctx->data; + AVCUDADeviceContext *device_hwctx = out_frames_ctx->device_ctx->hwctx; + NppStreamContext npp_ctx; @@ -174,7 +153,7 @@ index 0c38987..f90b291 100644 + if (ret < 0) + return ret; - switch (out_frames_ctx->sw_format) { +@@ -754,7 +777,8 @@ static int nppscale_interleave(AVFilterContext *ctx, NPPScaleStageContext *stage case AV_PIX_FMT_NV12: - err = nppiYCbCr420_8u_P3P2R((const uint8_t**)in->data, - in->linesize, @@ -188,27 +167,19 @@ index 0c38987..f90b291 100644 + (NppiSize){ in->width, in->height }, + npp_ctx); break; - default: - return AVERROR_BUG; diff --git a/libavfilter/vf_sharpen_npp.c b/libavfilter/vf_sharpen_npp.c -index 4989126..da086fb 100644 +index 3ec74f8..e3f3782 100644 --- a/libavfilter/vf_sharpen_npp.c +++ b/libavfilter/vf_sharpen_npp.c -@@ -24,6 +24,7 @@ - #include - #include +@@ -26,2 +26,3 @@ +#include "cuda/npp_compat.h" #include "filters.h" - #include "libavutil/pixdesc.h" - #include "libavutil/cuda_check.h" -@@ -157,17 +158,24 @@ static int nppsharpen_sharpen(AVFilterContext* ctx, AVFrame* out, AVFrame* in) - { - FilterLink *inl = ff_filter_link(ctx->inputs[0]); +@@ -161,2 +162,3 @@ static int nppsharpen_sharpen(AVFilterContext* ctx, AVFrame* out, AVFrame* in) AVHWFramesContext* in_ctx = (AVHWFramesContext*)inl->hw_frames_ctx->data; + AVCUDADeviceContext *device_hwctx = in_ctx->device_ctx->hwctx; NPPSharpenContext* s = ctx->priv; - +@@ -164,2 +166,8 @@ static int nppsharpen_sharpen(AVFilterContext* ctx, AVFrame* out, AVFrame* in) const AVPixFmtDescriptor *desc = av_pix_fmt_desc_get(in_ctx->sw_format); + NppStreamContext npp_ctx; + int ret; @@ -217,9 +188,7 @@ index 4989126..da086fb 100644 + if (ret < 0) + return ret; - for (int i = 0; i < FF_ARRAY_ELEMS(in->data) && in->data[i]; i++) { - int ow = AV_CEIL_RSHIFT(in->width, (i == 1 || i == 2) ? desc->log2_chroma_w : 0); - int oh = AV_CEIL_RSHIFT(in->height, (i == 1 || i == 2) ? desc->log2_chroma_h : 0); +@@ -169,5 +177,5 @@ static int nppsharpen_sharpen(AVFilterContext* ctx, AVFrame* out, AVFrame* in) - NppStatus err = nppiFilterSharpenBorder_8u_C1R( + NppStatus err = nppiFilterSharpenBorder_8u_C1R_Ctx( @@ -227,23 +196,15 @@ index 4989126..da086fb 100644 - out->data[i], out->linesize[i], (NppiSize){ow, oh}, s->border_type); + out->data[i], out->linesize[i], (NppiSize){ow, oh}, s->border_type, npp_ctx); if (err != NPP_SUCCESS) { - av_log(ctx, AV_LOG_ERROR, "NPP sharpen error: %d\n", err); - return AVERROR_EXTERNAL; diff --git a/libavfilter/vf_transpose_npp.c b/libavfilter/vf_transpose_npp.c -index 1706267..e24dc53 100644 +index 2315b10..ad43356 100644 --- a/libavfilter/vf_transpose_npp.c +++ b/libavfilter/vf_transpose_npp.c -@@ -29,6 +29,7 @@ - #include "libavutil/pixdesc.h" - +@@ -31,2 +31,3 @@ #include "avfilter.h" +#include "cuda/npp_compat.h" #include "filters.h" - #include "formats.h" - #include "video.h" -@@ -292,9 +293,16 @@ static int npptranspose_config_props(AVFilterLink *outlink) - static int npptranspose_rotate(AVFilterContext *ctx, NPPTransposeStageContext *stage, - AVFrame *out, AVFrame *in) +@@ -296,5 +297,12 @@ static int npptranspose_rotate(AVFilterContext *ctx, NPPTransposeStageContext *s { + AVHWFramesContext *in_frames_ctx = (AVHWFramesContext*)in->hw_frames_ctx->data; + AVCUDADeviceContext *device_hwctx = in_frames_ctx->device_ctx->hwctx; @@ -257,11 +218,7 @@ index 1706267..e24dc53 100644 + if (ret < 0) + return ret; - for (i = 0; i < FF_ARRAY_ELEMS(stage->planes_in) && i < FF_ARRAY_ELEMS(in->data) && in->data[i]; i++) { - int iw = stage->planes_in[i].width; -@@ -309,11 +317,11 @@ static int npptranspose_rotate(AVFilterContext *ctx, NPPTransposeStageContext *s - int shiftw = (s->dir == NPP_TRANSPOSE_CLOCK || s->dir == NPP_TRANSPOSE_CLOCK_FLIP) ? ow - 1 : 0; - int shifth = (s->dir == NPP_TRANSPOSE_CCLOCK || s->dir == NPP_TRANSPOSE_CLOCK_FLIP) ? oh - 1 : 0; +@@ -313,7 +321,7 @@ static int npptranspose_rotate(AVFilterContext *ctx, NPPTransposeStageContext *s - err = nppiRotate_8u_C1R(in->data[i], (NppiSize){ iw, ih }, - in->linesize[i], (NppiRect){ 0, 0, iw, ih }, @@ -274,11 +231,7 @@ index 1706267..e24dc53 100644 + (NppiRect){ 0, 0, ow, oh }, + angle, shiftw, shifth, NPPI_INTER_NN, npp_ctx); if (err != NPP_SUCCESS) { - av_log(ctx, AV_LOG_ERROR, "NPP rotate error: %d\n", err); - return AVERROR_UNKNOWN; -@@ -326,16 +334,23 @@ static int npptranspose_rotate(AVFilterContext *ctx, NPPTransposeStageContext *s - static int npptranspose_transpose(AVFilterContext *ctx, NPPTransposeStageContext *stage, - AVFrame *out, AVFrame *in) +@@ -330,4 +338,11 @@ static int npptranspose_transpose(AVFilterContext *ctx, NPPTransposeStageContext { + AVHWFramesContext *in_frames_ctx = (AVHWFramesContext*)in->hw_frames_ctx->data; + AVCUDADeviceContext *device_hwctx = in_frames_ctx->device_ctx->hwctx; @@ -291,9 +244,7 @@ index 1706267..e24dc53 100644 + if (ret < 0) + return ret; - for (i = 0; i < FF_ARRAY_ELEMS(stage->planes_in) && i < FF_ARRAY_ELEMS(in->data) && in->data[i]; i++) { - int iw = stage->planes_in[i].width; - int ih = stage->planes_in[i].height; +@@ -337,5 +352,5 @@ static int npptranspose_transpose(AVFilterContext *ctx, NPPTransposeStageContext - err = nppiTranspose_8u_C1R(in->data[i], in->linesize[i], - out->data[i], out->linesize[i], @@ -302,5 +253,6 @@ index 1706267..e24dc53 100644 + out->data[i], out->linesize[i], + (NppiSize){ iw, ih }, npp_ctx); if (err != NPP_SUCCESS) { - av_log(ctx, AV_LOG_ERROR, "NPP transpose error: %d\n", err); - return AVERROR_UNKNOWN; +-- +2.55.0 + diff --git a/deps/ffmpeg/8/0004-avcodec-nvdec-intra.patch b/deps/ffmpeg/8/0004-avcodec-nvdec-intra.patch new file mode 100644 index 00000000..dcdeb64b --- /dev/null +++ b/deps/ffmpeg/8/0004-avcodec-nvdec-intra.patch @@ -0,0 +1,27 @@ +From 5bcee64dd15453ff5c0bb88ce7b9d189a82c93a3 Mon Sep 17 00:00:00 2001 +From: Jan Pietek +Date: Tue, 15 Sep 2026 09:37:43 +0200 +Subject: [PATCH 4/7] avcodec/nvdec: preserve intra-only stream handling on 8.1 + +--- + libavcodec/cuviddec.c | 7 +++++++ + 1 file changed, 7 insertions(+) + +diff --git a/libavcodec/cuviddec.c b/libavcodec/cuviddec.c +index be183fce35..8a1b7e3042 100644 +--- a/libavcodec/cuviddec.c ++++ b/libavcodec/cuviddec.c +@@ -398,6 +398,13 @@ static int CUDAAPI cuvid_handle_video_sequence(void *opaque, CUVIDEOFORMAT* form + cuinfo.bitDepthMinus8 = format->bit_depth_luma_minus8; + cuinfo.DeinterlaceMode = ctx->deint_mode_current; + ++ if(avctx->gop_size == 100000) { ++ av_log(avctx, AV_LOG_INFO, "init decoder intra only\n"); ++ cuinfo.ulIntraDecodeOnly = 1; ++ } ++ ++ av_log(avctx, AV_LOG_INFO, "init decoder %d %d\n", ctx->nb_surfaces, avctx->refs); ++ + ctx->internal_error = CHECK_CU(ctx->cvdl->cuvidCreateDecoder(&ctx->cudecoder, &cuinfo)); + if (ctx->internal_error < 0) + return 0; diff --git a/deps/ffmpeg-patches/0005-avformat-rtp-rfc4175.patch b/deps/ffmpeg/8/0005-avformat-rtp-rfc4175.patch similarity index 65% rename from deps/ffmpeg-patches/0005-avformat-rtp-rfc4175.patch rename to deps/ffmpeg/8/0005-avformat-rtp-rfc4175.patch index 4380e8ac..8c094a34 100644 --- a/deps/ffmpeg-patches/0005-avformat-rtp-rfc4175.patch +++ b/deps/ffmpeg/8/0005-avformat-rtp-rfc4175.patch @@ -1,11 +1,8 @@ -From 1a90f8bb489004be416498b9bd3c1f76e498a675 Mon Sep 17 00:00:00 2001 -From: Jan Pietek -Date: Wed, 12 Aug 2026 15:54:13 +0000 -Subject: [PATCH 5/7] avformat/rtp: extend RFC 4175 frame handling +From 01084d45b0b91ee8d14118b44eeaa84f042c6026 Mon Sep 17 00:00:00 2001 +From: x +Date: Wed, 16 Sep 2026 19:56:02 +0200 +Subject: [PATCH] 0005-avformat-rtp-rfc4175.patch -Support the legacy 4:2:0 payload layout and skip incomplete RFC 4175 frames without emitting corrupted output. - -Original-commits: e4b4dbdccaa82b007bf3f320d7fd761f1cbd9a14 364d9773b677feedc39ed4a3e30a15e95878b811 --- libavformat/rtpdec_rfc4175.c | 88 ++++++++++++++++++++++++++++++++---- 1 file changed, 80 insertions(+), 8 deletions(-) @@ -14,53 +11,31 @@ diff --git a/libavformat/rtpdec_rfc4175.c b/libavformat/rtpdec_rfc4175.c index b49fc55..2f0dff0 100644 --- a/libavformat/rtpdec_rfc4175.c +++ b/libavformat/rtpdec_rfc4175.c -@@ -43,6 +43,9 @@ struct PayloadContext { - unsigned int frame_size; - unsigned int pgroup; /* size of the pixel group in bytes */ +@@ -45,2 +45,5 @@ struct PayloadContext { unsigned int xinc; + int is_yuv420; + int next_offset; + int next_line; - uint32_t timestamp; - }; -@@ -53,6 +56,7 @@ static int rfc4175_parse_format(AVStream *stream, PayloadContext *data) - int tag; - const AVPixFmtDescriptor *desc; +@@ -55,2 +58,3 @@ static int rfc4175_parse_format(AVStream *stream, PayloadContext *data) + data->is_yuv420 = 0; if (!strncmp(data->sampling, "YCbCr-4:2:2", 11)) { - tag = MKTAG('U', 'Y', 'V', 'Y'); - data->xinc = 2; -@@ -71,6 +75,7 @@ static int rfc4175_parse_format(AVStream *stream, PayloadContext *data) - } else if (!strncmp(data->sampling, "YCbCr-4:2:0", 11)) { - tag = MKTAG('I', '4', '2', '0'); +@@ -73,2 +77,3 @@ static int rfc4175_parse_format(AVStream *stream, PayloadContext *data) data->xinc = 4; + data->is_yuv420 = 1; - if (data->depth == 8) { - data->pgroup = 6; -@@ -108,6 +113,8 @@ static int rfc4175_parse_format(AVStream *stream, PayloadContext *data) - stream->codecpar->codec_tag = tag; - stream->codecpar->bits_per_coded_sample = av_get_bits_per_pixel(desc); +@@ -110,2 +115,4 @@ static int rfc4175_parse_format(AVStream *stream, PayloadContext *data) data->frame_size = data->width * data->height * data->pgroup / data->xinc; + data->next_offset = 0; + data->next_line = 0; - if (data->interlaced) - stream->codecpar->field_order = AV_FIELD_TT; -@@ -225,6 +232,8 @@ static int rfc4175_finalize_packet(PayloadContext *data, AVPacket *pkt, - av_freep(&data->frame); - } +@@ -227,2 +234,4 @@ static int rfc4175_finalize_packet(PayloadContext *data, AVPacket *pkt, data->frame = NULL; + data->next_offset = 0; + data->next_line = 0; } - - data->field = 0; -@@ -232,6 +241,14 @@ static int rfc4175_finalize_packet(PayloadContext *data, AVPacket *pkt, - return ret; - } +@@ -234,2 +243,10 @@ static int rfc4175_finalize_packet(PayloadContext *data, AVPacket *pkt, +static void rfc4175_discard_packet(PayloadContext *data) +{ @@ -71,11 +46,7 @@ index b49fc55..2f0dff0 100644 +} + static int rfc4175_handle_packet(AVFormatContext *ctx, PayloadContext *data, - AVStream *st, AVPacket *pkt, uint32_t *timestamp, - const uint8_t * buf, int len, -@@ -255,8 +272,12 @@ static int rfc4175_handle_packet(AVFormatContext *ctx, PayloadContext *data, - * previous frame (or pair of fields) anyway by filling the AVPacket. - */ +@@ -257,4 +274,8 @@ static int rfc4175_handle_packet(AVFormatContext *ctx, PayloadContext *data, av_log(ctx, AV_LOG_ERROR, "Missed previous RTP Marker\n"); - missed_last_packet = 1; - rfc4175_finalize_packet(data, pkt, st->index); @@ -86,11 +57,7 @@ index b49fc55..2f0dff0 100644 + rfc4175_finalize_packet(data, pkt, st->index); + } } - - if (!data->frame) -@@ -270,6 +291,11 @@ static int rfc4175_handle_packet(AVFormatContext *ctx, PayloadContext *data, - } - } +@@ -272,2 +293,7 @@ static int rfc4175_handle_packet(AVFormatContext *ctx, PayloadContext *data, + if (!data->frame) { + /* buffer already freed by discard, or duplicate packet */ @@ -98,11 +65,7 @@ index b49fc55..2f0dff0 100644 + } + /* - * looks for the 'Continuation bit' in scan lines' headers - * to find where data start -@@ -310,13 +336,59 @@ static int rfc4175_handle_packet(AVFormatContext *ctx, PayloadContext *data, - if (line >= data->height) - return AVERROR_INVALIDDATA; +@@ -312,9 +338,55 @@ static int rfc4175_handle_packet(AVFormatContext *ctx, PayloadContext *data, - /* prevent ill-formed packets to write after buffer's end */ - copy_offset = (line * data->width + offset) * data->pgroup / data->xinc; @@ -164,5 +127,6 @@ index b49fc55..2f0dff0 100644 + memcpy(dest, payload, length); + } - payload += length; - payload_len -= length; +-- +2.55.0 + diff --git a/deps/ffmpeg-patches/0006-avdevice-v4l2-compat.patch b/deps/ffmpeg/8/0006-avdevice-v4l2-compat.patch similarity index 90% rename from deps/ffmpeg-patches/0006-avdevice-v4l2-compat.patch rename to deps/ffmpeg/8/0006-avdevice-v4l2-compat.patch index e59e29ae..f8b073f8 100644 --- a/deps/ffmpeg-patches/0006-avdevice-v4l2-compat.patch +++ b/deps/ffmpeg/8/0006-avdevice-v4l2-compat.patch @@ -1,4 +1,4 @@ -From c89a55032f642c8c6afd7df0c52febe16a3c23d0 Mon Sep 17 00:00:00 2001 +From 74c71ce7e961dd4d99bbcbd14fc084e7a4d7a0c0 Mon Sep 17 00:00:00 2001 From: Jan Pietek Date: Wed, 12 Aug 2026 15:54:13 +0000 Subject: [PATCH 6/7] avdevice/v4l2: retain source timestamps @@ -11,7 +11,7 @@ Original-commit: 961ec93cbc352178ae5572c8d7f03f508edc3147 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/libavdevice/v4l2.c b/libavdevice/v4l2.c -index 0ae6872..644752f 100644 +index c38ecbb378..11203f8fec 100644 --- a/libavdevice/v4l2.c +++ b/libavdevice/v4l2.c @@ -584,7 +584,7 @@ static int mmap_read_frame(AVFormatContext *ctx, AVPacket *pkt) diff --git a/deps/ffmpeg-patches/0007-avdevice-ndi-v5.patch b/deps/ffmpeg/8/0007-avdevice-ndi-v5.patch similarity index 63% rename from deps/ffmpeg-patches/0007-avdevice-ndi-v5.patch rename to deps/ffmpeg/8/0007-avdevice-ndi-v5.patch index 6ad66c6c..8a0191f9 100644 --- a/deps/ffmpeg-patches/0007-avdevice-ndi-v5.patch +++ b/deps/ffmpeg/8/0007-avdevice-ndi-v5.patch @@ -1,11 +1,8 @@ -From 2d74ed07e99ba7a0f3a09ccf2d20cddbb9f39b71 Mon Sep 17 00:00:00 2001 -From: Jan Pietek -Date: Wed, 12 Aug 2026 15:54:13 +0000 -Subject: [PATCH 7/7] avdevice/ndi: restore NDI v5 integration +From 7fa87c315b2896edc23547c7dc867c3516524331 Mon Sep 17 00:00:00 2001 +From: x +Date: Wed, 16 Sep 2026 19:56:02 +0200 +Subject: [PATCH] 0007-avdevice-ndi-v5.patch -Register the existing NDI v5 input and output devices, add configure detection, and restore their user documentation. - -Original-commit: 6f4d2c58f2f6dd88e0612621ede61d51bc12980c --- configure | 7 +++++ doc/indevs.texi | 65 ++++++++++++++++++++++++++++++++++++++++ @@ -15,51 +12,33 @@ Original-commit: 6f4d2c58f2f6dd88e0612621ede61d51bc12980c 5 files changed, 123 insertions(+) diff --git a/configure b/configure -index 2cd7884..5ae4594 100755 +index 584e1df..f086594 100755 --- a/configure +++ b/configure -@@ -316,6 +316,7 @@ External library support: - --enable-lv2 enable LV2 audio filtering [no] - --disable-lzma disable lzma [autodetect] +@@ -325,2 +325,3 @@ External library support: --enable-decklink enable Blackmagic DeckLink I/O support [no] + --enable-libndi_newtek enable Newteck NDI I/O support [no] --enable-mbedtls enable mbedTLS, needed for https support - if openssl, gnutls or libtls is not used [no] - --enable-mediacodec enable Android MediaCodec support [no] -@@ -1878,6 +1879,7 @@ EXTERNAL_LIBRARY_GPL_LIST=" - - EXTERNAL_LIBRARY_NONFREE_LIST=" +@@ -2005,2 +2006,3 @@ EXTERNAL_LIBRARY_NONFREE_LIST=" decklink + libndi_newtek libfdk_aac - libtls - " -@@ -3726,6 +3728,10 @@ decklink_indev_suggest="libzvbi" - decklink_outdev_deps="decklink threads" - decklink_outdev_suggest="libklvanc" +@@ -3985,2 +3987,6 @@ decklink_outdev_suggest="libklvanc" decklink_outdev_extralibs="-lstdc++" +libndi_newtek_indev_deps="libndi_newtek" +libndi_newtek_indev_extralibs="-lndi" +libndi_newtek_outdev_deps="libndi_newtek" +libndi_newtek_outdev_extralibs="-lndi" dshow_indev_deps="IBaseFilter" - dshow_indev_extralibs="-lpsapi -lole32 -lstrmiids -luuid -loleaut32 -lshlwapi" - fbdev_indev_deps="linux_fb_h" -@@ -6883,6 +6889,7 @@ enabled chromaprint && { check_pkg_config chromaprint libchromaprint "chro - require chromaprint chromaprint.h chromaprint_get_version -lchromaprint; } - enabled decklink && { require_headers DeckLinkAPI.h && +@@ -7235,2 +7241,3 @@ enabled decklink && { require_headers DeckLinkAPI.h && { test_cpp_condition DeckLinkAPIVersion.h "BLACKMAGIC_DECKLINK_API_VERSION >= 0x0a0b0000" || die "ERROR: Decklink API version must be >= 10.11"; } } +enabled libndi_newtek && require libndi_newtek Processing.NDI.Lib.h NDIlib_initialize -lndi enabled frei0r && require_headers "frei0r.h" - enabled gmp && require gmp gmp.h mpz_export -lgmp - enabled gnutls && require_pkg_config gnutls gnutls gnutls/gnutls.h gnutls_global_init diff --git a/doc/indevs.texi b/doc/indevs.texi -index cdf44a6..30679b0 100644 +index 8822e07..5ef5d4c 100644 --- a/doc/indevs.texi +++ b/doc/indevs.texi -@@ -1156,6 +1156,71 @@ Set the video size given as a string such as @code{640x480} or @code{hd720}. - Default is @code{qvga}. - @end table +@@ -1123,2 +1123,67 @@ Default is @code{qvga}. +@section libndi_newtek + @@ -127,15 +106,11 @@ index cdf44a6..30679b0 100644 +@end itemize + @section openal - - The OpenAL input device provides audio capture on all systems with a diff --git a/doc/outdevs.texi b/doc/outdevs.texi -index 9ee8575..71f467d 100644 +index 86c78f3..62d342f 100644 --- a/doc/outdevs.texi +++ b/doc/outdevs.texi -@@ -301,6 +301,51 @@ ffmpeg -re -i INPUT -c:v rawvideo -pix_fmt bgra -f fbdev /dev/fb0 - - See also @url{http://linux-fbdev.sourceforge.net/}, and fbset(1). +@@ -303,2 +303,47 @@ See also @url{http://linux-fbdev.sourceforge.net/}, and fbset(1). +@section libndi_newtek + @@ -182,41 +157,30 @@ index 9ee8575..71f467d 100644 + +@end itemize + - @section opengl - OpenGL output device. Deprecated and will be removed. - + @section oss diff --git a/libavdevice/Makefile b/libavdevice/Makefile -index c304492..290ae8b 100644 +index a226368..272c80d 100644 --- a/libavdevice/Makefile +++ b/libavdevice/Makefile -@@ -22,6 +22,8 @@ OBJS-$(CONFIG_BKTR_INDEV) += bktr.o - OBJS-$(CONFIG_CACA_OUTDEV) += caca.o - OBJS-$(CONFIG_DECKLINK_OUTDEV) += decklink_enc.o decklink_enc_c.o decklink_common.o +@@ -23,2 +23,4 @@ OBJS-$(CONFIG_DECKLINK_OUTDEV) += decklink_enc.o decklink_enc_c.o deck OBJS-$(CONFIG_DECKLINK_INDEV) += decklink_dec.o decklink_dec_c.o decklink_common.o +OBJS-$(CONFIG_LIBNDI_NEWTEK_OUTDEV) += libndi_newtek_enc.o +OBJS-$(CONFIG_LIBNDI_NEWTEK_INDEV) += libndi_newtek_dec.o OBJS-$(CONFIG_DSHOW_INDEV) += dshow_crossbar.o dshow.o dshow_enummediatypes.o \ - dshow_enumpins.o dshow_filter.o \ - dshow_pin.o dshow_common.o -@@ -65,6 +67,8 @@ SHLIBOBJS-$(HAVE_GNU_WINDRES) += avdeviceres.o - SKIPHEADERS += decklink_common.h - SKIPHEADERS-$(CONFIG_DECKLINK) += decklink_enc.h decklink_dec.h \ +@@ -64,2 +66,4 @@ SKIPHEADERS-$(CONFIG_DECKLINK) += decklink_enc.h decklink_dec.h \ decklink_common_c.h +SKIPHEADERS-$(CONFIG_LIBNDI_NEWTEK_INDEV) += libndi_newtek_common.h +SKIPHEADERS-$(CONFIG_LIBNDI_NEWTEK_OUTDEV) += libndi_newtek_common.h SKIPHEADERS-$(CONFIG_DSHOW_INDEV) += dshow_capture.h - SKIPHEADERS-$(CONFIG_FBDEV_INDEV) += fbdev_common.h - SKIPHEADERS-$(CONFIG_FBDEV_OUTDEV) += fbdev_common.h diff --git a/libavdevice/alldevices.c b/libavdevice/alldevices.c -index 9b9a914..ac942ec 100644 +index 573595f..3b516e6 100644 --- a/libavdevice/alldevices.c +++ b/libavdevice/alldevices.c -@@ -36,6 +36,8 @@ extern const FFInputFormat ff_bktr_demuxer; - extern const FFOutputFormat ff_caca_muxer; - extern const FFInputFormat ff_decklink_demuxer; +@@ -37,2 +37,4 @@ extern const FFInputFormat ff_decklink_demuxer; extern const FFOutputFormat ff_decklink_muxer; +extern const FFInputFormat ff_libndi_newtek_demuxer; +extern const FFOutputFormat ff_libndi_newtek_muxer; extern const FFInputFormat ff_dshow_demuxer; - extern const FFInputFormat ff_fbdev_demuxer; - extern const FFOutputFormat ff_fbdev_muxer; +-- +2.55.0 + diff --git a/deps/ffmpeg/8/bases.env b/deps/ffmpeg/8/bases.env new file mode 100644 index 00000000..ea60d972 --- /dev/null +++ b/deps/ffmpeg/8/bases.env @@ -0,0 +1,5 @@ +n80_commit=140fd653aed8cad774f991ba083e2d01e86420c7 +n80_tree=6ff5e4f32a3acb868c6fd4a2ae9e16c25115a6f2 +n81_commit=9047fa1b084f76b1b4d065af2d743df1b40dfb56 +n81_tree=3f4f8ba98c493fe9c0e0bfae5a61cd8b83fc30ca +patch_count=7 diff --git a/deps/ffmpeg/README.md b/deps/ffmpeg/README.md new file mode 100644 index 00000000..9ec7429f --- /dev/null +++ b/deps/ffmpeg/README.md @@ -0,0 +1,67 @@ +# FFmpeg patch series (FFmpeg 8.x) + +One ordered series in `8/` applies to upstream **n8.0 and n8.1** from a single +copy; there are no per-version directories. `8/bases.env` pins each base +commit and the exact patched tree; `patch_count` guards against stray files. + +```bash +deps/ffmpeg/apply.sh # git am 8/*.patch + configure fix-up +deps/ffmpeg/verify.sh # isolated worktree, whole series, tree check +``` + +The demo Dockerfiles (`demos/mixer`, `demos/cuda-overlay`, +`demos/dmabuf-browser/consumer`) call `apply.sh`; `FFMPEG_TAG` defaults to +`n8.1` and may be set to `n8.0`. avcpp and avplumber must be rebuilt against +the selected FFmpeg libraries. + +## How one series serves both bases + +The composition-suite filters are new files, so they carry no base-tree +context; the remaining patches are generated with minimal (`-U1`) context so +they apply where 8.0 and 8.1 differ only cosmetically. The single genuine +8.0/8.1 difference is a libnpp `configure` check that 8.1 added and that fails +on CUDA 13 (the legacy `nppiYCbCr420_8u_P2P3R` symbol is gone). `apply.sh` +rewrites that check to probe the stream-context API; on 8.0, which has no such +check, the rewrite is a no-op. Everything else in the NPP patch (`npp_compat.h` +and the filter changes) is base-independent. + +## The series + +1. **swscale aarch64 ARGB→YUVA420P** — uses `SwsInternal`/`opts` and the + unscaled callback signature; algorithms unchanged. +2. **CUDA composition suite** — `convert_cuda`, `crop_cuda`, `overlay_many_cuda`, + `pad_cuda` (intentionally replaces upstream's), `transition_cuda`, plus the + scaling-edge/overlay-context fixes and YUVA444P CUDA frames. Upstream 8.x + `scale_cuda` (with its expanded pixel formats) is kept. +3. **NPP CUDA 13 compatibility** — stream-context helper and filter changes. +4. **NVDEC intra-only handling.** +5. **RFC 4175 RTP frame handling.** +6. **V4L2 source timestamps.** +7. **NDI v5 registration** — registers the optional integration only; no SDK + or device implementation is supplied, and NDI stays disabled in demo builds. + +## FFmpeg 8 notes + +- `AVFrame.pkt_pos` is gone: `transition_cuda`'s legacy `pos` expression + variable evaluates to `NAN`; `crop_cuda` guards it by API version. +- The `C` command-support marker left `-filters` output; the mixer Dockerfile + checks the transition's `mode` option in filter help instead. +- FFmpeg 8 validates CUDA input formats at filter init, so the AVP filter node + attaches `hw_frames_ctx` / `hw_device_ctx` between allocation and + initialisation (segmented graph parser); the metadata-driven crop node does + the same. +- When reusing a build tree with different GL flags, rebuild + `deps/cuda_loader/cuda_drvapi_dynlink.o`; a loader built without GL lacks the + EGL function-pointer variables, and linking `libcuda` directly to paper over + that crashes during DMA-BUF import. + +## Validation + +Compilation, linking and filter registration are the acceptance criteria for +the series itself; runtime behaviour is covered by the demo and `tests/cuda` +suites. Validated on a T4 with FFmpeg 8.1: the 8-bit mixer at 30/60 fps with +video and DMA-BUF browser inputs, wipe-cache loads, cut/fade/wipe, NVENC to a +WebRTC browser; the `demos/cuda-overlay` 45-case pixel-reference matrix; and +the live recorder surviving SRT disconnects (`ignore_eof` on the pre-sentinel +format nodes, opt-in). NPP, NDI and AArch64 paths are not compiled in the demo +images. diff --git a/deps/ffmpeg/apply.sh b/deps/ffmpeg/apply.sh new file mode 100755 index 00000000..0959298a --- /dev/null +++ b/deps/ffmpeg/apply.sh @@ -0,0 +1,19 @@ +#!/usr/bin/env bash +# Apply the FFmpeg 8.x patch series to a checkout of upstream n8.0 or n8.1. +set -euo pipefail +[[ $# -eq 1 ]] || { echo "usage: $0 /path/to/FFmpeg" >&2; exit 2; } +dir=$(CDPATH= cd -- "$(dirname -- "$0")" && pwd) +src=$1 +git -C "$src" -c user.name="avplumber patches" -c user.email="patches@local" \ + am --whitespace=nowarn "$dir"/8/*.patch +# FFmpeg 8.1 added a libnpp configure check that fails on CUDA 13 (the legacy +# nppiYCbCr420_8u_P2P3R symbol is gone); probe the stream-context API instead. +# 8.0 has no such check, so this is a no-op there. +sed -i.bak \ + -e 's/check_func_headers "nppi.h" nppiYCbCr420_8u_P2P3R \$libnpp_extralibs/check_func_headers "nppi.h" nppiYCbCr420_8u_P2P3R_Ctx $libnpp_extralibs/' \ + -e 's/libnpp support is deprecated, version 13.0 and up are not supported/libnpp stream context APIs not found/' \ + "$src/configure" && rm -f "$src/configure.bak" +if ! git -C "$src" diff --quiet -- configure; then + git -C "$src" -c user.name="avplumber patches" -c user.email="patches@local" \ + commit -qam "configure: probe the libnpp stream-context API (CUDA 13)" +fi diff --git a/deps/ffmpeg/verify.sh b/deps/ffmpeg/verify.sh new file mode 100755 index 00000000..252b1d01 --- /dev/null +++ b/deps/ffmpeg/verify.sh @@ -0,0 +1,19 @@ +#!/usr/bin/env bash +# Verify the series reproduces the pinned tree for an upstream base (n8.0 or n8.1). +set -euo pipefail +[[ $# -eq 2 ]] || { echo "usage: $0 /path/to/FFmpeg" >&2; exit 2; } +dir=$(CDPATH= cd -- "$(dirname -- "$0")" && pwd) +source "$dir/8/bases.env" +case "$1" in n8.0) k=n80 ;; n8.1) k=n81 ;; *) echo "unsupported base: $1" >&2; exit 2 ;; esac +v="${k}_commit"; base_commit=${!v}; v="${k}_tree"; expected_tree=${!v} +repo=$(git -C "$2" rev-parse --show-toplevel) +git -C "$repo" cat-file -e "${base_commit}^{commit}" || { echo "checkout lacks $1 ($base_commit)" >&2; exit 2; } +shopt -s nullglob; patches=("$dir"/8/*.patch) +[[ ${#patches[@]} -eq $patch_count ]] || { echo "expected $patch_count patches, found ${#patches[@]}" >&2; exit 1; } +root=$(mktemp -d "${TMPDIR:-/tmp}/avplumber-ffmpeg-verify.XXXXXX"); wt="$root/ffmpeg" +trap 'git -C "$repo" worktree remove --force "$wt" >/dev/null 2>&1 || true; rmdir "$root" 2>/dev/null || true' EXIT +git -C "$repo" worktree add --detach "$wt" "$base_commit" >/dev/null +"$dir/apply.sh" "$wt" >/dev/null +actual=$(git -C "$wt" rev-parse 'HEAD^{tree}') +[[ "$actual" == "$expected_tree" ]] || { echo "unexpected patched tree for $1: $actual (expected $expected_tree)" >&2; exit 1; } +echo "FFmpeg $1 patch series verified: $actual" diff --git a/doc/NODES.md b/doc/NODES.md index d90e1e9f..6bfc408e 100644 --- a/doc/NODES.md +++ b/doc/NODES.md @@ -284,6 +284,13 @@ when real metadata aren't available yet. 1 input, 1 output: `av::VideoFrame` or `av::AudioSamples` +Shared parameter (also supported by `fake_video_format` / `fake_audio_metadata`): + +- `ignore_eof` (bool, default `false`) - discard upstream EOF markers and keep + processing subsequent frames. Enable only at a live reconnect boundary before + a sentinel. The default forwards EOF and finishes; explicit shutdown and errors + are unaffected. + Parameters for video: - `width` (int) - default 1920 - `height` (int) - default 1080 @@ -425,11 +432,24 @@ encoder option in FFmpeg, works with non-integer FPS. - `interval_sec` (int / float / string of rational) - optional, keyframe interval, in seconds. If omitted, periodic forcing is disabled and the node acts as a pass-through until triggered externally. +- `min_interval_ms` (non-negative int, default `0`) - minimum media-PTS + spacing between forced keyframes, shared by periodic and external requests. + `0` preserves unrestricted forcing. Requires valid PTS when enabled; a + backwards jump before the last forced PTS resets the limiter. The Janus + mixer output defaults to `150` (9 frames at 60 fps), configurable with + `--keyframe-min-interval-ms` or `JanusVideoConfig.keyframe_min_interval_ms`. + +While rate-limited, ordinary frames pass through without waiting. Requests +coalesce and remain pending until the first eligible frame, including requests +from RTCP PLI/FIR; these may wait up to the minimum interval plus frame rounding. +Periodic forcing cannot bypass the limit. This limits the node's forced frames, +not keyframes independently inserted by an encoder, and is not a wall-clock RTP +packet pacer. Runtime control (via `node.object.set` / `node.object.get`): - `node.object.set trigger true` — request one keyframe on the next - processed frame (one-shot, edge-triggered). Also accepted as key `force` + eligible frame (one-shot, edge-triggered). Also accepted as key `force` or `request`. - `node.object.get status` — returns a JSON object: ```json @@ -439,7 +459,8 @@ Runtime control (via `node.object.set` / `node.object.get`): "pending": false, "triggered_frames": 2, "periodic_frames": 60, - "interval_enabled": true + "interval_enabled": true, + "min_interval_ms": 150 } ``` - `node.object.get pending` — bool, `true` if a trigger has been diff --git a/src/nodes/assume_metadata.cpp b/src/nodes/assume_metadata.cpp index be3bf67f..b9e55e2b 100644 --- a/src/nodes/assume_metadata.cpp +++ b/src/nodes/assume_metadata.cpp @@ -1,6 +1,24 @@ #include "node_common.hpp" -class AssumeAudioFormat: public TransparentNode, public IAudioMetadataSource, public ITimeBaseSource { +// Live format declarations must remain available while upstream reconnects. +// Keep ordinary EOF propagation as the default for finite graphs. +template class FormatDeclaration: public TransparentNode { +public: + using TransparentNode::TransparentNode; + bool ignore_eof = false; + + bool consumeEofIfPresent() override { + if (!ignore_eof) return TransparentNode::consumeEofIfPresent(); + T* frame = this->source_->peek(0); + if (frame && isEofMarker(*frame)) { + this->source_->pop(); + logstream << "Ignoring upstream EOF at live format boundary"; + } + return false; + } +}; + +class AssumeAudioFormat: public FormatDeclaration, public IAudioMetadataSource, public ITimeBaseSource { private: int sample_rate_; av::SampleFormat sample_format_; @@ -19,7 +37,7 @@ class AssumeAudioFormat: public TransparentNode, public IAudio return {1, sample_rate_}; } AssumeAudioFormat(std::unique_ptr> &&source, std::unique_ptr> &&sink, const int sample_rate, const av::SampleFormat sample_format, const uint64_t channel_layout): - TransparentNode(std::move(source), std::move(sink)), + FormatDeclaration(std::move(source), std::move(sink)), sample_rate_(sample_rate), sample_format_(sample_format), channel_layout_(channel_layout) { } static std::shared_ptr create(NodeCreationInfo &nci) { @@ -31,13 +49,15 @@ class AssumeAudioFormat: public TransparentNode, public IAudio if (params.count("sample_rate")) sr = params["sample_rate"]; if (params.count("sample_format")) fmt = av::SampleFormat(params["sample_format"].get()); if (params.count("channel_layout")) chl = stringToChannelLayout(params["channel_layout"].get().c_str()); - return NodeSISO::template createCommon(edges, params, sr, fmt, chl); + auto node = NodeSISO::template createCommon(edges, params, sr, fmt, chl); + node->ignore_eof = params.value("ignore_eof", false); + return node; } virtual ~AssumeAudioFormat() { } }; -class AssumeVideoFormat: public TransparentNode, public IVideoFormatSource { +class AssumeVideoFormat: public FormatDeclaration, public IVideoFormatSource { private: int width_, height_; av::PixelFormat pix_fmt_; @@ -57,7 +77,7 @@ class AssumeVideoFormat: public TransparentNode, public IVideoFo } AssumeVideoFormat(std::unique_ptr> &&source, std::unique_ptr> &&sink, const int width, const int height, const av::PixelFormat pixfmt, const av::PixelFormat real_pixfmt): - TransparentNode(std::move(source), std::move(sink)), + FormatDeclaration(std::move(source), std::move(sink)), width_(width), height_(height), pix_fmt_(pixfmt), real_pix_fmt_(real_pixfmt) { } static std::shared_ptr create(NodeCreationInfo &nci) { @@ -71,7 +91,9 @@ class AssumeVideoFormat: public TransparentNode, public IVideoFo if (params.count("height")) height = params["height"]; if (params.count("pixel_format")) pixel_format = av::PixelFormat(params["pixel_format"].get()); if (params.count("real_pixel_format")) real_pixel_format = av::PixelFormat(params["real_pixel_format"].get()); - return NodeSISO::template createCommon(edges, params, width, height, pixel_format, real_pixel_format); + auto node = NodeSISO::template createCommon(edges, params, width, height, pixel_format, real_pixel_format); + node->ignore_eof = params.value("ignore_eof", false); + return node; } }; diff --git a/src/nodes/crop_metadata_cuda.cpp b/src/nodes/crop_metadata_cuda.cpp index f52ccfd9..fb49f566 100644 --- a/src/nodes/crop_metadata_cuda.cpp +++ b/src/nodes/crop_metadata_cuda.cpp @@ -152,9 +152,9 @@ class MetadataDrivenCudaCrop: public NodeSISO, } const std::string in_args = buildSourceArgsString(); - int ret = avfilter_graph_create_filter(&buffersrc_ctx_, buffersrc, "in", in_args.c_str(), nullptr, filter_graph_); - if (ret < 0) { - throw Error("crop_metadata_cuda: couldn't create buffer source"); + buffersrc_ctx_ = avfilter_graph_alloc_filter(filter_graph_, buffersrc, "in"); + if (!buffersrc_ctx_) { + throw Error("crop_metadata_cuda: couldn't allocate buffer source"); } AVBufferSrcParameters *src_params = av_buffersrc_parameters_alloc(); @@ -162,13 +162,24 @@ class MetadataDrivenCudaCrop: public NodeSISO, throw Error("crop_metadata_cuda: av_buffersrc_parameters_alloc failed"); } src_params->hw_frames_ctx = av_buffer_ref(initial_hw_frames_ctx_); - ret = av_buffersrc_parameters_set(buffersrc_ctx_, src_params); + if (!src_params->hw_frames_ctx) { + av_freep(&src_params); + throw Error("crop_metadata_cuda: couldn't reference buffer source CUDA frames"); + } + int ret = av_buffersrc_parameters_set(buffersrc_ctx_, src_params); av_buffer_unref(&src_params->hw_frames_ctx); av_freep(&src_params); if (ret < 0) { throw Error("crop_metadata_cuda: av_buffersrc_parameters_set failed"); } + // FFmpeg 8.1 validates CUDA sources during init, so the frames context + // must be attached before initializing, not just before graph config. + ret = avfilter_init_str(buffersrc_ctx_, in_args.c_str()); + if (ret < 0) { + throw Error("crop_metadata_cuda: couldn't initialize buffer source: " + av::error2string(ret)); + } + std::stringstream crop_args; crop_args << "w=" << dst_width_ << ":h=" << dst_height_ << ":x=" << last_crop_x_ << ":y=" << last_crop_y_; diff --git a/src/nodes/filters.cpp b/src/nodes/filters.cpp index d3a20611..1a231a20 100644 --- a/src/nodes/filters.cpp +++ b/src/nodes/filters.cpp @@ -191,16 +191,13 @@ template class FilterNode: p logstream << "Unable to init source filter " << name << ": in_args_ not initialized"; } - // create buffersrc filter const AVFilter* buffersrc = avfilter_get_by_name(ms_.source_filter_name); - int ret = avfilter_graph_create_filter(&ctx_, buffersrc, name.c_str(), in_args_.c_str(), nullptr, filter_graph); - if (ret < 0) { - throw Error("Couldn't create buffer source"); + if (!buffersrc) { + throw Error("Couldn't find buffer source filter"); } - - ret = avfilter_link(ctx_, 0, dst->filter_ctx, dst->pad_idx); - if (ret != 0) { - throw Error("Couldn't link " + name); + ctx_ = avfilter_graph_alloc_filter(filter_graph, buffersrc, name.c_str()); + if (!ctx_) { + throw Error("Couldn't allocate buffer source"); } // Prefer copying hw_frames_ctx from the first frame (if captured), @@ -238,6 +235,17 @@ template class FilterNode: p av_buffer_unref(¶ms->hw_frames_ctx); av_freep(¶ms); } + + // FFmpeg 8.1 validates hardware inputs during initialization, so the + // buffersrc must already have its hw_frames_ctx at this point. + int ret = avfilter_init_str(ctx_, in_args_.c_str()); + if (ret < 0) { + throw Error("Couldn't initialize buffer source: " + av::error2string(ret)); + } + ret = avfilter_link(ctx_, 0, dst->filter_ctx, dst->pad_idx); + if (ret != 0) { + throw Error("Couldn't link " + name); + } } void initSinkFilter(const int index, AVFilterGraph *filter_graph, AVFilterInOut *src) { std::string name = "out" + std::to_string(index); @@ -352,6 +360,35 @@ template class FilterNode: p sinks_.resize(this->sink_edges_.size()); input_eof_.resize(this->source_edges_.size(), false); } + int parseFilterGraph(AVFilterInOut **inputs, AVFilterInOut **outputs) { + #ifdef AVFILTER_FLAG_HWDEVICE + if (hwaccel_) { + AVFilterGraphSegment *raw_segment = nullptr; + int ret = avfilter_graph_segment_parse(filter_graph_, graph_desc_.c_str(), + 0, &raw_segment); + auto free_segment = [](AVFilterGraphSegment *segment) { + avfilter_graph_segment_free(&segment); + }; + std::unique_ptr + segment(raw_segment, free_segment); + if (ret < 0) return ret; + ret = avfilter_graph_segment_create_filters(segment.get(), 0); + if (ret < 0) return ret; + + // FFmpeg 8.1's hwupload requires the device during init, not just + // format negotiation. Attach it before segment_apply initializes. + for (unsigned j = 0; j < filter_graph_->nb_filters; j++) { + AVFilterContext *fctx = filter_graph_->filters[j]; + if (fctx->filter->flags & AVFILTER_FLAG_HWDEVICE) { + fctx->hw_device_ctx = hwaccel_->refDeviceContext(); + if (!fctx->hw_device_ctx) return AVERROR(ENOMEM); + } + } + return avfilter_graph_segment_apply(segment.get(), 0, inputs, outputs); + } + #endif + return avfilter_graph_parse2(filter_graph_, graph_desc_.c_str(), inputs, outputs); + } bool maybeInitFilterGraph(bool frame_waiting = false) { if (filter_graph_ != nullptr) { freeFilterGraph(); @@ -384,27 +421,12 @@ template class FilterNode: p AVFilterInOut* outputs = nullptr; int ret; - ret = avfilter_graph_parse2(filter_graph_, graph_desc_.c_str(), &inputs, &outputs); + ret = parseFilterGraph(&inputs, &outputs); if (ret < 0) { log_init("error_parse", elapsed_ms()); throw Error("Couldn't parse filter graph"); } - #ifdef AVFILTER_FLAG_HWDEVICE - // Provide the hardware device to any filter in the graph that requests it - // (e.g. hwupload_cuda, hwupload, scale_cuda). This lets those filters - // allocate output HW frames even when the buffersrc is a software format. - // ffmpeg 6.1+ is required for this to work. - if (hwaccel_) { - for (unsigned j = 0; j < filter_graph_->nb_filters; j++) { - AVFilterContext *fctx = filter_graph_->filters[j]; - if (fctx->filter->flags & AVFILTER_FLAG_HWDEVICE) { - fctx->hw_device_ctx = hwaccel_->refDeviceContext(); - } - } - } - #endif - auto forEachInOut = [](AVFilterInOut *inout, std::function cb) { for (; inout != nullptr; inout = inout->next) { cb(inout); diff --git a/src/nodes/force_keyframe.cpp b/src/nodes/force_keyframe.cpp index 97089ca5..f8f725e8 100644 --- a/src/nodes/force_keyframe.cpp +++ b/src/nodes/force_keyframe.cpp @@ -1,5 +1,6 @@ #include "node_common.hpp" #include +#include class ForceKeyFrame: public NodeSISO, public IInputsObjects, @@ -8,6 +9,8 @@ class ForceKeyFrame: public NodeSISO, av::Rational interval_sec_; bool interval_enabled_ = false; int64_t last_result_ = -(1L<<62); + int min_interval_ms_ = 0; + av::Timestamp last_forced_pts_; std::atomic requested_generation_{0}; std::atomic forced_generation_{0}; std::atomic triggered_frames_{0}; @@ -45,11 +48,11 @@ class ForceKeyFrame: public NodeSISO, return true; } - // Coalescing semantics: any number of triggers received between two frames - // results in *one* forced keyframe. forced_generation_ is set to the latest - // requested value, so callers polling getObject("pending") cannot distinguish + // Coalescing semantics: any number of triggers received before the next + // eligible frame results in *one* forced keyframe. forced_generation_ is set + // to the latest requested value, so callers polling getObject("pending") cannot distinguish // their individual trigger from triggers emitted by other clients in the - // same inter-frame window. This is intentional — emitting one keyframe per + // same rate-limit window. This is intentional — emitting one keyframe per // received trigger would let a misbehaving controller spike the bitrate // arbitrarily. bool shouldForceTriggered() { @@ -73,13 +76,32 @@ class ForceKeyFrame: public NodeSISO, return; } av::VideoFrame frm = *ptr; - const bool force_triggered = shouldForceTriggered(); - const bool force_periodic = shouldForcePeriodic(frm); + const auto pts = frm.pts(); + if (min_interval_ms_ > 0) { + if (!pts.isValid() || pts.timebase().getNumerator() <= 0 || + pts.timebase().getDenominator() <= 0) { + throw Error("force_keyframe: min_interval_ms requires valid frame PTS"); + } + // A seek/reconnect starts a new media timeline. Do not wait for the + // old timeline to catch up before allowing a recovery keyframe. + if (last_forced_pts_.isValid() && pts < last_forced_pts_) { + last_forced_pts_ = av::Timestamp(); + last_result_ = -(1L<<62); + } + } + const bool eligible = min_interval_ms_ == 0 || !last_forced_pts_.isValid() || + pts - last_forced_pts_ >= av::Timestamp(min_interval_ms_, av::Rational(1, 1000)); + // Do not acknowledge triggers or advance the periodic bucket while + // throttled: both remain due until one eligible frame satisfies them. + const bool force_triggered = eligible && shouldForceTriggered(); + const bool force_periodic = eligible && shouldForcePeriodic(frm); if (force_triggered || force_periodic) { frm.setPictureType(AV_PICTURE_TYPE_I); frm.setKeyFrame(true); + last_forced_pts_ = pts; } else { frm.setPictureType(AV_PICTURE_TYPE_NONE); + frm.setKeyFrame(false); } this->sink_->put(frm); this->source_->pop(); @@ -113,6 +135,7 @@ class ForceKeyFrame: public NodeSISO, status["triggered_frames"] = triggered_frames_.load(std::memory_order_relaxed); status["periodic_frames"] = periodic_frames_.load(std::memory_order_relaxed); status["interval_enabled"] = interval_enabled_; + status["min_interval_ms"] = min_interval_ms_; return status; } if (key == "pending") { @@ -128,10 +151,12 @@ class ForceKeyFrame: public NodeSISO, std::unique_ptr> &&source, std::unique_ptr> &&sink, bool interval_enabled, - const av::Rational interval_sec + const av::Rational interval_sec, + int min_interval_ms ): NodeSISO(std::move(source), std::move(sink)), interval_sec_(interval_sec), - interval_enabled_(interval_enabled) { + interval_enabled_(interval_enabled), + min_interval_ms_(min_interval_ms) { } static std::shared_ptr create(NodeCreationInfo &nci) { @@ -139,6 +164,15 @@ class ForceKeyFrame: public NodeSISO, const Parameters ¶ms = nci.params; bool interval_enabled = false; av::Rational interval_sec(0, 1); + int min_interval_ms = 0; + if (params.contains("min_interval_ms")) { + const auto& value = params["min_interval_ms"]; + if (!value.is_number_integer() || value.get() < 0 || + value.get() > std::numeric_limits::max()) { + throw Error("min_interval_ms must be a non-negative integer in milliseconds"); + } + min_interval_ms = value.get(); + } if (params.count("interval_sec") > 0) { interval_sec = parseInterval(params["interval_sec"]); if (interval_sec.getNumerator() <= 0 || interval_sec.getDenominator() <= 0) { @@ -150,9 +184,10 @@ class ForceKeyFrame: public NodeSISO, edges, params, interval_enabled, - interval_sec + interval_sec, + min_interval_ms ); } }; -DECLNODE(force_keyframe, ForceKeyFrame); +DECLNODE(force_keyframe, ForceKeyFrame) diff --git a/tests/cuda/smoke_crop_filter_chain.py b/tests/cuda/smoke_crop_filter_chain.py new file mode 100644 index 00000000..9876b2d5 --- /dev/null +++ b/tests/cuda/smoke_crop_filter_chain.py @@ -0,0 +1,122 @@ +"""Bounded NVDEC -> metadata crop -> pad/scale/format conversion -> NVENC smoke. + +Run on an isolated NVIDIA host with the patched FFmpeg libraries and a video +fixture at least 640x360. Frames stay on CUDA throughout the processing graph; +the independent software decode below only validates the encoded output. +""" +import argparse +import json +from pathlib import Path +import subprocess +import time + +from pyplumber import AVPlumber +from pyplumber.node import ( + AssumeVideoFormat, CropMetadataCuda, DecVideo, Demux, EncVideo, + FilterVideo, ForceFPS, InputRec, +) + + +def run(args): + args.output.parent.mkdir(parents=True, exist_ok=True) + fixture = args.output.with_suffix(".input.mkv") + # Stream-copy a bounded encoded fixture; no pixel conversion is involved. + subprocess.run([args.ffmpeg, "-v", "error", "-y", "-i", args.input, + "-map", "0:v:0", "-c", "copy", "-frames:v", str(args.frames + 16), + str(fixture)], check=True, timeout=30) + avp = AVPlumber() + errors = [] + avp.on_exception = lambda *error: errors.append(tuple(map(str, error))) + avp.executeCommandsFromString('hwaccel.init {"name":"filter_chain_gpu","type":"cuda"}') + avp.edges.planCapacity("*", 4) + graph = ( + f"pad_cuda=672:384:16:12:color=black,{args.scaler}=640:360," + "convert_cuda=format=yuv420p,convert_cuda=format=nv12" + ) + if args.band_blur: + graph += ",band_blur_cuda" + nodes = [ + InputRec({"name": "input", "url": str(fixture), "loop": False, "dst": "packets"}), + Demux({"name": "demux", "src": "packets", "routing": {"v:0": "video_packets"}}), + DecVideo({"name": "decode", "src": "video_packets", "dst": "decoded", + "hwaccel": "filter_chain_gpu", "pixel_format": "cuda"}), + # Empty metadata selects the crop node's documented center fallback. + # The FFmpeg metadata filter edits properties, not CUDA pixel data. + FilterVideo({"name": "metadata", "src": "decoded", "dst": "marked", + "graph": "metadata=mode=add:key=reframer_bbox:value='{}'"}), + CropMetadataCuda({"name": "crop", "src": "marked", "dst": "cropped", + "dst_width": 640, "dst_height": 360, + "offset_log_path": str(args.output.with_suffix(".crop.log"))}), + FilterVideo({"name": "chain", "src": "cropped", "dst": "filtered", "graph": graph, + "hwaccel": "filter_chain_gpu", "dst_width": 640, "dst_height": 360, + "dst_pixel_format": "cuda", "dst_frame_rate": "25/1", + "defer_preliminary_init": True}), + # Match the recorder's static timebase boundary before NVENC. The + # deferred CUDA filter has no output timebase until its first frame. + ForceFPS({"name": "fps", "src": "filtered", "dst": "paced", "fps": "25/1"}), + AssumeVideoFormat({"name": "format", "src": "paced", "dst": "nv12", + "width": 640, "height": 360, "pixel_format": "cuda", + "real_pixel_format": "nv12"}), + EncVideo({"name": "encode", "src": "nv12", "dst": "encoded", + "hwaccel": "filter_chain_gpu", "codec": "h264_nvenc", + "options": {"preset": "p1", "tune": "ull", "bf": 0, + "rc-lookahead": 0, "delay": 0, "g": 25}}), + ] + for node in nodes: + node.parameters.update({"group": "test", "auto_restart": "off"}) + avp.addNode(node) + del node + encoded = avp.getEdge("encoded", "Packet") + packets = [] + try: + avp.group("test").startNodes() + deadline = time.monotonic() + args.timeout + finished = False + while time.monotonic() < deadline and not errors: + packet = encoded.tryGet(100) + if packet is None: + continue + # AVPlumber's packet EOF marker has a one-byte payload and NOPTS. + if packet.pts.timestamp == -(1 << 63): + finished = True + break + if packet.size > 0: + if len(packets) < args.frames: + packets.append(packet.data) + assert not errors, errors + assert finished, "encoded stream did not reach EOF" + assert len(packets) == args.frames, f"only {len(packets)}/{args.frames} encoded frames" + # The encoder can enqueue EOF before its worker (or an upstream + # decoder) has unwound. Await those EOF-driven workers before shutdown; + # they deliberately have no interface for an abrupt mid-frame stop. + eof_workers = [node.parameters["name"] for node in nodes + if node.parameters["name"] != "fps"] + while any(avp.node(name).isWorking for name in eof_workers): + assert time.monotonic() < deadline, "EOF workers did not finish" + time.sleep(0.01) + finally: + # The crop node is EOF-driven, so use a finite fixture and drain the + # stream instead of forcing a stop midway through its filter graph. + nodes.clear() + avp.shutdown() + args.output.write_bytes(b"".join(packets)) + decoded = subprocess.run([ + args.ffmpeg, "-v", "error", "-threads", "1", "-f", "h264", "-i", str(args.output), + "-vf", "scale=2:2", "-pix_fmt", "gray", "-fps_mode", "passthrough", + "-f", "rawvideo", "pipe:1", + ], check=True, capture_output=True, timeout=30).stdout + assert len(decoded) == args.frames * 4, "independent decode frame count differs" + print(json.dumps({"status": "passed", "frames": args.frames, "scaler": args.scaler, + "graph": graph, "output": str(args.output)}), flush=True) + + +if __name__ == "__main__": + parser = argparse.ArgumentParser(description=__doc__) + parser.add_argument("--input", required=True) + parser.add_argument("--output", required=True, type=Path) + parser.add_argument("--frames", type=int, default=60) + parser.add_argument("--timeout", type=float, default=30) + parser.add_argument("--scaler", choices=("scale_cuda", "scale_npp"), default="scale_cuda") + parser.add_argument("--band-blur", action="store_true", help="also check the reframer's optional band_blur_cuda patch") + parser.add_argument("--ffmpeg", default="/usr/local/bin/ffmpeg") + run(parser.parse_args()) diff --git a/tests/cuda/smoke_keyframe_limit.py b/tests/cuda/smoke_keyframe_limit.py new file mode 100644 index 00000000..a6abfe7d --- /dev/null +++ b/tests/cuda/smoke_keyframe_limit.py @@ -0,0 +1,148 @@ +"""Isolated 16-input, 60-fps mixer/NVENC keyframe-spam regression. + +Uses the demo's actual Janus encoder graph but terminates at an encoded-packet +reader, never a live Janus mountpoint. CPU decoding is only the post-encode +pixel oracle. No CPU/GPU round trip is added to the mixer processing graph. +Run separate processes with --minimum-ms 0 and 150 for the A/B comparison. +""" + +import argparse +import json +from pathlib import Path +import statistics +import subprocess +import sys +import threading +import time + +sys.path.insert(0, str(Path(__file__).resolve().parents[2] / "demos" / "mixer")) +from mixer import GraphOptions, build_application + + +def run(args): + stop = threading.Event() + app = build_application(GraphOptions( + inputs=tuple(args.inputs[i % 2] for i in range(args.source_count)), + fps=60, loop_inputs=True, janus_output=True, + janus_video_bitrate_kbps=2700, keyframe_min_interval_ms=args.minimum_ms, + remote_control_port=args.port, wipe_cache_mb=0, + prewarm_cut_scenes=("*",), + )) + # The owned test ends at the encoder; leave RTP/BSF lifecycle to its own + # tests, and never publish this graph to the demo's streaming mountpoint. + app.avp.executeCommandsFromString("node.delete janus_rtp_output\n" + "node.delete janus_mux\nnode.delete janus_repeat_headers") + app.rtcp_feedback_listener = None + errors, frames, cuts = [], [], [] + app.avp.on_exception = lambda *error: errors.append(error) + + def capture(packet): + if packet.size <= 0 or packet.pts.timestamp == -(1 << 63): + return + tb = packet.pts.timebase + frames.append({"at": time.monotonic(), "pts": round(packet.pts.timestamp * tb.num / tb.den * 60), + "key": bool(packet.flags & 1), "data": packet.data}) + + # A queue reader avoids Python callbacks from a native encoder thread + # during shutdown (shutdown holds the GIL while joining native workers). + encoded = app.avp.getEdge("janus_encoded", "Packet") + + def read_encoded(): + while not stop.is_set(): + packet = encoded.tryGet(100) + if packet is not None: + capture(packet) + + packet_reader = threading.Thread(target=read_encoded, daemon=True) + packet_reader.start() + + def wait_for(predicate, label, timeout=10): + deadline = time.monotonic() + timeout + while not predicate(): + assert not errors, errors + assert time.monotonic() < deadline, label + time.sleep(.005) + + try: + app.start() + app.avp.registerWithWebUI(args.webui, "keyframe-limit-smoke", "") + wait_for(lambda: len(frames) >= 60, "encoder startup") + state = app.avp.node("janus_force_keyframe").getObject("status") + assert state["min_interval_ms"] == args.minimum_ms, state + begin = time.monotonic() + for index in range(args.cuts): + target = 1 - index % 2 + cuts.append({"at": time.monotonic(), "target": target}) + app.mixer.cut(f"fullscreen_{target}") + time.sleep(.05) + # Let the final cut/keyframe request finish after spam stops. + time.sleep(.4) + assert not app.avp.node("janus_force_keyframe").getObject("pending") + cut_end = time.monotonic() + # Exercise interrupted crossfades, then a fully completed fade. + for index in range(12): + app.mixer.fade(f"fullscreen_{index % 2}", duration_sec=.25) + time.sleep(.075) + app.mixer.fade("fullscreen_0", duration_sec=.25) + time.sleep(.6) + end = time.monotonic() + state = app.avp.node("janus_force_keyframe").getObject("status") + assert not state["pending"], state + assert not errors, errors + captured = tuple(frames) + measured = [f for f in captured if begin <= f["at"] <= end] + pts = [f["pts"] for f in measured] + assert all(b - a == 1 for a, b in zip(pts, pts[1:])), "output frame gap or duplicate PTS" + keys = [f["pts"] for f in captured if f["key"]] + key_spacing = [b - a for a, b in zip(keys, keys[1:])] + assert len(key_spacing) >= 3, "insufficient keyframes" + assert all(g * 1000 >= args.minimum_ms * 60 for g in key_spacing), key_spacing + gaps = [(b["at"] - a["at"]) * 1000 for a, b in zip(measured, measured[1:])] + assert max(gaps) < 150, ("encoded output stalled", max(gaps)) + + # Red/blue solid input fixtures make each cut's actual first picture + # distinguishable. This also proves that cuts need not wait for an IDR. + decoded = subprocess.run([ + args.ffmpeg, "-v", "error", "-threads", "1", "-f", "h264", "-i", "pipe:0", + "-vf", "crop=2:2:540:960", "-pix_fmt", "gray", "-fps_mode", "passthrough", + "-f", "rawvideo", "pipe:1", + ], input=b"".join(f["data"] for f in captured), capture_output=True, check=True, timeout=30).stdout + assert len(decoded) == len(captured) * 4, (len(decoded), len(captured)) + for i, frame in enumerate(captured): + y = decoded[i * 4] + frame["source"] = 0 if 60 <= y <= 95 else 1 if 15 <= y <= 45 else None + latencies, non_key_cuts = [], 0 + for i, cut in enumerate(cuts): + deadline = cuts[i + 1]["at"] if i + 1 < len(cuts) else cut_end + matches = [f for f in captured if cut["at"] <= f["at"] < deadline + and f["source"] == cut["target"]] + assert matches, ("cut picture did not arrive before the next cut", i, cut) + latencies.append((matches[0]["at"] - cut["at"]) * 1000) + non_key_cuts += not matches[0]["key"] + if args.minimum_ms: + assert non_key_cuts > len(cuts) // 2, "cuts appear to wait for keyframes" + print("KEYFRAME_LIMIT_RESULT " + json.dumps({ + "minimum_ms": args.minimum_ms, "inputs": args.source_count, + "cuts": len(cuts), "cut_pictures_without_keyframe": non_key_cuts, + "cut_latency_ms": {"median": statistics.median(latencies), "max": max(latencies)}, + "frames": len(measured), "fps": (len(measured) - 1) / (measured[-1]["at"] - measured[0]["at"]), + "encoded_gap_ms": {"p99": sorted(gaps)[int(.99 * (len(gaps) - 1))], "max": max(gaps)}, + "keyframes": len(keys), "min_keyframe_spacing_frames": min(key_spacing), + "status": state, + }), flush=True) + finally: + app.stop() + stop.set() + packet_reader.join(timeout=2) + + +if __name__ == "__main__": + parser = argparse.ArgumentParser(description=__doc__) + parser.add_argument("--inputs", nargs=2, required=True, metavar=("RED_H264", "BLUE_H264")) + parser.add_argument("--port", type=int, required=True, help="unused control port for this isolated graph") + parser.add_argument("--webui", required=True, help="existing WebUI backend") + parser.add_argument("--source-count", type=int, choices=(2, 16), default=16) + parser.add_argument("--minimum-ms", type=int, choices=(0, 100, 150, 200), default=150) + parser.add_argument("--cuts", type=int, default=80) + parser.add_argument("--ffmpeg", default="/usr/local/bin/ffmpeg") + run(parser.parse_args()) diff --git a/tests/smoke_format_eof.py b/tests/smoke_format_eof.py new file mode 100644 index 00000000..ce9c3e43 --- /dev/null +++ b/tests/smoke_format_eof.py @@ -0,0 +1,77 @@ +"""Native regression: decode a short A/V fixture, reach EOF, then reconnect. + +Run in the built avplumber environment: + python3 tests/smoke_format_eof.py +""" + +import sys +import time + +from pyplumber import AVPlumber +from pyplumber import node as n + + +def run(path, media, live): + avp = AVPlumber() + errors = [] + avp.on_exception = lambda *error: errors.append(error) + video = media == "video" + data_type = "VideoFrame" if video else "AudioSamples" + decoder = n.DecVideo if video else n.DecAudio + declaration = n.FakeVideoFormat if video else n.FakeAudioMetadata + avp.edges.planCapacity("*", 4096) + for cls, params in ( + (n.InputRec, {"name": "input", "url": path, "dst": "packets"}), + (n.Demux, {"name": "demux", "src": "packets", "routing": { + "v:0" if video else "a:0": "selected"}}), + (decoder, {"name": "decode", "src": "selected", "dst": "decoded"}), + ): + avp.addNode(cls({"group": "input", **params})) + params = {"name": "format", "group": "output", "src": "decoded", "dst": "result"} + if live: + params["ignore_eof"] = True + avp.addNode(declaration(params)) + source = avp.getEdge("decoded", data_type) + result = avp.getEdge("result", data_type) + passes = [] + try: + avp.group("output").startNodes() + for attempt in range(2 if live else 1): + previous_input = source.enqueued_total + if attempt: + avp.group("input").restartNodes() + else: + avp.group("input").startNodes() + deadline = time.monotonic() + 15 + while True: + assert not errors, errors + assert time.monotonic() < deadline, "decode/EOF did not complete" + if (source.enqueued_total > previous_input + and not avp.node("decode").isWorking + and source.occupied == 0): + if live or not avp.node("format").isWorking: + break + time.sleep(0.01) + rows = [] + eof = 0 + while result.occupied: + frame = result.wait_dequeue() + if frame.pts.timestamp == -(1 << 63): + eof += 1 + else: + rows.append((frame.pts.timestamp, frame.pts.timebase.num, + frame.pts.timebase.den)) + assert rows, "no decoded media reached the output" + assert eof == (0 if live else 1), (live, eof) + assert avp.node("format").isWorking == live + passes.append(rows) + if live: + assert passes[0] == passes[1], "reconnect changed or lost decoded frames" + print(f"PASS {media} {'live reconnect' if live else 'default finite EOF'}: " + f"{len(passes[0])} frames per pass", flush=True) + finally: + avp.shutdown() + + +if __name__ == "__main__": + run(sys.argv[1], sys.argv[2], sys.argv[3] == "live") diff --git a/tests/test_force_keyframe_native.py b/tests/test_force_keyframe_native.py new file mode 100644 index 00000000..41f3a7c1 --- /dev/null +++ b/tests/test_force_keyframe_native.py @@ -0,0 +1,144 @@ +"""Native node tests; run with the built pyplumber extension on the test host.""" + +from contextlib import contextmanager +from fractions import Fraction +import subprocess + +import pytest + +pytest.importorskip("_avplumber") + +from pyplumber import AVPlumber +from pyplumber.node import DecVideo, Demux, ForceKeyFrame, Input + + +@pytest.fixture(scope="module") +def frames(tmp_path_factory): + def generate(fps="60/1", count=125): + path = tmp_path_factory.mktemp("keyframe-frames") / "frames.nut" + subprocess.run([ + "ffmpeg", "-v", "error", "-n", "-f", "lavfi", "-i", + f"testsrc2=size=64x64:rate={fps}", "-frames:v", str(count), + "-c:v", "rawvideo", "-f", "nut", str(path), + ], check=True, capture_output=True, timeout=15) + avp = AVPlumber() + errors = [] + avp.on_exception = lambda *error: errors.append(error) + for cls, params in ( + (Input, {"url": str(path), "dst": "packets"}), + (Demux, {"src": "packets", "routing": {"v:0": "video"}}), + (DecVideo, {"src": "video", "dst": "frames"}), + ): + avp.addNode(cls({"group": "fixture", **params})) + edge = avp.getEdge("frames", "VideoFrame") + result = [] + try: + avp.group("fixture").startNodes() + for _ in range(count): + frame = edge.get(5000) + assert not errors, errors + result.append(frame) + return result + finally: + avp.shutdown() + return generate + + +@contextmanager +def limiter(**params): + avp = AVPlumber() + errors = [] + avp.on_exception = lambda *error: errors.append(error) + try: + avp.addNode(ForceKeyFrame({ + "name": "limit", "src": "input", "dst": "output", **params, + }), early_create=True) + source = avp.getEdge("input", "VideoFrame") + output = avp.getEdge("output", "VideoFrame") + avp.node("limit").start() + + def step(frame, requests=0): + for _ in range(requests): + avp.executeCommandsFromString("node.object.set limit trigger true") + source.enqueue(frame) + result = output.get(5000) + assert not errors, errors + assert result.pts.timestamp == frame.pts.timestamp + assert result.data == frame.data + return result.keyFrame + + yield step, lambda: avp.node("limit").getObject("status") + finally: + avp.shutdown() + + +@pytest.mark.parametrize("fps,minimum,spacing", [ + ("60/1", 100, 6), ("60/1", 150, 9), ("60/1", 200, 12), + ("60000/1001", 150, 9), ("60000/1001", 200, 12), ("24000/1001", 200, 5), +]) +def test_spam_is_coalesced_without_losing_frames(frames, fps, minimum, spacing): + footage = frames(fps=fps) + with limiter(min_interval_ms=minimum, interval_sec="1/1") as (step, status): + keys = [i for i, frame in enumerate(footage) if step(frame, requests=5)] + assert keys == list(range(0, len(footage), spacing)) + state = status() + assert state["requested_generation"] == len(footage) * 5 + assert state["triggered_frames"] == len(keys) + assert state["min_interval_ms"] == minimum + assert all(Fraction(b - a, 1) / Fraction(fps) >= Fraction(minimum, 1000) + for a, b in zip(keys, keys[1:])) + + +def test_last_request_survives_the_cooldown_and_is_not_repeated(frames): + with limiter(min_interval_ms=200) as (step, status): + keys = [] + for i, frame in enumerate(frames(count=30)): + if step(frame, requests=1 if i in (0, 1) else 0): + keys.append(i) + if 1 <= i < 12: + assert status()["pending"] + assert status()["forced_generation"] == 1 + if i >= 12: + assert not status()["pending"] + assert keys == [0, 12] + assert status()["forced_generation"] == 2 + + +def test_periodic_keyframe_cannot_bypass_a_recent_trigger(frames): + with limiter(min_interval_ms=200, interval_sec="1/1") as (step, status): + keys = [i for i, frame in enumerate(frames()) if step(frame, requests=int(i == 59))] + assert keys == [0, 59, 71, 120] + assert status()["periodic_frames"] == 3 + assert status()["triggered_frames"] == 1 + + +@pytest.mark.parametrize("params", [{}, {"min_interval_ms": 0}]) +def test_default_preserves_unlimited_triggering(frames, params): + with limiter(**params) as (step, status): + assert all(step(frame, requests=1) for frame in frames(count=15)) + assert status()["triggered_frames"] == 15 + assert not status()["pending"] + + +def test_periodic_forcing_can_be_disabled(frames): + with limiter(min_interval_ms=200) as (step, status): + assert not any(step(frame) for frame in frames()) + assert status()["periodic_frames"] == 0 + + +def test_backwards_pts_restarts_the_limit_and_duplicate_pts_do_not(frames): + footage = frames(count=30) + with limiter(min_interval_ms=200) as (step, _): + assert step(footage[24], requests=1) + assert not step(footage[24], requests=1) + assert step(footage[0]) # Pending request survives the timeline reset. + assert not step(footage[0], requests=1) + assert not step(footage[11]) + assert step(footage[12]) + + +@pytest.mark.parametrize("value", [-1, -0.1, 0.2, True, "200", 2**31]) +def test_invalid_interval_is_rejected(value): + with pytest.raises(RuntimeError, match="min_interval_ms"): + with limiter(min_interval_ms=value): + pytest.fail("invalid interval was accepted")