PR #24727 opened by rnissl URL: https://code.ffmpeg.org/FFmpeg/FFmpeg/pulls/24727 Patch URL: https://code.ffmpeg.org/FFmpeg/FFmpeg/pulls/24727.patch
This is 4/9 in the GoPro Max review order. It applies on current master by itself. Runtime dlopen of nvJPEG (13 then 12 then unversioned). Depends on ffnvcodec and cuda-nvcc or cuda-llvm. No --enable-nonfree. Preferred review order, each request on its own: 1/9 avutil/log: flush stderr and match Windows warning yellow 2/9 avfilter: add yuvrgb_cuda 3/9 avfilter: add gopromax2_cuda GoPro Max 2 EAC stitch 4/9 avcodec: add mjpeg_nvjpeg NVIDIA JPEG encoder 5/9 avcodec/bsf: add jpegapp to inject JPEG APP segments 6/9 avfilter: add gopromax2 GoPro Max 2 EAC stitch 7/9 avfilter: add gopromax2_opencl GoPro Max 2 EAC stitch 8/9 avfilter: add gopromax_opencl GoPro Max 1 stitch 9/9 avcodec/mjpegenc: encode packed RGB via YUV conversion After 1-5, mapillary_tools processes GoPro Max 2 files GPU-only in my setup: scale_cuda=yuv444p10, yuvrgb_cuda=rgb0, gopromax2_cuda, mjpeg_nvjpeg. jpegapp inserts the APP1 segment while encoding, so the JPEG ffmpeg writes needs no postprocessing. After 6, the Max 2 CPU fallback works. After 7, Max 2 via OpenCL works. After 8, original GoPro Max (Max 1) via OpenCL works. 9/9 is needed when 4/9 cannot be used, for example because no nvjpeg library is present, and encoding falls back to the software mjpeg encoder. That encoder then receives packed RGB (rgb24). Upstream mjpegenc does not accept it; this change converts it to YUV first. The six patches I sent to ffmpeg-devel on 20 September are obsolete. Once the filter requests are merged, I will send one short documentation request that only adds the cross-references between these sections. From 9f750a9188d5f7bae64723524c1c6da4ea91ae2c Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Reinhard=20Ni=C3=9Fl?= <[email protected]> Date: Sat, 26 Sep 2026 21:26:00 +0200 Subject: [PATCH] avcodec: add mjpeg_nvjpeg NVIDIA JPEG encoder MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Runtime dlopen of nvJPEG (13 then 12 then unversioned). Depends on ffnvcodec and cuda-nvcc or cuda-llvm. No --enable-nonfree. Co-authored-by: Cursor <[email protected]> Signed-off-by: Reinhard Nißl <[email protected]> Co-authored-by: Cursor <[email protected]> --- Changelog | 3 + configure | 2 + doc/encoders.texi | 35 +++ libavcodec/Makefile | 1 + libavcodec/allcodecs.c | 1 + libavcodec/nvjpegenc.c | 584 ++++++++++++++++++++++++++++++++++++++++ libavcodec/nvjpegenc.cu | 61 +++++ 7 files changed, 687 insertions(+) create mode 100644 libavcodec/nvjpegenc.c create mode 100644 libavcodec/nvjpegenc.cu diff --git a/Changelog b/Changelog index 1cfa9db1b6..dd08217835 100644 --- a/Changelog +++ b/Changelog @@ -15,6 +15,9 @@ version <next>: - H.264 data partitioning support - DSD (dsd_msbf) encoder +- jpegapp bitstream filter +- MJPEG encoder packed RGB input +- mjpeg_nvjpeg NVIDIA JPEG encoder version 9.0: - Extend AMF Color Converter (vf_vpp_amf) HDR capabilities diff --git a/configure b/configure index f686b727cf..564f8d371e 100755 --- a/configure +++ b/configure @@ -3659,6 +3659,8 @@ mjpeg_qsv_encoder_deps="libmfx" mjpeg_qsv_encoder_select="qsvenc" mjpeg_vaapi_encoder_deps="VAEncPictureParameterBufferJPEG" mjpeg_vaapi_encoder_select="cbs_jpeg jpegtables vaapi_encode" +mjpeg_nvjpeg_encoder_deps="ffnvcodec" +mjpeg_nvjpeg_encoder_deps_any="cuda_nvcc cuda_llvm" mp3_mf_encoder_deps="mediafoundation" mp3_mediacodec_decoder_deps="mediacodec" mp3_mediacodec_decoder_select="mpegaudioheader" diff --git a/doc/encoders.texi b/doc/encoders.texi index a3ce15fcff..d49c6d5393 100644 --- a/doc/encoders.texi +++ b/doc/encoders.texi @@ -1268,6 +1268,41 @@ Compute and use optimal huffman tables. @end table @end table +@example +-c:v mjpeg -quality 95 -matrix jpeg -chroma 444 +@end example + +@anchor{mjpeg_nvjpeg} +@section mjpeg_nvjpeg + +NVIDIA nvJPEG encoder. Encodes CUDA frames in place (no hwdownload). +The library is loaded at runtime: @file{nvjpeg64_13} then @file{nvjpeg64_12} +then unversioned @file{nvjpeg}, paired with the matching @file{cudart64_*} +major. Those DLLs come from the CUDA Toolkit, not the display driver. + +Input is @code{AV_PIX_FMT_CUDA} with software format @code{rgb0} or 8-bit +@code{yuv444p}. Packed @code{rgb0} is converted to planar YCbCr on the GPU +(ITU T.871 by default). 10-bit YUV is not supported. + +@table @option +@item quality +JPEG quality @code{1}--@code{100}. Default 26. + +@item chroma +@samp{444}, @samp{422}, or @samp{420} (default) for RGB input. + +@item matrix +@samp{jpeg} (default, BT.601 full range) or @samp{bt709}. + +@item huffman +Optimized Huffman tables (default 1). +@end table + +@example +-c:v mjpeg_nvjpeg -quality 83 -matrix jpeg -chroma 420 +@end example + + @anchor{wavpackenc} @section wavpack diff --git a/libavcodec/Makefile b/libavcodec/Makefile index e2c60e9051..06eb295ae3 100644 --- a/libavcodec/Makefile +++ b/libavcodec/Makefile @@ -542,6 +542,7 @@ OBJS-$(CONFIG_MJPEGB_DECODER) += mjpegbdec.o OBJS-$(CONFIG_MJPEG_CUVID_DECODER) += cuviddec.o OBJS-$(CONFIG_MJPEG_QSV_ENCODER) += qsvenc_jpeg.o OBJS-$(CONFIG_MJPEG_VAAPI_ENCODER) += vaapi_encode_mjpeg.o +OBJS-$(CONFIG_MJPEG_NVJPEG_ENCODER) += nvjpegenc.o nvjpegenc.ptx.o OBJS-$(CONFIG_MLP_DECODER) += mlpdec.o mlpdsp.o OBJS-$(CONFIG_MLP_ENCODER) += mlpenc.o mlp.o OBJS-$(CONFIG_MMVIDEO_DECODER) += mmvideo.o diff --git a/libavcodec/allcodecs.c b/libavcodec/allcodecs.c index ca20c76fe6..2f454ca6ac 100644 --- a/libavcodec/allcodecs.c +++ b/libavcodec/allcodecs.c @@ -899,6 +899,7 @@ extern const FFCodec ff_mjpeg_cuvid_decoder; extern const FFCodec ff_mjpeg_qsv_encoder; extern const FFCodec ff_mjpeg_qsv_decoder; extern const FFCodec ff_mjpeg_vaapi_encoder; +extern const FFCodec ff_mjpeg_nvjpeg_encoder; extern const FFCodec ff_mp3_mediacodec_decoder; extern const FFCodec ff_mp3_mf_encoder; extern const FFCodec ff_mpeg1_cuvid_decoder; diff --git a/libavcodec/nvjpegenc.c b/libavcodec/nvjpegenc.c new file mode 100644 index 0000000000..fd8509b1d9 --- /dev/null +++ b/libavcodec/nvjpegenc.c @@ -0,0 +1,584 @@ +/* + * NVIDIA nvJPEG encoder (CUDA). + * + * Copyright (c) 2026 Reinhard Nißl + * + * This file is part of FFmpeg. + * + * FFmpeg is free software; you can redistribute it and/or + * modify it under the terms of the GNU Lesser General Public + * License as published by the Free Software Foundation; either + * version 2.1 of the License, or (at your option) any later version. + * + * FFmpeg is distributed in the hope that it will be useful, + * but WITHOUT ANY WARRANTY; without even the implied warranty of + * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU + * Lesser General Public License for more details. + * + * You should have received a copy of the GNU Lesser General Public + * License along with FFmpeg; if not, write to the Free Software + * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA + */ + +#include "config.h" + +#include <stdint.h> +#include <string.h> + +#include "libavutil/common.h" +#include "libavutil/cuda_check.h" +#include "libavutil/hwcontext.h" +#include "libavutil/hwcontext_cuda_internal.h" +#include "libavutil/internal.h" +#include "libavutil/mem.h" +#include "libavutil/opt.h" +#include "libavutil/pixdesc.h" + +#include "avcodec.h" +#include "codec_internal.h" +#include "encode.h" +#include "hwconfig.h" + +#include "compat/w32dlfcn.h" + +#define CHECK_CU(x) FF_CUDA_CHECK_DL(avctx, s->hwctx->internal->cuda_dl, x) +#define DIV_UP(a, b) (((a) + (b) - 1) / (b)) +#define BLOCKX 32 +#define BLOCKY 16 + +enum { + NVJPEG_STATUS_SUCCESS = 0, +}; + +enum { + NVJPEG_CSS_444 = 0, + NVJPEG_CSS_422 = 1, + NVJPEG_CSS_420 = 2, +}; + +enum { + NVJPEG_ENCODING_BASELINE_DCT = 0xc0, +}; + +enum { + NVJPEG_MATRIX_JPEG = 0, + NVJPEG_MATRIX_BT709 = 1, +}; + +typedef struct NVJPEGImage { + unsigned char *channel[4]; + size_t pitch[4]; +} NVJPEGImage; + +typedef void *nvjpegHandle; +typedef void *nvjpegEncoderState; +typedef void *nvjpegEncoderParams; +typedef int nvjpegStatus; + +typedef nvjpegStatus (*pfn_nvjpegCreateSimple)(nvjpegHandle *); +typedef nvjpegStatus (*pfn_nvjpegDestroy)(nvjpegHandle); +typedef nvjpegStatus (*pfn_nvjpegEncoderStateCreate)(nvjpegHandle, nvjpegEncoderState *, void *); +typedef nvjpegStatus (*pfn_nvjpegEncoderStateDestroy)(nvjpegEncoderState); +typedef nvjpegStatus (*pfn_nvjpegEncoderParamsCreate)(nvjpegHandle, nvjpegEncoderParams *, void *); +typedef nvjpegStatus (*pfn_nvjpegEncoderParamsDestroy)(nvjpegEncoderParams); +typedef nvjpegStatus (*pfn_nvjpegEncoderParamsSetQuality)(nvjpegEncoderParams, const int, void *); +typedef nvjpegStatus (*pfn_nvjpegEncoderParamsSetEncoding)(nvjpegEncoderParams, int, void *); +typedef nvjpegStatus (*pfn_nvjpegEncoderParamsSetOptimizedHuffman)(nvjpegEncoderParams, const int, void *); +typedef nvjpegStatus (*pfn_nvjpegEncoderParamsSetSamplingFactors)(nvjpegEncoderParams, int, void *); +typedef nvjpegStatus (*pfn_nvjpegEncodeYUV)(nvjpegHandle, nvjpegEncoderState, const nvjpegEncoderParams, + const NVJPEGImage *, int, int, int, void *); +typedef nvjpegStatus (*pfn_nvjpegEncodeRetrieveBitstream)(nvjpegHandle, nvjpegEncoderState, + unsigned char *, size_t *, void *); + +typedef struct NVJPEGEncContext { + const AVClass *class; + + AVCUDADeviceContext *hwctx; + CUstream cu_stream; + CUmodule cu_module; + CUfunction cu_rgb0_to_yuv444; + CUdeviceptr yuv; + size_t yuv_size; + int yuv_pitch; + + void *lib; + nvjpegHandle handle; + nvjpegEncoderState state; + nvjpegEncoderParams params; + pfn_nvjpegCreateSimple CreateSimple; + pfn_nvjpegDestroy Destroy; + pfn_nvjpegEncoderStateCreate EncoderStateCreate; + pfn_nvjpegEncoderStateDestroy EncoderStateDestroy; + pfn_nvjpegEncoderParamsCreate EncoderParamsCreate; + pfn_nvjpegEncoderParamsDestroy EncoderParamsDestroy; + pfn_nvjpegEncoderParamsSetQuality EncoderParamsSetQuality; + pfn_nvjpegEncoderParamsSetEncoding EncoderParamsSetEncoding; + pfn_nvjpegEncoderParamsSetOptimizedHuffman EncoderParamsSetOptimizedHuffman; + pfn_nvjpegEncoderParamsSetSamplingFactors EncoderParamsSetSamplingFactors; + pfn_nvjpegEncodeYUV EncodeYUV; + pfn_nvjpegEncodeRetrieveBitstream EncodeRetrieveBitstream; + + int quality; + int chroma; + int huffman; + int matrix; + enum AVPixelFormat sw_format; +} NVJPEGEncContext; + +static const AVCodecHWConfigInternal *const nvjpeg_hw_configs[] = { + HW_CONFIG_ENCODER_FRAMES(CUDA, CUDA), + HW_CONFIG_ENCODER_DEVICE(NONE, CUDA), + NULL +}; + +#define CHECK_NVJPEG(avctx, call) do { \ + nvjpegStatus st = (call); \ + if (st != NVJPEG_STATUS_SUCCESS) { \ + av_log(avctx, AV_LOG_ERROR, "%s failed: %d\n", #call, st); \ + return AVERROR_EXTERNAL; \ + } \ +} while (0) + +static void *nvjpeg_dlopen(void) +{ +#ifdef _WIN32 + void *lib = dlopen("nvjpeg64_13.dll", 0); + if (lib) + return lib; + lib = dlopen("nvjpeg64_12.dll", 0); + if (lib) + return lib; + return dlopen("nvjpeg.dll", 0); +#else + void *lib = dlopen("libnvjpeg.so.13", 0); + if (lib) + return lib; + lib = dlopen("libnvjpeg.so.12", 0); + if (lib) + return lib; + return dlopen("libnvjpeg.so", 0); +#endif +} + +static av_cold int nvjpeg_load_dll(AVCodecContext *avctx) +{ + NVJPEGEncContext *s = avctx->priv_data; + + s->lib = nvjpeg_dlopen(); + if (!s->lib) { + av_log(avctx, AV_LOG_ERROR, + "Cannot load nvjpeg64_13.dll or nvjpeg64_12.dll (place it next to ffmpeg.exe)\n"); + return AVERROR(ENOSYS); + } + +#define LOAD(field, name) do { \ + s->field = (void *)dlsym(s->lib, name); \ + if (!s->field) { \ + av_log(avctx, AV_LOG_ERROR, "Missing %s\n", name); \ + return AVERROR(ENOSYS); \ + } \ +} while (0) + + LOAD(CreateSimple, "nvjpegCreateSimple"); + LOAD(Destroy, "nvjpegDestroy"); + LOAD(EncoderStateCreate, "nvjpegEncoderStateCreate"); + LOAD(EncoderStateDestroy, "nvjpegEncoderStateDestroy"); + LOAD(EncoderParamsCreate, "nvjpegEncoderParamsCreate"); + LOAD(EncoderParamsDestroy, "nvjpegEncoderParamsDestroy"); + LOAD(EncoderParamsSetQuality, "nvjpegEncoderParamsSetQuality"); + LOAD(EncoderParamsSetEncoding, "nvjpegEncoderParamsSetEncoding"); + LOAD(EncoderParamsSetOptimizedHuffman, "nvjpegEncoderParamsSetOptimizedHuffman"); + LOAD(EncoderParamsSetSamplingFactors, "nvjpegEncoderParamsSetSamplingFactors"); + LOAD(EncodeYUV, "nvjpegEncodeYUV"); + LOAD(EncodeRetrieveBitstream, "nvjpegEncodeRetrieveBitstream"); +#undef LOAD + return 0; +} + +static int nvjpeg_css(int chroma) +{ + switch (chroma) { + case NVJPEG_CSS_420: + return NVJPEG_CSS_420; + case NVJPEG_CSS_422: + return NVJPEG_CSS_422; + default: + return NVJPEG_CSS_444; + } +} + +static av_cold int nvjpeg_load_module(AVCodecContext *avctx) +{ + NVJPEGEncContext *s = avctx->priv_data; + CudaFunctions *cu = s->hwctx->internal->cuda_dl; + CUcontext dummy; + int ret, pushed = 0; + + extern const unsigned char ff_nvjpegenc_ptx_data[]; + extern const unsigned int ff_nvjpegenc_ptx_len; + + ret = CHECK_CU(cu->cuCtxPushCurrent(s->hwctx->cuda_ctx)); + if (ret < 0) + return ret; + pushed = 1; + + ret = CHECK_CU(cu->cuModuleLoadData(&s->cu_module, ff_nvjpegenc_ptx_data)); + if (ret < 0) + goto fail; + ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_rgb0_to_yuv444, s->cu_module, + "rgb0_to_yuv444")); +fail: + if (pushed) + CHECK_CU(cu->cuCtxPopCurrent(&dummy)); + return ret; +} + +static av_cold int nvjpeg_encode_init(AVCodecContext *avctx) +{ + NVJPEGEncContext *s = avctx->priv_data; + AVHWFramesContext *frames_ctx = NULL; + CudaFunctions *cu; + CUcontext dummy; + void *stream; + int ret, css, pushed = 0; + + ret = nvjpeg_load_dll(avctx); + if (ret < 0) + return ret; + + if (avctx->hw_frames_ctx) { + frames_ctx = (AVHWFramesContext *)avctx->hw_frames_ctx->data; + if (frames_ctx->format != AV_PIX_FMT_CUDA) { + av_log(avctx, AV_LOG_ERROR, "hw_frames_ctx must be CUDA\n"); + return AVERROR(EINVAL); + } + s->hwctx = frames_ctx->device_ctx->hwctx; + s->sw_format = frames_ctx->sw_format; + } else if (avctx->hw_device_ctx) { + AVHWDeviceContext *dev = (AVHWDeviceContext *)avctx->hw_device_ctx->data; + if (dev->type != AV_HWDEVICE_TYPE_CUDA) { + av_log(avctx, AV_LOG_ERROR, "hw_device_ctx must be CUDA\n"); + return AVERROR(EINVAL); + } + s->hwctx = dev->hwctx; + s->sw_format = AV_PIX_FMT_NONE; + } else { + av_log(avctx, AV_LOG_ERROR, "CUDA frames or device context required\n"); + return AVERROR(EINVAL); + } + + s->cu_stream = s->hwctx->stream; + stream = (void *)(uintptr_t)s->cu_stream; + + if (s->sw_format != AV_PIX_FMT_NONE && + s->sw_format != AV_PIX_FMT_RGB0 && + s->sw_format != AV_PIX_FMT_YUV444P) { + av_log(avctx, AV_LOG_ERROR, + "nvJPEG needs CUDA rgb0 or yuv444p (got %s)\n", + av_get_pix_fmt_name(s->sw_format)); + return AVERROR(ENOSYS); + } + + ret = nvjpeg_load_module(avctx); + if (ret < 0) + return ret; + + cu = s->hwctx->internal->cuda_dl; + ret = CHECK_CU(cu->cuCtxPushCurrent(s->hwctx->cuda_ctx)); + if (ret < 0) + return ret; + pushed = 1; + + if (s->CreateSimple(&s->handle) != NVJPEG_STATUS_SUCCESS || + s->EncoderStateCreate(s->handle, &s->state, stream) != NVJPEG_STATUS_SUCCESS || + s->EncoderParamsCreate(s->handle, &s->params, stream) != NVJPEG_STATUS_SUCCESS) { + av_log(avctx, AV_LOG_ERROR, "nvjpeg encoder create failed\n"); + ret = AVERROR_EXTERNAL; + goto fail; + } + css = nvjpeg_css(s->chroma); + if (s->EncoderParamsSetQuality(s->params, s->quality, stream) != NVJPEG_STATUS_SUCCESS || + s->EncoderParamsSetEncoding(s->params, NVJPEG_ENCODING_BASELINE_DCT, stream) != NVJPEG_STATUS_SUCCESS || + s->EncoderParamsSetOptimizedHuffman(s->params, s->huffman, stream) != NVJPEG_STATUS_SUCCESS || + s->EncoderParamsSetSamplingFactors(s->params, css, stream) != NVJPEG_STATUS_SUCCESS) { + av_log(avctx, AV_LOG_ERROR, "nvjpeg encoder params failed\n"); + ret = AVERROR_EXTERNAL; + goto fail; + } + + ret = CHECK_CU(cu->cuCtxPopCurrent(&dummy)); + pushed = 0; + if (ret < 0) + return ret; + + avctx->color_range = AVCOL_RANGE_JPEG; + av_log(avctx, AV_LOG_VERBOSE, "mjpeg_nvjpeg quality=%d chroma=%d matrix=%s huffman=%d\n", + s->quality, css, s->matrix == NVJPEG_MATRIX_BT709 ? "bt709" : "jpeg", + s->huffman); + return 0; + +fail: + if (pushed) + cu->cuCtxPopCurrent(&dummy); + return ret; +} + +static av_cold int nvjpeg_encode_close(AVCodecContext *avctx) +{ + NVJPEGEncContext *s = avctx->priv_data; + CUcontext dummy; + + if (s->hwctx && s->hwctx->internal && s->hwctx->internal->cuda_dl) { + CudaFunctions *cu = s->hwctx->internal->cuda_dl; + int pushed = 0; + if (s->hwctx->cuda_ctx && + cu->cuCtxPushCurrent(s->hwctx->cuda_ctx) == CUDA_SUCCESS) + pushed = 1; + if (s->yuv) + cu->cuMemFree(s->yuv); + s->yuv = 0; + if (s->cu_module) + cu->cuModuleUnload(s->cu_module); + s->cu_module = NULL; + if (pushed) + cu->cuCtxPopCurrent(&dummy); + } + + if (s->params && s->EncoderParamsDestroy) + s->EncoderParamsDestroy(s->params); + s->params = NULL; + if (s->state && s->EncoderStateDestroy) + s->EncoderStateDestroy(s->state); + s->state = NULL; + if (s->handle && s->Destroy) + s->Destroy(s->handle); + s->handle = NULL; + if (s->lib) + dlclose(s->lib); + s->lib = NULL; + return 0; +} + +static int nvjpeg_ensure_yuv(AVCodecContext *avctx, int width, int height) +{ + NVJPEGEncContext *s = avctx->priv_data; + CudaFunctions *cu = s->hwctx->internal->cuda_dl; + size_t need = (size_t)width * height * 3; + int ret; + + if (s->yuv && s->yuv_size >= need) + return 0; + + if (s->yuv) { + ret = CHECK_CU(cu->cuMemFree(s->yuv)); + if (ret < 0) + return ret; + s->yuv = 0; + } + ret = CHECK_CU(cu->cuMemAlloc(&s->yuv, need)); + if (ret < 0) + return ret; + s->yuv_size = need; + s->yuv_pitch = width; + return 0; +} + +static int nvjpeg_rgb0_to_yuv444(AVCodecContext *avctx, const AVFrame *frame) +{ + NVJPEGEncContext *s = avctx->priv_data; + CudaFunctions *cu = s->hwctx->internal->cuda_dl; + CUdeviceptr src = (CUdeviceptr)frame->data[0]; + CUdeviceptr y, u, v; + int spitch = frame->linesize[0]; + int w = frame->width, h = frame->height; + int yp, up, vp, matrix; + void *args[11]; + int ret; + + ret = nvjpeg_ensure_yuv(avctx, w, h); + if (ret < 0) + return ret; + + y = s->yuv; + u = s->yuv + (size_t)w * h; + v = s->yuv + 2 * (size_t)w * h; + yp = up = vp = s->yuv_pitch; + matrix = s->matrix; + args[0] = &src; + args[1] = &spitch; + args[2] = &y; + args[3] = &u; + args[4] = &v; + args[5] = &yp; + args[6] = &up; + args[7] = &vp; + args[8] = &w; + args[9] = &h; + args[10] = &matrix; + return CHECK_CU(cu->cuLaunchKernel(s->cu_rgb0_to_yuv444, + DIV_UP(w, BLOCKX), DIV_UP(h, BLOCKY), 1, + BLOCKX, BLOCKY, 1, 0, s->cu_stream, args, NULL)); +} + +static int nvjpeg_encode_frame(AVCodecContext *avctx, AVPacket *pkt, + const AVFrame *frame, int *got_packet) +{ + NVJPEGEncContext *s = avctx->priv_data; + CudaFunctions *cu; + CUcontext dummy; + NVJPEGImage img = { 0 }; + void *stream; + size_t length = 0; + int ret, pushed = 0, st; + enum AVPixelFormat sw; + + if (!frame) + return 0; + + if (!frame->hw_frames_ctx) { + av_log(avctx, AV_LOG_ERROR, "CUDA frame required\n"); + return AVERROR(EINVAL); + } + + { + AVHWFramesContext *fc = (AVHWFramesContext *)frame->hw_frames_ctx->data; + sw = fc->sw_format; + if (!s->hwctx) + s->hwctx = fc->device_ctx->hwctx; + if (s->sw_format == AV_PIX_FMT_NONE) + s->sw_format = sw; + } + + if (sw != AV_PIX_FMT_RGB0 && sw != AV_PIX_FMT_YUV444P) { + av_log(avctx, AV_LOG_ERROR, "Unsupported CUDA sw_format %s\n", + av_get_pix_fmt_name(sw)); + return AVERROR(ENOSYS); + } + + cu = s->hwctx->internal->cuda_dl; + stream = (void *)(uintptr_t)s->cu_stream; + + ret = CHECK_CU(cu->cuCtxPushCurrent(s->hwctx->cuda_ctx)); + if (ret < 0) + return ret; + pushed = 1; + + if (sw == AV_PIX_FMT_RGB0) { + CUdeviceptr y, u, v; + size_t plane; + ret = nvjpeg_rgb0_to_yuv444(avctx, frame); + if (ret < 0) + goto fail; + plane = (size_t)frame->width * frame->height; + y = s->yuv; + u = s->yuv + plane; + v = s->yuv + 2 * plane; + img.channel[0] = (unsigned char *)y; + img.channel[1] = (unsigned char *)u; + img.channel[2] = (unsigned char *)v; + img.pitch[0] = s->yuv_pitch; + img.pitch[1] = s->yuv_pitch; + img.pitch[2] = s->yuv_pitch; + } else { + img.channel[0] = frame->data[0]; + img.channel[1] = frame->data[1]; + img.channel[2] = frame->data[2]; + img.pitch[0] = frame->linesize[0]; + img.pitch[1] = frame->linesize[1]; + img.pitch[2] = frame->linesize[2]; + } + st = s->EncodeYUV(s->handle, s->state, s->params, &img, + NVJPEG_CSS_444, frame->width, frame->height, stream); + if (st != NVJPEG_STATUS_SUCCESS) { + av_log(avctx, AV_LOG_ERROR, "nvjpeg encode failed: %d\n", st); + ret = AVERROR_EXTERNAL; + goto fail; + } + + st = s->EncodeRetrieveBitstream(s->handle, s->state, NULL, &length, stream); + if (st != NVJPEG_STATUS_SUCCESS) { + av_log(avctx, AV_LOG_ERROR, "nvjpeg size query failed: %d\n", st); + ret = AVERROR_EXTERNAL; + goto fail; + } + ret = CHECK_CU(cu->cuStreamSynchronize(s->cu_stream)); + if (ret < 0) + goto fail; + + ret = ff_get_encode_buffer(avctx, pkt, length, 0); + if (ret < 0) + goto fail; + + st = s->EncodeRetrieveBitstream(s->handle, s->state, pkt->data, &length, stream); + if (st != NVJPEG_STATUS_SUCCESS) { + av_log(avctx, AV_LOG_ERROR, "nvjpeg retrieve failed: %d\n", st); + ret = AVERROR_EXTERNAL; + goto fail; + } + ret = CHECK_CU(cu->cuStreamSynchronize(s->cu_stream)); + if (ret < 0) + goto fail; + + ret = CHECK_CU(cu->cuCtxPopCurrent(&dummy)); + if (ret < 0) + goto fail; + pushed = 0; + + pkt->size = (int)length; + pkt->flags |= AV_PKT_FLAG_KEY; + *got_packet = 1; + return 0; + +fail: + if (pushed) + CHECK_CU(cu->cuCtxPopCurrent(&dummy)); + return ret < 0 ? ret : AVERROR_EXTERNAL; +} + +#define OFFSET(x) offsetof(NVJPEGEncContext, x) +#define VE AV_OPT_FLAG_VIDEO_PARAM | AV_OPT_FLAG_ENCODING_PARAM +static const AVOption options[] = { + { "quality", "JPEG quality 1-100 (26 ≈ mjpeg default, 83 ≈ -q:v 2)", OFFSET(quality), AV_OPT_TYPE_INT, + { .i64 = 26 }, 1, 100, VE }, + { "chroma", "output chroma subsampling (RGB input)", OFFSET(chroma), AV_OPT_TYPE_INT, + { .i64 = NVJPEG_CSS_420 }, 0, 2, VE, .unit = "chroma" }, + { "444", "4:4:4", 0, AV_OPT_TYPE_CONST, { .i64 = NVJPEG_CSS_444 }, 0, 0, VE, .unit = "chroma" }, + { "422", "4:2:2 (MCU 2x1)", 0, AV_OPT_TYPE_CONST, { .i64 = NVJPEG_CSS_422 }, 0, 0, VE, .unit = "chroma" }, + { "420", "4:2:0", 0, AV_OPT_TYPE_CONST, { .i64 = NVJPEG_CSS_420 }, 0, 0, VE, .unit = "chroma" }, + { "matrix", "RGB→YCbCr matrix (RGB input)", OFFSET(matrix), AV_OPT_TYPE_INT, + { .i64 = NVJPEG_MATRIX_JPEG }, 0, NVJPEG_MATRIX_BT709, VE, .unit = "matrix" }, + { "jpeg", "BT.601 full range (ITU T.871)", 0, AV_OPT_TYPE_CONST, + { .i64 = NVJPEG_MATRIX_JPEG }, 0, 0, VE, .unit = "matrix" }, + { "bt601", "Same as jpeg", 0, AV_OPT_TYPE_CONST, + { .i64 = NVJPEG_MATRIX_JPEG }, 0, 0, VE, .unit = "matrix" }, + { "bt709", "BT.709 full range", 0, AV_OPT_TYPE_CONST, + { .i64 = NVJPEG_MATRIX_BT709 }, 0, 0, VE, .unit = "matrix" }, + { "huffman", "optimized Huffman tables", OFFSET(huffman), AV_OPT_TYPE_BOOL, { .i64 = 1 }, 0, 1, VE }, + { NULL }, +}; + +static const AVClass nvjpeg_class = { + .class_name = "mjpeg_nvjpeg", + .item_name = av_default_item_name, + .option = options, + .version = LIBAVUTIL_VERSION_INT, +}; + +const FFCodec ff_mjpeg_nvjpeg_encoder = { + .p.name = "mjpeg_nvjpeg", + CODEC_LONG_NAME("NVIDIA nvJPEG encoder"), + .p.type = AVMEDIA_TYPE_VIDEO, + .p.id = AV_CODEC_ID_MJPEG, + .priv_data_size = sizeof(NVJPEGEncContext), + .init = nvjpeg_encode_init, + FF_CODEC_ENCODE_CB(nvjpeg_encode_frame), + .close = nvjpeg_encode_close, + .p.capabilities = AV_CODEC_CAP_DR1 | AV_CODEC_CAP_HARDWARE | + AV_CODEC_CAP_ENCODER_REORDERED_OPAQUE, + .caps_internal = FF_CODEC_CAP_NOT_INIT_THREADSAFE | FF_CODEC_CAP_INIT_CLEANUP, + CODEC_PIXFMTS(AV_PIX_FMT_CUDA), + .p.priv_class = &nvjpeg_class, + .p.wrapper_name = "nvjpeg", + .hw_configs = nvjpeg_hw_configs, + .color_ranges = AVCOL_RANGE_JPEG, +}; diff --git a/libavcodec/nvjpegenc.cu b/libavcodec/nvjpegenc.cu new file mode 100644 index 0000000000..7f735181b3 --- /dev/null +++ b/libavcodec/nvjpegenc.cu @@ -0,0 +1,61 @@ +/* + * RGB0 → planar YCbCr 4:4:4 for nvJPEG (same integer T.871 / BT.709 + * as libavcodec/mjpegenc.c, so GoPro Player/CPU JPEG color matches). + * + * Copyright (c) 2026 Reinhard Nißl + * + * This file is part of FFmpeg. + * + * FFmpeg is free software; you can redistribute it and/or + * modify it under the terms of the GNU Lesser General Public + * License as published by the Free Software Foundation; either + * version 2.1 of the License, or (at your option) any later version. + * + * FFmpeg is distributed in the hope that it will be useful, + * but WITHOUT ANY WARRANTY; without even the implied warranty of + * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU + * Lesser General Public License for more details. + * + * You should have received a copy of the GNU Lesser General Public + * License along with FFmpeg; if not, write to the Free Software + * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA + */ + +extern "C" { + +__device__ static inline unsigned char clip_u8(int v) +{ + return (unsigned char)(v < 0 ? 0 : (v > 255 ? 255 : v)); +} + +__global__ void rgb0_to_yuv444(const unsigned char *src, int spitch, + unsigned char *ypl, unsigned char *upl, + unsigned char *vpl, + int ypitch, int upitch, int vpitch, + int width, int height, int matrix) +{ + int x = blockIdx.x * blockDim.x + threadIdx.x; + int y = blockIdx.y * blockDim.y + threadIdx.y; + if (x >= width || y >= height) + return; + + const unsigned char *s = src + y * spitch + x * 4; + int r = s[0], g = s[1], b = s[2]; + int Y, Cb, Cr; + + if (matrix) { + Y = ( 54 * r + 183 * g + 18 * b + 128) >> 8; + Cb = ((-30 * r - 98 * g + 128 * b + 128) >> 8) + 128; + Cr = ((128 * r - 116 * g - 12 * b + 128) >> 8) + 128; + } else { + Y = ( 77 * r + 150 * g + 29 * b + 128) >> 8; + Cb = ((-43 * r - 85 * g + 128 * b + 128) >> 8) + 128; + Cr = ((128 * r - 107 * g - 21 * b + 128) >> 8) + 128; + } + + ypl[y * ypitch + x] = clip_u8(Y); + upl[y * upitch + x] = clip_u8(Cb); + vpl[y * vpitch + x] = clip_u8(Cr); +} + +} -- 2.52.0 _______________________________________________ ffmpeg-devel mailing list -- [email protected] To unsubscribe send an email to [email protected]
