diff --git a/.gitignore b/.gitignore index 605ab2f5..3784c0cf 100644 --- a/.gitignore +++ b/.gitignore @@ -37,6 +37,7 @@ vite.config.d.ts electron/native/wgc-capture/build/ electron/native/cursor-monitor/build/ electron/native/gpu-export-probe/build/ +electron/native/nvidia-cuda-compositor/build/ electron/native/bin/*/whisper-* electron/native/bin/*/whisper-runtime.json diff --git a/electron/ipc/export/native-video.test.ts b/electron/ipc/export/native-video.test.ts index 3e004f1b..91f11c01 100644 --- a/electron/ipc/export/native-video.test.ts +++ b/electron/ipc/export/native-video.test.ts @@ -14,6 +14,7 @@ vi.mock("../ffmpeg/binary", () => ({ import { buildNativeVideoAudioMuxArgs, + getNvidiaCudaAudioExportSkipReason, normalizeNativeStaticLayoutBackground, parseFfmpegDurationSeconds, parseFfmpegFrameRate, @@ -21,8 +22,29 @@ import { parseNvidiaCudaExportSummary, parseWindowsGpuExportProgressLine, parseWindowsGpuExportSummary, + validateNvidiaCudaExportSummary, } from "./native-video"; +function withNvidiaCudaAudioOverride(value: string | undefined, callback: () => T) { + const envName = "RECORDLY_NVIDIA_CUDA_ALLOW_AUDIO_EXPORT"; + const originalValue = process.env[envName]; + if (value === undefined) { + delete process.env[envName]; + } else { + process.env[envName] = value; + } + + try { + return callback(); + } finally { + if (originalValue === undefined) { + delete process.env[envName]; + } else { + process.env[envName] = originalValue; + } + } +} + describe("normalizeNativeStaticLayoutBackground", () => { it("falls back to a solid background when the configured image file is missing", async () => { const normalized = await normalizeNativeStaticLayoutBackground({ @@ -46,6 +68,35 @@ describe("normalizeNativeStaticLayoutBackground", () => { }); }); +describe("getNvidiaCudaAudioExportSkipReason", () => { + it("allows video-only CUDA exports by default", () => { + withNvidiaCudaAudioOverride(undefined, () => { + expect(getNvidiaCudaAudioExportSkipReason(undefined)).toBeNull(); + expect(getNvidiaCudaAudioExportSkipReason("none")).toBeNull(); + }); + }); + + it("guards CUDA for audio exports unless explicitly overridden", () => { + withNvidiaCudaAudioOverride(undefined, () => { + expect(getNvidiaCudaAudioExportSkipReason("copy-source")).toBe( + "audio-mode:copy-source", + ); + expect(getNvidiaCudaAudioExportSkipReason("trim-source")).toBe( + "audio-mode:trim-source", + ); + expect(getNvidiaCudaAudioExportSkipReason("edited-track")).toBe( + "audio-mode:edited-track", + ); + }); + }); + + it("allows CUDA audio exports only for explicit lab overrides", () => { + withNvidiaCudaAudioOverride("1", () => { + expect(getNvidiaCudaAudioExportSkipReason("copy-source")).toBeNull(); + }); + }); +}); + describe("buildNativeVideoAudioMuxArgs", () => { it("stream-copies source audio and preserves the requested video duration", () => { const args = buildNativeVideoAudioMuxArgs("video.mp4", "source.mp4", "out.mp4", { @@ -151,6 +202,87 @@ describe("parseNvidiaCudaExportSummary", () => { }); }); +describe("validateNvidiaCudaExportSummary", () => { + it("accepts CUDA output when frames and stream durations match the export target", () => { + const issues = validateNvidiaCudaExportSummary( + { + success: true, + targetFrames: 300, + durationSec: 10, + nativeSummary: { success: true, frames: 300 }, + outputVideo: { duration: "9.999900", nb_frames: "300" }, + outputAudio: { duration: "10.005000" }, + }, + { durationSec: 10, targetFrames: 300 }, + ); + + expect(issues).toEqual([]); + }); + + it("rejects CUDA output that reports too few frames or a short video stream", () => { + const issues = validateNvidiaCudaExportSummary( + { + success: true, + targetFrames: 300, + durationSec: 10, + nativeSummary: { success: true, frames: 144 }, + outputVideo: { duration: "4.799952", nb_frames: "144" }, + outputAudio: { duration: "10.005000" }, + }, + { durationSec: 10, targetFrames: 300 }, + ); + + expect(issues).toEqual([ + "native frames 144 below expected minimum 285", + "output video frames 144 below expected minimum 285", + "output video duration 4.800s differs from expected 10.000s", + ]); + }); + + it("rejects audio CUDA output unless the helper reports timestamp-aligned frame selection", () => { + const issues = validateNvidiaCudaExportSummary( + { + success: true, + targetFrames: 300, + durationSec: 10, + nativeSummary: { + success: true, + frames: 300, + selectionStage: "decoder-policy-mapped-callback", + }, + outputVideo: { duration: "9.999900", nb_frames: "300" }, + outputAudio: { duration: "10.005000" }, + }, + { durationSec: 10, targetFrames: 300, requiresTimelineSync: true }, + ); + + expect(issues).toEqual([ + "CUDA timeline mode is not timestamp-aligned for audio export", + ]); + }); + + it("accepts audio CUDA output when the helper reports PTS-aligned selection", () => { + const issues = validateNvidiaCudaExportSummary( + { + success: true, + targetFrames: 300, + durationSec: 10, + nativeSummary: { + success: true, + frames: 300, + sourceTimestampMode: "pts", + selectionStage: "timestamp-mapped-callback", + }, + outputVideo: { duration: "9.999900", nb_frames: "300" }, + outputAudio: { duration: "10.005000" }, + }, + { durationSec: 10, targetFrames: 300, requiresTimelineSync: true }, + ); + + expect(issues).toEqual([]); + }); +}); + describe("parseWindowsGpuExportProgressLine", () => { it("parses bounded helper progress lines", () => { expect( diff --git a/electron/ipc/export/native-video.ts b/electron/ipc/export/native-video.ts index d7598f28..97964ada 100644 --- a/electron/ipc/export/native-video.ts +++ b/electron/ipc/export/native-video.ts @@ -14,6 +14,7 @@ import type { NativeStaticLayoutBackend, NativeStaticLayoutExportArgsConfig, NativeVideoAudioMuxMetrics, + NativeVideoExportAudioMode, NativeVideoExportFinishOptions, } from "../nativeVideoExport"; import { @@ -40,6 +41,7 @@ const getNowMs = () => performance.now(); const formatFfmpegSeconds = (milliseconds: number) => (milliseconds / 1000).toFixed(3); const MISSING_NATIVE_STATIC_BACKGROUND_COLOR = "#ffffff"; const NATIVE_EXPORT_HIGH_PRIORITY = os.constants.priority.PRIORITY_HIGH; +const NVIDIA_CUDA_ALLOW_AUDIO_EXPORT_ENV = "RECORDLY_NVIDIA_CUDA_ALLOW_AUDIO_EXPORT"; export type NativeVideoExportSession = { ffmpegProcess: ChildProcessByStdio; @@ -186,18 +188,24 @@ export interface NvidiaCudaExportSummary { bitrateMbps?: number; durationSec?: number; targetFrames?: number; + sourcePtsFrames?: number; + sourcePtsSource?: string; timingsMs?: { demux?: number; backgroundConvert?: number; cursorAtlas?: number; webcamConvert?: number; webcamDemux?: number; + sourcePtsProbe?: number; nativeEncode?: number; mux?: number; endToEnd?: number; }; nativeSummary?: { success?: boolean; + selectionStage?: string; + sourceTimestampMode?: string; + timelineMode?: string; frames?: number; totalMs?: number; fps?: number; @@ -332,6 +340,111 @@ export function parseNvidiaCudaExportSummary(stdout: string): NvidiaCudaExportSu } } +function getFiniteNumber(value: unknown) { + const numberValue = typeof value === "string" ? Number(value) : value; + return typeof numberValue === "number" && Number.isFinite(numberValue) + ? numberValue + : null; +} + +function getNvidiaCudaOutputStreamNumber(stream: unknown, property: string) { + if (!stream || typeof stream !== "object") { + return null; + } + + return getFiniteNumber((stream as Record)[property]); +} + +export function validateNvidiaCudaExportSummary( + summary: NvidiaCudaExportSummary, + expected: { + durationSec: number; + targetFrames: number; + requiresTimelineSync?: boolean; + }, +) { + const issues: string[] = []; + const expectedFrames = Math.max(1, Math.round(expected.targetFrames)); + const minimumFrames = Math.max(1, Math.floor(expectedFrames * 0.95)); + const expectedDurationSec = Math.max(0, expected.durationSec); + const durationToleranceSec = Math.min( + 2, + Math.max(0.5, expectedDurationSec * 0.02), + ); + const nativeFrames = getFiniteNumber(summary.nativeSummary?.frames); + const outputVideoFrames = getNvidiaCudaOutputStreamNumber( + summary.outputVideo, + "nb_frames", + ); + const outputVideoDurationSec = getNvidiaCudaOutputStreamNumber( + summary.outputVideo, + "duration", + ); + const outputAudioDurationSec = getNvidiaCudaOutputStreamNumber( + summary.outputAudio, + "duration", + ); + + if (!summary.outputVideo) { + issues.push("missing output video probe"); + } + if ( + expected.requiresTimelineSync && + !isNvidiaCudaTimestampAlignedSummary(summary) + ) { + issues.push( + "CUDA timeline mode is not timestamp-aligned for audio export", + ); + } + if (nativeFrames === null) { + issues.push("missing native frame count"); + } else if (nativeFrames < minimumFrames) { + issues.push( + `native frames ${nativeFrames} below expected minimum ${minimumFrames}`, + ); + } + if (outputVideoFrames !== null && outputVideoFrames < minimumFrames) { + issues.push( + `output video frames ${outputVideoFrames} below expected minimum ${minimumFrames}`, + ); + } + if ( + outputVideoDurationSec !== null && + Math.abs(outputVideoDurationSec - expectedDurationSec) > durationToleranceSec + ) { + issues.push( + `output video duration ${outputVideoDurationSec.toFixed( + 3, + )}s differs from expected ${expectedDurationSec.toFixed(3)}s`, + ); + } + if ( + outputAudioDurationSec !== null && + Math.abs(outputAudioDurationSec - expectedDurationSec) > durationToleranceSec + ) { + issues.push( + `output audio duration ${outputAudioDurationSec.toFixed( + 3, + )}s differs from expected ${expectedDurationSec.toFixed(3)}s`, + ); + } + + return issues; +} + +function isNvidiaCudaTimestampAlignedSummary(summary: NvidiaCudaExportSummary) { + const nativeSummary = summary.nativeSummary; + const timelineFields = [ + nativeSummary?.sourceTimestampMode, + nativeSummary?.timelineMode, + nativeSummary?.selectionStage, + ]; + return timelineFields.some((value) => { + const normalized = String(value ?? "").toLowerCase(); + return normalized.includes("pts") || normalized.includes("timestamp"); + }); +} + function shouldPersistNvidiaCudaExportDiagnostics() { return ( process.env.RECORDLY_NVIDIA_CUDA_EXPORT_DIAGNOSTICS === "1" || @@ -1048,12 +1161,34 @@ async function resolveExperimentalWindowsGpuExporterPath() { return null; } -function isExperimentalNvidiaCudaExportEnabled(options: NativeStaticLayoutExportOptions) { - return Boolean( - process.platform === "win32" && - process.env.RECORDLY_EXPERIMENTAL_NVIDIA_CUDA_EXPORT === "1" && - options.experimentalWindowsGpuCompositor, - ); +export function getNvidiaCudaAudioExportSkipReason( + audioMode: NativeVideoExportAudioMode | undefined, +) { + const resolvedAudioMode = audioMode ?? "none"; + if (resolvedAudioMode === "none") { + return null; + } + if (process.env[NVIDIA_CUDA_ALLOW_AUDIO_EXPORT_ENV] === "1") { + return null; + } + + return `audio-mode:${resolvedAudioMode}`; +} + +function getExperimentalNvidiaCudaExportSkipReason( + options: NativeStaticLayoutExportOptions, +) { + if (process.platform !== "win32") { + return "not-windows"; + } + if (process.env.RECORDLY_EXPERIMENTAL_NVIDIA_CUDA_EXPORT !== "1") { + return "env-disabled"; + } + if (!options.experimentalWindowsGpuCompositor) { + return "windows-gpu-compositor-disabled"; + } + + return getNvidiaCudaAudioExportSkipReason(options.audioOptions?.audioMode); } async function resolveExperimentalNvidiaCudaExportScriptPath() { @@ -1067,6 +1202,28 @@ async function resolveExperimentalNvidiaCudaExportScriptPath() { candidates.push(configuredPath); } candidates.push( + path.join( + process.cwd(), + "electron", + "native", + "nvidia-cuda-compositor", + "run-mp4-pipeline.mjs", + ), + path.join( + app.getAppPath(), + "electron", + "native", + "nvidia-cuda-compositor", + "run-mp4-pipeline.mjs", + ), + path.join( + process.resourcesPath, + "app.asar.unpacked", + "electron", + "native", + "nvidia-cuda-compositor", + "run-mp4-pipeline.mjs", + ), path.join(process.cwd(), ".tmp", "nvdec-nvenc-probe", "run-mp4-pipeline.mjs"), path.join(app.getAppPath(), ".tmp", "nvdec-nvenc-probe", "run-mp4-pipeline.mjs"), ); @@ -1664,6 +1821,7 @@ async function runExperimentalNvidiaCudaStaticLayoutExport( let stdout = ""; let stderr = ""; let stderrLineBuffer = ""; + let lastProgressPercentage = 0; let settled = false; const timeout = setTimeout(() => { if (settled) return; @@ -1693,8 +1851,13 @@ async function runExperimentalNvidiaCudaStaticLayoutExport( : elapsedMs > 0 && progress.currentFrame > 0 ? (progress.currentFrame * 1000) / elapsedMs : undefined; + lastProgressPercentage = Math.max( + lastProgressPercentage, + progress.percentage, + ); onProgress?.({ ...progress, + percentage: lastProgressPercentage, sessionId: options.sessionId, backend: "nvidia-cuda-compositor", elapsedMs, @@ -2030,7 +2193,23 @@ export async function exportNativeStaticLayoutVideo( } } let experimentalNvidiaCudaOptions = experimentalGpuOptions; - let shouldTryNvidiaCuda = isExperimentalNvidiaCudaExportEnabled(options); + const nvidiaCudaSkipReason = + getExperimentalNvidiaCudaExportSkipReason(options); + let shouldTryNvidiaCuda = nvidiaCudaSkipReason === null; + if ( + !shouldTryNvidiaCuda && + process.env.RECORDLY_EXPERIMENTAL_NVIDIA_CUDA_EXPORT === "1" && + nvidiaCudaSkipReason !== "env-disabled" + ) { + console.warn( + "[native-static-layout-export] Skipping NVIDIA CUDA compositor; falling back to Windows GPU compositor", + { + reason: nvidiaCudaSkipReason, + audioMode: options.audioOptions?.audioMode ?? "none", + overrideEnv: NVIDIA_CUDA_ALLOW_AUDIO_EXPORT_ENV, + }, + ); + } if (shouldTryNvidiaCuda && options.cursorTelemetry?.length) { const cursorTelemetryPath = await prepareNvidiaCudaCursorTelemetry( options, @@ -2064,6 +2243,21 @@ export async function exportNativeStaticLayoutVideo( session, onProgress, ); + const cudaValidationIssues = validateNvidiaCudaExportSummary( + cudaResult.summary, + { + durationSec: options.durationSec, + targetFrames: Math.ceil(options.durationSec * options.frameRate), + requiresTimelineSync: + (experimentalNvidiaCudaOptions.audioOptions?.audioMode ?? + "none") !== "none", + }, + ); + if (cudaValidationIssues.length > 0) { + throw new Error( + `Experimental NVIDIA CUDA compositor produced an invalid output: ${cudaValidationIssues.join("; ")}`, + ); + } console.info( "[native-static-layout-export] NVIDIA CUDA compositor completed", { @@ -2092,9 +2286,9 @@ export async function exportNativeStaticLayoutVideo( webcamOverlay: cudaResult.summary.nativeSummary?.webcamOverlay, cursorAtlas: cudaResult.summary.nativeSummary?.cursorAtlas, zoomOverlay: cudaResult.summary.nativeSummary?.zoomOverlay, - zoomSamples: cudaResult.summary.nativeSummary?.zoomSamples, - }, - ); + zoomSamples: cudaResult.summary.nativeSummary?.zoomSamples, + }, + ); const outputStat = await fs.stat(videoOnlyPath); metrics.chunkCount = 1; metrics.chunkDurationSec = options.durationSec; diff --git a/electron/ipc/register/recording.ts b/electron/ipc/register/recording.ts index 8ca4720f..a019a341 100644 --- a/electron/ipc/register/recording.ts +++ b/electron/ipc/register/recording.ts @@ -20,8 +20,8 @@ import { startNativeCursorMonitor, stopNativeCursorMonitor } from "../cursor/mon import { clamp, pauseCursorCapture, - resumeCursorCapture, resetCursorCaptureClock, + resumeCursorCapture, sampleCursorPoint, snapshotCursorTelemetryForPersistence, startCursorSampling, @@ -1082,6 +1082,18 @@ export function registerRecordingHandlers( try { return await finalizeStoredVideo(videoPath); } catch { + try { + await validateRecordedVideo(videoPath); + return { + success: false, + path: videoPath, + message: "Failed to mux native Windows recording", + error: String(error), + }; + } catch { + // The fallback path is not safely playable; surface the original mux error. + } + return { success: false, message: "Failed to mux native Windows recording", diff --git a/electron/native/bin/win32-x64/helpers-manifest.json b/electron/native/bin/win32-x64/helpers-manifest.json index 6b3d48aa..e2c5313f 100644 --- a/electron/native/bin/win32-x64/helpers-manifest.json +++ b/electron/native/bin/win32-x64/helpers-manifest.json @@ -23,6 +23,13 @@ "sourceDir": "electron/native/gpu-export-probe", "sourceFingerprint": "75bf080c4a5cbbcb1d42fb088a1545130c613e42d9e7681cc23f9151d1c8072b", "updatedAt": "2026-05-03T15:39:47.938Z" + }, + "recordly-nvidia-cuda-compositor": { + "binaryName": "recordly-nvidia-cuda-compositor.exe", + "binarySha256": "f274039c46445a3b8e57c9faa6307238d6facc04d5c88f7bb84c43563feda649", + "sourceDir": "electron/native/nvidia-cuda-compositor", + "sourceFingerprint": "bc6f657ff3226f8c3bcdc89354de4118d0ef0f3ac1b5fde6d496bda9b77b65c6", + "updatedAt": "2026-05-04T02:54:12.432Z" } } } diff --git a/electron/native/bin/win32-x64/recordly-nvidia-cuda-compositor.exe b/electron/native/bin/win32-x64/recordly-nvidia-cuda-compositor.exe new file mode 100644 index 00000000..2dfd3e71 Binary files /dev/null and b/electron/native/bin/win32-x64/recordly-nvidia-cuda-compositor.exe differ diff --git a/electron/native/nvidia-cuda-compositor/CMakeLists.txt b/electron/native/nvidia-cuda-compositor/CMakeLists.txt new file mode 100644 index 00000000..e79d9630 --- /dev/null +++ b/electron/native/nvidia-cuda-compositor/CMakeLists.txt @@ -0,0 +1,47 @@ +cmake_minimum_required(VERSION 3.24) + +project(recordly_nvidia_cuda_compositor LANGUAGES CXX CUDA) + +set(CMAKE_CXX_STANDARD 17) +set(CMAKE_CXX_STANDARD_REQUIRED ON) +set(CMAKE_CUDA_STANDARD 17) +set(CMAKE_CUDA_STANDARD_REQUIRED ON) + +set(RECORDLY_NVIDIA_VIDEO_CODEC_SDK_ROOT + "${CMAKE_CURRENT_LIST_DIR}/../../../.tmp/video-sdk-samples" + CACHE PATH + "Path to the NVIDIA Video Codec SDK samples checkout" +) +set(NVIDIA_SAMPLES_DIR "${RECORDLY_NVIDIA_VIDEO_CODEC_SDK_ROOT}/Samples") +set(NVCODEC_DIR "${NVIDIA_SAMPLES_DIR}/NvCodec") + +if(NOT EXISTS "${NVCODEC_DIR}") + message(FATAL_ERROR "NVIDIA Video Codec SDK NvCodec directory not found: ${NVCODEC_DIR}. Set RECORDLY_NVIDIA_VIDEO_CODEC_SDK_ROOT.") +endif() + +add_executable(recordly-nvidia-cuda-compositor + src/main.cu + "${NVCODEC_DIR}/NvDecoder/NvDecoder.cpp" + "${NVCODEC_DIR}/NvEncoder/NvEncoder.cpp" + "${NVCODEC_DIR}/NvEncoder/NvEncoderCuda.cpp" +) + +target_include_directories(recordly-nvidia-cuda-compositor PRIVATE + "${NVIDIA_SAMPLES_DIR}" + "${NVCODEC_DIR}" +) + +target_compile_definitions(recordly-nvidia-cuda-compositor PRIVATE + NOMINMAX + WIN32_LEAN_AND_MEAN +) + +find_package(CUDAToolkit REQUIRED) +target_link_libraries(recordly-nvidia-cuda-compositor PRIVATE + CUDA::cuda_driver + "${NVCODEC_DIR}/Lib/x64/nvcuvid.lib" +) + +set_target_properties(recordly-nvidia-cuda-compositor PROPERTIES + CUDA_SEPARABLE_COMPILATION OFF +) diff --git a/electron/native/nvidia-cuda-compositor/render-tahoe-cursor-atlas.cjs b/electron/native/nvidia-cuda-compositor/render-tahoe-cursor-atlas.cjs new file mode 100644 index 00000000..233ed252 --- /dev/null +++ b/electron/native/nvidia-cuda-compositor/render-tahoe-cursor-atlas.cjs @@ -0,0 +1,144 @@ +const { app, BrowserWindow, ipcMain } = require("electron"); +const fs = require("node:fs"); +const path = require("node:path"); + +const drawHeight = 256; +const padding = 2; +const cursorTypes = [ + "arrow", + "text", + "pointer", + "crosshair", + "open-hand", + "closed-hand", + "resize-ew", + "resize-ns", + "not-allowed", +]; +const tahoeAssets = { + arrow: ["pointer-1__14-6.svg", 0.14, 0.06], + text: ["ibeam-1__50-44.svg", 0.5, 0.44], + pointer: ["pointinghand-1__40-10.svg", 0.4, 0.1], + crosshair: ["crosshair-1__50-50.svg", 0.5, 0.5], + "open-hand": ["openhand-1__55-57.svg", 0.55, 0.57], + "closed-hand": ["closedhand-1__50-46.svg", 0.5, 0.46], + "resize-ew": ["resizeeastwest-1__50-50.svg", 0.5, 0.5], + "resize-ns": ["resizenorthsouth-1__50-49.svg", 0.5, 0.49], + "not-allowed": ["notallowed-1__23-0.svg", 0.23, 0], +}; + +function arg(name, fallback = "") { + const index = process.argv.indexOf(name); + return index >= 0 && index + 1 < process.argv.length ? process.argv[index + 1] : fallback; +} + +const repoRoot = arg("--repo-root"); +const atlasRgbaPath = arg("--output-rgba"); +const atlasMetadataPath = arg("--output-metadata"); + +if (!repoRoot || !atlasRgbaPath || !atlasMetadataPath) { + console.error("Usage: electron render-tahoe-cursor-atlas.cjs --repo-root --output-rgba --output-metadata "); + process.exit(1); +} + +const assets = cursorTypes.map((type, index) => { + const [fileName, anchorX, anchorY] = tahoeAssets[type]; + return { + type, + index, + filePath: path.join(repoRoot, "src", "assets", "cursors", "tahoe", fileName), + anchorX, + anchorY, + }; +}); + +app.disableHardwareAcceleration(); + +app.whenReady().then(async () => { + const window = new BrowserWindow({ + show: false, + width: 1, + height: 1, + webPreferences: { + nodeIntegration: true, + contextIsolation: false, + backgroundThrottling: false, + }, + }); + + ipcMain.once("atlas-ready", (_event, result) => { + console.log(JSON.stringify(result)); + app.quit(); + }); + + const html = ` + + +`; + + await window.loadURL("data:text/html;charset=utf-8," + encodeURIComponent(html)); +}); diff --git a/electron/native/nvidia-cuda-compositor/run-mp4-pipeline.mjs b/electron/native/nvidia-cuda-compositor/run-mp4-pipeline.mjs new file mode 100644 index 00000000..88f7bd4f --- /dev/null +++ b/electron/native/nvidia-cuda-compositor/run-mp4-pipeline.mjs @@ -0,0 +1,1104 @@ +import { spawn, spawnSync } from "node:child_process"; +import { existsSync, mkdirSync, readFileSync, writeFileSync } from "node:fs"; +import { createRequire } from "node:module"; +import os from "node:os"; +import { basename, dirname, join, resolve } from "node:path"; +import { fileURLToPath } from "node:url"; + +const scriptDir = dirname(fileURLToPath(import.meta.url)); +const require = createRequire(import.meta.url); +const repoRoot = resolve(scriptDir, "..", "..", ".."); + +const cursorTypes = [ + "arrow", + "text", + "pointer", + "crosshair", + "open-hand", + "closed-hand", + "resize-ew", + "resize-ns", + "not-allowed", +]; +const cursorTypeIndexes = new Map(cursorTypes.map((type, index) => [type, index])); + +function fail(message) { + throw new Error(message); +} + +function getArg(name, fallback = "") { + const index = process.argv.indexOf(name); + if (index === -1) { + return fallback; + } + if (index + 1 >= process.argv.length) { + fail(`Missing value for ${name}`); + } + return process.argv[index + 1]; +} + +function getNumberArg(name, fallback) { + const value = getArg(name, ""); + if (!value) { + return fallback; + } + const parsed = Number(value); + if (!Number.isFinite(parsed) || parsed <= 0) { + fail(`Invalid ${name}: ${value}`); + } + return parsed; +} + +function getNonNegativeNumberArg(name, fallback) { + const value = getArg(name, ""); + if (!value) { + return fallback; + } + const parsed = Number(value); + if (!Number.isFinite(parsed) || parsed < 0) { + fail(`Invalid ${name}: ${value}`); + } + return parsed; +} + +function hasArg(name) { + return process.argv.includes(name); +} + +function shouldRaiseChildPriority() { + return process.env.RECORDLY_NVIDIA_CUDA_EXPORT_HIGH_PRIORITY !== "0"; +} + +function raiseChildPriority(child, label) { + if (!shouldRaiseChildPriority() || !child.pid) { + return false; + } + + try { + os.setPriority(child.pid, os.constants.priority.PRIORITY_HIGH); + return true; + } catch (error) { + console.warn(`[nvidia-cuda-export] Failed to raise ${label} priority: ${error}`); + return false; + } +} + +function run(command, args, options = {}) { + const startedAt = performance.now(); + const result = spawnSync(command, args, { + encoding: "utf8", + maxBuffer: 64 * 1024 * 1024, + windowsHide: true, + ...options, + }); + const elapsedMs = performance.now() - startedAt; + if (result.error) { + fail(`${command} failed to start: ${result.error.message}`); + } + if (result.status !== 0) { + const stderr = result.stderr?.trim(); + const stdout = result.stdout?.trim(); + fail( + `${command} exited with ${result.status}` + + (stderr ? `\nSTDERR:\n${stderr}` : "") + + (stdout ? `\nSTDOUT:\n${stdout}` : ""), + ); + } + return { + elapsedMs, + stdout: result.stdout ?? "", + stderr: result.stderr ?? "", + }; +} + +function emitPreparationProgress(totalFrames, percentage) { + const payload = { + currentFrame: 0, + totalFrames: Math.max(1, Math.floor(totalFrames)), + percentage: Number(Math.min(99, Math.max(0, percentage)).toFixed(2)), + }; + process.stderr.write(`PROGRESS ${JSON.stringify(payload)}\n`); +} + +function parseFfmpegStatsFrameCount(stderr) { + const matches = [...String(stderr ?? "").matchAll(/frame=\s*(\d+)/g)]; + if (!matches.length) { + return 0; + } + const frameCount = Number(matches[matches.length - 1][1]); + return Number.isFinite(frameCount) && frameCount > 0 ? Math.floor(frameCount) : 0; +} + +function sampleGpu() { + const result = spawnSync( + "nvidia-smi", + [ + "--query-gpu=timestamp,temperature.gpu,power.draw,pstate,utilization.gpu,utilization.decoder,utilization.encoder,clocks.sm,clocks.mem", + "--format=csv,noheader,nounits", + ], + { encoding: "utf8", windowsHide: true }, + ); + if (result.status !== 0 || !result.stdout.trim()) { + return null; + } + const parts = result.stdout.trim().split(",").map((part) => part.trim()); + return { + timestamp: parts[0], + temperatureC: Number(parts[1]), + powerW: Number(parts[2]), + pstate: parts[3], + gpuUtilizationPct: Number(parts[4]), + decoderUtilizationPct: Number(parts[5]), + encoderUtilizationPct: Number(parts[6]), + smClockMhz: Number(parts[7]), + memoryClockMhz: Number(parts[8]), + }; +} + +function summarizeGpuSamples(samples) { + if (!samples.length) { + return null; + } + const numeric = [ + "temperatureC", + "powerW", + "gpuUtilizationPct", + "decoderUtilizationPct", + "encoderUtilizationPct", + "smClockMhz", + "memoryClockMhz", + ]; + const summary = { + samples: samples.length, + pstateValues: [...new Set(samples.map((sample) => sample.pstate))], + }; + for (const key of numeric) { + const values = samples.map((sample) => sample[key]).filter(Number.isFinite); + if (!values.length) { + continue; + } + summary[key] = { + min: Math.min(...values), + max: Math.max(...values), + avg: Number((values.reduce((sum, value) => sum + value, 0) / values.length).toFixed(2)), + }; + } + return summary; +} + +function cursorBounceScale(interactionType, ageMs, durationMs = 180) { + if (!["click", "double-click", "right-click", "middle-click"].includes(interactionType)) { + return 1; + } + if (ageMs < 0 || ageMs > durationMs) { + return 1; + } + const progress = 1 - ageMs / durationMs; + return Math.max(0.72, 1 - Math.sin(progress * Math.PI) * 0.08); +} + +function latestClickSample(samples, sampleIndex) { + for (let index = sampleIndex; index >= 0; index -= 1) { + const sample = samples[index]; + if (["click", "double-click", "right-click", "middle-click"].includes(sample?.interactionType)) { + return sample; + } + } + return null; +} + +function writeCursorSamples(cursorPayload, outputPath) { + const samples = Array.isArray(cursorPayload.samples) ? cursorPayload.samples : []; + const cursorLines = samples + .map((sample, index) => { + if ( + !Number.isFinite(sample?.timeMs) || + !Number.isFinite(sample?.cx) || + !Number.isFinite(sample?.cy) + ) { + return null; + } + const clickSample = latestClickSample(samples, index); + const bounceScale = Number.isFinite(sample.bounceScale) + ? sample.bounceScale + : clickSample + ? cursorBounceScale(clickSample.interactionType, sample.timeMs - clickSample.timeMs) + : 1; + return [ + sample.timeMs, + sample.cx, + sample.cy, + cursorTypeIndexes.get(sample.cursorType) ?? + (Number.isFinite(sample.cursorTypeIndex) ? Math.max(0, Math.min(8, Math.round(sample.cursorTypeIndex))) : 0), + Number(bounceScale.toFixed(4)), + ].join("\t"); + }) + .filter(Boolean) + .join("\n"); + writeFileSync(outputPath, cursorLines ? `${cursorLines}\n` : ""); + return samples.length; +} + +function renderTahoeCursorAtlas(workDir) { + const rgbaPath = join(workDir, "tahoe-cursor-atlas.rgba"); + const metadataPath = join(workDir, "tahoe-cursor-atlas.tsv"); + const electronPath = require("electron"); + const render = run(electronPath, [ + join(scriptDir, "render-tahoe-cursor-atlas.cjs"), + "--repo-root", + repoRoot, + "--output-rgba", + rgbaPath, + "--output-metadata", + metadataPath, + ], { + env: { + ...process.env, + ELECTRON_DISABLE_SECURITY_WARNINGS: "1", + }, + }); + const resultLine = render.stdout.trim().split(/\r?\n/).filter(Boolean).pop(); + if (!resultLine) { + fail("Cursor atlas renderer did not report a result"); + } + const result = JSON.parse(resultLine); + if (result.error) { + fail(`Cursor atlas renderer failed: ${result.error}`); + } + return { + rgbaPath, + metadataPath, + width: result.width, + height: result.height, + entries: result.entries, + elapsedMs: render.elapsedMs, + }; +} + +function prepareExternalCursorAtlas(workDir, pngPath, metadataPath) { + const startedAt = performance.now(); + const resolvedPngPath = resolve(pngPath); + const resolvedMetadataPath = resolve(metadataPath); + if (!existsSync(resolvedPngPath)) { + fail(`Cursor atlas PNG does not exist: ${resolvedPngPath}`); + } + if (!existsSync(resolvedMetadataPath)) { + fail(`Cursor atlas metadata does not exist: ${resolvedMetadataPath}`); + } + + const json = ffprobeJson([ + "-select_streams", + "v:0", + "-show_entries", + "stream=width,height", + resolvedPngPath, + ]); + const stream = json.streams?.[0]; + const width = Number(stream?.width); + const height = Number(stream?.height); + if (!Number.isFinite(width) || !Number.isFinite(height) || width <= 0 || height <= 0) { + fail(`Invalid cursor atlas dimensions: ${resolvedPngPath}`); + } + + const rgbaPath = join(workDir, "external-cursor-atlas.rgba"); + run("ffmpeg", [ + "-y", + "-hide_banner", + "-loglevel", + "error", + "-i", + resolvedPngPath, + "-f", + "rawvideo", + "-pix_fmt", + "rgba", + rgbaPath, + ]); + const entries = readFileSync(resolvedMetadataPath, "utf8") + .split(/\r?\n/) + .map((line) => line.trim()) + .filter(Boolean).length; + if (entries === 0) { + fail(`Cursor atlas metadata is empty: ${resolvedMetadataPath}`); + } + + return { + rgbaPath, + metadataPath: resolvedMetadataPath, + width, + height, + entries, + elapsedMs: performance.now() - startedAt, + }; +} + +async function runWithGpuMonitor(command, args, sampleIntervalMs) { + const samples = []; + const startedAt = performance.now(); + const shouldSampleGpu = Number.isFinite(sampleIntervalMs) && sampleIntervalMs > 0; + let stdout = ""; + let stderr = ""; + + const child = spawn(command, args, { + windowsHide: true, + stdio: ["ignore", "pipe", "pipe"], + }); + const priorityBoosted = raiseChildPriority(child, "native CUDA encoder"); + + const collectSample = () => { + if (!shouldSampleGpu) { + return; + } + const sample = sampleGpu(); + if (sample) { + samples.push({ + elapsedMs: Number((performance.now() - startedAt).toFixed(2)), + ...sample, + }); + } + }; + collectSample(); + const interval = shouldSampleGpu ? setInterval(collectSample, sampleIntervalMs) : null; + + child.stdout.on("data", (chunk) => { + stdout += chunk.toString(); + }); + child.stderr.on("data", (chunk) => { + const text = chunk.toString(); + stderr += text; + process.stderr.write(text); + }); + + const status = await new Promise((resolve, reject) => { + child.on("error", reject); + child.on("close", resolve); + }); + if (interval) { + clearInterval(interval); + } + collectSample(); + + const elapsedMs = performance.now() - startedAt; + if (status !== 0) { + fail( + `${command} exited with ${status}` + + (stderr.trim() ? `\nSTDERR:\n${stderr.trim()}` : "") + + (stdout.trim() ? `\nSTDOUT:\n${stdout.trim()}` : ""), + ); + } + return { + elapsedMs, + stdout, + stderr, + gpuSamples: samples, + gpuSummary: summarizeGpuSamples(samples), + priorityBoosted, + }; +} + +function ffprobeJson(args) { + const result = run("ffprobe", ["-v", "error", ...args, "-of", "json"]); + return JSON.parse(result.stdout); +} + +function ffprobeCsv(args) { + const result = run("ffprobe", ["-v", "error", ...args, "-of", "csv=p=0"]); + return result.stdout; +} + +function getVideoInfo(inputPath) { + const json = ffprobeJson([ + "-select_streams", + "v:0", + "-show_entries", + "stream=codec_name,width,height,duration,avg_frame_rate,nb_frames", + inputPath, + ]); + let stream = json.streams?.[0]; + if (!stream) { + fail(`No video stream found in ${inputPath}`); + } + if (stream.codec_name !== "h264") { + fail(`The NVIDIA CUDA compositor currently expects H.264 input, got ${stream.codec_name}`); + } + const durationSec = Number(stream.duration); + if (!Number.isFinite(durationSec) || durationSec <= 0) { + fail("ffprobe did not return a valid video duration"); + } + let sourceFrames = Number(stream.nb_frames); + if (!Number.isFinite(sourceFrames) || sourceFrames <= 0) { + const countedJson = ffprobeJson([ + "-count_frames", + "-select_streams", + "v:0", + "-show_entries", + "stream=codec_name,width,height,duration,avg_frame_rate,nb_frames,nb_read_frames", + inputPath, + ]); + stream = countedJson.streams?.[0] ?? stream; + sourceFrames = Number(stream.nb_read_frames || stream.nb_frames); + } + if (!Number.isFinite(sourceFrames) || sourceFrames <= 0) { + fail("ffprobe did not return a valid source frame count"); + } + return { + codec: stream.codec_name, + width: Number(stream.width), + height: Number(stream.height), + durationSec, + sourceFrames, + avgFrameRate: stream.avg_frame_rate, + }; +} + +function normalizeMonotonicTimestamps(timestamps) { + if (timestamps.length < 2) { + return []; + } + + const first = timestamps[0]; + const normalized = timestamps + .map((value) => Math.max(0, value - first)) + .filter((value, index, values) => index === 0 || value >= values[index - 1]); + return normalized.length === timestamps.length ? normalized : []; +} + +function parseTimestampCsv(csv) { + return csv + .split(/\r?\n/) + .map((line) => line.trim()) + .filter(Boolean) + .map((line) => { + const columns = line.split(","); + const value = columns + .map((column) => Number(column.trim())) + .find((candidate) => Number.isFinite(candidate)); + return value ?? Number.NaN; + }) + .filter((value) => Number.isFinite(value)); +} + +function getVideoPacketPts(inputPath, durationSec) { + const csv = ffprobeCsv([ + "-select_streams", + "v:0", + "-read_intervals", + `%+${durationSec}`, + "-show_packets", + "-show_entries", + "packet=pts_time,dts_time", + inputPath, + ]); + return normalizeMonotonicTimestamps(parseTimestampCsv(csv)); +} + +function getVideoFramePts(inputPath, durationSec) { + const csv = ffprobeCsv([ + "-select_streams", + "v:0", + "-read_intervals", + `%+${durationSec}`, + "-show_frames", + "-show_entries", + "frame=best_effort_timestamp_time", + inputPath, + ]); + return normalizeMonotonicTimestamps(parseTimestampCsv(csv)); +} + +function writeFramePtsSidecar(inputPath, durationSec, outputPath) { + const startedAt = performance.now(); + let source = "packet-pts"; + let timestamps = getVideoPacketPts(inputPath, durationSec); + if (timestamps.length === 0) { + source = "frame-pts"; + timestamps = getVideoFramePts(inputPath, durationSec); + } + const elapsedMs = performance.now() - startedAt; + if (timestamps.length === 0) { + return { path: null, frames: 0, elapsedMs, source: "none" }; + } + + writeFileSync(outputPath, timestamps.map((value) => value.toFixed(9)).join("\n")); + return { path: outputPath, frames: timestamps.length, elapsedMs, source }; +} + +function roundedRectMaskExpression({ x, y, width, height, radius }) { + const right = x + width; + const bottom = y + height; + const cornerRight = right - radius; + const cornerBottom = bottom - radius; + const radiusSquared = radius * radius; + const centerBand = `between(X,${x + radius},${cornerRight})*between(Y,${y},${bottom})`; + const middleBand = `between(X,${x},${right})*between(Y,${y + radius},${cornerBottom})`; + const topLeft = `lte((X-${x + radius})*(X-${x + radius})+(Y-${y + radius})*(Y-${y + radius}),${radiusSquared})*lte(X,${x + radius})*lte(Y,${y + radius})`; + const topRight = `lte((X-${cornerRight})*(X-${cornerRight})+(Y-${y + radius})*(Y-${y + radius}),${radiusSquared})*gte(X,${cornerRight})*lte(Y,${y + radius})`; + const bottomLeft = `lte((X-${x + radius})*(X-${x + radius})+(Y-${cornerBottom})*(Y-${cornerBottom}),${radiusSquared})*lte(X,${x + radius})*gte(Y,${cornerBottom})`; + const bottomRight = `lte((X-${cornerRight})*(X-${cornerRight})+(Y-${cornerBottom})*(Y-${cornerBottom}),${radiusSquared})*gte(X,${cornerRight})*gte(Y,${cornerBottom})`; + return `${centerBand}+${middleBand}+${topLeft}+${topRight}+${bottomLeft}+${bottomRight}`; +} + +function createBackgroundFilter(videoInfo, shadowOptions) { + const base = `[0:v]scale=${videoInfo.width}:${videoInfo.height}:force_original_aspect_ratio=increase,crop=${videoInfo.width}:${videoInfo.height},format=rgba[bg]`; + if (!shadowOptions) { + return { + filterArgs: [ + "-vf", + `scale=${videoInfo.width}:${videoInfo.height}:force_original_aspect_ratio=increase,crop=${videoInfo.width}:${videoInfo.height},format=nv12`, + ], + bakedShadow: false, + }; + } + + const shadowAlpha = Math.round(255 * Math.min(0.5, shadowOptions.intensityPct / 200)); + const shadowBlur = Math.max(12, Math.round(shadowOptions.radius * 1.5)); + const mask = roundedRectMaskExpression({ + x: shadowOptions.x, + y: shadowOptions.y, + width: shadowOptions.width, + height: shadowOptions.height, + radius: shadowOptions.radius, + }); + const shadow = `[1:v]format=rgba,geq=r='0':g='0':b='0':a='if(${mask},${shadowAlpha},0)',boxblur=luma_radius=${shadowBlur}:luma_power=1:chroma_radius=${shadowBlur}:chroma_power=1:alpha_radius=${shadowBlur}:alpha_power=1[shadow]`; + return { + filterArgs: [ + "-f", + "lavfi", + "-i", + `color=c=black@0.0:s=${videoInfo.width}x${videoInfo.height}:d=1`, + "-filter_complex", + `${base};${shadow};[bg][shadow]overlay=format=auto,format=nv12[out]`, + "-map", + "[out]", + ], + bakedShadow: true, + shadowAlpha, + shadowBlur, + }; +} + +function parseProbeSummary(stdout) { + const lines = stdout.split(/\r?\n/).map((line) => line.trim()).filter(Boolean); + const jsonLine = lines.find((line) => line.startsWith("{") && line.includes("\"success\"")); + if (!jsonLine) { + fail(`Native probe did not emit a JSON summary:\n${stdout}`); + } + try { + return JSON.parse(jsonLine); + } catch { + // The helper currently prints raw Windows paths with backslashes. + // Keep parsing resilient while this remains a throwaway benchmark tool. + return JSON.parse(jsonLine.replace(/,"outputPath":".*"\}$/, "}")); + } +} + +const inputPath = resolve(getArg("--input")); +const outputPath = resolve(getArg("--output", join(scriptDir, "recordly-nvdec-nvenc-mp4-output.mp4"))); +const fps = Math.round(getNumberArg("--fps", 30)); +const bitrateMbps = Math.round(getNumberArg("--bitrate-mbps", 18)); +const workDir = resolve(getArg("--work-dir", join(scriptDir, "mp4-work"))); +const reuseIntermediates = hasArg("--reuse-intermediates"); +const reuseDemux = hasArg("--reuse-demux") || reuseIntermediates; +const sampleGpuDuringEncode = hasArg("--sample-gpu"); +const gpuSampleIntervalMs = Math.round(getNumberArg("--gpu-sample-interval-ms", 1000)); +const streamSync = hasArg("--stream-sync"); +const prewarmMs = Math.round(getNumberArg("--prewarm-ms", 0)); +const maxOutputFrames = Math.round(getNumberArg("--max-output-frames", 0)); +const requestedDurationSec = getNumberArg("--duration-sec", 0); +const chunkMb = Math.round(getNumberArg("--chunk-mb", 4)); +const skipMux = hasArg("--skip-mux"); +const videoOnly = hasArg("--video-only"); +const contentX = Math.round(getNonNegativeNumberArg("--content-x", 0)); +const contentY = Math.round(getNonNegativeNumberArg("--content-y", 0)); +const contentWidth = Math.round(getNumberArg("--content-width", 0)); +const contentHeight = Math.round(getNumberArg("--content-height", 0)); +const radius = Math.round(getNonNegativeNumberArg("--radius", 0)); +const backgroundY = Math.round(getNonNegativeNumberArg("--background-y", 16)); +const backgroundU = Math.round(getNonNegativeNumberArg("--background-u", 128)); +const backgroundV = Math.round(getNonNegativeNumberArg("--background-v", 128)); +const backgroundImage = getArg("--background-image", ""); +const backgroundNv12 = getArg("--background-nv12", ""); +const shadowOffsetY = Math.round(getNonNegativeNumberArg("--shadow-offset-y", 0)); +const shadowIntensityPct = Math.round(getNonNegativeNumberArg("--shadow-intensity-pct", 0)); +const webcamInput = getArg("--webcam-input", ""); +const webcamX = Math.round(getNonNegativeNumberArg("--webcam-x", 0)); +const webcamY = Math.round(getNonNegativeNumberArg("--webcam-y", 0)); +const webcamSize = Math.round(getNumberArg("--webcam-size", 0)); +const webcamRadius = Math.round(getNonNegativeNumberArg("--webcam-radius", 0)); +const webcamMirror = hasArg("--webcam-mirror"); +const webcamStream = hasArg("--webcam-stream"); +const cursorJson = getArg("--cursor-json", ""); +const cursorHeight = Math.round(getNumberArg("--cursor-height", 0)); +const cursorStyle = getArg("--cursor-style", "vector"); +const cursorAtlasPng = getArg("--cursor-atlas-png", ""); +const cursorAtlasMetadata = getArg("--cursor-atlas-metadata", ""); +const zoomTelemetry = getArg("--zoom-telemetry", ""); + +if (!existsSync(inputPath)) { + fail(`Input does not exist: ${inputPath}`); +} +mkdirSync(workDir, { recursive: true }); +mkdirSync(dirname(outputPath), { recursive: true }); + +function resolveNativeProbePath() { + const configuredPath = process.env.RECORDLY_NVIDIA_CUDA_EXPORT_EXE; + const platformArch = process.arch === "arm64" ? "win32-arm64" : "win32-x64"; + const candidates = [ + configuredPath, + join(scriptDir, "build", "Release", "recordly-nvidia-cuda-compositor.exe"), + join(repoRoot, "electron", "native", "bin", platformArch, "recordly-nvidia-cuda-compositor.exe"), + // Backward-compatible legacy helper path while old work dirs are being retired. + join(scriptDir, "build", "Release", "recordly-nvdec-nvenc-probe.exe"), + ].filter(Boolean); + + for (const candidate of candidates) { + if (existsSync(candidate)) { + return candidate; + } + } + + fail(`Native NVIDIA CUDA compositor is not built. Checked: ${candidates.join(", ")}`); +} +const nativeProbe = resolveNativeProbePath(); + +const baseName = basename(inputPath).replace(/\.[^.]+$/, ""); +const webcamBaseName = webcamInput + ? basename(webcamInput).replace(/\.[^.]+$/, "") + : `${baseName}.webcam`; +const annexBPath = join(workDir, `${baseName}.annexb.h264`); +const webcamAnnexBPath = join(workDir, `${webcamBaseName}.annexb.h264`); +const cursorSamplesPath = join(workDir, `${baseName}.cursor.tsv`); +const encodedPath = join(workDir, `${baseName}.mapped-callback.h264`); +const shouldBakeStaticShadow = + Boolean(backgroundImage) && + contentWidth > 0 && + contentHeight > 0 && + shadowOffsetY > 0 && + shadowIntensityPct > 0; +const backgroundSuffix = shouldBakeStaticShadow + ? `.shadow-${shadowOffsetY}-${shadowIntensityPct}` + : ""; +const generatedBackgroundNv12Path = join( + workDir, + `${baseName}${backgroundSuffix}.background.nv12`, +); +const generatedWebcamNv12Path = join( + workDir, + `${baseName}.webcam-${webcamSize}${webcamMirror ? "-mirror" : ""}.nv12`, +); +const sourcePtsPath = join(workDir, `${baseName}.source-pts.csv`); + +const videoInfo = getVideoInfo(inputPath); +const webcamInfo = webcamInput ? getVideoInfo(webcamInput) : null; +const durationSec = + requestedDurationSec > 0 ? Math.min(videoInfo.durationSec, requestedDurationSec) : videoInfo.durationSec; +const targetFrames = Math.ceil(durationSec * fps); +emitPreparationProgress(targetFrames, 1); +let sourceWindowFrames = Math.max( + 1, + Math.min(videoInfo.sourceFrames, Math.ceil((videoInfo.sourceFrames * durationSec) / videoInfo.durationSec)), +); +let webcamSourceWindowFrames = webcamInfo + ? Math.max( + 1, + Math.min(webcamInfo.sourceFrames, Math.ceil((webcamInfo.sourceFrames * durationSec) / webcamInfo.durationSec)), + ) + : 0; +const backgroundNv12Path = backgroundImage + ? generatedBackgroundNv12Path + : backgroundNv12; +const backgroundFilter = createBackgroundFilter( + videoInfo, + shouldBakeStaticShadow + ? { + x: contentX, + y: contentY + shadowOffsetY, + width: contentWidth, + height: contentHeight, + radius: radius + 8, + intensityPct: shadowIntensityPct, + } + : null, +); + +const backgroundConvert = + backgroundImage && !(reuseIntermediates && existsSync(backgroundNv12Path)) + ? run("ffmpeg", [ + "-y", + "-hide_banner", + "-loglevel", + "error", + "-i", + resolve(backgroundImage), + ...backgroundFilter.filterArgs, + "-frames:v", + "1", + "-f", + "rawvideo", + backgroundNv12Path, + ]) + : { elapsedMs: 0 }; + +const webcamConvert = + webcamInput && !webcamStream && webcamSize > 0 && !(reuseIntermediates && existsSync(generatedWebcamNv12Path)) + ? run("ffmpeg", [ + "-y", + "-hide_banner", + "-loglevel", + "error", + "-i", + resolve(webcamInput), + "-vf", + `${webcamMirror ? "hflip," : ""}scale=${webcamSize}:${webcamSize}:force_original_aspect_ratio=increase,crop=${webcamSize}:${webcamSize},format=nv12`, + "-frames:v", + "1", + "-f", + "rawvideo", + generatedWebcamNv12Path, + ]) + : { elapsedMs: 0 }; + +const webcamDemux = + webcamInput && webcamStream && !(reuseDemux && existsSync(webcamAnnexBPath)) + ? run("ffmpeg", [ + "-y", + "-hide_banner", + "-loglevel", + "error", + "-stats", + "-i", + resolve(webcamInput), + "-t", + String(durationSec), + "-map", + "0:v:0", + "-c:v", + "copy", + "-bsf:v", + "h264_mp4toannexb", + "-an", + webcamAnnexBPath, + ]) + : { elapsedMs: 0 }; + +if (cursorJson) { + const cursorPayload = JSON.parse(readFileSync(resolve(cursorJson), "utf8")); + writeCursorSamples(cursorPayload, cursorSamplesPath); +} +const cursorAtlas = + cursorJson && cursorHeight > 0 && cursorAtlasPng && cursorAtlasMetadata + ? prepareExternalCursorAtlas(workDir, cursorAtlasPng, cursorAtlasMetadata) + : cursorJson && cursorHeight > 0 && cursorStyle === "tahoe" + ? renderTahoeCursorAtlas(workDir) + : null; + +const demux = reuseDemux && existsSync(annexBPath) + ? { elapsedMs: 0 } + : run("ffmpeg", [ + "-y", + "-hide_banner", + "-loglevel", + "error", + "-stats", + "-i", + inputPath, + "-t", + String(durationSec), + "-map", + "0:v:0", + "-c:v", + "copy", + "-bsf:v", + "h264_mp4toannexb", + "-an", + annexBPath, + ]); +const demuxFrameCount = parseFfmpegStatsFrameCount(demux.stderr); +if (demuxFrameCount > 0) { + sourceWindowFrames = demuxFrameCount; +} +const webcamDemuxFrameCount = parseFfmpegStatsFrameCount(webcamDemux.stderr); +if (webcamDemuxFrameCount > 0) { + webcamSourceWindowFrames = webcamDemuxFrameCount; +} +emitPreparationProgress(targetFrames, 2); +const sourcePts = writeFramePtsSidecar(inputPath, durationSec, sourcePtsPath); +emitPreparationProgress(targetFrames, 3); + +const encodeArgs = [ + "--input", + annexBPath, + "--output", + encodedPath, + "--fps", + String(fps), + "--input-frames", + String(sourceWindowFrames), + "--target-frames", + String(targetFrames), + "--bitrate-mbps", + String(bitrateMbps), + "--callback-encode", + "--chunk-mb", + String(chunkMb), +]; +if (sourcePts.path && sourcePts.frames >= sourceWindowFrames) { + encodeArgs.push("--source-pts", sourcePts.path); +} +if (maxOutputFrames > 0) { + encodeArgs.push("--max-frames", String(maxOutputFrames)); +} +if (cursorJson && cursorHeight > 0) { + encodeArgs.push("--cursor-samples", cursorSamplesPath, "--cursor-height", String(cursorHeight)); + if (cursorAtlas) { + encodeArgs.push( + "--cursor-atlas-rgba", + cursorAtlas.rgbaPath, + "--cursor-atlas-metadata", + cursorAtlas.metadataPath, + "--cursor-atlas-width", + String(cursorAtlas.width), + "--cursor-atlas-height", + String(cursorAtlas.height), + ); + } +} +if (streamSync) { + encodeArgs.push("--stream-sync"); +} +if (prewarmMs > 0) { + encodeArgs.push("--prewarm-ms", String(prewarmMs)); +} +if (contentWidth > 0 && contentHeight > 0) { + encodeArgs.push( + "--content-x", + String(contentX), + "--content-y", + String(contentY), + "--content-width", + String(contentWidth), + "--content-height", + String(contentHeight), + "--radius", + String(radius), + "--background-y", + String(backgroundY), + "--background-u", + String(backgroundU), + "--background-v", + String(backgroundV), + ); + if (backgroundNv12Path) { + encodeArgs.push("--background-nv12", backgroundNv12Path); + } + if (!shouldBakeStaticShadow && shadowOffsetY > 0 && shadowIntensityPct > 0) { + encodeArgs.push( + "--shadow-offset-y", + String(shadowOffsetY), + "--shadow-intensity-pct", + String(shadowIntensityPct), + ); + } + if (webcamInput && webcamSize > 0) { + encodeArgs.push( + "--webcam-x", + String(webcamX), + "--webcam-y", + String(webcamY), + "--webcam-size", + String(webcamSize), + "--webcam-radius", + String(webcamRadius), + ); + if (webcamMirror) { + encodeArgs.push("--webcam-mirror"); + } + if (webcamStream) { + encodeArgs.push( + "--webcam-annexb", + webcamAnnexBPath, + "--webcam-input-frames", + String(webcamSourceWindowFrames), + "--webcam-target-frames", + String(targetFrames), + "--webcam-source-width", + String(webcamInfo.width), + "--webcam-source-height", + String(webcamInfo.height), + ); + } else { + encodeArgs.push("--webcam-nv12", generatedWebcamNv12Path); + } + } +} +if (zoomTelemetry) { + encodeArgs.push("--zoom-samples", resolve(zoomTelemetry)); +} +const encode = reuseIntermediates && existsSync(encodedPath) + ? { elapsedMs: 0, stdout: "", gpuSummary: null } + : await runWithGpuMonitor( + nativeProbe, + encodeArgs, + sampleGpuDuringEncode ? gpuSampleIntervalMs : 0, + ); +const nativeSummary = encode.stdout ? parseProbeSummary(encode.stdout) : null; + +const mux = skipMux + ? { elapsedMs: 0 } + : videoOnly + ? run("ffmpeg", [ + "-y", + "-hide_banner", + "-loglevel", + "error", + "-framerate", + String(fps), + "-i", + encodedPath, + "-map", + "0:v:0", + "-c:v", + "copy", + outputPath, + ]) + : run("ffmpeg", [ + "-y", + "-hide_banner", + "-loglevel", + "error", + "-framerate", + String(fps), + "-i", + encodedPath, + "-i", + inputPath, + "-map", + "0:v:0", + "-map", + "1:a?", + "-c:v", + "copy", + "-c:a", + "copy", + "-t", + String(durationSec), + outputPath, +]); + +const outputInfo = skipMux + ? { streams: [] } + : ffprobeJson([ + "-show_entries", + "stream=index,codec_type,codec_name,width,height,duration,avg_frame_rate,nb_frames", + outputPath, + ]); +const outputStreams = outputInfo.streams ?? []; +const outputVideo = outputStreams.find((stream) => stream.codec_type === "video") ?? null; +const outputAudio = outputStreams.find((stream) => stream.codec_type === "audio") ?? null; + +console.log( + JSON.stringify( + { + success: true, + inputPath, + outputPath, + fps, + bitrateMbps, + streamSync, + prewarmMs, + maxOutputFrames, + chunkMb, + skipMux, + videoOnly, + durationSec, + staticLayout: + contentWidth > 0 && contentHeight > 0 + ? { + contentX, + contentY, + contentWidth, + contentHeight, + radius, + backgroundY, + backgroundU, + backgroundV, + shadowOffsetY, + shadowIntensityPct, + shadowBakedIntoBackground: shouldBakeStaticShadow, + backgroundShadowAlpha: backgroundFilter.shadowAlpha ?? null, + backgroundShadowBlur: backgroundFilter.shadowBlur ?? null, + webcam: + webcamInput && webcamSize > 0 + ? { + inputPath: resolve(webcamInput), + x: webcamX, + y: webcamY, + size: webcamSize, + radius: webcamRadius, + mirror: webcamMirror, + staticFrameOnly: !webcamStream, + stream: webcamStream, + } + : null, + cursor: + cursorJson && cursorHeight > 0 + ? { + inputPath: resolve(cursorJson), + height: cursorHeight, + style: cursorStyle, + atlas: cursorAtlas + ? { + width: cursorAtlas.width, + height: cursorAtlas.height, + entries: cursorAtlas.entries, + } + : null, + } + : null, + zoom: zoomTelemetry + ? { + inputPath: resolve(zoomTelemetry), + } + : null, + } + : null, + gpuSampleIntervalMs: sampleGpuDuringEncode ? gpuSampleIntervalMs : null, + videoInfo, + sourceWindowFrames, + sourcePtsFrames: sourcePts.frames, + sourcePtsSource: sourcePts.source, + targetFrames, + timingsMs: { + demux: Number(demux.elapsedMs.toFixed(2)), + backgroundConvert: Number(backgroundConvert.elapsedMs.toFixed(2)), + cursorAtlas: Number((cursorAtlas?.elapsedMs ?? 0).toFixed(2)), + webcamConvert: Number(webcamConvert.elapsedMs.toFixed(2)), + webcamDemux: Number(webcamDemux.elapsedMs.toFixed(2)), + sourcePtsProbe: Number(sourcePts.elapsedMs.toFixed(2)), + nativeEncode: Number(encode.elapsedMs.toFixed(2)), + mux: Number(mux.elapsedMs.toFixed(2)), + endToEnd: Number( + ( + backgroundConvert.elapsedMs + + (cursorAtlas?.elapsedMs ?? 0) + + webcamConvert.elapsedMs + + webcamDemux.elapsedMs + + demux.elapsedMs + + sourcePts.elapsedMs + + encode.elapsedMs + + mux.elapsedMs + ).toFixed(2), + ), + }, + nativeSummary, + nativeProcessPriorityBoosted: encode.priorityBoosted ?? false, + gpuSamples: encode.gpuSamples ?? [], + gpuSummary: encode.gpuSummary ?? null, + outputVideo, + outputAudio, + outputStreams, + }, + null, + 2, + ), +); diff --git a/electron/native/nvidia-cuda-compositor/src/main.cu b/electron/native/nvidia-cuda-compositor/src/main.cu new file mode 100644 index 00000000..385dc18b --- /dev/null +++ b/electron/native/nvidia-cuda-compositor/src/main.cu @@ -0,0 +1,3006 @@ +#include +#include + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include "NvDecoder/NvDecoder.h" +#include "NvEncoder/NvEncoderCuda.h" +#include "Utils/Logger.h" + +simplelogger::Logger* logger = simplelogger::LoggerFactory::CreateConsoleLogger(ERROR); + +namespace { + +struct Options { + std::string inputPath; + std::string outputPath = "recordly-nvidia-cuda-compositor.h264"; + std::string sourcePtsPath; + int fps = 30; + int maxFrames = 0; + int inputFrames = 0; + int targetFrames = 0; + int bitrateMbps = 18; + bool postSelect = false; + bool callbackEncode = false; + bool streamSync = false; + int prewarmMs = 0; + int chunkMb = 4; + int contentX = 0; + int contentY = 0; + int contentWidth = 0; + int contentHeight = 0; + int radius = 0; + int backgroundY = 16; + int backgroundU = 128; + int backgroundV = 128; + std::string backgroundNv12Path; + int shadowOffsetY = 0; + int shadowIntensityPct = 0; + std::string webcamNv12Path; + std::string webcamAnnexbPath; + int webcamInputFrames = 0; + int webcamTargetFrames = 0; + int webcamSourceWidth = 0; + int webcamSourceHeight = 0; + int webcamX = 0; + int webcamY = 0; + int webcamSize = 0; + int webcamRadius = 0; + bool webcamMirror = false; + std::string cursorSamplesPath; + int cursorHeight = 0; + std::string cursorAtlasRgbaPath; + std::string cursorAtlasMetadataPath; + int cursorAtlasWidth = 0; + int cursorAtlasHeight = 0; + std::string zoomSamplesPath; +}; + +constexpr int kMaxCursorAtlasEntries = 16; +constexpr int kWebcamPrefetchOutputFrames = 900; + +[[noreturn]] void fail(const std::string& message) { + throw std::runtime_error(message); +} + +void checkCuda(cudaError_t status, const char* expression) { + if (status != cudaSuccess) { + std::ostringstream stream; + stream << expression << " failed: " << cudaGetErrorString(status); + fail(stream.str()); + } +} + +void checkCu(CUresult status, const char* expression) { + if (status != CUDA_SUCCESS) { + const char* name = nullptr; + const char* message = nullptr; + cuGetErrorName(status, &name); + cuGetErrorString(status, &message); + std::ostringstream stream; + stream << expression << " failed: " << (name ? name : "CUDA_ERROR") + << " (" << (message ? message : "no detail") << ")"; + fail(stream.str()); + } +} + +int parsePositiveInt(const char* value, const char* name) { + char* end = nullptr; + const long parsed = std::strtol(value, &end, 10); + if (!end || *end != '\0' || parsed <= 0 || parsed > 1000000) { + std::ostringstream stream; + stream << "Invalid " << name << ": " << value; + fail(stream.str()); + } + return static_cast(parsed); +} + +int parseNonNegativeInt(const char* value, const char* name) { + char* end = nullptr; + const long parsed = std::strtol(value, &end, 10); + if (!end || *end != '\0' || parsed < 0 || parsed > 1000000) { + std::ostringstream stream; + stream << "Invalid " << name << ": " << value; + fail(stream.str()); + } + return static_cast(parsed); +} + +Options parseOptions(int argc, char** argv) { + Options options; + for (int index = 1; index < argc; ++index) { + const std::string arg = argv[index]; + auto requireValue = [&](const char* name) -> const char* { + if (index + 1 >= argc) { + std::ostringstream stream; + stream << "Missing value for " << name; + fail(stream.str()); + } + return argv[++index]; + }; + + if (arg == "--input") { + options.inputPath = requireValue("--input"); + } else if (arg == "--output") { + options.outputPath = requireValue("--output"); + } else if (arg == "--source-pts") { + options.sourcePtsPath = requireValue("--source-pts"); + } else if (arg == "--fps") { + options.fps = parsePositiveInt(requireValue("--fps"), "--fps"); + } else if (arg == "--max-frames") { + options.maxFrames = parsePositiveInt(requireValue("--max-frames"), "--max-frames"); + } else if (arg == "--input-frames") { + options.inputFrames = parsePositiveInt(requireValue("--input-frames"), "--input-frames"); + } else if (arg == "--target-frames") { + options.targetFrames = parsePositiveInt(requireValue("--target-frames"), "--target-frames"); + } else if (arg == "--bitrate-mbps") { + options.bitrateMbps = parsePositiveInt(requireValue("--bitrate-mbps"), "--bitrate-mbps"); + } else if (arg == "--post-select") { + options.postSelect = true; + } else if (arg == "--callback-encode") { + options.callbackEncode = true; + } else if (arg == "--stream-sync") { + options.streamSync = true; + } else if (arg == "--prewarm-ms") { + options.prewarmMs = parsePositiveInt(requireValue("--prewarm-ms"), "--prewarm-ms"); + } else if (arg == "--chunk-mb") { + options.chunkMb = parsePositiveInt(requireValue("--chunk-mb"), "--chunk-mb"); + } else if (arg == "--content-x") { + options.contentX = parseNonNegativeInt(requireValue("--content-x"), "--content-x"); + } else if (arg == "--content-y") { + options.contentY = parseNonNegativeInt(requireValue("--content-y"), "--content-y"); + } else if (arg == "--content-width") { + options.contentWidth = parsePositiveInt(requireValue("--content-width"), "--content-width"); + } else if (arg == "--content-height") { + options.contentHeight = parsePositiveInt(requireValue("--content-height"), "--content-height"); + } else if (arg == "--radius") { + options.radius = parseNonNegativeInt(requireValue("--radius"), "--radius"); + } else if (arg == "--background-y") { + options.backgroundY = parseNonNegativeInt(requireValue("--background-y"), "--background-y"); + } else if (arg == "--background-u") { + options.backgroundU = parseNonNegativeInt(requireValue("--background-u"), "--background-u"); + } else if (arg == "--background-v") { + options.backgroundV = parseNonNegativeInt(requireValue("--background-v"), "--background-v"); + } else if (arg == "--background-nv12") { + options.backgroundNv12Path = requireValue("--background-nv12"); + } else if (arg == "--shadow-offset-y") { + options.shadowOffsetY = parseNonNegativeInt(requireValue("--shadow-offset-y"), "--shadow-offset-y"); + } else if (arg == "--shadow-intensity-pct") { + options.shadowIntensityPct = parseNonNegativeInt(requireValue("--shadow-intensity-pct"), "--shadow-intensity-pct"); + } else if (arg == "--webcam-nv12") { + options.webcamNv12Path = requireValue("--webcam-nv12"); + } else if (arg == "--webcam-annexb") { + options.webcamAnnexbPath = requireValue("--webcam-annexb"); + } else if (arg == "--webcam-input-frames") { + options.webcamInputFrames = parsePositiveInt(requireValue("--webcam-input-frames"), "--webcam-input-frames"); + } else if (arg == "--webcam-target-frames") { + options.webcamTargetFrames = + parsePositiveInt(requireValue("--webcam-target-frames"), "--webcam-target-frames"); + } else if (arg == "--webcam-source-width") { + options.webcamSourceWidth = parsePositiveInt(requireValue("--webcam-source-width"), "--webcam-source-width"); + } else if (arg == "--webcam-source-height") { + options.webcamSourceHeight = + parsePositiveInt(requireValue("--webcam-source-height"), "--webcam-source-height"); + } else if (arg == "--webcam-x") { + options.webcamX = parseNonNegativeInt(requireValue("--webcam-x"), "--webcam-x"); + } else if (arg == "--webcam-y") { + options.webcamY = parseNonNegativeInt(requireValue("--webcam-y"), "--webcam-y"); + } else if (arg == "--webcam-size") { + options.webcamSize = parsePositiveInt(requireValue("--webcam-size"), "--webcam-size"); + } else if (arg == "--webcam-radius") { + options.webcamRadius = parseNonNegativeInt(requireValue("--webcam-radius"), "--webcam-radius"); + } else if (arg == "--webcam-mirror") { + options.webcamMirror = true; + } else if (arg == "--cursor-samples") { + options.cursorSamplesPath = requireValue("--cursor-samples"); + } else if (arg == "--cursor-height") { + options.cursorHeight = parsePositiveInt(requireValue("--cursor-height"), "--cursor-height"); + } else if (arg == "--cursor-atlas-rgba") { + options.cursorAtlasRgbaPath = requireValue("--cursor-atlas-rgba"); + } else if (arg == "--cursor-atlas-metadata") { + options.cursorAtlasMetadataPath = requireValue("--cursor-atlas-metadata"); + } else if (arg == "--cursor-atlas-width") { + options.cursorAtlasWidth = parsePositiveInt(requireValue("--cursor-atlas-width"), "--cursor-atlas-width"); + } else if (arg == "--cursor-atlas-height") { + options.cursorAtlasHeight = + parsePositiveInt(requireValue("--cursor-atlas-height"), "--cursor-atlas-height"); + } else if (arg == "--zoom-samples") { + options.zoomSamplesPath = requireValue("--zoom-samples"); + } else if (arg == "--help") { + std::cout << "Usage: recordly-nvidia-cuda-compositor --input input.annexb.h264 " + "[--output out.h264] [--source-pts source-pts.csv] [--fps 30] " + "[--max-frames N] [--bitrate-mbps N] " + "[--post-select] [--callback-encode] [--stream-sync] [--prewarm-ms N] [--chunk-mb N] " + "[--content-x N --content-y N --content-width N --content-height N --radius N] " + "[--background-nv12 background.nv12] [--shadow-offset-y N --shadow-intensity-pct N] " + "[--webcam-nv12 webcam.nv12 --webcam-x N --webcam-y N --webcam-size N --webcam-radius N] " + "[--webcam-annexb webcam.h264 --webcam-input-frames N --webcam-target-frames N] " + "[--cursor-samples cursor.tsv --cursor-height N] " + "[--cursor-atlas-rgba cursor.rgba --cursor-atlas-metadata cursor.tsv " + "--cursor-atlas-width N --cursor-atlas-height N] " + "[--zoom-samples zoom.csv]\n"; + std::exit(0); + } else { + std::ostringstream stream; + stream << "Unknown argument: " << arg; + fail(stream.str()); + } + } + if (options.inputPath.empty()) { + fail("--input is required"); + } + return options; +} + +bool shouldEncodeFrame(int sourceFrameIndex, int encodedFrames, const Options& options) { + if (options.inputFrames <= 0 || options.targetFrames <= 0) { + return true; + } + if (encodedFrames >= options.targetFrames) { + return false; + } + + const int expectedEncodedFrames = + ((sourceFrameIndex + 1) * options.targetFrames + options.inputFrames - 1) / options.inputFrames; + return encodedFrames < expectedEncodedFrames; +} + +bool hasStaticLayout(const Options& options) { + return options.contentWidth > 0 && options.contentHeight > 0; +} + +bool hasWebcamOverlay(const Options& options) { + return (!options.webcamNv12Path.empty() || !options.webcamAnnexbPath.empty()) && options.webcamSize > 0; +} + +int webcamFrameIndexForOutputFrame(int outputFrameIndex, const Options& options) { + if (options.webcamInputFrames > 0 && options.webcamTargetFrames > 0) { + return static_cast( + (static_cast(outputFrameIndex) * options.webcamInputFrames) / options.webcamTargetFrames); + } + return outputFrameIndex; +} + +std::vector loadFramePts(const std::string& path) { + std::vector timestamps; + if (path.empty()) { + return timestamps; + } + + std::ifstream input(path); + if (!input) { + fail("Failed to open source PTS sidecar: " + path); + } + + std::string line; + double lastTimestamp = -std::numeric_limits::infinity(); + while (std::getline(input, line)) { + if (line.empty()) { + continue; + } + char* end = nullptr; + const double timestamp = std::strtod(line.c_str(), &end); + if (!end || *end != '\0' || !std::isfinite(timestamp) || timestamp < lastTimestamp) { + fail("Invalid source PTS sidecar entry: " + line); + } + timestamps.push_back(timestamp); + lastTimestamp = timestamp; + } + + return timestamps; +} + +int maxSelectedFramesForTimeline(int targetFrames, int maxFrames) { + if (targetFrames <= 0) { + return maxFrames > 0 ? maxFrames : std::numeric_limits::max(); + } + return maxFrames > 0 ? std::min(maxFrames, targetFrames) : targetFrames; +} + +int expectedOutputFramesForSourceFrame( + int sourceFrameIndex, + int inputFrames, + int targetFrames, + int maxFrames, + int fps, + const std::vector* sourcePts) { + const int maxOutputFrames = maxSelectedFramesForTimeline(targetFrames, maxFrames); + if (sourcePts && sourceFrameIndex >= 0 && sourceFrameIndex < static_cast(sourcePts->size())) { + if (inputFrames > 0 && sourceFrameIndex + 1 >= inputFrames) { + return maxOutputFrames; + } + const double frameTimeSec = std::max(0.0, (*sourcePts)[sourceFrameIndex]); + const int64_t expected = static_cast(std::floor(frameTimeSec * fps)) + 1; + return static_cast(std::min(std::max(1, expected), maxOutputFrames)); + } + if (inputFrames <= 0 || targetFrames <= 0) { + return maxOutputFrames; + } + + const int64_t expected = + (static_cast(sourceFrameIndex + 1) * targetFrames + inputFrames - 1) / inputFrames; + return static_cast(std::min(expected, maxOutputFrames)); +} + +unsigned char clampByte(int value) { + return static_cast(std::max(0, std::min(255, value))); +} + +struct FrameSelectionState { + int inputFrames = 0; + int targetFrames = 0; + int maxFrames = 0; + int sourceFrames = 0; + int selectedFrames = 0; + int fps = 30; + const std::vector* sourcePts = nullptr; +}; + +bool shouldCopyDisplayFrame(int displayFrameIndex, void* userData) { + auto* state = static_cast(userData); + state->sourceFrames = displayFrameIndex + 1; + + const int maxSelectedFrames = + state->maxFrames > 0 ? std::min(state->maxFrames, state->targetFrames) : state->targetFrames; + if (state->selectedFrames >= maxSelectedFrames) { + return false; + } + + // maxFrames is a smoke-test stop cap; it must not spread the sampled frames + // across the full source because that hides the true first-window performance. + const int expectedSelectedFrames = expectedOutputFramesForSourceFrame( + displayFrameIndex, + state->inputFrames, + state->targetFrames, + state->maxFrames, + state->fps, + state->sourcePts); + if (state->selectedFrames >= expectedSelectedFrames) { + return false; + } + + ++state->selectedFrames; + return true; +} + +double elapsedMs(std::chrono::steady_clock::time_point start, std::chrono::steady_clock::time_point end); +struct ProgressCounters { + double decodeWallMs = 0.0; + double encodeMs = 0.0; + double compositeMs = 0.0; + double nvencMs = 0.0; + double packetWriteMs = 0.0; + double webcamDecodeMs = 0.0; + double webcamCopyMs = 0.0; + int roiCompositeFrames = 0; + int monolithicCompositeFrames = 0; + int copyCompositeFrames = 0; +}; + +struct ProgressReportState { + std::chrono::steady_clock::time_point startedAt; + std::chrono::steady_clock::time_point lastReportAt; + int lastReportedFrame = 0; + ProgressCounters lastCounters; +}; + +void reportEncodingProgress( + int encodedFrames, + int totalFrames, + ProgressReportState& state, + const ProgressCounters& counters, + bool force = false); + +struct WebcamFrameCache { + std::vector frames; + double decodeMs = 0.0; + double copyMs = 0.0; + int sourceFrames = 0; + int baseFrameIndex = 0; + int decodedFrames = 0; + int peakFrames = 0; + int width = 0; + int height = 0; + + ~WebcamFrameCache() { + for (unsigned char* frame : frames) { + cudaFree(frame); + } + } + + void pushFrame(unsigned char* frame) { + frames.push_back(frame); + decodedFrames = baseFrameIndex + static_cast(frames.size()); + sourceFrames = decodedFrames; + peakFrames = std::max(peakFrames, static_cast(frames.size())); + } + + void dropBefore(int minFrameIndex) { + const int dropCount = std::min( + std::max(0, minFrameIndex - baseFrameIndex), + std::max(0, static_cast(frames.size()) - 1)); + if (dropCount <= 0) { + return; + } + for (int index = 0; index < dropCount; ++index) { + cudaFree(frames[index]); + } + frames.erase(frames.begin(), frames.begin() + dropCount); + baseFrameIndex += dropCount; + } + + const unsigned char* frameAt(int frameIndex) const { + if (frames.empty()) { + return nullptr; + } + const int clampedFrameIndex = + std::max(baseFrameIndex, std::min(frameIndex, baseFrameIndex + static_cast(frames.size()) - 1)); + return frames[clampedFrameIndex - baseFrameIndex]; + } +}; + +struct CursorSample { + double timeMs = 0.0; + double cx = 0.0; + double cy = 0.0; + int typeIndex = 0; + double bounceScale = 1.0; +}; + +struct CursorPosition { + bool visible = false; + double cx = 0.0; + double cy = 0.0; + int typeIndex = 0; + double bounceScale = 1.0; +}; + +struct CursorTrack { + std::vector samples; + + CursorPosition positionAt(double timeMs) const { + if (samples.empty()) { + return {}; + } + if (timeMs <= samples.front().timeMs) { + return {true, samples.front().cx, samples.front().cy, samples.front().typeIndex, samples.front().bounceScale}; + } + if (timeMs >= samples.back().timeMs) { + return {true, samples.back().cx, samples.back().cy, samples.back().typeIndex, samples.back().bounceScale}; + } + + int low = 0; + int high = static_cast(samples.size()) - 1; + while (low < high - 1) { + const int mid = (low + high) / 2; + if (samples[mid].timeMs <= timeMs) { + low = mid; + } else { + high = mid; + } + } + + const CursorSample& left = samples[low]; + const CursorSample& right = samples[high]; + const double span = right.timeMs - left.timeMs; + if (span <= 0.0) { + return {true, left.cx, left.cy, left.typeIndex, left.bounceScale}; + } + + const double t = (timeMs - left.timeMs) / span; + return { + true, + left.cx + (right.cx - left.cx) * t, + left.cy + (right.cy - left.cy) * t, + t < 0.5 ? left.typeIndex : right.typeIndex, + left.bounceScale + (right.bounceScale - left.bounceScale) * t, + }; + } +}; + +std::unique_ptr loadCursorTrack(const Options& options) { + if (options.cursorSamplesPath.empty()) { + return nullptr; + } + if (options.cursorHeight <= 0) { + fail("--cursor-height is required with --cursor-samples"); + } + + std::ifstream input(options.cursorSamplesPath); + if (!input) { + fail("Failed to open cursor samples: " + options.cursorSamplesPath); + } + + auto track = std::make_unique(); + std::string line; + while (std::getline(input, line)) { + if (line.empty()) { + continue; + } + std::istringstream row(line); + CursorSample sample; + if (!(row >> sample.timeMs >> sample.cx >> sample.cy)) { + continue; + } + if (!(row >> sample.typeIndex)) { + sample.typeIndex = 0; + } + if (!(row >> sample.bounceScale)) { + sample.bounceScale = 1.0; + } + if (sample.cx < -1.0 || sample.cx > 2.0 || sample.cy < -1.0 || sample.cy > 2.0) { + continue; + } + sample.typeIndex = std::max(0, std::min(kMaxCursorAtlasEntries - 1, sample.typeIndex)); + sample.bounceScale = std::max(0.5, std::min(2.0, sample.bounceScale)); + track->samples.push_back(sample); + } + if (track->samples.empty()) { + fail("No cursor samples were loaded: " + options.cursorSamplesPath); + } + return track; +} + +struct ZoomSample { + double timeMs = 0.0; + double scale = 1.0; + double x = 0.0; + double y = 0.0; +}; + +struct ZoomTrack { + std::vector samples; + + ZoomSample sampleAt(double timeMs) const { + if (samples.empty()) { + return {}; + } + if (timeMs <= samples.front().timeMs) { + return samples.front(); + } + if (timeMs >= samples.back().timeMs) { + return samples.back(); + } + + int low = 0; + int high = static_cast(samples.size()) - 1; + while (low < high - 1) { + const int mid = (low + high) / 2; + if (samples[mid].timeMs <= timeMs) { + low = mid; + } else { + high = mid; + } + } + + const ZoomSample& left = samples[low]; + const ZoomSample& right = samples[high]; + const double span = right.timeMs - left.timeMs; + if (span <= 0.0) { + return left; + } + + const double t = (timeMs - left.timeMs) / span; + return { + timeMs, + left.scale + (right.scale - left.scale) * t, + left.x + (right.x - left.x) * t, + left.y + (right.y - left.y) * t, + }; + } +}; + +std::unique_ptr loadZoomTrack(const Options& options) { + if (options.zoomSamplesPath.empty()) { + return nullptr; + } + + std::ifstream input(options.zoomSamplesPath); + if (!input) { + fail("Failed to open zoom samples: " + options.zoomSamplesPath); + } + + auto track = std::make_unique(); + std::string line; + while (std::getline(input, line)) { + if (line.empty()) { + continue; + } + std::replace(line.begin(), line.end(), ',', ' '); + std::istringstream row(line); + ZoomSample sample; + if (!(row >> sample.timeMs >> sample.scale >> sample.x >> sample.y)) { + continue; + } + if (!std::isfinite(sample.timeMs) || !std::isfinite(sample.scale) || + !std::isfinite(sample.x) || !std::isfinite(sample.y)) { + continue; + } + sample.timeMs = std::max(0.0, sample.timeMs); + sample.scale = std::max(0.01, sample.scale); + track->samples.push_back(sample); + } + if (track->samples.empty()) { + fail("No zoom samples were loaded: " + options.zoomSamplesPath); + } + std::sort(track->samples.begin(), track->samples.end(), [](const auto& left, const auto& right) { + return left.timeMs < right.timeMs; + }); + return track; +} + +struct CursorAtlasEntry { + int x = 0; + int y = 0; + int width = 0; + int height = 0; + double anchorX = 0.0; + double anchorY = 0.0; + double aspectRatio = 0.0; + bool valid = false; +}; + +struct WebcamCacheState { + WebcamFrameCache* cache = nullptr; +}; + +void cacheMappedWebcamFrame( + CUdeviceptr dpSrcFrame, + unsigned int nSrcPitch, + int width, + int height, + int surfaceHeight, + int64_t, + void* userData) { + auto* state = static_cast(userData); + if (state->cache->width == 0) { + state->cache->width = width; + state->cache->height = height; + } + if (width != state->cache->width || height != state->cache->height) { + std::ostringstream stream; + stream << "Decoded webcam frame size changed from " << state->cache->width << "x" << state->cache->height + << " to " << width << "x" << height; + fail(stream.str()); + } + + const auto copyStart = std::chrono::steady_clock::now(); + const size_t expectedBytes = static_cast(width) * static_cast(height) * 3 / 2; + unsigned char* frame = nullptr; + checkCuda(cudaMalloc(&frame, expectedBytes), "cudaMalloc webcam cached frame"); + + CUDA_MEMCPY2D copy = {}; + copy.srcMemoryType = CU_MEMORYTYPE_DEVICE; + copy.srcDevice = dpSrcFrame; + copy.srcPitch = nSrcPitch; + copy.dstMemoryType = CU_MEMORYTYPE_DEVICE; + copy.dstDevice = reinterpret_cast(frame); + copy.dstPitch = width; + copy.WidthInBytes = width; + copy.Height = height; + checkCu(cuMemcpy2D(©), "cuMemcpy2D webcam luma"); + + copy.srcDevice = dpSrcFrame + nSrcPitch * surfaceHeight; + copy.dstDevice = reinterpret_cast(frame + width * height); + copy.Height = height / 2; + checkCu(cuMemcpy2D(©), "cuMemcpy2D webcam chroma"); + + state->cache->pushFrame(frame); + const auto copyEnd = std::chrono::steady_clock::now(); + state->cache->copyMs += elapsedMs(copyStart, copyEnd); +} + +class WebcamStreamDecoder { +public: + WebcamStreamDecoder(CUcontext context, const Options& options) + : options_(options), + chunk_(static_cast(options.chunkMb) * 1024 * 1024) { + if (options_.webcamInputFrames <= 0 || options_.webcamTargetFrames <= 0) { + fail("--webcam-input-frames and --webcam-target-frames are required with --webcam-annexb"); + } + if (options_.webcamSourceWidth <= 0 || options_.webcamSourceHeight <= 0) { + fail("--webcam-source-width and --webcam-source-height are required with --webcam-annexb"); + } + + const int cropSide = std::min(options_.webcamSourceWidth, options_.webcamSourceHeight) & ~1; + const int cropLeft = ((options_.webcamSourceWidth - cropSide) / 2) & ~1; + const int cropTop = ((options_.webcamSourceHeight - cropSide) / 2) & ~1; + crop_ = Rect{cropLeft, cropTop, cropLeft + cropSide, cropTop + cropSide}; + + cacheState_.cache = &cache_; + decoder_ = + std::make_unique(context, 0, 0, true, cudaVideoCodec_H264, nullptr, true, true, &crop_, nullptr); + decoder_->SetMappedFrameHandler(cacheMappedWebcamFrame, &cacheState_); + + input_.open(options_.webcamAnnexbPath, std::ios::binary); + if (!input_) { + fail("Failed to open webcam input: " + options_.webcamAnnexbPath); + } + } + + WebcamFrameCache* cache() { + return &cache_; + } + + void ensureFrame(int frameIndex) { + if (frameIndex < 0) { + return; + } + while (!flushed_ && cache_.decodedFrames <= frameIndex) { + input_.read(reinterpret_cast(chunk_.data()), static_cast(chunk_.size())); + const int bytesRead = static_cast(input_.gcount()); + const auto decodeStart = std::chrono::steady_clock::now(); + if (bytesRead > 0) { + decoder_->Decode(chunk_.data(), bytesRead, &frames_, &returnedFrames_); + } else { + decoder_->Decode(nullptr, 0, &frames_, &returnedFrames_); + flushed_ = true; + } + const auto decodeEnd = std::chrono::steady_clock::now(); + cache_.decodeMs += elapsedMs(decodeStart, decodeEnd); + } + if (cache_.frames.empty()) { + fail("No webcam frames were decoded"); + } + } + + void dropBefore(int frameIndex) { + cache_.dropBefore(frameIndex); + } + +private: + const Options& options_; + WebcamFrameCache cache_; + WebcamCacheState cacheState_{}; + Rect crop_{}; + std::unique_ptr decoder_; + std::ifstream input_; + std::vector chunk_; + uint8_t** frames_ = nullptr; + int returnedFrames_ = 0; + bool flushed_ = false; +}; + +std::unique_ptr createWebcamStreamDecoder(CUcontext context, const Options& options) { + if (options.webcamAnnexbPath.empty()) { + return nullptr; + } + return std::make_unique(context, options); +} + +__global__ void copyNv12Kernel( + const unsigned char* src, + int srcPitch, + int srcWidth, + int srcHeight, + int srcSurfaceHeight, + unsigned char* dst, + int dstPitch, + int dstChromaOffset, + int dstWidth, + int dstHeight) { + const int x = blockIdx.x * blockDim.x + threadIdx.x; + const int y = blockIdx.y * blockDim.y + threadIdx.y; + if (x >= dstWidth || y >= dstHeight) { + return; + } + + const int sx = min(srcWidth - 1, (x * srcWidth) / dstWidth); + const int sy = min(srcHeight - 1, (y * srcHeight) / dstHeight); + dst[y * dstPitch + x] = src[sy * srcPitch + sx]; + + if ((x % 2) == 0 && (y % 2) == 0) { + const int suvX = min(srcWidth - 2, ((x * srcWidth) / dstWidth) & ~1); + const int suvY = min((srcHeight / 2) - 1, (y * srcHeight / dstHeight) / 2); + const unsigned char* srcUv = src + srcPitch * srcSurfaceHeight + suvY * srcPitch + suvX; + unsigned char* dstUv = dst + dstChromaOffset + (y / 2) * dstPitch + x; + dstUv[0] = srcUv[0]; + dstUv[1] = srcUv[1]; + } +} + +__global__ void fillNv12Kernel( + unsigned char* dst, + int dstPitch, + int dstChromaOffset, + int dstWidth, + int dstHeight, + unsigned char yValue, + unsigned char uValue, + unsigned char vValue) { + const int x = blockIdx.x * blockDim.x + threadIdx.x; + const int y = blockIdx.y * blockDim.y + threadIdx.y; + if (x >= dstWidth || y >= dstHeight) { + return; + } + + dst[y * dstPitch + x] = yValue; + if ((x % 2) == 0 && (y % 2) == 0) { + unsigned char* dstUv = dst + dstChromaOffset + (y / 2) * dstPitch + x; + dstUv[0] = uValue; + dstUv[1] = vValue; + } +} + +__device__ bool isInsideRoundedRect( + int x, + int y, + int left, + int top, + int width, + int height, + int radius) { + if (x < left || y < top || x >= left + width || y >= top + height) { + return false; + } + if (radius <= 0) { + return true; + } + + const int right = left + width - 1; + const int bottom = top + height - 1; + const int innerLeft = left + radius; + const int innerRight = right - radius; + const int innerTop = top + radius; + const int innerBottom = bottom - radius; + if ((x >= innerLeft && x <= innerRight) || (y >= innerTop && y <= innerBottom)) { + return true; + } + + const int cx = x < innerLeft ? innerLeft : innerRight; + const int cy = y < innerTop ? innerTop : innerBottom; + const int dx = x - cx; + const int dy = y - cy; + return dx * dx + dy * dy <= radius * radius; +} + +__global__ void overlayContentRectNv12Kernel( + const unsigned char* src, + int srcPitch, + int srcWidth, + int srcHeight, + int srcSurfaceHeight, + unsigned char* dst, + int dstPitch, + int dstChromaOffset, + int dstWidth, + int dstHeight, + int contentX, + int contentY, + int contentWidth, + int contentHeight) { + const int localX = blockIdx.x * blockDim.x + threadIdx.x; + const int localY = blockIdx.y * blockDim.y + threadIdx.y; + if (localX >= contentWidth || localY >= contentHeight) { + return; + } + + const int x = contentX + localX; + const int y = contentY + localY; + if (x < 0 || y < 0 || x >= dstWidth || y >= dstHeight) { + return; + } + + const int srcX = min(srcWidth - 1, (localX * srcWidth) / contentWidth); + const int srcY = min(srcHeight - 1, (localY * srcHeight) / contentHeight); + dst[y * dstPitch + x] = src[srcY * srcPitch + srcX]; + + if ((x % 2) == 0 && (y % 2) == 0) { + const int localUvX = max(0, min(contentWidth - 1, localX + 1)); + const int localUvY = max(0, min(contentHeight - 1, localY + 1)); + const int srcUvX = min(srcWidth - 2, ((localUvX * srcWidth) / contentWidth) & ~1); + const int srcUvY = min((srcHeight / 2) - 1, ((localUvY * srcHeight) / contentHeight) / 2); + const unsigned char* srcUv = src + srcPitch * srcSurfaceHeight + srcUvY * srcPitch + srcUvX; + unsigned char* dstUv = dst + dstChromaOffset + (y / 2) * dstPitch + x; + dstUv[0] = srcUv[0]; + dstUv[1] = srcUv[1]; + } +} + +__global__ void overlayContentTransformNv12Kernel( + const unsigned char* src, + int srcPitch, + int srcWidth, + int srcHeight, + int srcSurfaceHeight, + unsigned char* dst, + int dstPitch, + int dstChromaOffset, + int dstWidth, + int dstHeight, + int regionX, + int regionY, + int regionWidth, + int regionHeight, + int contentX, + int contentY, + int contentWidth, + int contentHeight, + int radius, + float zoomScale, + float invZoomScale, + float srcScaleX, + float srcScaleY, + float zoomX, + float zoomY) { + const int localX = blockIdx.x * blockDim.x + threadIdx.x; + const int localY = blockIdx.y * blockDim.y + threadIdx.y; + if (localX >= regionWidth || localY >= regionHeight) { + return; + } + + const int x = regionX + localX; + const int y = regionY + localY; + if (x < 0 || y < 0 || x >= dstWidth || y >= dstHeight) { + return; + } + + const float layoutXf = (static_cast(x) - zoomX) * invZoomScale; + const float layoutYf = (static_cast(y) - zoomY) * invZoomScale; + const int layoutX = __float2int_rd(layoutXf); + const int layoutY = __float2int_rd(layoutYf); + if (!isInsideRoundedRect(layoutX, layoutY, contentX, contentY, contentWidth, contentHeight, radius)) { + return; + } + + const float localContentX = + fminf(static_cast(contentWidth - 1), fmaxf(0.0f, layoutXf - contentX)); + const float localContentY = + fminf(static_cast(contentHeight - 1), fmaxf(0.0f, layoutYf - contentY)); + const int sx = min(srcWidth - 1, __float2int_rd(localContentX * srcScaleX)); + const int sy = min(srcHeight - 1, __float2int_rd(localContentY * srcScaleY)); + dst[y * dstPitch + x] = src[sy * srcPitch + sx]; + + if ((x % 2) == 0 && (y % 2) == 0 && x + 1 < dstWidth && y + 1 < dstHeight) { + const float uvLayoutXf = (static_cast(x + 1) - zoomX) * invZoomScale; + const float uvLayoutYf = (static_cast(y + 1) - zoomY) * invZoomScale; + const int uvLayoutX = __float2int_rd(uvLayoutXf); + const int uvLayoutY = __float2int_rd(uvLayoutYf); + if (isInsideRoundedRect( + uvLayoutX, + uvLayoutY, + contentX, + contentY, + contentWidth, + contentHeight, + radius)) { + const float uvLocalContentX = + fminf(static_cast(contentWidth - 1), fmaxf(0.0f, uvLayoutXf - contentX)); + const float uvLocalContentY = + fminf(static_cast(contentHeight - 1), fmaxf(0.0f, uvLayoutYf - contentY)); + const int suvX = + min(srcWidth - 2, __float2int_rd(uvLocalContentX * srcScaleX) & ~1); + const int suvY = + min((srcHeight / 2) - 1, __float2int_rd(uvLocalContentY * srcScaleY) / 2); + const unsigned char* srcUv = src + srcPitch * srcSurfaceHeight + suvY * srcPitch + suvX; + unsigned char* dstUv = dst + dstChromaOffset + (y / 2) * dstPitch + x; + dstUv[0] = srcUv[0]; + dstUv[1] = srcUv[1]; + } + } +} + +__global__ void restoreRoundedContentCornersNv12Kernel( + unsigned char* dst, + int dstPitch, + int dstChromaOffset, + int dstWidth, + int dstHeight, + int contentX, + int contentY, + int contentWidth, + int contentHeight, + int radius, + unsigned char backgroundY, + unsigned char backgroundU, + unsigned char backgroundV, + const unsigned char* background) { + const int localX = blockIdx.x * blockDim.x + threadIdx.x; + const int localY = blockIdx.y * blockDim.y + threadIdx.y; + if (localX >= radius || localY >= radius) { + return; + } + + const int corner = blockIdx.z; + const bool right = corner == 1 || corner == 3; + const bool bottom = corner >= 2; + const int x = right ? contentX + contentWidth - radius + localX : contentX + localX; + const int y = bottom ? contentY + contentHeight - radius + localY : contentY + localY; + if (x < 0 || y < 0 || x >= dstWidth || y >= dstHeight) { + return; + } + + if (!isInsideRoundedRect(x, y, contentX, contentY, contentWidth, contentHeight, radius)) { + dst[y * dstPitch + x] = background ? background[y * dstWidth + x] : backgroundY; + } + + if ((x % 2) == 0 && (y % 2) == 0 && + !isInsideRoundedRect(x + 1, y + 1, contentX, contentY, contentWidth, contentHeight, radius)) { + unsigned char* dstUv = dst + dstChromaOffset + (y / 2) * dstPitch + x; + if (background) { + const unsigned char* bgUv = background + dstWidth * dstHeight + (y / 2) * dstWidth + x; + dstUv[0] = bgUv[0]; + dstUv[1] = bgUv[1]; + } else { + dstUv[0] = backgroundU; + dstUv[1] = backgroundV; + } + } +} + +__device__ bool pointInCursorPolygon(float x, float y, bool inner) { + constexpr int kCount = 7; + const float outerX[kCount] = {2.0f, 61.0f, 45.0f, 57.0f, 38.0f, 27.0f, 13.0f}; + const float outerY[kCount] = {2.0f, 61.0f, 63.0f, 91.0f, 95.0f, 67.0f, 79.0f}; + const float innerX[kCount] = {10.0f, 52.0f, 37.0f, 49.0f, 40.0f, 28.0f, 18.0f}; + const float innerY[kCount] = {11.0f, 53.0f, 53.0f, 78.0f, 83.0f, 57.0f, 66.0f}; + const float* px = inner ? innerX : outerX; + const float* py = inner ? innerY : outerY; + + bool inside = false; + for (int index = 0, previous = kCount - 1; index < kCount; previous = index++) { + const bool crosses = ((py[index] > y) != (py[previous] > y)) && + (x < (px[previous] - px[index]) * (y - py[index]) / (py[previous] - py[index]) + px[index]); + if (crosses) { + inside = !inside; + } + } + return inside; +} + +__device__ int cursorMaskAt( + int x, + int y, + int cursorX, + int cursorY, + int cursorWidth, + int cursorHeight) { + if (cursorWidth <= 0 || cursorHeight <= 0 || x < cursorX || y < cursorY || + x >= cursorX + cursorWidth || y >= cursorY + cursorHeight) { + return 0; + } + + const float localX = static_cast(x - cursorX) * 64.0f / static_cast(cursorWidth); + const float localY = static_cast(y - cursorY) * 96.0f / static_cast(cursorHeight); + if (pointInCursorPolygon(localX, localY, true)) { + return 2; + } + if (pointInCursorPolygon(localX, localY, false)) { + return 1; + } + return 0; +} + +__device__ unsigned char clampByteDevice(int value) { + return static_cast(min(255, max(0, value))); +} + +__device__ unsigned char blendByte(unsigned char base, unsigned char overlay, int alpha) { + return static_cast( + (static_cast(base) * (255 - alpha) + static_cast(overlay) * alpha + 127) / 255); +} + +__device__ bool sampleCursorAtlasNv12( + const unsigned char* atlas, + int atlasWidth, + int atlasHeight, + int entryX, + int entryY, + int entryWidth, + int entryHeight, + int cursorX, + int cursorY, + int cursorWidth, + int cursorHeight, + int x, + int y, + unsigned char* outY, + unsigned char* outU, + unsigned char* outV, + int* outAlpha) { + if (!atlas || cursorWidth <= 0 || cursorHeight <= 0 || entryWidth <= 0 || entryHeight <= 0 || + x < cursorX || y < cursorY || x >= cursorX + cursorWidth || y >= cursorY + cursorHeight) { + return false; + } + + const int localX = max(0, min(cursorWidth - 1, x - cursorX)); + const int localY = max(0, min(cursorHeight - 1, y - cursorY)); + const int sampleX = entryX + min(entryWidth - 1, (localX * entryWidth) / cursorWidth); + const int sampleY = entryY + min(entryHeight - 1, (localY * entryHeight) / cursorHeight); + if (sampleX < 0 || sampleY < 0 || sampleX >= atlasWidth || sampleY >= atlasHeight) { + return false; + } + + const int offset = (sampleY * atlasWidth + sampleX) * 4; + const int alpha = atlas[offset + 3]; + if (alpha <= 0) { + return false; + } + + const int r = atlas[offset]; + const int g = atlas[offset + 1]; + const int b = atlas[offset + 2]; + *outY = clampByteDevice(((66 * r + 129 * g + 25 * b + 128) >> 8) + 16); + *outU = clampByteDevice(((-38 * r - 74 * g + 112 * b + 128) >> 8) + 128); + *outV = clampByteDevice(((112 * r - 94 * g - 18 * b + 128) >> 8) + 128); + *outAlpha = alpha; + return true; +} + +__device__ int sampleCursorAtlasAlpha( + const unsigned char* atlas, + int atlasWidth, + int atlasHeight, + int entryX, + int entryY, + int entryWidth, + int entryHeight, + int cursorX, + int cursorY, + int cursorWidth, + int cursorHeight, + int x, + int y) { + if (!atlas || cursorWidth <= 0 || cursorHeight <= 0 || entryWidth <= 0 || entryHeight <= 0 || + x < cursorX || y < cursorY || x >= cursorX + cursorWidth || y >= cursorY + cursorHeight) { + return 0; + } + + const int localX = max(0, min(cursorWidth - 1, x - cursorX)); + const int localY = max(0, min(cursorHeight - 1, y - cursorY)); + const int sampleX = entryX + min(entryWidth - 1, (localX * entryWidth) / cursorWidth); + const int sampleY = entryY + min(entryHeight - 1, (localY * entryHeight) / cursorHeight); + if (sampleX < 0 || sampleY < 0 || sampleX >= atlasWidth || sampleY >= atlasHeight) { + return 0; + } + + return atlas[(sampleY * atlasWidth + sampleX) * 4 + 3]; +} + +__device__ int sampleCursorAtlasShadowAlpha( + const unsigned char* atlas, + int atlasWidth, + int atlasHeight, + int entryX, + int entryY, + int entryWidth, + int entryHeight, + int cursorX, + int cursorY, + int cursorWidth, + int cursorHeight, + int x, + int y) { + int weightedAlpha = 0; + weightedAlpha += sampleCursorAtlasAlpha( + atlas, + atlasWidth, + atlasHeight, + entryX, + entryY, + entryWidth, + entryHeight, + cursorX, + cursorY + 2, + cursorWidth, + cursorHeight, + x, + y) * 20; + weightedAlpha += sampleCursorAtlasAlpha( + atlas, + atlasWidth, + atlasHeight, + entryX, + entryY, + entryWidth, + entryHeight, + cursorX - 2, + cursorY + 2, + cursorWidth, + cursorHeight, + x, + y) * 6; + weightedAlpha += sampleCursorAtlasAlpha( + atlas, + atlasWidth, + atlasHeight, + entryX, + entryY, + entryWidth, + entryHeight, + cursorX + 2, + cursorY + 2, + cursorWidth, + cursorHeight, + x, + y) * 6; + weightedAlpha += sampleCursorAtlasAlpha( + atlas, + atlasWidth, + atlasHeight, + entryX, + entryY, + entryWidth, + entryHeight, + cursorX, + cursorY, + cursorWidth, + cursorHeight, + x, + y) * 4; + weightedAlpha += sampleCursorAtlasAlpha( + atlas, + atlasWidth, + atlasHeight, + entryX, + entryY, + entryWidth, + entryHeight, + cursorX, + cursorY + 4, + cursorWidth, + cursorHeight, + x, + y) * 4; + return min(255, weightedAlpha / 100); +} + +__global__ void compositeStaticNv12Kernel( + const unsigned char* src, + int srcPitch, + int srcWidth, + int srcHeight, + int srcSurfaceHeight, + unsigned char* dst, + int dstPitch, + int dstChromaOffset, + int dstWidth, + int dstHeight, + int contentX, + int contentY, + int contentWidth, + int contentHeight, + int radius, + unsigned char backgroundY, + unsigned char backgroundU, + unsigned char backgroundV, + const unsigned char* background, + int shadowOffsetY, + int shadowIntensityPct, + const unsigned char* webcam, + int webcamX, + int webcamY, + int webcamSize, + int webcamFrameWidth, + int webcamFrameHeight, + int webcamRadius, + bool webcamMirror, + bool cursorVisible, + int cursorX, + int cursorY, + int cursorWidth, + int cursorHeight, + const unsigned char* cursorAtlasRgba, + int cursorAtlasWidth, + int cursorAtlasHeight, + int cursorAtlasEntryX, + int cursorAtlasEntryY, + int cursorAtlasEntryWidth, + int cursorAtlasEntryHeight, + bool zoomEnabled, + float zoomScale, + float zoomX, + float zoomY) { + const int x = blockIdx.x * blockDim.x + threadIdx.x; + const int y = blockIdx.y * blockDim.y + threadIdx.y; + if (x >= dstWidth || y >= dstHeight) { + return; + } + + const bool zoomActive = zoomEnabled && zoomScale > 0.01f; + const float safeZoomScale = fmaxf(zoomScale, 0.01f); + const float layoutXf = zoomActive ? (static_cast(x) - zoomX) / safeZoomScale : static_cast(x); + const float layoutYf = zoomActive ? (static_cast(y) - zoomY) / safeZoomScale : static_cast(y); + const int layoutX = static_cast(floorf(layoutXf)); + const int layoutY = static_cast(floorf(layoutYf)); + + const bool inside = isInsideRoundedRect(layoutX, layoutY, contentX, contentY, contentWidth, contentHeight, radius); + unsigned char outY = background ? background[y * dstWidth + x] : backgroundY; + if (inside) { + const float localX = fminf(static_cast(contentWidth - 1), fmaxf(0.0f, layoutXf - contentX)); + const float localY = fminf(static_cast(contentHeight - 1), fmaxf(0.0f, layoutYf - contentY)); + const int sx = min(srcWidth - 1, static_cast((localX * srcWidth) / contentWidth)); + const int sy = min(srcHeight - 1, static_cast((localY * srcHeight) / contentHeight)); + outY = src[sy * srcPitch + sx]; + } else { + const bool shadowInside = + shadowIntensityPct > 0 && + isInsideRoundedRect( + layoutX, + layoutY, + contentX, + contentY + shadowOffsetY, + contentWidth, + contentHeight, + radius + 8); + if (shadowInside) { + const int darkenPct = min(75, max(0, shadowIntensityPct / 2)); + outY = static_cast((static_cast(outY) * (100 - darkenPct)) / 100); + } + } + if (webcam && isInsideRoundedRect(x, y, webcamX, webcamY, webcamSize, webcamSize, webcamRadius)) { + const int localX = max(0, min(webcamSize - 1, x - webcamX)); + const int localY = max(0, min(webcamSize - 1, y - webcamY)); + const int sampleX = min(webcamFrameWidth - 1, (localX * webcamFrameWidth) / webcamSize); + const int sampleY = min(webcamFrameHeight - 1, (localY * webcamFrameHeight) / webcamSize); + const int mirroredX = webcamMirror ? webcamFrameWidth - 1 - sampleX : sampleX; + outY = webcam[sampleY * webcamFrameWidth + mirroredX]; + } + unsigned char cursorYValue = 0; + unsigned char cursorUValue = 128; + unsigned char cursorVValue = 128; + int cursorAlpha = 0; + const int cursorShadowAlpha = cursorVisible && cursorAtlasRgba + ? sampleCursorAtlasShadowAlpha( + cursorAtlasRgba, + cursorAtlasWidth, + cursorAtlasHeight, + cursorAtlasEntryX, + cursorAtlasEntryY, + cursorAtlasEntryWidth, + cursorAtlasEntryHeight, + cursorX, + cursorY, + cursorWidth, + cursorHeight, + x, + y) + : 0; + if (cursorShadowAlpha > 0) { + outY = blendByte(outY, 16, cursorShadowAlpha); + } + const bool cursorAtlasHit = + cursorVisible && + sampleCursorAtlasNv12( + cursorAtlasRgba, + cursorAtlasWidth, + cursorAtlasHeight, + cursorAtlasEntryX, + cursorAtlasEntryY, + cursorAtlasEntryWidth, + cursorAtlasEntryHeight, + cursorX, + cursorY, + cursorWidth, + cursorHeight, + x, + y, + &cursorYValue, + &cursorUValue, + &cursorVValue, + &cursorAlpha); + if (cursorAtlasHit) { + outY = blendByte(outY, cursorYValue, cursorAlpha); + } else { + const int cursorMask = cursorVisible && !cursorAtlasRgba + ? cursorMaskAt(x, y, cursorX, cursorY, cursorWidth, cursorHeight) + : 0; + if (cursorMask == 1) { + outY = 235; + } else if (cursorMask == 2) { + outY = 16; + } + } + dst[y * dstPitch + x] = outY; + + if ((x % 2) == 0 && (y % 2) == 0) { + unsigned char* dstUv = dst + dstChromaOffset + (y / 2) * dstPitch + x; + const float uvLayoutXf = + zoomActive ? (static_cast(x + 1) - zoomX) / safeZoomScale : static_cast(x + 1); + const float uvLayoutYf = + zoomActive ? (static_cast(y + 1) - zoomY) / safeZoomScale : static_cast(y + 1); + const int uvLayoutX = static_cast(floorf(uvLayoutXf)); + const int uvLayoutY = static_cast(floorf(uvLayoutYf)); + const bool uvInside = isInsideRoundedRect( + uvLayoutX, + uvLayoutY, + contentX, + contentY, + contentWidth, + contentHeight, + radius); + if (uvInside) { + const float localX = fminf(static_cast(contentWidth - 1), fmaxf(0.0f, uvLayoutXf - contentX)); + const float localY = fminf(static_cast(contentHeight - 1), fmaxf(0.0f, uvLayoutYf - contentY)); + const int suvX = min(srcWidth - 2, (static_cast((localX * srcWidth) / contentWidth)) & ~1); + const int suvY = min((srcHeight / 2) - 1, static_cast(localY * srcHeight / contentHeight) / 2); + const unsigned char* srcUv = src + srcPitch * srcSurfaceHeight + suvY * srcPitch + suvX; + dstUv[0] = srcUv[0]; + dstUv[1] = srcUv[1]; + } else { + if (background) { + const unsigned char* bgUv = background + dstWidth * dstHeight + (y / 2) * dstWidth + x; + dstUv[0] = bgUv[0]; + dstUv[1] = bgUv[1]; + } else { + dstUv[0] = backgroundU; + dstUv[1] = backgroundV; + } + } + if (webcam && + isInsideRoundedRect( + x + 1, + y + 1, + webcamX, + webcamY, + webcamSize, + webcamSize, + webcamRadius)) { + const int localX = max(0, min(webcamSize - 1, x + 1 - webcamX)); + const int localY = max(0, min(webcamSize - 1, y + 1 - webcamY)); + const int sampleX = min(webcamFrameWidth - 1, (localX * webcamFrameWidth) / webcamSize); + const int sampleY = min(webcamFrameHeight - 1, (localY * webcamFrameHeight) / webcamSize); + const int mirroredX = webcamMirror ? webcamFrameWidth - 1 - sampleX : sampleX; + const int webcamUvX = min(webcamFrameWidth - 2, mirroredX & ~1); + const int webcamUvY = min((webcamFrameHeight / 2) - 1, sampleY / 2); + const unsigned char* webcamUv = + webcam + webcamFrameWidth * webcamFrameHeight + webcamUvY * webcamFrameWidth + webcamUvX; + dstUv[0] = webcamUv[0]; + dstUv[1] = webcamUv[1]; + } + unsigned char cursorUvY = 0; + unsigned char cursorUvU = 128; + unsigned char cursorUvV = 128; + int cursorUvAlpha = 0; + const int cursorUvShadowAlpha = cursorVisible && cursorAtlasRgba + ? sampleCursorAtlasShadowAlpha( + cursorAtlasRgba, + cursorAtlasWidth, + cursorAtlasHeight, + cursorAtlasEntryX, + cursorAtlasEntryY, + cursorAtlasEntryWidth, + cursorAtlasEntryHeight, + cursorX, + cursorY, + cursorWidth, + cursorHeight, + x + 1, + y + 1) + : 0; + if (cursorUvShadowAlpha > 0) { + dstUv[0] = blendByte(dstUv[0], 128, cursorUvShadowAlpha); + dstUv[1] = blendByte(dstUv[1], 128, cursorUvShadowAlpha); + } + const bool cursorAtlasUvHit = + cursorVisible && + sampleCursorAtlasNv12( + cursorAtlasRgba, + cursorAtlasWidth, + cursorAtlasHeight, + cursorAtlasEntryX, + cursorAtlasEntryY, + cursorAtlasEntryWidth, + cursorAtlasEntryHeight, + cursorX, + cursorY, + cursorWidth, + cursorHeight, + x + 1, + y + 1, + &cursorUvY, + &cursorUvU, + &cursorUvV, + &cursorUvAlpha); + if (cursorAtlasUvHit) { + dstUv[0] = blendByte(dstUv[0], cursorUvU, cursorUvAlpha); + dstUv[1] = blendByte(dstUv[1], cursorUvV, cursorUvAlpha); + } else { + const int cursorUvMask = + cursorVisible && !cursorAtlasRgba + ? cursorMaskAt(x + 1, y + 1, cursorX, cursorY, cursorWidth, cursorHeight) + : 0; + if (cursorUvMask > 0) { + dstUv[0] = 128; + dstUv[1] = 128; + } + } + } +} + +__global__ void overlayWebcamNv12Kernel( + unsigned char* dst, + int dstPitch, + int dstChromaOffset, + int dstWidth, + int dstHeight, + const unsigned char* webcam, + int regionX, + int regionY, + int regionWidth, + int regionHeight, + int webcamX, + int webcamY, + int webcamSize, + int webcamFrameWidth, + int webcamFrameHeight, + int webcamRadius, + bool webcamMirror) { + const int localX = blockIdx.x * blockDim.x + threadIdx.x; + const int localY = blockIdx.y * blockDim.y + threadIdx.y; + if (!webcam || localX >= regionWidth || localY >= regionHeight) { + return; + } + + const int x = regionX + localX; + const int y = regionY + localY; + if (x < 0 || y < 0 || x >= dstWidth || y >= dstHeight) { + return; + } + + if (isInsideRoundedRect(x, y, webcamX, webcamY, webcamSize, webcamSize, webcamRadius)) { + const int webcamLocalX = max(0, min(webcamSize - 1, x - webcamX)); + const int webcamLocalY = max(0, min(webcamSize - 1, y - webcamY)); + const int sampleX = min(webcamFrameWidth - 1, (webcamLocalX * webcamFrameWidth) / webcamSize); + const int sampleY = min(webcamFrameHeight - 1, (webcamLocalY * webcamFrameHeight) / webcamSize); + const int mirroredX = webcamMirror ? webcamFrameWidth - 1 - sampleX : sampleX; + dst[y * dstPitch + x] = webcam[sampleY * webcamFrameWidth + mirroredX]; + } + + if ((x % 2) == 0 && (y % 2) == 0 && x + 1 < dstWidth && y + 1 < dstHeight && + isInsideRoundedRect(x + 1, y + 1, webcamX, webcamY, webcamSize, webcamSize, webcamRadius)) { + const int uvLocalX = max(0, min(webcamSize - 1, x + 1 - webcamX)); + const int uvLocalY = max(0, min(webcamSize - 1, y + 1 - webcamY)); + const int uvSampleX = min(webcamFrameWidth - 1, (uvLocalX * webcamFrameWidth) / webcamSize); + const int uvSampleY = min(webcamFrameHeight - 1, (uvLocalY * webcamFrameHeight) / webcamSize); + const int uvMirroredX = webcamMirror ? webcamFrameWidth - 1 - uvSampleX : uvSampleX; + const int webcamUvX = min(webcamFrameWidth - 2, uvMirroredX & ~1); + const int webcamUvY = min((webcamFrameHeight / 2) - 1, uvSampleY / 2); + const unsigned char* webcamUv = + webcam + webcamFrameWidth * webcamFrameHeight + webcamUvY * webcamFrameWidth + webcamUvX; + unsigned char* dstUv = dst + dstChromaOffset + (y / 2) * dstPitch + x; + dstUv[0] = webcamUv[0]; + dstUv[1] = webcamUv[1]; + } +} + +__global__ void overlayCursorNv12Kernel( + unsigned char* dst, + int dstPitch, + int dstChromaOffset, + int dstWidth, + int dstHeight, + int regionX, + int regionY, + int regionWidth, + int regionHeight, + bool cursorVisible, + int cursorX, + int cursorY, + int cursorWidth, + int cursorHeight, + const unsigned char* cursorAtlasRgba, + int cursorAtlasWidth, + int cursorAtlasHeight, + int cursorAtlasEntryX, + int cursorAtlasEntryY, + int cursorAtlasEntryWidth, + int cursorAtlasEntryHeight) { + const int localX = blockIdx.x * blockDim.x + threadIdx.x; + const int localY = blockIdx.y * blockDim.y + threadIdx.y; + if (!cursorVisible || localX >= regionWidth || localY >= regionHeight) { + return; + } + + const int x = regionX + localX; + const int y = regionY + localY; + if (x < 0 || y < 0 || x >= dstWidth || y >= dstHeight) { + return; + } + + unsigned char outY = dst[y * dstPitch + x]; + unsigned char cursorYValue = 0; + unsigned char cursorUValue = 128; + unsigned char cursorVValue = 128; + int cursorAlpha = 0; + const int cursorShadowAlpha = cursorAtlasRgba + ? sampleCursorAtlasShadowAlpha( + cursorAtlasRgba, + cursorAtlasWidth, + cursorAtlasHeight, + cursorAtlasEntryX, + cursorAtlasEntryY, + cursorAtlasEntryWidth, + cursorAtlasEntryHeight, + cursorX, + cursorY, + cursorWidth, + cursorHeight, + x, + y) + : 0; + if (cursorShadowAlpha > 0) { + outY = blendByte(outY, 16, cursorShadowAlpha); + } + const bool cursorAtlasHit = + sampleCursorAtlasNv12( + cursorAtlasRgba, + cursorAtlasWidth, + cursorAtlasHeight, + cursorAtlasEntryX, + cursorAtlasEntryY, + cursorAtlasEntryWidth, + cursorAtlasEntryHeight, + cursorX, + cursorY, + cursorWidth, + cursorHeight, + x, + y, + &cursorYValue, + &cursorUValue, + &cursorVValue, + &cursorAlpha); + if (cursorAtlasHit) { + outY = blendByte(outY, cursorYValue, cursorAlpha); + } else { + const int cursorMask = !cursorAtlasRgba + ? cursorMaskAt(x, y, cursorX, cursorY, cursorWidth, cursorHeight) + : 0; + if (cursorMask == 1) { + outY = 235; + } else if (cursorMask == 2) { + outY = 16; + } + } + dst[y * dstPitch + x] = outY; + + if ((x % 2) == 0 && (y % 2) == 0 && x + 1 < dstWidth && y + 1 < dstHeight) { + unsigned char* dstUv = dst + dstChromaOffset + (y / 2) * dstPitch + x; + unsigned char cursorUvY = 0; + unsigned char cursorUvU = 128; + unsigned char cursorUvV = 128; + int cursorUvAlpha = 0; + const int cursorUvShadowAlpha = cursorAtlasRgba + ? sampleCursorAtlasShadowAlpha( + cursorAtlasRgba, + cursorAtlasWidth, + cursorAtlasHeight, + cursorAtlasEntryX, + cursorAtlasEntryY, + cursorAtlasEntryWidth, + cursorAtlasEntryHeight, + cursorX, + cursorY, + cursorWidth, + cursorHeight, + x + 1, + y + 1) + : 0; + if (cursorUvShadowAlpha > 0) { + dstUv[0] = blendByte(dstUv[0], 128, cursorUvShadowAlpha); + dstUv[1] = blendByte(dstUv[1], 128, cursorUvShadowAlpha); + } + const bool cursorAtlasUvHit = + sampleCursorAtlasNv12( + cursorAtlasRgba, + cursorAtlasWidth, + cursorAtlasHeight, + cursorAtlasEntryX, + cursorAtlasEntryY, + cursorAtlasEntryWidth, + cursorAtlasEntryHeight, + cursorX, + cursorY, + cursorWidth, + cursorHeight, + x + 1, + y + 1, + &cursorUvY, + &cursorUvU, + &cursorUvV, + &cursorUvAlpha); + if (cursorAtlasUvHit) { + dstUv[0] = blendByte(dstUv[0], cursorUvU, cursorUvAlpha); + dstUv[1] = blendByte(dstUv[1], cursorUvV, cursorUvAlpha); + } else { + const int cursorUvMask = + !cursorAtlasRgba ? cursorMaskAt(x + 1, y + 1, cursorX, cursorY, cursorWidth, cursorHeight) : 0; + if (cursorUvMask > 0) { + dstUv[0] = 128; + dstUv[1] = 128; + } + } + } +} + +__global__ void prewarmKernel(unsigned int* state, unsigned int seed) { + const unsigned int index = blockIdx.x * blockDim.x + threadIdx.x; + unsigned int value = seed ^ (index * 747796405u + 2891336453u); + for (int iteration = 0; iteration < 256; ++iteration) { + value = value * 1664525u + 1013904223u; + value ^= value >> 13; + } + state[index] = value; +} + +void prewarmCuda(int durationMs) { + if (durationMs <= 0) { + return; + } + + constexpr int blockSize = 256; + constexpr int blockCount = 256; + unsigned int* state = nullptr; + checkCuda(cudaMalloc(&state, blockSize * blockCount * sizeof(unsigned int)), "cudaMalloc prewarm"); + + const auto start = std::chrono::steady_clock::now(); + int iteration = 0; + while (elapsedMs(start, std::chrono::steady_clock::now()) < durationMs) { + prewarmKernel<<>>(state, static_cast(iteration++)); + checkCuda(cudaGetLastError(), "prewarmKernel"); + checkCuda(cudaDeviceSynchronize(), "cudaDeviceSynchronize prewarm"); + } + + checkCuda(cudaFree(state), "cudaFree prewarm"); +} + +class NvencSink { +public: + NvencSink( + CUcontext context, + int width, + int height, + int fps, + uint32_t bitrate, + const std::string& outputPath, + bool streamSync, + Options layoutOptions, + const WebcamFrameCache* webcamCache, + const CursorTrack* cursorTrack, + const ZoomTrack* zoomTrack) + : encoder_(context, width, height, NV_ENC_BUFFER_FORMAT_NV12), + width_(width), + height_(height), + fps_(fps), + streamSync_(streamSync), + layoutOptions_(layoutOptions), + webcamCache_(webcamCache), + cursorTrack_(cursorTrack), + zoomTrack_(zoomTrack) { + loadBackgroundFrame(); + loadWebcamFrame(); + loadCursorAtlas(); + if (streamSync_) { + checkCuda(cudaStreamCreateWithFlags(©Stream_, cudaStreamNonBlocking), "cudaStreamCreateWithFlags"); + } + + NV_ENC_INITIALIZE_PARAMS initializeParams = {NV_ENC_INITIALIZE_PARAMS_VER}; + NV_ENC_CONFIG encodeConfig = {NV_ENC_CONFIG_VER}; + initializeParams.encodeConfig = &encodeConfig; + encoder_.CreateDefaultEncoderParams(&initializeParams, NV_ENC_CODEC_H264_GUID, NV_ENC_PRESET_HP_GUID); + + initializeParams.frameRateNum = static_cast(fps); + initializeParams.frameRateDen = 1; + initializeParams.enableEncodeAsync = 1; + encodeConfig.profileGUID = NV_ENC_H264_PROFILE_HIGH_GUID; + encodeConfig.gopLength = static_cast(fps * 2); + encodeConfig.frameIntervalP = 1; + encodeConfig.rcParams.rateControlMode = NV_ENC_PARAMS_RC_VBR; + encodeConfig.rcParams.averageBitRate = bitrate; + encodeConfig.rcParams.maxBitRate = static_cast( + std::min(0xffffffffu, static_cast(bitrate) * 3 / 2)); + encodeConfig.rcParams.vbvBufferSize = bitrate; + encodeConfig.rcParams.vbvInitialDelay = bitrate; + encodeConfig.encodeCodecConfig.h264Config.idrPeriod = encodeConfig.gopLength; + encoder_.CreateEncoder(&initializeParams); + + output_.open(outputPath, std::ios::binary); + if (!output_) { + fail("Failed to open output: " + outputPath); + } + } + + void encodeFrame( + const unsigned char* srcFrame, + int srcPitch, + int srcWidth, + int srcHeight, + int srcSurfaceHeight, + int outputFrameIndex) { + const NvEncInputFrame* inputFrame = encoder_.GetNextInputFrame(); + const dim3 block(16, 16); + const dim3 grid((width_ + block.x - 1) / block.x, (height_ + block.y - 1) / block.y); + const unsigned char* webcamFrame = selectWebcamFrame(outputFrameIndex); + const double frameTimeMs = static_cast(outputFrameIndex) * 1000.0 / static_cast(fps_); + const ZoomSample zoomSample = zoomTrack_ ? zoomTrack_->sampleAt(frameTimeMs) : ZoomSample{}; + const bool zoomEnabled = zoomTrack_ && zoomSample.scale > 0.01; + const CursorPosition cursorPosition = cursorTrack_ + ? cursorTrack_->positionAt(frameTimeMs) + : CursorPosition{}; + const CursorAtlasEntry* cursorEntry = cursorAtlasEntryFor(cursorPosition.typeIndex); + const bool useCursorAtlas = cursorEntry && cursorAtlasDevice_; + const int cursorHeight = layoutOptions_.cursorHeight > 0 + ? std::max(1, static_cast(std::round(layoutOptions_.cursorHeight * cursorPosition.bounceScale))) + : 0; + const double cursorAspectRatio = useCursorAtlas ? cursorEntry->aspectRatio : (618.0 / 958.0); + const int cursorWidth = + cursorHeight > 0 ? std::max(1, static_cast(std::round(cursorHeight * cursorAspectRatio))) : 0; + const int cursorHotspotX = useCursorAtlas + ? static_cast(std::round(cursorWidth * cursorEntry->anchorX)) + : cursorWidth * 14 / 100; + const int cursorHotspotY = useCursorAtlas + ? static_cast(std::round(cursorHeight * cursorEntry->anchorY)) + : cursorHeight * 6 / 100; + const double cursorHotspotContentX = + layoutOptions_.contentX + cursorPosition.cx * layoutOptions_.contentWidth; + const double cursorHotspotContentY = + layoutOptions_.contentY + cursorPosition.cy * layoutOptions_.contentHeight; + const double cursorHotspotOutputX = + zoomEnabled ? cursorHotspotContentX * zoomSample.scale + zoomSample.x : cursorHotspotContentX; + const double cursorHotspotOutputY = + zoomEnabled ? cursorHotspotContentY * zoomSample.scale + zoomSample.y : cursorHotspotContentY; + const int cursorX = cursorPosition.visible + ? static_cast(std::round(cursorHotspotOutputX)) - cursorHotspotX + : 0; + const int cursorY = cursorPosition.visible + ? static_cast(std::round(cursorHotspotOutputY)) - cursorHotspotY + : 0; + const bool zoomChangesLayout = + zoomTrack_ && + (std::abs(zoomSample.scale - 1.0) > 0.001 || + std::abs(zoomSample.x) > 0.5 || + std::abs(zoomSample.y) > 0.5); + const bool useFastRoiComposite = + canUseFastRoiComposite(zoomChangesLayout); + const bool useLayeredStaticRoiComposite = + !useFastRoiComposite && canUseLayeredStaticRoiComposite(zoomChangesLayout); + const auto compositeStart = std::chrono::steady_clock::now(); + if (useFastRoiComposite) { + copyNv12Kernel<<>>( + srcFrame, + srcPitch, + srcWidth, + srcHeight, + srcSurfaceHeight, + static_cast(inputFrame->inputPtr), + static_cast(inputFrame->pitch), + static_cast(inputFrame->chromaOffsets[0]), + width_, + height_); + checkCuda(cudaGetLastError(), "copyNv12Kernel fast ROI base"); + + if (webcamFrame && layoutOptions_.webcamSize > 0) { + const int webcamRegionX = std::max(0, layoutOptions_.webcamX - 1); + const int webcamRegionY = std::max(0, layoutOptions_.webcamY - 1); + const int webcamRegionRight = std::min( + width_, + layoutOptions_.webcamX + layoutOptions_.webcamSize); + const int webcamRegionBottom = std::min( + height_, + layoutOptions_.webcamY + layoutOptions_.webcamSize); + const int webcamRegionWidth = webcamRegionRight - webcamRegionX; + const int webcamRegionHeight = webcamRegionBottom - webcamRegionY; + if (webcamRegionWidth > 0 && webcamRegionHeight > 0) { + const dim3 webcamGrid( + (webcamRegionWidth + block.x - 1) / block.x, + (webcamRegionHeight + block.y - 1) / block.y); + overlayWebcamNv12Kernel<<>>( + static_cast(inputFrame->inputPtr), + static_cast(inputFrame->pitch), + static_cast(inputFrame->chromaOffsets[0]), + width_, + height_, + webcamFrame, + webcamRegionX, + webcamRegionY, + webcamRegionWidth, + webcamRegionHeight, + layoutOptions_.webcamX, + layoutOptions_.webcamY, + layoutOptions_.webcamSize, + webcamFrameWidth(), + webcamFrameHeight(), + layoutOptions_.webcamRadius, + layoutOptions_.webcamMirror); + checkCuda(cudaGetLastError(), "overlayWebcamNv12Kernel"); + } + } + + if (cursorPosition.visible && cursorWidth > 0 && cursorHeight > 0) { + const int cursorPadding = useCursorAtlas ? 4 : 2; + const int regionX = std::max(0, cursorX - cursorPadding); + const int regionY = std::max(0, cursorY - cursorPadding); + const int regionRight = std::min(width_, cursorX + cursorWidth + cursorPadding); + const int regionBottom = std::min(height_, cursorY + cursorHeight + cursorPadding); + const int regionWidth = regionRight - regionX; + const int regionHeight = regionBottom - regionY; + if (regionWidth > 0 && regionHeight > 0) { + const dim3 cursorGrid( + (regionWidth + block.x - 1) / block.x, + (regionHeight + block.y - 1) / block.y); + overlayCursorNv12Kernel<<>>( + static_cast(inputFrame->inputPtr), + static_cast(inputFrame->pitch), + static_cast(inputFrame->chromaOffsets[0]), + width_, + height_, + regionX, + regionY, + regionWidth, + regionHeight, + cursorPosition.visible, + cursorX, + cursorY, + cursorWidth, + cursorHeight, + useCursorAtlas ? cursorAtlasDevice_ : nullptr, + cursorAtlasWidth_, + cursorAtlasHeight_, + useCursorAtlas ? cursorEntry->x : 0, + useCursorAtlas ? cursorEntry->y : 0, + useCursorAtlas ? cursorEntry->width : 0, + useCursorAtlas ? cursorEntry->height : 0); + checkCuda(cudaGetLastError(), "overlayCursorNv12Kernel"); + } + } + ++roiCompositeFrames_; + } else if (useLayeredStaticRoiComposite) { + if (backgroundDevice_) { + checkCuda( + cudaMemcpy2DAsync( + static_cast(inputFrame->inputPtr), + static_cast(inputFrame->pitch), + backgroundDevice_, + static_cast(width_), + static_cast(width_), + static_cast(height_), + cudaMemcpyDeviceToDevice, + copyStream_), + "cudaMemcpy2DAsync layered ROI background Y"); + checkCuda( + cudaMemcpy2DAsync( + static_cast(inputFrame->inputPtr) + + static_cast(inputFrame->chromaOffsets[0]), + static_cast(inputFrame->pitch), + backgroundDevice_ + width_ * height_, + static_cast(width_), + static_cast(width_), + static_cast(height_ / 2), + cudaMemcpyDeviceToDevice, + copyStream_), + "cudaMemcpy2DAsync layered ROI background UV"); + } else { + fillNv12Kernel<<>>( + static_cast(inputFrame->inputPtr), + static_cast(inputFrame->pitch), + static_cast(inputFrame->chromaOffsets[0]), + width_, + height_, + clampByte(layoutOptions_.backgroundY), + clampByte(layoutOptions_.backgroundU), + clampByte(layoutOptions_.backgroundV)); + checkCuda(cudaGetLastError(), "fillNv12Kernel layered ROI background"); + } + + const float safeZoomScale = std::max(0.01f, static_cast(zoomSample.scale)); + const float invZoomScale = 1.0f / safeZoomScale; + const float srcScaleX = + static_cast(srcWidth) / static_cast(std::max(1, layoutOptions_.contentWidth)); + const float srcScaleY = + static_cast(srcHeight) / static_cast(std::max(1, layoutOptions_.contentHeight)); + const int transformedContentX = zoomChangesLayout + ? static_cast(std::floor(layoutOptions_.contentX * safeZoomScale + zoomSample.x)) + : layoutOptions_.contentX; + const int transformedContentY = zoomChangesLayout + ? static_cast(std::floor(layoutOptions_.contentY * safeZoomScale + zoomSample.y)) + : layoutOptions_.contentY; + const int transformedContentRight = zoomChangesLayout + ? static_cast( + std::ceil( + (layoutOptions_.contentX + layoutOptions_.contentWidth) * + safeZoomScale + + zoomSample.x)) + : layoutOptions_.contentX + layoutOptions_.contentWidth; + const int transformedContentBottom = zoomChangesLayout + ? static_cast( + std::ceil( + (layoutOptions_.contentY + layoutOptions_.contentHeight) * + safeZoomScale + + zoomSample.y)) + : layoutOptions_.contentY + layoutOptions_.contentHeight; + const int contentRegionLeft = std::max(0, transformedContentX); + const int contentRegionTop = std::max(0, transformedContentY); + const int contentRegionRight = std::min(width_, transformedContentRight); + const int contentRegionBottom = std::min(height_, transformedContentBottom); + const int contentRegionWidth = contentRegionRight - contentRegionLeft; + const int contentRegionHeight = contentRegionBottom - contentRegionTop; + if (contentRegionWidth > 0 && contentRegionHeight > 0) { + const dim3 contentGrid( + (contentRegionWidth + block.x - 1) / block.x, + (contentRegionHeight + block.y - 1) / block.y); + if (zoomChangesLayout) { + overlayContentTransformNv12Kernel<<>>( + srcFrame, + srcPitch, + srcWidth, + srcHeight, + srcSurfaceHeight, + static_cast(inputFrame->inputPtr), + static_cast(inputFrame->pitch), + static_cast(inputFrame->chromaOffsets[0]), + width_, + height_, + contentRegionLeft, + contentRegionTop, + contentRegionWidth, + contentRegionHeight, + layoutOptions_.contentX, + layoutOptions_.contentY, + layoutOptions_.contentWidth, + layoutOptions_.contentHeight, + layoutOptions_.radius, + safeZoomScale, + invZoomScale, + srcScaleX, + srcScaleY, + static_cast(zoomSample.x), + static_cast(zoomSample.y)); + checkCuda(cudaGetLastError(), "overlayContentTransformNv12Kernel"); + } else { + overlayContentRectNv12Kernel<<>>( + srcFrame, + srcPitch, + srcWidth, + srcHeight, + srcSurfaceHeight, + static_cast(inputFrame->inputPtr), + static_cast(inputFrame->pitch), + static_cast(inputFrame->chromaOffsets[0]), + width_, + height_, + layoutOptions_.contentX, + layoutOptions_.contentY, + layoutOptions_.contentWidth, + layoutOptions_.contentHeight); + checkCuda(cudaGetLastError(), "overlayContentRectNv12Kernel"); + + const int cornerRadius = std::min( + layoutOptions_.radius, + std::min(layoutOptions_.contentWidth, layoutOptions_.contentHeight) / 2); + if (cornerRadius > 0) { + const dim3 cornerGrid( + (cornerRadius + block.x - 1) / block.x, + (cornerRadius + block.y - 1) / block.y, + 4); + restoreRoundedContentCornersNv12Kernel<<>>( + static_cast(inputFrame->inputPtr), + static_cast(inputFrame->pitch), + static_cast(inputFrame->chromaOffsets[0]), + width_, + height_, + layoutOptions_.contentX, + layoutOptions_.contentY, + layoutOptions_.contentWidth, + layoutOptions_.contentHeight, + cornerRadius, + clampByte(layoutOptions_.backgroundY), + clampByte(layoutOptions_.backgroundU), + clampByte(layoutOptions_.backgroundV), + backgroundDevice_); + checkCuda(cudaGetLastError(), "restoreRoundedContentCornersNv12Kernel"); + } + } + } + + if (webcamFrame && layoutOptions_.webcamSize > 0) { + const int webcamRegionX = std::max(0, layoutOptions_.webcamX - 1); + const int webcamRegionY = std::max(0, layoutOptions_.webcamY - 1); + const int webcamRegionRight = std::min( + width_, + layoutOptions_.webcamX + layoutOptions_.webcamSize); + const int webcamRegionBottom = std::min( + height_, + layoutOptions_.webcamY + layoutOptions_.webcamSize); + const int webcamRegionWidth = webcamRegionRight - webcamRegionX; + const int webcamRegionHeight = webcamRegionBottom - webcamRegionY; + if (webcamRegionWidth > 0 && webcamRegionHeight > 0) { + const dim3 webcamGrid( + (webcamRegionWidth + block.x - 1) / block.x, + (webcamRegionHeight + block.y - 1) / block.y); + overlayWebcamNv12Kernel<<>>( + static_cast(inputFrame->inputPtr), + static_cast(inputFrame->pitch), + static_cast(inputFrame->chromaOffsets[0]), + width_, + height_, + webcamFrame, + webcamRegionX, + webcamRegionY, + webcamRegionWidth, + webcamRegionHeight, + layoutOptions_.webcamX, + layoutOptions_.webcamY, + layoutOptions_.webcamSize, + webcamFrameWidth(), + webcamFrameHeight(), + layoutOptions_.webcamRadius, + layoutOptions_.webcamMirror); + checkCuda(cudaGetLastError(), "overlayWebcamNv12Kernel layered ROI"); + } + } + + if (cursorPosition.visible && cursorWidth > 0 && cursorHeight > 0) { + const int cursorPadding = useCursorAtlas ? 4 : 2; + const int regionX = std::max(0, cursorX - cursorPadding); + const int regionY = std::max(0, cursorY - cursorPadding); + const int regionRight = std::min(width_, cursorX + cursorWidth + cursorPadding); + const int regionBottom = std::min(height_, cursorY + cursorHeight + cursorPadding); + const int regionWidth = regionRight - regionX; + const int regionHeight = regionBottom - regionY; + if (regionWidth > 0 && regionHeight > 0) { + const dim3 cursorGrid( + (regionWidth + block.x - 1) / block.x, + (regionHeight + block.y - 1) / block.y); + overlayCursorNv12Kernel<<>>( + static_cast(inputFrame->inputPtr), + static_cast(inputFrame->pitch), + static_cast(inputFrame->chromaOffsets[0]), + width_, + height_, + regionX, + regionY, + regionWidth, + regionHeight, + cursorPosition.visible, + cursorX, + cursorY, + cursorWidth, + cursorHeight, + useCursorAtlas ? cursorAtlasDevice_ : nullptr, + cursorAtlasWidth_, + cursorAtlasHeight_, + useCursorAtlas ? cursorEntry->x : 0, + useCursorAtlas ? cursorEntry->y : 0, + useCursorAtlas ? cursorEntry->width : 0, + useCursorAtlas ? cursorEntry->height : 0); + checkCuda(cudaGetLastError(), "overlayCursorNv12Kernel layered ROI"); + } + } + ++roiCompositeFrames_; + } else if (hasStaticLayout(layoutOptions_)) { + compositeStaticNv12Kernel<<>>( + srcFrame, + srcPitch, + srcWidth, + srcHeight, + srcSurfaceHeight, + static_cast(inputFrame->inputPtr), + static_cast(inputFrame->pitch), + static_cast(inputFrame->chromaOffsets[0]), + width_, + height_, + layoutOptions_.contentX, + layoutOptions_.contentY, + layoutOptions_.contentWidth, + layoutOptions_.contentHeight, + layoutOptions_.radius, + clampByte(layoutOptions_.backgroundY), + clampByte(layoutOptions_.backgroundU), + clampByte(layoutOptions_.backgroundV), + backgroundDevice_, + layoutOptions_.shadowOffsetY, + layoutOptions_.shadowIntensityPct, + webcamFrame, + layoutOptions_.webcamX, + layoutOptions_.webcamY, + layoutOptions_.webcamSize, + webcamFrameWidth(), + webcamFrameHeight(), + layoutOptions_.webcamRadius, + layoutOptions_.webcamMirror, + cursorPosition.visible, + cursorX, + cursorY, + cursorWidth, + cursorHeight, + useCursorAtlas ? cursorAtlasDevice_ : nullptr, + cursorAtlasWidth_, + cursorAtlasHeight_, + useCursorAtlas ? cursorEntry->x : 0, + useCursorAtlas ? cursorEntry->y : 0, + useCursorAtlas ? cursorEntry->width : 0, + useCursorAtlas ? cursorEntry->height : 0, + zoomEnabled, + static_cast(zoomSample.scale), + static_cast(zoomSample.x), + static_cast(zoomSample.y)); + checkCuda(cudaGetLastError(), "compositeStaticNv12Kernel"); + ++monolithicCompositeFrames_; + } else { + copyNv12Kernel<<>>( + srcFrame, + srcPitch, + srcWidth, + srcHeight, + srcSurfaceHeight, + static_cast(inputFrame->inputPtr), + static_cast(inputFrame->pitch), + static_cast(inputFrame->chromaOffsets[0]), + width_, + height_); + checkCuda(cudaGetLastError(), "copyNv12Kernel"); + ++copyCompositeFrames_; + } + if (streamSync_) { + checkCuda(cudaStreamSynchronize(copyStream_), "cudaStreamSynchronize copy"); + } else { + checkCuda(cudaDeviceSynchronize(), "cudaDeviceSynchronize"); + } + const auto compositeEnd = std::chrono::steady_clock::now(); + compositeMs_ += elapsedMs(compositeStart, compositeEnd); + + std::vector> packets; + const auto nvencStart = std::chrono::steady_clock::now(); + encoder_.EncodeFrame(packets); + const auto nvencEnd = std::chrono::steady_clock::now(); + nvencMs_ += elapsedMs(nvencStart, nvencEnd); + const auto writeStart = std::chrono::steady_clock::now(); + writePackets(packets); + const auto writeEnd = std::chrono::steady_clock::now(); + packetWriteMs_ += elapsedMs(writeStart, writeEnd); + ++frames_; + } + + void finish() { + std::vector> packets; + encoder_.EndEncode(packets); + writePackets(packets); + encoder_.DestroyEncoder(); + output_.close(); + if (copyStream_) { + checkCuda(cudaStreamDestroy(copyStream_), "cudaStreamDestroy"); + copyStream_ = nullptr; + } + if (backgroundDevice_) { + checkCuda(cudaFree(backgroundDevice_), "cudaFree backgroundDevice"); + backgroundDevice_ = nullptr; + } + if (webcamDevice_) { + checkCuda(cudaFree(webcamDevice_), "cudaFree webcamDevice"); + webcamDevice_ = nullptr; + } + if (cursorAtlasDevice_) { + checkCuda(cudaFree(cursorAtlasDevice_), "cudaFree cursorAtlasDevice"); + cursorAtlasDevice_ = nullptr; + } + } + + uint64_t outputBytes() const { + return outputBytes_; + } + + double compositeMs() const { + return compositeMs_; + } + + double nvencMs() const { + return nvencMs_; + } + + double packetWriteMs() const { + return packetWriteMs_; + } + + int roiCompositeFrames() const { + return roiCompositeFrames_; + } + + int monolithicCompositeFrames() const { + return monolithicCompositeFrames_; + } + + int copyCompositeFrames() const { + return copyCompositeFrames_; + } + +private: + bool canUseFastRoiComposite(bool zoomChangesLayout) const { + return hasStaticLayout(layoutOptions_) && + layoutOptions_.contentX == 0 && + layoutOptions_.contentY == 0 && + layoutOptions_.contentWidth == width_ && + layoutOptions_.contentHeight == height_ && + layoutOptions_.radius == 0 && + layoutOptions_.shadowIntensityPct == 0 && + backgroundDevice_ == nullptr && + !zoomChangesLayout; + } + + bool canUseLayeredStaticRoiComposite(bool zoomChangesLayout) const { + return hasStaticLayout(layoutOptions_) && + layoutOptions_.contentWidth > 0 && + layoutOptions_.contentHeight > 0 && + layoutOptions_.contentX < width_ && + layoutOptions_.contentY < height_ && + layoutOptions_.contentX + layoutOptions_.contentWidth > 0 && + layoutOptions_.contentY + layoutOptions_.contentHeight > 0 && + layoutOptions_.shadowIntensityPct == 0; + } + + const unsigned char* selectWebcamFrame(int outputFrameIndex) const { + if (webcamCache_ && !webcamCache_->frames.empty()) { + return webcamCache_->frameAt(webcamFrameIndexForOutputFrame(outputFrameIndex, layoutOptions_)); + } + return webcamDevice_; + } + + int webcamFrameWidth() const { + return webcamCache_ ? webcamCache_->width : layoutOptions_.webcamSize; + } + + int webcamFrameHeight() const { + return webcamCache_ ? webcamCache_->height : layoutOptions_.webcamSize; + } + + void loadBackgroundFrame() { + if (layoutOptions_.backgroundNv12Path.empty()) { + return; + } + + const size_t expectedBytes = static_cast(width_) * static_cast(height_) * 3 / 2; + std::vector bytes(expectedBytes); + std::ifstream input(layoutOptions_.backgroundNv12Path, std::ios::binary); + if (!input) { + fail("Failed to open background NV12: " + layoutOptions_.backgroundNv12Path); + } + input.read(reinterpret_cast(bytes.data()), static_cast(bytes.size())); + if (static_cast(input.gcount()) != expectedBytes) { + fail("Background NV12 has an unexpected size: " + layoutOptions_.backgroundNv12Path); + } + + checkCuda(cudaMalloc(&backgroundDevice_, expectedBytes), "cudaMalloc backgroundDevice"); + checkCuda( + cudaMemcpy(backgroundDevice_, bytes.data(), expectedBytes, cudaMemcpyHostToDevice), + "cudaMemcpy backgroundDevice"); + } + + void loadWebcamFrame() { + if (layoutOptions_.webcamNv12Path.empty()) { + return; + } + + const size_t expectedBytes = + static_cast(layoutOptions_.webcamSize) * static_cast(layoutOptions_.webcamSize) * 3 / 2; + std::vector bytes(expectedBytes); + std::ifstream input(layoutOptions_.webcamNv12Path, std::ios::binary); + if (!input) { + fail("Failed to open webcam NV12: " + layoutOptions_.webcamNv12Path); + } + input.read(reinterpret_cast(bytes.data()), static_cast(bytes.size())); + if (static_cast(input.gcount()) != expectedBytes) { + fail("Webcam NV12 has an unexpected size: " + layoutOptions_.webcamNv12Path); + } + + checkCuda(cudaMalloc(&webcamDevice_, expectedBytes), "cudaMalloc webcamDevice"); + checkCuda( + cudaMemcpy(webcamDevice_, bytes.data(), expectedBytes, cudaMemcpyHostToDevice), + "cudaMemcpy webcamDevice"); + } + + const CursorAtlasEntry* cursorAtlasEntryFor(int typeIndex) const { + if (typeIndex < 0 || typeIndex >= kMaxCursorAtlasEntries) { + return nullptr; + } + const CursorAtlasEntry& entry = cursorAtlasEntries_[typeIndex]; + return entry.valid ? &entry : nullptr; + } + + void loadCursorAtlas() { + if (layoutOptions_.cursorAtlasRgbaPath.empty()) { + return; + } + if (layoutOptions_.cursorAtlasMetadataPath.empty() || + layoutOptions_.cursorAtlasWidth <= 0 || + layoutOptions_.cursorAtlasHeight <= 0) { + fail("Cursor atlas requires metadata, width, and height"); + } + + std::ifstream metadata(layoutOptions_.cursorAtlasMetadataPath); + if (!metadata) { + fail("Failed to open cursor atlas metadata: " + layoutOptions_.cursorAtlasMetadataPath); + } + + int loadedEntries = 0; + int index = 0; + CursorAtlasEntry entry; + while (metadata >> index >> entry.x >> entry.y >> entry.width >> entry.height >> + entry.anchorX >> entry.anchorY >> entry.aspectRatio) { + if (index < 0 || index >= kMaxCursorAtlasEntries || entry.width <= 0 || entry.height <= 0) { + continue; + } + entry.valid = true; + cursorAtlasEntries_[index] = entry; + ++loadedEntries; + } + if (loadedEntries == 0) { + fail("No cursor atlas entries were loaded: " + layoutOptions_.cursorAtlasMetadataPath); + } + + cursorAtlasWidth_ = layoutOptions_.cursorAtlasWidth; + cursorAtlasHeight_ = layoutOptions_.cursorAtlasHeight; + const size_t expectedBytes = + static_cast(cursorAtlasWidth_) * static_cast(cursorAtlasHeight_) * 4; + std::vector bytes(expectedBytes); + std::ifstream input(layoutOptions_.cursorAtlasRgbaPath, std::ios::binary); + if (!input) { + fail("Failed to open cursor atlas RGBA: " + layoutOptions_.cursorAtlasRgbaPath); + } + input.read(reinterpret_cast(bytes.data()), static_cast(bytes.size())); + if (static_cast(input.gcount()) != expectedBytes) { + fail("Cursor atlas RGBA has an unexpected size: " + layoutOptions_.cursorAtlasRgbaPath); + } + + checkCuda(cudaMalloc(&cursorAtlasDevice_, expectedBytes), "cudaMalloc cursorAtlasDevice"); + checkCuda( + cudaMemcpy(cursorAtlasDevice_, bytes.data(), expectedBytes, cudaMemcpyHostToDevice), + "cudaMemcpy cursorAtlasDevice"); + } + + void writePackets(const std::vector>& packets) { + for (const auto& packet : packets) { + if (packet.empty()) { + continue; + } + output_.write(reinterpret_cast(packet.data()), static_cast(packet.size())); + outputBytes_ += packet.size(); + } + } + + NvEncoderCuda encoder_; + std::ofstream output_; + int width_ = 0; + int height_ = 0; + int fps_ = 30; + int frames_ = 0; + uint64_t outputBytes_ = 0; + double compositeMs_ = 0.0; + double nvencMs_ = 0.0; + double packetWriteMs_ = 0.0; + int roiCompositeFrames_ = 0; + int monolithicCompositeFrames_ = 0; + int copyCompositeFrames_ = 0; + bool streamSync_ = false; + Options layoutOptions_; + unsigned char* backgroundDevice_ = nullptr; + unsigned char* webcamDevice_ = nullptr; + unsigned char* cursorAtlasDevice_ = nullptr; + CursorAtlasEntry cursorAtlasEntries_[kMaxCursorAtlasEntries]; + int cursorAtlasWidth_ = 0; + int cursorAtlasHeight_ = 0; + const WebcamFrameCache* webcamCache_ = nullptr; + const CursorTrack* cursorTrack_ = nullptr; + const ZoomTrack* zoomTrack_ = nullptr; + cudaStream_t copyStream_ = nullptr; +}; + +struct CallbackEncodeState { + CUcontext context = nullptr; + const Options* options = nullptr; + uint32_t bitrate = 0; + std::unique_ptr* sink = nullptr; + double* decodeMs = nullptr; + double* encodeMs = nullptr; + int* encodedFrames = nullptr; + int displayFrameIndex = 0; + const WebcamFrameCache* webcamCache = nullptr; + const CursorTrack* cursorTrack = nullptr; + const ZoomTrack* zoomTrack = nullptr; + ProgressReportState* progress = nullptr; + bool oneFramePerMappedDisplayFrame = false; + int mappedFrames = 0; + const std::vector* sourcePts = nullptr; +}; + +ProgressCounters collectProgressCounters( + const NvencSink* sink, + const WebcamFrameCache* webcamCache, + double decodeWallMs, + double encodeMs) { + ProgressCounters counters; + counters.decodeWallMs = decodeWallMs; + counters.encodeMs = encodeMs; + if (sink) { + counters.compositeMs = sink->compositeMs(); + counters.nvencMs = sink->nvencMs(); + counters.packetWriteMs = sink->packetWriteMs(); + counters.roiCompositeFrames = sink->roiCompositeFrames(); + counters.monolithicCompositeFrames = sink->monolithicCompositeFrames(); + counters.copyCompositeFrames = sink->copyCompositeFrames(); + } + if (webcamCache) { + counters.webcamDecodeMs = webcamCache->decodeMs; + counters.webcamCopyMs = webcamCache->copyMs; + } + return counters; +} + +int maxCallbackOutputFrames(const Options& options) { + int maxOutputFrames = options.targetFrames > 0 + ? options.targetFrames + : std::numeric_limits::max(); + if (options.maxFrames > 0) { + maxOutputFrames = std::min(maxOutputFrames, options.maxFrames); + } + return maxOutputFrames; +} + +bool shouldContinueEncoding(int encodedFrames, const Options& options) { + return encodedFrames < maxCallbackOutputFrames(options); +} + +int expectedCallbackOutputFramesForSourceFrame( + int sourceFrameIndex, + const Options& options, + const std::vector* sourcePts) { + return expectedOutputFramesForSourceFrame( + sourceFrameIndex, + options.inputFrames, + options.targetFrames, + options.maxFrames, + options.fps, + sourcePts); +} + +void encodeMappedDisplayFrame( + CUdeviceptr dpSrcFrame, + unsigned int nSrcPitch, + int width, + int height, + int surfaceHeight, + int64_t, + void* userData) { + auto* state = static_cast(userData); + ++state->mappedFrames; + const int sourceFrameIndex = state->displayFrameIndex++; + const int maxOutputFrames = maxCallbackOutputFrames(*state->options); + if (*state->encodedFrames >= maxOutputFrames) { + return; + } + const int expectedOutputFrames = state->oneFramePerMappedDisplayFrame + ? *state->encodedFrames + 1 + : expectedCallbackOutputFramesForSourceFrame( + sourceFrameIndex, + *state->options, + state->sourcePts); + if (*state->encodedFrames >= expectedOutputFrames) { + return; + } + + if (!*state->sink) { + *state->sink = std::make_unique( + state->context, + width, + height, + state->options->fps, + state->bitrate, + state->options->outputPath, + state->options->streamSync, + *state->options, + state->webcamCache, + state->cursorTrack, + state->zoomTrack); + } + + while (*state->encodedFrames < expectedOutputFrames && *state->encodedFrames < maxOutputFrames) { + const auto encodeStart = std::chrono::steady_clock::now(); + (*state->sink)->encodeFrame( + reinterpret_cast(dpSrcFrame), + static_cast(nSrcPitch), + width, + height, + surfaceHeight, + *state->encodedFrames); + const auto encodeEnd = std::chrono::steady_clock::now(); + *state->encodeMs += elapsedMs(encodeStart, encodeEnd); + ++*state->encodedFrames; + if (state->progress) { + const NvencSink* activeSink = state->sink && *state->sink ? state->sink->get() : nullptr; + reportEncodingProgress( + *state->encodedFrames, + maxOutputFrames, + *state->progress, + collectProgressCounters( + activeSink, + state->webcamCache, + state->decodeMs ? *state->decodeMs : 0.0, + state->encodeMs ? *state->encodeMs : 0.0)); + } + } +} + +double elapsedMs(std::chrono::steady_clock::time_point start, std::chrono::steady_clock::time_point end) { + return std::chrono::duration(end - start).count(); +} + +void reportEncodingProgress( + int encodedFrames, + int totalFrames, + ProgressReportState& state, + const ProgressCounters& counters, + bool force) { + if (totalFrames <= 0) { + return; + } + + const auto now = std::chrono::steady_clock::now(); + if (!force && encodedFrames < totalFrames && encodedFrames % 30 != 0 && elapsedMs(state.lastReportAt, now) < 500.0) { + return; + } + + const double percentage = std::min(100.0, std::max(0.0, static_cast(encodedFrames) * 100.0 / totalFrames)); + const double elapsedSeconds = std::max(elapsedMs(state.startedAt, now) / 1000.0, 0.001); + const double averageFps = encodedFrames > 0 ? static_cast(encodedFrames) / elapsedSeconds : 0.0; + const double intervalMs = std::max(elapsedMs(state.lastReportAt, now), 0.0); + const int intervalFrames = std::max(0, encodedFrames - state.lastReportedFrame); + const double instantFps = + intervalMs > 0.0 && intervalFrames > 0 ? static_cast(intervalFrames) * 1000.0 / intervalMs : 0.0; + const double intervalEncodeMs = std::max(0.0, counters.encodeMs - state.lastCounters.encodeMs); + const double intervalCompositeMs = std::max(0.0, counters.compositeMs - state.lastCounters.compositeMs); + const double intervalNvencMs = std::max(0.0, counters.nvencMs - state.lastCounters.nvencMs); + const double intervalPacketWriteMs = std::max(0.0, counters.packetWriteMs - state.lastCounters.packetWriteMs); + const double intervalWebcamDecodeMs = std::max(0.0, counters.webcamDecodeMs - state.lastCounters.webcamDecodeMs); + const double intervalWebcamCopyMs = std::max(0.0, counters.webcamCopyMs - state.lastCounters.webcamCopyMs); + const double intervalDecodeWallMs = std::max(0.0, counters.decodeWallMs - state.lastCounters.decodeWallMs); + const double intervalPipelineWaitMs = std::max(0.0, intervalMs - intervalEncodeMs); + const int intervalRoiCompositeFrames = + std::max(0, counters.roiCompositeFrames - state.lastCounters.roiCompositeFrames); + const int intervalMonolithicCompositeFrames = + std::max(0, counters.monolithicCompositeFrames - state.lastCounters.monolithicCompositeFrames); + const int intervalCopyCompositeFrames = + std::max(0, counters.copyCompositeFrames - state.lastCounters.copyCompositeFrames); + std::cerr << std::fixed << std::setprecision(2) + << "PROGRESS {\"currentFrame\":" << encodedFrames + << ",\"totalFrames\":" << totalFrames + << ",\"percentage\":" << percentage + << ",\"averageFps\":" << averageFps + << ",\"instantFps\":" << instantFps + << ",\"intervalMs\":" << intervalMs + << ",\"intervalFrames\":" << intervalFrames + << ",\"intervalDecodeWallMs\":" << intervalDecodeWallMs + << ",\"intervalEncodeMs\":" << intervalEncodeMs + << ",\"intervalPipelineWaitMs\":" << intervalPipelineWaitMs + << ",\"intervalCompositeMs\":" << intervalCompositeMs + << ",\"intervalNvencMs\":" << intervalNvencMs + << ",\"intervalPacketWriteMs\":" << intervalPacketWriteMs + << ",\"intervalWebcamDecodeMs\":" << intervalWebcamDecodeMs + << ",\"intervalWebcamCopyMs\":" << intervalWebcamCopyMs + << ",\"intervalRoiCompositeFrames\":" << intervalRoiCompositeFrames + << ",\"intervalMonolithicCompositeFrames\":" << intervalMonolithicCompositeFrames + << ",\"intervalCopyCompositeFrames\":" << intervalCopyCompositeFrames + << "}" << std::endl; + state.lastReportAt = now; + state.lastReportedFrame = encodedFrames; + state.lastCounters = counters; +} + +} // namespace + +int main(int argc, char** argv) { + try { + const Options options = parseOptions(argc, argv); + const uint32_t bitrate = static_cast(options.bitrateMbps) * 1000U * 1000U; + + checkCuda(cudaSetDevice(0), "cudaSetDevice"); + checkCu(cuInit(0), "cuInit"); + CUdevice device = 0; + checkCu(cuDeviceGet(&device, 0), "cuDeviceGet"); + CUcontext context = nullptr; + checkCu(cuCtxCreate(&context, 0, device), "cuCtxCreate"); + checkCu(cuCtxSetCurrent(context), "cuCtxSetCurrent"); + prewarmCuda(options.prewarmMs); + + std::ifstream input(options.inputPath, std::ios::binary); + if (!input) { + fail("Failed to open input: " + options.inputPath); + } + + std::unique_ptr webcamStream = createWebcamStreamDecoder(context, options); + const WebcamFrameCache* webcamCachePtr = webcamStream ? webcamStream->cache() : nullptr; + std::unique_ptr cursorTrack = loadCursorTrack(options); + const CursorTrack* cursorTrackPtr = cursorTrack.get(); + std::unique_ptr zoomTrack = loadZoomTrack(options); + const ZoomTrack* zoomTrackPtr = zoomTrack.get(); + const std::vector sourcePts = loadFramePts(options.sourcePtsPath); + const bool useSourcePts = + options.inputFrames > 0 && + sourcePts.size() >= static_cast(options.inputFrames); + auto decoder = std::make_unique(context, 0, 0, true, cudaVideoCodec_H264, nullptr, true, true); + std::unique_ptr sink; + const bool useDecoderFramePolicy = + options.inputFrames > 0 && + options.targetFrames > 0 && + options.inputFrames >= options.targetFrames && + !useSourcePts && + !options.postSelect; + FrameSelectionState selectionState{ + options.inputFrames, + options.targetFrames, + options.maxFrames, + 0, + 0, + options.fps, + useSourcePts ? &sourcePts : nullptr, + }; + if (useDecoderFramePolicy) { + decoder->SetDisplayFramePolicy(shouldCopyDisplayFrame, &selectionState); + } + + std::vector chunk(static_cast(options.chunkMb) * 1024 * 1024); + uint8_t** frames = nullptr; + int returnedFrames = 0; + int sourceFrames = 0; + int encodedFrames = 0; + const auto totalStart = std::chrono::steady_clock::now(); + double decodeMs = 0.0; + double encodeMs = 0.0; + ProgressReportState progressState; + progressState.startedAt = std::chrono::steady_clock::now(); + progressState.lastReportAt = progressState.startedAt; + const int progressTotalFrames = maxCallbackOutputFrames(options); + reportEncodingProgress(0, progressTotalFrames, progressState, ProgressCounters{}, true); + CallbackEncodeState callbackState{ + context, + &options, + bitrate, + &sink, + &decodeMs, + &encodeMs, + &encodedFrames, + 0, + webcamCachePtr, + cursorTrackPtr, + zoomTrackPtr, + &progressState, + useDecoderFramePolicy, + 0, + useSourcePts ? &sourcePts : nullptr, + }; + if (options.callbackEncode) { + decoder->SetMappedFrameHandler(encodeMappedDisplayFrame, &callbackState); + } + + auto prepareWebcamFrames = [&]() { + if (!webcamStream) { + return; + } + int outputFrameIndex = encodedFrames + kWebcamPrefetchOutputFrames; + if (options.maxFrames > 0) { + outputFrameIndex = std::min(outputFrameIndex, options.maxFrames - 1); + } + if (options.targetFrames > 0) { + outputFrameIndex = std::min(outputFrameIndex, options.targetFrames - 1); + } + webcamStream->ensureFrame(webcamFrameIndexForOutputFrame(outputFrameIndex, options)); + + const int keepFromOutputFrame = std::max(0, encodedFrames - 8); + webcamStream->dropBefore(webcamFrameIndexForOutputFrame(keepFromOutputFrame, options)); + }; + + while (input && shouldContinueEncoding(encodedFrames, options)) { + prepareWebcamFrames(); + input.read(reinterpret_cast(chunk.data()), static_cast(chunk.size())); + const int bytesRead = static_cast(input.gcount()); + if (bytesRead <= 0) { + break; + } + + const auto decodeStart = std::chrono::steady_clock::now(); + decoder->Decode(chunk.data(), bytesRead, &frames, &returnedFrames); + const auto decodeEnd = std::chrono::steady_clock::now(); + decodeMs += elapsedMs(decodeStart, decodeEnd); + + for (int index = 0; index < returnedFrames; ++index) { + if (!shouldContinueEncoding(encodedFrames, options)) { + break; + } + const int sourceFrameIndex = sourceFrames++; + if (!useDecoderFramePolicy && !shouldEncodeFrame(sourceFrameIndex, encodedFrames, options)) { + continue; + } + if (!sink) { + sink = std::make_unique( + context, + decoder->GetWidth(), + decoder->GetHeight(), + options.fps, + bitrate, + options.outputPath, + options.streamSync, + options, + webcamCachePtr, + cursorTrackPtr, + zoomTrackPtr); + } + const auto encodeStart = std::chrono::steady_clock::now(); + sink->encodeFrame( + frames[index], + decoder->GetDeviceFramePitch(), + decoder->GetWidth(), + decoder->GetHeight(), + decoder->GetHeight(), + encodedFrames); + const auto encodeEnd = std::chrono::steady_clock::now(); + encodeMs += elapsedMs(encodeStart, encodeEnd); + ++encodedFrames; + reportEncodingProgress( + encodedFrames, + progressTotalFrames, + progressState, + collectProgressCounters(sink.get(), webcamCachePtr, decodeMs, encodeMs)); + } + } + + if (shouldContinueEncoding(encodedFrames, options)) { + const auto decodeStart = std::chrono::steady_clock::now(); + decoder->Decode(nullptr, 0, &frames, &returnedFrames); + const auto decodeEnd = std::chrono::steady_clock::now(); + decodeMs += elapsedMs(decodeStart, decodeEnd); + for (int index = 0; index < returnedFrames; ++index) { + if (!shouldContinueEncoding(encodedFrames, options)) { + break; + } + const int sourceFrameIndex = sourceFrames++; + if (!useDecoderFramePolicy && !shouldEncodeFrame(sourceFrameIndex, encodedFrames, options)) { + continue; + } + if (!sink) { + sink = std::make_unique( + context, + decoder->GetWidth(), + decoder->GetHeight(), + options.fps, + bitrate, + options.outputPath, + options.streamSync, + options, + webcamCachePtr, + cursorTrackPtr, + zoomTrackPtr); + } + const auto encodeStart = std::chrono::steady_clock::now(); + sink->encodeFrame( + frames[index], + decoder->GetDeviceFramePitch(), + decoder->GetWidth(), + decoder->GetHeight(), + decoder->GetHeight(), + encodedFrames); + const auto encodeEnd = std::chrono::steady_clock::now(); + encodeMs += elapsedMs(encodeStart, encodeEnd); + ++encodedFrames; + reportEncodingProgress( + encodedFrames, + progressTotalFrames, + progressState, + collectProgressCounters(sink.get(), webcamCachePtr, decodeMs, encodeMs)); + } + } + + if (!sink) { + fail("No decoded frames were produced"); + } + const auto flushStart = std::chrono::steady_clock::now(); + sink->finish(); + const auto flushEnd = std::chrono::steady_clock::now(); + const auto totalEnd = std::chrono::steady_clock::now(); + reportEncodingProgress( + encodedFrames, + progressTotalFrames, + progressState, + collectProgressCounters(sink.get(), webcamCachePtr, decodeMs, encodeMs), + true); + + const double totalMs = elapsedMs(totalStart, totalEnd); + const double mediaMs = static_cast(encodedFrames) * 1000.0 / options.fps; + const double measuredFps = static_cast(encodedFrames) / (totalMs / 1000.0); + const double realtime = mediaMs / totalMs; + const int reportedSourceFrames = + useDecoderFramePolicy ? selectionState.sourceFrames : + (options.callbackEncode ? decoder->GetDisplayFrameCount() : sourceFrames); + const int mappedDisplayFrames = options.callbackEncode ? callbackState.mappedFrames : encodedFrames; + const int selectedDisplayFrames = + useDecoderFramePolicy ? selectionState.selectedFrames : mappedDisplayFrames; + const int skippedDisplayFrames = + useDecoderFramePolicy ? std::max(0, selectionState.sourceFrames - selectionState.selectedFrames) : 0; + const double decodeOnlyApproxMs = options.callbackEncode ? std::max(0.0, decodeMs - encodeMs) : decodeMs; + + std::cout << std::fixed << std::setprecision(2) + << "{" + << "\"success\":true," + << "\"mode\":\"nvdec-cuda-nvenc-annexb\"," + << "\"selectionStage\":\"" + << (options.callbackEncode + ? (useDecoderFramePolicy ? "decoder-policy-mapped-callback" : "mapped-callback") + : (useDecoderFramePolicy ? "decoder" : "post")) + << "\"," + << "\"sourceTimestampMode\":\"" << (useSourcePts ? "pts" : "ordinal") << "\"," + << "\"syncMode\":\"" << (options.streamSync ? "stream" : "device") << "\"," + << "\"prewarmMs\":" << options.prewarmMs << "," + << "\"chunkMb\":" << options.chunkMb << "," + << "\"width\":" << decoder->GetWidth() << "," + << "\"height\":" << decoder->GetHeight() << "," + << "\"fps\":" << options.fps << "," + << "\"staticLayout\":" << (hasStaticLayout(options) ? "true" : "false") << "," + << "\"contentX\":" << options.contentX << "," + << "\"contentY\":" << options.contentY << "," + << "\"contentWidth\":" << options.contentWidth << "," + << "\"contentHeight\":" << options.contentHeight << "," + << "\"radius\":" << options.radius << "," + << "\"backgroundImage\":" << (!options.backgroundNv12Path.empty() ? "true" : "false") << "," + << "\"shadowOffsetY\":" << options.shadowOffsetY << "," + << "\"shadowIntensityPct\":" << options.shadowIntensityPct << "," + << "\"webcamOverlay\":" << (hasWebcamOverlay(options) ? "true" : "false") << "," + << "\"webcamX\":" << options.webcamX << "," + << "\"webcamY\":" << options.webcamY << "," + << "\"webcamSize\":" << options.webcamSize << "," + << "\"webcamRadius\":" << options.webcamRadius << "," + << "\"webcamStream\":" << (!options.webcamAnnexbPath.empty() ? "true" : "false") << "," + << "\"webcamMirror\":" << (options.webcamMirror ? "true" : "false") << "," + << "\"webcamCachedFrames\":" << (webcamCachePtr ? webcamCachePtr->frames.size() : 0) << "," + << "\"webcamPeakCachedFrames\":" << (webcamCachePtr ? webcamCachePtr->peakFrames : 0) << "," + << "\"webcamCacheBaseFrame\":" << (webcamCachePtr ? webcamCachePtr->baseFrameIndex : 0) << "," + << "\"webcamDecodedFrames\":" << (webcamCachePtr ? webcamCachePtr->decodedFrames : 0) << "," + << "\"webcamDecodeMs\":" << (webcamCachePtr ? webcamCachePtr->decodeMs : 0.0) << "," + << "\"webcamCopyMs\":" << (webcamCachePtr ? webcamCachePtr->copyMs : 0.0) << "," + << "\"cursorOverlay\":" << (cursorTrackPtr ? "true" : "false") << "," + << "\"cursorSamples\":" << (cursorTrackPtr ? cursorTrackPtr->samples.size() : 0) << "," + << "\"cursorHeight\":" << options.cursorHeight << "," + << "\"cursorAtlas\":" << (!options.cursorAtlasRgbaPath.empty() ? "true" : "false") << "," + << "\"zoomOverlay\":" << (zoomTrackPtr ? "true" : "false") << "," + << "\"zoomSamples\":" << (zoomTrackPtr ? zoomTrackPtr->samples.size() : 0) << "," + << "\"sourceFrames\":" << reportedSourceFrames << "," + << "\"mappedDisplayFrames\":" << mappedDisplayFrames << "," + << "\"selectedDisplayFrames\":" << selectedDisplayFrames << "," + << "\"skippedDisplayFrames\":" << skippedDisplayFrames << "," + << "\"frames\":" << encodedFrames << "," + << "\"totalMs\":" << totalMs << "," + << "\"decodeMs\":" << decodeOnlyApproxMs << "," + << "\"decodeWallMs\":" << decodeMs << "," + << "\"encodeMs\":" << encodeMs << "," + << "\"compositeMs\":" << sink->compositeMs() << "," + << "\"roiCompositeFrames\":" << sink->roiCompositeFrames() << "," + << "\"monolithicCompositeFrames\":" << sink->monolithicCompositeFrames() << "," + << "\"copyCompositeFrames\":" << sink->copyCompositeFrames() << "," + << "\"nvencMs\":" << sink->nvencMs() << "," + << "\"packetWriteMs\":" << sink->packetWriteMs() << "," + << "\"flushMs\":" << elapsedMs(flushStart, flushEnd) << "," + << "\"measuredFps\":" << measuredFps << "," + << "\"realtimeMultiplier\":" << realtime << "," + << "\"outputBytes\":" << sink->outputBytes() << "," + << "\"outputPath\":\"" << options.outputPath << "\"" + << "}" << std::endl; + + sink.reset(); + decoder.reset(); + webcamStream.reset(); + checkCu(cuCtxDestroy(context), "cuCtxDestroy"); + return 0; + } catch (const std::exception& error) { + std::cerr << "{\"success\":false,\"error\":\"" << error.what() << "\"}" << std::endl; + return 1; + } +} diff --git a/package.json b/package.json index aee7e919..f2dcda9f 100644 --- a/package.json +++ b/package.json @@ -25,8 +25,9 @@ "rebuild:native": "node ./node_modules/@electron/rebuild/lib/cli.js --force --only uiohook-napi", "build:native-helpers": "node scripts/build-native-helpers.mjs", "build:whisper-runtime": "node scripts/build-whisper-runtime.mjs", - "build:platform-native-helpers": "npm run build:native-helpers && npm run build:windows-capture && npm run build:windows-gpu-export && npm run build:cursor-monitor && npm run build:whisper-runtime", + "build:platform-native-helpers": "npm run build:native-helpers && npm run build:windows-capture && npm run build:windows-gpu-export && npm run build:nvidia-cuda-compositor && npm run build:cursor-monitor && npm run build:whisper-runtime", "build:windows-gpu-export": "node scripts/build-windows-gpu-export.mjs", + "build:nvidia-cuda-compositor": "node scripts/build-nvidia-cuda-compositor.mjs", "build:windows-capture": "node scripts/build-windows-capture.mjs", "build:cursor-monitor": "node scripts/build-cursor-monitor.mjs", "build:mac": "npm run build:platform-native-helpers && tsc && vite build --config vite.config.ts && electron-builder --mac", diff --git a/scripts/build-nvidia-cuda-compositor.mjs b/scripts/build-nvidia-cuda-compositor.mjs new file mode 100644 index 00000000..cdccbc12 --- /dev/null +++ b/scripts/build-nvidia-cuda-compositor.mjs @@ -0,0 +1,193 @@ +import { execSync } from "node:child_process"; +import { copyFileSync, existsSync, mkdirSync, rmSync } from "node:fs"; +import path from "node:path"; + +import { + formatNativeHelperManifestWarning, + updateNativeHelperManifest, + verifyNativeHelperManifest, +} from "./native-helper-manifest.mjs"; + +const projectRoot = process.cwd(); +const sourceDir = path.join(projectRoot, "electron", "native", "nvidia-cuda-compositor"); +const buildDir = path.join(sourceDir, "build"); +const bundledDir = path.join( + projectRoot, + "electron", + "native", + "bin", + process.arch === "arm64" ? "win32-arm64" : "win32-x64", +); +const bundledExePath = path.join(bundledDir, "recordly-nvidia-cuda-compositor.exe"); +const helperId = "recordly-nvidia-cuda-compositor"; +const generatorArch = process.arch === "arm64" ? "ARM64" : "x64"; +const videoCodecSdkRoot = + process.env.RECORDLY_NVIDIA_VIDEO_CODEC_SDK_ROOT?.trim() || + path.join(projectRoot, ".tmp", "video-sdk-samples"); + +if (process.platform !== "win32") { + console.log("[build-nvidia-cuda-compositor] Skipping NVIDIA CUDA compositor build."); + process.exit(0); +} + +if (!existsSync(path.join(sourceDir, "CMakeLists.txt"))) { + console.error("[build-nvidia-cuda-compositor] CMakeLists.txt not found at", sourceDir); + process.exit(1); +} + +function fallbackToBundledHelperOrExit(reason) { + if (existsSync(bundledExePath)) { + const verification = verifyNativeHelperManifest({ + projectRoot, + helperId, + sourceDir, + binaryPath: bundledExePath, + binaryName: "recordly-nvidia-cuda-compositor.exe", + }); + if (!verification.ok) { + console.warn( + formatNativeHelperManifestWarning("build-nvidia-cuda-compositor", verification), + ); + } + console.log(`[build-nvidia-cuda-compositor] ${reason}`); + console.log(`[build-nvidia-cuda-compositor] Using bundled helper: ${bundledExePath}`); + process.exit(0); + } + + console.error(`[build-nvidia-cuda-compositor] ${reason}`); + console.error( + "[build-nvidia-cuda-compositor] No bundled helper is available; install CUDA Toolkit + NVIDIA Video Codec SDK or provide a staged helper.", + ); + process.exit(1); +} + +function findCmake() { + try { + execSync("cmake --version", { stdio: "pipe" }); + return "cmake"; + } catch { + // Continue probing common Windows install locations. + } + + const standaloneCmakePaths = [ + path.join("C:", "Program Files", "CMake", "bin", "cmake.exe"), + path.join("C:", "Program Files (x86)", "CMake", "bin", "cmake.exe"), + ]; + for (const cmakePath of standaloneCmakePaths) { + if (existsSync(cmakePath)) { + return `"${cmakePath}"`; + } + } + + const vsRoots = [ + path.join("C:", "Program Files", "Microsoft Visual Studio"), + path.join("C:", "Program Files (x86)", "Microsoft Visual Studio"), + ]; + const vsEditions = ["Preview", "Community", "Professional", "Enterprise", "BuildTools"]; + const vsVersions = ["2022", "2019"]; + for (const root of vsRoots) { + for (const version of vsVersions) { + for (const edition of vsEditions) { + const cmakePath = path.join( + root, + version, + edition, + "Common7", + "IDE", + "CommonExtensions", + "Microsoft", + "CMake", + "CMake", + "bin", + "cmake.exe", + ); + if (existsSync(cmakePath)) { + return `"${cmakePath}"`; + } + } + } + } + + return null; +} + +if (!existsSync(path.join(videoCodecSdkRoot, "Samples", "NvCodec"))) { + fallbackToBundledHelperOrExit( + `NVIDIA Video Codec SDK samples not found at ${videoCodecSdkRoot}. Set RECORDLY_NVIDIA_VIDEO_CODEC_SDK_ROOT to build from source.`, + ); +} + +const cmake = findCmake(); +if (!cmake) { + fallbackToBundledHelperOrExit( + "CMake not found. Install Visual Studio with C++ CMake tools or standalone CMake.", + ); +} + +mkdirSync(buildDir, { recursive: true }); + +function clearCmakeCache() { + rmSync(path.join(buildDir, "CMakeCache.txt"), { force: true }); + rmSync(path.join(buildDir, "CMakeFiles"), { recursive: true, force: true }); +} + +console.log("[build-nvidia-cuda-compositor] Configuring CMake..."); +try { + clearCmakeCache(); + execSync( + `${cmake} .. -G "Visual Studio 17 2022" -A ${generatorArch} -DRECORDLY_NVIDIA_VIDEO_CODEC_SDK_ROOT="${videoCodecSdkRoot}"`, + { + cwd: buildDir, + stdio: "inherit", + timeout: 120000, + }, + ); +} catch { + console.log("[build-nvidia-cuda-compositor] VS 2022 generator not found, trying VS 2019..."); + try { + clearCmakeCache(); + execSync( + `${cmake} .. -G "Visual Studio 16 2019" -A ${generatorArch} -DRECORDLY_NVIDIA_VIDEO_CODEC_SDK_ROOT="${videoCodecSdkRoot}"`, + { + cwd: buildDir, + stdio: "inherit", + timeout: 120000, + }, + ); + } catch (error) { + fallbackToBundledHelperOrExit( + `CMake configure failed: ${error instanceof Error ? error.message : String(error)}`, + ); + } +} + +console.log("[build-nvidia-cuda-compositor] Building NVIDIA CUDA compositor..."); +try { + execSync(`${cmake} --build . --config Release`, { + cwd: buildDir, + stdio: "inherit", + timeout: 300000, + }); +} catch (error) { + fallbackToBundledHelperOrExit( + `Build failed: ${error instanceof Error ? error.message : String(error)}`, + ); +} + +const exePath = path.join(buildDir, "Release", "recordly-nvidia-cuda-compositor.exe"); +if (!existsSync(exePath)) { + console.error("[build-nvidia-cuda-compositor] Expected exe not found at", exePath); + process.exit(1); +} + +mkdirSync(bundledDir, { recursive: true }); +copyFileSync(exePath, bundledExePath); +console.log(`[build-nvidia-cuda-compositor] Staged bundled helper: ${bundledExePath}`); +const manifestPath = updateNativeHelperManifest({ + projectRoot, + helperId, + sourceDir, + binaryPath: bundledExePath, + binaryName: "recordly-nvidia-cuda-compositor.exe", +}); +console.log(`[build-nvidia-cuda-compositor] Updated helper manifest: ${manifestPath}`); diff --git a/scripts/launch-recordly-cuda-auto.ps1 b/scripts/launch-recordly-cuda-auto.ps1 new file mode 100644 index 00000000..02d3dff2 --- /dev/null +++ b/scripts/launch-recordly-cuda-auto.ps1 @@ -0,0 +1,80 @@ +param( + [string]$AppPath, + [string]$CudaScriptPath, + [switch]$CloseExisting, + [switch]$NoDiagnostics, + [switch]$AllowCudaAudio +) + +$ErrorActionPreference = "Stop" + +$repoRoot = Resolve-Path (Join-Path $PSScriptRoot "..") + +if (-not $AppPath) { + $packagedAppPath = Join-Path $repoRoot "release\win-unpacked\Recordly.exe" + $installedAppPath = Join-Path $env:LOCALAPPDATA "Programs\recordly\Recordly.exe" + if (Test-Path $packagedAppPath) { + $AppPath = $packagedAppPath + } elseif (Test-Path $installedAppPath) { + $AppPath = $installedAppPath + } else { + throw "Recordly.exe was not found. Build the Windows package first or pass -AppPath." + } +} + +if (-not $CudaScriptPath) { + $CudaScriptPath = Join-Path $repoRoot "electron\native\nvidia-cuda-compositor\run-mp4-pipeline.mjs" +} + +if (-not (Test-Path $AppPath)) { + throw "Recordly app not found: $AppPath" +} + +if (-not (Test-Path $CudaScriptPath)) { + throw "NVIDIA CUDA/NVENC wrapper script not found: $CudaScriptPath" +} + +$existingRecordly = @(Get-Process -Name "Recordly" -ErrorAction SilentlyContinue) +if ($existingRecordly.Count -gt 0) { + if (-not $CloseExisting) { + Write-Host "Recordly is already running. Close it first, or rerun with -CloseExisting so the CUDA env is inherited by the new app process." + $existingRecordly | Select-Object Id, ProcessName, Path | Format-Table -AutoSize + exit 2 + } + + $existingRecordly | Stop-Process -Force + Start-Sleep -Milliseconds 500 +} + +$env:RECORDLY_EXPERIMENTAL_NVIDIA_CUDA_EXPORT = "1" +$env:RECORDLY_NVIDIA_CUDA_EXPORT_SCRIPT = (Resolve-Path $CudaScriptPath).Path +$env:RECORDLY_NVIDIA_CUDA_EXPORT_HIGH_PRIORITY = "1" +$env:RECORDLY_NVIDIA_CUDA_SAMPLE_GPU = "1" + +if ($AllowCudaAudio) { + $env:RECORDLY_NVIDIA_CUDA_ALLOW_AUDIO_EXPORT = "1" +} else { + Remove-Item Env:\RECORDLY_NVIDIA_CUDA_ALLOW_AUDIO_EXPORT -ErrorAction SilentlyContinue +} + +if ($NoDiagnostics) { + Remove-Item Env:\RECORDLY_NVIDIA_CUDA_EXPORT_DIAGNOSTICS -ErrorAction SilentlyContinue +} else { + $env:RECORDLY_NVIDIA_CUDA_EXPORT_DIAGNOSTICS = "1" +} + +$resolvedAppPath = (Resolve-Path $AppPath).Path +$appDirectory = Split-Path $resolvedAppPath -Parent +if ($AllowCudaAudio) { + $cudaAudioMode = "candidate only; app still requires timestamp-aligned CUDA output" +} else { + $cudaAudioMode = "guarded; audio exports fall back to Windows D3D11" +} + +Write-Host "Launching Recordly with guarded NVIDIA CUDA/NVENC auto export enabled:" +Write-Host " App: $resolvedAppPath" +Write-Host " CUDA wrapper: $env:RECORDLY_NVIDIA_CUDA_EXPORT_SCRIPT" +Write-Host " CUDA audio exports: $cudaAudioMode" +Write-Host " Diagnostics: $($env:RECORDLY_NVIDIA_CUDA_EXPORT_DIAGNOSTICS -eq '1')" + +Start-Process -FilePath $resolvedAppPath -WorkingDirectory $appDirectory diff --git a/scripts/smoke-packaged-binaries.mjs b/scripts/smoke-packaged-binaries.mjs index 745f787b..97e5eb65 100644 --- a/scripts/smoke-packaged-binaries.mjs +++ b/scripts/smoke-packaged-binaries.mjs @@ -171,6 +171,11 @@ function getExpectedNativeHelperFiles(archTag) { label: "Windows GPU export helper", executable: true, }, + { + name: "recordly-nvidia-cuda-compositor.exe", + label: "NVIDIA CUDA compositor helper", + executable: true, + }, { name: "helpers-manifest.json", label: "Windows helper manifest" }, { name: "whisper-cli.exe", label: "Whisper CLI runtime", executable: true }, { name: "whisper-runtime.json", label: "Whisper runtime manifest" }, diff --git a/src/hooks/useScreenRecorder.ts b/src/hooks/useScreenRecorder.ts index 1de21039..0f12fca3 100644 --- a/src/hooks/useScreenRecorder.ts +++ b/src/hooks/useScreenRecorder.ts @@ -731,16 +731,7 @@ export function useScreenRecorder(): UseScreenRecorderReturn { await window.electronAPI.muxNativeWindowsRecording(pauseSegments); if (!muxResult?.success || !muxResult.path) { void logNativeCaptureDiagnostics("mux-native-windows-recording"); - if (!muxResult?.path) { - const failureMessage = await buildNativeCaptureFailureMessage( - "mux-native-windows-recording", - muxResult?.message || - "Failed to finalize the Windows recording, so the editor was not opened.", - ); - await notifyRecordingFinalizationFailure(failureMessage); - return; - } - + const fallbackPath = muxResult?.path ?? finalPath; const warningMessage = muxResult?.error || muxResult?.message || @@ -749,8 +740,10 @@ export function useScreenRecorder(): UseScreenRecorderReturn { `${warningMessage}. Recording was saved, but audio playback or export may be incomplete.`, { id: SOURCE_AUDIO_MUX_TOAST_ID, duration: 10000 }, ); + finalPath = fallbackPath; + } else { + finalPath = muxResult.path; } - finalPath = muxResult.path; } await storeMicrophoneSidecar( @@ -1186,7 +1179,7 @@ export function useScreenRecorder(): UseScreenRecorderReturn { maxFrameRate: TARGET_FRAME_RATE, }, }, - } as any); + } as unknown as MediaStreamConstraints); } }; diff --git a/src/lib/exporter/modernVideoExporter.ts b/src/lib/exporter/modernVideoExporter.ts index 17a2f99b..bb374f12 100644 --- a/src/lib/exporter/modernVideoExporter.ts +++ b/src/lib/exporter/modernVideoExporter.ts @@ -1601,8 +1601,14 @@ export class ModernVideoExporter { return; } - const nativeCurrentFrame = Math.floor(progress.currentFrame); const nativeTotalFrames = Math.max(1, Math.floor(progress.totalFrames)); + const progressPercentFrame = Number.isFinite(progress.percentage) + ? Math.floor((nativeTotalFrames * progress.percentage) / 100) + : 0; + const nativeCurrentFrame = Math.max( + Math.floor(progress.currentFrame), + progressPercentFrame, + ); const nativeFramesComplete = nativeCurrentFrame >= nativeTotalFrames; const maxExtractingFrame = Math.max( 0,