diff --git a/src/torchcodec/_core/BetaCudaDeviceInterface.cpp b/src/torchcodec/_core/BetaCudaDeviceInterface.cpp index 804c3ba78..38298a02b 100644 --- a/src/torchcodec/_core/BetaCudaDeviceInterface.cpp +++ b/src/torchcodec/_core/BetaCudaDeviceInterface.cpp @@ -92,20 +92,21 @@ pfnDisplayPictureCallback(void* pUserData, CUVIDPARSERDISPINFO* dispInfo) { return decoder->frameReadyInDisplayOrder(dispInfo); } -static UniqueCUvideodecoder createDecoder(CUVIDEOFORMAT* videoFormat) { +static UniqueCUvideodecoder createDecoder( + CUVIDEOFORMAT* videoFormat, + bool useP016) { // Decoder creation parameters, most are taken from DALI CUVIDDECODECREATEINFO decoderParams = {}; decoderParams.bitDepthMinus8 = videoFormat->bit_depth_luma_minus8; decoderParams.ChromaFormat = videoFormat->chroma_format; - // We explicitly request NV12 format, which means 10bit videos will be - // automatically converted to 8bits by NVDEC itself. That is, the raw frames - // we get back from cuvidMapVideoFrame will already be in 8bit format. We - // won't need to do the conversion ourselves, so that's a lot easier. - // In the ffmpeg CUDA interface, we have to do the 10 -> 8bits conversion - // ourselves later in convertAVFrameToFrameOutput(), because FFmpeg explicitly - // requests 10 or 16bits output formats for >8-bit videos! - // https://github.com/FFmpeg/FFmpeg/blob/e05f8acabff468c1382277c1f31fa8e9d90c3202/libavcodec/nvdec.c#L376-L403 - decoderParams.OutputFormat = cudaVideoSurfaceFormat_NV12; + // For >8-bit videos when the user requests high bit depth output (float32 or + // auto), we use P016 to preserve the full bit depth. Otherwise we use NV12, + // which lets NVDEC handle the 10->8 bit conversion internally. + if (useP016) { + decoderParams.OutputFormat = cudaVideoSurfaceFormat_P016; + } else { + decoderParams.OutputFormat = cudaVideoSurfaceFormat_NV12; + } decoderParams.ulCreationFlags = cudaVideoCreate_Default; decoderParams.CodecType = videoFormat->codec; decoderParams.ulHeight = videoFormat->coded_height; @@ -194,7 +195,8 @@ std::optional validateCodecSupport(AVCodecID codecId) { bool nativeNVDECSupport( const StableDevice& device, - const SharedAVCodecContext& codecContext) { + const SharedAVCodecContext& codecContext, + OutputDtype outputDtype) { // Return true iff the input video stream is supported by our NVDEC // implementation. @@ -241,15 +243,23 @@ bool nativeNVDECSupport( return false; } - // We'll set the decoderParams.OutputFormat to NV12, so we need to make - // sure it's actually supported. - // TODO: If this fail, we could consider decoding to something else than NV12 - // (like cudaVideoSurfaceFormat_P016) instead of falling back to CPU. This is - // what FFmpeg does. - bool supportsNV12Output = - (caps.nOutputFormatMask >> cudaVideoSurfaceFormat_NV12) & 1; - if (!supportsNV12Output) { - return false; + // Check that the hardware supports the output format we need. + // For >8-bit content with FLOAT32/AUTO output, we need P016 to preserve + // full bit depth. Otherwise NV12 suffices (NVDEC converts 10->8 bit + // internally). + bool needsP016 = (bitDepthMinus8 > 0) && (outputDtype != OutputDtype::UINT8); + if (needsP016) { + bool supportsP016Output = + (caps.nOutputFormatMask >> cudaVideoSurfaceFormat_P016) & 1; + if (!supportsP016Output) { + return false; + } + } else { + bool supportsNV12Output = + (caps.nOutputFormatMask >> cudaVideoSurfaceFormat_NV12) & 1; + if (!supportsNV12Output) { + return false; + } } return true; @@ -285,7 +295,7 @@ BetaCudaDeviceInterface::~BetaCudaDeviceInterface() { flush(); unmapPreviousFrame(); NVDECCache::getCache(device_).returnDecoder( - &videoFormat_, std::move(decoder_)); + &videoFormat_, surfaceFormat_, std::move(decoder_)); } if (videoParser_) { @@ -296,13 +306,32 @@ BetaCudaDeviceInterface::~BetaCudaDeviceInterface() { returnNppStreamContextToCache(device_, std::move(nppCtx_)); } +void BetaCudaDeviceInterface::initializeVideo( + const VideoStreamOptions& videoStreamOptions, + const std::vector>& transforms, + const std::optional& resizedOutputDims) { + if (cpuFallback_) { + cpuFallback_->initializeVideo( + videoStreamOptions, transforms, resizedOutputDims); + } +} + void BetaCudaDeviceInterface::initialize( const AVStream* avStream, const UniqueDecodingAVFormatContext& avFormatCtx, - [[maybe_unused]] const SharedAVCodecContext& codecContext) { + const SharedAVCodecContext& codecContext, + OutputDtype outputDtype) { STD_TORCH_CHECK(avStream != nullptr, "AVStream cannot be null"); + outputDtype_ = outputDtype; + + // Set bitDepth_ from the codec context so it's correct on both paths. On + // the native NVDEC path, streamPropertyChange used to also set this from the + // parser; we now rely on this initialization for both paths. + bitDepth_ = getBitDepthFromAVPixelFormat(codecContext->pix_fmt); + rotation_ = rotationFromDegrees(getRotationFromStream(avStream)); - if (!nvcuvidAvailable_ || !nativeNVDECSupport(device_, codecContext)) { + if (!nvcuvidAvailable_ || + !nativeNVDECSupport(device_, codecContext, outputDtype)) { cpuFallback_ = createDeviceInterface(kStableCPU); STD_TORCH_CHECK( cpuFallback_ != nullptr, "Failed to create CPU device interface"); @@ -441,6 +470,7 @@ int BetaCudaDeviceInterface::streamPropertyChange(CUVIDEOFORMAT* videoFormat) { STD_TORCH_CHECK(videoFormat != nullptr, "Invalid video format"); videoFormat_ = *videoFormat; + // bitDepth_ is already set in initialize() from codecContext->pix_fmt. if (videoFormat_.min_num_decode_surfaces == 0) { // Same as DALI's fallback @@ -448,13 +478,17 @@ int BetaCudaDeviceInterface::streamPropertyChange(CUVIDEOFORMAT* videoFormat) { } if (!decoder_) { - decoder_ = NVDECCache::getCache(device_).getDecoder(videoFormat); + bool useP016 = resolvedBitDepth(bitDepth_, outputDtype_) > 8; + surfaceFormat_ = + useP016 ? cudaVideoSurfaceFormat_P016 : cudaVideoSurfaceFormat_NV12; + decoder_ = + NVDECCache::getCache(device_).getDecoder(videoFormat, surfaceFormat_); if (!decoder_) { // TODONVDEC P2: consider re-configuring an existing decoder instead of // re-creating one. See docs, see DALI. Re-configuration doesn't seem to // be enabled in DALI by default. - decoder_ = createDecoder(videoFormat); + decoder_ = createDecoder(videoFormat, useP016); } STD_TORCH_CHECK(decoder_, "Failed to get or create decoder"); @@ -868,38 +902,172 @@ UniqueAVFrame BetaCudaDeviceInterface::transferCpuFrameToGpuNV12( return gpuFrame; } +UniqueAVFrame BetaCudaDeviceInterface::transferCpuFrameToGpuP016( + UniqueAVFrame& cpuFrame) { + // Similar to transferCpuFrameToGpuNV12, but for 10-bit content. + // Converts CPU frame to P016 (16-bit YUV 4:2:0) and uploads to GPU. + STD_TORCH_CHECK(cpuFrame != nullptr, "CPU frame cannot be null"); + + int width = cpuFrame->width; + int height = cpuFrame->height; + + // Intermediate P016 CPU frame + UniqueAVFrame p016CpuFrame(av_frame_alloc()); + STD_TORCH_CHECK(p016CpuFrame != nullptr, "Failed to allocate P016 CPU frame"); + + p016CpuFrame->format = AV_PIX_FMT_P016LE; + p016CpuFrame->width = width; + p016CpuFrame->height = height; + + int ret = av_frame_get_buffer(p016CpuFrame.get(), 0); + STD_TORCH_CHECK( + ret >= 0, + "Failed to allocate P016 CPU frame buffer: ", + getFFMPEGErrorStringFromErrorCode(ret)); + + SwsConfig swsConfig( + width, + height, + static_cast(cpuFrame->format), + cpuFrame->colorspace, + width, + height); + + if (!swsContext_ || prevSwsConfig_ != swsConfig) { + swsContext_ = createSwsContext(swsConfig, AV_PIX_FMT_P016LE, SWS_BILINEAR); + prevSwsConfig_ = swsConfig; + } + + int convertedHeight = sws_scale( + swsContext_.get(), + cpuFrame->data, + cpuFrame->linesize, + 0, + height, + p016CpuFrame->data, + p016CpuFrame->linesize); + STD_TORCH_CHECK( + convertedHeight == height, "sws_scale failed for CPU->P016 conversion"); + + // P016: Y plane is width * height * 2 bytes (16-bit per sample) + // UV plane is width * (height/2) * 2 bytes (interleaved 16-bit U/V, half + // height) + int ySize = width * height * 2; + STD_TORCH_CHECK( + height % 2 == 0, + "height must be even for P016. Please report on TorchCodec repo."); + int uvSize = width * (height / 2) * 2; + size_t totalSize = static_cast(ySize + uvSize); + + uint8_t* cudaBuffer = nullptr; + cudaError_t err = + cudaMalloc(reinterpret_cast(&cudaBuffer), totalSize); + STD_TORCH_CHECK( + err == cudaSuccess, + "Failed to allocate CUDA memory: ", + cudaGetErrorString(err)); + + UniqueAVFrame gpuFrame(av_frame_alloc()); + STD_TORCH_CHECK(gpuFrame != nullptr, "Failed to allocate GPU AVFrame"); + + gpuFrame->format = AV_PIX_FMT_CUDA; + gpuFrame->width = width; + gpuFrame->height = height; + gpuFrame->data[0] = cudaBuffer; + gpuFrame->data[1] = cudaBuffer + ySize; + // For P016, linesize is width * 2 (2 bytes per sample) + gpuFrame->linesize[0] = width * 2; + gpuFrame->linesize[1] = width * 2; + + // Copy Y plane + err = cudaMemcpy2D( + gpuFrame->data[0], + gpuFrame->linesize[0], + p016CpuFrame->data[0], + p016CpuFrame->linesize[0], + width * 2, // width in bytes + height, + cudaMemcpyHostToDevice); + STD_TORCH_CHECK( + err == cudaSuccess, + "Failed to copy Y plane to GPU: ", + cudaGetErrorString(err)); + + // Copy UV plane + err = cudaMemcpy2D( + gpuFrame->data[1], + gpuFrame->linesize[1], + p016CpuFrame->data[1], + p016CpuFrame->linesize[1], + width * 2, // width in bytes (interleaved U/V, same width as Y) + height / 2, + cudaMemcpyHostToDevice); + STD_TORCH_CHECK( + err == cudaSuccess, + "Failed to copy UV plane to GPU: ", + cudaGetErrorString(err)); + + ret = av_frame_copy_props(gpuFrame.get(), cpuFrame.get()); + STD_TORCH_CHECK( + ret >= 0, + "Failed to copy frame properties: ", + getFFMPEGErrorStringFromErrorCode(ret)); + + gpuFrame->opaque_ref = + av_buffer_create(nullptr, 0, cudaBufferFreeCallback, cudaBuffer, 0); + STD_TORCH_CHECK( + gpuFrame->opaque_ref != nullptr, + "Failed to create GPU memory cleanup reference"); + + return gpuFrame; +} + void BetaCudaDeviceInterface::convertAVFrameToFrameOutput( UniqueAVFrame& avFrame, FrameOutput& frameOutput, std::optional preAllocatedOutputTensor) { - UniqueAVFrame gpuFrame = - cpuFallback_ ? transferCpuFrameToGpuNV12(avFrame) : std::move(avFrame); + bool is10Bit = resolvedBitDepth(bitDepth_, outputDtype_) > 8; + + UniqueAVFrame gpuFrame; + if (cpuFallback_) { + gpuFrame = is10Bit ? transferCpuFrameToGpuP016(avFrame) + : transferCpuFrameToGpuNV12(avFrame); + } else { + gpuFrame = std::move(avFrame); + } - // TODONVDEC P2: we may need to handle 10bit videos the same way the CUDA - // ffmpeg interface does it with maybeConvertAVFrameToNV12OrRGB24(). STD_TORCH_CHECK( gpuFrame->format == AV_PIX_FMT_CUDA, "Expected CUDA format frame from BETA CUDA interface"); cudaStream_t nvdecStream = getCurrentCudaStream(device_.index()); - if (rotation_ == Rotation::NONE) { - validatePreAllocatedTensorShape(preAllocatedOutputTensor, gpuFrame); - frameOutput.data = convertNV12FrameToRGB( - gpuFrame, device_, nppCtx_, nvdecStream, preAllocatedOutputTensor); + if (is10Bit) { + // P016 path: convert to uint16 RGB using NPP 16-bit ColorTwist. + if (rotation_ == Rotation::NONE) { + validatePreAllocatedTensorShape(preAllocatedOutputTensor, gpuFrame); + frameOutput.data = convertP016FrameToRGB( + gpuFrame, device_, nppCtx_, nvdecStream, preAllocatedOutputTensor); + } else { + frameOutput.data = + convertP016FrameToRGB(gpuFrame, device_, nppCtx_, nvdecStream); + applyRotation(frameOutput, preAllocatedOutputTensor); + } } else { - // preAllocatedOutputTensor has post-rotation dimensions, but NV12->RGB - // conversion outputs pre-rotation dimensions, so we can't use it as the - // conversion destination or validate it against the frame shape. - // Once we support native transforms on the beta CUDA interface, rotation - // should be handled as part of the transform pipeline instead. - frameOutput.data = convertNV12FrameToRGB( - gpuFrame, - device_, - nppCtx_, - nvdecStream, - /*preAllocatedOutputTensor=*/std::nullopt); - applyRotation(frameOutput, preAllocatedOutputTensor); + // NV12 path: existing 8-bit conversion using NPP. + if (rotation_ == Rotation::NONE) { + validatePreAllocatedTensorShape(preAllocatedOutputTensor, gpuFrame); + frameOutput.data = convertNV12FrameToRGB( + gpuFrame, device_, nppCtx_, nvdecStream, preAllocatedOutputTensor); + } else { + frameOutput.data = convertNV12FrameToRGB( + gpuFrame, + device_, + nppCtx_, + nvdecStream, + /*preAllocatedOutputTensor=*/std::nullopt); + applyRotation(frameOutput, preAllocatedOutputTensor); + } } } diff --git a/src/torchcodec/_core/BetaCudaDeviceInterface.h b/src/torchcodec/_core/BetaCudaDeviceInterface.h index 8f44dbda1..826e87e31 100644 --- a/src/torchcodec/_core/BetaCudaDeviceInterface.h +++ b/src/torchcodec/_core/BetaCudaDeviceInterface.h @@ -41,7 +41,8 @@ class BetaCudaDeviceInterface : public DeviceInterface { void initialize( const AVStream* avStream, const UniqueDecodingAVFormatContext& avFormatCtx, - const SharedAVCodecContext& codecContext) override; + const SharedAVCodecContext& codecContext, + OutputDtype outputDtype = OutputDtype::UINT8) override; void convertAVFrameToFrameOutput( UniqueAVFrame& avFrame, @@ -58,6 +59,11 @@ class BetaCudaDeviceInterface : public DeviceInterface { int frameReadyForDecoding(CUVIDPICPARAMS* picParams); int frameReadyInDisplayOrder(CUVIDPARSERDISPINFO* dispInfo); + void initializeVideo( + const VideoStreamOptions& videoStreamOptions, + const std::vector>& transforms, + const std::optional& resizedOutputDims) override; + std::string getDetails() override; private: @@ -81,6 +87,7 @@ class BetaCudaDeviceInterface : public DeviceInterface { const CUVIDPARSERDISPINFO& dispInfo); UniqueAVFrame transferCpuFrameToGpuNV12(UniqueAVFrame& cpuFrame); + UniqueAVFrame transferCpuFrameToGpuP016(UniqueAVFrame& cpuFrame); void applyRotation( FrameOutput& frameOutput, @@ -108,6 +115,13 @@ class BetaCudaDeviceInterface : public DeviceInterface { SwsConfig prevSwsConfig_; Rotation rotation_ = Rotation::NONE; + + // Bit depth of the source video. 8 for standard, 10+ for HDR. + int bitDepth_ = 8; + // User-requested output dtype. + OutputDtype outputDtype_ = OutputDtype::UINT8; + // NVDEC output surface format (NV12 for 8-bit, P016 for >8-bit HDR output). + cudaVideoSurfaceFormat surfaceFormat_ = cudaVideoSurfaceFormat_NV12; }; } // namespace facebook::torchcodec diff --git a/src/torchcodec/_core/CUDACommon.cpp b/src/torchcodec/_core/CUDACommon.cpp index 2dd11b171..51bf5a9b3 100644 --- a/src/torchcodec/_core/CUDACommon.cpp +++ b/src/torchcodec/_core/CUDACommon.cpp @@ -282,6 +282,82 @@ const Npp32f bt2020LimitedRangeColorTwist[3][4] = { {1.16438356f, 2.14177232f, 0.0f, -292.7770f}}; #endif +// 16-bit (P016) color twist matrices for nppiNV12ToRGB_16u_ColorTwist32f. +// P016 data has values in [0, 65535] (10-bit values left-shifted by 6). +// The 3x3 coefficients are the same as the 8-bit matrices, but the 4th column +// offsets are scaled by 256 since the data values are 256x larger. +// +// The CUDA version convention for offsets is the same as for 8-bit: +// CUDA >= 13: NPP internally centers U/V by subtracting 32768 (128*256). +// CUDA < 13: offsets must include the full Cb/Cr centering contribution. + +// BT.601 (kr=0.299, kb=0.114) -- used as default for unspecified colorspace +// Limited range: Y in [16*256, 235*256], UV in [16*256, 240*256] +// Y_scale = 255/(235-16) = 1.16438356 +// R = Y_scale*(Y-16*256) + 1.59602679*(Cr-128*256) +// etc. +#if CUDART_VERSION >= 13000 +const Npp32f bt601LimitedRange16ColorTwist[3][4] = { + {1.16438356f, 0.0f, 1.59602679f, -4096.0f}, + {1.16438356f, -0.391762290f, -0.812967647f, -32768.0f}, + {1.16438356f, 2.01723214f, 0.0f, -32768.0f}}; + +const Npp32f bt709FullRange16ColorTwist[3][4] = { + {1.0f, 0.0f, 1.5748f, 0.0f}, + {1.0f, -0.187324273f, -0.468124273f, -32768.0f}, + {1.0f, 1.8556f, 0.0f, -32768.0f}}; + +const Npp32f bt709LimitedRange16ColorTwist[3][4] = { + {1.16438356f, 0.0f, 1.79274107f, -4096.0f}, + {1.16438356f, -0.213248614f, -0.532909329f, -32768.0f}, + {1.16438356f, 2.11240179f, 0.0f, -32768.0f}}; + +const Npp32f bt2020FullRange16ColorTwist[3][4] = { + {1.0f, 0.0f, 1.4746f, 0.0f}, + {1.0f, -0.164553127f, -0.571353127f, -32768.0f}, + {1.0f, 1.8814f, 0.0f, -32768.0f}}; + +const Npp32f bt2020LimitedRange16ColorTwist[3][4] = { + {1.16438356f, 0.0f, 1.67867411f, -4096.0f}, + {1.16438356f, -0.187326105f, -0.650424319f, -32768.0f}, + {1.16438356f, 2.14177232f, 0.0f, -32768.0f}}; +#else +// CUDA < 13: expand centering into the offset column. +// For 16-bit, each offset = -(coeff_Y * Y_offset) - (coeff_Cb * 32768) - +// (coeff_Cr * 32768) where Y_offset = 4096 for limited, 0 for full, and +// 32768 = 128*256 is the UV centering offset. + +// BT.601 limited range +const Npp32f bt601LimitedRange16ColorTwist[3][4] = { + {1.16438356f, 0.0f, 1.59602679f, -57067.92f}, + {1.16438356f, -0.391762290f, -0.812967647f, 34707.28f}, + {1.16438356f, 2.01723214f, 0.0f, -70869.98f}}; + +// BT.709 full range +const Npp32f bt709FullRange16ColorTwist[3][4] = { + {1.0f, 0.0f, 1.5748f, -51603.04f}, + {1.0f, -0.187324273f, -0.468124273f, 21477.73f}, + {1.0f, 1.8556f, 0.0f, -60804.30f}}; + +// BT.709 limited range +const Npp32f bt709LimitedRange16ColorTwist[3][4] = { + {1.16438356f, 0.0f, 1.79274107f, -63513.85f}, + {1.16438356f, -0.213248614f, -0.532909329f, 19680.79f}, + {1.16438356f, 2.11240179f, 0.0f, -73988.50f}}; + +// BT.2020 full range +const Npp32f bt2020FullRange16ColorTwist[3][4] = { + {1.0f, 0.0f, 1.4746f, -48319.69f}, + {1.0f, -0.164553127f, -0.571353127f, 24114.17f}, + {1.0f, 1.8814f, 0.0f, -61649.73f}}; + +// BT.2020 limited range +const Npp32f bt2020LimitedRange16ColorTwist[3][4] = { + {1.16438356f, 0.0f, 1.67867411f, -59776.11f}, + {1.16438356f, -0.187326105f, -0.650424319f, 22682.09f}, + {1.16438356f, 2.14177232f, 0.0f, -74950.91f}}; +#endif + torch::stable::Tensor convertNV12FrameToRGB( UniqueAVFrame& avFrame, const StableDevice& device, @@ -407,6 +483,79 @@ torch::stable::Tensor convertNV12FrameToRGB( return dst; } +torch::stable::Tensor convertP016FrameToRGB( + UniqueAVFrame& avFrame, + const StableDevice& device, + const UniqueNppContext& nppCtx, + cudaStream_t nvdecStream, + std::optional preAllocatedOutputTensor) { + auto frameDims = FrameDims(avFrame->height, avFrame->width); + torch::stable::Tensor dst; + if (preAllocatedOutputTensor.has_value()) { + dst = preAllocatedOutputTensor.value(); + } else { + dst = allocateEmptyHWCTensor(frameDims, device, /*bitDepth=*/16); + } + + // Sync streams: make sure NVDEC has finished before we run NPP. + cudaStream_t nppStream = getCurrentCudaStream(device.index()); + syncStreams(/*runningStream=*/nvdecStream, /*waitingStream=*/nppStream); + + nppCtx->hStream = nppStream; + cudaError_t err = cudaStreamGetFlags(nppCtx->hStream, &nppCtx->nStreamFlags); + STD_TORCH_CHECK( + err == cudaSuccess, + "cudaStreamGetFlags failed: ", + cudaGetErrorString(err)); + + NppiSize oSizeROI = {frameDims.width, frameDims.height}; + // P016 data is 16-bit. The _16u_ NPP function takes Npp16u* pointers. + const Npp16u* yuvData[2] = { + reinterpret_cast(avFrame->data[0]), + reinterpret_cast(avFrame->data[1])}; + + NppStatus status; + + // nppiNV12ToRGB_16u_ColorTwist32f_P2C3R_Ctx works with P016 data (16-bit + // semi-planar YUV 4:2:0) and outputs 16-bit RGB. We always use ColorTwist + // for 16-bit since there are no pre-defined colorspace-specific 16-bit + // functions in NPP. + int srcStep[2] = {avFrame->linesize[0], avFrame->linesize[1]}; + // Output stride is in bytes. uint16 tensor stride(0) is in elements. + int dstStep = static_cast(dst.stride(0)) * 2; + + // Select the appropriate 16-bit color twist matrix based on colorspace and + // range. This mirrors the 8-bit dispatch logic in convertNV12FrameToRGB. + const Npp32f(*matrix)[4] = nullptr; + + if (avFrame->colorspace == AVColorSpace::AVCOL_SPC_BT709) { + matrix = (avFrame->color_range == AVColorRange::AVCOL_RANGE_JPEG) + ? bt709FullRange16ColorTwist + : bt709LimitedRange16ColorTwist; + } else if ( + avFrame->colorspace == AVColorSpace::AVCOL_SPC_BT2020_NCL || + avFrame->colorspace == AVColorSpace::AVCOL_SPC_BT2020_CL) { + matrix = (avFrame->color_range == AVColorRange::AVCOL_RANGE_JPEG) + ? bt2020FullRange16ColorTwist + : bt2020LimitedRange16ColorTwist; + } else { + // Default: BT.601 limited range (matches the 8-bit default behavior). + matrix = bt601LimitedRange16ColorTwist; + } + + status = nppiNV12ToRGB_16u_ColorTwist32f_P2C3R_Ctx( + yuvData, + srcStep, + reinterpret_cast(dst.mutable_data_ptr()), + dstStep, + oSizeROI, + matrix, + *nppCtx); + STD_TORCH_CHECK(status == NPP_SUCCESS, "Failed to convert P016 frame."); + + return dst; +} + UniqueNppContext getNppStreamContext(const StableDevice& device) { int deviceIndex = getDeviceIndex(device); diff --git a/src/torchcodec/_core/CUDACommon.h b/src/torchcodec/_core/CUDACommon.h index 9e50a4e25..fc2d677e5 100644 --- a/src/torchcodec/_core/CUDACommon.h +++ b/src/torchcodec/_core/CUDACommon.h @@ -38,6 +38,16 @@ torch::stable::Tensor convertNV12FrameToRGB( std::optional preAllocatedOutputTensor = std::nullopt); +// Convert a P016 (16-bit YUV 4:2:0) GPU frame to a uint16 RGB tensor. +// Used for 10-bit video decoding on the Beta CUDA path. +torch::stable::Tensor convertP016FrameToRGB( + UniqueAVFrame& avFrame, + const StableDevice& device, + const UniqueNppContext& nppCtx, + cudaStream_t nvdecStream, + std::optional preAllocatedOutputTensor = + std::nullopt); + UniqueNppContext getNppStreamContext(const StableDevice& device); void returnNppStreamContextToCache( const StableDevice& device, diff --git a/src/torchcodec/_core/CpuDeviceInterface.cpp b/src/torchcodec/_core/CpuDeviceInterface.cpp index d0925fc21..c2711fc18 100644 --- a/src/torchcodec/_core/CpuDeviceInterface.cpp +++ b/src/torchcodec/_core/CpuDeviceInterface.cpp @@ -7,6 +7,24 @@ #include "CpuDeviceInterface.h" namespace facebook::torchcodec { + +namespace { + +// Returns the appropriate RGB output format based on source bit depth. +// RGB24 is 8 bits per channel, RGB48 is 16 bits per channel. For >8-bit +// sources (e.g. 10-bit HDR), we need RGB48 to preserve the full bit depth; +// RGB24 would lose the extra precision. +AVPixelFormat getOutputPixelFormat(int bitDepth) { + return (bitDepth > 8) ? AV_PIX_FMT_RGB48 : AV_PIX_FMT_RGB24; +} + +// Returns the format filter string for the given output pixel format. +std::string getFormatFilterString(AVPixelFormat outputFormat) { + return (outputFormat == AV_PIX_FMT_RGB48) ? "format=rgb48," : "format=rgb24,"; +} + +} // namespace + namespace { static bool g_cpu = registerDeviceInterface( @@ -25,7 +43,8 @@ CpuDeviceInterface::CpuDeviceInterface(const StableDevice& device) void CpuDeviceInterface::initialize( const AVStream* avStream, [[maybe_unused]] const UniqueDecodingAVFormatContext& avFormatCtx, - const SharedAVCodecContext& codecContext) { + const SharedAVCodecContext& codecContext, + [[maybe_unused]] OutputDtype outputDtype) { STD_TORCH_CHECK(avStream != nullptr, "avStream is null"); codecContext_ = codecContext; timeBase_ = avStream->time_base; @@ -86,16 +105,25 @@ void CpuDeviceInterface::initializeVideo( if (!transforms.empty()) { // Note [Transform and Format Conversion Order] // We have to ensure that all user filters happen AFTER the explicit format - // conversion. That is, we want the filters to be applied in RGB24, not the - // pixel format of the input frame. + // conversion. That is, we want the filters to be applied in the output RGB + // color space, not the pixel format of the input frame. // - // The ouput frame will always be in RGB24, as we specify the sink node with - // AV_PIX_FORMAT_RGB24. Filtergraph will automatically insert a filter - // conversion to ensure the output frame matches the pixel format - // specified in the sink. But by default, it will insert it after the user - // filters. We need an explicit format conversion to get the behavior we - // want. - filters_ = "format=rgb24," + filters.str(); + // The output frame will always be in the output RGB format (RGB24 or + // RGB48), as we specify the sink node with the appropriate format. + // Filtergraph will automatically insert a format conversion to ensure the + // output frame matches the pixel format specified in the sink. But by + // default, it will insert it after the user filters. We need an explicit + // format conversion to get the behavior we want. + // + // Build the final filters_ string with the format prefix based on the + // resolved output bit depth. The format prefix ensures user transforms + // run in the correct output color space (RGB24 or RGB48). + int sourceBitDepth = getBitDepthFromAVPixelFormat( + static_cast(codecContext_->pix_fmt)); + int bitDepth = + resolvedBitDepth(sourceBitDepth, videoStreamOptions.outputDtype); + AVPixelFormat outputPixelFormat = getOutputPixelFormat(bitDepth); + filters_ = getFormatFilterString(outputPixelFormat) + filters.str(); } initialized_ = true; @@ -179,7 +207,14 @@ void CpuDeviceInterface::convertVideoAVFrameToFrameOutput( // FrameBatchOutputs based on the the stream metadata. But single-frame APIs // can still work in such situations, so they should. auto inputDims = FrameDims(avFrame->height, avFrame->width); - auto outputDims = resizedOutputDims_.value_or(inputDims); + auto avFrameFormat = static_cast(avFrame->format); + int sourceBitDepth = getBitDepthFromAVPixelFormat(avFrameFormat); + int bitDepth = + resolvedBitDepth(sourceBitDepth, videoStreamOptions_.outputDtype); + AVPixelFormat outputPixelFormat = getOutputPixelFormat(bitDepth); + + auto outputDims = + resizedOutputDims_.value_or(FrameDims(avFrame->height, avFrame->width)); if (preAllocatedOutputTensor.has_value()) { auto shape = preAllocatedOutputTensor.value().sizes(); @@ -200,10 +235,7 @@ void CpuDeviceInterface::convertVideoAVFrameToFrameOutput( if (colorConversionLibrary == ColorConversionLibrary::SWSCALE) { outputTensor = preAllocatedOutputTensor.value_or( - allocateEmptyHWCTensor(outputDims, kStableCPU)); - - enum AVPixelFormat avFrameFormat = - static_cast(avFrame->format); + allocateEmptyHWCTensor(outputDims, kStableCPU, bitDepth)); SwsConfig swsConfig( avFrame->width, @@ -213,8 +245,10 @@ void CpuDeviceInterface::convertVideoAVFrameToFrameOutput( outputDims.width, outputDims.height); - if (!swScale_ || swScale_->getConfig() != swsConfig) { - swScale_ = std::make_unique(swsConfig, swsFlags_); + if (!swScale_ || swScale_->getConfig() != swsConfig || + swScale_->getOutputFormat() != outputPixelFormat) { + swScale_ = + std::make_unique(swsConfig, outputPixelFormat, swsFlags_); } int resultHeight = swScale_->convert(avFrame, outputTensor); @@ -231,7 +265,8 @@ void CpuDeviceInterface::convertVideoAVFrameToFrameOutput( frameOutput.data = outputTensor; } else if (colorConversionLibrary == ColorConversionLibrary::FILTERGRAPH) { - outputTensor = convertAVFrameToTensorUsingFilterGraph(avFrame, outputDims); + outputTensor = + convertAVFrameToTensorUsingFilterGraph(avFrame, outputDims, bitDepth); // Similarly to above, if this check fails it means the frame wasn't // reshaped to its expected dimensions by filtergraph. @@ -265,9 +300,10 @@ void CpuDeviceInterface::convertVideoAVFrameToFrameOutput( torch::stable::Tensor CpuDeviceInterface::convertAVFrameToTensorUsingFilterGraph( const UniqueAVFrame& avFrame, - const FrameDims& outputDims) { - enum AVPixelFormat avFrameFormat = - static_cast(avFrame->format); + const FrameDims& outputDims, + int bitDepth) { + auto avFrameFormat = static_cast(avFrame->format); + AVPixelFormat outputFormat = getOutputPixelFormat(bitDepth); FiltersConfig filtersConfig( avFrame->width, @@ -276,7 +312,7 @@ CpuDeviceInterface::convertAVFrameToTensorUsingFilterGraph( avFrame->sample_aspect_ratio, outputDims.width, outputDims.height, - /*outputFormat=*/AV_PIX_FMT_RGB24, + /*outputFormat=*/outputFormat, filters_, timeBase_); diff --git a/src/torchcodec/_core/CpuDeviceInterface.h b/src/torchcodec/_core/CpuDeviceInterface.h index 7cec64a43..b8e8279bc 100644 --- a/src/torchcodec/_core/CpuDeviceInterface.h +++ b/src/torchcodec/_core/CpuDeviceInterface.h @@ -28,7 +28,8 @@ class CpuDeviceInterface : public DeviceInterface { virtual void initialize( const AVStream* avStream, const UniqueDecodingAVFormatContext& avFormatCtx, - const SharedAVCodecContext& codecContext) override; + const SharedAVCodecContext& codecContext, + OutputDtype outputDtype = OutputDtype::UINT8) override; virtual void initializeVideo( const VideoStreamOptions& videoStreamOptions, @@ -65,7 +66,8 @@ class CpuDeviceInterface : public DeviceInterface { torch::stable::Tensor convertAVFrameToTensorUsingFilterGraph( const UniqueAVFrame& avFrame, - const FrameDims& outputDims); + const FrameDims& outputDims, + int bitDepth); ColorConversionLibrary getColorConversionLibrary( const FrameDims& inputDims, diff --git a/src/torchcodec/_core/CudaDeviceInterface.cpp b/src/torchcodec/_core/CudaDeviceInterface.cpp index ad664f206..a993fec7b 100644 --- a/src/torchcodec/_core/CudaDeviceInterface.cpp +++ b/src/torchcodec/_core/CudaDeviceInterface.cpp @@ -111,7 +111,8 @@ CudaDeviceInterface::~CudaDeviceInterface() { void CudaDeviceInterface::initialize( const AVStream* avStream, const UniqueDecodingAVFormatContext& avFormatCtx, - const SharedAVCodecContext& codecContext) { + const SharedAVCodecContext& codecContext, + [[maybe_unused]] OutputDtype outputDtype) { STD_TORCH_CHECK(avStream != nullptr, "avStream is null"); codecContext_ = codecContext; timeBase_ = avStream->time_base; diff --git a/src/torchcodec/_core/CudaDeviceInterface.h b/src/torchcodec/_core/CudaDeviceInterface.h index d660559a1..73896a275 100644 --- a/src/torchcodec/_core/CudaDeviceInterface.h +++ b/src/torchcodec/_core/CudaDeviceInterface.h @@ -25,7 +25,8 @@ class CudaDeviceInterface : public DeviceInterface { void initialize( const AVStream* avStream, const UniqueDecodingAVFormatContext& avFormatCtx, - const SharedAVCodecContext& codecContext) override; + const SharedAVCodecContext& codecContext, + OutputDtype outputDtype = OutputDtype::UINT8) override; void initializeVideo( const VideoStreamOptions& videoStreamOptions, diff --git a/src/torchcodec/_core/DeviceInterface.cpp b/src/torchcodec/_core/DeviceInterface.cpp index 8aa7536b8..fbc9c9ec9 100644 --- a/src/torchcodec/_core/DeviceInterface.cpp +++ b/src/torchcodec/_core/DeviceInterface.cpp @@ -9,6 +9,10 @@ #include #include "StableABICompat.h" +extern "C" { +#include +} + namespace facebook::torchcodec { namespace { @@ -117,25 +121,44 @@ std::unique_ptr createDeviceInterface( } torch::stable::Tensor rgbAVFrameToTensor(const UniqueAVFrame& avFrame) { - STD_TORCH_CHECK(avFrame->format == AV_PIX_FMT_RGB24, "Expected RGB24 format"); + auto format = static_cast(avFrame->format); + STD_TORCH_CHECK( + format == AV_PIX_FMT_RGB24 || format == AV_PIX_FMT_RGB48, + "Expected RGB24 or RGB48 format, got ", + av_get_pix_fmt_name(format)); int height = avFrame->height; int width = avFrame->width; - std::vector shape = {height, width, 3}; - std::vector strides = {avFrame->linesize[0], 3, 1}; AVFrame* avFrameClone = av_frame_clone(avFrame.get()); auto deleter = [avFrameClone](void*) { UniqueAVFrame avFrameToDelete(avFrameClone); }; - return torch::stable::from_blob( - avFrameClone->data[0], - shape, - strides, - StableDevice(kStableCPU), - kStableUInt8, - deleter); + std::vector shape = {height, width, 3}; + + if (format == AV_PIX_FMT_RGB48) { + // RGB48: 6 bytes per pixel (3 channels x 2 bytes each). + // Strides are in uint16 elements: linesize is in bytes, so divide by 2. + std::vector strides = {avFrameClone->linesize[0] / 2, 3, 1}; + return torch::stable::from_blob( + avFrameClone->data[0], + shape, + strides, + StableDevice(kStableCPU), + kStableUInt16, + deleter); + } else { + // RGB24: 3 bytes per pixel. + std::vector strides = {avFrameClone->linesize[0], 3, 1}; + return torch::stable::from_blob( + avFrameClone->data[0], + shape, + strides, + StableDevice(kStableCPU), + kStableUInt8, + deleter); + } } } // namespace facebook::torchcodec diff --git a/src/torchcodec/_core/DeviceInterface.h b/src/torchcodec/_core/DeviceInterface.h index 6b5388d26..f607de346 100644 --- a/src/torchcodec/_core/DeviceInterface.h +++ b/src/torchcodec/_core/DeviceInterface.h @@ -53,10 +53,14 @@ class DeviceInterface { }; // Initialize the device with parameters generic to all kinds of decoding. + // outputDtype is passed here (rather than in initializeVideo) so that + // implementations can use it during initialization, e.g. to decide whether + // hardware decode supports the required output format. virtual void initialize( const AVStream* avStream, const UniqueDecodingAVFormatContext& avFormatCtx, - const SharedAVCodecContext& codecContext) = 0; + const SharedAVCodecContext& codecContext, + OutputDtype outputDtype = OutputDtype::UINT8) = 0; // Initialize the device with parameters specific to video decoding. There is // a default empty implementation. @@ -184,6 +188,8 @@ std::unique_ptr createDeviceInterface( const StableDevice& device, const std::string_view variant = "ffmpeg"); +// Wraps an RGB AVFrame (RGB24 or RGB48) as a torch tensor without copying. +// For RGB24: returns uint8 tensor. For RGB48: returns uint16 tensor. torch::stable::Tensor rgbAVFrameToTensor(const UniqueAVFrame& avFrame); } // namespace facebook::torchcodec diff --git a/src/torchcodec/_core/FFMPEGCommon.cpp b/src/torchcodec/_core/FFMPEGCommon.cpp index e7b3efa35..f9ab11b0f 100644 --- a/src/torchcodec/_core/FFMPEGCommon.cpp +++ b/src/torchcodec/_core/FFMPEGCommon.cpp @@ -11,6 +11,7 @@ extern "C" { #include #include +#include } namespace facebook::torchcodec { @@ -538,6 +539,14 @@ UniqueAVFrame convertAudioAVFrameSamples( return convertedAVFrame; } +int getBitDepthFromAVPixelFormat(AVPixelFormat format) { + const AVPixFmtDescriptor* desc = av_pix_fmt_desc_get(format); + if (desc && desc->nb_components > 0) { + return desc->comp[0].depth; + } + return 8; +} + void setFFmpegLogLevel() { auto logLevel = AV_LOG_QUIET; const char* logLevelEnvPtr = std::getenv("TORCHCODEC_FFMPEG_LOG_LEVEL"); diff --git a/src/torchcodec/_core/FFMPEGCommon.h b/src/torchcodec/_core/FFMPEGCommon.h index 77738c3c6..52f97ecd2 100644 --- a/src/torchcodec/_core/FFMPEGCommon.h +++ b/src/torchcodec/_core/FFMPEGCommon.h @@ -318,4 +318,8 @@ UniqueSwsContext createSwsContext( AVPixelFormat outputFormat, int swsFlags); +// Returns the bit depth per channel for the given pixel format. Returns 8 if +// the format descriptor is unavailable (e.g. AV_PIX_FMT_NONE). +int getBitDepthFromAVPixelFormat(AVPixelFormat format); + } // namespace facebook::torchcodec diff --git a/src/torchcodec/_core/Frame.cpp b/src/torchcodec/_core/Frame.cpp index 233305c27..f12b830a6 100644 --- a/src/torchcodec/_core/Frame.cpp +++ b/src/torchcodec/_core/Frame.cpp @@ -17,35 +17,35 @@ FrameDims::FrameDims(int height, int width) : height(height), width(width) { FrameBatchOutput::FrameBatchOutput( int64_t numFrames, const FrameDims& outputDims, - const StableDevice& device) + const StableDevice& device, + int bitDepth) : ptsSeconds(torch::stable::empty({numFrames}, kStableFloat64)), durationSeconds(torch::stable::empty({numFrames}, kStableFloat64)) { - data = allocateEmptyHWCTensor(outputDims, device, numFrames); + data = allocateEmptyHWCTensor(outputDims, device, bitDepth, numFrames); } torch::stable::Tensor allocateEmptyHWCTensor( const FrameDims& frameDims, const StableDevice& device, + int bitDepth, std::optional numFrames) { STD_TORCH_CHECK( frameDims.height > 0, "height must be > 0, got: ", frameDims.height); STD_TORCH_CHECK( frameDims.width > 0, "width must be > 0, got: ", frameDims.width); + auto dtype = bitDepth > 8 ? kStableUInt16 : kStableUInt8; if (numFrames.has_value()) { auto numFramesValue = numFrames.value(); STD_TORCH_CHECK( numFramesValue >= 0, "numFrames must be >= 0, got: ", numFramesValue); return torch::stable::empty( {numFramesValue, frameDims.height, frameDims.width, 3}, - kStableUInt8, + dtype, std::nullopt, device); } else { return torch::stable::empty( - {frameDims.height, frameDims.width, 3}, - kStableUInt8, - std::nullopt, - device); + {frameDims.height, frameDims.width, 3}, dtype, std::nullopt, device); } } diff --git a/src/torchcodec/_core/Frame.h b/src/torchcodec/_core/Frame.h index 5e8baa07c..7574a4d76 100644 --- a/src/torchcodec/_core/Frame.h +++ b/src/torchcodec/_core/Frame.h @@ -46,7 +46,8 @@ struct FrameBatchOutput { FrameBatchOutput( int64_t numFrames, const FrameDims& outputDims, - const StableDevice& device); + const StableDevice& device, + int bitDepth = 8); }; struct AudioFramesOutput { @@ -67,6 +68,7 @@ struct AudioFramesOutput { torch::stable::Tensor allocateEmptyHWCTensor( const FrameDims& frameDims, const StableDevice& device, + int bitDepth = 8, std::optional numFrames = std::nullopt); } // namespace facebook::torchcodec diff --git a/src/torchcodec/_core/NVDECCache.cpp b/src/torchcodec/_core/NVDECCache.cpp index 13b841072..8127287f5 100644 --- a/src/torchcodec/_core/NVDECCache.cpp +++ b/src/torchcodec/_core/NVDECCache.cpp @@ -29,8 +29,10 @@ NVDECCache& NVDECCache::getCache(const StableDevice& device) { return getCacheInstances()[getDeviceIndex(device)]; } -UniqueCUvideodecoder NVDECCache::getDecoder(CUVIDEOFORMAT* videoFormat) { - CacheKey key(videoFormat); +UniqueCUvideodecoder NVDECCache::getDecoder( + CUVIDEOFORMAT* videoFormat, + cudaVideoSurfaceFormat surfaceFormat) { + CacheKey key(videoFormat, surfaceFormat); std::lock_guard lock(cacheLock_); // Find an entry with matching key @@ -61,10 +63,11 @@ void NVDECCache::evictLRUEntry() { void NVDECCache::returnDecoder( CUVIDEOFORMAT* videoFormat, + cudaVideoSurfaceFormat surfaceFormat, UniqueCUvideodecoder decoder) { STD_TORCH_CHECK(decoder != nullptr, "decoder must not be null"); - CacheKey key(videoFormat); + CacheKey key(videoFormat, surfaceFormat); std::lock_guard lock(cacheLock_); int capacity = getNVDECCacheCapacity(); diff --git a/src/torchcodec/_core/NVDECCache.h b/src/torchcodec/_core/NVDECCache.h index 0d373e998..55a653be4 100644 --- a/src/torchcodec/_core/NVDECCache.h +++ b/src/torchcodec/_core/NVDECCache.h @@ -52,10 +52,15 @@ class NVDECCache { static NVDECCache& getCache(const StableDevice& device); // Get decoder from cache - returns nullptr if none available. - UniqueCUvideodecoder getDecoder(CUVIDEOFORMAT* videoFormat); + UniqueCUvideodecoder getDecoder( + CUVIDEOFORMAT* videoFormat, + cudaVideoSurfaceFormat surfaceFormat); // Return decoder to cache using LRU eviction. - void returnDecoder(CUVIDEOFORMAT* videoFormat, UniqueCUvideodecoder decoder); + void returnDecoder( + CUVIDEOFORMAT* videoFormat, + cudaVideoSurfaceFormat surfaceFormat, + UniqueCUvideodecoder decoder); // Iterates all per-device cache instances and evicts LRU entries until each // cache's size is at most capacity. Called from setNVDECCacheCapacity(). @@ -74,10 +79,11 @@ class NVDECCache { cudaVideoChromaFormat chromaFormat; uint32_t bitDepthLumaMinus8; uint8_t numDecodeSurfaces; + cudaVideoSurfaceFormat outputFormat; CacheKey() = delete; - explicit CacheKey(CUVIDEOFORMAT* videoFormat) { + CacheKey(CUVIDEOFORMAT* videoFormat, cudaVideoSurfaceFormat surfaceFormat) { STD_TORCH_CHECK(videoFormat != nullptr, "videoFormat must not be null"); codecType = videoFormat->codec; width = videoFormat->coded_width; @@ -85,6 +91,7 @@ class NVDECCache { chromaFormat = videoFormat->chroma_format; bitDepthLumaMinus8 = videoFormat->bit_depth_luma_minus8; numDecodeSurfaces = videoFormat->min_num_decode_surfaces; + outputFormat = surfaceFormat; } CacheKey(const CacheKey&) = default; @@ -97,14 +104,16 @@ class NVDECCache { height, chromaFormat, bitDepthLumaMinus8, - numDecodeSurfaces) < + numDecodeSurfaces, + outputFormat) < std::tie( other.codecType, other.width, other.height, other.chromaFormat, other.bitDepthLumaMinus8, - other.numDecodeSurfaces); + other.numDecodeSurfaces, + other.outputFormat); } }; diff --git a/src/torchcodec/_core/SingleStreamDecoder.cpp b/src/torchcodec/_core/SingleStreamDecoder.cpp index 49c321658..360b2c9cb 100644 --- a/src/torchcodec/_core/SingleStreamDecoder.cpp +++ b/src/torchcodec/_core/SingleStreamDecoder.cpp @@ -451,7 +451,8 @@ void SingleStreamDecoder::addStream( AVMediaType mediaType, const StableDevice& device, const std::string_view deviceVariant, - std::optional ffmpegThreadCount) { + std::optional ffmpegThreadCount, + OutputDtype outputDtype) { STD_TORCH_CHECK( activeStreamIndex_ == NO_ACTIVE_STREAM, "Can only add one single stream."); @@ -523,7 +524,7 @@ void SingleStreamDecoder::addStream( // Initialize the device interface with the codec context deviceInterface_->initialize( - streamInfo.stream, formatContext_, streamInfo.codecContext); + streamInfo.stream, formatContext_, streamInfo.codecContext, outputDtype); containerMetadata_.allStreamMetadata[activeStreamIndex_].codecName = std::string(avcodec_get_name(streamInfo.codecContext->codec_id)); @@ -553,7 +554,8 @@ void SingleStreamDecoder::addVideoStream( AVMEDIA_TYPE_VIDEO, videoStreamOptions.device, videoStreamOptions.deviceVariant, - videoStreamOptions.ffmpegThreadCount); + videoStreamOptions.ffmpegThreadCount, + videoStreamOptions.outputDtype); auto& streamMetadata = containerMetadata_.allStreamMetadata[activeStreamIndex_]; @@ -580,6 +582,10 @@ void SingleStreamDecoder::addVideoStream( // Set preRotationDims_ for the active stream. These are the raw encoded // dimensions from FFmpeg, used as a fallback for tensor pre-allocation when // no resize/rotation transforms are applied. + int sourceBitDepth = getBitDepthFromAVPixelFormat( + static_cast(streamInfo.stream->codecpar->format)); + outputBitDepth_ = + resolvedBitDepth(sourceBitDepth, videoStreamOptions.outputDtype); preRotationDims_ = FrameDims( streamInfo.stream->codecpar->height, streamInfo.stream->codecpar->width); @@ -748,7 +754,10 @@ FrameBatchOutput SingleStreamDecoder::getFramesAtIndices( const auto& streamInfo = streamInfos_[activeStreamIndex_]; const auto& videoStreamOptions = streamInfo.videoStreamOptions; FrameBatchOutput frameBatchOutput( - frameIndices.numel(), getOutputDims(), videoStreamOptions.device); + frameIndices.numel(), + getOutputDims(), + videoStreamOptions.device, + outputBitDepth_); auto frameBatchOutputPtsSeconds = mutableAccessor(frameBatchOutput.ptsSeconds); @@ -813,7 +822,10 @@ FrameBatchOutput SingleStreamDecoder::getFramesInRange( int64_t numOutputFrames = std::ceil((stop - start) / double(step)); const auto& videoStreamOptions = streamInfo.videoStreamOptions; FrameBatchOutput frameBatchOutput( - numOutputFrames, getOutputDims(), videoStreamOptions.device); + numOutputFrames, + getOutputDims(), + videoStreamOptions.device, + outputBitDepth_); auto frameBatchOutputPtsSeconds = mutableAccessor(frameBatchOutput.ptsSeconds); @@ -950,7 +962,7 @@ FrameBatchOutput SingleStreamDecoder::getFramesPlayedInRange( // below. Hence, we need this special case below. if (startSeconds == stopSeconds) { FrameBatchOutput frameBatchOutput( - 0, getOutputDims(), videoStreamOptions.device); + 0, getOutputDims(), videoStreamOptions.device, outputBitDepth_); frameBatchOutput.data = maybePermuteHWC2CHW(frameBatchOutput.data); return frameBatchOutput; } @@ -993,7 +1005,10 @@ FrameBatchOutput SingleStreamDecoder::getFramesPlayedInRange( int64_t numOutputFrames = static_cast(std::round(product)); FrameBatchOutput frameBatchOutput( - numOutputFrames, getOutputDims(), videoStreamOptions.device); + numOutputFrames, + getOutputDims(), + videoStreamOptions.device, + outputBitDepth_); auto frameBatchOutputPtsSeconds = mutableAccessor(frameBatchOutput.ptsSeconds); @@ -1039,7 +1054,7 @@ FrameBatchOutput SingleStreamDecoder::getFramesPlayedInRange( int64_t numFrames = stopFrameIndex - startFrameIndex; FrameBatchOutput frameBatchOutput( - numFrames, getOutputDims(), videoStreamOptions.device); + numFrames, getOutputDims(), videoStreamOptions.device, outputBitDepth_); auto frameBatchOutputPtsSeconds = mutableAccessor(frameBatchOutput.ptsSeconds); auto frameBatchOutputDurationSeconds = diff --git a/src/torchcodec/_core/SingleStreamDecoder.h b/src/torchcodec/_core/SingleStreamDecoder.h index 33264e545..7f849d898 100644 --- a/src/torchcodec/_core/SingleStreamDecoder.h +++ b/src/torchcodec/_core/SingleStreamDecoder.h @@ -319,7 +319,8 @@ class FORCE_PUBLIC_VISIBILITY SingleStreamDecoder { AVMediaType mediaType, const StableDevice& device = StableDevice(kStableCPU), const std::string_view deviceVariant = "ffmpeg", - std::optional ffmpegThreadCount = std::nullopt); + std::optional ffmpegThreadCount = std::nullopt, + OutputDtype outputDtype = OutputDtype::UINT8); // Returns the "best" stream index for a given media type. The "best" is // determined by various heuristics in FFMPEG. @@ -400,6 +401,9 @@ class FORCE_PUBLIC_VISIBILITY SingleStreamDecoder { std::vector> transforms_; std::optional resizedOutputDims_; FrameDims preRotationDims_; + // Resolved output bit depth (8 for SDR/uint8, >8 for HDR paths). + // Separate from FrameDims since transforms are purely geometric. + int outputBitDepth_ = 8; // Whether or not we have already scanned all streams to update the metadata. bool scannedAllStreams_ = false; diff --git a/src/torchcodec/_core/StableABICompat.h b/src/torchcodec/_core/StableABICompat.h index cac3d278d..0a756a935 100644 --- a/src/torchcodec/_core/StableABICompat.h +++ b/src/torchcodec/_core/StableABICompat.h @@ -62,6 +62,7 @@ constexpr auto kStableXPU = torch::headeronly::DeviceType::XPU; // Scalar type constants constexpr auto kStableUInt8 = torch::headeronly::ScalarType::Byte; +constexpr auto kStableUInt16 = torch::headeronly::ScalarType::UInt16; constexpr auto kStableInt32 = torch::headeronly::ScalarType::Int; constexpr auto kStableInt64 = torch::headeronly::ScalarType::Long; constexpr auto kStableFloat32 = torch::headeronly::ScalarType::Float; diff --git a/src/torchcodec/_core/StreamOptions.h b/src/torchcodec/_core/StreamOptions.h index 6cab3c8e8..9af7541d9 100644 --- a/src/torchcodec/_core/StreamOptions.h +++ b/src/torchcodec/_core/StreamOptions.h @@ -21,6 +21,27 @@ enum ColorConversionLibrary { SWSCALE }; +// Controls the dtype of decoded frame tensors. +// UINT8: Always output uint8 tensors (default, backward compatible). +// FLOAT32: Always output float32 tensors normalized to [0, 1]. +// AUTO: Output uint8 for SDR (<=8-bit) sources, float32 for HDR (>8-bit). +enum class OutputDtype { UINT8, FLOAT32, AUTO }; + +// Returns the effective output bit depth given the source bit depth and the +// user's OutputDtype setting. +// UINT8: always 8. FLOAT32: preserve source. AUTO: 8 for <=8-bit, else source. +inline int resolvedBitDepth(int sourceBitDepth, OutputDtype outputDtype) { + switch (outputDtype) { + case OutputDtype::UINT8: + return 8; + case OutputDtype::FLOAT32: + return sourceBitDepth; + case OutputDtype::AUTO: + return (sourceBitDepth <= 8) ? 8 : sourceBitDepth; + } + return 8; +} + struct VideoStreamOptions { VideoStreamOptions() {} @@ -47,6 +68,9 @@ struct VideoStreamOptions { // Device variant (e.g., "ffmpeg", "beta", etc.) std::string_view deviceVariant = "ffmpeg"; + // Controls the dtype of decoded frame tensors. + OutputDtype outputDtype = OutputDtype::UINT8; + // Encoding options std::optional codec; // Optional pixel format for video encoding (e.g., "yuv420p", "yuv444p") diff --git a/src/torchcodec/_core/SwScale.cpp b/src/torchcodec/_core/SwScale.cpp index ca748f170..f444f37cd 100644 --- a/src/torchcodec/_core/SwScale.cpp +++ b/src/torchcodec/_core/SwScale.cpp @@ -9,13 +9,18 @@ namespace facebook::torchcodec { -SwScale::SwScale(const SwsConfig& config, int swsFlags) - : config_(config), swsFlags_(swsFlags) { +SwScale::SwScale( + const SwsConfig& config, + AVPixelFormat outputFormat, + int swsFlags) + : config_(config), outputFormat_(outputFormat), swsFlags_(swsFlags) { needsResize_ = (config_.inputHeight != config_.outputHeight || config_.inputWidth != config_.outputWidth); - // Create color conversion context (input format -> RGB24). + bytesPerPixel_ = (outputFormat_ == AV_PIX_FMT_RGB48) ? 6 : 3; + + // Create color conversion context (input format -> output RGB format). // Color conversion always outputs at the input resolution. // When no resize is needed, input and output resolutions are the same. SwsConfig colorConversionFrameConfig( @@ -30,25 +35,25 @@ SwScale::SwScale(const SwsConfig& config, int swsFlags) colorConversionFrameConfig, // See [Transform and Format Conversion Order] for more on the output // pixel format. - /*outputFormat=*/AV_PIX_FMT_RGB24, + /*outputFormat=*/outputFormat_, // No flags for color conversion. When resizing is needed, we use a // separate swscale context with the appropriate resize flags. /*swsFlags=*/0); - // Create resize context if needed (RGB24 at input resolution -> RGB24 at - // output resolution). + // Create resize context if needed (output RGB at input resolution -> + // output RGB at output resolution). if (needsResize_) { SwsConfig resizeFrameConfig( config_.inputWidth, config_.inputHeight, - AV_PIX_FMT_RGB24, + outputFormat_, AVCOL_SPC_RGB, config_.outputWidth, config_.outputHeight); resizeSwsContext_ = createSwsContext( resizeFrameConfig, - /*outputFormat=*/AV_PIX_FMT_RGB24, + /*outputFormat=*/outputFormat_, /*swsFlags=*/swsFlags_); } } @@ -56,25 +61,30 @@ SwScale::SwScale(const SwsConfig& config, int swsFlags) int SwScale::convert( const UniqueAVFrame& avFrame, torch::stable::Tensor& outputTensor) { - // When resizing is needed, we do sws_scale twice: first convert to RGB24 at - // original resolution, then resize in RGB24 space. This ensures transforms - // happen in the output color space (RGB24) rather than the input color space - // (YUV). + // When resizing is needed, we do sws_scale twice: first convert to output + // RGB at original resolution, then resize in output RGB space. This ensures + // transforms happen in the output color space (RGB) rather than the input + // color space (YUV). // // When no resize is needed, we do color conversion directly into the output // tensor. + int bitDepth = (outputFormat_ == AV_PIX_FMT_RGB48) ? 16 : 8; torch::stable::Tensor colorConvertedTensor = needsResize_ ? allocateEmptyHWCTensor( - FrameDims(config_.inputHeight, config_.inputWidth), kStableCPU) + FrameDims(config_.inputHeight, config_.inputWidth), + kStableCPU, + bitDepth) : outputTensor; + // sws_scale always takes uint8_t* pointers regardless of actual bit depth. uint8_t* colorConvertedPointers[4] = { - colorConvertedTensor.mutable_data_ptr(), + static_cast(colorConvertedTensor.mutable_data_ptr()), nullptr, nullptr, nullptr}; int colorConvertedWidth = static_cast(colorConvertedTensor.sizes()[1]); - int colorConvertedLinesizes[4] = {colorConvertedWidth * 3, 0, 0, 0}; + int colorConvertedLinesizes[4] = { + colorConvertedWidth * bytesPerPixel_, 0, 0, 0}; int colorConvertedHeight = sws_scale( colorConversionSwsContext_.get(), @@ -94,16 +104,19 @@ int SwScale::convert( if (needsResize_) { uint8_t* srcPointers[4] = { - colorConvertedTensor.mutable_data_ptr(), + static_cast(colorConvertedTensor.mutable_data_ptr()), nullptr, nullptr, nullptr}; - int srcLinesizes[4] = {config_.inputWidth * 3, 0, 0, 0}; + int srcLinesizes[4] = {config_.inputWidth * bytesPerPixel_, 0, 0, 0}; uint8_t* dstPointers[4] = { - outputTensor.mutable_data_ptr(), nullptr, nullptr, nullptr}; + static_cast(outputTensor.mutable_data_ptr()), + nullptr, + nullptr, + nullptr}; int expectedOutputWidth = static_cast(outputTensor.sizes()[1]); - int dstLinesizes[4] = {expectedOutputWidth * 3, 0, 0, 0}; + int dstLinesizes[4] = {expectedOutputWidth * bytesPerPixel_, 0, 0, 0}; colorConvertedHeight = sws_scale( resizeSwsContext_.get(), diff --git a/src/torchcodec/_core/SwScale.h b/src/torchcodec/_core/SwScale.h index 8dd994365..5be73d1df 100644 --- a/src/torchcodec/_core/SwScale.h +++ b/src/torchcodec/_core/SwScale.h @@ -14,17 +14,22 @@ namespace facebook::torchcodec { struct FrameDims; // SwScale uses a double swscale path: -// 1. Color conversion (e.g., YUV -> RGB24) at the original frame resolution -// 2. Resize in RGB24 space (if resizing is needed) +// 1. Color conversion (e.g., YUV -> RGB24/RGB48) at the original frame +// resolution +// 2. Resize in output RGB space (if resizing is needed) // // This approach ensures that transforms happen in the output color space -// (RGB24) rather than the input color space (YUV). +// (RGB) rather than the input color space (YUV). // // The caller is responsible for caching SwScale instances and recreating them // when the context changes, similar to how FilterGraph is managed. class SwScale { public: - SwScale(const SwsConfig& config, int swsFlags = SWS_BILINEAR); + // outputFormat: AV_PIX_FMT_RGB24 for 8-bit, AV_PIX_FMT_RGB48 for >8-bit + SwScale( + const SwsConfig& config, + AVPixelFormat outputFormat = AV_PIX_FMT_RGB24, + int swsFlags = SWS_BILINEAR); int convert( const UniqueAVFrame& avFrame, @@ -34,15 +39,24 @@ class SwScale { return config_; } + AVPixelFormat getOutputFormat() const { + return outputFormat_; + } + private: SwsConfig config_; + AVPixelFormat outputFormat_; int swsFlags_; bool needsResize_; - // Color conversion context (input format -> RGB24 at original resolution). + // Bytes per pixel for the output format (3 for RGB24, 6 for RGB48). + int bytesPerPixel_; + + // Color conversion context (input format -> output RGB at original + // resolution). UniqueSwsContext colorConversionSwsContext_; - // Resize context (RGB24 -> RGB24 at output resolution). + // Resize context (output RGB at input res -> output RGB at output res). // May be null if no resize is needed. UniqueSwsContext resizeSwsContext_; }; diff --git a/src/torchcodec/_core/_decoder_utils.py b/src/torchcodec/_core/_decoder_utils.py index f1e6d0eba..f451dab91 100644 --- a/src/torchcodec/_core/_decoder_utils.py +++ b/src/torchcodec/_core/_decoder_utils.py @@ -166,6 +166,7 @@ def create_video_decoder( device_variant: str = "ffmpeg", transforms: Sequence[DecoderTransform | nn.Module] | None = None, custom_frame_mappings: tuple[Tensor, Tensor, Tensor] | None = None, + output_dtype: str | None = None, ) -> tuple[Tensor, int, VideoStreamMetadata]: decoder = create_decoder(source=source, seek_mode=seek_mode) @@ -191,6 +192,7 @@ def create_video_decoder( device_variant=device_variant, transform_specs=transform_specs, custom_frame_mappings=custom_frame_mappings, + output_dtype=output_dtype, ) return (decoder, stream_index, metadata) diff --git a/src/torchcodec/_core/custom_ops.cpp b/src/torchcodec/_core/custom_ops.cpp index 7746175bd..58bcb3846 100644 --- a/src/torchcodec/_core/custom_ops.cpp +++ b/src/torchcodec/_core/custom_ops.cpp @@ -57,9 +57,9 @@ STABLE_TORCH_LIBRARY(torchcodec_ns, m) { m.def( "_create_from_file_like(int file_like_context, str? seek_mode=None) -> Tensor"); m.def( - "_add_video_stream(Tensor(a!) decoder, *, int? num_threads=None, str? dimension_order=None, int? stream_index=None, str device=\"cpu\", str device_variant=\"ffmpeg\", str transform_specs=\"\", Tensor? custom_frame_mappings_pts=None, Tensor? custom_frame_mappings_duration=None, Tensor? custom_frame_mappings_keyframe_indices=None, str? color_conversion_library=None) -> ()"); + "_add_video_stream(Tensor(a!) decoder, *, int? num_threads=None, str? dimension_order=None, int? stream_index=None, str device=\"cpu\", str device_variant=\"ffmpeg\", str transform_specs=\"\", Tensor? custom_frame_mappings_pts=None, Tensor? custom_frame_mappings_duration=None, Tensor? custom_frame_mappings_keyframe_indices=None, str? color_conversion_library=None, str? output_dtype=None) -> ()"); m.def( - "add_video_stream(Tensor(a!) decoder, *, int? num_threads=None, str? dimension_order=None, int? stream_index=None, str device=\"cpu\", str device_variant=\"ffmpeg\", str transform_specs=\"\", Tensor? custom_frame_mappings_pts=None, Tensor? custom_frame_mappings_duration=None, Tensor? custom_frame_mappings_keyframe_indices=None) -> ()"); + "add_video_stream(Tensor(a!) decoder, *, int? num_threads=None, str? dimension_order=None, int? stream_index=None, str device=\"cpu\", str device_variant=\"ffmpeg\", str transform_specs=\"\", Tensor? custom_frame_mappings_pts=None, Tensor? custom_frame_mappings_duration=None, Tensor? custom_frame_mappings_keyframe_indices=None, str? output_dtype=None) -> ()"); m.def( "add_audio_stream(Tensor(a!) decoder, *, int? stream_index=None, int? sample_rate=None, int? num_channels=None) -> ()"); m.def("seek_to_pts(Tensor(a!) decoder, float seconds) -> ()"); @@ -507,10 +507,28 @@ void _add_video_stream( std::nullopt, std::optional custom_frame_mappings_keyframe_indices = std::nullopt, - std::optional color_conversion_library = std::nullopt) { + std::optional color_conversion_library = std::nullopt, + std::optional output_dtype = std::nullopt) { VideoStreamOptions videoStreamOptions; videoStreamOptions.ffmpegThreadCount = num_threads; + if (output_dtype.has_value()) { + const std::string& val = *output_dtype; + if (val == "uint8") { + videoStreamOptions.outputDtype = OutputDtype::UINT8; + } else if (val == "float32") { + videoStreamOptions.outputDtype = OutputDtype::FLOAT32; + } else if (val == "auto") { + videoStreamOptions.outputDtype = OutputDtype::AUTO; + } else { + STD_TORCH_CHECK( + false, + "Invalid output_dtype=", + val, + ". Supported values are: uint8, float32, auto."); + } + } + if (dimension_order.has_value()) { STD_TORCH_CHECK( *dimension_order == "NHWC" || *dimension_order == "NCHW", @@ -580,7 +598,8 @@ void add_video_stream( std::optional custom_frame_mappings_duration = std::nullopt, std::optional - custom_frame_mappings_keyframe_indices = std::nullopt) { + custom_frame_mappings_keyframe_indices = std::nullopt, + std::optional output_dtype = std::nullopt) { _add_video_stream( decoder, num_threads, @@ -591,7 +610,9 @@ void add_video_stream( std::move(transform_specs), std::move(custom_frame_mappings_pts), std::move(custom_frame_mappings_duration), - std::move(custom_frame_mappings_keyframe_indices)); + std::move(custom_frame_mappings_keyframe_indices), + /*color_conversion_library=*/std::nullopt, + std::move(output_dtype)); } void add_audio_stream( diff --git a/src/torchcodec/_core/ops.py b/src/torchcodec/_core/ops.py index fd67ac4cb..71cf143d5 100644 --- a/src/torchcodec/_core/ops.py +++ b/src/torchcodec/_core/ops.py @@ -84,6 +84,7 @@ def add_video_stream( custom_frame_mappings: ( tuple[torch.Tensor, torch.Tensor, torch.Tensor] | None ) = None, + output_dtype: str | None = None, ) -> None: custom_frame_mappings_pts: torch.Tensor | None = None custom_frame_mappings_keyframe_indices: torch.Tensor | None = None @@ -105,6 +106,7 @@ def add_video_stream( custom_frame_mappings_pts=custom_frame_mappings_pts, custom_frame_mappings_duration=custom_frame_mappings_duration, custom_frame_mappings_keyframe_indices=custom_frame_mappings_keyframe_indices, + output_dtype=output_dtype, ) diff --git a/src/torchcodec/_frame.py b/src/torchcodec/_frame.py index 2ceb890b7..7d57cd73a 100644 --- a/src/torchcodec/_frame.py +++ b/src/torchcodec/_frame.py @@ -70,7 +70,7 @@ class FrameBatch(Iterable): """ data: Tensor - """The frames data (``torch.Tensor`` of uint8).""" + """The frames data (``torch.Tensor`` of uint8 for 8-bit sources, uint16 for 10-bit HDR sources).""" pts_seconds: Tensor """The :term:`pts` of the frame, in seconds (``torch.Tensor`` of floats).""" duration_seconds: Tensor diff --git a/src/torchcodec/decoders/_video_decoder.py b/src/torchcodec/decoders/_video_decoder.py index 19b415c2c..f7db1eae1 100644 --- a/src/torchcodec/decoders/_video_decoder.py +++ b/src/torchcodec/decoders/_video_decoder.py @@ -169,6 +169,7 @@ def __init__( custom_frame_mappings: ( str | bytes | io.RawIOBase | io.BufferedReader | None ) = None, + output_dtype: "torch.dtype | str | None" = None, ): torch._C._log_api_usage_once("torchcodec.decoders.VideoDecoder") allowed_seek_modes = ("exact", "approximate") @@ -203,6 +204,18 @@ def __init__( if num_ffmpeg_threads is None: raise ValueError(f"{num_ffmpeg_threads = } should be an int.") + if output_dtype is None or output_dtype == torch.uint8: + _output_dtype_str = "uint8" + elif output_dtype == torch.float32: + _output_dtype_str = "float32" + elif output_dtype == "auto": + _output_dtype_str = "auto" + else: + raise ValueError( + f"Invalid output_dtype ({output_dtype}). " + f'Supported values are "auto", torch.uint8, and torch.float32.' + ) + device_variant = _get_cuda_backend() if device is None: device = str(torch.get_default_device()) @@ -222,6 +235,7 @@ def __init__( device_variant=device_variant, transforms=transforms, custom_frame_mappings=custom_frame_mappings_data, + output_dtype=_output_dtype_str, ) assert self.metadata.begin_stream_seconds is not None # mypy. @@ -232,6 +246,10 @@ def __init__( self._end_stream_seconds = self.metadata.end_stream_seconds self._num_frames = self.metadata.num_frames + # Store whether we need to convert C++ uint16 output to float32. + # C++ returns uint8 or uint16; we convert to float32 in Python. + self._output_dtype = output_dtype + self._cpu_fallback = CpuFallbackStatus() if device.startswith("cuda"): if device_variant == "beta": @@ -271,11 +289,20 @@ def cpu_fallback(self) -> CpuFallbackStatus: return self._cpu_fallback + def _maybe_to_float32(self, tensor: Tensor) -> Tensor: + """Convert uint8/uint16 tensor to float32 [0, 1] if output_dtype requires it.""" + if self._output_dtype == torch.float32 or ( + self._output_dtype == "auto" and tensor.dtype == torch.uint16 + ): + max_val = 65535.0 if tensor.dtype == torch.uint16 else 255.0 + return tensor.to(torch.float32).div_(max_val) + return tensor + def _getitem_int(self, key: int) -> Tensor: assert isinstance(key, int) frame_data, *_ = core.get_frame_at_index(self._decoder, frame_index=key) - return frame_data + return self._maybe_to_float32(frame_data) def _getitem_slice(self, key: slice) -> Tensor: assert isinstance(key, slice) @@ -287,7 +314,7 @@ def _getitem_slice(self, key: slice) -> Tensor: stop=stop, step=step, ) - return frame_data + return self._maybe_to_float32(frame_data) def __getitem__(self, key: numbers.Integral | slice) -> Tensor: """Return frame or frames as tensors, at the given index or range. @@ -341,7 +368,7 @@ def get_frame_at(self, index: int) -> Frame: self._decoder, frame_index=index ) return Frame( - data=data, + data=self._maybe_to_float32(data), pts_seconds=pts_seconds.item(), duration_seconds=duration_seconds.item(), ) @@ -361,7 +388,7 @@ def get_frames_at(self, indices: torch.Tensor | list[int]) -> FrameBatch: ) return FrameBatch( - data=data, + data=self._maybe_to_float32(data), pts_seconds=pts_seconds, duration_seconds=duration_seconds, ) @@ -382,13 +409,17 @@ def get_frames_in_range(self, start: int, stop: int, step: int = 1) -> FrameBatc """ # Adjust start / stop indices to enable indexing semantics, ex. [-10, 1000] returns the last 10 frames start, stop, step = slice(start, stop, step).indices(self._num_frames) - frames = core.get_frames_in_range( + data, pts_seconds, duration_seconds = core.get_frames_in_range( self._decoder, start=start, stop=stop, step=step, ) - return FrameBatch(*frames) + return FrameBatch( + data=self._maybe_to_float32(data), + pts_seconds=pts_seconds, + duration_seconds=duration_seconds, + ) def get_frame_played_at(self, seconds: float) -> Frame: """Return a single frame played at the given timestamp in seconds. @@ -418,7 +449,7 @@ def get_frame_played_at(self, seconds: float) -> Frame: self._decoder, seconds ) return Frame( - data=data, + data=self._maybe_to_float32(data), pts_seconds=pts_seconds.item(), duration_seconds=duration_seconds.item(), ) @@ -437,7 +468,7 @@ def get_frames_played_at(self, seconds: torch.Tensor | list[float]) -> FrameBatc self._decoder, timestamps=seconds ) return FrameBatch( - data=data, + data=self._maybe_to_float32(data), pts_seconds=pts_seconds, duration_seconds=duration_seconds, ) @@ -478,13 +509,17 @@ def get_frames_played_in_range( f"Invalid stop seconds: {stop_seconds}. " f"It must be less than or equal to {self._end_stream_seconds}." ) - frames = core.get_frames_by_pts_in_range( + data, pts_seconds, duration_seconds = core.get_frames_by_pts_in_range( self._decoder, start_seconds=start_seconds, stop_seconds=stop_seconds, fps=fps, ) - return FrameBatch(*frames) + return FrameBatch( + data=self._maybe_to_float32(data), + pts_seconds=pts_seconds, + duration_seconds=duration_seconds, + ) def get_all_frames(self, fps: float | None = None) -> FrameBatch: """Returns all frames in the video. diff --git a/test/generate_reference_resources.py b/test/generate_reference_resources.py index 7d28e9993..ef4dde120 100644 --- a/test/generate_reference_resources.py +++ b/test/generate_reference_resources.py @@ -12,7 +12,15 @@ import torch from PIL import Image -from .utils import AV1_VIDEO, H265_VIDEO, NASA_VIDEO, TestVideo +from .utils import ( + AV1_VIDEO, + H265_VIDEO, + NASA_VIDEO, + NASA_VIDEO_HDR, + TEST_SRC_2_12BIT_HDR, + TEST_SRC_2_720P_HDR, + TestVideo, +) # Run this script to update the resources used in unit tests. The resources are all derived # from source media already checked into the repo. @@ -80,6 +88,47 @@ def generate_frame_by_index( convert_image_to_tensor(output_bmp) +def generate_frame_by_index_rgb48( + video: TestVideo, + *, + frame_index: int, + stream_index: int, +) -> None: + base_path = video.get_base_path_by_index(frame_index, stream_index=stream_index) + output_raw = f"{base_path}.rgb48.raw" + + select = f"select='eq(n\\,{frame_index})'" + filtergraph = ",".join([select, "format=rgb48"]) + + cmd = [ + "ffmpeg", + "-y", + "-i", + video.path, + "-map", + f"0:{stream_index}", + "-vf", + filtergraph, + "-fps_mode", + "passthrough", + "-f", + "rawvideo", + "-pix_fmt", + "rgb48", + "-frames:v", + "1", + output_raw, + ] + subprocess.run(cmd, check=True) + + height = video.get_height(stream_index=stream_index) + width = video.get_width(stream_index=stream_index) + data = np.fromfile(output_raw, dtype=np.uint16).reshape(height, width, 3) + tensor = torch.from_numpy(data.copy()) + torch.save(tensor, f"{base_path}.rgb48.pt", _use_new_zipfile_serialization=True) + Path(output_raw).unlink() + + def generate_frame_by_timestamp( video_path: str, timestamp: float, output_path: str ) -> None: @@ -165,10 +214,18 @@ def generate_av1_video_references(): generate_frame_by_index(AV1_VIDEO, frame_index=frame, stream_index=0) +def generate_hdr_references_rgb48(): + frames = [0, 5, 10] + for video in (NASA_VIDEO_HDR, TEST_SRC_2_720P_HDR, TEST_SRC_2_12BIT_HDR): + for frame in frames: + generate_frame_by_index_rgb48(video, frame_index=frame, stream_index=0) + + def main(): generate_nasa_13013_references() generate_h265_video_references() generate_av1_video_references() + generate_hdr_references_rgb48() if __name__ == "__main__": diff --git a/test/resources/nasa_13013_hdr.mp4 b/test/resources/nasa_13013_hdr.mp4 new file mode 100644 index 000000000..671a869e5 Binary files /dev/null and b/test/resources/nasa_13013_hdr.mp4 differ diff --git a/test/resources/nasa_13013_hdr.mp4.stream0.frame000000.rgb48.pt b/test/resources/nasa_13013_hdr.mp4.stream0.frame000000.rgb48.pt new file mode 100644 index 000000000..ef9401936 Binary files /dev/null and b/test/resources/nasa_13013_hdr.mp4.stream0.frame000000.rgb48.pt differ diff --git a/test/resources/nasa_13013_hdr.mp4.stream0.frame000005.rgb48.pt b/test/resources/nasa_13013_hdr.mp4.stream0.frame000005.rgb48.pt new file mode 100644 index 000000000..2ca12d76c Binary files /dev/null and b/test/resources/nasa_13013_hdr.mp4.stream0.frame000005.rgb48.pt differ diff --git a/test/resources/nasa_13013_hdr.mp4.stream0.frame000010.rgb48.pt b/test/resources/nasa_13013_hdr.mp4.stream0.frame000010.rgb48.pt new file mode 100644 index 000000000..e4f5e8003 Binary files /dev/null and b/test/resources/nasa_13013_hdr.mp4.stream0.frame000010.rgb48.pt differ diff --git a/test/resources/testsrc2_12bit_hdr.mp4 b/test/resources/testsrc2_12bit_hdr.mp4 new file mode 100644 index 000000000..e99b0cda7 Binary files /dev/null and b/test/resources/testsrc2_12bit_hdr.mp4 differ diff --git a/test/resources/testsrc2_12bit_hdr.mp4.stream0.frame000000.rgb48.pt b/test/resources/testsrc2_12bit_hdr.mp4.stream0.frame000000.rgb48.pt new file mode 100644 index 000000000..36dc4bbe8 Binary files /dev/null and b/test/resources/testsrc2_12bit_hdr.mp4.stream0.frame000000.rgb48.pt differ diff --git a/test/resources/testsrc2_12bit_hdr.mp4.stream0.frame000005.rgb48.pt b/test/resources/testsrc2_12bit_hdr.mp4.stream0.frame000005.rgb48.pt new file mode 100644 index 000000000..384981045 Binary files /dev/null and b/test/resources/testsrc2_12bit_hdr.mp4.stream0.frame000005.rgb48.pt differ diff --git a/test/resources/testsrc2_12bit_hdr.mp4.stream0.frame000010.rgb48.pt b/test/resources/testsrc2_12bit_hdr.mp4.stream0.frame000010.rgb48.pt new file mode 100644 index 000000000..1b8b13d98 Binary files /dev/null and b/test/resources/testsrc2_12bit_hdr.mp4.stream0.frame000010.rgb48.pt differ diff --git a/test/resources/testsrc2_hdr.mp4 b/test/resources/testsrc2_hdr.mp4 new file mode 100644 index 000000000..8635c437c Binary files /dev/null and b/test/resources/testsrc2_hdr.mp4 differ diff --git a/test/resources/testsrc2_hdr.mp4.stream0.frame000000.rgb48.pt b/test/resources/testsrc2_hdr.mp4.stream0.frame000000.rgb48.pt new file mode 100644 index 000000000..0d3f5ad9b Binary files /dev/null and b/test/resources/testsrc2_hdr.mp4.stream0.frame000000.rgb48.pt differ diff --git a/test/resources/testsrc2_hdr.mp4.stream0.frame000005.rgb48.pt b/test/resources/testsrc2_hdr.mp4.stream0.frame000005.rgb48.pt new file mode 100644 index 000000000..2b92bcb52 Binary files /dev/null and b/test/resources/testsrc2_hdr.mp4.stream0.frame000005.rgb48.pt differ diff --git a/test/resources/testsrc2_hdr.mp4.stream0.frame000010.rgb48.pt b/test/resources/testsrc2_hdr.mp4.stream0.frame000010.rgb48.pt new file mode 100644 index 000000000..9f5f16361 Binary files /dev/null and b/test/resources/testsrc2_hdr.mp4.stream0.frame000010.rgb48.pt differ diff --git a/test/test_decoders.py b/test/test_decoders.py index bf252abb0..3eb27f498 100644 --- a/test/test_decoders.py +++ b/test/test_decoders.py @@ -38,14 +38,15 @@ get_ffmpeg_minor_version, get_python_version, H264_10BITS, - H265_10BITS, H265_VIDEO, in_fbcode, + IS_WINDOWS, make_video_decoder, NASA_AUDIO, NASA_AUDIO_MP3, NASA_AUDIO_MP3_44100, NASA_VIDEO, + NASA_VIDEO_HDR, NASA_VIDEO_ROTATED, needs_cuda, needs_ffmpeg_cli, @@ -56,8 +57,10 @@ SINE_MONO_S32_44100, SINE_MONO_S32_8000, TEST_NON_ZERO_START, + TEST_SRC_2_12BIT_HDR, TEST_SRC_2_720P, TEST_SRC_2_720P_H265, + TEST_SRC_2_720P_HDR, TEST_SRC_2_720P_MPEG4, TEST_SRC_2_720P_VP8, TEST_SRC_2_720P_VP9, @@ -184,6 +187,11 @@ def test_create_fails(self): with pytest.raises(ValueError, match="Invalid seek mode"): VideoDecoder(NASA_VIDEO.path, seek_mode="blah") + @pytest.mark.parametrize("output_dtype", [torch.int32, "blah", 123]) + def test_create_fails_invalid_output_dtype(self, output_dtype): + with pytest.raises(ValueError, match="Invalid output_dtype"): + VideoDecoder(NASA_VIDEO.path, output_dtype=output_dtype) + @pytest.mark.parametrize("num_ffmpeg_threads", (1, 4)) @pytest.mark.parametrize("device", all_supported_devices()) @pytest.mark.parametrize("seek_mode", ("exact", "approximate")) @@ -1520,8 +1528,6 @@ def test_10bit_gpu_fallsback_to_cpu(self): # do the color conversion on the CPU. # Here we just assert that the GPU results are the same as the CPU # results. - # TODO see other TODO below in test_10bit_videos_cpu: we should validate - # the frames against a reference. # We know from previous tests that the H264_10BITS video isn't supported # by NVDEC, so NVDEC decodes it on the CPU. @@ -1545,14 +1551,167 @@ def test_10bit_gpu_fallsback_to_cpu(self): frames_cpu = decoder_cpu.get_frames_at(frame_indices).data assert_frames_equal(frames_gpu.cpu(), frames_cpu) - @pytest.mark.parametrize("device", all_supported_devices()) - @pytest.mark.parametrize("asset", (H264_10BITS, H265_10BITS)) - def test_10bit_videos(self, device, asset): - # This just validates that we can decode 10-bit videos. - # TODO validate against the ref that the decoded frames are correct + @pytest.mark.parametrize( + "device", + ( + "cpu", + # Note: the FFmpeg CUDA batch API pre-allocates uint8 tensors, + # which doesn't work for 10-bit content (uint16). This is fine + # since we're moving to Beta CUDA. + # pytest.param("cuda", marks=pytest.mark.needs_cuda), + pytest.param("cuda:beta", marks=pytest.mark.needs_cuda), + ), + ) + @pytest.mark.parametrize( + "asset", + ( + pytest.param( + NASA_VIDEO_HDR, + marks=pytest.mark.skipif( + IS_WINDOWS and ffmpeg_major_version < 5, + reason="swscale YUV10->RGB48 differs on Windows + FFmpeg 4", + ), + ), + pytest.param( + TEST_SRC_2_720P_HDR, + marks=pytest.mark.skipif( + IS_WINDOWS and ffmpeg_major_version < 5, + reason="swscale YUV10->RGB48 differs on Windows + FFmpeg 4", + ), + ), + ), + ) + def test_10bit_batch_apis(self, device, asset): + # Validate batch APIs on 10-bit float32 content against ffmpeg rgb48 + # references. + decoder, device = make_video_decoder( + asset.path, device=device, output_dtype=torch.float32 + ) + + def assert_frame_correct(frame, frame_index): + frame_as_uint16 = (frame.cpu() * 65535).round().to(torch.uint16) + ref = asset.get_frame_data_by_index_rgb48(frame_index) + if device == "cpu": + torch.testing.assert_close(frame_as_uint16, ref, rtol=0, atol=0) + else: + frame_10bit = frame_as_uint16.to(torch.int32) >> 6 + ref_10bit = ref.to(torch.int32) >> 6 + assert_tensor_close_on_at_least( + frame_10bit, ref_10bit, percentage=88, atol=3 + ) + + # __getitem__ + assert_frame_correct(decoder[0], 0) + + # get_frame_at + assert_frame_correct(decoder.get_frame_at(5).data, 5) + + # get_frame_played_at + pts = decoder.metadata.begin_stream_seconds + assert_frame_correct(decoder.get_frame_played_at(pts).data, 0) + + # get_frames_at + indices = [0, 5, 10] + frames = decoder.get_frames_at(indices) + for i, idx in enumerate(indices): + assert_frame_correct(frames.data[i], idx) + + # get_frames_in_range — validate frames 5 and 10 against references + frames_range = decoder.get_frames_in_range(start=5, stop=11) + assert_frame_correct(frames_range.data[0], 5) + assert_frame_correct(frames_range.data[5], 10) + + def test_output_dtype_auto(self): + # "auto" should produce uint8 for SDR (<=8-bit) sources and float32 + # for HDR (>8-bit) sources. Values are validated against ffmpeg CLI + # references (rgb24 for SDR, rgb48 for HDR). + + # SDR source: auto should produce uint8 matching ffmpeg rgb24 + decoder_auto = VideoDecoder(NASA_VIDEO.path, output_dtype="auto") + + for frame_index in [0, 1, 9]: + frame = decoder_auto.get_frame_at(frame_index).data + assert frame.dtype == torch.uint8 + ref = NASA_VIDEO.get_frame_data_by_index(frame_index) + assert_frames_equal(frame, ref) + + # HDR source: auto should produce float32 matching ffmpeg rgb48 + decoder_auto_hdr = VideoDecoder(NASA_VIDEO_HDR.path, output_dtype="auto") + + for frame_index in [0, 5, 10]: + frame = decoder_auto_hdr.get_frame_at(frame_index).data + assert frame.dtype == torch.float32 - decoder, _ = make_video_decoder(asset.path, device=device) - decoder.get_frame_at(10) + if IS_WINDOWS and ffmpeg_major_version < 5: + # swscale YUV10->RGB48 differs on Windows + FFmpeg 4 + continue + + frame_as_uint16 = (frame * 65535).round().to(torch.uint16) + ref = NASA_VIDEO_HDR.get_frame_data_by_index_rgb48(frame_index) + torch.testing.assert_close(frame_as_uint16, ref, rtol=0, atol=0) + + @pytest.mark.parametrize( + "device", + ( + "cpu", + pytest.param("cuda:beta", marks=pytest.mark.needs_cuda), + ), + ) + @pytest.mark.parametrize( + "asset", + ( + pytest.param( + NASA_VIDEO_HDR, + marks=pytest.mark.skipif( + IS_WINDOWS and ffmpeg_major_version < 5, + reason="swscale YUV10->RGB48 differs on Windows + FFmpeg 4", + ), + ), + pytest.param( + TEST_SRC_2_720P_HDR, + marks=pytest.mark.skipif( + IS_WINDOWS and ffmpeg_major_version < 5, + reason="swscale YUV10->RGB48 differs on Windows + FFmpeg 4", + ), + ), + pytest.param( + TEST_SRC_2_12BIT_HDR, + marks=pytest.mark.skipif( + IS_WINDOWS and ffmpeg_major_version < 5, + reason="swscale YUV10->RGB48 differs on Windows + FFmpeg 4", + ), + ), + ), + ) + @pytest.mark.parametrize("seek_mode", ("exact", "approximate")) + def test_10bit_float32_against_ffmpeg_rgb48(self, asset, device, seek_mode): + # Validate float32 HDR decode against ffmpeg CLI's rgb48 output. + # CPU uses the same swscale path as ffmpeg, so we expect exact match. + # GPU (NVDEC + NPP) uses a different color conversion pipeline, so we + # allow a small tolerance in 10-bit space. + decoder, device = make_video_decoder( + asset.path, + device=device, + output_dtype=torch.float32, + seek_mode=seek_mode, + ) + + for frame_index in [0, 5, 10]: + frame = decoder.get_frame_at(frame_index).data.cpu() + assert frame.dtype == torch.float32 + + frame_as_uint16 = (frame * 65535).round().to(torch.uint16) + ref = asset.get_frame_data_by_index_rgb48(frame_index) + + if device == "cpu": + torch.testing.assert_close(frame_as_uint16, ref, rtol=0, atol=0) + else: + # Compare in 10-bit space (right-shift by 6) + frame_10bit = frame_as_uint16.to(torch.int32) >> 6 + ref_10bit = ref.to(torch.int32) >> 6 + assert_tensor_close_on_at_least( + frame_10bit, ref_10bit, percentage=88, atol=3 + ) def setup_frame_mappings(tmp_path, file, stream_index): json_path = tmp_path / "custom_frame_mappings.json" diff --git a/test/test_transform_ops.py b/test/test_transform_ops.py index 369d55959..bfaac4f4b 100644 --- a/test/test_transform_ops.py +++ b/test/test_transform_ops.py @@ -5,20 +5,16 @@ # LICENSE file in the root directory of this source tree. import contextlib - import json import os import subprocess import pytest - import torch import torchcodec - from torchcodec._core import get_frame_at_index, get_json_metadata from torchcodec._core.ops import _add_video_stream, add_video_stream, create_from_file from torchcodec.decoders import VideoDecoder - from torchvision.transforms import v2 from .utils import ( @@ -28,9 +24,11 @@ get_ffmpeg_minor_version, H265_VIDEO, NASA_VIDEO, + NASA_VIDEO_HDR, needs_cuda, TEST_NON_ZERO_START as NON_32_ALIGNED_WIDTH_VIDEO, TEST_SRC_2_720P, + TEST_SRC_2_720P_HDR, ) @@ -39,23 +37,54 @@ class TestPublicVideoDecoderTransformOps: "height_scaling_factor, width_scaling_factor", ((1.5, 1.31), (0.5, 0.71), (0.7, 1.31), (1.5, 0.71), (1.0, 1.0), (2.0, 2.0)), ) - @pytest.mark.parametrize("video", [NASA_VIDEO, TEST_SRC_2_720P]) + @pytest.mark.parametrize( + "video", + [ + NASA_VIDEO, + TEST_SRC_2_720P, + # TODO: On FFmpeg 4, 10-bit HDR (BT.2020 + SMPTE2084) videos produce + # different decoded results when transforms are applied. Plain 10-bit + # without HDR metadata and all 8-bit videos are unaffected. The root + # cause is unknown. + pytest.param( + NASA_VIDEO_HDR, + marks=pytest.mark.skipif( + torchcodec.ffmpeg_major_version < 5, + reason="10-bit HDR produces different results with transforms on FFmpeg 4", + ), + ), + pytest.param( + TEST_SRC_2_720P_HDR, + marks=pytest.mark.skipif( + torchcodec.ffmpeg_major_version < 5, + reason="10-bit HDR produces different results with transforms on FFmpeg 4", + ), + ), + ], + ) + @pytest.mark.parametrize("output_dtype", [None, torch.float32]) def test_resize_torchvision( - self, video, height_scaling_factor, width_scaling_factor + self, video, output_dtype, height_scaling_factor, width_scaling_factor ): height = int(video.get_height() * height_scaling_factor) width = int(video.get_width() * width_scaling_factor) + dtype_kwargs = dict(output_dtype=output_dtype) if output_dtype else {} + # We're using both the TorchCodec object and the TorchVision object to # ensure that they specify exactly the same thing. decoder_resize = VideoDecoder( - video.path, transforms=[torchcodec.transforms.Resize(size=(height, width))] + video.path, + transforms=[torchcodec.transforms.Resize(size=(height, width))], + **dtype_kwargs, ) decoder_resize_tv = VideoDecoder( - video.path, transforms=[v2.Resize(size=(height, width))] + video.path, + transforms=[v2.Resize(size=(height, width))], + **dtype_kwargs, ) - decoder_full = VideoDecoder(video.path) + decoder_full = VideoDecoder(video.path, **dtype_kwargs) num_frames = len(decoder_resize) assert num_frames == len(decoder_full) @@ -87,21 +116,67 @@ def test_resize_torchvision( assert frame_tv.shape == expected_shape assert frame_tv_no_antialias.shape == expected_shape + is_hdr = video in (NASA_VIDEO_HDR, TEST_SRC_2_720P_HDR) + + if output_dtype == torch.float32 and is_hdr: + # float32 HDR: compare in uint16 space for cleaner tolerances. + # atol=256 in uint16 space ≈ atol=1 in uint8 space. + close_pct, close_atol, max_atol = 99.5, 256, 12 * 256 + frame_resize_cmp = (frame_resize * 65535).round().to(torch.int32) + frame_tv_cmp = (frame_tv * 65535).round().to(torch.int32) + frame_tv_no_antialias_cmp = ( + (frame_tv_no_antialias * 65535).round().to(torch.int32) + ) + elif output_dtype == torch.float32: + # float32 SDR: compare in uint8 space. + close_pct, close_atol, max_atol = 99.8, 1, 6 + frame_resize_cmp = (frame_resize * 255).round().to(torch.int32) + frame_tv_cmp = (frame_tv * 255).round().to(torch.int32) + frame_tv_no_antialias_cmp = ( + (frame_tv_no_antialias * 255).round().to(torch.int32) + ) + elif is_hdr: + # 10-bit HDR content decoded to uint8 has slightly larger + # swscale vs torchvision resize diffs than native 8-bit + # content due to the 10->8 bit quantization in swscale + close_pct, close_atol, max_atol = 99.8, 1, 10 + frame_resize_cmp = frame_resize + frame_tv_cmp = frame_tv + frame_tv_no_antialias_cmp = frame_tv_no_antialias + else: + close_pct, close_atol, max_atol = 99.8, 1, 6 + frame_resize_cmp = frame_resize + frame_tv_cmp = frame_tv + frame_tv_no_antialias_cmp = frame_tv_no_antialias + assert_tensor_close_on_at_least( - frame_resize, frame_tv, percentage=99.8, atol=1 + frame_resize_cmp, frame_tv_cmp, percentage=close_pct, atol=close_atol + ) + torch.testing.assert_close( + frame_resize_cmp, frame_tv_cmp, rtol=0, atol=max_atol ) - torch.testing.assert_close(frame_resize, frame_tv, rtol=0, atol=6) if height_scaling_factor < 1 or width_scaling_factor < 1: # Antialias only relevant when down-scaling! - with pytest.raises(AssertionError, match="Expected at least"): - assert_tensor_close_on_at_least( - frame_resize, frame_tv_no_antialias, percentage=99, atol=1 - ) - with pytest.raises(AssertionError, match="Tensor-likes are not close"): - torch.testing.assert_close( - frame_resize, frame_tv_no_antialias, rtol=0, atol=6 - ) + # For HDR content with mild downscale factors, the + # antialias difference may be too small to detect. + if not is_hdr: + with pytest.raises(AssertionError, match="Expected at least"): + assert_tensor_close_on_at_least( + frame_resize_cmp, + frame_tv_no_antialias_cmp, + percentage=99, + atol=close_atol, + ) + with pytest.raises( + AssertionError, match="Tensor-likes are not close" + ): + torch.testing.assert_close( + frame_resize_cmp, + frame_tv_no_antialias_cmp, + rtol=0, + atol=max_atol, + ) def test_resize_fails(self): with pytest.raises( @@ -161,25 +236,56 @@ def test_resize_non_32_aligned_input_width(self): "height_scaling_factor, width_scaling_factor", ((0.5, 0.5), (0.25, 0.1), (1.0, 1.0), (0.15, 0.75)), ) - @pytest.mark.parametrize("video", [NASA_VIDEO, TEST_SRC_2_720P]) + @pytest.mark.parametrize( + "video", + [ + NASA_VIDEO, + TEST_SRC_2_720P, + # TODO: On FFmpeg 4, 10-bit HDR (BT.2020 + SMPTE2084) videos produce + # different decoded results when transforms are applied. Plain 10-bit + # without HDR metadata and all 8-bit videos are unaffected. The root + # cause is unknown. + pytest.param( + NASA_VIDEO_HDR, + marks=pytest.mark.skipif( + torchcodec.ffmpeg_major_version < 5, + reason="10-bit HDR produces different results with transforms on FFmpeg 4", + ), + ), + pytest.param( + TEST_SRC_2_720P_HDR, + marks=pytest.mark.skipif( + torchcodec.ffmpeg_major_version < 5, + reason="10-bit HDR produces different results with transforms on FFmpeg 4", + ), + ), + ], + ) + @pytest.mark.parametrize("output_dtype", [None, torch.float32]) def test_center_crop_torchvision( self, height_scaling_factor, width_scaling_factor, video, + output_dtype, ): height = int(video.get_height() * height_scaling_factor) width = int(video.get_width() * width_scaling_factor) + dtype_kwargs = dict(output_dtype=output_dtype) if output_dtype else {} + tc_center_crop = torchcodec.transforms.CenterCrop(size=(height, width)) - decoder_center_crop = VideoDecoder(video.path, transforms=[tc_center_crop]) + decoder_center_crop = VideoDecoder( + video.path, transforms=[tc_center_crop], **dtype_kwargs + ) decoder_center_crop_tv = VideoDecoder( video.path, transforms=[v2.CenterCrop(size=(height, width))], + **dtype_kwargs, ) - decoder_full = VideoDecoder(video.path) + decoder_full = VideoDecoder(video.path, **dtype_kwargs) num_frames = len(decoder_center_crop_tv) assert num_frames == len(decoder_full) @@ -216,25 +322,55 @@ def test_center_crop_fails(self): "height_scaling_factor, width_scaling_factor", ((0.5, 0.5), (0.25, 0.1), (1.0, 1.0), (0.15, 0.75)), ) - @pytest.mark.parametrize("video", [NASA_VIDEO, TEST_SRC_2_720P]) + @pytest.mark.parametrize( + "video", + [ + NASA_VIDEO, + TEST_SRC_2_720P, + # TODO: On FFmpeg 4, 10-bit HDR (BT.2020 + SMPTE2084) videos produce + # different decoded results when transforms are applied. Plain 10-bit + # without HDR metadata and all 8-bit videos are unaffected. The root + # cause is unknown. + pytest.param( + NASA_VIDEO_HDR, + marks=pytest.mark.skipif( + torchcodec.ffmpeg_major_version < 5, + reason="10-bit HDR produces different results with transforms on FFmpeg 4", + ), + ), + pytest.param( + TEST_SRC_2_720P_HDR, + marks=pytest.mark.skipif( + torchcodec.ffmpeg_major_version < 5, + reason="10-bit HDR produces different results with transforms on FFmpeg 4", + ), + ), + ], + ) @pytest.mark.parametrize("seed", [0, 1234]) + @pytest.mark.parametrize("output_dtype", [None, torch.float32]) def test_random_crop_torchvision( self, height_scaling_factor, width_scaling_factor, video, seed, + output_dtype, ): height = int(video.get_height() * height_scaling_factor) width = int(video.get_width() * width_scaling_factor) + dtype_kwargs = dict(output_dtype=output_dtype) if output_dtype else {} + # We want both kinds of RandomCrop objects to get arrive at the same # locations to crop, so we need to make sure they get the same random # seed. It's used in RandomCrop's _make_transform_spec() method, called # by the VideoDecoder. torch.manual_seed(seed) tc_random_crop = torchcodec.transforms.RandomCrop(size=(height, width)) - decoder_random_crop = VideoDecoder(video.path, transforms=[tc_random_crop]) + decoder_random_crop = VideoDecoder( + video.path, transforms=[tc_random_crop], **dtype_kwargs + ) # Resetting manual seed for when TorchCodec's RandomCrop, created from # the TorchVision RandomCrop, is used inside of the VideoDecoder. It @@ -243,9 +379,10 @@ def test_random_crop_torchvision( decoder_random_crop_tv = VideoDecoder( video.path, transforms=[v2.RandomCrop(size=(height, width))], + **dtype_kwargs, ) - decoder_full = VideoDecoder(video.path) + decoder_full = VideoDecoder(video.path, **dtype_kwargs) num_frames = len(decoder_random_crop_tv) assert num_frames == len(decoder_full) @@ -353,15 +490,21 @@ def test_random_crop_reusable_objects(self, seed): (v2.Resize, v2.RandomCrop), ], ) - def test_transform_pipeline(self, resize, random_crop): + @pytest.mark.parametrize( + "video", [TEST_SRC_2_720P, NASA_VIDEO_HDR, TEST_SRC_2_720P_HDR] + ) + @pytest.mark.parametrize("output_dtype", [None, torch.float32]) + def test_transform_pipeline(self, resize, random_crop, video, output_dtype): + dtype_kwargs = dict(output_dtype=output_dtype) if output_dtype else {} decoder = VideoDecoder( - TEST_SRC_2_720P.path, + video.path, transforms=[ # resized to bigger than original resize(size=(2160, 3840)), # crop to smaller than the resize, but still bigger than original random_crop(size=(1080, 1920)), ], + **dtype_kwargs, ) num_frames = len(decoder) @@ -373,7 +516,7 @@ def test_transform_pipeline(self, resize, random_crop): num_frames - 1, ]: frame = decoder[frame_index] - assert frame.shape == (TEST_SRC_2_720P.get_num_color_channels(), 1080, 1920) + assert frame.shape == (video.get_num_color_channels(), 1080, 1920) def test_transform_fails(self): with pytest.raises( @@ -393,7 +536,14 @@ def get_num_frames_core_ops(self, video): assert num_frames is not None return num_frames - @pytest.mark.parametrize("video", [NASA_VIDEO, H265_VIDEO, AV1_VIDEO]) + @pytest.mark.parametrize( + "video", + [ + NASA_VIDEO, + H265_VIDEO, + AV1_VIDEO, + ], + ) def test_color_conversion_library(self, video): num_frames = self.get_num_frames_core_ops(video) diff --git a/test/utils.py b/test/utils.py index 1d95c3d7c..ccc3cce8a 100644 --- a/test/utils.py +++ b/test/utils.py @@ -406,6 +406,19 @@ def get_frame_data_by_index( tensor_file_path = f"{base_path}.pt" return torch.load(tensor_file_path, weights_only=True).permute(2, 0, 1) + def get_frame_data_by_index_rgb48( + self, + idx: int, + *, + stream_index: int | None = None, + ) -> torch.Tensor: + if stream_index is None: + stream_index = self.default_stream_index + + base_path = self.get_base_path_by_index(idx, stream_index=stream_index) + tensor_file_path = f"{base_path}.rgb48.pt" + return torch.load(tensor_file_path, weights_only=True).permute(2, 0, 1) + def get_frame_data_by_range( self, start: int, @@ -626,6 +639,46 @@ def get_empty_chw_tensor(self, *, stream_index: int) -> torch.Tensor: frames={0: {}}, # Not needed for now ) +# HDR re-encode of NASA video (10-bit H265 with BT.2020 + PQ), generated with: +# ffmpeg -i test/resources/nasa_13013.mp4 -map 0:v:0 -c:v libx265 -pix_fmt yuv420p10le \ +# -x265-params "colorprim=bt2020:transfer=smpte2084:colormatrix=bt2020nc:range=limited" \ +# -preset fast -crf 23 test/resources/nasa_13013_hdr.mp4 +NASA_VIDEO_HDR = TestVideo( + filename="nasa_13013_hdr.mp4", + default_stream_index=0, + stream_infos={ + 0: TestVideoStreamInfo(width=320, height=180, num_color_channels=3), + }, + frames={0: {}}, # Not needed for now +) + +# HDR re-encode of testsrc2 (10-bit H265 with BT.2020 + PQ), generated with: +# ffmpeg -i test/resources/testsrc2.mp4 -c:v libx265 -pix_fmt yuv420p10le \ +# -x265-params "colorprim=bt2020:transfer=smpte2084:colormatrix=bt2020nc:range=limited" \ +# -preset fast -crf 23 test/resources/testsrc2_hdr.mp4 +TEST_SRC_2_720P_HDR = TestVideo( + filename="testsrc2_hdr.mp4", + default_stream_index=0, + stream_infos={ + 0: TestVideoStreamInfo(width=320, height=180, num_color_channels=3), + }, + frames={0: {}}, # Not needed for now +) + +# 12-bit HDR testsrc2 (H265 with BT.2020 + PQ), generated with: +# ffmpeg -f lavfi -i testsrc2=duration=2:size=320x180:rate=30 -c:v libx265 +# -pix_fmt yuv420p12le -x265-params +# "colorprim=bt2020:transfer=smpte2084:colormatrix=bt2020nc:range=limited" +# -preset fast -crf 23 test/resources/testsrc2_12bit_hdr.mp4 +TEST_SRC_2_12BIT_HDR = TestVideo( + filename="testsrc2_12bit_hdr.mp4", + default_stream_index=0, + stream_infos={ + 0: TestVideoStreamInfo(width=320, height=180, num_color_channels=3), + }, + frames={0: {}}, +) + # ffmpeg -f lavfi -i testsrc2=duration=2:size=1280x720:rate=30 -c:v libx264 -profile:v baseline -level 3.1 -pix_fmt yuv420p -b:v 2500k -r 30 -movflags +faststart output_720p_2s.mp4 TEST_SRC_2_720P = TestVideo( filename="testsrc2.mp4",