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]> --- Changelog | 1 + configure | 2 + doc/encoders.texi | 30 +++ libavcodec/Makefile | 1 + libavcodec/allcodecs.c | 1 + libavcodec/nvjpegenc.c | 584 ++++++++++++++++++++++++++++++++++++++++ libavcodec/nvjpegenc.cu | 61 +++++ 7 files changed, 680 insertions(+) create mode 100644 libavcodec/nvjpegenc.c create mode 100644 libavcodec/nvjpegenc.cu diff --git a/Changelog b/Changelog index 0c6341e570..dd08217835 100644 --- a/Changelog +++ b/Changelog @@ -17,6 +17,7 @@ version <next>: - 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 035d4d090a..5621ff7622 100755 --- a/configure +++ b/configure @@ -3663,6 +3663,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 b908dd8bd8..6f9f5445ed 100644 --- a/doc/encoders.texi +++ b/doc/encoders.texi @@ -1285,6 +1285,36 @@ Chroma subsampling for RGB input: @samp{444} (default), @samp{422}, or -c:v mjpeg -quality 95 -matrix jpeg -chroma 444 @end example +@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.43.0 _______________________________________________ ffmpeg-devel mailing list -- [email protected] To unsubscribe send an email to [email protected]
