PR #25153 opened by rnissl URL: https://code.ffmpeg.org/FFmpeg/FFmpeg/pulls/25153 Patch URL: https://code.ffmpeg.org/FFmpeg/FFmpeg/pulls/25153.patch
mjpeg_cuda encodes CUDA yuv444p frames to baseline JPEG inside libavcodec. The device code is PTX loaded through the driver API. The encoder does not link the CUDA runtime and does not depend on --enable-nonfree. It needs ffnvcodec and cuda_llvm. The forward DCT is the IJG integer DCT already used by the mjpeg encoder. The warp Huffman follows the BSD-2 GPUJPEG encoder and writes the standard Huffman tables. The PTX targets sm_75. When cuda_nvcc is also enabled, nvcc emits the PTX. Otherwise clang does. On 106 frames of 7680x3840 4:2:0 at quality 83, the nvcc PTX run is 262416444 bytes in 17.145 s, frame-0 RGB PSNR 44.426 dB. The clang PTX of the same source decodes, and that codegen is not tuned yet. | encoder | bytes | wall time | frame-0 RGB PSNR | | --- | ---: | ---: | ---: | | mjpeg qscale 2 | 291564015 | 25.638 s | 47.008 dB | | nvJPEG quality 87 | 279686107 | 23.979 s | 44.691 dB | | GPUJPEG quality 84 | 287545054 | 16.773 s | 44.458 dB | | mjpeg_cuda nvcc, quality 83 | 262416444 | 17.145 s | 44.426 dB | Head a0f44d528a on master d975849594. libavcodec minor goes from 14 to 15. ## Test plan - [ ] CI on this pull request - [ ] configure with ffnvcodec and cuda_llvm, without --enable-nonfree - [ ] encode one CUDA yuv444p frame and decode the JPEG - [ ] 4:2:0 output at the default restart interval of 8 From a0f44d528ac213a4e5dc44bfce2c98d304e24ab6 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Reinhard=20Ni=C3=9Fl?= <[email protected]> Date: Sat, 17 Oct 2026 14:48:52 +0200 Subject: [PATCH] avcodec: add a CUDA baseline JPEG encoder MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit mjpeg_cuda encodes CUDA yuv444p frames to baseline JPEG. The device code is PTX loaded through the driver API, so the encoder does not link the CUDA runtime and does not depend on --enable-nonfree. It needs ffnvcodec and cuda_llvm. The forward DCT is the IJG integer DCT already used by the mjpeg encoder. The warp Huffman follows the BSD-2 GPUJPEG encoder and writes the standard Huffman tables. The PTX targets sm_75. When cuda_nvcc is also enabled, nvcc emits it. Otherwise clang does. On 106 frames of 7680x3840 4:2:0 at quality 83, the nvcc PTX run is 262416444 bytes in 17.145 s, with frame-0 RGB PSNR 44.426 dB. The clang PTX of the same source decodes, and its codegen is slower. Co-authored-by: Cursor <[email protected]> Signed-off-by: Reinhard Nißl <[email protected]> --- Changelog | 1 + configure | 2 + doc/encoders.texi | 29 ++ libavcodec/Makefile | 11 + libavcodec/allcodecs.c | 1 + libavcodec/mjpegenc_cuda.c | 739 ++++++++++++++++++++++++++++++++++++ libavcodec/mjpegenc_cuda.cu | 600 +++++++++++++++++++++++++++++ libavcodec/version.h | 2 +- 8 files changed, 1384 insertions(+), 1 deletion(-) create mode 100644 libavcodec/mjpegenc_cuda.c create mode 100644 libavcodec/mjpegenc_cuda.cu diff --git a/Changelog b/Changelog index 1cfa9db1b6..f32933f458 100644 --- a/Changelog +++ b/Changelog @@ -14,6 +14,7 @@ version <next>: - NVIDIA optical flow accelerated interpolation filter (vf_fruc_vulkan) - H.264 data partitioning support - DSD (dsd_msbf) encoder +- mjpeg_cuda encoder version 9.0: diff --git a/configure b/configure index f686b727cf..6810626526 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_cuda_encoder_deps="ffnvcodec cuda_llvm" +mjpeg_cuda_encoder_select="jpegtables" 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..1daa7cb2e9 100644 --- a/doc/encoders.texi +++ b/doc/encoders.texi @@ -1268,6 +1268,35 @@ Compute and use optimal huffman tables. @end table @end table +@anchor{mjpeg_cuda} +@section mjpeg_cuda + +Baseline JPEG encoder for CUDA frames. The input software format is +@code{yuv444p}. The device code is PTX loaded through the CUDA driver +API, so the encoder does not link the CUDA runtime. Build it with +@code{cuda_llvm} and @code{ffnvcodec}. It does not require +@code{--enable-nonfree}. + +The forward DCT is the Independent JPEG Group integer DCT used by the +mjpeg encoder. The warp Huffman follows the BSD 2-Clause GPUJPEG +encoder and writes the standard Huffman tables. The PTX targets +@code{sm_75}. + +@subsection Options + +@table @option +@item quality +JPEG quality from 1 to 100. The default is 83. + +@item chroma +Output subsampling. @code{420} is the default. @code{444} is also +supported. 4:2:2 is rejected. + +@item restart +Restart interval in MCUs. The default is 8. The interval must be at +least 1. +@end table + @anchor{wavpackenc} @section wavpack diff --git a/libavcodec/Makefile b/libavcodec/Makefile index e2c60e9051..134539a590 100644 --- a/libavcodec/Makefile +++ b/libavcodec/Makefile @@ -538,6 +538,7 @@ OBJS-$(CONFIG_MJPEG_DECODER) += mjpegdec.o mjpegdec_common.o OBJS-$(CONFIG_MJPEG_QSV_DECODER) += qsvdec.o OBJS-$(CONFIG_MJPEG_ENCODER) += mjpegenc.o mjpegenc_common.o \ mjpegenc_huffman.o +OBJS-$(CONFIG_MJPEG_CUDA_ENCODER) += mjpegenc_cuda.o mjpegenc_cuda.ptx.o OBJS-$(CONFIG_MJPEGB_DECODER) += mjpegbdec.o OBJS-$(CONFIG_MJPEG_CUVID_DECODER) += cuviddec.o OBJS-$(CONFIG_MJPEG_QSV_ENCODER) += qsvenc_jpeg.o @@ -1453,3 +1454,13 @@ $(SUBDIR)pcm.o: $(SUBDIR)pcm_tables.h $(SUBDIR)qdm2.o: $(SUBDIR)qdm2_tables.h $(SUBDIR)sinewin.o: $(SUBDIR)sinewin_tables.h endif + +# Warp ballot needs sm_70. PTX only, loaded through the driver API. +# nvcc emits this PTX when cuda_nvcc is enabled. Otherwise clang does. +ifeq ($(CONFIG_CUDA_NVCC),yes) +$(SUBDIR)mjpegenc_cuda.ptx: $(SRC_PATH)/libavcodec/mjpegenc_cuda.cu + $(NVCC) -gencode arch=compute_75,code=sm_75 -O2 -std=c++14 -m64 -ptx -o $@ $< +else +$(SUBDIR)mjpegenc_cuda.ptx: $(SRC_PATH)/libavcodec/mjpegenc_cuda.cu $(SRC_PATH)/compat/cuda/cuda_runtime.h + $(NVCC) --cuda-gpu-arch=sm_75 -O2 -std=c++14 -m64 -S -nocudalib -nocudainc --cuda-device-only -Wno-c++11-narrowing -Wno-unknown-cuda-version -include $(SRC_PATH)/compat/cuda/cuda_runtime.h -o $@ $< +endif diff --git a/libavcodec/allcodecs.c b/libavcodec/allcodecs.c index ca20c76fe6..2930c50a96 100644 --- a/libavcodec/allcodecs.c +++ b/libavcodec/allcodecs.c @@ -198,6 +198,7 @@ extern const FFCodec ff_mdec_decoder; extern const FFCodec ff_media100_decoder; extern const FFCodec ff_mimic_decoder; extern const FFCodec ff_mjpeg_encoder; +extern const FFCodec ff_mjpeg_cuda_encoder; extern const FFCodec ff_mjpeg_decoder; extern const FFCodec ff_mjpegb_decoder; extern const FFCodec ff_mmvideo_decoder; diff --git a/libavcodec/mjpegenc_cuda.c b/libavcodec/mjpegenc_cuda.c new file mode 100644 index 0000000000..9af1baefb4 --- /dev/null +++ b/libavcodec/mjpegenc_cuda.c @@ -0,0 +1,739 @@ +/* + * Baseline JPEG encoder on CUDA frames. + * + * Device code is PTX loaded through the driver API. The DCT is the IJG + * integer DCT. The warp Huffman follows CESNET GPUJPEG (BSD 2-Clause). + * This software is based in part on the work of the Independent JPEG Group. + * + * 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/avassert.h" +#include "libavutil/cuda_check.h" +#include "libavutil/hwcontext.h" +#include "libavutil/hwcontext_cuda_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 "jpegquanttables.h" +#include "jpegtables.h" + +#define CHECK_CU(x) FF_CUDA_CHECK_DL(avctx, s->hwctx->internal->cuda_dl, x) + +enum { + MJCU_444 = 0, + MJCU_420 = 2, +}; + +typedef struct MJCudaSeg { + int codewords; + int bytes; +} MJCudaSeg; + +typedef struct MJCudaContext { + const AVClass *class; + + AVCUDADeviceContext *hwctx; + CUstream cu_stream; + CUmodule module; + CUfunction fdct, huffman, serialize, prefix, compact; + + int quality; + int chroma; + int restart; + + int width, height, chroma420; + int mcu_x, mcu_y, nseg; + int blocks_x_y, blocks_y_y, blocks_x_c, blocks_y_c; + int cw_stride, max_bytes; + + CUdeviceptr d_qy, d_qc, d_order, d_decomp, d_lut; + CUdeviceptr d_coef_y, d_coef_u, d_coef_v; + CUdeviceptr d_cw, d_bytes, d_segs, d_off, d_compact, d_total, d_status; + + int warned; +} MJCudaContext; + +static const AVCodecHWConfigInternal *const mjpeg_cuda_hw_configs[] = { + HW_CONFIG_ENCODER_FRAMES(CUDA, CUDA), + HW_CONFIG_ENCODER_DEVICE(NONE, CUDA), + NULL +}; + +extern const unsigned char ff_mjpegenc_cuda_ptx_data[]; +extern const unsigned int ff_mjpegenc_cuda_ptx_len; +extern const uint8_t ff_zigzag_direct[64]; + +static void w8(uint8_t **p, int v) +{ + *(*p)++ = (uint8_t)v; +} + +static void w16(uint8_t **p, int v) +{ + w8(p, v >> 8); + w8(p, v); +} + +static int symbol_count(const uint8_t *bits) +{ + int n = 0, i; + + for (i = 1; i <= 16; i++) + n += bits[i]; + return n; +} + +static void write_huff(uint8_t **p, int cls, int id, const uint8_t *bits, + const uint8_t *val) +{ + int i, n = symbol_count(bits); + + w8(p, (cls << 4) | id); + for (i = 1; i <= 16; i++) + w8(p, bits[i]); + for (i = 0; i < n; i++) + w8(p, val[i]); +} + +static void pack_table(uint32_t *dest, const uint8_t *bits, const uint8_t *val, + int is_ac) +{ + uint8_t size[256]; + uint16_t code[256]; + int k = 0, c = 0, i, j; + + memset(size, 0, sizeof(size)); + memset(code, 0, sizeof(code)); + for (i = 1; i <= 16; i++) { + for (j = 0; j < bits[i]; j++) { + int sym = val[k++]; + + size[sym] = (uint8_t)i; + code[sym] = (uint16_t)c++; + } + c <<= 1; + } + for (i = 0; i <= 256; i++) { + int sz = size[i & 255]; + + dest[i] = sz ? ((uint32_t)code[i & 255] << (32 - sz)) | (uint32_t)sz : 0; + } + if (is_ac) + dest[0] = 0; +} + +static void scale_quant(uint16_t *dst, const uint8_t *src, int quality) +{ + int s = quality < 50 ? 5000 / quality : 200 - 2 * quality; + int i; + + for (i = 0; i < 64; i++) { + int v = (s * src[i] + 50) / 100; + + if (v < 1) + v = 1; + if (v > 255) + v = 255; + dst[i] = (uint16_t)v; + } +} + +static int header_write(uint8_t *dst, int dst_size, int w, int h, int chroma420, + int restart, const uint16_t *qy, const uint16_t *qc) +{ + uint8_t *p = dst; + uint8_t *dht_len; + int i, y_samp = chroma420 ? 0x22 : 0x11; + int nbits_ac_y = symbol_count(ff_mjpeg_bits_ac_luminance); + int nbits_ac_c = symbol_count(ff_mjpeg_bits_ac_chrominance); + int dht_bytes = 2 + (1 + 16 + 12) * 2 + (1 + 16 + nbits_ac_y) + (1 + 16 + nbits_ac_c); + + if (dst_size < 1024) + return AVERROR(EINVAL); + + w8(&p, 0xff); w8(&p, 0xd8); + w8(&p, 0xff); w8(&p, 0xe0); + w16(&p, 16); + w8(&p, 'J'); w8(&p, 'F'); w8(&p, 'I'); w8(&p, 'F'); w8(&p, 0); + w8(&p, 1); w8(&p, 2); + w8(&p, 0); + w16(&p, 1); w16(&p, 1); + w8(&p, 0); w8(&p, 0); + + w8(&p, 0xff); w8(&p, 0xdb); + w16(&p, 2 + 65 * 2); + w8(&p, 0x00); + for (i = 0; i < 64; i++) + w8(&p, qy[ff_zigzag_direct[i]]); + w8(&p, 0x01); + for (i = 0; i < 64; i++) + w8(&p, qc[ff_zigzag_direct[i]]); + + w8(&p, 0xff); w8(&p, 0xc0); + w16(&p, 17); + w8(&p, 8); + w16(&p, h); + w16(&p, w); + w8(&p, 3); + w8(&p, 1); w8(&p, y_samp); w8(&p, 0); + w8(&p, 2); w8(&p, 0x11); w8(&p, 1); + w8(&p, 3); w8(&p, 0x11); w8(&p, 1); + + w8(&p, 0xff); w8(&p, 0xc4); + dht_len = p; + w16(&p, dht_bytes); + write_huff(&p, 0, 0, ff_mjpeg_bits_dc_luminance, ff_mjpeg_val_dc); + write_huff(&p, 1, 0, ff_mjpeg_bits_ac_luminance, ff_mjpeg_val_ac_luminance); + write_huff(&p, 0, 1, ff_mjpeg_bits_dc_chrominance, ff_mjpeg_val_dc); + write_huff(&p, 1, 1, ff_mjpeg_bits_ac_chrominance, ff_mjpeg_val_ac_chrominance); + av_assert0(p - dht_len == dht_bytes); + + w8(&p, 0xff); w8(&p, 0xdd); + w16(&p, 4); + w16(&p, restart); + + w8(&p, 0xff); w8(&p, 0xda); + w16(&p, 12); + w8(&p, 3); + w8(&p, 1); w8(&p, 0x00); + w8(&p, 2); w8(&p, 0x11); + w8(&p, 3); w8(&p, 0x11); + w8(&p, 0); w8(&p, 63); w8(&p, 0); + + return (int)(p - dst); +} + +static void free_frame(MJCudaContext *s) +{ + CudaFunctions *cu = s->hwctx ? s->hwctx->internal->cuda_dl : NULL; + + if (!cu) + return; +#define FREE(p) do { if (p) { cu->cuMemFree(p); p = 0; } } while (0) + FREE(s->d_coef_y); FREE(s->d_coef_u); FREE(s->d_coef_v); + FREE(s->d_cw); FREE(s->d_bytes); FREE(s->d_segs); FREE(s->d_off); + FREE(s->d_compact); FREE(s->d_total); FREE(s->d_status); +#undef FREE +} + +static void free_dev(MJCudaContext *s) +{ + CudaFunctions *cu = s->hwctx ? s->hwctx->internal->cuda_dl : NULL; + + free_frame(s); + if (!cu) + return; +#define FREE(p) do { if (p) { cu->cuMemFree(p); p = 0; } } while (0) + FREE(s->d_qy); FREE(s->d_qc); FREE(s->d_order); FREE(s->d_decomp); FREE(s->d_lut); +#undef FREE +} + +static int alloc_bufs(AVCodecContext *avctx, int w, int h) +{ + MJCudaContext *s = avctx->priv_data; + CudaFunctions *cu = s->hwctx->internal->cuda_dl; + int chroma420 = s->chroma == MJCU_420; + int ret; + + if (w <= 0 || h <= 0) + return AVERROR(EINVAL); + if (chroma420) { + if ((w & 15) || (h & 15)) { + av_log(avctx, AV_LOG_ERROR, "4:2:0 needs width and height divisible by 16\n"); + return AVERROR(EINVAL); + } + } else if ((w & 7) || (h & 7)) { + av_log(avctx, AV_LOG_ERROR, "4:4:4 needs width and height divisible by 8\n"); + return AVERROR(EINVAL); + } + + /* Quant, zigzag and Huffman tables stay allocated for the context. + * free_dev() here used to drop them, and the DCT then read address 0. */ + free_frame(s); + s->width = w; + s->height = h; + s->chroma420 = chroma420; + s->blocks_x_y = w / 8; + s->blocks_y_y = h / 8; + s->blocks_x_c = chroma420 ? w / 16 : w / 8; + s->blocks_y_c = chroma420 ? h / 16 : h / 8; + s->mcu_x = s->blocks_x_c; + s->mcu_y = s->blocks_y_c; + s->nseg = (s->mcu_x * s->mcu_y + s->restart - 1) / s->restart; + s->cw_stride = s->restart * (chroma420 ? 6 : 3) * 64 + 128; + s->max_bytes = s->restart * (chroma420 ? 6 : 3) * 192; + + ret = CHECK_CU(cu->cuMemAlloc(&s->d_coef_y, (size_t)s->blocks_x_y * s->blocks_y_y * 64 * sizeof(short))); + if (ret < 0) + return ret; + ret = CHECK_CU(cu->cuMemAlloc(&s->d_coef_u, (size_t)s->blocks_x_c * s->blocks_y_c * 64 * sizeof(short))); + if (ret < 0) + return ret; + ret = CHECK_CU(cu->cuMemAlloc(&s->d_coef_v, (size_t)s->blocks_x_c * s->blocks_y_c * 64 * sizeof(short))); + if (ret < 0) + return ret; + ret = CHECK_CU(cu->cuMemAlloc(&s->d_cw, (size_t)s->nseg * s->cw_stride * sizeof(unsigned))); + if (ret < 0) + return ret; + ret = CHECK_CU(cu->cuMemAlloc(&s->d_bytes, (size_t)s->nseg * s->max_bytes)); + if (ret < 0) + return ret; + ret = CHECK_CU(cu->cuMemAlloc(&s->d_segs, (size_t)s->nseg * sizeof(MJCudaSeg))); + if (ret < 0) + return ret; + ret = CHECK_CU(cu->cuMemAlloc(&s->d_off, (size_t)s->nseg * sizeof(unsigned))); + if (ret < 0) + return ret; + ret = CHECK_CU(cu->cuMemAlloc(&s->d_compact, (size_t)s->nseg * s->max_bytes)); + if (ret < 0) + return ret; + ret = CHECK_CU(cu->cuMemAlloc(&s->d_total, sizeof(unsigned))); + if (ret < 0) + return ret; + ret = CHECK_CU(cu->cuMemAlloc(&s->d_status, sizeof(int))); + if (ret < 0) + return ret; + return 0; +} + +static int launch(AVCodecContext *avctx, CUfunction fn, + unsigned gx, unsigned gy, unsigned bx, unsigned by, void **args) +{ + MJCudaContext *s = avctx->priv_data; + + return CHECK_CU(s->hwctx->internal->cuda_dl->cuLaunchKernel( + fn, gx, gy, 1, bx, by, 1, 0, s->cu_stream, args, NULL)); +} + +static int fdct_plane(AVCodecContext *avctx, CUdeviceptr src, int stride, + int blocks_x, int blocks_y, int step, CUdeviceptr quant, + CUdeviceptr dst) +{ + MJCudaContext *s = avctx->priv_data; + void *args[] = { + &src, &stride, &blocks_x, &blocks_y, &step, &quant, &dst + }; + unsigned n = (unsigned)blocks_x * (unsigned)blocks_y; + + if (!n) + return 0; + return launch(avctx, s->fdct, n, 1, 32, 1, args); +} + +static av_cold int mjpeg_cuda_init(AVCodecContext *avctx) +{ + MJCudaContext *s = avctx->priv_data; + AVHWFramesContext *frames = NULL; + CudaFunctions *cu; + CUcontext dummy; + uint16_t qy[64], qc[64]; + uint32_t lut[257 * 4]; + uint32_t decomp[8192]; + int order[64]; + int ret, pushed = 0, i; + + if (s->quality < 1 || s->quality > 100 || s->restart < 1) + return AVERROR(EINVAL); + if (s->chroma != MJCU_444 && s->chroma != MJCU_420) { + av_log(avctx, AV_LOG_ERROR, "chroma must be 444 or 420\n"); + return AVERROR(EINVAL); + } + + if (avctx->hw_frames_ctx) { + frames = (AVHWFramesContext *)avctx->hw_frames_ctx->data; + if (frames->format != AV_PIX_FMT_CUDA) + return AVERROR(EINVAL); + s->hwctx = frames->device_ctx->hwctx; + } else if (avctx->hw_device_ctx) { + AVHWDeviceContext *dev = (AVHWDeviceContext *)avctx->hw_device_ctx->data; + + if (dev->type != AV_HWDEVICE_TYPE_CUDA) + return AVERROR(EINVAL); + s->hwctx = dev->hwctx; + } else { + av_log(avctx, AV_LOG_ERROR, "CUDA frames or device context required\n"); + return AVERROR(EINVAL); + } + s->cu_stream = s->hwctx->stream; + cu = s->hwctx->internal->cuda_dl; + + ret = CHECK_CU(cu->cuCtxPushCurrent(s->hwctx->cuda_ctx)); + if (ret < 0) + return ret; + pushed = 1; + + ret = CHECK_CU(cu->cuModuleLoadData(&s->module, ff_mjpegenc_cuda_ptx_data)); + if (ret < 0) + goto fail; + ret = CHECK_CU(cu->cuModuleGetFunction(&s->fdct, s->module, "mjpeg_cuda_fdct")); + if (ret < 0) + goto fail; + ret = CHECK_CU(cu->cuModuleGetFunction(&s->huffman, s->module, "mjpeg_cuda_huffman")); + if (ret < 0) + goto fail; + ret = CHECK_CU(cu->cuModuleGetFunction(&s->serialize, s->module, "mjpeg_cuda_serialize")); + if (ret < 0) + goto fail; + ret = CHECK_CU(cu->cuModuleGetFunction(&s->prefix, s->module, "mjpeg_cuda_prefix")); + if (ret < 0) + goto fail; + ret = CHECK_CU(cu->cuModuleGetFunction(&s->compact, s->module, "mjpeg_cuda_compact")); + if (ret < 0) + goto fail; + + scale_quant(qy, ff_mjpeg_std_luminance_quant_tbl, s->quality); + scale_quant(qc, ff_mjpeg_std_chrominance_quant_tbl, s->quality); + for (i = 0; i < 64; i++) + order[i] = ff_zigzag_direct[i]; + for (i = -4096; i <= 4095; i++) { + unsigned code = (unsigned)i; + int mag = i; + int nbits = 0; + + if (i < 0) { + mag = -mag; + code--; + } + while (mag) { + nbits++; + mag >>= 1; + } + decomp[i + 4096] = nbits ? (nbits | (code << (32 - nbits))) : 0; + } + pack_table(lut + 257 * 0, ff_mjpeg_bits_ac_luminance, ff_mjpeg_val_ac_luminance, 1); + pack_table(lut + 257 * 1, ff_mjpeg_bits_dc_luminance, ff_mjpeg_val_dc, 0); + pack_table(lut + 257 * 2, ff_mjpeg_bits_ac_chrominance, ff_mjpeg_val_ac_chrominance, 1); + pack_table(lut + 257 * 3, ff_mjpeg_bits_dc_chrominance, ff_mjpeg_val_dc, 0); + + ret = CHECK_CU(cu->cuMemAlloc(&s->d_qy, sizeof(qy))); + if (ret < 0) + goto fail; + ret = CHECK_CU(cu->cuMemAlloc(&s->d_qc, sizeof(qc))); + if (ret < 0) + goto fail; + ret = CHECK_CU(cu->cuMemAlloc(&s->d_order, sizeof(order))); + if (ret < 0) + goto fail; + ret = CHECK_CU(cu->cuMemAlloc(&s->d_decomp, sizeof(decomp))); + if (ret < 0) + goto fail; + ret = CHECK_CU(cu->cuMemAlloc(&s->d_lut, sizeof(lut))); + if (ret < 0) + goto fail; + ret = CHECK_CU(cu->cuMemcpyHtoDAsync(s->d_qy, qy, sizeof(qy), s->cu_stream)); + if (ret < 0) + goto fail; + ret = CHECK_CU(cu->cuMemcpyHtoDAsync(s->d_qc, qc, sizeof(qc), s->cu_stream)); + if (ret < 0) + goto fail; + ret = CHECK_CU(cu->cuMemcpyHtoDAsync(s->d_order, order, sizeof(order), s->cu_stream)); + if (ret < 0) + goto fail; + ret = CHECK_CU(cu->cuMemcpyHtoDAsync(s->d_decomp, decomp, sizeof(decomp), s->cu_stream)); + if (ret < 0) + goto fail; + ret = CHECK_CU(cu->cuMemcpyHtoDAsync(s->d_lut, lut, sizeof(lut), s->cu_stream)); + if (ret < 0) + goto fail; + + if (avctx->width && avctx->height) { + ret = alloc_bufs(avctx, avctx->width, avctx->height); + if (ret < 0) + goto fail; + } + + ret = CHECK_CU(cu->cuStreamSynchronize(s->cu_stream)); + if (ret < 0) + goto fail; + ret = CHECK_CU(cu->cuCtxPopCurrent(&dummy)); + if (ret < 0) + return ret; + + avctx->color_range = AVCOL_RANGE_JPEG; + return 0; + +fail: + free_dev(s); + if (s->module) { + cu->cuModuleUnload(s->module); + s->module = NULL; + } + if (pushed) + cu->cuCtxPopCurrent(&dummy); + return ret < 0 ? ret : AVERROR_EXTERNAL; +} + +static av_cold int mjpeg_cuda_close(AVCodecContext *avctx) +{ + MJCudaContext *s = avctx->priv_data; + CUcontext dummy; + CudaFunctions *cu; + + if (!s->hwctx || !s->hwctx->internal || !s->hwctx->internal->cuda_dl) + return 0; + cu = s->hwctx->internal->cuda_dl; + if (cu->cuCtxPushCurrent(s->hwctx->cuda_ctx) != CUDA_SUCCESS) { + av_log(avctx, AV_LOG_ERROR, "Cannot push CUDA context; encoder state kept\n"); + return AVERROR_EXTERNAL; + } + free_dev(s); + if (s->module) + cu->cuModuleUnload(s->module); + s->module = NULL; + cu->cuCtxPopCurrent(&dummy); + return 0; +} + +static int mjpeg_cuda_encode(AVCodecContext *avctx, AVPacket *pkt, + const AVFrame *frame, int *got_packet) +{ + MJCudaContext *s = avctx->priv_data; + CudaFunctions *cu; + CUcontext dummy; + CUdeviceptr y, u, v; + uint8_t *scan = NULL, *hdr = NULL; + unsigned total = 0, status = 0; + int ret, pushed = 0, hdr_len, stride_y, stride_c; + unsigned groups, gx, gy, sx; + void *hargs[16], *sargs[8], *pargs[4], *cargs[6]; + + if (!frame) + return 0; + if (!frame->hw_frames_ctx) + return AVERROR(EINVAL); + { + AVHWFramesContext *fc = (AVHWFramesContext *)frame->hw_frames_ctx->data; + + if (!s->hwctx) + s->hwctx = fc->device_ctx->hwctx; + if (fc->sw_format != AV_PIX_FMT_YUV444P) { + av_log(avctx, AV_LOG_ERROR, "CUDA sw_format must be yuv444p\n"); + return AVERROR(ENOSYS); + } + } + + cu = s->hwctx->internal->cuda_dl; + ret = CHECK_CU(cu->cuCtxPushCurrent(s->hwctx->cuda_ctx)); + if (ret < 0) + return ret; + pushed = 1; + + if (s->width != frame->width || s->height != frame->height) { + ret = alloc_bufs(avctx, frame->width, frame->height); + if (ret < 0) + goto fail; + } + + y = (CUdeviceptr)(uintptr_t)frame->data[0]; + u = (CUdeviceptr)(uintptr_t)frame->data[1]; + v = (CUdeviceptr)(uintptr_t)frame->data[2]; + stride_y = frame->linesize[0]; + stride_c = frame->linesize[1]; + if (stride_y < frame->width || stride_c < frame->width || frame->linesize[2] != stride_c) { + av_log(avctx, AV_LOG_ERROR, "unexpected CUDA linesize\n"); + ret = AVERROR(EINVAL); + goto fail; + } + + ret = CHECK_CU(cu->cuMemsetD8Async(s->d_status, 0, sizeof(int), s->cu_stream)); + if (ret < 0) + goto fail; + ret = CHECK_CU(cu->cuMemsetD8Async(s->d_segs, 0, (size_t)s->nseg * sizeof(MJCudaSeg), s->cu_stream)); + if (ret < 0) + goto fail; + + ret = fdct_plane(avctx, y, stride_y, s->blocks_x_y, s->blocks_y_y, 1, s->d_qy, s->d_coef_y); + if (ret < 0) + goto fail; + ret = fdct_plane(avctx, u, stride_c, s->blocks_x_c, s->blocks_y_c, s->chroma420 ? 2 : 1, s->d_qc, s->d_coef_u); + if (ret < 0) + goto fail; + ret = fdct_plane(avctx, v, frame->linesize[2], s->blocks_x_c, s->blocks_y_c, s->chroma420 ? 2 : 1, s->d_qc, s->d_coef_v); + if (ret < 0) + goto fail; + + groups = (unsigned)((s->nseg + 3) / 4); + gx = groups > 65535 ? 65535 : groups; + gy = (groups + gx - 1) / gx; + hargs[0] = &s->d_segs; hargs[1] = &s->nseg; hargs[2] = &gx; hargs[3] = &s->restart; + hargs[4] = &s->mcu_x; hargs[5] = &s->mcu_y; hargs[6] = &s->chroma420; + hargs[7] = &s->d_coef_y; hargs[8] = &s->blocks_x_y; + hargs[9] = &s->d_coef_u; hargs[10] = &s->d_coef_v; hargs[11] = &s->blocks_x_c; + hargs[12] = &s->d_order; hargs[13] = &s->d_decomp; hargs[14] = &s->d_lut; + hargs[15] = &s->d_cw; + { + void *args[18]; + int si; + + for (si = 0; si < 16; si++) + args[si] = hargs[si]; + args[16] = &s->cw_stride; + args[17] = &s->d_status; + ret = launch(avctx, s->huffman, gx, gy, 128, 1, args); + if (ret < 0) + goto fail; + } + + sx = 128; + groups = (unsigned)((s->nseg + (int)sx - 1) / (int)sx); + gx = groups > 65535 ? 65535 : groups; + gy = (groups + gx - 1) / gx; + /* serialize is 1D in the kernel; fold gy into gx when it fits */ + if (gy != 1) { + av_log(avctx, AV_LOG_ERROR, "too many restart segments\n"); + ret = AVERROR(EINVAL); + goto fail; + } + sargs[0] = &s->d_segs; sargs[1] = &s->nseg; sargs[2] = &s->d_cw; + sargs[3] = &s->cw_stride; sargs[4] = &s->d_bytes; sargs[5] = &s->max_bytes; + sargs[6] = &s->d_status; + ret = launch(avctx, s->serialize, gx, 1, sx, 1, sargs); + if (ret < 0) + goto fail; + + pargs[0] = &s->d_segs; pargs[1] = &s->nseg; pargs[2] = &s->d_off; pargs[3] = &s->d_total; + ret = launch(avctx, s->prefix, 1, 1, 1, 1, pargs); + if (ret < 0) + goto fail; + cargs[0] = &s->d_segs; cargs[1] = &s->d_off; cargs[2] = &s->nseg; + cargs[3] = &s->d_bytes; cargs[4] = &s->max_bytes; cargs[5] = &s->d_compact; + ret = launch(avctx, s->compact, gx, 1, sx, 1, cargs); + if (ret < 0) + goto fail; + + ret = CHECK_CU(cu->cuStreamSynchronize(s->cu_stream)); + if (ret < 0) + goto fail; + ret = CHECK_CU(cu->cuMemcpyDtoH(&status, s->d_status, sizeof(status))); + if (ret < 0) + goto fail; + if (status) { + av_log(avctx, AV_LOG_ERROR, "entropy buffer too small\n"); + ret = AVERROR(ERANGE); + goto fail; + } + ret = CHECK_CU(cu->cuMemcpyDtoH(&total, s->d_total, sizeof(total))); + if (ret < 0) + goto fail; + if (total < 2 || total > (unsigned)s->nseg * (unsigned)s->max_bytes) { + ret = AVERROR(ERANGE); + goto fail; + } + + scan = av_malloc(total); + hdr = av_malloc(2048); + if (!scan || !hdr) { + ret = AVERROR(ENOMEM); + goto fail; + } + ret = CHECK_CU(cu->cuMemcpyDtoH(scan, s->d_compact, total)); + if (ret < 0) + goto fail; + + { + uint16_t qy[64], qc[64]; + + scale_quant(qy, ff_mjpeg_std_luminance_quant_tbl, s->quality); + scale_quant(qc, ff_mjpeg_std_chrominance_quant_tbl, s->quality); + hdr_len = header_write(hdr, 2048, frame->width, frame->height, s->chroma420, + s->restart, qy, qc); + } + if (hdr_len < 0) { + ret = hdr_len; + goto fail; + } + if (scan[total - 2] != 0xff || scan[total - 1] < 0xd0 || scan[total - 1] > 0xd7) { + av_log(avctx, AV_LOG_ERROR, "missing restart marker at end of scan\n"); + ret = AVERROR_INVALIDDATA; + goto fail; + } + scan[total - 1] = 0xd9; + + ret = ff_get_encode_buffer(avctx, pkt, hdr_len + (int)total, 0); + if (ret < 0) + goto fail; + memcpy(pkt->data, hdr, hdr_len); + memcpy(pkt->data + hdr_len, scan, total); + + ret = CHECK_CU(cu->cuCtxPopCurrent(&dummy)); + if (ret < 0) + goto fail; + pushed = 0; + + pkt->flags |= AV_PKT_FLAG_KEY; + *got_packet = 1; + av_freep(&scan); + av_freep(&hdr); + return 0; + +fail: + av_freep(&scan); + av_freep(&hdr); + if (pushed) + cu->cuCtxPopCurrent(&dummy); + return ret < 0 ? ret : AVERROR_EXTERNAL; +} + +#define OFFSET(x) offsetof(MJCudaContext, x) +#define VE AV_OPT_FLAG_VIDEO_PARAM | AV_OPT_FLAG_ENCODING_PARAM +static const AVOption options[] = { + { "quality", "JPEG quality 1-100", OFFSET(quality), AV_OPT_TYPE_INT, { .i64 = 83 }, 1, 100, VE }, + { "chroma", "output chroma subsampling", OFFSET(chroma), AV_OPT_TYPE_INT, { .i64 = MJCU_420 }, 0, 2, VE, .unit = "chroma" }, + { "444", "4:4:4", 0, AV_OPT_TYPE_CONST, { .i64 = MJCU_444 }, 0, 0, VE, .unit = "chroma" }, + { "420", "4:2:0", 0, AV_OPT_TYPE_CONST, { .i64 = MJCU_420 }, 0, 0, VE, .unit = "chroma" }, + { "restart", "restart interval in MCUs", OFFSET(restart), AV_OPT_TYPE_INT, { .i64 = 8 }, 1, 65535, VE }, + { NULL } +}; + +static const AVClass mjpeg_cuda_class = { + .class_name = "mjpeg_cuda", + .item_name = av_default_item_name, + .option = options, + .version = LIBAVUTIL_VERSION_INT, +}; + +const FFCodec ff_mjpeg_cuda_encoder = { + .p.name = "mjpeg_cuda", + CODEC_LONG_NAME("CUDA baseline JPEG"), + .p.type = AVMEDIA_TYPE_VIDEO, + .p.id = AV_CODEC_ID_MJPEG, + .priv_data_size = sizeof(MJCudaContext), + .init = mjpeg_cuda_init, + FF_CODEC_ENCODE_CB(mjpeg_cuda_encode), + .close = mjpeg_cuda_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 = &mjpeg_cuda_class, + .p.wrapper_name = "cuda", + .hw_configs = mjpeg_cuda_hw_configs, + .color_ranges = AVCOL_RANGE_JPEG, +}; diff --git a/libavcodec/mjpegenc_cuda.cu b/libavcodec/mjpegenc_cuda.cu new file mode 100644 index 0000000000..2c3ce6535d --- /dev/null +++ b/libavcodec/mjpegenc_cuda.cu @@ -0,0 +1,600 @@ +/* + * CUDA baseline JPEG: IJG integer DCT and a GPU Huffman. + * + * The forward DCT is the Independent JPEG Group slow integer DCT + * (jfdctint, 8-bit), the same transform FFmpeg's mjpeg encoder uses. + * This software is based in part on the work of the Independent JPEG Group. + * + * The warp Huffman is adapted from CESNET GPUJPEG + * (src/gpujpeg_huffman_gpu_encoder.cu): + * Copyright (c) 2011-2025, CESNET + * Copyright (c) 2011, Silicon Genome, LLC. + * BSD 2-Clause. Runs of 16 or more zeros take a serial fallback; + * the warp path only emits a single run symbol. + * + * 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 + */ + +#define CONST_BITS 13 +#define PASS1_BITS 4 +#define OUT_SHIFT PASS1_BITS + +#define FIX_0_298631336 ((int) 2446) +#define FIX_0_390180644 ((int) 3196) +#define FIX_0_541196100 ((int) 4433) +#define FIX_0_765366865 ((int) 6270) +#define FIX_0_899976223 ((int) 7373) +#define FIX_1_175875602 ((int) 9633) +#define FIX_1_501321110 ((int) 12299) +#define FIX_1_847759065 ((int) 15137) +#define FIX_1_961570560 ((int) 16069) +#define FIX_2_053119869 ((int) 16819) +#define FIX_2_562915447 ((int) 20995) +#define FIX_3_072711026 ((int) 25172) + +#define DESCALE(x, n) (((x) + (1 << ((n) - 1))) >> (n)) +#define MULTIPLY(var, cnst) ((int)((var) * (cnst))) + +/* Clang's device compile through compat/cuda/cuda_runtime.h does not + * provide the CUDA __shared__ keyword. nvcc already has it. */ +#ifdef __clang__ +#define __shared__ __attribute__((shared)) +#endif + +struct MJCudaSeg { + int codewords; + int bytes; +}; + +__device__ static void +islow_fdct_pass1(int s0, int s1, int s2, int s3, int s4, int s5, int s6, int s7, + int &d0, int &d1, int &d2, int &d3, int &d4, int &d5, int &d6, int &d7) +{ + int tmp0 = s0 + s7; + int tmp7 = s0 - s7; + int tmp1 = s1 + s6; + int tmp6 = s1 - s6; + int tmp2 = s2 + s5; + int tmp5 = s2 - s5; + int tmp3 = s3 + s4; + int tmp4 = s3 - s4; + int tmp10 = tmp0 + tmp3; + int tmp13 = tmp0 - tmp3; + int tmp11 = tmp1 + tmp2; + int tmp12 = tmp1 - tmp2; + int z1, z1o, z2, z3, z4, z5; + + d0 = (int)(short)((tmp10 + tmp11) << PASS1_BITS); + d4 = (int)(short)((tmp10 - tmp11) << PASS1_BITS); + + z1 = MULTIPLY(tmp12 + tmp13, FIX_0_541196100); + d2 = (int)(short)DESCALE(z1 + MULTIPLY(tmp13, FIX_0_765366865), CONST_BITS - PASS1_BITS); + d6 = (int)(short)DESCALE(z1 + MULTIPLY(tmp12, -FIX_1_847759065), CONST_BITS - PASS1_BITS); + + z1o = tmp4 + tmp7; + z2 = tmp5 + tmp6; + z3 = tmp4 + tmp6; + z4 = tmp5 + tmp7; + z5 = MULTIPLY(z3 + z4, FIX_1_175875602); + + tmp4 = MULTIPLY(tmp4, FIX_0_298631336); + tmp5 = MULTIPLY(tmp5, FIX_2_053119869); + tmp6 = MULTIPLY(tmp6, FIX_3_072711026); + tmp7 = MULTIPLY(tmp7, FIX_1_501321110); + z1o = MULTIPLY(z1o, -FIX_0_899976223); + z2 = MULTIPLY(z2, -FIX_2_562915447); + z3 = MULTIPLY(z3, -FIX_1_961570560); + z4 = MULTIPLY(z4, -FIX_0_390180644); + z3 += z5; + z4 += z5; + + d7 = (int)(short)DESCALE(tmp4 + z1o + z3, CONST_BITS - PASS1_BITS); + d5 = (int)(short)DESCALE(tmp5 + z2 + z4, CONST_BITS - PASS1_BITS); + d3 = (int)(short)DESCALE(tmp6 + z2 + z3, CONST_BITS - PASS1_BITS); + d1 = (int)(short)DESCALE(tmp7 + z1o + z4, CONST_BITS - PASS1_BITS); +} + +__device__ static void +islow_fdct_pass2(int s0, int s1, int s2, int s3, int s4, int s5, int s6, int s7, + int &d0, int &d1, int &d2, int &d3, int &d4, int &d5, int &d6, int &d7) +{ + int tmp0 = s0 + s7; + int tmp7 = s0 - s7; + int tmp1 = s1 + s6; + int tmp6 = s1 - s6; + int tmp2 = s2 + s5; + int tmp5 = s2 - s5; + int tmp3 = s3 + s4; + int tmp4 = s3 - s4; + int tmp10 = tmp0 + tmp3; + int tmp13 = tmp0 - tmp3; + int tmp11 = tmp1 + tmp2; + int tmp12 = tmp1 - tmp2; + int z1, z1o, z2, z3, z4, z5; + + d0 = DESCALE(tmp10 + tmp11, OUT_SHIFT); + d4 = DESCALE(tmp10 - tmp11, OUT_SHIFT); + + z1 = MULTIPLY(tmp12 + tmp13, FIX_0_541196100); + d2 = DESCALE(z1 + MULTIPLY(tmp13, FIX_0_765366865), CONST_BITS + OUT_SHIFT); + d6 = DESCALE(z1 + MULTIPLY(tmp12, -FIX_1_847759065), CONST_BITS + OUT_SHIFT); + + z1o = tmp4 + tmp7; + z2 = tmp5 + tmp6; + z3 = tmp4 + tmp6; + z4 = tmp5 + tmp7; + z5 = MULTIPLY(z3 + z4, FIX_1_175875602); + + tmp4 = MULTIPLY(tmp4, FIX_0_298631336); + tmp5 = MULTIPLY(tmp5, FIX_2_053119869); + tmp6 = MULTIPLY(tmp6, FIX_3_072711026); + tmp7 = MULTIPLY(tmp7, FIX_1_501321110); + z1o = MULTIPLY(z1o, -FIX_0_899976223); + z2 = MULTIPLY(z2, -FIX_2_562915447); + z3 = MULTIPLY(z3, -FIX_1_961570560); + z4 = MULTIPLY(z4, -FIX_0_390180644); + z3 += z5; + z4 += z5; + + d7 = DESCALE(tmp4 + z1o + z3, CONST_BITS + OUT_SHIFT); + d5 = DESCALE(tmp5 + z2 + z4, CONST_BITS + OUT_SHIFT); + d3 = DESCALE(tmp6 + z2 + z3, CONST_BITS + OUT_SHIFT); + d1 = DESCALE(tmp7 + z1o + z4, CONST_BITS + OUT_SHIFT); +} + +__device__ static int quantize(int coef, int quant) +{ + int div = quant << 3; + int mag = coef < 0 ? -coef : coef; + int q = (mag + (div >> 1)) / div; + + return coef < 0 ? -q : q; +} + +__device__ static int +sample(const unsigned char *src, int stride, int x, int y, int step) +{ + int s; + + if (step == 1) + return (int)src[y * stride + x] - 128; + s = src[(y * 2) * stride + x * 2] + src[(y * 2) * stride + x * 2 + 1] + + src[(y * 2 + 1) * stride + x * 2] + src[(y * 2 + 1) * stride + x * 2 + 1]; + return ((s + 2) >> 2) - 128; +} + +extern "C" __global__ void +mjpeg_cuda_fdct(const unsigned char *src, int stride, int blocks_x, int blocks_y, + int step, const unsigned short *quant, short *dst) +{ + int block = (int)blockIdx.x; + int col = (int)threadIdx.x; + int bx, by, row; + int s0, s1, s2, s3, s4, s5, s6, s7; + int c0, c1, c2, c3, c4, c5, c6, c7; + int r0, r1, r2, r3, r4, r5, r6, r7; + int d0, d1, d2, d3, d4, d5, d6, d7; + __shared__ int sh[8 * 9]; + + if (block >= blocks_x * blocks_y) + return; + bx = block % blocks_x; + by = block / blocks_x; + + s0 = s1 = s2 = s3 = s4 = s5 = s6 = s7 = 0; + c0 = c1 = c2 = c3 = c4 = c5 = c6 = c7 = 0; + /* Same order as jfdctint: rows first, then columns. */ + if (col < 8) { + int y = by * 8 + col; + int x = bx * 8; + + s0 = sample(src, stride, x + 0, y, step); + s1 = sample(src, stride, x + 1, y, step); + s2 = sample(src, stride, x + 2, y, step); + s3 = sample(src, stride, x + 3, y, step); + s4 = sample(src, stride, x + 4, y, step); + s5 = sample(src, stride, x + 5, y, step); + s6 = sample(src, stride, x + 6, y, step); + s7 = sample(src, stride, x + 7, y, step); + islow_fdct_pass1(s0, s1, s2, s3, s4, s5, s6, s7, c0, c1, c2, c3, c4, c5, c6, c7); + sh[col * 9 + 0] = c0; + sh[col * 9 + 1] = c1; + sh[col * 9 + 2] = c2; + sh[col * 9 + 3] = c3; + sh[col * 9 + 4] = c4; + sh[col * 9 + 5] = c5; + sh[col * 9 + 6] = c6; + sh[col * 9 + 7] = c7; + } + asm volatile("bar.warp.sync 0xffffffff;"); + if (col >= 8) + return; + + row = col; + r0 = sh[0 * 9 + row]; + r1 = sh[1 * 9 + row]; + r2 = sh[2 * 9 + row]; + r3 = sh[3 * 9 + row]; + r4 = sh[4 * 9 + row]; + r5 = sh[5 * 9 + row]; + r6 = sh[6 * 9 + row]; + r7 = sh[7 * 9 + row]; + islow_fdct_pass2(r0, r1, r2, r3, r4, r5, r6, r7, d0, d1, d2, d3, d4, d5, d6, d7); + + { + short *out = dst + (by * blocks_x + bx) * 64; + + out[0 * 8 + row] = (short)quantize(d0, quant[0 * 8 + row]); + out[1 * 8 + row] = (short)quantize(d1, quant[1 * 8 + row]); + out[2 * 8 + row] = (short)quantize(d2, quant[2 * 8 + row]); + out[3 * 8 + row] = (short)quantize(d3, quant[3 * 8 + row]); + out[4 * 8 + row] = (short)quantize(d4, quant[4 * 8 + row]); + out[5 * 8 + row] = (short)quantize(d5, quant[5 * 8 + row]); + out[6 * 8 + row] = (short)quantize(d6, quant[6 * 8 + row]); + out[7 * 8 + row] = (short)quantize(d7, quant[7 * 8 + row]); + } +} + +__device__ static unsigned mj_ballot(unsigned pred) +{ + unsigned ret; + + asm volatile("{\n\t.reg .pred p;\n\tsetp.ne.u32 p, %1, 0;\n\t" + "vote.sync.ballot.b32 %0, p, 0xffffffff;\n\t}" + : "=r"(ret) : "r"(pred)); + return ret; +} + +__device__ static void syncw(void) +{ + asm volatile("bar.warp.sync 0xffffffff;"); +} + +__device__ static int clz32(unsigned x) +{ + unsigned r; + + asm("clz.b32 %0, %1;" : "=r"(r) : "r"(x)); + return (int)r; +} + +__device__ static int popc32(unsigned x) +{ + unsigned r; + + asm("popc.b32 %0, %1;" : "=r"(r) : "r"(x)); + return (int)r; +} + +__device__ static unsigned +enc_val(const unsigned *decomp, const unsigned *lut, int run, int coef, int lut_off) +{ + unsigned packed = decomp[4096 + coef]; + int nbits = (int)(packed & 0xf); + unsigned val = packed & ~0xf; + unsigned prefix = lut[lut_off + run * 16 + nbits]; + unsigned pn = prefix & 31; + + return (prefix + (unsigned)nbits) | (val >> pn); +} + +__device__ static int +serial_block(const short *block, const int *order, const unsigned *decomp, + const unsigned *lut, unsigned *s_out, int base, int lut_off, int comp) +{ + int n = 0; + int dc = block[order[0]]; + int diff, run, i; + + diff = dc - (int)s_out[256 + comp]; + s_out[256 + comp] = (unsigned)dc; + s_out[base + n++] = enc_val(decomp, lut, 0, diff, lut_off + 257); + run = 0; + for (i = 1; i < 64; i++) { + int c = block[order[i]]; + + if (!c) { + run++; + continue; + } + while (run >= 16) { + s_out[base + n++] = enc_val(decomp, lut, 15, 0, lut_off); + run -= 16; + } + s_out[base + n++] = enc_val(decomp, lut, run, c, lut_off); + run = 0; + } + if (run) + s_out[base + n++] = enc_val(decomp, lut, 16, 0, lut_off); + return n; +} + +__device__ static void +flush_full(unsigned *s_out, unsigned **outp, int *remaining, int tid) +{ + unsigned *out; + int i; + + if (*remaining <= 128) + return; + syncw(); + out = *outp + tid * 4; + for (i = 0; i < 4; i++) + out[i] = s_out[tid * 4 + i]; + *outp += 128; + syncw(); + for (i = 0; i < 4; i++) + s_out[tid * 4 + i] = s_out[128 + tid * 4 + i]; + *remaining -= 128; +} + +__device__ static void +encode_block(const short *block, const int *order, const unsigned *decomp, + const unsigned *lut, unsigned *s_out, unsigned **outp, int *remaining, + int tid, int lut_off, int comp, int *scratch) +{ + int load = tid * 2; + int in_even = block[order[load]]; + int in_odd = block[order[load + 1]]; + unsigned nonzero_mask = (tid == 0) ? 0 : (0xffffffffu >> (32 - tid)); + unsigned bitmap0 = 1u | mj_ballot(in_even); + unsigned bitmap1 = mj_ballot(in_odd); + unsigned pairs = bitmap0 | bitmap1; + int zero_pairs = clz32(pairs & nonzero_mask); + int ze = 2 * (zero_pairs + tid - 32); + int zo, added, even_off, mask, follows; + unsigned even_code, odd_code, even_pres, odd_pres; + int even_sz, odd_sz, total, off; + + if ((0x80000000u >> zero_pairs) > (bitmap1 & nonzero_mask)) + ze += 1; + zo = in_even || !tid ? 0 : ze + 1; + + if (mj_ballot((unsigned)(ze >= 16 || zo >= 16))) { + if (*remaining > 96) + flush_full(s_out, outp, remaining, tid); + added = 0; + if (tid == 0) + added = serial_block(block, order, decomp, lut, s_out, *remaining, lut_off, comp); + if (tid == 0) + scratch[0] = added; + syncw(); + *remaining += scratch[0]; + flush_full(s_out, outp, remaining, tid); + return; + } + + follows = pairs & ~nonzero_mask; + mask = follows ? 0xf : 0; + ze &= mask; + zo &= mask; + + even_off = lut_off; + if (tid == 0) { + int original = in_even; + + even_off += 257; + in_even = original - (int)s_out[256 + comp]; + s_out[256 + comp] = (unsigned)original; + } + if (tid == 31 && in_odd == 0) + zo = 16; + + even_code = enc_val(decomp, lut, ze, in_even, even_off); + odd_code = enc_val(decomp, lut, zo, in_odd, lut_off); + even_sz = (int)(even_code & 31); + odd_sz = (int)(odd_code & 31); + total = even_sz + odd_sz; + if (total <= 27) { + even_code = (unsigned)total | ((odd_code & ~31u) >> even_sz) | (even_code & ~31u); + odd_code = 0; + } + + even_pres = mj_ballot(even_code); + odd_pres = mj_ballot(odd_code); + off = popc32(nonzero_mask & even_pres) + popc32(nonzero_mask & odd_pres); + if (even_code) { + s_out[*remaining + off] = even_code; + if (odd_code) + s_out[*remaining + off + 1] = odd_code; + } + *remaining += popc32(odd_pres) + popc32(even_pres); + flush_full(s_out, outp, remaining, tid); +} + +__device__ static const short * +block_ptr(const short *base, int blocks_x, int bx, int by) +{ + return base + ((by * blocks_x + bx) * 64); +} + +extern "C" __global__ void +mjpeg_cuda_huffman(MJCudaSeg *segs, int nseg, int grid_x, int restart, int mcu_x, int mcu_y, + int chroma420, const short *coef_y, int blocks_x_y, + const short *coef_u, const short *coef_v, int blocks_x_c, + const int *order, const unsigned *decomp, const unsigned *lut, + unsigned *cw, int cw_stride, int *status) +{ + int warp = (int)threadIdx.x >> 5; + int tid = (int)threadIdx.x & 31; + int seg = ((int)blockIdx.y * grid_x + (int)blockIdx.x) * 4 + warp; + __shared__ unsigned shared_out[4][260]; + __shared__ int scratch[4]; + unsigned *s_out = shared_out[warp]; + unsigned *out, *start; + int remaining = 0; + int mcu0, mcu_n, m, mx, my; + + if (seg >= nseg) + return; + + out = cw + seg * cw_stride; + start = out; + if (tid < 3) + s_out[256 + tid] = 0; + + mcu0 = seg * restart; + mcu_n = restart; + if (mcu0 >= mcu_x * mcu_y) + return; + if (mcu0 + mcu_n > mcu_x * mcu_y) + mcu_n = mcu_x * mcu_y - mcu0; + + for (m = 0; m < mcu_n; m++) { + int mcu = mcu0 + m; + + mx = mcu % mcu_x; + my = mcu / mcu_x; + if (chroma420) { + encode_block(block_ptr(coef_y, blocks_x_y, mx * 2, my * 2), order, decomp, lut, s_out, &out, &remaining, tid, 0, 0, scratch + warp); + encode_block(block_ptr(coef_y, blocks_x_y, mx * 2 + 1, my * 2), order, decomp, lut, s_out, &out, &remaining, tid, 0, 0, scratch + warp); + encode_block(block_ptr(coef_y, blocks_x_y, mx * 2, my * 2 + 1), order, decomp, lut, s_out, &out, &remaining, tid, 0, 0, scratch + warp); + encode_block(block_ptr(coef_y, blocks_x_y, mx * 2 + 1, my * 2 + 1), order, decomp, lut, s_out, &out, &remaining, tid, 0, 0, scratch + warp); + encode_block(block_ptr(coef_u, blocks_x_c, mx, my), order, decomp, lut, s_out, &out, &remaining, tid, 514, 1, scratch + warp); + encode_block(block_ptr(coef_v, blocks_x_c, mx, my), order, decomp, lut, s_out, &out, &remaining, tid, 514, 2, scratch + warp); + } else { + encode_block(block_ptr(coef_y, blocks_x_y, mx, my), order, decomp, lut, s_out, &out, &remaining, tid, 0, 0, scratch + warp); + encode_block(block_ptr(coef_u, blocks_x_c, mx, my), order, decomp, lut, s_out, &out, &remaining, tid, 514, 1, scratch + warp); + encode_block(block_ptr(coef_v, blocks_x_c, mx, my), order, decomp, lut, s_out, &out, &remaining, tid, 514, 2, scratch + warp); + } + if ((int)(out - start) + remaining + 128 > cw_stride) { + if (tid == 0) + *status = 1; + segs[seg].codewords = -1; + return; + } + } + + syncw(); + if (tid == 0) { + int i; + + for (i = 0; i < remaining; i++) + out[i] = s_out[i]; + segs[seg].codewords = (int)(out - start) + remaining; + } +} + +__device__ static int +emit_bits(unsigned *remaining_bits, int *byte_count, int *bit_count, + unsigned char *out, unsigned packed, int max_bytes) +{ + unsigned code_word = packed & ~31u; + unsigned code_bit_size = packed & 31u; + + *remaining_bits |= code_word >> *bit_count; + *bit_count += (int)code_bit_size; + if (*bit_count >= 8) { + do { + unsigned out_byte = *remaining_bits >> 24; + + if (*byte_count >= max_bytes) + return -1; + out[(*byte_count)++] = (unsigned char)out_byte; + if (out_byte == 0xff) { + if (*byte_count >= max_bytes) + return -1; + out[(*byte_count)++] = 0; + } + *remaining_bits <<= 8; + *bit_count -= 8; + } while (*bit_count >= 8); + *remaining_bits = code_word << (code_bit_size - *bit_count); + *remaining_bits &= 0xfffffffe << (31 - *bit_count); + } + return 0; +} + +extern "C" __global__ void +mjpeg_cuda_serialize(MJCudaSeg *segs, int nseg, const unsigned *cw, int cw_stride, + unsigned char *bytes, int max_bytes, int *status) +{ + int seg = (int)blockIdx.x * (int)blockDim.x + (int)threadIdx.x; + const unsigned *in; + unsigned char *out; + int nw, bc, bits, i; + unsigned rem; + + if (seg >= nseg) + return; + if (segs[seg].codewords < 0) { + segs[seg].bytes = -1; + return; + } + in = cw + seg * cw_stride; + out = bytes + seg * max_bytes; + nw = segs[seg].codewords; + bc = 0; + bits = 0; + rem = 0; + for (i = 0; i < nw; i++) { + if (emit_bits(&rem, &bc, &bits, out, in[i], max_bytes) < 0) { + segs[seg].bytes = -1; + *status = 1; + return; + } + } + if (emit_bits(&rem, &bc, &bits, out, 0xfe000007u, max_bytes) < 0 || + bc + 2 > max_bytes) { + segs[seg].bytes = -1; + *status = 1; + return; + } + out[bc++] = 0xff; + out[bc++] = (unsigned char)(0xd0 + (seg & 7)); + segs[seg].bytes = bc; +} + +extern "C" __global__ void +mjpeg_cuda_prefix(const MJCudaSeg *segs, int nseg, unsigned *offsets, unsigned *total) +{ + unsigned sum = 0; + int i; + + if (threadIdx.x || blockIdx.x) + return; + for (i = 0; i < nseg; i++) { + int b = segs[i].bytes; + + offsets[i] = sum; + if (b > 0) + sum += (unsigned)b; + } + *total = sum; +} + +extern "C" __global__ void +mjpeg_cuda_compact(const MJCudaSeg *segs, const unsigned *offsets, int nseg, + const unsigned char *src, int max_bytes, unsigned char *dst) +{ + int seg = (int)blockIdx.x * (int)blockDim.x + (int)threadIdx.x; + int b, i; + const unsigned char *in; + unsigned char *out; + + if (seg >= nseg) + return; + b = segs[seg].bytes; + if (b <= 0) + return; + in = src + seg * max_bytes; + out = dst + offsets[seg]; + for (i = 0; i < b; i++) + out[i] = in[i]; +} diff --git a/libavcodec/version.h b/libavcodec/version.h index 7531c6c42a..9b8c267529 100644 --- a/libavcodec/version.h +++ b/libavcodec/version.h @@ -29,7 +29,7 @@ #include "version_major.h" -#define LIBAVCODEC_VERSION_MINOR 14 +#define LIBAVCODEC_VERSION_MINOR 15 #define LIBAVCODEC_VERSION_MICRO 100 #define LIBAVCODEC_VERSION_INT AV_VERSION_INT(LIBAVCODEC_VERSION_MAJOR, \ -- 2.54.0 _______________________________________________ ffmpeg-devel mailing list -- [email protected] To unsubscribe send an email to [email protected]
