diff --git a/CMakeLists.txt b/CMakeLists.txt index 96c996c4c9..561fc7c3d9 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -73,6 +73,7 @@ option(BUILD_PLUGIN_WUJI_GLOVE "Build Wuji glove plugin (requires the wuji_sdk C option(BUILD_PLUGIN_SHARPA_AVATAR "Build Sharpa AvatarExo Glove plugin (Linux; skips if the Avatar SDK is missing)" OFF) option(BUILD_PLUGIN_SENSING "Build SENSING GMSL camera plugin (Jetson only; requires SIPL and NVENC)" OFF) option(BUILD_VIDEO_STAGE "Build the video stage: the Orin-to-PC video transport and the PC relay" OFF) +option(BUILD_VIDEO_STAGE_GPU "Add the optional GPU stage to the video stage (needs a CUDA toolkit; its tests need a CUDA GPU)" OFF) option(BUILD_EXAMPLES "Build examples" ON) option(BUILD_EXAMPLE_TELEOP_ROS2 "Build only the teleop_ros2 ROS 2 reference integration (e.g. for Docker)" OFF) option(BUILD_TESTING "Build unit tests" ON) diff --git a/src/video_stage/CMakeLists.txt b/src/video_stage/CMakeLists.txt index d7bb15b103..e51853431d 100644 --- a/src/video_stage/CMakeLists.txt +++ b/src/video_stage/CMakeLists.txt @@ -2,20 +2,53 @@ # SPDX-License-Identifier: Apache-2.0 # ============================================================================== -# Video stage: the Orin-to-PC video transport and the PC relay to LiveKit +# Video stage: the Orin-to-PC video transport, the PC relay to LiveKit, and an optional PC GPU stage # ============================================================================== # wire/ video_stage_wire -- AVS1 records, clock exchange, H.264 access-unit helpers, frame correlator # sender/ video_stage_sender -- the Orin-side sender of AVS1 records over a direct link # app/ video_stage_worker -- the PC relay: receives the Orin's records and publishes them to LiveKit or a file # -# No CUDA. The LiveKit SDK (src/livekit_publisher) is needed only by the relay. +# core/ video_stage_core -- (GPU) the operation interface and the example sine-wave warp +# codec/ video_stage_codec -- (GPU) NVDEC decoder and NVENC encoder +# +# No CUDA unless BUILD_VIDEO_STAGE_GPU is ON: the relay forwards the Orin's access units untouched. The LiveKit SDK +# (src/livekit_publisher) is needed only by the relay. # # Configure with: -# cmake -B build -DBUILD_VIDEO_STAGE=ON +# cmake -B build -DBUILD_VIDEO_STAGE=ON [-DBUILD_VIDEO_STAGE_GPU=ON] # ============================================================================== add_subdirectory(wire) add_subdirectory(sender) + +if(BUILD_VIDEO_STAGE_GPU) + enable_language(CUDA) + + if(NOT DEFINED CMAKE_CUDA_ARCHITECTURES) + # Real code for the common desktop architectures plus PTX for 90, which newer GPUs JIT-compile. + set(CMAKE_CUDA_ARCHITECTURES "80-real;86-real;89-real;90") + endif() + + find_package(CUDAToolkit REQUIRED) + + # NVDEC/NVENC API headers (MIT licence) with run-time loaders for libnvcuvid and libnvidia-encode, pinned by hash. + include(FetchContent) + FetchContent_Declare( + nv_codec_headers + URL https://github.com/FFmpeg/nv-codec-headers/archive/refs/tags/n12.2.72.0.tar.gz + URL_HASH SHA256=dbeaec433d93b850714760282f1d0992b1254fc3b5a6cb7d76fc1340a1e47563 + ) + FetchContent_GetProperties(nv_codec_headers) + if(NOT nv_codec_headers_POPULATED) + FetchContent_Populate(nv_codec_headers) + endif() + add_library(video_stage_nv_codec_headers INTERFACE) + target_include_directories(video_stage_nv_codec_headers SYSTEM INTERFACE "${nv_codec_headers_SOURCE_DIR}/include") + + add_subdirectory(core) + add_subdirectory(codec) +endif() + add_subdirectory(app) if(BUILD_TESTING) diff --git a/src/video_stage/README.md b/src/video_stage/README.md index 120aa12eab..dc7df3e4d3 100644 --- a/src/video_stage/README.md +++ b/src/video_stage/README.md @@ -6,8 +6,8 @@ SPDX-License-Identifier: Apache-2.0 # Video stage Moves camera video from the Orin to the PC and on to LiveKit. The Orin sends its encoded H.264 access units to the PC -over a direct link; the PC relay (`video_stage_worker`) hands them, untouched, to LiveKit or a file. There is no decode, -encode or GPU in the relay, and it builds without CUDA. +over a direct link; the PC relay (`video_stage_worker`) hands them, untouched, to LiveKit or a file. By default there is no decode, +encode or GPU in the relay, and it builds without CUDA; an optional GPU stage is described below. ```text Orin producer --DATA (AVS1 records, H.264)--> PC relay --> LiveKit track (or an Annex-B file) @@ -62,6 +62,49 @@ hello order, retry, keyframe start, bounded queue with drop-to-keyframe, keyfram offset clock), and feeds it Jetson-style H.264 (no AUDs, IDR every 5 s, constant bitrate). One `--port` per camera. Run it against a relay writing a file, then check the file with `ffprobe`. +## The GPU stage (optional) + +Built only with `-DBUILD_VIDEO_STAGE_GPU=ON` (needs a CUDA toolkit; the tests need a CUDA GPU). `video_stage_worker --operation=sine-warp` then decodes each picture on NVDEC, keeps it in GPU memory through one GPU operation, and re-encodes it on NVENC, instead of forwarding the Orin's access units untouched. A build without it rejects `--operation=sine-warp` with a clear message and relays exactly as before. + +### The operation contract + +`core/inc/video_stage/gpu_operation.hpp` defines `GpuOperation`: NV12 in, NV12 out, fixed geometry. + +- `prepare(input_geometry)` runs once, outside the frame loop. Build and upload anything per-geometry here (a + rectification map, for example). Nothing in `enqueue` allocates or uploads. +- `enqueue(in, out, stream)` only queues kernels on `stream` and returns. It does not synchronise the device or the + stream. After it returns, the caller records a CUDA event on `stream`; whoever reads `out` waits on that event. +- `in` and `out` are different surfaces (the operation is not in-place). `out` has `output_geometry(in.geometry)` and + keeps the input's color description. Both are `Nv12View`s: device pointers, byte pitches, size. +- Frame identity and timestamps are not the operation's concern; they travel with the frame. + +### The example: `SineWarp` + +`core/sine_warp.cu` is a horizontal sine ripple, `source_x = x + A*sin(2*pi*y/wavelength + phase)`, +`source_y = y`, sampled bilinearly, with the border fill for the declared range (Y=16 limited, 0 full; chroma 128). +Amplitude 0 is an identity. It is an example, not calibrated rectification: the warped picture does not keep the +source camera's projection. + +The stage assumes the decoder's NV12 uses the H.264 default chroma siting described below; a stream that signals a +different siting must be rejected at the decoder, not warped with the wrong sample positions. + +Chroma is not sampled at the luma coordinate divided by two. A chroma sample of H.264 NV12 sits on an even luma column +and half way between two luma rows, so the kernel maps that luma-space position, then converts back to chroma-plane +coordinates. Keep that when substituting an operation. + +### Substituting the fisheye rectification + +Replace the mapping, keep the rest: a class deriving from `GpuOperation` whose `prepare` generates the destination-to- +source maps once (on the CPU, uploaded once) and whose `enqueue` launches a kernel that reads `map_x/map_y` instead of +computing the sine. Sample luma at the map position and chroma at the position derived from the same map with the +siting above. `output_geometry` returns the rectified output size. Nothing else in the stage changes. + +### Identity and overload + +Each picture carries a key through NVDEC and NVENC, and the source frame number and stamp are looked up by that key, so +metadata follows its own frame through loss. A picture arriving while every encoder input surface is busy is dropped and +counted, never queued. + ## Build ```bash @@ -70,4 +113,6 @@ cmake -B build -DBUILD_VIDEO_STAGE=ON -DBUILD_PLUGINS=OFF -DBUILD_VIZ=OFF -DBUIL cmake --build build && ctest --test-dir build -R 'video_stage|avs1|clock' ``` +Add `-DBUILD_VIDEO_STAGE_GPU=ON` for the GPU stage and its tests. + The LiveKit SDK is fetched, or taken from `LIVEKIT_SDK_DIR`. diff --git a/src/video_stage/app/CMakeLists.txt b/src/video_stage/app/CMakeLists.txt index 0d29961af5..5703571b5a 100644 --- a/src/video_stage/app/CMakeLists.txt +++ b/src/video_stage/app/CMakeLists.txt @@ -1,8 +1,8 @@ # SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. # SPDX-License-Identifier: Apache-2.0 -# The PC relay: receives AVS1 access units from the Orin over its direct link and hands them, untouched, to LiveKit or -# a file. +# The PC worker: receives AVS1 access units from the Orin and hands them, untouched, to LiveKit or a file. With +# BUILD_VIDEO_STAGE_GPU it can also decode on NVDEC, apply a GPU operation and re-encode on NVENC first. add_executable(video_stage_worker main.cpp control_channel.cpp @@ -16,3 +16,9 @@ target_link_libraries(video_stage_worker video_stage::wire sensing::livekit ) + +if(BUILD_VIDEO_STAGE_GPU) + target_sources(video_stage_worker PRIVATE pipeline.cpp) + target_compile_definitions(video_stage_worker PRIVATE VIDEO_STAGE_GPU) + target_link_libraries(video_stage_worker PRIVATE video_stage::core video_stage::codec) +endif() diff --git a/src/video_stage/app/main.cpp b/src/video_stage/app/main.cpp index 4088c4e677..73d71426f7 100644 --- a/src/video_stage/app/main.cpp +++ b/src/video_stage/app/main.cpp @@ -4,6 +4,10 @@ #include "livekit_output.hpp" #include "stream_worker.hpp" +#ifdef VIDEO_STAGE_GPU +# include +#endif + #include #include #include @@ -32,11 +36,13 @@ void usage(const char* program) std::cout << "Usage: " << program << " --stream=PORT:TRACK ... --livekit-url=URL --livekit-token-file=PATH\n" << " " << program << " --stream=PORT:file=PATH.h264 ...\n" - << " [--bitrate=BPS] [--max-upstream-uncertainty-ms=MS] [--stats-interval=S] [--once]\n" + << " [--operation=none|sine-warp] [--amplitude=PX] [--bitrate=BPS] [--idr-period=N] [--max-upstream-uncertainty-ms=MS] [--stats-interval=S] [--once]\n" << "One --stream per camera. Each listens on PORT for the Orin's direct connection (AVS1 H.264 records plus a\n" << "control link) and publishes the access units as LiveKit track TRACK (all tracks under one participant) or\n" - << "writes them to a file. The Orin's access units are forwarded untouched, with no decode, encode or GPU. Serves\n" - << "one Orin session after another unless --once. Prints a summary at the end.\n"; + << "writes them to a file. By default (--operation=none) the Orin's access units are forwarded untouched, with no\n" + << "decode, encode or GPU. --operation=sine-warp decodes on NVDEC, applies the example sine warp and re-encodes on\n" + << "NVENC, as a test of the GPU stage (a build with -DBUILD_VIDEO_STAGE_GPU=ON). Serves one Orin\n" + << "session after another unless --once. Prints per-stage percentiles at the end.\n"; } std::string read_token(const std::string& path) @@ -87,8 +93,33 @@ try stats_interval_s = std::stoi(arg.substr(17)); else if (arg == "--once") settings.once = true; + else if (arg.rfind("--operation=", 0) == 0) + { + const std::string name = arg.substr(12); + if (name == "none") + settings.operation = Operation::None; + else if (name == "sine-warp") + settings.operation = Operation::SineWarp; + else + { + std::cerr << "--operation must be none or sine-warp\n"; + return 1; + } + } else if (arg.rfind("--bitrate=", 0) == 0) - settings.livekit.max_bitrate_bps = static_cast(std::stoul(arg.substr(10))); + { + const auto bitrate = static_cast(std::stoul(arg.substr(10))); + settings.livekit.max_bitrate_bps = bitrate; +#ifdef VIDEO_STAGE_GPU + settings.pipeline.bitrate_bps = bitrate; +#endif + } +#ifdef VIDEO_STAGE_GPU + else if (arg.rfind("--amplitude=", 0) == 0) + settings.warp.amplitude_px = std::stof(arg.substr(12)); + else if (arg.rfind("--idr-period=", 0) == 0) + settings.pipeline.idr_period = static_cast(std::stoul(arg.substr(13))); +#endif else if (arg.rfind("--max-upstream-uncertainty-ms=", 0) == 0) settings.livekit.max_upstream_uncertainty_ns = static_cast(std::stod(arg.substr(30)) * 1e6); else @@ -111,14 +142,36 @@ try return 1; } + // Only an operation that uses the GPU needs a CUDA context; the default forwards the Orin's access units untouched. +#ifdef VIDEO_STAGE_GPU + std::unique_ptr context; + if (settings.operation != Operation::None) + { + context = std::make_unique(); + std::cout << "GPU: " << context->device_name() << "\n"; + } +#else + if (settings.operation != Operation::None) + { + std::cerr << "this build has no GPU stage; configure with -DBUILD_VIDEO_STAGE_GPU=ON\n"; + return 1; + } + CudaContext* const context = nullptr; +#endif // One LiveKit participant for every stream, as the Orin producer joins once for all its cameras. std::unique_ptr livekit; if (any_livekit) livekit = std::make_unique(livekit_url, read_token(token_file)); +#ifdef VIDEO_STAGE_GPU + CudaContext* const context_ptr = context.get(); +#else + CudaContext* const context_ptr = context; +#endif std::vector> workers; for (const StreamSpec& spec : specs) - workers.push_back(std::make_unique(spec, settings, spec.file.empty() ? livekit.get() : nullptr)); + workers.push_back( + std::make_unique(context_ptr, spec, settings, spec.file.empty() ? livekit.get() : nullptr)); std::signal(SIGINT, on_signal); std::signal(SIGTERM, on_signal); diff --git a/src/video_stage/app/pipeline.cpp b/src/video_stage/app/pipeline.cpp new file mode 100644 index 0000000000..92846781c5 --- /dev/null +++ b/src/video_stage/app/pipeline.cpp @@ -0,0 +1,237 @@ +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: Apache-2.0 + +#include "pipeline.hpp" + +#include + +#include +#include +#include +#include +#include +#include +#include + +namespace video_stage +{ + +namespace +{ +constexpr size_t kMetaRing = 64; +} + +struct Pipeline::Impl +{ + CudaContext& context; + GpuOperation& operation; + PacketSink& sink; + PipelineConfig config; + cudaStream_t stream = nullptr; + + std::array meta{}; + uint64_t next_key = 0; + bool have_previous_frame = false; + uint64_t previous_frame = 0; + bool need_idr = true; + std::atomic idr_requested{ false }; + + Geometry input_geometry; + Geometry output_geometry; + std::unique_ptr encoder; + std::unique_ptr decoder; + PipelineStats stats; + int64_t decode_start_ns = 0; + std::string last_sps; + uint32_t source_fps = 0; ///< the frame rate the source declared in its last record + + Impl(CudaContext& ctx, GpuOperation& op, PacketSink& s, const PipelineConfig& cfg) + : context(ctx), operation(op), sink(s), config(cfg) + { + context.make_current(); + if (cudaStreamCreateWithFlags(&stream, cudaStreamNonBlocking) != cudaSuccess) + throw std::runtime_error("Pipeline: cannot create the CUDA stream"); + decoder = std::make_unique( + context, stream, [this](DecodedFrame&& f) { on_frame(std::move(f)); }, + [this](const StreamFormat& f) { on_format(f); }); + } + + // A new source session: nothing the previous one left in the codec may reach the new one. The decoder is rebuilt + // (its parser and any pictures it still holds are dropped with it), the encoder is released so the new session's + // first sequence header creates a matching one that opens with an IDR, and the identity ring forgets every key. + void reset_session() + { + context.make_current(); + cudaStreamSynchronize(stream); + decoder.reset(); + encoder.reset(); + decoder = std::make_unique( + context, stream, [this](DecodedFrame&& f) { on_frame(std::move(f)); }, + [this](const StreamFormat& f) { on_format(f); }); + meta.fill(FrameMeta{}); + next_key = 0; + input_geometry = Geometry{}; + output_geometry = Geometry{}; + last_sps.clear(); + have_previous_frame = false; + previous_frame = 0; + need_idr = true; + idr_requested = false; + } + + ~Impl() + { + decoder.reset(); + encoder.reset(); + context.make_current(); + cudaStreamDestroy(stream); + } + + // The input geometry changed: prepare the operation and a matching encoder before the next picture. The first + // output after a change is an IDR. + void on_format(const StreamFormat& format) + { + const Geometry g = format.geometry; + input_geometry = g; + operation.prepare(g); + output_geometry = operation.output_geometry(g); + EncoderConfig enc; + enc.geometry = output_geometry; + enc.full_range = format.color.full_range; // the encoder signals the range the pictures actually have + enc.bitrate_bps = config.bitrate_bps; + enc.idr_period = config.idr_period; + enc.fps_numerator = source_fps != 0 ? source_fps : config.fps_numerator; + encoder.reset(); + encoder = std::make_unique(context, enc); + need_idr = true; + } + + void on_frame(DecodedFrame&& frame) + { + ++stats.pictures; + const int64_t decoded_ns = monotonic_ns(); + const FrameMeta m = meta[frame.key % kMetaRing]; + stats.decode_ns.add(decoded_ns - m.submitted_ns); + stats.decode_lag.add(static_cast(next_key - 1 - frame.key)); + if (m.key != frame.key) + { + ++stats.metadata_lost; + return; // identity of this picture is unknown: never publish it under another frame's metadata + } + + NvencEncoder::Surface surface; + if (!encoder || !encoder->acquire(surface)) + { + ++stats.dropped_no_surface; + return; + } + EncodedPacket packet; + try + { + surface.view.color = frame.view.color; + operation.enqueue(frame.view, surface.view, stream); + const bool idr = need_idr || idr_requested.exchange(false); + packet = encoder->encode(surface, stream, frame.key, idr); + need_idr = false; + } + catch (...) + { + // Exceptional path only: launches already queued may still read the decoder surface or write the encoder's + // one. Drain this stream, then hand the encoder surface back; the decoder lease releases on unwind. + cudaStreamSynchronize(stream); + encoder->release(surface); + throw; + } + frame.lease.reset(); // the encoder has consumed the output; the decoder surface can be reused + const int64_t done_ns = monotonic_ns(); + stats.process_encode_ns.add(done_ns - decoded_ns); + stats.receive_to_output_ns.add(done_ns - m.received_ns); + const FrameMeta& out_meta = meta[packet.key % kMetaRing]; + if (out_meta.key != packet.key) + { + ++stats.metadata_lost; + return; + } + sink.write(packet, out_meta, output_geometry); + } +}; + +Pipeline::Pipeline(CudaContext& context, GpuOperation& operation, PacketSink& sink, const PipelineConfig& config) + : m_impl(std::make_unique(context, operation, sink, config)) +{ +} + +Pipeline::~Pipeline() = default; + +void Pipeline::process(const avs1::Record& record, int64_t received_ns) +{ + Impl& p = *m_impl; + ++p.stats.records; + const uint64_t frame = record.sample.frame_number; + // Frame ids need not be consecutive (the camera captures faster than it publishes); they must increase. + if (p.have_previous_frame && frame <= p.previous_frame) + ++p.stats.gaps; + p.previous_frame = frame; + p.have_previous_frame = true; + + const uint64_t key = p.next_key++; + FrameMeta& slot = p.meta[key % kMetaRing]; + slot.frame_number = frame; + slot.source_timestamp_ns = record.sample.frame_timestamp_ns; + slot.received_ns = received_ns; + slot.key = key; + p.source_fps = record.sample.fps_numerator; + p.decode_start_ns = monotonic_ns(); + slot.submitted_ns = p.decode_start_ns; + p.stats.queue_age_ns.add(p.decode_start_ns - received_ns); + + // A sequence header that does not say "no reordering" makes NVDEC hold back pictures; say it when the source's + // contract is that it has no B-frames, and report what the stream declared. + const uint8_t* payload = record.payload; + size_t payload_size = record.payload_size; + std::vector rewritten; + if (const auto sps = h264::parse_sps(payload, payload_size)) + { + // The operation samples chroma at the default siting; a stream that declares another would be warped with + // its colour misplaced, so it is refused here rather than processed. + if (sps->chroma_loc_present && (sps->chroma_loc_top != 0 || sps->chroma_loc_bottom != 0)) + throw std::runtime_error("the stream declares chroma sample locations (" + + std::to_string(sps->chroma_loc_top) + ", " + + std::to_string(sps->chroma_loc_bottom) + "); only the default siting is supported"); + const std::string described = sps->describe(); + if (described != p.last_sps) + { + p.last_sps = described; + std::cerr << "source SPS: " << described << std::endl; + } + const bool needs = !sps->bitstream_restriction || sps->max_num_reorder_frames != 0 || + sps->max_dec_frame_buffering > std::max(sps->max_num_ref_frames, 1); + if (p.config.source_has_no_b_frames && needs) + { + if (auto fixed = h264::declare_no_reordering(payload, payload_size)) + { + rewritten = std::move(*fixed); + payload = rewritten.data(); + payload_size = rewritten.size(); + } + } + } + p.decoder->decode(payload, payload_size, key); +} + +void Pipeline::reset_session() +{ + m_impl->reset_session(); +} + +void Pipeline::request_idr() +{ + m_impl->idr_requested = true; +} + +const PipelineStats& Pipeline::stats() const +{ + return m_impl->stats; +} + +} // namespace video_stage diff --git a/src/video_stage/app/pipeline.hpp b/src/video_stage/app/pipeline.hpp new file mode 100644 index 0000000000..52b864ea46 --- /dev/null +++ b/src/video_stage/app/pipeline.hpp @@ -0,0 +1,65 @@ +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: Apache-2.0 + +#pragma once + +#include "stats.hpp" + +#include +#include +#include +#include +#include +#include + +#include +#include +#include +#include +#include +#include + +namespace video_stage +{ + +struct PipelineConfig +{ + /// The source never emits B-frames (the Orin producer's encoder contract). A decoder that cannot see that holds + /// back pictures to be safe, so the sequence header is rewritten to declare it. + bool source_has_no_b_frames = true; + uint32_t bitrate_bps = 12'000'000; + uint32_t idr_period = 30; + uint32_t fps_numerator = 30; +}; + +/** + * @brief Decode -> GPU operation -> encode on one thread, frame by frame. + * + * Pictures stay in GPU memory from NVDEC to NVENC. Identity travels by key: each access unit's key indexes a small + * ring of FrameMeta, and NVDEC and NVENC hand the same key back, so metadata follows its own frame through any loss. + * A picture arriving while every encoder input surface is busy is dropped and counted, never queued. + */ +class Pipeline +{ +public: + Pipeline(CudaContext& context, GpuOperation& operation, PacketSink& sink, const PipelineConfig& config); + ~Pipeline(); + + /// Process one received record. Throws std::runtime_error on a decoder or encoder failure. + void process(const avs1::Record& record, int64_t received_ns); + + /// A new source session begins: forget the previous one's frame numbering and start the encoder at an IDR. Call + /// between sessions, from the thread that calls process(). + void reset_session(); + + /// Make the next encoded picture an IDR (a subscriber joined or the receiving side asked). Thread-safe. + void request_idr(); + + const PipelineStats& stats() const; + +private: + struct Impl; + std::unique_ptr m_impl; +}; + +} // namespace video_stage diff --git a/src/video_stage/app/stats.hpp b/src/video_stage/app/stats.hpp index 1be197b6c3..8666099f5a 100644 --- a/src/video_stage/app/stats.hpp +++ b/src/video_stage/app/stats.hpp @@ -4,6 +4,7 @@ #pragma once #include +#include #include #include #include @@ -62,4 +63,20 @@ class DurationStats std::vector m_samples; }; +struct PipelineStats +{ + // Written by the processing thread and read by the statistics thread, so the counters are atomic. + std::atomic records{ 0 }; + std::atomic pictures{ 0 }; + std::atomic dropped_no_surface{ 0 }; + std::atomic gaps{ 0 }; ///< records whose frame number did not increase + std::atomic metadata_lost{ 0 }; ///< pictures whose metadata entry had been overwritten + DurationStats queue_age_ns; ///< from intake to the start of processing + DurationStats decode_ns; ///< from handing an access unit to the decoder to its picture appearing + DurationStats decode_lag; ///< records handed to the decoder after this picture's own record, before it appeared + ///< (frames) + DurationStats process_encode_ns; + DurationStats receive_to_output_ns; +}; + } // namespace video_stage diff --git a/src/video_stage/app/stream_worker.cpp b/src/video_stage/app/stream_worker.cpp index 38a371ac95..09f94a9f68 100644 --- a/src/video_stage/app/stream_worker.cpp +++ b/src/video_stage/app/stream_worker.cpp @@ -108,6 +108,13 @@ class RecordQueue bool m_closed = false; }; +// What a stream without a GPU pipeline reports for the pipeline's own stages: nothing. +const PipelineStats& no_pipeline_stats() +{ + static const PipelineStats empty; + return empty; +} + void print_stage(std::ostream& out, const char* name, const DurationStats& s) { char line[160]; @@ -120,14 +127,21 @@ void print_stage(std::ostream& out, const char* name, const DurationStats& s) struct StreamWorker::Impl { + CudaContext* context; StreamSpec spec; StreamSettings settings; SharedLiveKit* livekit; +#ifdef VIDEO_STAGE_GPU + std::unique_ptr operation; ///< null for Operation::None +#endif std::unique_ptr receiver; std::atomic active_control{ nullptr }; std::unique_ptr sink; LiveKitOutput* livekit_output = nullptr; +#ifdef VIDEO_STAGE_GPU + std::unique_ptr pipeline; +#endif std::thread thread; std::atomic stop{ false }; @@ -140,12 +154,23 @@ struct StreamWorker::Impl std::atomic clock_rejected{ 0 }; std::atomic clock_discontinuities{ 0 }; DurationStats backlog_bytes; - // The picture size from the last sequence header, and a count. + // Pass-through state (Operation::None): the picture size from the last sequence header, and a count. Geometry passthrough_geometry; std::atomic passthrough_records{ 0 }; - Impl(const StreamSpec& s, const StreamSettings& st, SharedLiveKit* lk) : spec(s), settings(st), livekit(lk) + Impl(CudaContext* ctx, const StreamSpec& s, const StreamSettings& st, SharedLiveKit* lk) + : context(ctx), spec(s), settings(st), livekit(lk) { + if (settings.operation == Operation::SineWarp) + { +#ifdef VIDEO_STAGE_GPU + if (context == nullptr) + throw std::runtime_error("the sine warp needs a GPU context"); + operation = std::make_unique(settings.warp); +#else + throw std::runtime_error("this build has no GPU stage; configure with -DBUILD_VIDEO_STAGE_GPU=ON"); +#endif + } if (livekit != nullptr) { LiveKitOutputConfig out = settings.livekit; @@ -159,7 +184,14 @@ struct StreamWorker::Impl }, [this] { - // A subscriber needs a keyframe: only the Orin's encoder can make one. + // A subscriber needs a keyframe: ask this stage's encoder, or, with nothing in between, the Orin's. +#ifdef VIDEO_STAGE_GPU + if (pipeline) + { + pipeline->request_idr(); + return; + } +#endif receiver->request_keyframe(); }); livekit_output = output.get(); @@ -169,11 +201,33 @@ struct StreamWorker::Impl { sink = std::make_unique(spec.file); } +#ifdef VIDEO_STAGE_GPU + if (operation) + pipeline = std::make_unique(*context, *operation, *sink, settings.pipeline); +#endif receiver = std::make_unique(spec.port); } - // The access unit goes to the sink untouched. The picture size comes from the stream's own sequence header; - // nothing is published before the first one. + bool has_pipeline() const + { +#ifdef VIDEO_STAGE_GPU + return pipeline != nullptr; +#else + return false; +#endif + } + + const PipelineStats& pipeline_stats() const + { +#ifdef VIDEO_STAGE_GPU + if (pipeline) + return pipeline->stats(); +#endif + return no_pipeline_stats(); + } + + // Operation::None: the access unit goes to the sink untouched. The picture size comes from the stream's own + // sequence header; nothing is published before the first one. void forward(const avs1::Record& record, int64_t arrived_ns) { const bool keyframe = avs1::contains_idr(record.payload, record.payload_size); @@ -258,6 +312,13 @@ struct StreamWorker::Impl const auto record = avs1::parse(item.bytes.data(), item.bytes.size(), error); if (!record) throw std::runtime_error("bad AVS1 record: " + error); +#ifdef VIDEO_STAGE_GPU + if (pipeline) + { + pipeline->process(*record, item.arrived_ns); + continue; + } +#endif forward(*record, item.arrived_ns); } } @@ -289,6 +350,10 @@ struct StreamWorker::Impl // over. if (livekit_output) livekit_output->end_session(); +#ifdef VIDEO_STAGE_GPU + if (pipeline) + pipeline->reset_session(); +#endif passthrough_geometry = Geometry{}; epoch_base += 1'000'000; ControlChannel control(*receiver, epoch_base + 1); @@ -316,8 +381,11 @@ struct StreamWorker::Impl } }; -StreamWorker::StreamWorker(const StreamSpec& spec, const StreamSettings& settings, SharedLiveKit* livekit) - : m_impl(std::make_unique(spec, settings, livekit)) +StreamWorker::StreamWorker(CudaContext* context, + const StreamSpec& spec, + const StreamSettings& settings, + SharedLiveKit* livekit) + : m_impl(std::make_unique(context, spec, settings, livekit)) { } @@ -349,9 +417,17 @@ bool StreamWorker::join() void StreamWorker::print_live(std::ostream& out) const { const Impl& d = *m_impl; + const PipelineStats& s = d.pipeline_stats(); char line[512]; - int n = std::snprintf(line, sizeof(line), "[%u] backlog p99 %lld B", static_cast(d.spec.port), - static_cast(d.backlog_bytes.percentile(99))); + int n = std::snprintf( + line, sizeof(line), + "[%u] queue age p50 %.1f p99 %.1f ms | backlog p99 %lld B | decoder lag p50 %lld p99 %lld frames | " + "submit->picture p50 %.1f p99 %.1f ms | warp+encode p50 %.1f p99 %.1f ms", + static_cast(d.spec.port), s.queue_age_ns.percentile(50) / 1e6, s.queue_age_ns.percentile(99) / 1e6, + static_cast(d.backlog_bytes.percentile(99)), static_cast(s.decode_lag.percentile(50)), + static_cast(s.decode_lag.percentile(99)), s.decode_ns.percentile(50) / 1e6, + s.decode_ns.percentile(99) / 1e6, s.process_encode_ns.percentile(50) / 1e6, + s.process_encode_ns.percentile(99) / 1e6); if (d.livekit_output && n > 0 && n < static_cast(sizeof(line))) { std::snprintf( @@ -372,15 +448,18 @@ void StreamWorker::print_live(std::ostream& out) const void StreamWorker::print_summary(std::ostream& out) const { const Impl& d = *m_impl; + const PipelineStats& s = d.pipeline_stats(); char line[256]; std::snprintf( line, sizeof(line), - "stream %u: sessions %llu, records %llu\n" + "stream %u: sessions %llu, records %llu, pictures %llu, dropped (no surface) %llu, gaps %llu, metadata lost %llu\n" " overflows %llu, records discarded %llu, clock exchanges %llu (rejected %llu, discontinuities %llu)\n", static_cast(d.spec.port), static_cast(d.sessions.load()), - static_cast(d.passthrough_records.load()), - static_cast(d.overflows.load()), static_cast(d.discarded.load()), - static_cast(d.clock_exchanges.load()), + static_cast(d.has_pipeline() ? s.records.load() : d.passthrough_records.load()), + static_cast(d.has_pipeline() ? s.pictures.load() : d.passthrough_records.load()), + static_cast(s.dropped_no_surface.load()), static_cast(s.gaps.load()), + static_cast(s.metadata_lost.load()), static_cast(d.overflows.load()), + static_cast(d.discarded.load()), static_cast(d.clock_exchanges.load()), static_cast(d.clock_rejected.load()), static_cast(d.clock_discontinuities.load())); out << line; @@ -407,6 +486,18 @@ void StreamWorker::print_summary(std::ostream& out) const print_stage(out, "Orin -> PC arrival", d.livekit_output->uplink_ns()); print_stage(out, "Orin -> publish", d.livekit_output->source_to_publish_ns()); } + print_stage(out, "queue age", s.queue_age_ns); + { + char lag[160]; + std::snprintf(lag, sizeof(lag), " decoder lag (frames) p50 %lld p99 %lld max %lld\n", + static_cast(s.decode_lag.percentile(50)), + static_cast(s.decode_lag.percentile(99)), + static_cast(s.decode_lag.percentile(100))); + out << lag; + } + print_stage(out, "decode", s.decode_ns); + print_stage(out, "warp + encode", s.process_encode_ns); + print_stage(out, "receive -> output", s.receive_to_output_ns); } } // namespace video_stage diff --git a/src/video_stage/app/stream_worker.hpp b/src/video_stage/app/stream_worker.hpp index afad93c4fa..6be5e003c4 100644 --- a/src/video_stage/app/stream_worker.hpp +++ b/src/video_stage/app/stream_worker.hpp @@ -5,6 +5,13 @@ #include "livekit_output.hpp" +#ifdef VIDEO_STAGE_GPU +# include "pipeline.hpp" + +# include +# include +#endif + #include #include #include @@ -14,6 +21,10 @@ namespace video_stage { +#ifndef VIDEO_STAGE_GPU +class CudaContext; // only ever null in a build without the GPU stage +#endif + /// One camera's stream: where the Orin connects, and where the processed video goes. struct StreamSpec { @@ -24,8 +35,24 @@ struct StreamSpec std::string file; }; +/// What happens to each picture between the Orin and LiveKit. +enum class Operation +{ + /// Nothing: the Orin's encoded access units go to LiveKit as they are. No decode, no encode, no GPU; the lowest + /// latency and no quality loss. The default. + None, + /// Decode on NVDEC, apply the example sine-wave warp, re-encode on NVENC. For testing the GPU stage and as the + /// template for the fisheye rectification. + SineWarp, +}; + struct StreamSettings { + Operation operation = Operation::None; +#ifdef VIDEO_STAGE_GPU + PipelineConfig pipeline; + SineWarpParams warp; +#endif LiveKitOutputConfig livekit; bool once = false; ///< serve one Orin session, then finish }; @@ -33,13 +60,14 @@ struct StreamSettings /** * @brief One stream's thread: accept an Orin session, run it to its end, accept the next. * - * Owns that stream's receiver, control channel and output. When publishing, streams share one LiveKit participant and - * otherwise run independently, so one camera's trouble stays its own. The Orin's access units go on as they are. + * Owns that stream's receiver, control channel, optional GPU pipeline and output. Streams share the CUDA context and, + * when publishing, one LiveKit participant, and otherwise run independently so one camera's trouble stays its own. */ class StreamWorker { public: - StreamWorker(const StreamSpec& spec, const StreamSettings& settings, SharedLiveKit* livekit); + /// `context` is needed only for an operation that uses the GPU (null for Operation::None). + StreamWorker(CudaContext* context, const StreamSpec& spec, const StreamSettings& settings, SharedLiveKit* livekit); ~StreamWorker(); StreamWorker(const StreamWorker&) = delete; diff --git a/src/video_stage/codec/CMakeLists.txt b/src/video_stage/codec/CMakeLists.txt new file mode 100644 index 0000000000..e755cc46e6 --- /dev/null +++ b/src/video_stage/codec/CMakeLists.txt @@ -0,0 +1,27 @@ +# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# SPDX-License-Identifier: Apache-2.0 + +# NVDEC decoder and NVENC encoder adapters over a shared CUDA context. +add_library(video_stage_codec STATIC + cuda_context.cpp + nvdec_decoder.cpp + nvenc_encoder.cpp +) +add_library(video_stage::codec ALIAS video_stage_codec) + +set_target_properties(video_stage_codec PROPERTIES POSITION_INDEPENDENT_CODE ON) + +target_include_directories(video_stage_codec + INTERFACE + $ +) + +target_link_libraries(video_stage_codec + PUBLIC + video_stage::core + CUDA::cudart + CUDA::cuda_driver + PRIVATE + video_stage_nv_codec_headers + ${CMAKE_DL_LIBS} +) diff --git a/src/video_stage/codec/cuda_context.cpp b/src/video_stage/codec/cuda_context.cpp new file mode 100644 index 0000000000..714165b063 --- /dev/null +++ b/src/video_stage/codec/cuda_context.cpp @@ -0,0 +1,50 @@ +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: Apache-2.0 + +#include "inc/video_stage/cuda_context.hpp" + +#include + +namespace video_stage +{ + +void check_cu(CUresult result, const char* what) +{ + if (result == CUDA_SUCCESS) + return; + const char* message = nullptr; + cuGetErrorString(result, &message); + throw std::runtime_error(std::string(what) + ": " + (message ? message : "unknown CUDA driver error")); +} + +CudaContext::CudaContext(int device_index) +{ + check_cu(cuInit(0), "cuInit"); + check_cu(cuDeviceGet(&m_device, device_index), "cuDeviceGet"); + char name[128] = {}; + check_cu(cuDeviceGetName(name, sizeof(name), m_device), "cuDeviceGetName"); + m_name = name; + check_cu(cuDevicePrimaryCtxRetain(&m_context, m_device), "cuDevicePrimaryCtxRetain"); + make_current(); + // Bind the runtime API to the same primary context. + if (cudaSetDevice(device_index) != cudaSuccess) + throw std::runtime_error("cudaSetDevice failed"); +} + +CudaContext::~CudaContext() +{ + if (m_context) + cuDevicePrimaryCtxRelease(m_device); +} + +void CudaContext::make_current_noexcept() const noexcept +{ + cuCtxSetCurrent(m_context); +} + +void CudaContext::make_current() const +{ + check_cu(cuCtxSetCurrent(m_context), "cuCtxSetCurrent"); +} + +} // namespace video_stage diff --git a/src/video_stage/codec/inc/video_stage/cuda_context.hpp b/src/video_stage/codec/inc/video_stage/cuda_context.hpp new file mode 100644 index 0000000000..5410a3e939 --- /dev/null +++ b/src/video_stage/codec/inc/video_stage/cuda_context.hpp @@ -0,0 +1,49 @@ +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: Apache-2.0 + +#pragma once + +#include +#include +#include + +namespace video_stage +{ + +/// Throws std::runtime_error naming `what` when a CUDA driver call failed. +void check_cu(CUresult result, const char* what); + +/** + * @brief The device's primary CUDA context, shared by decode, the GPU operation and encode. + * + * The runtime API (kernels, streams) and the driver API (NVDEC, NVENC) both use the primary context, so surfaces + * are visible to all of them. Every thread that touches CUDA calls make_current() once. + */ +class CudaContext +{ +public: + explicit CudaContext(int device_index = 0); + ~CudaContext(); + + CudaContext(const CudaContext&) = delete; + CudaContext& operator=(const CudaContext&) = delete; + + void make_current() const; + /// For cleanup paths that must not throw: the result is ignored. + void make_current_noexcept() const noexcept; + CUcontext handle() const + { + return m_context; + } + const std::string& device_name() const + { + return m_name; + } + +private: + CUdevice m_device = 0; + CUcontext m_context = nullptr; + std::string m_name; +}; + +} // namespace video_stage diff --git a/src/video_stage/codec/inc/video_stage/nvdec_decoder.hpp b/src/video_stage/codec/inc/video_stage/nvdec_decoder.hpp new file mode 100644 index 0000000000..fc41280e12 --- /dev/null +++ b/src/video_stage/codec/inc/video_stage/nvdec_decoder.hpp @@ -0,0 +1,79 @@ +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: Apache-2.0 + +#pragma once + +#include "cuda_context.hpp" + +#include + +#include +#include +#include +#include +#include + +namespace video_stage +{ + +/// What a decoded stream looks like: its size and its quantisation range. A change of either at a sequence header +/// recreates the decoder and is reported to the consumer. +struct StreamFormat +{ + Geometry geometry; + ColorInfo color; + + bool operator==(const StreamFormat& other) const + { + return geometry == other.geometry && color.full_range == other.color.full_range; + } +}; + +/// A decoded picture in device memory, still mapped to the decoder's output surface. Keep the lease until the work +/// that reads `view` has completed; releasing it returns the surface to the decoder. `key` is the value given with +/// the access unit that produced the picture. +struct DecodedFrame +{ + Nv12View view; + uint64_t key = 0; + std::shared_ptr lease; +}; + +/** + * @brief H.264 access units to NV12 device surfaces on NVDEC. + * + * decode() feeds one complete Annex-B access unit and returns after the parser has called back. With display delay 0 + * and no B-frames, a picture is delivered from the call that supplied its access unit. Pictures arrive in display + * order through the callback, on the calling thread. Every picture carries the key of its own access unit, so + * identity survives any reordering and a lost access unit cannot shift later keys. + * + * A change of coded size or format in the stream recreates the decoder at that access unit; on_format_change() + * tells the consumer, which must reprepare its operation for the new geometry before the next picture. + */ +class NvdecDecoder +{ +public: + using FrameCallback = std::function; + using FormatCallback = std::function; + + /// `stream` is the CUDA stream the decoded surfaces are made ready on; work enqueued on it afterwards is ordered + /// after the decode. + NvdecDecoder(CudaContext& context, cudaStream_t stream, FrameCallback on_frame, FormatCallback on_format_change); + ~NvdecDecoder(); + + NvdecDecoder(const NvdecDecoder&) = delete; + NvdecDecoder& operator=(const NvdecDecoder&) = delete; + + /// Throws std::runtime_error if the decoder rejects the data. + void decode(const uint8_t* access_unit, size_t size, uint64_t key); + + /// Pictures delivered so far and access units that produced none; the difference shows decoder-side loss. + uint64_t pictures() const; + uint64_t access_units() const; + +private: + struct Impl; + std::unique_ptr m_impl; +}; + +} // namespace video_stage diff --git a/src/video_stage/codec/inc/video_stage/nvenc_encoder.hpp b/src/video_stage/codec/inc/video_stage/nvenc_encoder.hpp new file mode 100644 index 0000000000..18caf62cf1 --- /dev/null +++ b/src/video_stage/codec/inc/video_stage/nvenc_encoder.hpp @@ -0,0 +1,72 @@ +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: Apache-2.0 + +#pragma once + +#include "cuda_context.hpp" + +#include + +#include +#include +#include +#include +#include + +namespace video_stage +{ + +struct EncoderConfig +{ + Geometry geometry; + uint32_t fps_numerator = 30; + uint32_t fps_denominator = 1; + uint32_t bitrate_bps = 12'000'000; + /// Frames between IDRs; a stream recovers from loss only at an IDR, so keep it short. + uint32_t idr_period = 30; + bool full_range = false; + /// Reusable input surfaces, registered with the encoder once at startup. + uint32_t input_surfaces = 3; +}; + +/** + * @brief NV12 device surfaces to H.264 access units on NVENC, tuned for low latency. + * + * The encoder owns a small pool of CUDA surfaces registered once; the GPU operation writes its output straight + * into one, so there is no copy and no per-frame registration. Low-latency settings: the ultra-low-latency tuning, + * no B-frames, no look-ahead, constant bitrate with a one-frame buffer, parameter sets repeated before every IDR. + * encode() returns the access unit for the picture it was given. + */ +class NvencEncoder +{ +public: + /// A pooled input surface. Write the picture into `view`, then pass the same `slot` to encode(). + struct Surface + { + Nv12View view; + uint32_t slot = 0; + }; + + NvencEncoder(CudaContext& context, const EncoderConfig& config); + ~NvencEncoder(); + + NvencEncoder(const NvencEncoder&) = delete; + NvencEncoder& operator=(const NvencEncoder&) = delete; + + /// The next free input surface, or nullopt-equivalent `false` when all are in flight. + bool acquire(Surface& out); + + /// Give back a surface acquired but not encoded (an error path). Idempotent, and encode() releases its own. + void release(const Surface& surface); + + /// Encode the picture written to `surface` after `stream` has completed the work that wrote it, and release the + /// surface. `key` comes back on the packet. `force_idr` makes this picture an IDR with parameter sets. + /// Throws std::runtime_error on an encoder error. + EncodedPacket encode(const Surface& surface, cudaStream_t stream, uint64_t key, bool force_idr); + +private: + struct Impl; + std::unique_ptr m_impl; +}; + +} // namespace video_stage diff --git a/src/video_stage/codec/nvdec_decoder.cpp b/src/video_stage/codec/nvdec_decoder.cpp new file mode 100644 index 0000000000..c10117afe7 --- /dev/null +++ b/src/video_stage/codec/nvdec_decoder.cpp @@ -0,0 +1,231 @@ +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: Apache-2.0 + +#include "inc/video_stage/nvdec_decoder.hpp" + +#include + +#include +#include +#include +#include + +namespace video_stage +{ + +namespace +{ + +void check_cu_v(CUresult r, const char* what) +{ + check_cu(r, what); +} + +} // namespace + +struct NvdecDecoder::Impl +{ + CudaContext& context; + cudaStream_t stream; + FrameCallback on_frame; + FormatCallback on_format_change; + + CuvidFunctions* cvdl = nullptr; + CUvideoparser parser = nullptr; + CUvideodecoder decoder = nullptr; + CUVIDEOFORMAT format{}; + StreamFormat stream_format; + uint64_t pictures = 0; + uint64_t access_units = 0; + // Set by the callbacks so decode() can surface an error thrown inside the parser's C callbacks. + std::string callback_error; + + Impl(CudaContext& ctx, cudaStream_t s, FrameCallback f, FormatCallback fc) + : context(ctx), stream(s), on_frame(std::move(f)), on_format_change(std::move(fc)) + { + if (cuvid_load_functions(&cvdl, nullptr) < 0 || cvdl == nullptr) + throw std::runtime_error("NvdecDecoder: libnvcuvid is not available"); + context.make_current(); + CUVIDPARSERPARAMS params{}; + params.CodecType = cudaVideoCodec_H264; + // The parser only needs a minimum here; the sequence callback reports the real requirement. + params.ulMaxNumDecodeSurfaces = 1; + params.ulClockRate = 0; + params.ulErrorThreshold = 100; + params.ulMaxDisplayDelay = 0; // output as soon as decoded; the stream has no B-frames + params.pUserData = this; + params.pfnSequenceCallback = &Impl::sequence_cb; + params.pfnDecodePicture = &Impl::decode_cb; + params.pfnDisplayPicture = &Impl::display_cb; + check_cu_v(cvdl->cuvidCreateVideoParser(&parser, ¶ms), "cuvidCreateVideoParser"); + } + + ~Impl() + { + context.make_current(); + if (parser) + cvdl->cuvidDestroyVideoParser(parser); + destroy_decoder(); + cuvid_free_functions(&cvdl); + } + + void destroy_decoder() + { + if (decoder) + { + cvdl->cuvidDestroyDecoder(decoder); + decoder = nullptr; + } + } + + static int CUDAAPI sequence_cb(void* user, CUVIDEOFORMAT* fmt) + { + auto* self = static_cast(user); + try + { + return self->on_sequence(fmt); + } + catch (const std::exception& e) + { + self->callback_error = e.what(); + return 0; // 0 = failure: the parser stops + } + } + + int on_sequence(CUVIDEOFORMAT* fmt) + { + if (fmt->codec != cudaVideoCodec_H264 || fmt->chroma_format != cudaVideoChromaFormat_420 || + fmt->bit_depth_luma_minus8 != 0) + throw std::runtime_error("NvdecDecoder: only 8-bit 4:2:0 H.264 is supported"); + const unsigned width = static_cast(fmt->display_area.right - fmt->display_area.left); + const unsigned height = static_cast(fmt->display_area.bottom - fmt->display_area.top); + StreamFormat next; + next.geometry = Geometry{ width, height }; + next.color.full_range = fmt->video_signal_description.video_full_range_flag != 0; + const bool same = decoder != nullptr && format.coded_width == fmt->coded_width && + format.coded_height == fmt->coded_height && stream_format == next && + format.min_num_decode_surfaces == fmt->min_num_decode_surfaces; + if (same) + return static_cast(format.min_num_decode_surfaces); + + destroy_decoder(); + CUVIDDECODECREATEINFO info{}; + info.ulWidth = fmt->coded_width; + info.ulHeight = fmt->coded_height; + info.ulMaxWidth = fmt->coded_width; + info.ulMaxHeight = fmt->coded_height; + // The codec's own minimum plus margin so a consumer holding a surface does not stall the parser. + info.ulNumDecodeSurfaces = fmt->min_num_decode_surfaces + 2; + info.CodecType = fmt->codec; + info.ChromaFormat = fmt->chroma_format; + info.bitDepthMinus8 = fmt->bit_depth_luma_minus8; + info.ulCreationFlags = cudaVideoCreate_PreferCUVID; + info.OutputFormat = cudaVideoSurfaceFormat_NV12; + info.DeinterlaceMode = cudaVideoDeinterlaceMode_Weave; + info.ulTargetWidth = width; + info.ulTargetHeight = height; + info.display_area.left = static_cast(fmt->display_area.left); + info.display_area.top = static_cast(fmt->display_area.top); + info.display_area.right = static_cast(fmt->display_area.right); + info.display_area.bottom = static_cast(fmt->display_area.bottom); + info.ulNumOutputSurfaces = 3; + context.make_current(); + check_cu_v(cvdl->cuvidCreateDecoder(&decoder, &info), "cuvidCreateDecoder"); + format = *fmt; + stream_format = next; + if (on_format_change) + on_format_change(next); + return static_cast(info.ulNumDecodeSurfaces); + } + + static int CUDAAPI decode_cb(void* user, CUVIDPICPARAMS* pic) + { + auto* self = static_cast(user); + if (self->cvdl->cuvidDecodePicture(self->decoder, pic) != CUDA_SUCCESS) + { + self->callback_error = "cuvidDecodePicture failed"; + return 0; + } + return 1; + } + + static int CUDAAPI display_cb(void* user, CUVIDPARSERDISPINFO* disp) + { + auto* self = static_cast(user); + try + { + self->on_display(disp); + return 1; + } + catch (const std::exception& e) + { + self->callback_error = e.what(); + return 0; + } + } + + void on_display(CUVIDPARSERDISPINFO* disp) + { + CUVIDPROCPARAMS proc{}; + proc.progressive_frame = disp->progressive_frame; + proc.second_field = disp->repeat_first_field + 1; + proc.top_field_first = disp->top_field_first; + proc.unpaired_field = disp->repeat_first_field < 0; + proc.output_stream = stream; + CUdeviceptr dptr = 0; + unsigned pitch = 0; + check_cu_v(cvdl->cuvidMapVideoFrame(decoder, disp->picture_index, &dptr, &pitch, &proc), "cuvidMapVideoFrame"); + + DecodedFrame frame; + frame.key = static_cast(disp->timestamp); + frame.view.geometry = stream_format.geometry; + frame.view.y = reinterpret_cast(dptr); + // The decoder's NV12 output surface holds the chroma plane right after the luma plane, `target height` rows on. + frame.view.uv = frame.view.y + static_cast(pitch) * stream_format.geometry.height; + frame.view.y_pitch = pitch; + frame.view.uv_pitch = pitch; + frame.view.color = stream_format.color; + CuvidFunctions* cvdl_copy = cvdl; + CUvideodecoder dec = decoder; + frame.lease = std::shared_ptr(reinterpret_cast(dptr), [cvdl_copy, dec](void* p) + { cvdl_copy->cuvidUnmapVideoFrame(dec, reinterpret_cast(p)); }); + ++pictures; + on_frame(std::move(frame)); + } +}; + +NvdecDecoder::NvdecDecoder(CudaContext& context, cudaStream_t stream, FrameCallback on_frame, FormatCallback on_format_change) + : m_impl(std::make_unique(context, stream, std::move(on_frame), std::move(on_format_change))) +{ +} + +NvdecDecoder::~NvdecDecoder() = default; + +void NvdecDecoder::decode(const uint8_t* access_unit, size_t size, uint64_t key) +{ + Impl& d = *m_impl; + d.context.make_current(); + d.callback_error.clear(); + CUVIDSOURCEDATAPACKET packet{}; + packet.payload = access_unit; + packet.payload_size = size; + packet.flags = CUVID_PKT_TIMESTAMP | CUVID_PKT_ENDOFPICTURE; + packet.timestamp = static_cast(key); + ++d.access_units; + const CUresult r = d.cvdl->cuvidParseVideoData(d.parser, &packet); + if (!d.callback_error.empty()) + throw std::runtime_error("NvdecDecoder: " + d.callback_error); + check_cu_v(r, "cuvidParseVideoData"); +} + +uint64_t NvdecDecoder::pictures() const +{ + return m_impl->pictures; +} + +uint64_t NvdecDecoder::access_units() const +{ + return m_impl->access_units; +} + +} // namespace video_stage diff --git a/src/video_stage/codec/nvenc_encoder.cpp b/src/video_stage/codec/nvenc_encoder.cpp new file mode 100644 index 0000000000..bd19ce55c6 --- /dev/null +++ b/src/video_stage/codec/nvenc_encoder.cpp @@ -0,0 +1,284 @@ +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: Apache-2.0 + +#include "inc/video_stage/nvenc_encoder.hpp" + +#include + +#include +#include +#include + +namespace video_stage +{ + +namespace +{ + +void check_nv(NVENCSTATUS status, const char* what) +{ + if (status != NV_ENC_SUCCESS) + throw std::runtime_error(std::string("NvencEncoder: ") + what + " failed with status " + std::to_string(status)); +} + +} // namespace + +struct NvencEncoder::Impl +{ + CudaContext& context; + EncoderConfig config; + + NvencFunctions* nvenc = nullptr; + NV_ENCODE_API_FUNCTION_LIST fn{}; + void* encoder = nullptr; + NV_ENC_OUTPUT_PTR bitstream = nullptr; + + struct Slot + { + CUdeviceptr ptr = 0; + size_t pitch = 0; + NV_ENC_REGISTERED_PTR registered = nullptr; + bool in_use = false; + }; + std::vector slots; + + Impl(CudaContext& ctx, const EncoderConfig& cfg) : context(ctx), config(cfg) + { + if (config.geometry.width == 0 || config.geometry.height == 0 || (config.geometry.width % 2) != 0 || + (config.geometry.height % 2) != 0) + throw std::invalid_argument("NvencEncoder: NV12 needs non-zero even dimensions"); + try + { + context.make_current(); + if (nvenc_load_functions(&nvenc, nullptr) < 0 || nvenc == nullptr) + throw std::runtime_error("NvencEncoder: libnvidia-encode is not available"); + fn.version = NV_ENCODE_API_FUNCTION_LIST_VER; + check_nv(nvenc->NvEncodeAPICreateInstance(&fn), "NvEncodeAPICreateInstance"); + + NV_ENC_OPEN_ENCODE_SESSION_EX_PARAMS open{}; + open.version = NV_ENC_OPEN_ENCODE_SESSION_EX_PARAMS_VER; + open.deviceType = NV_ENC_DEVICE_TYPE_CUDA; + open.device = context.handle(); + open.apiVersion = NVENCAPI_VERSION; + check_nv(fn.nvEncOpenEncodeSessionEx(&open, &encoder), "nvEncOpenEncodeSessionEx"); + + configure(); + allocate_surfaces(); + + NV_ENC_CREATE_BITSTREAM_BUFFER out{}; + out.version = NV_ENC_CREATE_BITSTREAM_BUFFER_VER; + check_nv(fn.nvEncCreateBitstreamBuffer(encoder, &out), "nvEncCreateBitstreamBuffer"); + bitstream = out.bitstreamBuffer; + } + catch (...) + { + // A constructor that throws does not run the destructor, so a failure part way releases what exists. + shutdown(); + throw; + } + } + + ~Impl() + { + shutdown(); + } + + // Releases the session, the registered surfaces and the loaded library; safe to call on a partly built object. + void shutdown() noexcept + { + context.make_current_noexcept(); + if (encoder) + { + for (Slot& s : slots) + { + if (s.registered) + fn.nvEncUnregisterResource(encoder, s.registered); + if (s.ptr) + cuMemFree(s.ptr); + s = Slot{}; + } + if (bitstream) + fn.nvEncDestroyBitstreamBuffer(encoder, bitstream); + fn.nvEncDestroyEncoder(encoder); + encoder = nullptr; + bitstream = nullptr; + } + if (nvenc) + nvenc_free_functions(&nvenc); + } + + void configure() + { + const GUID codec = NV_ENC_CODEC_H264_GUID; + const GUID preset = NV_ENC_PRESET_P1_GUID; + const NV_ENC_TUNING_INFO tuning = NV_ENC_TUNING_INFO_ULTRA_LOW_LATENCY; + + NV_ENC_PRESET_CONFIG preset_config{}; + preset_config.version = NV_ENC_PRESET_CONFIG_VER; + preset_config.presetCfg.version = NV_ENC_CONFIG_VER; + check_nv(fn.nvEncGetEncodePresetConfigEx(encoder, codec, preset, tuning, &preset_config), + "nvEncGetEncodePresetConfigEx"); + + NV_ENC_CONFIG enc_config = preset_config.presetCfg; + enc_config.version = NV_ENC_CONFIG_VER; + enc_config.gopLength = config.idr_period; + enc_config.frameIntervalP = 1; // no B-frames + enc_config.rcParams.rateControlMode = NV_ENC_PARAMS_RC_CBR; + enc_config.rcParams.averageBitRate = config.bitrate_bps; + enc_config.rcParams.maxBitRate = config.bitrate_bps; + // A buffer of one frame's worth of bits keeps the output close to the instantaneous rate. + const uint32_t frame_bits = static_cast(static_cast(config.bitrate_bps) * + config.fps_denominator / config.fps_numerator); + enc_config.rcParams.vbvBufferSize = frame_bits; + enc_config.rcParams.vbvInitialDelay = frame_bits; + enc_config.rcParams.lookaheadDepth = 0; + NV_ENC_CONFIG_H264& h264 = enc_config.encodeCodecConfig.h264Config; + h264.idrPeriod = config.idr_period; + h264.repeatSPSPPS = 1; + h264.chromaFormatIDC = 1; + h264.h264VUIParameters.videoSignalTypePresentFlag = 1; + h264.h264VUIParameters.videoFormat = NV_ENC_VUI_VIDEO_FORMAT_UNSPECIFIED; + h264.h264VUIParameters.videoFullRangeFlag = config.full_range ? 1 : 0; + + NV_ENC_INITIALIZE_PARAMS init{}; + init.version = NV_ENC_INITIALIZE_PARAMS_VER; + init.encodeGUID = codec; + init.presetGUID = preset; + init.tuningInfo = tuning; + init.encodeWidth = config.geometry.width; + init.encodeHeight = config.geometry.height; + init.darWidth = config.geometry.width; + init.darHeight = config.geometry.height; + init.maxEncodeWidth = config.geometry.width; + init.maxEncodeHeight = config.geometry.height; + init.frameRateNum = config.fps_numerator; + init.frameRateDen = config.fps_denominator; + init.enablePTD = 1; + init.encodeConfig = &enc_config; + check_nv(fn.nvEncInitializeEncoder(encoder, &init), "nvEncInitializeEncoder"); + } + + void allocate_surfaces() + { + slots.resize(config.input_surfaces); + const size_t rows = static_cast(config.geometry.height) * 3 / 2; + for (Slot& s : slots) + { + // One allocation: luma rows then chroma rows, the layout NVENC reads for NV12. + check_cu(cuMemAllocPitch(&s.ptr, &s.pitch, config.geometry.width, rows, 16), "cuMemAllocPitch"); + NV_ENC_REGISTER_RESOURCE reg{}; + reg.version = NV_ENC_REGISTER_RESOURCE_VER; + reg.resourceType = NV_ENC_INPUT_RESOURCE_TYPE_CUDADEVICEPTR; + reg.width = config.geometry.width; + reg.height = config.geometry.height; + reg.pitch = static_cast(s.pitch); + reg.resourceToRegister = reinterpret_cast(s.ptr); + reg.bufferFormat = NV_ENC_BUFFER_FORMAT_NV12; + reg.bufferUsage = NV_ENC_INPUT_IMAGE; + check_nv(fn.nvEncRegisterResource(encoder, ®), "nvEncRegisterResource"); + s.registered = reg.registeredResource; + } + } +}; + +NvencEncoder::NvencEncoder(CudaContext& context, const EncoderConfig& config) + : m_impl(std::make_unique(context, config)) +{ +} + +NvencEncoder::~NvencEncoder() = default; + +void NvencEncoder::release(const Surface& surface) +{ + m_impl->slots.at(surface.slot).in_use = false; +} + +bool NvencEncoder::acquire(Surface& out) +{ + Impl& e = *m_impl; + for (uint32_t i = 0; i < e.slots.size(); ++i) + { + Impl::Slot& s = e.slots[i]; + if (s.in_use) + continue; + s.in_use = true; + out.slot = i; + out.view.geometry = e.config.geometry; + out.view.color.full_range = e.config.full_range; + out.view.y = reinterpret_cast(s.ptr); + out.view.y_pitch = s.pitch; + out.view.uv = reinterpret_cast(s.ptr) + s.pitch * e.config.geometry.height; + out.view.uv_pitch = s.pitch; + return true; + } + return false; +} + +EncodedPacket NvencEncoder::encode(const Surface& surface, cudaStream_t stream, uint64_t key, bool force_idr) +{ + Impl& e = *m_impl; + Impl::Slot& slot = e.slots.at(surface.slot); + EncodedPacket packet; + try + { + e.context.make_current(); + // NVENC reads the surface directly; wait for this stream's work (the warp) to finish writing it. + if (cudaStreamSynchronize(stream) != cudaSuccess) + throw std::runtime_error("NvencEncoder: the GPU work feeding the encoder failed"); + + NV_ENC_MAP_INPUT_RESOURCE map{}; + map.version = NV_ENC_MAP_INPUT_RESOURCE_VER; + map.registeredResource = slot.registered; + check_nv(e.fn.nvEncMapInputResource(e.encoder, &map), "nvEncMapInputResource"); + // Unmap on every exit path: a registered surface that stays mapped cannot be reused or freed. + struct MapGuard + { + Impl& impl; + NV_ENC_INPUT_PTR mapped; + ~MapGuard() + { + impl.fn.nvEncUnmapInputResource(impl.encoder, mapped); + } + } map_guard{ e, map.mappedResource }; + + NV_ENC_PIC_PARAMS pic{}; + pic.version = NV_ENC_PIC_PARAMS_VER; + pic.inputBuffer = map.mappedResource; + pic.bufferFmt = map.mappedBufferFmt; + pic.inputWidth = e.config.geometry.width; + pic.inputHeight = e.config.geometry.height; + pic.inputPitch = static_cast(slot.pitch); + pic.outputBitstream = e.bitstream; + pic.pictureStruct = NV_ENC_PIC_STRUCT_FRAME; + pic.inputTimeStamp = key; + if (force_idr) + pic.encodePicFlags = NV_ENC_PIC_FLAG_FORCEIDR | NV_ENC_PIC_FLAG_OUTPUT_SPSPPS; + check_nv(e.fn.nvEncEncodePicture(e.encoder, &pic), "nvEncEncodePicture"); + + NV_ENC_LOCK_BITSTREAM lock{}; + lock.version = NV_ENC_LOCK_BITSTREAM_VER; + lock.outputBitstream = e.bitstream; + check_nv(e.fn.nvEncLockBitstream(e.encoder, &lock), "nvEncLockBitstream"); + struct LockGuard + { + Impl& impl; + ~LockGuard() + { + impl.fn.nvEncUnlockBitstream(impl.encoder, impl.bitstream); + } + } lock_guard{ e }; + const auto* bytes = static_cast(lock.bitstreamBufferPtr); + packet.data.assign(bytes, bytes + lock.bitstreamSizeInBytes); + packet.key = lock.outputTimeStamp; + packet.keyframe = lock.pictureType == NV_ENC_PIC_TYPE_IDR; + } + catch (...) + { + slot.in_use = false; + throw; + } + slot.in_use = false; + return packet; +} + +} // namespace video_stage diff --git a/src/video_stage/core/CMakeLists.txt b/src/video_stage/core/CMakeLists.txt new file mode 100644 index 0000000000..8935189570 --- /dev/null +++ b/src/video_stage/core/CMakeLists.txt @@ -0,0 +1,23 @@ +# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# SPDX-License-Identifier: Apache-2.0 + +add_library(video_stage_core STATIC + sine_warp.cu +) +add_library(video_stage::core ALIAS video_stage_core) + +set_target_properties(video_stage_core PROPERTIES + CUDA_SEPARABLE_COMPILATION OFF + POSITION_INDEPENDENT_CODE ON +) + +target_include_directories(video_stage_core + INTERFACE + $ +) + +target_link_libraries(video_stage_core + PUBLIC + video_stage::wire + CUDA::cudart +) diff --git a/src/video_stage/core/inc/video_stage/gpu_operation.hpp b/src/video_stage/core/inc/video_stage/gpu_operation.hpp new file mode 100644 index 0000000000..62f8008718 --- /dev/null +++ b/src/video_stage/core/inc/video_stage/gpu_operation.hpp @@ -0,0 +1,42 @@ +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: Apache-2.0 + +#pragma once + +#include "nv12.hpp" + +#include + +namespace video_stage +{ + +/** + * @brief One per-frame GPU image operation: NV12 in, NV12 out, fixed geometry. + * + * Contract: + * - prepare() runs once per input geometry, outside the frame loop. It may allocate and upload; enqueue() may not. + * - enqueue() only queues work on `stream` and returns. It never synchronises the device or the stream, and + * allocates nothing. When it returns, the output is not yet written: the caller records a CUDA event on + * `stream` after the call, and whoever consumes `out` waits on that event. + * - `in` and `out` are distinct surfaces; the operation is not in-place. `in` must stay valid until the work + * completes. `out` has output_geometry() and carries the input's color description. + * - Frame identity, timestamps and metadata are not an operation's business: they travel with the frame. + * - If enqueue() throws, earlier launches of that call may already be in flight on `stream`. The caller drains + * the stream (or keeps the surfaces leased) before it recycles `in` or `out`. + */ +class GpuOperation +{ +public: + virtual ~GpuOperation() = default; + + /// The geometry this operation produces for `input`. Constant for a given operation and input. + virtual Geometry output_geometry(const Geometry& input) const = 0; + + /// Prepare resources for frames of this input geometry. Throws std::invalid_argument if unsupported. + virtual void prepare(const Geometry& input) = 0; + + /// Queue the operation on `stream`. Throws std::logic_error if not prepared for in.geometry. + virtual void enqueue(const Nv12View& in, const Nv12View& out, cudaStream_t stream) = 0; +}; + +} // namespace video_stage diff --git a/src/video_stage/core/inc/video_stage/nv12.hpp b/src/video_stage/core/inc/video_stage/nv12.hpp new file mode 100644 index 0000000000..0d31f451b1 --- /dev/null +++ b/src/video_stage/core/inc/video_stage/nv12.hpp @@ -0,0 +1,36 @@ +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: Apache-2.0 + +#pragma once + +#include + +#include +#include + +namespace video_stage +{ + +/// Color description carried unchanged from input to output. Only the quantisation range matters to an +/// operation here: it decides the value of a border fill (black is Y=16 limited-range, Y=0 full-range). +struct ColorInfo +{ + bool full_range = false; +}; + +/// A device-memory NV12 image: a luma plane followed by one interleaved chroma (U,V) plane at half +/// resolution in both directions. Pitches are in bytes and may exceed the row size. +/// +/// A view does not own its memory. The owner decides how long the planes stay valid; an operation only +/// reads or writes them on the stream it is given. +struct Nv12View +{ + uint8_t* y = nullptr; + uint8_t* uv = nullptr; + size_t y_pitch = 0; + size_t uv_pitch = 0; + Geometry geometry; + ColorInfo color; +}; + +} // namespace video_stage diff --git a/src/video_stage/core/inc/video_stage/sine_warp.hpp b/src/video_stage/core/inc/video_stage/sine_warp.hpp new file mode 100644 index 0000000000..76f6719441 --- /dev/null +++ b/src/video_stage/core/inc/video_stage/sine_warp.hpp @@ -0,0 +1,49 @@ +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: Apache-2.0 + +#pragma once + +#include "gpu_operation.hpp" + +namespace video_stage +{ + +/// Parameters of the example warp. The defaults are a visible but modest ripple. +struct SineWarpParams +{ + /// Peak horizontal displacement in luma pixels. Zero makes the operation an identity. + float amplitude_px = 8.0F; + /// Rows per full sine period. + float wavelength_px = 128.0F; + /// Phase in radians, applied at row 0. + float phase_rad = 0.0F; +}; + +/** + * @brief EXAMPLE operation: a horizontal sine ripple, not calibrated rectification. + * + * Destination-to-source mapping, in luma coordinates: + * source_x = x + amplitude * sin(2*pi * y / wavelength + phase), source_y = y + * Output pixel (x, y) is the bilinear sample of the input at (source_x, source_y). A source position outside + * the image gives the border fill (black for the declared range). Output size and color are unchanged. The warped + * picture no longer has the source camera's calibrated projection. + * + * Chroma is sampled from its own plane at the position the same mapping gives for the chroma sample's luma-space + * location, with the H.264 default siting: chroma sits on even luma columns and half way between luma rows. A + * replacement operation (the fisheye remap) keeps the same enqueue contract and swaps only the mapping. + */ +class SineWarp final : public GpuOperation +{ +public: + explicit SineWarp(const SineWarpParams& params = {}); + + Geometry output_geometry(const Geometry& input) const override; + void prepare(const Geometry& input) override; + void enqueue(const Nv12View& in, const Nv12View& out, cudaStream_t stream) override; + +private: + SineWarpParams m_params; + Geometry m_prepared; +}; + +} // namespace video_stage diff --git a/src/video_stage/core/sine_warp.cu b/src/video_stage/core/sine_warp.cu new file mode 100644 index 0000000000..d039512155 --- /dev/null +++ b/src/video_stage/core/sine_warp.cu @@ -0,0 +1,151 @@ +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: Apache-2.0 + +#include "inc/video_stage/sine_warp.hpp" + +#include +#include +#include +#include +#include + +namespace video_stage +{ + +namespace +{ + +constexpr float kTwoPi = 6.28318530717958647692F; +constexpr unsigned kBlockX = 32; +constexpr unsigned kBlockY = 8; + +struct WarpArgs +{ + float amplitude; + float cycles_per_row; // 1 / wavelength + float phase; +}; + +__device__ __forceinline__ float displaced_x(float x, float y, const WarpArgs& a) +{ + return x + a.amplitude * sinf(kTwoPi * y * a.cycles_per_row + a.phase); +} + +// Bilinear sample of one byte channel. `stride` is the byte distance between neighbouring samples of the channel +// (1 for luma, 2 for interleaved chroma). The caller has checked that (sx, sy) lies inside [0,w-1]x[0,h-1]. +__device__ __forceinline__ float bilinear(const uint8_t* plane, size_t pitch, int stride, int w, int h, float sx, float sy) +{ + const int x0 = static_cast(floorf(sx)); + const int y0 = static_cast(floorf(sy)); + const int x1 = min(x0 + 1, w - 1); + const int y1 = min(y0 + 1, h - 1); + const float fx = sx - static_cast(x0); + const float fy = sy - static_cast(y0); + const uint8_t* r0 = plane + static_cast(y0) * pitch; + const uint8_t* r1 = plane + static_cast(y1) * pitch; + const float top = (1.0F - fx) * r0[x0 * stride] + fx * r0[x1 * stride]; + const float bottom = (1.0F - fx) * r1[x0 * stride] + fx * r1[x1 * stride]; + return (1.0F - fy) * top + fy * bottom; +} + +__device__ __forceinline__ uint8_t to_byte(float v) +{ + return static_cast(fminf(fmaxf(v + 0.5F, 0.0F), 255.0F)); +} + +__global__ void warp_luma( + const uint8_t* in, size_t in_pitch, uint8_t* out, size_t out_pitch, int w, int h, WarpArgs args, uint8_t fill) +{ + const int x = blockIdx.x * blockDim.x + threadIdx.x; + const int y = blockIdx.y * blockDim.y + threadIdx.y; + if (x >= w || y >= h) + return; + const float sx = displaced_x(static_cast(x), static_cast(y), args); + const float sy = static_cast(y); + const bool inside = sx >= 0.0F && sx <= static_cast(w - 1); + out[static_cast(y) * out_pitch + x] = inside ? to_byte(bilinear(in, in_pitch, 1, w, h, sx, sy)) : fill; +} + +// Chroma sample (cx, cy) sits at luma position (2*cx, 2*cy + 0.5): co-sited with even luma columns, half way between +// luma rows. The mapping is applied there and the result is converted back to chroma-plane coordinates. +__global__ void warp_chroma(const uint8_t* in, size_t in_pitch, uint8_t* out, size_t out_pitch, int cw, int ch, WarpArgs args) +{ + const int cx = blockIdx.x * blockDim.x + threadIdx.x; + const int cy = blockIdx.y * blockDim.y + threadIdx.y; + if (cx >= cw || cy >= ch) + return; + const float luma_x = 2.0F * static_cast(cx); + const float luma_y = 2.0F * static_cast(cy) + 0.5F; + const float src_luma_x = displaced_x(luma_x, luma_y, args); + const float scx = 0.5F * src_luma_x; + const float scy = 0.5F * (luma_y - 0.5F); + uint8_t* dst = out + static_cast(cy) * out_pitch + 2 * cx; + // A positive predicate: a NaN coordinate fails it and takes the fill path instead of reaching the sampler. + if (!(scx >= 0.0F && scx <= static_cast(cw - 1))) + { + dst[0] = 128; + dst[1] = 128; + return; + } + dst[0] = to_byte(bilinear(in, in_pitch, 2, cw, ch, scx, scy)); + dst[1] = to_byte(bilinear(in + 1, in_pitch, 2, cw, ch, scx, scy)); +} + +void check(cudaError_t status, const char* what) +{ + if (status != cudaSuccess) + throw std::runtime_error(std::string("SineWarp: ") + what + ": " + cudaGetErrorString(status)); +} + +} // namespace + +SineWarp::SineWarp(const SineWarpParams& params) : m_params(params) +{ + const bool finite = + std::isfinite(params.amplitude_px) && std::isfinite(params.wavelength_px) && std::isfinite(params.phase_rad); + if (!finite || !(params.wavelength_px > 0.0F) || !std::isfinite(1.0F / params.wavelength_px)) + throw std::invalid_argument( + "SineWarp: amplitude, phase and wavelength must be finite and the wavelength positive"); +} + +Geometry SineWarp::output_geometry(const Geometry& input) const +{ + return input; +} + +void SineWarp::prepare(const Geometry& input) +{ + if (input.width == 0 || input.height == 0 || (input.width % 2) != 0 || (input.height % 2) != 0) + throw std::invalid_argument("SineWarp: NV12 needs non-zero even dimensions"); + m_prepared = input; +} + +void SineWarp::enqueue(const Nv12View& in, const Nv12View& out, cudaStream_t stream) +{ + if (m_prepared.width == 0 || !(in.geometry == m_prepared)) + throw std::logic_error("SineWarp: not prepared for this input geometry"); + if (!(out.geometry == output_geometry(in.geometry))) + throw std::invalid_argument("SineWarp: output geometry does not match"); + const auto bad = [](const Nv12View& v) + { return v.y == nullptr || v.uv == nullptr || v.y_pitch < v.geometry.width || v.uv_pitch < v.geometry.width; }; + if (bad(in) || bad(out)) + throw std::invalid_argument("SineWarp: surface has a null plane or a pitch below the row size"); + + const int w = static_cast(in.geometry.width); + const int h = static_cast(in.geometry.height); + const WarpArgs args{ m_params.amplitude_px, 1.0F / m_params.wavelength_px, m_params.phase_rad }; + const uint8_t fill = in.color.full_range ? 0 : 16; + + const dim3 block(kBlockX, kBlockY); + const dim3 luma_grid((w + kBlockX - 1) / kBlockX, (h + kBlockY - 1) / kBlockY); + warp_luma<<>>(in.y, in.y_pitch, out.y, out.y_pitch, w, h, args, fill); + check(cudaGetLastError(), "luma launch"); + + const int cw = w / 2; + const int ch = h / 2; + const dim3 chroma_grid((cw + kBlockX - 1) / kBlockX, (ch + kBlockY - 1) / kBlockY); + warp_chroma<<>>(in.uv, in.uv_pitch, out.uv, out.uv_pitch, cw, ch, args); + check(cudaGetLastError(), "chroma launch"); +} + +} // namespace video_stage diff --git a/src/video_stage/tests/CMakeLists.txt b/src/video_stage/tests/CMakeLists.txt index 272ec56770..9f7a5b4f9d 100644 --- a/src/video_stage/tests/CMakeLists.txt +++ b/src/video_stage/tests/CMakeLists.txt @@ -1,7 +1,18 @@ # SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. # SPDX-License-Identifier: Apache-2.0 -# The relay's receiver compiled into its own test, so the test links no LiveKit. +if(BUILD_VIDEO_STAGE_GPU) + add_executable(video_stage_tests + test_sine_warp.cpp + ) + + target_link_libraries(video_stage_tests PRIVATE video_stage::core Catch2::Catch2WithMain) + + catch_discover_tests(video_stage_tests) +endif() + +# The receiver is part of the worker executable, not a library, so its source is compiled in directly. It needs +# only POSIX sockets and the AVS1 codec: no CUDA device. add_executable(video_stage_tcp_receiver_tests test_tcp_receiver.cpp ${CMAKE_CURRENT_SOURCE_DIR}/../app/tcp_receiver.cpp diff --git a/src/video_stage/tests/test_sine_warp.cpp b/src/video_stage/tests/test_sine_warp.cpp new file mode 100644 index 0000000000..49e4f089af --- /dev/null +++ b/src/video_stage/tests/test_sine_warp.cpp @@ -0,0 +1,540 @@ +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: Apache-2.0 + +#include +#include + +#include +#include +#include +#include +#include +#include +#include +#include +#include + +using video_stage::ColorInfo; +using video_stage::Geometry; +using video_stage::Nv12View; +using video_stage::SineWarp; +using video_stage::SineWarpParams; + +namespace +{ + +constexpr uint8_t kPadding = 0xAB; +constexpr float kTwoPi = 6.28318530717958647692F; + +void cuda_ok(cudaError_t status) +{ + if (status != cudaSuccess) + FAIL("CUDA error: " << cudaGetErrorString(status)); +} + +bool have_cuda_device() +{ + int count = 0; + return cudaGetDeviceCount(&count) == cudaSuccess && count > 0; +} + +// Host-side NV12 image with tightly packed rows. +struct HostNv12 +{ + uint32_t w = 0; + uint32_t h = 0; + std::vector y; // w * h + std::vector uv; // w * (h / 2), interleaved U,V +}; + +HostNv12 make_image(uint32_t w, uint32_t h) +{ + HostNv12 img; + img.w = w; + img.h = h; + img.y.resize(static_cast(w) * h); + img.uv.resize(static_cast(w) * (h / 2)); + for (uint32_t y = 0; y < h; ++y) + for (uint32_t x = 0; x < w; ++x) + img.y[y * w + x] = static_cast(20 + (3 * x + 2 * y + ((x * y) % 7)) % 200); + for (uint32_t cy = 0; cy < h / 2; ++cy) + for (uint32_t cx = 0; cx < w / 2; ++cx) + { + img.uv[cy * w + 2 * cx] = static_cast(30 + (5 * cx + cy) % 190); // U + img.uv[cy * w + 2 * cx + 1] = static_cast(220 - (2 * cx + 4 * cy) % 180); // V + } + return img; +} + +// Device NV12 surface with pitch deliberately larger than the row size. +class DeviceSurface +{ +public: + DeviceSurface(uint32_t w, uint32_t h, bool full_range = false) : m_w(w), m_h(h) + { + cuda_ok(cudaMallocPitch(reinterpret_cast(&m_y), &m_y_pitch, w + 64, h)); + cuda_ok(cudaMallocPitch(reinterpret_cast(&m_uv), &m_uv_pitch, w + 64, h / 2)); + REQUIRE(m_y_pitch > w); + REQUIRE(m_uv_pitch > w); + cuda_ok(cudaMemset2D(m_y, m_y_pitch, kPadding, m_y_pitch, h)); + cuda_ok(cudaMemset2D(m_uv, m_uv_pitch, kPadding, m_uv_pitch, h / 2)); + m_color.full_range = full_range; + } + ~DeviceSurface() + { + cudaFree(m_y); + cudaFree(m_uv); + } + DeviceSurface(const DeviceSurface&) = delete; + DeviceSurface& operator=(const DeviceSurface&) = delete; + + void upload(const HostNv12& img) + { + cuda_ok(cudaMemcpy2D(m_y, m_y_pitch, img.y.data(), m_w, m_w, m_h, cudaMemcpyHostToDevice)); + cuda_ok(cudaMemcpy2D(m_uv, m_uv_pitch, img.uv.data(), m_w, m_w, m_h / 2, cudaMemcpyHostToDevice)); + } + + // Copies the visible area out and checks that the row padding was left untouched. + HostNv12 download() const + { + HostNv12 img; + img.w = m_w; + img.h = m_h; + img.y.resize(static_cast(m_w) * m_h); + img.uv.resize(static_cast(m_w) * (m_h / 2)); + cuda_ok(cudaMemcpy2D(img.y.data(), m_w, m_y, m_y_pitch, m_w, m_h, cudaMemcpyDeviceToHost)); + cuda_ok(cudaMemcpy2D(img.uv.data(), m_w, m_uv, m_uv_pitch, m_w, m_h / 2, cudaMemcpyDeviceToHost)); + check_padding(m_y, m_y_pitch, m_h); + check_padding(m_uv, m_uv_pitch, m_h / 2); + return img; + } + + Nv12View view() const + { + Nv12View v; + v.y = m_y; + v.uv = m_uv; + v.y_pitch = m_y_pitch; + v.uv_pitch = m_uv_pitch; + v.geometry = { m_w, m_h }; + v.color = m_color; + return v; + } + +private: + void check_padding(const uint8_t* plane, size_t pitch, uint32_t rows) const + { + std::vector pad(pitch - m_w); + for (uint32_t r = 0; r < rows; ++r) + { + cuda_ok(cudaMemcpy(pad.data(), plane + r * pitch + m_w, pad.size(), cudaMemcpyDeviceToHost)); + for (uint8_t b : pad) + REQUIRE(b == kPadding); + } + } + + uint32_t m_w, m_h; + uint8_t* m_y = nullptr; + uint8_t* m_uv = nullptr; + size_t m_y_pitch = 0, m_uv_pitch = 0; + ColorInfo m_color; +}; + +// CPU reference of the documented mapping. +float bilinear(const std::vector& p, int stride, int ofs, int row_bytes, int w, int h, float sx, float sy) +{ + const int x0 = static_cast(std::floor(sx)); + const int y0 = static_cast(std::floor(sy)); + const int x1 = std::min(x0 + 1, w - 1); + const int y1 = std::min(y0 + 1, h - 1); + const float fx = sx - static_cast(x0); + const float fy = sy - static_cast(y0); + auto at = [&](int x, int y) { return static_cast(p[y * row_bytes + x * stride + ofs]); }; + const float top = (1.0F - fx) * at(x0, y0) + fx * at(x1, y0); + const float bottom = (1.0F - fx) * at(x0, y1) + fx * at(x1, y1); + return (1.0F - fy) * top + fy * bottom; +} + +uint8_t to_byte(float v) +{ + return static_cast(std::min(std::max(v + 0.5F, 0.0F), 255.0F)); +} + +float displaced_x(float x, float y, const SineWarpParams& p) +{ + return x + p.amplitude_px * std::sin(kTwoPi * y * (1.0F / p.wavelength_px) + p.phase_rad); +} + +struct Reference +{ + HostNv12 img; + std::vector luma_filled; // per luma pixel + std::vector chroma_filled; // per chroma sample pair +}; + +Reference reference_warp(const HostNv12& in, const SineWarpParams& p, bool full_range) +{ + const int w = static_cast(in.w); + const int h = static_cast(in.h); + const int cw = w / 2; + const int ch = h / 2; + Reference ref; + ref.img.w = in.w; + ref.img.h = in.h; + ref.img.y.resize(in.y.size()); + ref.img.uv.resize(in.uv.size()); + ref.luma_filled.assign(in.y.size(), false); + ref.chroma_filled.assign(static_cast(cw) * ch, false); + const uint8_t fill = full_range ? 0 : 16; + for (int y = 0; y < h; ++y) + for (int x = 0; x < w; ++x) + { + const float sx = displaced_x(static_cast(x), static_cast(y), p); + if (sx >= 0.0F && sx <= static_cast(w - 1)) + ref.img.y[y * w + x] = to_byte(bilinear(in.y, 1, 0, w, w, h, sx, static_cast(y))); + else + { + ref.img.y[y * w + x] = fill; + ref.luma_filled[y * w + x] = true; + } + } + for (int cy = 0; cy < ch; ++cy) + for (int cx = 0; cx < cw; ++cx) + { + const float luma_y = 2.0F * static_cast(cy) + 0.5F; + const float scx = 0.5F * displaced_x(2.0F * static_cast(cx), luma_y, p); + const float scy = 0.5F * (luma_y - 0.5F); + if (scx < 0.0F || scx > static_cast(cw - 1)) + { + ref.img.uv[cy * w + 2 * cx] = 128; + ref.img.uv[cy * w + 2 * cx + 1] = 128; + ref.chroma_filled[cy * cw + cx] = true; + } + else + { + ref.img.uv[cy * w + 2 * cx] = to_byte(bilinear(in.uv, 2, 0, w, cw, ch, scx, scy)); + ref.img.uv[cy * w + 2 * cx + 1] = to_byte(bilinear(in.uv, 2, 1, w, cw, ch, scx, scy)); + } + } + return ref; +} + +void run_once(SineWarp& op, const DeviceSurface& in, const DeviceSurface& out) +{ + cudaStream_t s = nullptr; + cuda_ok(cudaStreamCreate(&s)); + op.enqueue(in.view(), out.view(), s); + cuda_ok(cudaStreamSynchronize(s)); + cuda_ok(cudaStreamDestroy(s)); +} + +// Compares two images within `tol` levels; returns the worst difference. +int max_diff(const std::vector& a, const std::vector& b) +{ + REQUIRE(a.size() == b.size()); + int worst = 0; + for (size_t i = 0; i < a.size(); ++i) + worst = std::max(worst, std::abs(static_cast(a[i]) - static_cast(b[i]))); + return worst; +} + +void check_against_reference(uint32_t w, uint32_t h, const SineWarpParams& params, bool full_range = false) +{ + const HostNv12 src = make_image(w, h); + DeviceSurface in(w, h, full_range), out(w, h, full_range); + in.upload(src); + SineWarp op(params); + op.prepare({ w, h }); + run_once(op, in, out); + const HostNv12 got = out.download(); + const Reference ref = reference_warp(src, params, full_range); + CHECK(max_diff(got.y, ref.img.y) <= 1); + CHECK(max_diff(got.uv, ref.img.uv) <= 1); + // Border-filled pixels are exact. + for (size_t i = 0; i < ref.luma_filled.size(); ++i) + if (ref.luma_filled[i]) + REQUIRE(got.y[i] == (full_range ? 0 : 16)); + const size_t cw = w / 2; + for (size_t i = 0; i < ref.chroma_filled.size(); ++i) + if (ref.chroma_filled[i]) + { + const size_t row = i / cw, col = i % cw; + REQUIRE(got.uv[row * w + 2 * col] == 128); + REQUIRE(got.uv[row * w + 2 * col + 1] == 128); + } +} + +struct Size +{ + uint32_t w, h; +}; + +} // namespace + +#define REQUIRE_CUDA_OR_SKIP() \ + do \ + { \ + if (!have_cuda_device()) \ + SKIP("No CUDA device available"); \ + } while (false) + +TEST_CASE("SineWarp zero amplitude is a bit-exact identity", "[video_stage][gpu]") +{ + REQUIRE_CUDA_OR_SKIP(); + for (const Size sz : { Size{ 64, 48 }, Size{ 130, 66 } }) + { + SECTION(std::to_string(sz.w) + "x" + std::to_string(sz.h)) + { + const HostNv12 src = make_image(sz.w, sz.h); + DeviceSurface in(sz.w, sz.h), out(sz.w, sz.h); + in.upload(src); + SineWarp op(SineWarpParams{ 0.0F, 37.0F, 0.7F }); + op.prepare({ sz.w, sz.h }); + run_once(op, in, out); + const HostNv12 got = out.download(); + CHECK(got.y == src.y); + CHECK(got.uv == src.uv); + } + } +} + +TEST_CASE("SineWarp matches the CPU reference", "[video_stage][gpu]") +{ + REQUIRE_CUDA_OR_SKIP(); + for (const Size sz : { Size{ 64, 48 }, Size{ 130, 66 } }) + { + SECTION("default-ish parameters " + std::to_string(sz.w) + "x" + std::to_string(sz.h)) + { + check_against_reference(sz.w, sz.h, SineWarpParams{ 5.0F, 32.0F, 0.0F }); + } + SECTION("phase and wavelength " + std::to_string(sz.w) + "x" + std::to_string(sz.h)) + { + check_against_reference(sz.w, sz.h, SineWarpParams{ 3.5F, 21.0F, 1.1F }); + } + SECTION("large amplitude leaves the image " + std::to_string(sz.w) + "x" + std::to_string(sz.h)) + { + check_against_reference(sz.w, sz.h, SineWarpParams{ 40.0F, 24.0F, 0.4F }); + } + SECTION("full-range border fill " + std::to_string(sz.w) + "x" + std::to_string(sz.h)) + { + check_against_reference(sz.w, sz.h, SineWarpParams{ 40.0F, 24.0F, 0.4F }, true); + } + } +} + +TEST_CASE("SineWarp large amplitude actually produces border fill", "[video_stage][gpu]") +{ + REQUIRE_CUDA_OR_SKIP(); + const SineWarpParams params{ 40.0F, 24.0F, 0.4F }; + const HostNv12 src = make_image(64, 48); + const Reference ref = reference_warp(src, params, false); + CHECK(std::count(ref.luma_filled.begin(), ref.luma_filled.end(), true) > 0); + CHECK(std::count(ref.chroma_filled.begin(), ref.chroma_filled.end(), true) > 0); +} + +TEST_CASE("SineWarp keeps constant chroma constant in the interior", "[video_stage][gpu]") +{ + REQUIRE_CUDA_OR_SKIP(); + const uint32_t w = 130, h = 66; + HostNv12 src = make_image(w, h); + for (uint32_t cy = 0; cy < h / 2; ++cy) + for (uint32_t cx = 0; cx < w / 2; ++cx) + { + src.uv[cy * w + 2 * cx] = 90; + src.uv[cy * w + 2 * cx + 1] = 200; + } + const SineWarpParams params{ 4.0F, 30.0F, 0.3F }; + DeviceSurface in(w, h), out(w, h); + in.upload(src); + SineWarp op(params); + op.prepare({ w, h }); + run_once(op, in, out); + const HostNv12 got = out.download(); + const Reference ref = reference_warp(src, params, false); + size_t interior = 0; + for (uint32_t cy = 0; cy < h / 2; ++cy) + for (uint32_t cx = 0; cx < w / 2; ++cx) + { + if (ref.chroma_filled[cy * (w / 2) + cx]) + continue; + ++interior; + REQUIRE(got.uv[cy * w + 2 * cx] == 90); + REQUIRE(got.uv[cy * w + 2 * cx + 1] == 200); + } + CHECK(interior > 0); +} + +TEST_CASE("SineWarp chroma follows the siting, not luma/2", "[video_stage][gpu]") +{ + REQUIRE_CUDA_OR_SKIP(); + // Vertical chroma edge: U steps from 40 to 220 at chroma column 16. Compared with the reference the output + // pins the half-luma-row siting of the displacement. + const uint32_t w = 64, h = 48; + HostNv12 src = make_image(w, h); + for (uint32_t cy = 0; cy < h / 2; ++cy) + for (uint32_t cx = 0; cx < w / 2; ++cx) + { + src.uv[cy * w + 2 * cx] = cx < 16 ? 40 : 220; + src.uv[cy * w + 2 * cx + 1] = 128; + } + const SineWarpParams params{ 6.0F, 16.0F, 0.0F }; + DeviceSurface in(w, h), out(w, h); + in.upload(src); + SineWarp op(params); + op.prepare({ w, h }); + run_once(op, in, out); + const HostNv12 got = out.download(); + const Reference ref = reference_warp(src, params, false); + CHECK(max_diff(got.uv, ref.img.uv) <= 1); + + // The naive mapping (sample at luma_y = 2*cy, no half-row offset) must differ somewhere on this edge. + SineWarpParams p = params; + int naive_worst = 0; + for (uint32_t cy = 0; cy < h / 2; ++cy) + for (uint32_t cx = 0; cx < w / 2; ++cx) + { + const float scx = 0.5F * displaced_x(2.0F * static_cast(cx), 2.0F * static_cast(cy), p); + if (scx < 0.0F || scx > static_cast(w / 2 - 1)) + continue; + const uint8_t naive = to_byte(bilinear(src.uv, 2, 0, w, w / 2, h / 2, scx, static_cast(cy))); + naive_worst = + std::max(naive_worst, std::abs(static_cast(naive) - static_cast(got.uv[cy * w + 2 * cx]))); + } + CHECK(naive_worst > 1); +} + +TEST_CASE("SineWarp contract violations throw", "[video_stage][gpu]") +{ + SECTION("non-positive wavelength throws in the constructor") + { + CHECK_THROWS_AS(SineWarp(SineWarpParams{ 1.0F, 0.0F, 0.0F }), std::invalid_argument); + CHECK_THROWS_AS(SineWarp(SineWarpParams{ 1.0F, -5.0F, 0.0F }), std::invalid_argument); + } + SECTION("prepare rejects odd or zero dimensions") + { + SineWarp op; + CHECK_THROWS_AS(op.prepare({ 0, 48 }), std::invalid_argument); + CHECK_THROWS_AS(op.prepare({ 64, 0 }), std::invalid_argument); + CHECK_THROWS_AS(op.prepare({ 63, 48 }), std::invalid_argument); + CHECK_THROWS_AS(op.prepare({ 64, 47 }), std::invalid_argument); + } + + REQUIRE_CUDA_OR_SKIP(); + DeviceSurface in(64, 48), out(64, 48); + SineWarp op; + + SECTION("enqueue before prepare throws std::logic_error") + { + CHECK_THROWS_AS(op.enqueue(in.view(), out.view(), nullptr), std::logic_error); + } + SECTION("enqueue with a different input geometry throws std::logic_error") + { + op.prepare({ 128, 96 }); + CHECK_THROWS_AS(op.enqueue(in.view(), out.view(), nullptr), std::logic_error); + } + + op.prepare({ 64, 48 }); + + SECTION("mismatched output geometry") + { + DeviceSurface other(32, 48); + CHECK_THROWS_AS(op.enqueue(in.view(), other.view(), nullptr), std::invalid_argument); + } + SECTION("null planes") + { + Nv12View bad = out.view(); + bad.y = nullptr; + CHECK_THROWS_AS(op.enqueue(in.view(), bad, nullptr), std::invalid_argument); + bad = out.view(); + bad.uv = nullptr; + CHECK_THROWS_AS(op.enqueue(in.view(), bad, nullptr), std::invalid_argument); + bad = in.view(); + bad.y = nullptr; + CHECK_THROWS_AS(op.enqueue(bad, out.view(), nullptr), std::invalid_argument); + } + SECTION("pitch below the row size") + { + Nv12View bad = out.view(); + bad.y_pitch = 63; + CHECK_THROWS_AS(op.enqueue(in.view(), bad, nullptr), std::invalid_argument); + bad = out.view(); + bad.uv_pitch = 63; + CHECK_THROWS_AS(op.enqueue(in.view(), bad, nullptr), std::invalid_argument); + bad = in.view(); + bad.y_pitch = 10; + CHECK_THROWS_AS(op.enqueue(bad, out.view(), nullptr), std::invalid_argument); + } +} + +TEST_CASE("SineWarp enqueue is asynchronous and correct after the caller syncs", "[video_stage][gpu]") +{ + REQUIRE_CUDA_OR_SKIP(); + const SineWarpParams params{ 5.0F, 32.0F, 0.2F }; + const HostNv12 src = make_image(130, 66); + DeviceSurface in(130, 66), out(130, 66); + in.upload(src); + SineWarp op(params); + op.prepare({ 130, 66 }); + cudaStream_t s = nullptr; + cuda_ok(cudaStreamCreate(&s)); + cudaEvent_t done = nullptr; + cuda_ok(cudaEventCreate(&done)); + op.enqueue(in.view(), out.view(), s); + cuda_ok(cudaEventRecord(done, s)); + cuda_ok(cudaEventSynchronize(done)); + const HostNv12 got = out.download(); + const Reference ref = reference_warp(src, params, false); + CHECK(max_diff(got.y, ref.img.y) <= 1); + CHECK(max_diff(got.uv, ref.img.uv) <= 1); + cuda_ok(cudaEventDestroy(done)); + cuda_ok(cudaStreamDestroy(s)); +} + +TEST_CASE("SineWarp prepared operation is reusable across surfaces and streams", "[video_stage][gpu]") +{ + REQUIRE_CUDA_OR_SKIP(); + const SineWarpParams params{ 5.0F, 32.0F, 0.2F }; + const uint32_t w = 130, h = 66; + const HostNv12 a = make_image(w, h); + HostNv12 b = make_image(w, h); + for (auto& v : b.y) + v = static_cast(255 - v); + for (auto& v : b.uv) + v = static_cast(255 - v); + + DeviceSurface in_a(w, h), in_b(w, h), out_a(w, h), out_b(w, h); + in_a.upload(a); + in_b.upload(b); + SineWarp op(params); + op.prepare({ w, h }); + const Reference ref_a = reference_warp(a, params, false); + const Reference ref_b = reference_warp(b, params, false); + + SECTION("two surfaces back-to-back on one stream") + { + cudaStream_t s = nullptr; + cuda_ok(cudaStreamCreate(&s)); + op.enqueue(in_a.view(), out_a.view(), s); + op.enqueue(in_b.view(), out_b.view(), s); + cuda_ok(cudaStreamSynchronize(s)); + cuda_ok(cudaStreamDestroy(s)); + } + SECTION("two surfaces on two streams") + { + cudaStream_t s1 = nullptr, s2 = nullptr; + cuda_ok(cudaStreamCreate(&s1)); + cuda_ok(cudaStreamCreate(&s2)); + op.enqueue(in_a.view(), out_a.view(), s1); + op.enqueue(in_b.view(), out_b.view(), s2); + cuda_ok(cudaStreamSynchronize(s1)); + cuda_ok(cudaStreamSynchronize(s2)); + cuda_ok(cudaStreamDestroy(s1)); + cuda_ok(cudaStreamDestroy(s2)); + } + const HostNv12 got_a = out_a.download(); + const HostNv12 got_b = out_b.download(); + CHECK(max_diff(got_a.y, ref_a.img.y) <= 1); + CHECK(max_diff(got_a.uv, ref_a.img.uv) <= 1); + CHECK(max_diff(got_b.y, ref_b.img.y) <= 1); + CHECK(max_diff(got_b.uv, ref_b.img.uv) <= 1); +}