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 | 1 + 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, 545 insertions(+) create mode 100644 libavfilter/vf_yuvrgb_cuda.c create mode 100644 libavfilter/vf_yuvrgb_cuda.cu diff --git a/Changelog b/Changelog index dd08217835..531fd76145 100644 --- a/Changelog +++ b/Changelog @@ -18,6 +18,7 @@ version <next>: - 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 5621ff7622..83a4570b8c 100755 --- a/configure +++ b/configure @@ -3562,6 +3562,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 0407c46917..861c0b1053 100644 --- a/doc/filters.texi +++ b/doc/filters.texi @@ -27644,6 +27644,26 @@ JPEG (full) range @anchor{overlay_cuda} @section overlay_cuda +@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 + Overlay one video on top of another. 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.43.0 _______________________________________________ ffmpeg-devel mailing list -- [email protected] To unsubscribe send an email to [email protected]
