PR #24729 opened by rnissl URL: https://code.ffmpeg.org/FFmpeg/FFmpeg/pulls/24729 Patch URL: https://code.ffmpeg.org/FFmpeg/FFmpeg/pulls/24729.patch
This is 2/9 in the GoPro Max review order. It applies on current master by itself. CUDA conversion between planar YUV444 and packed RGB with JPEG or BT.709 full-range matrix. 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 a6dcc98e0b33adb6fa2587ab24ae8b8a370a4314 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Reinhard=20Ni=C3=9Fl?= <[email protected]> Date: Sat, 26 Sep 2026 21:34:53 +0200 Subject: [PATCH] avfilter: add yuvrgb_cuda MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit CUDA conversion between planar YUV444 and packed RGB with JPEG or BT.709 full-range matrix. Co-authored-by: Cursor <[email protected]> Signed-off-by: Reinhard Nißl <[email protected]> --- Changelog | 4 + configure | 2 + doc/filters.texi | 20 ++ libavfilter/Makefile | 2 + libavfilter/allfilters.c | 1 + libavfilter/vf_yuvrgb_cuda.c | 383 ++++++++++++++++++++++++++++++++++ libavfilter/vf_yuvrgb_cuda.cu | 136 ++++++++++++ 7 files changed, 548 insertions(+) create mode 100644 libavfilter/vf_yuvrgb_cuda.c create mode 100644 libavfilter/vf_yuvrgb_cuda.cu diff --git a/Changelog b/Changelog index 1cfa9db1b6..531fd76145 100644 --- a/Changelog +++ b/Changelog @@ -15,6 +15,10 @@ version <next>: - H.264 data partitioning support - DSD (dsd_msbf) encoder +- jpegapp bitstream filter +- MJPEG encoder packed RGB input +- mjpeg_nvjpeg NVIDIA JPEG encoder +- yuvrgb_cuda filter version 9.0: - Extend AMF Color Converter (vf_vpp_amf) HDR capabilities diff --git a/configure b/configure index f686b727cf..351583e918 100755 --- a/configure +++ b/configure @@ -3558,6 +3558,8 @@ transpose_cuda_filter_deps="ffnvcodec" transpose_cuda_filter_deps_any="cuda_nvcc cuda_llvm" overlay_cuda_filter_deps="ffnvcodec" overlay_cuda_filter_deps_any="cuda_nvcc cuda_llvm" +yuvrgb_cuda_filter_deps="ffnvcodec" +yuvrgb_cuda_filter_deps_any="cuda_nvcc cuda_llvm" pad_cuda_filter_deps="ffnvcodec" pad_cuda_filter_deps_any="cuda_nvcc cuda_llvm" diff --git a/doc/filters.texi b/doc/filters.texi index 1a9bbeff7e..d73bf7df1a 100644 --- a/doc/filters.texi +++ b/doc/filters.texi @@ -27849,6 +27849,26 @@ scale_cuda=passthrough=0 @end example @end itemize +@anchor{yuvrgb_cuda} +@section yuvrgb_cuda + +Convert CUDA @code{yuv444p} / @code{yuv444p10} to packed @code{rgb0}, or +@code{rgb0} back to @code{yuv444p}. @ref{scale_cuda} cannot do this +conversion. + +@table @option +@item format +Output software format: @code{rgb0} or @code{yuv444p}. +@item matrix +@samp{jpeg} / @samp{bt601} (ITU T.871) or @samp{bt709}. Default @samp{jpeg}. +@item range +@samp{jpeg}/@samp{pc} (full, default) or @samp{mpeg}/@samp{tv}. +@end table + +@example +scale_cuda=format=yuv444p10,yuvrgb_cuda=format=rgb0:matrix=bt709:range=pc +@end example + @section thumbnail_cuda Select the most representative frame in a given sequence of consecutive frames using CUDA. diff --git a/libavfilter/Makefile b/libavfilter/Makefile index ceea328e67..7ceddf552c 100644 --- a/libavfilter/Makefile +++ b/libavfilter/Makefile @@ -596,6 +596,8 @@ OBJS-$(CONFIG_YADIF_VIDEOTOOLBOX_FILTER) += vf_yadif_videotoolbox.o \ metal/utils.o \ yadif_common.o OBJS-$(CONFIG_YAEPBLUR_FILTER) += vf_yaepblur.o +OBJS-$(CONFIG_YUVRGB_CUDA_FILTER) += vf_yuvrgb_cuda.o vf_yuvrgb_cuda.ptx.o \ + cuda/load_helper.o OBJS-$(CONFIG_ZMQ_FILTER) += f_zmq.o OBJS-$(CONFIG_ZOOMPAN_FILTER) += vf_zoompan.o OBJS-$(CONFIG_ZSCALE_FILTER) += vf_zscale.o diff --git a/libavfilter/allfilters.c b/libavfilter/allfilters.c index 3b93fe12f6..82a3f001e9 100644 --- a/libavfilter/allfilters.c +++ b/libavfilter/allfilters.c @@ -556,6 +556,7 @@ extern const FFFilter ff_vf_yadif; extern const FFFilter ff_vf_yadif_cuda; extern const FFFilter ff_vf_yadif_videotoolbox; extern const FFFilter ff_vf_yaepblur; +extern const FFFilter ff_vf_yuvrgb_cuda; extern const FFFilter ff_vf_zmq; extern const FFFilter ff_vf_zoompan; extern const FFFilter ff_vf_zscale; diff --git a/libavfilter/vf_yuvrgb_cuda.c b/libavfilter/vf_yuvrgb_cuda.c new file mode 100644 index 0000000000..49740f37b2 --- /dev/null +++ b/libavfilter/vf_yuvrgb_cuda.c @@ -0,0 +1,383 @@ +/* + * CUDA packed RGB <-> planar YUV444 conversion. + * + * 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 "libavutil/common.h" +#include "libavutil/internal.h" +#include "libavutil/cuda_check.h" +#include "libavutil/hwcontext.h" +#include "libavutil/hwcontext_cuda_internal.h" +#include "libavutil/opt.h" +#include "libavutil/pixdesc.h" + +#include "avfilter.h" +#include "cuda/load_helper.h" +#include "filters.h" +#include "video.h" + +#define CHECK_CU(x) FF_CUDA_CHECK_DL(ctx, s->hwctx->internal->cuda_dl, x) +#define DIV_UP(a, b) (((a) + (b) - 1) / (b)) +#define BLOCKX 32 +#define BLOCKY 16 + +enum { + YUVRGB_MATRIX_JPEG = 0, + YUVRGB_MATRIX_BT709, +}; + +typedef struct YUVRGBCUDAContext { + const AVClass *class; + + AVCUDADeviceContext *hwctx; + AVBufferRef *hw_frames_ctx; + CUmodule cu_module; + CUfunction cu_yuv_to_rgb0; + CUfunction cu_rgb0_to_yuv8; + CUstream cu_stream; + + int matrix; + int range; + enum AVPixelFormat format; + enum AVPixelFormat in_fmt; + enum AVPixelFormat out_fmt; + int depth; + int packed_in; + int packed_out; +} YUVRGBCUDAContext; + +static int format_ok(enum AVPixelFormat fmt) +{ + switch (fmt) { + case AV_PIX_FMT_YUV444P: + case AV_PIX_FMT_YUV444P10: + case AV_PIX_FMT_RGB0: + return 1; + default: + return 0; + } +} + +static av_cold void yuvrgb_cuda_uninit(AVFilterContext *ctx) +{ + YUVRGBCUDAContext *s = ctx->priv; + + if (s->hwctx && s->cu_module) { + CudaFunctions *cu = s->hwctx->internal->cuda_dl; + CUcontext dummy; + CHECK_CU(cu->cuCtxPushCurrent(s->hwctx->cuda_ctx)); + CHECK_CU(cu->cuModuleUnload(s->cu_module)); + s->cu_module = NULL; + CHECK_CU(cu->cuCtxPopCurrent(&dummy)); + } + av_buffer_unref(&s->hw_frames_ctx); +} + +static av_cold int yuvrgb_cuda_load(AVFilterContext *ctx) +{ + YUVRGBCUDAContext *s = ctx->priv; + CudaFunctions *cu = s->hwctx->internal->cuda_dl; + CUcontext dummy, cuda_ctx = s->hwctx->cuda_ctx; + int ret; + + extern const unsigned char ff_vf_yuvrgb_cuda_ptx_data[]; + extern const unsigned int ff_vf_yuvrgb_cuda_ptx_len; + + ret = CHECK_CU(cu->cuCtxPushCurrent(cuda_ctx)); + if (ret < 0) + return ret; + + ret = ff_cuda_load_module(ctx, s->hwctx, &s->cu_module, + ff_vf_yuvrgb_cuda_ptx_data, + ff_vf_yuvrgb_cuda_ptx_len); + if (ret < 0) + goto fail; + + ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_yuv_to_rgb0, s->cu_module, + "yuv_to_rgb0")); + if (ret < 0) + goto fail; + ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_rgb0_to_yuv8, s->cu_module, + "rgb0_to_yuv8")); + if (ret < 0) + goto fail; + +fail: + CHECK_CU(cu->cuCtxPopCurrent(&dummy)); + return ret; +} + +static int init_hwframe_ctx(AVFilterContext *ctx, AVBufferRef *device_ref, + int width, int height) +{ + YUVRGBCUDAContext *s = ctx->priv; + AVBufferRef *out_ref; + AVHWFramesContext *out_ctx; + int ret; + + out_ref = av_hwframe_ctx_alloc(device_ref); + if (!out_ref) + return AVERROR(ENOMEM); + out_ctx = (AVHWFramesContext *)out_ref->data; + out_ctx->format = AV_PIX_FMT_CUDA; + out_ctx->sw_format = s->out_fmt; + out_ctx->width = FFALIGN(width, 32); + out_ctx->height = FFALIGN(height, 32); + + ret = av_hwframe_ctx_init(out_ref); + if (ret < 0) { + av_buffer_unref(&out_ref); + return ret; + } + + av_buffer_unref(&s->hw_frames_ctx); + s->hw_frames_ctx = out_ref; + return 0; +} + +static av_cold int yuvrgb_cuda_config_output(AVFilterLink *outlink) +{ + AVFilterContext *ctx = outlink->src; + AVFilterLink *inlink = ctx->inputs[0]; + FilterLink *inl = ff_filter_link(inlink); + FilterLink *outl = ff_filter_link(outlink); + YUVRGBCUDAContext *s = ctx->priv; + AVHWFramesContext *in_fc; + const AVPixFmtDescriptor *desc; + int ret; + + if (!inl->hw_frames_ctx) { + av_log(ctx, AV_LOG_ERROR, "CUDA frames context required\n"); + return AVERROR(EINVAL); + } + + in_fc = (AVHWFramesContext *)inl->hw_frames_ctx->data; + s->in_fmt = in_fc->sw_format; + if (!format_ok(s->in_fmt)) { + av_log(ctx, AV_LOG_ERROR, "Unsupported input sw_format %s\n", + av_get_pix_fmt_name(s->in_fmt)); + return AVERROR(ENOSYS); + } + + if (s->format == AV_PIX_FMT_NONE) { + if (s->in_fmt == AV_PIX_FMT_RGB0) + s->out_fmt = AV_PIX_FMT_YUV444P; + else + s->out_fmt = AV_PIX_FMT_RGB0; + } else { + s->out_fmt = s->format; + } + + if (!format_ok(s->out_fmt)) { + av_log(ctx, AV_LOG_ERROR, "Unsupported output format %s\n", + av_get_pix_fmt_name(s->out_fmt)); + return AVERROR(ENOSYS); + } + + s->packed_in = s->in_fmt == AV_PIX_FMT_RGB0; + s->packed_out = s->out_fmt == AV_PIX_FMT_RGB0; + if (s->packed_in == s->packed_out) { + av_log(ctx, AV_LOG_ERROR, "Need a YUV444 <-> rgb0 conversion, got %s -> %s\n", + av_get_pix_fmt_name(s->in_fmt), av_get_pix_fmt_name(s->out_fmt)); + return AVERROR(EINVAL); + } + if (s->packed_in && s->out_fmt != AV_PIX_FMT_YUV444P) { + av_log(ctx, AV_LOG_ERROR, "rgb0 can only convert to yuv444p\n"); + return AVERROR(ENOSYS); + } + + desc = av_pix_fmt_desc_get(s->packed_in ? s->out_fmt : s->in_fmt); + s->depth = desc->comp[0].depth; + + outlink->w = inlink->w; + outlink->h = inlink->h; + outlink->sample_aspect_ratio = inlink->sample_aspect_ratio; + + ret = init_hwframe_ctx(ctx, in_fc->device_ref, outlink->w, outlink->h); + if (ret < 0) + return ret; + + outl->hw_frames_ctx = av_buffer_ref(s->hw_frames_ctx); + if (!outl->hw_frames_ctx) + return AVERROR(ENOMEM); + + s->hwctx = in_fc->device_ctx->hwctx; + s->cu_stream = s->hwctx->stream; + + ret = yuvrgb_cuda_load(ctx); + if (ret < 0) + return ret; + + av_log(ctx, AV_LOG_VERBOSE, "yuvrgb_cuda %s -> %s matrix=%s range=%s\n", + av_get_pix_fmt_name(s->in_fmt), av_get_pix_fmt_name(s->out_fmt), + s->matrix == YUVRGB_MATRIX_BT709 ? "bt709" : "jpeg", + s->range == AVCOL_RANGE_MPEG ? "mpeg" : "jpeg"); + return 0; +} + +static void matrix_kr_kb(int matrix, float *kr, float *kb) +{ + if (matrix == YUVRGB_MATRIX_BT709) { + *kr = 0.2126f; + *kb = 0.0722f; + } else { + *kr = 0.299f; + *kb = 0.114f; + } +} + +static int yuvrgb_cuda_filter_frame(AVFilterLink *inlink, AVFrame *in) +{ + AVFilterContext *ctx = inlink->dst; + AVFilterLink *outlink = ctx->outputs[0]; + YUVRGBCUDAContext *s = ctx->priv; + CudaFunctions *cu = s->hwctx->internal->cuda_dl; + AVFrame *out = NULL; + CUcontext dummy; + float kr, kb; + int ret, pushed = 0, full_range, w, h; + + matrix_kr_kb(s->matrix, &kr, &kb); + full_range = s->range != AVCOL_RANGE_MPEG; + w = inlink->w; + h = inlink->h; + + out = ff_get_video_buffer(outlink, outlink->w, outlink->h); + if (!out) { + ret = AVERROR(ENOMEM); + goto fail; + } + + ret = CHECK_CU(cu->cuCtxPushCurrent(s->hwctx->cuda_ctx)); + if (ret < 0) + goto fail; + pushed = 1; + + if (s->packed_out) { + CUdeviceptr y = (CUdeviceptr)in->data[0]; + CUdeviceptr u = (CUdeviceptr)in->data[1]; + CUdeviceptr v = (CUdeviceptr)in->data[2]; + CUdeviceptr d = (CUdeviceptr)out->data[0]; + int yp = in->linesize[0], up = in->linesize[1], vp = in->linesize[2]; + int dp = out->linesize[0]; + void *args[] = { + &y, &u, &v, &yp, &up, &vp, &d, &dp, + &w, &h, &s->depth, &full_range, &kr, &kb + }; + ret = CHECK_CU(cu->cuLaunchKernel(s->cu_yuv_to_rgb0, + DIV_UP(w, BLOCKX), DIV_UP(h, BLOCKY), 1, + BLOCKX, BLOCKY, 1, 0, s->cu_stream, args, NULL)); + } else { + CUdeviceptr src = (CUdeviceptr)in->data[0]; + CUdeviceptr y = (CUdeviceptr)out->data[0]; + CUdeviceptr u = (CUdeviceptr)out->data[1]; + CUdeviceptr v = (CUdeviceptr)out->data[2]; + int sp = in->linesize[0]; + int yp = out->linesize[0], up = out->linesize[1], vp = out->linesize[2]; + void *args[] = { + &src, &sp, &y, &u, &v, &yp, &up, &vp, &w, &h, &kr, &kb + }; + ret = CHECK_CU(cu->cuLaunchKernel(s->cu_rgb0_to_yuv8, + DIV_UP(w, BLOCKX), DIV_UP(h, BLOCKY), 1, + BLOCKX, BLOCKY, 1, 0, s->cu_stream, args, NULL)); + } + if (ret < 0) + goto fail; + + ret = CHECK_CU(cu->cuCtxPopCurrent(&dummy)); + if (ret < 0) + goto fail; + pushed = 0; + + ret = av_frame_copy_props(out, in); + if (ret < 0) + goto fail; + + if (s->packed_out) { + out->color_range = AVCOL_RANGE_JPEG; + out->colorspace = AVCOL_SPC_RGB; + } else { + out->color_range = AVCOL_RANGE_JPEG; + out->colorspace = s->matrix == YUVRGB_MATRIX_BT709 + ? AVCOL_SPC_BT709 : AVCOL_SPC_BT470BG; + } + + av_frame_free(&in); + return ff_filter_frame(outlink, out); + +fail: + if (pushed) + CHECK_CU(cu->cuCtxPopCurrent(&dummy)); + av_frame_free(&in); + av_frame_free(&out); + return ret; +} + +#define OFFSET(x) offsetof(YUVRGBCUDAContext, x) +#define FLAGS (AV_OPT_FLAG_FILTERING_PARAM | AV_OPT_FLAG_VIDEO_PARAM) +static const AVOption yuvrgb_cuda_options[] = { + { "format", "output sw format (rgb0 or yuv444p)", OFFSET(format), + AV_OPT_TYPE_PIXEL_FMT, { .i64 = AV_PIX_FMT_NONE }, -1, INT_MAX, FLAGS }, + { "matrix", "YCbCr matrix", OFFSET(matrix), AV_OPT_TYPE_INT, + { .i64 = YUVRGB_MATRIX_JPEG }, 0, YUVRGB_MATRIX_BT709, FLAGS, .unit = "matrix" }, + { "jpeg", "BT.601 full range (ITU T.871)", 0, AV_OPT_TYPE_CONST, + { .i64 = YUVRGB_MATRIX_JPEG }, 0, 0, FLAGS, .unit = "matrix" }, + { "bt601", "Same as jpeg", 0, AV_OPT_TYPE_CONST, + { .i64 = YUVRGB_MATRIX_JPEG }, 0, 0, FLAGS, .unit = "matrix" }, + { "bt709", "BT.709", 0, AV_OPT_TYPE_CONST, + { .i64 = YUVRGB_MATRIX_BT709 }, 0, 0, FLAGS, .unit = "matrix" }, + { "range", "YUV range (YUV input or output)", OFFSET(range), AV_OPT_TYPE_INT, + { .i64 = AVCOL_RANGE_JPEG }, 0, AVCOL_RANGE_NB - 1, FLAGS, .unit = "range" }, + { "jpeg", "full range", 0, AV_OPT_TYPE_CONST, { .i64 = AVCOL_RANGE_JPEG }, 0, 0, FLAGS, .unit = "range" }, + { "pc", "full range", 0, AV_OPT_TYPE_CONST, { .i64 = AVCOL_RANGE_JPEG }, 0, 0, FLAGS, .unit = "range" }, + { "mpeg", "limited range", 0, AV_OPT_TYPE_CONST, { .i64 = AVCOL_RANGE_MPEG }, 0, 0, FLAGS, .unit = "range" }, + { "tv", "limited range", 0, AV_OPT_TYPE_CONST, { .i64 = AVCOL_RANGE_MPEG }, 0, 0, FLAGS, .unit = "range" }, + { NULL }, +}; + +AVFILTER_DEFINE_CLASS(yuvrgb_cuda); + +static const AVFilterPad yuvrgb_cuda_inputs[] = { + { + .name = "default", + .type = AVMEDIA_TYPE_VIDEO, + .filter_frame = yuvrgb_cuda_filter_frame, + }, +}; + +static const AVFilterPad yuvrgb_cuda_outputs[] = { + { + .name = "default", + .type = AVMEDIA_TYPE_VIDEO, + .config_props = yuvrgb_cuda_config_output, + }, +}; + +const FFFilter ff_vf_yuvrgb_cuda = { + .p.name = "yuvrgb_cuda", + .p.description = NULL_IF_CONFIG_SMALL("CUDA YUV444 <-> rgb0 conversion"), + .p.priv_class = &yuvrgb_cuda_class, + .priv_size = sizeof(YUVRGBCUDAContext), + .uninit = yuvrgb_cuda_uninit, + FILTER_INPUTS(yuvrgb_cuda_inputs), + FILTER_OUTPUTS(yuvrgb_cuda_outputs), + FILTER_SINGLE_PIXFMT(AV_PIX_FMT_CUDA), + .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE, +}; diff --git a/libavfilter/vf_yuvrgb_cuda.cu b/libavfilter/vf_yuvrgb_cuda.cu new file mode 100644 index 0000000000..410374b322 --- /dev/null +++ b/libavfilter/vf_yuvrgb_cuda.cu @@ -0,0 +1,136 @@ +/* + * CUDA YUV <-> packed RGB conversion (JPEG / BT.601 / BT.709). + * + * 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 sat_u8(float v) +{ + return (unsigned char)fminf(fmaxf(v + 0.5f, 0.f), 255.f); +} + +__device__ static inline void ycbcr_to_rgb(float y, float cb, float cr, + float kr, float kb, + unsigned char *r, unsigned char *g, + unsigned char *b) +{ + float kg = 1.f - kr - kb; + float cb2 = cb - 128.f; + float cr2 = cr - 128.f; + float rf = y + 2.f * (1.f - kr) * cr2; + float bf = y + 2.f * (1.f - kb) * cb2; + float gf = y - (2.f * kb * (1.f - kb) / kg) * cb2 + - (2.f * kr * (1.f - kr) / kg) * cr2; + *r = sat_u8(rf); + *g = sat_u8(gf); + *b = sat_u8(bf); +} + +__device__ static inline void rgb_to_ycbcr(float r, float g, float b, + float kr, float kb, + unsigned char *y, unsigned char *cb, + unsigned char *cr) +{ + float kg = 1.f - kr - kb; + float yf = kr * r + kg * g + kb * b; + *y = sat_u8(yf); + *cb = sat_u8(128.f + (b - yf) / (2.f * (1.f - kb))); + *cr = sat_u8(128.f + (r - yf) / (2.f * (1.f - kr))); +} + +/* Full-range n-bit → 8-bit: (v + 1<<(d-9)) >> (d-8). 10-bit 512 → 128, + * not 512*255/1023 ≈ 127.56, which pulled R/B down ~0.8 on grass. */ +__device__ static inline float full_to_u8(float v, int depth) +{ + int sh = depth - 8; + if (sh <= 0) + return v; + return floorf((v + (float)(1 << (sh - 1))) / (float)(1 << sh)); +} + +__device__ static inline float luma_to_full(float y, int depth, int full_range) +{ + if (full_range) + return full_to_u8(y, depth); + float offset = (float)(16 << (depth - 8)); + float scale = (float)(219 << (depth - 8)); + return (y - offset) * 255.f / scale; +} + +__device__ static inline float chroma_to_full(float c, int depth, int full_range) +{ + if (full_range) + return full_to_u8(c, depth); + float mid = (float)(1 << (depth - 1)); + float scale = (float)(224 << (depth - 8)); + return (c - mid) * 255.f / scale + 128.f; +} + +__global__ void yuv_to_rgb0(const unsigned char *ypl, const unsigned char *upl, + const unsigned char *vpl, + int ypitch, int upitch, int vpitch, + unsigned char *dst, int dpitch, + int width, int height, int depth, int full_range, + float kr, float kb) +{ + int x = blockIdx.x * blockDim.x + threadIdx.x; + int y = blockIdx.y * blockDim.y + threadIdx.y; + if (x >= width || y >= height) + return; + + float Y, Cb, Cr; + if (depth > 8) { + const unsigned short *ys = (const unsigned short *)(ypl + y * ypitch); + const unsigned short *us = (const unsigned short *)(upl + y * upitch); + const unsigned short *vs = (const unsigned short *)(vpl + y * vpitch); + Y = luma_to_full((float)ys[x], depth, full_range); + Cb = chroma_to_full((float)us[x], depth, full_range); + Cr = chroma_to_full((float)vs[x], depth, full_range); + } else { + Y = luma_to_full((float)ypl[y * ypitch + x], depth, full_range); + Cb = chroma_to_full((float)upl[y * upitch + x], depth, full_range); + Cr = chroma_to_full((float)vpl[y * vpitch + x], depth, full_range); + } + + unsigned char *p = dst + y * dpitch + x * 4; + ycbcr_to_rgb(Y, Cb, Cr, kr, kb, &p[0], &p[1], &p[2]); + p[3] = 0; +} + +__global__ void rgb0_to_yuv8(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, float kr, float kb) +{ + 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 *p = src + y * spitch + x * 4; + rgb_to_ycbcr((float)p[0], (float)p[1], (float)p[2], kr, kb, + &ypl[y * ypitch + x], + &upl[y * upitch + x], + &vpl[y * vpitch + x]); +} + +} -- 2.52.0 _______________________________________________ ffmpeg-devel mailing list -- [email protected] To unsubscribe send an email to [email protected]
