diff --git a/examples/camera_streamer/README.md b/examples/camera_streamer/README.md index 2d8fc367ab..4685b51c6e 100644 --- a/examples/camera_streamer/README.md +++ b/examples/camera_streamer/README.md @@ -37,6 +37,25 @@ argus_sender/build/argus_sender --config "configs/${PROFILE}.yaml" The sender listens for a TCP receiver on port 5000 by default. +## Workstation NVDEC decoder + +On a Linux workstation with an NVIDIA GPU, CUDA toolkit, CMake and `uv`, run from +`examples/camera_streamer`: + +```bash +uv venv .venv --python 3.12 +uv pip install --python .venv/bin/python pybind11 numpy +cmake -S camera_receiver/codec -B camera_receiver/codec/build \ + -DCMAKE_BUILD_TYPE=Release -DPython3_EXECUTABLE="$PWD/.venv/bin/python" +cmake --build camera_receiver/codec/build --parallel 2 +``` + +The build downloads the pinned NVIDIA Video Codec SDK. CPU-side Python tests: + +```bash +uv run --project tests --extra dev python -m pytest tests/python -q +``` + ## Licensing This example is Apache-2.0 licensed. Existing upstream notices and external SDK diff --git a/examples/camera_streamer/camera_receiver/.gitignore b/examples/camera_streamer/camera_receiver/.gitignore new file mode 100644 index 0000000000..216bcdfefa --- /dev/null +++ b/examples/camera_streamer/camera_receiver/.gitignore @@ -0,0 +1,27 @@ +# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# SPDX-License-Identifier: Apache-2.0 + +# Python environment and caches +.venv/ +__pycache__/ +*.py[cod] +*$py.class +.pytest_cache/ +.mypy_cache/ +.ruff_cache/ +.coverage +htmlcov/ + +# Native codec build outputs +codec/build/ +codec/_camera_receiver_codec*.so + +# Local logs and runtime artifacts +*.log +*.pid +*.sock +*.avs1 +*.hvs2 +*.nv12 +*.h265 +*.hevc diff --git a/examples/camera_streamer/camera_receiver/codec/CMakeLists.txt b/examples/camera_streamer/camera_receiver/codec/CMakeLists.txt new file mode 100644 index 0000000000..c659529f60 --- /dev/null +++ b/examples/camera_streamer/camera_receiver/codec/CMakeLists.txt @@ -0,0 +1,123 @@ +# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# SPDX-License-Identifier: Apache-2.0 +# +# Native H.264/HEVC NVDEC pybind11 module for camera_receiver. The .so +# lands next to __init__.py so the package imports without any +# install step. + +cmake_minimum_required(VERSION 3.20) + +# Ampere / Ada / Hopper / Orin / Thor. Override with +# CMAKE_CUDA_ARCHITECTURES on the configure line for fast iteration. +if(NOT DEFINED CMAKE_CUDA_ARCHITECTURES) + set(CMAKE_CUDA_ARCHITECTURES "80;86;87;89;90") +endif() + +project(camera_receiver_codec LANGUAGES C 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(CMAKE_POSITION_INDEPENDENT_CODE ON) + +# --------------------------------------------------------------------------- +# Dependencies +# --------------------------------------------------------------------------- +find_package(CUDAToolkit REQUIRED) +find_package(Python3 REQUIRED COMPONENTS Interpreter Development.Module) + +# pybind11 is pip-installed into the active camera_receiver Python environment. +execute_process( + COMMAND ${Python3_EXECUTABLE} -c "import pybind11; print(pybind11.get_cmake_dir())" + OUTPUT_VARIABLE pybind11_DIR + OUTPUT_STRIP_TRAILING_WHITESPACE + RESULT_VARIABLE _pybind11_rc +) +if(NOT _pybind11_rc EQUAL 0) + message(FATAL_ERROR "pybind11 not found. `uv pip install pybind11` into the active venv.") +endif() +find_package(pybind11 REQUIRED CONFIG) + +# NVDEC entry-point library (ships with the NVIDIA driver). +if(CMAKE_SYSTEM_PROCESSOR MATCHES "aarch64|arm64") + set(_lib_arch_dir "/usr/lib/aarch64-linux-gnu/nvidia") +else() + set(_lib_arch_dir "/usr/lib/x86_64-linux-gnu") +endif() + +find_library(NVCUVID_LIBRARY + NAMES libnvcuvid.so.1 libnvcuvid.so nvcuvid + HINTS ${CUDA_TOOLKIT_ROOT_DIR}/lib64 + PATHS "${_lib_arch_dir}" +) +if(NOT NVCUVID_LIBRARY) + message(FATAL_ERROR "libnvcuvid not found. Install the NVIDIA driver.") +endif() +message(STATUS "camera_receiver codec: NVCUVID=${NVCUVID_LIBRARY}") + +# --------------------------------------------------------------------------- +# NVIDIA Video Codec SDK headers + reference decoder sources. +# --------------------------------------------------------------------------- +set(_nvc_sdk_url + "https://edge.urm.nvidia.com/artifactory/sw-holoscan-thirdparty-generic-local/holohub/operators/holohub_nvc_13.0.19.zip" +) +set(_nvc_sdk_zip "${CMAKE_BINARY_DIR}/holohub_nvc_13.0.19.zip") +set(_nvc_sdk_vendor_zip "${CMAKE_CURRENT_SOURCE_DIR}/third_party/holohub_nvc_13.0.19.zip") +set(_nvc_sdk_dir "${CMAKE_BINARY_DIR}/_deps/nvc_sdk") + +if(EXISTS ${_nvc_sdk_vendor_zip} AND NOT EXISTS ${_nvc_sdk_zip}) + message(STATUS "Using vendored NVIDIA Video Codec SDK zip...") + file(COPY ${_nvc_sdk_vendor_zip} DESTINATION ${CMAKE_BINARY_DIR}) +endif() +if(NOT EXISTS ${_nvc_sdk_zip}) + message(STATUS "Downloading NVIDIA Video Codec SDK...") + file(DOWNLOAD ${_nvc_sdk_url} ${_nvc_sdk_zip} + SHOW_PROGRESS + EXPECTED_MD5 "a08c33c282bc92469ab8cf863aca9087" + ) +endif() +if(NOT EXISTS ${_nvc_sdk_dir}/NvDecoder) + message(STATUS "Extracting NVIDIA Video Codec SDK...") + file(MAKE_DIRECTORY ${_nvc_sdk_dir}) + execute_process( + COMMAND ${CMAKE_COMMAND} -E tar xzf ${_nvc_sdk_zip} + WORKING_DIRECTORY ${_nvc_sdk_dir} + ) +endif() + +# --------------------------------------------------------------------------- +# pybind11 module +# --------------------------------------------------------------------------- +pybind11_add_module(_camera_receiver_codec MODULE + codec_bindings.cpp + nv_video_decoder.cpp + nv12_to_rgba.cu + ${_nvc_sdk_dir}/NvDecoder/NvDecoder.cpp + ${_nvc_sdk_dir}/Logger.cpp +) + +target_include_directories(_camera_receiver_codec PRIVATE + ${_nvc_sdk_dir} + ${_nvc_sdk_dir}/NvDecoder +) + +target_compile_definitions(_camera_receiver_codec PRIVATE NV_WINDOWS=0 NV_LINUX=1) + +# NVCODEC SDK reference sources carry legacy warnings. +target_compile_options(_camera_receiver_codec PRIVATE + $<$:-Wno-reorder -Wno-uninitialized -Wno-unused-variable + -Wno-unused-parameter -Wno-missing-field-initializers -Wno-sign-compare> +) + +target_link_libraries(_camera_receiver_codec PRIVATE + CUDA::cudart + CUDA::cuda_driver + ${NVCUVID_LIBRARY} +) + +# Drop the .so next to __init__.py so the package imports straight +# from the source tree. +set_target_properties(_camera_receiver_codec PROPERTIES + LIBRARY_OUTPUT_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR} +) diff --git a/examples/camera_streamer/camera_receiver/codec/__init__.py b/examples/camera_streamer/camera_receiver/codec/__init__.py new file mode 100644 index 0000000000..f404495c38 --- /dev/null +++ b/examples/camera_streamer/camera_receiver/codec/__init__.py @@ -0,0 +1,23 @@ +# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# SPDX-License-Identifier: Apache-2.0 +"""Native NVDEC video decoder for camera_receiver.""" + +from __future__ import annotations + +try: + from ._camera_receiver_codec import ( # type: ignore[import-not-found] + Codec, + DecoderConfig, + NvVideoDecoder, + ) +except ImportError as e: + raise ImportError( + "camera_receiver native codec not built. Follow " + "the host receiver build instructions in examples/camera_streamer/README.md." + ) from e + +__all__ = [ + "Codec", + "DecoderConfig", + "NvVideoDecoder", +] diff --git a/examples/camera_streamer/camera_receiver/codec/codec_bindings.cpp b/examples/camera_streamer/camera_receiver/codec/codec_bindings.cpp new file mode 100644 index 0000000000..3a409dab76 --- /dev/null +++ b/examples/camera_streamer/camera_receiver/codec/codec_bindings.cpp @@ -0,0 +1,95 @@ +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: Apache-2.0 +// +// pybind11 bindings — public surface is ``camera_receiver.codec``. + +#include "cuda_buffers.hpp" +#include "nv_video_decoder.hpp" + +#include +#include + +#include +#include +#include + +namespace py = pybind11; +using camera_receiver::codec::Codec; +using camera_receiver::codec::DecoderConfig; +using camera_receiver::codec::NvVideoDecoder; + +PYBIND11_MODULE(_camera_receiver_codec, m) +{ + m.doc() = "Native NVDEC H.264/HEVC decoder for camera_receiver."; + + py::enum_(m, "Codec").value("H264", Codec::H264).value("HEVC", Codec::HEVC); + + py::class_(m, "DecoderConfig") + .def(py::init<>()) + .def_readwrite("width", &DecoderConfig::width) + .def_readwrite("height", &DecoderConfig::height) + .def_readwrite("full_range", &DecoderConfig::full_range) + .def_readwrite("gpu_id", &DecoderConfig::gpu_id) + .def_readwrite("codec", &DecoderConfig::codec); + + py::class_(m, "NvVideoDecoder") + .def(py::init(), py::arg("config")) + .def( + "decode", + [](NvVideoDecoder& self, const py::buffer& packet, const py::object& rgba_out, std::uint32_t width, + std::uint32_t height, int rotation) -> bool + { + py::buffer_info pkt = packet.request(); + if (pkt.ndim != 1) + { + throw std::runtime_error("packet must be a 1D byte buffer"); + } + if (pkt.itemsize != 1) + { + throw std::runtime_error("packet must be a uint8 buffer"); + } + // Destination capacity must match the decoder, not caller-supplied dimensions. + if (width != self.width() || height != self.height()) + throw std::runtime_error("width/height do not match the configured decoder size"); + const CudaBuf buf = extract_cuda_buffer(rgba_out, width, height); + bool produced = false; + { + py::gil_scoped_release release; + produced = self.decode(static_cast(pkt.ptr), + static_cast(pkt.size), buf.ptr, buf.row_pitch_bytes, rotation); + } + return produced; + }, + py::arg("packet"), py::arg("rgba_out"), py::arg("width"), py::arg("height"), py::arg("rotation") = 0, + "Feed one Annex-B AU. Returns True iff a frame was written to rgba_out.") + .def( + "decode_nv12", + [](NvVideoDecoder& self, const py::buffer& packet, const py::object& nv12_out, std::uint32_t width, + std::uint32_t height, int rotation) -> bool + { + py::buffer_info pkt = packet.request(); + if (pkt.ndim != 1) + { + throw std::runtime_error("packet must be a 1D byte buffer"); + } + if (pkt.itemsize != 1) + { + throw std::runtime_error("packet must be a uint8 buffer"); + } + // Destination capacity must match the decoder, not caller-supplied dimensions. + if (width != self.width() || height != self.height()) + throw std::runtime_error("width/height do not match the configured decoder size"); + const CudaBuf buf = extract_nv12_buffer(nv12_out, width, height); + bool produced = false; + { + py::gil_scoped_release release; + produced = + self.decode_nv12(static_cast(pkt.ptr), static_cast(pkt.size), + buf.ptr, buf.row_pitch_bytes, rotation); + } + return produced; + }, + py::arg("packet"), py::arg("nv12_out"), py::arg("width"), py::arg("height"), py::arg("rotation") = 0, + "Feed one Annex-B AU. Returns True iff a frame was copied to a tight NV12 GPU buffer.") + .def("reset", &NvVideoDecoder::reset, "Tear down NVDEC state. Use after long stream silence."); +} diff --git a/examples/camera_streamer/camera_receiver/codec/cuda_buffers.hpp b/examples/camera_streamer/camera_receiver/codec/cuda_buffers.hpp new file mode 100644 index 0000000000..936a84bf4a --- /dev/null +++ b/examples/camera_streamer/camera_receiver/codec/cuda_buffers.hpp @@ -0,0 +1,155 @@ +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: Apache-2.0 +#pragma once + +#include + +#include +#include +#include + +namespace py = pybind11; + +namespace +{ + +// Extracts device pointer + row pitch from an object exposing +// __cuda_array_interface__. Expects HxWx4 RGBA8 matching the configured +// (width, height). +struct CudaBuf +{ + std::uintptr_t ptr; + std::size_t row_pitch_bytes; +}; + +inline CudaBuf extract_cuda_buffer(const py::object& obj, std::uint32_t expected_width, std::uint32_t expected_height) +{ + if (!py::hasattr(obj, "__cuda_array_interface__")) + { + throw std::runtime_error("expected object exposing __cuda_array_interface__"); + } + py::dict iface = obj.attr("__cuda_array_interface__").cast(); + if (!iface.contains("shape") || !iface.contains("data")) + { + throw std::runtime_error("__cuda_array_interface__ missing required keys (shape/data)"); + } + + if (!iface.contains("typestr") || iface["typestr"].cast() != "|u1") + throw std::runtime_error("output buffer must have uint8 dtype (|u1)"); + py::tuple output_data = iface["data"].cast(); + if (output_data.size() != 2 || output_data[1].cast()) + throw std::runtime_error("output buffer must be writable"); + if (output_data[0].cast() == 0) + throw std::runtime_error("output buffer pointer is null"); + + py::tuple shape = iface["shape"].cast(); + if (shape.size() != 3) + { + throw std::runtime_error("expected 3D array (H, W, 4); got rank " + std::to_string(shape.size())); + } + const auto h = shape[0].cast(); + const auto w = shape[1].cast(); + const auto c = shape[2].cast(); + if (c != 4) + { + throw std::runtime_error("expected 4-channel RGBA; got " + std::to_string(c)); + } + if (h != expected_height || w != expected_width) + { + throw std::runtime_error("frame size " + std::to_string(w) + "x" + std::to_string(h) + + " does not match configured " + std::to_string(expected_width) + "x" + + std::to_string(expected_height)); + } + + py::tuple data = iface["data"].cast(); + const auto ptr = data[0].cast(); + + // Default to tightly packed RGBA8. + std::size_t row_pitch = static_cast(w) * 4u; + if (iface.contains("strides") && !iface["strides"].is_none()) + { + py::tuple strides = iface["strides"].cast(); + if (strides.size() != 3) + { + throw std::runtime_error("strides rank does not match shape rank"); + } + const auto row_stride = strides[0].cast(); + const auto pixel_stride = strides[1].cast(); + const auto chan_stride = strides[2].cast(); + if (pixel_stride != 4 || chan_stride != 1) + { + throw std::runtime_error("non-contiguous-per-pixel RGBA layout not supported"); + } + if (row_stride < static_cast(w * 4u)) + { + throw std::runtime_error("row stride smaller than W*4"); + } + row_pitch = static_cast(row_stride); + } + + return CudaBuf{ ptr, row_pitch }; +} + +inline CudaBuf extract_nv12_buffer(const py::object& obj, std::uint32_t expected_width, std::uint32_t expected_height) +{ + if (!py::hasattr(obj, "__cuda_array_interface__")) + { + throw std::runtime_error("expected object exposing __cuda_array_interface__"); + } + py::dict iface = obj.attr("__cuda_array_interface__").cast(); + if (!iface.contains("shape") || !iface.contains("data")) + { + throw std::runtime_error("__cuda_array_interface__ missing required keys (shape/data)"); + } + + if (!iface.contains("typestr") || iface["typestr"].cast() != "|u1") + throw std::runtime_error("output buffer must have uint8 dtype (|u1)"); + py::tuple output_data = iface["data"].cast(); + if (output_data.size() != 2 || output_data[1].cast()) + throw std::runtime_error("output buffer must be writable"); + if (output_data[0].cast() == 0) + throw std::runtime_error("output buffer pointer is null"); + + py::tuple shape = iface["shape"].cast(); + if (shape.size() != 2) + { + throw std::runtime_error("expected 2D NV12 array ((H*3)/2, W); got rank " + std::to_string(shape.size())); + } + const auto h = shape[0].cast(); + const auto w = shape[1].cast(); + const auto expected_rows = expected_height * 3u / 2u; + if (h != expected_rows || w != expected_width) + { + throw std::runtime_error("NV12 frame size " + std::to_string(w) + "x" + std::to_string(h) + + " does not match configured " + std::to_string(expected_width) + "x" + + std::to_string(expected_rows)); + } + + py::tuple data = iface["data"].cast(); + const auto ptr = data[0].cast(); + + std::size_t row_pitch = expected_width; + if (iface.contains("strides") && !iface["strides"].is_none()) + { + py::tuple strides = iface["strides"].cast(); + if (strides.size() != 2) + { + throw std::runtime_error("NV12 strides rank does not match shape rank"); + } + const auto row_stride = strides[0].cast(); + const auto col_stride = strides[1].cast(); + if (col_stride != 1) + { + throw std::runtime_error("non-contiguous NV12 columns are not supported"); + } + if (row_stride < static_cast(expected_width)) + { + throw std::runtime_error("NV12 row stride smaller than width"); + } + row_pitch = static_cast(row_stride); + } + + return CudaBuf{ ptr, row_pitch }; +} + +} // namespace diff --git a/examples/camera_streamer/camera_receiver/codec/nv12_to_rgba.cu b/examples/camera_streamer/camera_receiver/codec/nv12_to_rgba.cu new file mode 100644 index 0000000000..09b56ebe79 --- /dev/null +++ b/examples/camera_streamer/camera_receiver/codec/nv12_to_rgba.cu @@ -0,0 +1,155 @@ +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: Apache-2.0 +// +// NV12 → RGBA8 colour conversion in a single kernel launch. +// BT.709 limited-range (H.264 default) or BT.601 full-range (T.871). +// Alpha is always 255. + +#include "nv12_to_rgba.cuh" + +#include +#include +#include + +namespace +{ + +__global__ void nv12_to_rgba_kernel(const unsigned char* __restrict__ y_plane, + const unsigned char* __restrict__ uv_plane, + int y_pitch, + int uv_pitch, + int width, + int height, + unsigned char* __restrict__ rgba_out, + int rgba_row_bytes, + int full_range, + int rotation) +{ + const int x = blockIdx.x * blockDim.x + threadIdx.x; + const int y = blockIdx.y * blockDim.y + threadIdx.y; + if (x >= width || y >= height) + { + return; + } + + const int Y = y_plane[y * y_pitch + x]; + // UV is half-resolution, interleaved Cb,Cr per 2-pixel block. + const int uv_x = (x & ~1); + const int uv_y = y >> 1; + const int Cb = uv_plane[uv_y * uv_pitch + uv_x + 0]; + const int Cr = uv_plane[uv_y * uv_pitch + uv_x + 1]; + + float R, G, B; + if (full_range) + { + // BT.601 full-range (ITU-T T.871). + const float yf = (float)Y; + const float u = (float)Cb - 128.f; + const float v = (float)Cr - 128.f; + R = yf + 1.402f * v; + G = yf - 0.344136f * u - 0.714136f * v; + B = yf + 1.772f * u; + } + else + { + // BT.709 limited-range (luma 16-235). + const float yf = ((float)Y - 16.f) * 1.16438f; + const float u = (float)Cb - 128.f; + const float v = (float)Cr - 128.f; + R = yf + 1.79274f * v; + G = yf - 0.21325f * u - 0.53291f * v; + B = yf + 2.11240f * u; + } + + const int output_x = rotation == 180 ? width - 1 - x : x; + const int output_y = rotation == 180 ? height - 1 - y : y; + const int idx = output_y * rgba_row_bytes + output_x * 4; + rgba_out[idx + 0] = (unsigned char)(R < 0.f ? 0 : (R > 255.f ? 255 : R)); + rgba_out[idx + 1] = (unsigned char)(G < 0.f ? 0 : (G > 255.f ? 255 : G)); + rgba_out[idx + 2] = (unsigned char)(B < 0.f ? 0 : (B > 255.f ? 255 : B)); + rgba_out[idx + 3] = 255; +} + +__global__ void nv12_rotate_180_kernel(const unsigned char* __restrict__ y_plane, + const unsigned char* __restrict__ uv_plane, + int y_pitch, + int uv_pitch, + int width, + int height, + unsigned char* __restrict__ nv12_out, + int output_row_bytes) +{ + const int x = blockIdx.x * blockDim.x + threadIdx.x; + const int y = blockIdx.y * blockDim.y + threadIdx.y; + if (x < width && y < height) + { + nv12_out[(height - 1 - y) * output_row_bytes + width - 1 - x] = y_plane[y * y_pitch + x]; + } + + // One thread per interleaved Cb/Cr pair. Rotate pair locations while + // preserving byte order within each pair. + if (x < width / 2 && y < height / 2) + { + const int source_x = x * 2; + const int destination_x = width - 2 - source_x; + const int destination_y = height / 2 - 1 - y; + unsigned char* destination_uv = nv12_out + output_row_bytes * height; + destination_uv[destination_y * output_row_bytes + destination_x] = uv_plane[y * uv_pitch + source_x]; + destination_uv[destination_y * output_row_bytes + destination_x + 1] = uv_plane[y * uv_pitch + source_x + 1]; + } +} + +} // namespace + +namespace camera_receiver::codec +{ + +void launch_nv12_to_rgba(const uint8_t* y_plane, + const uint8_t* uv_plane, + int y_pitch, + int uv_pitch, + int width, + int height, + uint8_t* rgba_out, + int rgba_row_bytes, + bool full_range, + int rotation, + cudaStream_t stream) +{ + const dim3 block(16, 16, 1); + const dim3 grid((width + 15) / 16, (height + 15) / 16, 1); + nv12_to_rgba_kernel<<>>( + y_plane, uv_plane, y_pitch, uv_pitch, width, height, rgba_out, rgba_row_bytes, full_range ? 1 : 0, rotation); + const cudaError_t launch_err = cudaGetLastError(); + if (launch_err != cudaSuccess) + { + std::ostringstream os; + os << "NV12->RGBA kernel launch failed: " << cudaGetErrorString(launch_err); + throw std::runtime_error(os.str()); + } +} + +void launch_nv12_rotate_180(const uint8_t* y_plane, + const uint8_t* uv_plane, + int y_pitch, + int uv_pitch, + int width, + int height, + uint8_t* nv12_out, + int output_row_bytes, + cudaStream_t stream) +{ + const dim3 block(16, 16, 1); + const dim3 grid((width + 15) / 16, (height + 15) / 16, 1); + nv12_rotate_180_kernel<<>>( + y_plane, uv_plane, y_pitch, uv_pitch, width, height, nv12_out, output_row_bytes); + const cudaError_t launch_err = cudaGetLastError(); + if (launch_err != cudaSuccess) + { + std::ostringstream os; + os << "NV12 rotate-180 kernel launch failed: " << cudaGetErrorString(launch_err); + throw std::runtime_error(os.str()); + } +} + +} // namespace camera_receiver::codec diff --git a/examples/camera_streamer/camera_receiver/codec/nv12_to_rgba.cuh b/examples/camera_streamer/camera_receiver/codec/nv12_to_rgba.cuh new file mode 100644 index 0000000000..52b90f1b63 --- /dev/null +++ b/examples/camera_streamer/camera_receiver/codec/nv12_to_rgba.cuh @@ -0,0 +1,39 @@ +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: Apache-2.0 + +#pragma once + +#include +#include + +namespace camera_receiver::codec +{ + +// Launch the NV12 → RGBA8 kernel. +// full_range = true → BT.601 full-range (T.871) +// = false → BT.709 limited-range (H.264 default) +void launch_nv12_to_rgba(const uint8_t* y_plane, + const uint8_t* uv_plane, + int y_pitch, + int uv_pitch, + int width, + int height, + uint8_t* rgba_out, + int rgba_row_bytes, + bool full_range, + int rotation, + cudaStream_t stream); + +// Rotate pitched NV12 into the caller's output buffer while copying it out of +// NVDEC. UV pairs retain their Cb/Cr byte order. +void launch_nv12_rotate_180(const uint8_t* y_plane, + const uint8_t* uv_plane, + int y_pitch, + int uv_pitch, + int width, + int height, + uint8_t* nv12_out, + int output_row_bytes, + cudaStream_t stream); + +} // namespace camera_receiver::codec diff --git a/examples/camera_streamer/camera_receiver/codec/nv_video_decoder.cpp b/examples/camera_streamer/camera_receiver/codec/nv_video_decoder.cpp new file mode 100644 index 0000000000..dbf8dd4589 --- /dev/null +++ b/examples/camera_streamer/camera_receiver/codec/nv_video_decoder.cpp @@ -0,0 +1,374 @@ +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: Apache-2.0 + +#include "nv_video_decoder.hpp" + +#include "NvDecoder/NvDecoder.h" +#include "nv12_to_rgba.cuh" + +#include +#include +#include +#include +#include + +namespace camera_receiver::codec +{ + +namespace +{ + +inline void check_cu(CUresult result, const char* what) +{ + if (result != CUDA_SUCCESS) + { + const char* err = nullptr; + cuGetErrorString(result, &err); + std::ostringstream os; + os << "NvVideoDecoder: " << what << " failed: " << (err ? err : "unknown"); + throw std::runtime_error(os.str()); + } +} + +inline void check_rt(cudaError_t result, const char* what) +{ + if (result != cudaSuccess) + { + std::ostringstream os; + os << "NvVideoDecoder: " << what << " failed: " << cudaGetErrorString(result); + throw std::runtime_error(os.str()); + } +} + +inline cudaVideoCodec to_cuvid_codec(Codec c) +{ + return (c == Codec::HEVC) ? cudaVideoCodec_HEVC : cudaVideoCodec_H264; +} + +} // namespace + +struct NvVideoDecoder::Impl +{ + DecoderConfig cfg; + + CUdevice cu_device = 0; + CUcontext cu_context = nullptr; + cudaStream_t stream = nullptr; + std::unique_ptr decoder; + std::mutex decoder_mutex; + bool warned_size_mismatch = false; + + Impl(const DecoderConfig& c) : cfg(c) + { + if (cfg.width == 0 || cfg.height == 0) + { + throw std::runtime_error("NvVideoDecoder: width/height must be > 0"); + } + check_cu(cuInit(0), "cuInit"); + check_cu(cuDeviceGet(&cu_device, cfg.gpu_id), "cuDeviceGet"); + check_cu(cuDevicePrimaryCtxRetain(&cu_context, cu_device), "cuDevicePrimaryCtxRetain"); + + check_cu(cuCtxPushCurrent(cu_context), "cuCtxPushCurrent"); + try + { + check_rt(cudaStreamCreateWithFlags(&stream, cudaStreamNonBlocking), "cudaStreamCreate"); + // Zero-latency, decode-order output. Output stays on GPU. + decoder = std::make_unique(cu_context, + true, // bUseDeviceFrame + to_cuvid_codec(cfg.codec), + true, // bLowLatency + false, // bDeviceFramePitched + nullptr, // pCropRect + nullptr, // pResizeDim + false, // bExtractSEIMessage + 0, 0, // max width/height (auto) + 1000, // nClockRate + true); // bForceZeroLatency + } + catch (...) + { + if (stream) + { + cudaStreamDestroy(stream); + stream = nullptr; + } + cuCtxPopCurrent(nullptr); + if (cu_context) + { + cuDevicePrimaryCtxRelease(cu_device); + cu_context = nullptr; + } + throw; + } + check_cu(cuCtxPopCurrent(nullptr), "cuCtxPopCurrent"); + } + + ~Impl() + { + decoder.reset(); + if (stream) + { + cudaStreamDestroy(stream); + stream = nullptr; + } + if (cu_context) + { + cuDevicePrimaryCtxRelease(cu_device); + cu_context = nullptr; + } + } + + bool decode_one( + const uint8_t* packet, std::size_t packet_size, uintptr_t rgba_out, std::size_t rgba_row_bytes, int rotation) + { + std::lock_guard lock(decoder_mutex); + if (!decoder) + { + throw std::runtime_error("NvVideoDecoder: decoder unavailable"); + } + + if (rotation != 0 && rotation != 180) + { + throw std::runtime_error("NvVideoDecoder: rotation must be 0 or 180"); + } + check_cu(cuCtxPushCurrent(cu_context), "cuCtxPushCurrent"); + int n_frames = 0; + try + { + n_frames = decoder->Decode(packet, static_cast(packet_size)); + } + catch (...) + { + cuCtxPopCurrent(nullptr); + throw; + } + if (n_frames == 0) + { + check_cu(cuCtxPopCurrent(nullptr), "cuCtxPopCurrent"); + return false; + } + + uint8_t* nv12 = decoder->GetLockedFrame(); + if (!nv12) + { + check_cu(cuCtxPopCurrent(nullptr), "cuCtxPopCurrent"); + return false; + } + + const int w = decoder->GetWidth(); + const int h = decoder->GetHeight(); + const int pitch = decoder->GetDeviceFramePitch(); + const int luma_size = decoder->GetLumaPlaneSize(); + + if (w != static_cast(cfg.width) || h != static_cast(cfg.height)) + { + // Caller's RGBA buffer is fixed at construction; drop frames + // that don't match. Log once. + if (!warned_size_mismatch) + { + std::ostringstream os; + os << "NvVideoDecoder: stream is " << w << "x" << h << " but configured for " << cfg.width << "x" + << cfg.height << "; dropping frames"; + std::fputs(os.str().c_str(), stderr); + std::fputc('\n', stderr); + warned_size_mismatch = true; + } + decoder->UnlockFrame(&nv12); + check_cu(cuCtxPopCurrent(nullptr), "cuCtxPopCurrent"); + return false; + } + + cudaError_t sync_err = cudaSuccess; + try + { + if (decoder->GetOutputFormat() != cudaVideoSurfaceFormat_NV12) + throw std::runtime_error("NvVideoDecoder: only 8-bit NV12 decoded surfaces are supported"); + launch_nv12_to_rgba(nv12, nv12 + luma_size, pitch, pitch, w, h, reinterpret_cast(rgba_out), + static_cast(rgba_row_bytes), cfg.full_range, rotation, stream); + // UnlockFrame returns the surface to NVDEC's pool, so the + // kernel reading it must finish first. + sync_err = cudaStreamSynchronize(stream); + } + catch (...) + { + // Complete any submitted reads before returning the surface to NVDEC. + cudaStreamSynchronize(stream); + decoder->UnlockFrame(&nv12); + cuCtxPopCurrent(nullptr); + throw; + } + decoder->UnlockFrame(&nv12); + check_cu(cuCtxPopCurrent(nullptr), "cuCtxPopCurrent"); + if (sync_err != cudaSuccess) + { + std::ostringstream os; + os << "NvVideoDecoder: NV12→RGBA kernel sync failed: " << cudaGetErrorString(sync_err); + throw std::runtime_error(os.str()); + } + return true; + } + + bool decode_one_nv12( + const uint8_t* packet, std::size_t packet_size, uintptr_t nv12_out, std::size_t nv12_row_bytes, int rotation) + { + std::lock_guard lock(decoder_mutex); + if (!decoder) + { + throw std::runtime_error("NvVideoDecoder: decoder unavailable"); + } + if (nv12_row_bytes < cfg.width) + { + throw std::runtime_error("NvVideoDecoder: NV12 row stride smaller than width"); + } + if (rotation != 0 && rotation != 180) + { + throw std::runtime_error("NvVideoDecoder: rotation must be 0 or 180"); + } + + check_cu(cuCtxPushCurrent(cu_context), "cuCtxPushCurrent"); + int n_frames = 0; + try + { + n_frames = decoder->Decode(packet, static_cast(packet_size)); + } + catch (...) + { + cuCtxPopCurrent(nullptr); + throw; + } + if (n_frames == 0) + { + check_cu(cuCtxPopCurrent(nullptr), "cuCtxPopCurrent"); + return false; + } + + uint8_t* nv12 = decoder->GetLockedFrame(); + if (!nv12) + { + check_cu(cuCtxPopCurrent(nullptr), "cuCtxPopCurrent"); + return false; + } + + const int w = decoder->GetWidth(); + const int h = decoder->GetHeight(); + const int src_pitch = decoder->GetDeviceFramePitch(); + const int luma_size = decoder->GetLumaPlaneSize(); + + if (w != static_cast(cfg.width) || h != static_cast(cfg.height)) + { + if (!warned_size_mismatch) + { + std::ostringstream os; + os << "NvVideoDecoder: stream is " << w << "x" << h << " but configured for " << cfg.width << "x" + << cfg.height << "; dropping frames"; + std::fputs(os.str().c_str(), stderr); + std::fputc('\n', stderr); + warned_size_mismatch = true; + } + decoder->UnlockFrame(&nv12); + check_cu(cuCtxPopCurrent(nullptr), "cuCtxPopCurrent"); + return false; + } + + cudaError_t sync_err = cudaSuccess; + try + { + if (decoder->GetOutputFormat() != cudaVideoSurfaceFormat_NV12) + throw std::runtime_error("NvVideoDecoder: only 8-bit NV12 decoded surfaces are supported"); + auto* dst_y = reinterpret_cast(nv12_out); + auto* dst_uv = dst_y + nv12_row_bytes * static_cast(h); + if (rotation == 180) + { + launch_nv12_rotate_180(nv12, nv12 + luma_size, src_pitch, src_pitch, w, h, dst_y, + static_cast(nv12_row_bytes), stream); + } + else + { + check_rt( + cudaMemcpy2DAsync(dst_y, nv12_row_bytes, nv12, src_pitch, w, h, cudaMemcpyDeviceToDevice, stream), + "cudaMemcpy2DAsync NV12 Y"); + check_rt(cudaMemcpy2DAsync(dst_uv, nv12_row_bytes, nv12 + luma_size, src_pitch, w, h / 2, + cudaMemcpyDeviceToDevice, stream), + "cudaMemcpy2DAsync NV12 UV"); + } + sync_err = cudaStreamSynchronize(stream); + } + catch (...) + { + // The Y copy may already be queued when the UV copy fails. + cudaStreamSynchronize(stream); + decoder->UnlockFrame(&nv12); + cuCtxPopCurrent(nullptr); + throw; + } + decoder->UnlockFrame(&nv12); + check_cu(cuCtxPopCurrent(nullptr), "cuCtxPopCurrent"); + if (sync_err != cudaSuccess) + { + std::ostringstream os; + os << "NvVideoDecoder: NV12 copy sync failed: " << cudaGetErrorString(sync_err); + throw std::runtime_error(os.str()); + } + return true; + } + + void reset_decoder() + { + std::lock_guard lock(decoder_mutex); + check_cu(cuCtxPushCurrent(cu_context), "cuCtxPushCurrent"); + decoder.reset(); + try + { + decoder = std::make_unique( + cu_context, true, to_cuvid_codec(cfg.codec), true, false, nullptr, nullptr, false, 0, 0, 1000, true); + } + catch (...) + { + cuCtxPopCurrent(nullptr); + throw; + } + check_cu(cuCtxPopCurrent(nullptr), "cuCtxPopCurrent"); + warned_size_mismatch = false; + } +}; + +NvVideoDecoder::NvVideoDecoder(const DecoderConfig& cfg) : impl_(std::make_unique(cfg)) +{ +} + +NvVideoDecoder::~NvVideoDecoder() = default; + +uint32_t NvVideoDecoder::width() const +{ + return impl_->cfg.width; +} + +uint32_t NvVideoDecoder::height() const +{ + return impl_->cfg.height; +} + +bool NvVideoDecoder::decode(const uint8_t* packet, + std::size_t packet_size, + uintptr_t rgba_out_device_ptr, + std::size_t row_pitch_bytes, + int rotation) +{ + return impl_->decode_one(packet, packet_size, rgba_out_device_ptr, row_pitch_bytes, rotation); +} + +bool NvVideoDecoder::decode_nv12(const uint8_t* packet, + std::size_t packet_size, + uintptr_t nv12_out_device_ptr, + std::size_t row_pitch_bytes, + int rotation) +{ + return impl_->decode_one_nv12(packet, packet_size, nv12_out_device_ptr, row_pitch_bytes, rotation); +} + +void NvVideoDecoder::reset() +{ + impl_->reset_decoder(); +} + +} // namespace camera_receiver::codec diff --git a/examples/camera_streamer/camera_receiver/codec/nv_video_decoder.hpp b/examples/camera_streamer/camera_receiver/codec/nv_video_decoder.hpp new file mode 100644 index 0000000000..a0f3a06c3a --- /dev/null +++ b/examples/camera_streamer/camera_receiver/codec/nv_video_decoder.hpp @@ -0,0 +1,80 @@ +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: Apache-2.0 +// +// NVDEC H.264/HEVC video decoder. It can either copy decoded NV12 into a caller-owned +// tight NV12 buffer or convert into caller-owned RGBA8. Zero-latency +// (no reorder buffer, decode-order output). RGBA conversion supports BT.709 +// limited and BT.601 full range. + +#pragma once + +#include +#include + +namespace camera_receiver::codec +{ + +enum class Codec +{ + H264 = 0, + HEVC = 1, +}; + +struct DecoderConfig +{ + uint32_t width = 0; + uint32_t height = 0; + // BT.601 full-range (ITU-T T.871) when true; BT.709 limited (16-235) when false. + bool full_range = false; + int gpu_id = 0; + // Elementary-stream codec. Must match the sender. + Codec codec = Codec::H264; +}; + +class NvVideoDecoder +{ +public: + explicit NvVideoDecoder(const DecoderConfig& cfg); + ~NvVideoDecoder(); + + NvVideoDecoder(const NvVideoDecoder&) = delete; + NvVideoDecoder& operator=(const NvVideoDecoder&) = delete; + NvVideoDecoder(NvVideoDecoder&&) = delete; + NvVideoDecoder& operator=(NvVideoDecoder&&) = delete; + + uint32_t width() const; + uint32_t height() const; + + // Feed one Annex-B access unit. Returns true iff a frame was + // produced and converted into the caller's RGBA8 buffer. + // ``rgba_out_device_ptr`` points to the top-left of a HxWx4 GPU + // buffer; ``row_pitch_bytes`` is the byte stride between rows. + // Synchronous on an internal stream — buffer is safe to read from + // any stream once decode() returns. + bool decode(const uint8_t* packet, + std::size_t packet_size, + uintptr_t rgba_out_device_ptr, + std::size_t row_pitch_bytes, + int rotation = 0); + + // Feed one Annex-B access unit. Returns true iff a frame was produced and + // copied into a caller-owned tight-ish NV12 buffer shaped + // ((height * 3) / 2, width). ``row_pitch_bytes`` is the byte stride between + // output rows; the Y plane starts at ``nv12_out_device_ptr`` and the UV + // plane starts at row ``height``. + bool decode_nv12(const uint8_t* packet, + std::size_t packet_size, + uintptr_t nv12_out_device_ptr, + std::size_t row_pitch_bytes, + int rotation = 0); + + // Tear down NVDEC state. Use on stream-timeout / disconnect so the + // next packet sees a fresh DPB. + void reset(); + +private: + struct Impl; + std::unique_ptr impl_; +}; + +} // namespace camera_receiver::codec diff --git a/examples/camera_streamer/camera_receiver/nv_video_decode.py b/examples/camera_streamer/camera_receiver/nv_video_decode.py new file mode 100644 index 0000000000..60b9627aa7 --- /dev/null +++ b/examples/camera_streamer/camera_receiver/nv_video_decode.py @@ -0,0 +1,87 @@ +# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# SPDX-License-Identifier: Apache-2.0 +"""NVDEC video decoder wrapper for camera_receiver.""" + +from __future__ import annotations + +import logging + +logger = logging.getLogger("isaaccapture.camera_streamer." + __name__) + + +class NvVideoDecoder: + """Annex-B H.264 / HEVC access-unit decoder to GPU buffers.""" + + def __init__( + self, + width: int, + height: int, + full_range: bool = False, + gpu_id: int = 0, + low_latency: bool = True, + codec: str = "h264", + ) -> None: + del low_latency + + codec_norm = str(codec).lower() + if codec_norm in ("h265", "hevc"): + codec_norm = "hevc" + elif codec_norm in ("h264", "avc"): + codec_norm = "h264" + else: + raise ValueError(f"NvVideoDecoder: codec must be h264|hevc, got {codec!r}") + + self._width = width + self._height = height + self._full_range = full_range + self._gpu_id = gpu_id + self._codec = codec_norm + self._decoder = None + + def _ensure_initialized(self) -> None: + if self._decoder is not None: + return + try: + import codec + except ImportError as e: + raise RuntimeError( + "NvVideoDecoder requires the native codec. Follow " + "the host receiver build instructions in examples/camera_streamer/README.md." + ) from e + + cfg = codec.DecoderConfig() + cfg.width = self._width + cfg.height = self._height + cfg.full_range = self._full_range + cfg.gpu_id = self._gpu_id + try: + cfg.codec = codec.Codec.HEVC if self._codec == "hevc" else codec.Codec.H264 + except AttributeError as e: + raise RuntimeError( + "The native codec module is stale (no Codec enum). Rebuild it " + "using the host receiver build instructions in examples/camera_streamer/README.md." + ) from e + self._decoder = codec.NvVideoDecoder(cfg) + + def decode(self, packet: bytes, rgba_out, rotation: int = 0) -> bool: + """Feed one Annex-B AU. Returns True iff a frame was written.""" + self._ensure_initialized() + return self._decoder.decode( + packet, rgba_out, self._width, self._height, rotation + ) + + def decode_nv12(self, packet: bytes, nv12_out, rotation: int = 0) -> bool: + """Feed one Annex-B AU into a tight NV12 GPU buffer. + + ``nv12_out`` is expected to be a ``uint8`` array shaped + ``((height * 3) // 2, width)``. Returns True iff a frame was written. + """ + self._ensure_initialized() + return self._decoder.decode_nv12( + packet, nv12_out, self._width, self._height, rotation + ) + + def reset(self) -> None: + """Tear down NVDEC state. Use after stream timeout or disconnect.""" + if self._decoder is not None: + self._decoder.reset() diff --git a/examples/camera_streamer/tests/pyproject.toml b/examples/camera_streamer/tests/pyproject.toml new file mode 100644 index 0000000000..f45714fa26 --- /dev/null +++ b/examples/camera_streamer/tests/pyproject.toml @@ -0,0 +1,13 @@ +# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# SPDX-License-Identifier: Apache-2.0 + +[project] +name = "camera-streamer-tests" +version = "0.0.0" +requires-python = ">=3.10" + +[project.optional-dependencies] +dev = ["pytest>=8", "pybind11>=2.13", "numpy>=1.24,<3", "PyYAML>=6"] + +[tool.pytest.ini_options] +testpaths = ["python"] diff --git a/examples/camera_streamer/tests/python/conftest.py b/examples/camera_streamer/tests/python/conftest.py new file mode 100644 index 0000000000..f2adb71797 --- /dev/null +++ b/examples/camera_streamer/tests/python/conftest.py @@ -0,0 +1,7 @@ +# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# SPDX-License-Identifier: Apache-2.0 +import sys +from pathlib import Path + +# The receiver is also supported as a standalone script directory. +sys.path.insert(0, str(Path(__file__).resolve().parents[2] / "camera_receiver")) diff --git a/examples/camera_streamer/tests/python/test_cuda_buffers.py b/examples/camera_streamer/tests/python/test_cuda_buffers.py new file mode 100644 index 0000000000..015415b3b9 --- /dev/null +++ b/examples/camera_streamer/tests/python/test_cuda_buffers.py @@ -0,0 +1,77 @@ +# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# SPDX-License-Identifier: Apache-2.0 +"""Exercise the native buffer checks without a CUDA device or toolkit.""" + +import importlib.util +import subprocess +import sysconfig +from pathlib import Path +from types import SimpleNamespace + +import pybind11 +import pytest + + +@pytest.fixture(scope="module") +def validators(tmp_path_factory): + path = tmp_path_factory.mktemp("cuda_buffers") + source = path / "bindings.cpp" + source.write_text("""#include +PYBIND11_MODULE(_test_cuda_buffers, m) { + m.def("nv12", [](py::object o) { return extract_nv12_buffer(o, 4, 2).row_pitch_bytes; }); + m.def("rgba", [](py::object o) { return extract_cuda_buffer(o, 4, 2).row_pitch_bytes; }); +} +""") + extension = path / ("_test_cuda_buffers" + sysconfig.get_config_var("EXT_SUFFIX")) + codec = Path(__file__).resolve().parents[2] / "camera_receiver" / "codec" + subprocess.run( + [ + "c++", + "-shared", + "-fPIC", + "-std=c++17", + str(source), + "-o", + str(extension), + "-I" + str(codec), + "-I" + pybind11.get_include(), + "-I" + sysconfig.get_paths()["include"], + ], + check=True, + ) + spec = importlib.util.spec_from_file_location("_test_cuda_buffers", extension) + module = importlib.util.module_from_spec(spec) + spec.loader.exec_module(module) + return module + + +@pytest.mark.parametrize( + "kind,shape,pitch", [("nv12", (3, 4), 4), ("rgba", (2, 4, 4), 16)] +) +@pytest.mark.parametrize( + "typestr,readonly,error", + [ + ("|u1", False, None), + (" +namespace camera_receiver::codec { +struct NvVideoDecoder::Impl { DecoderConfig cfg; }; +NvVideoDecoder::NvVideoDecoder(const DecoderConfig& cfg) : impl_(new Impl{cfg}) {} +NvVideoDecoder::~NvVideoDecoder() = default; +uint32_t NvVideoDecoder::width() const { return impl_->cfg.width; } +uint32_t NvVideoDecoder::height() const { return impl_->cfg.height; } +bool NvVideoDecoder::decode(const uint8_t*, std::size_t, uintptr_t, std::size_t, int) { return true; } +bool NvVideoDecoder::decode_nv12(const uint8_t*, std::size_t, uintptr_t, std::size_t, int) { return true; } +void NvVideoDecoder::reset() {} +} +""") + codec = Path(__file__).resolve().parents[2] / "camera_receiver" / "codec" + extension = path / ( + "_camera_receiver_codec" + sysconfig.get_config_var("EXT_SUFFIX") + ) + subprocess.run( + [ + "c++", + "-shared", + "-fPIC", + "-std=c++17", + str(codec / "codec_bindings.cpp"), + str(stub), + "-o", + str(extension), + "-I" + str(codec), + "-I" + pybind11.get_include(), + "-I" + sysconfig.get_paths()["include"], + ], + check=True, + ) + spec = importlib.util.spec_from_file_location("_camera_receiver_codec", extension) + module = importlib.util.module_from_spec(spec) + spec.loader.exec_module(module) + return module + + +@pytest.mark.parametrize("method", ["decode", "decode_nv12"]) +@pytest.mark.parametrize( + "width,height", [(1280, 720), (640, 720), (1280, 360), (640, 360), (1920, 1080)] +) +def test_binding_checks_dimensions_against_decoder(bindings, method, width, height): + config = bindings.DecoderConfig() + config.width, config.height = 1280, 720 + decoder = bindings.NvVideoDecoder(config) + shape = (height, width, 4) if method == "decode" else (height * 3 // 2, width) + output = SimpleNamespace( + __cuda_array_interface__={ + "shape": shape, + "data": (1234, False), + "typestr": "|u1", + "strides": None, + } + ) + call = getattr(decoder, method) + if (width, height) == (config.width, config.height): + assert call(b"au", output, width, height) + else: + with pytest.raises(RuntimeError, match="configured decoder size"): + call(b"au", output, width, height) diff --git a/examples/camera_streamer/tests/python/test_nv_video_decode.py b/examples/camera_streamer/tests/python/test_nv_video_decode.py new file mode 100644 index 0000000000..4494cd7cce --- /dev/null +++ b/examples/camera_streamer/tests/python/test_nv_video_decode.py @@ -0,0 +1,46 @@ +# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. +# SPDX-License-Identifier: Apache-2.0 + +import unittest + +from nv_video_decode import NvVideoDecoder + + +class _FakeNativeDecoder: + def __init__(self) -> None: + self.calls = [] + + def decode(self, *args): + self.calls.append(("rgba", args)) + return True + + def decode_nv12(self, *args): + self.calls.append(("nv12", args)) + return True + + +class NvVideoDecoderRotationTest(unittest.TestCase): + def setUp(self) -> None: + self.decoder = NvVideoDecoder(2560, 1984, codec="hevc") + self.native = _FakeNativeDecoder() + self.decoder._decoder = self.native + + def test_rgba_rotation_is_forwarded(self) -> None: + output = object() + self.assertTrue(self.decoder.decode(b"au", output, rotation=180)) + self.assertEqual( + self.native.calls, + [("rgba", (b"au", output, 2560, 1984, 180))], + ) + + def test_nv12_rotation_is_forwarded(self) -> None: + output = object() + self.assertTrue(self.decoder.decode_nv12(b"au", output, rotation=180)) + self.assertEqual( + self.native.calls, + [("nv12", (b"au", output, 2560, 1984, 180))], + ) + + +if __name__ == "__main__": + unittest.main()