PR #25083 opened by rnissl URL: https://code.ffmpeg.org/FFmpeg/FFmpeg/pulls/25083 Patch URL: https://code.ffmpeg.org/FFmpeg/FFmpeg/pulls/25083.patch
Rebased scale_cuda YUV/RGB conversion so it can go through CI and review. #22340 remains Zhao Zhili's pull request and still conflicts with master in tests/fate/filter-video.mak. This branch is not pushed onto that head. Zhao Zhili's conversion and FATE commits and Timo Rothenpieler's documentation keep their authors. Two later commits round the YUV/RGB stores and describe the YUV matrix on the output link. Head is b3337e205a, based on 045dea2a34. ## Test plan - [ ] CI on this pull request - [ ] fate-filter-scale-cuda tests in the branch - [ ] Review of the rounding and the output color matrix From 9428be7cc8ca30f13493aab83c0099abd14f386e Mon Sep 17 00:00:00 2001 From: Zhao Zhili <[email protected]> Date: Mon, 2 Mar 2026 00:40:10 +0800 Subject: [PATCH 1/6] avfilter/vf_scale_cuda: add YUV to RGB color space conversion When the input colorspace is unspecified, BT.601 (SMPTE 170M) is used as the default, matching swscale behavior. Both limited and full range inputs are handled correctly. The conversion kernels use direct texture reads for luma and simple coordinate mapping for chroma upsampling, avoiding unnecessary Subsample_Nearest overhead since color conversion operates at the same resolution. Simultaneous resize and color conversion is rejected with an error directing users to chain separate scale_cuda instances. --- libavfilter/vf_scale_cuda.c | 119 +++++++++++++++++++++++-- libavfilter/vf_scale_cuda.cu | 165 +++++++++++++++++++++++++++++++++++ libavfilter/vf_scale_cuda.h | 8 ++ 3 files changed, 285 insertions(+), 7 deletions(-) diff --git a/libavfilter/vf_scale_cuda.c b/libavfilter/vf_scale_cuda.c index 747cc624b3..bab7b35cc5 100644 --- a/libavfilter/vf_scale_cuda.c +++ b/libavfilter/vf_scale_cuda.c @@ -21,11 +21,13 @@ */ #include <float.h> +#include <stdbool.h> #include <stdio.h> #include <string.h> #include "libavutil/avassert.h" #include "libavutil/common.h" +#include "libavutil/csp.h" #include "libavutil/hwcontext.h" #include "libavutil/hwcontext_cuda_internal.h" #include "libavutil/cuda_check.h" @@ -38,6 +40,7 @@ #include "libswscale/filters.h" #include "avfilter.h" +#include "colorspace.h" #include "filters.h" #include "scale_eval.h" #include "video.h" @@ -191,6 +194,11 @@ typedef struct CUDAScaleContext { int use_filters_opt; /* -1 for auto */ float param; + + bool yuv2rgb; + enum AVColorSpace colorspace; ///< resolved colorspace for conversion + enum AVColorRange color_range; ///< resolved color range for conversion + CUDAScaleColorMatrix color_matrix; } CUDAScaleContext; /* @@ -457,6 +465,38 @@ static const char* get_format_name(enum AVPixelFormat fmt) return NULL; } +static int compute_yuv2rgb_matrix(void *log_ctx, CUDAScaleColorMatrix *mat, + enum AVColorSpace colorspace, + bool limited_range) +{ + const AVLumaCoefficients *coeffs = av_csp_luma_coeffs_from_avcsp(colorspace); + float y_scale = 1.0f; + float uv_scale = 1.0f; + double rgb2yuv[3][3], yuv2rgb[3][3]; + + if (!coeffs) { + av_log(log_ctx, AV_LOG_ERROR, + "Unsupported colorspace %d for YUV->RGB conversion\n", colorspace); + return AVERROR(EINVAL); + } + + ff_fill_rgb2yuv_table(coeffs, rgb2yuv); + ff_matrix_invert_3x3(rgb2yuv, yuv2rgb); + + if (limited_range) { + y_scale = 255.0f / (235.0f - 16.0f); + uv_scale = 255.0f / (240.0f - 16.0f); + } + + for (int i = 0; i < 3; i++) { + mat->m[i][0] = (float)yuv2rgb[i][0] * y_scale; + mat->m[i][1] = (float)yuv2rgb[i][1] * uv_scale; + mat->m[i][2] = (float)yuv2rgb[i][2] * uv_scale; + } + + return 0; +} + static av_cold void set_format_info(AVFilterContext *ctx, enum AVPixelFormat in_format, enum AVPixelFormat out_format) { CUDAScaleContext *s = ctx->priv; @@ -517,6 +557,52 @@ static av_cold int init_processing_chain(AVFilterContext *ctx, int in_width, int } set_format_info(ctx, in_format, out_format); + + s->yuv2rgb = !(s->in_desc->flags & AV_PIX_FMT_FLAG_RGB) && + (s->out_desc->flags & AV_PIX_FMT_FLAG_RGB); + if (s->yuv2rgb) { + AVFilterLink *inlink = ctx->inputs[0]; + + if (in_width != out_width || in_height != out_height) { + av_log(ctx, AV_LOG_ERROR, + "Resizing is not supported during YUV->RGB conversion, " + "use a separate scale_cuda instance for resizing\n"); + return AVERROR(ENOSYS); + } + + enum AVColorSpace cs = inlink->colorspace; + bool limited_range; + + /* When colorspace is unspecified, default to BT.601 (SMPTE 170M) + * to match swscale behavior (SWS_CS_DEFAULT = ITU601). + */ + if (cs == AVCOL_SPC_UNSPECIFIED) + cs = AVCOL_SPC_SMPTE170M; + + limited_range = (inlink->color_range != AVCOL_RANGE_JPEG); + + ret = compute_yuv2rgb_matrix(ctx, &s->color_matrix, cs, limited_range); + if (ret < 0) + return ret; + + s->colorspace = cs; + s->color_range = limited_range ? AVCOL_RANGE_MPEG : AVCOL_RANGE_JPEG; + + /* YUV->RGB only supports Nearest for now. Forcing it here also makes + * the generic filter path get disabled below. */ + if (s->interp_algo != INTERP_ALGO_DEFAULT && + s->interp_algo != INTERP_ALGO_NEAREST) + av_log(ctx, AV_LOG_WARNING, + "YUV->RGB conversion only supports nearest interpolation, " + "ignoring interp_algo\n"); + s->interp_algo = INTERP_ALGO_NEAREST; + + av_log(ctx, AV_LOG_VERBOSE, + "YUV->RGB conversion enabled (%s, %s range)\n", + av_color_space_name(cs), + limited_range ? "limited" : "full"); + } + if (s->passthrough && in_width == out_width && in_height == out_height && in_format == out_format) { s->frames_ctx = av_buffer_ref(inl->hw_frames_ctx); if (!s->frames_ctx) @@ -671,12 +757,16 @@ static av_cold int cudascale_load_functions(AVFilterContext *ctx) if (ret < 0) goto unsupported; - snprintf(buf, sizeof(buf), "Subsample_%s_%s_%s_uv", fixed_infix, - in_fmt_name, out_fmt_name); - ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_fixed_uv, - s->cu_module, buf)); - if (ret < 0) - goto unsupported; + /* YUV->RGB produces all output components in a single kernel, + * there is no separate chroma pass. */ + if (!s->yuv2rgb) { + snprintf(buf, sizeof(buf), "Subsample_%s_%s_%s_uv", fixed_infix, + in_fmt_name, out_fmt_name); + ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_fixed_uv, + s->cu_module, buf)); + if (ret < 0) + goto unsupported; + } } goto fail; @@ -1146,6 +1236,7 @@ static int call_resize_kernel(AVFilterContext *ctx, CUfunction func, .src_depth = s->in_desc->comp[0].depth, .src_storage_max = ((1U << s->in_desc->comp[0].depth) - 1) << s->in_desc->comp[0].shift, + .color_matrix = s->color_matrix, }; if (filter) { @@ -1154,6 +1245,11 @@ static int call_resize_kernel(AVFilterContext *ctx, CUfunction func, params.filter_size = filter->filter_size; } + if (s->yuv2rgb) { + params.log2_chroma_w = s->in_desc->log2_chroma_w; + params.log2_chroma_h = s->in_desc->log2_chroma_h; + } + void *args[] = { ¶ms }; return CHECK_CU(cu->cuLaunchKernel(func, @@ -1169,7 +1265,11 @@ static int scalecuda_resize(AVFilterContext *ctx, const CUDAScaleFilter *filter = NULL; const CUDAScaleFilter *filter_uv = NULL; CUfunction func, func_uv; - int mpeg_range = in->color_range != AVCOL_RANGE_JPEG; + /* The color matrix was computed at config time, so for yuv2rgb use the + * resolved color range to keep mpeg_range consistent with it. For plain + * scaling, follow the frame the texture was mapped from. */ + int mpeg_range = s->yuv2rgb ? (s->color_range != AVCOL_RANGE_JPEG) + : (in->color_range != AVCOL_RANGE_JPEG); int ret; int out_planes = s->out_planes; @@ -1290,6 +1390,11 @@ static int cudascale_scale(AVFilterContext *ctx, AVFrame *out, AVFrame *in, if (ret < 0) goto fail; + if (s->yuv2rgb) { + out->colorspace = AVCOL_SPC_RGB; + out->color_range = AVCOL_RANGE_JPEG; + } + out->crop_top = 0; out->crop_bottom = 0; out->crop_left = 0; diff --git a/libavfilter/vf_scale_cuda.cu b/libavfilter/vf_scale_cuda.cu index 7243777651..2a13ffd906 100644 --- a/libavfilter/vf_scale_cuda.cu +++ b/libavfilter/vf_scale_cuda.cu @@ -1113,6 +1113,128 @@ __device__ static inline T Subsample_Nearest(cudaTextureObject_t tex, return tex2D<T>(tex, xi, yi); } +__device__ static inline float clamp_f(float x, float lo, float hi) +{ + return x < lo ? lo : (x > hi ? hi : x); +} + +// YUV to RGB CONVERSION +// R = m[0][0]*(Y-Yoff) + m[0][1]*(U-128) + m[0][2]*(V-128) +// G = m[1][0]*(Y-Yoff) + m[1][1]*(U-128) + m[1][2]*(V-128) +// B = m[2][0]*(Y-Yoff) + m[2][1]*(U-128) + m[2][2]*(V-128) +// channel_order: +// 0 = BGRA (B,G,R,A), 1 = RGBA (R,G,B,A), +// 2 = BGR0 (B,G,R,0), 3 = RGB0 (R,G,B,0) +__device__ static inline uchar4 yuv8_to_rgba_generic( + float fy, float fu, float fv, + const CUDAScaleColorMatrix &mat, + int channel_order) +{ + float r = mat.m[0][0] * fy + mat.m[0][1] * fu + mat.m[0][2] * fv; + float g = mat.m[1][0] * fy + mat.m[1][1] * fu + mat.m[1][2] * fv; + float b = mat.m[2][0] * fy + mat.m[2][1] * fu + mat.m[2][2] * fv; + + uchar rc = (uchar)clamp_f(r, 0.0f, 255.0f); + uchar gc = (uchar)clamp_f(g, 0.0f, 255.0f); + uchar bc = (uchar)clamp_f(b, 0.0f, 255.0f); + + switch (channel_order) { + case 0: return make_uchar4(bc, gc, rc, 0xFF); // BGRA + case 1: return make_uchar4(rc, gc, bc, 0xFF); // RGBA + case 2: return make_uchar4(bc, gc, rc, 0x00); // BGR0 + case 3: return make_uchar4(rc, gc, bc, 0x00); // RGB0 + default: return make_uchar4(bc, gc, rc, 0xFF); + } +} + +// YUV to RGB kernel for semiplanar input (NV12, P010, NV16, P210) +// No resize: src and dst have identical dimensions. +template<typename in_y_T, typename in_uv_T, int channel_order> +__device__ void YuvToRgb_semiplanar(CUDAScaleKernelParams params) +{ + int xo = blockIdx.x * blockDim.x + threadIdx.x; + int yo = blockIdx.y * blockDim.y + threadIdx.y; + if (yo >= params.dst_height || xo >= params.dst_width) return; + + // Direct 1:1 Y plane read + in_y_T y_raw = tex2D<in_y_T>( + params.src_tex[0], + xo + params.src_left + 0.5f, + yo + params.src_top + 0.5f); + + // Chroma: map luma coords to chroma coords for subsampled formats + float cx = ((float)xo + 0.5f) / (float)(1 << params.log2_chroma_w) + + (float)(params.src_left >> params.log2_chroma_w); + float cy = ((float)yo + 0.5f) / (float)(1 << params.log2_chroma_h) + + (float)(params.src_top >> params.log2_chroma_h); + in_uv_T uv_raw = tex2D<in_uv_T>(params.src_tex[1], cx, cy); + + float fy, fu, fv; + if (sizeof(in_y_T) == 1) { + fy = (float)(int)y_raw; + fu = (float)(int)uv_raw.x; + fv = (float)(int)uv_raw.y; + } else { + // 10/16-bit semiplanar: normalize to 8-bit range by dividing by 256 + fy = (float)(int)y_raw / 256.0f; + fu = (float)(int)uv_raw.x / 256.0f; + fv = (float)(int)uv_raw.y / 256.0f; + } + + if (params.mpeg_range) fy -= 16.0f; + fu -= 128.0f; + fv -= 128.0f; + + uchar4 *dst = (uchar4*)params.dst[0]; + int dst_pitch = params.dst_pitch / sizeof(uchar4); + dst[yo * dst_pitch + xo] = yuv8_to_rgba_generic(fy, fu, fv, params.color_matrix, channel_order); +} + +// YUV to RGB kernel for planar input (YUV420P, YUV422P, YUV444P, etc) +// No resize: src and dst have identical dimensions. +template<typename in_y_T, typename in_uv_T, int channel_order> +__device__ void YuvToRgb_planar(CUDAScaleKernelParams params) +{ + int xo = blockIdx.x * blockDim.x + threadIdx.x; + int yo = blockIdx.y * blockDim.y + threadIdx.y; + if (yo >= params.dst_height || xo >= params.dst_width) return; + + // Direct 1:1 Y plane read + in_y_T y_raw = tex2D<in_y_T>( + params.src_tex[0], + xo + params.src_left + 0.5f, + yo + params.src_top + 0.5f); + + // Chroma: map luma coords to chroma coords for subsampled formats + float cx = ((float)xo + 0.5f) / (float)(1 << params.log2_chroma_w) + + (float)(params.src_left >> params.log2_chroma_w); + float cy = ((float)yo + 0.5f) / (float)(1 << params.log2_chroma_h) + + (float)(params.src_top >> params.log2_chroma_h); + + in_uv_T u_raw = tex2D<in_uv_T>(params.src_tex[1], cx, cy); + in_uv_T v_raw = tex2D<in_uv_T>(params.src_tex[2], cx, cy); + + float fy, fu, fv; + if (sizeof(in_y_T) == 1) { + fy = (float)(int)y_raw; + fu = (float)(int)u_raw; + fv = (float)(int)v_raw; + } else { + // planar10: 10-bit native values, divide by 4 + fy = (float)(int)y_raw / 4.0f; + fu = (float)(int)u_raw / 4.0f; + fv = (float)(int)v_raw / 4.0f; + } + + if (params.mpeg_range) fy -= 16.0f; + fu -= 128.0f; + fv -= 128.0f; + + uchar4 *dst = (uchar4*)params.dst[0]; + int dst_pitch = params.dst_pitch / sizeof(uchar4); + dst[yo * dst_pitch + xo] = yuv8_to_rgba_generic(fy, fu, fv, params.color_matrix, channel_order); +} + template<typename T> __device__ static inline T Subsample_Bilinear(cudaTextureObject_t tex, int xo, int yo, @@ -1553,4 +1675,47 @@ GENERIC_FLOAT_KERNEL_Y(bgr0, uchar4, float4) GENERIC_FLOAT_KERNEL_Y(rgba, uchar4, float4) GENERIC_FLOAT_KERNEL_Y(bgra, uchar4, float4) + +// --- YUV to RGB kernel exports --- +// Single-pass kernels (no resize, color conversion only) + +// semiplanar8 (NV12, NV16) -> RGB outputs +__global__ void Subsample_Nearest_semiplanar8_bgra(CUDAScaleKernelParams p) +{ YuvToRgb_semiplanar<uchar, uchar2, 0>(p); } +__global__ void Subsample_Nearest_semiplanar8_rgba(CUDAScaleKernelParams p) +{ YuvToRgb_semiplanar<uchar, uchar2, 1>(p); } +__global__ void Subsample_Nearest_semiplanar8_bgr0(CUDAScaleKernelParams p) +{ YuvToRgb_semiplanar<uchar, uchar2, 2>(p); } +__global__ void Subsample_Nearest_semiplanar8_rgb0(CUDAScaleKernelParams p) +{ YuvToRgb_semiplanar<uchar, uchar2, 3>(p); } + +// semiplanar10 (P010, P210) -> RGB outputs +__global__ void Subsample_Nearest_semiplanar10_bgra(CUDAScaleKernelParams p) +{ YuvToRgb_semiplanar<ushort, ushort2, 0>(p); } +__global__ void Subsample_Nearest_semiplanar10_rgba(CUDAScaleKernelParams p) +{ YuvToRgb_semiplanar<ushort, ushort2, 1>(p); } +__global__ void Subsample_Nearest_semiplanar10_bgr0(CUDAScaleKernelParams p) +{ YuvToRgb_semiplanar<ushort, ushort2, 2>(p); } +__global__ void Subsample_Nearest_semiplanar10_rgb0(CUDAScaleKernelParams p) +{ YuvToRgb_semiplanar<ushort, ushort2, 3>(p); } + +// planar8 (YUV420P, YUV422P, YUV444P) -> RGB outputs +__global__ void Subsample_Nearest_planar8_bgra(CUDAScaleKernelParams p) +{ YuvToRgb_planar<uchar, uchar, 0>(p); } +__global__ void Subsample_Nearest_planar8_rgba(CUDAScaleKernelParams p) +{ YuvToRgb_planar<uchar, uchar, 1>(p); } +__global__ void Subsample_Nearest_planar8_bgr0(CUDAScaleKernelParams p) +{ YuvToRgb_planar<uchar, uchar, 2>(p); } +__global__ void Subsample_Nearest_planar8_rgb0(CUDAScaleKernelParams p) +{ YuvToRgb_planar<uchar, uchar, 3>(p); } + +// planar10 (YUV420P10, YUV422P10, YUV444P10) -> RGB outputs +__global__ void Subsample_Nearest_planar10_bgra(CUDAScaleKernelParams p) +{ YuvToRgb_planar<ushort, ushort, 0>(p); } +__global__ void Subsample_Nearest_planar10_rgba(CUDAScaleKernelParams p) +{ YuvToRgb_planar<ushort, ushort, 1>(p); } +__global__ void Subsample_Nearest_planar10_bgr0(CUDAScaleKernelParams p) +{ YuvToRgb_planar<ushort, ushort, 2>(p); } +__global__ void Subsample_Nearest_planar10_rgb0(CUDAScaleKernelParams p) +{ YuvToRgb_planar<ushort, ushort, 3>(p); } } diff --git a/libavfilter/vf_scale_cuda.h b/libavfilter/vf_scale_cuda.h index 7bf2fa002d..a2a2bcb9a5 100644 --- a/libavfilter/vf_scale_cuda.h +++ b/libavfilter/vf_scale_cuda.h @@ -32,6 +32,10 @@ typedef unsigned char* CUdeviceptr; #define SCALE_CUDA_PARAM_DEFAULT 999999.0f +typedef struct { + float m[3][3]; +} CUDAScaleColorMatrix; + typedef struct { CUtexObject src_tex[4]; CUdeviceptr dst[4]; @@ -51,6 +55,10 @@ typedef struct { CUdeviceptr weights; CUdeviceptr offsets; int filter_size; + + int log2_chroma_w; + int log2_chroma_h; + CUDAScaleColorMatrix color_matrix; } CUDAScaleKernelParams; #endif -- 2.52.0 From 804ef1823409536ebcf597f766bd4a43dd4a8204 Mon Sep 17 00:00:00 2001 From: Zhao Zhili <[email protected]> Date: Mon, 2 Mar 2026 00:40:31 +0800 Subject: [PATCH 2/6] avfilter/vf_scale_cuda: add RGB to YUV color space conversion The output color range is controllable via the out_range option. Color matrix selection follows the same rules as the YUV to RGB path, defaulting to BT.601 when colorspace is unspecified to match swscale behavior. --- libavfilter/vf_scale_cuda.c | 125 ++++++++++++- libavfilter/vf_scale_cuda.cu | 344 +++++++++++++++++++++++++++++++++++ 2 files changed, 464 insertions(+), 5 deletions(-) diff --git a/libavfilter/vf_scale_cuda.c b/libavfilter/vf_scale_cuda.c index bab7b35cc5..046d396122 100644 --- a/libavfilter/vf_scale_cuda.c +++ b/libavfilter/vf_scale_cuda.c @@ -42,6 +42,7 @@ #include "avfilter.h" #include "colorspace.h" #include "filters.h" +#include "formats.h" #include "scale_eval.h" #include "video.h" @@ -196,6 +197,8 @@ typedef struct CUDAScaleContext { float param; bool yuv2rgb; + bool rgb2yuv; + int out_range; ///< output color range for RGB->YUV (AVColorRange) enum AVColorSpace colorspace; ///< resolved colorspace for conversion enum AVColorRange color_range; ///< resolved color range for conversion CUDAScaleColorMatrix color_matrix; @@ -497,6 +500,38 @@ static int compute_yuv2rgb_matrix(void *log_ctx, CUDAScaleColorMatrix *mat, return 0; } +static int compute_rgb2yuv_matrix(void *log_ctx, CUDAScaleColorMatrix *mat, + enum AVColorSpace colorspace, + bool limited_range) +{ + const AVLumaCoefficients *coeffs = av_csp_luma_coeffs_from_avcsp(colorspace); + double rgb2yuv[3][3]; + float y_scale = 1.0f; + float uv_scale = 1.0f; + + if (!coeffs) { + av_log(log_ctx, AV_LOG_ERROR, + "Unsupported colorspace %d for RGB->YUV conversion\n", colorspace); + return AVERROR(EINVAL); + } + + ff_fill_rgb2yuv_table(coeffs, rgb2yuv); + + if (limited_range) { + y_scale = (235.0f - 16.0f) / 255.0f; + uv_scale = (240.0f - 16.0f) / 255.0f; + } + + for (int i = 0; i < 3; i++) { + float scale = (i == 0) ? y_scale : uv_scale; + mat->m[i][0] = (float)rgb2yuv[i][0] * scale; + mat->m[i][1] = (float)rgb2yuv[i][1] * scale; + mat->m[i][2] = (float)rgb2yuv[i][2] * scale; + } + + return 0; +} + static av_cold void set_format_info(AVFilterContext *ctx, enum AVPixelFormat in_format, enum AVPixelFormat out_format) { CUDAScaleContext *s = ctx->priv; @@ -603,6 +638,50 @@ static av_cold int init_processing_chain(AVFilterContext *ctx, int in_width, int limited_range ? "limited" : "full"); } + s->rgb2yuv = (s->in_desc->flags & AV_PIX_FMT_FLAG_RGB) && + !(s->out_desc->flags & AV_PIX_FMT_FLAG_RGB); + if (s->rgb2yuv) { + AVFilterLink *outlink = ctx->outputs[0]; + + if (in_width != out_width || in_height != out_height) { + av_log(ctx, AV_LOG_ERROR, + "Resizing is not supported during RGB->YUV conversion, " + "use a separate scale_cuda instance for resizing\n"); + return AVERROR(ENOSYS); + } + enum AVColorSpace cs = ctx->inputs[0]->colorspace; + bool limited_range; + + /* When colorspace is unspecified or RGB, default to BT.601 + * (SMPTE 170M) to match swscale behavior. + */ + if (cs == AVCOL_SPC_UNSPECIFIED || cs == AVCOL_SPC_RGB) + cs = AVCOL_SPC_SMPTE170M; + + limited_range = (outlink->color_range != AVCOL_RANGE_JPEG); + + ret = compute_rgb2yuv_matrix(ctx, &s->color_matrix, cs, limited_range); + if (ret < 0) + return ret; + + s->colorspace = cs; + s->color_range = outlink->color_range; + + /* RGB->YUV only supports Nearest for now. Forcing it here also makes + * the generic filter path get disabled below. */ + if (s->interp_algo != INTERP_ALGO_DEFAULT && + s->interp_algo != INTERP_ALGO_NEAREST) + av_log(ctx, AV_LOG_WARNING, + "RGB->YUV conversion only supports nearest interpolation, " + "ignoring interp_algo\n"); + s->interp_algo = INTERP_ALGO_NEAREST; + + av_log(ctx, AV_LOG_VERBOSE, + "RGB->YUV conversion enabled (%s, %s range)\n", + av_color_space_name(cs), + limited_range ? "limited" : "full"); + } + if (s->passthrough && in_width == out_width && in_height == out_height && in_format == out_format) { s->frames_ctx = av_buffer_ref(inl->hw_frames_ctx); if (!s->frames_ctx) @@ -1035,6 +1114,27 @@ static av_cold int cudascale_setup_filters(AVFilterContext *ctx, return ret; } +static int cudascale_query_formats(const AVFilterContext *ctx, + AVFilterFormatsConfig **cfg_in, + AVFilterFormatsConfig **cfg_out) +{ + const CUDAScaleContext *s = ctx->priv; + static const enum AVPixelFormat pix_fmts[] = { AV_PIX_FMT_CUDA, AV_PIX_FMT_NONE }; + AVFilterFormats *formats; + int ret; + + if ((ret = ff_set_common_formats_from_list2(ctx, cfg_in, cfg_out, pix_fmts)) < 0) + return ret; + + formats = s->out_range != AVCOL_RANGE_UNSPECIFIED + ? ff_make_formats_list_singleton(s->out_range) + : ff_all_color_ranges(); + if ((ret = ff_formats_ref(formats, &cfg_out[0]->color_ranges)) < 0) + return ret; + + return 0; +} + static av_cold int cudascale_config_props(AVFilterLink *outlink) { AVFilterContext *ctx = outlink->src; @@ -1248,6 +1348,9 @@ static int call_resize_kernel(AVFilterContext *ctx, CUfunction func, if (s->yuv2rgb) { params.log2_chroma_w = s->in_desc->log2_chroma_w; params.log2_chroma_h = s->in_desc->log2_chroma_h; + } else if (s->rgb2yuv) { + params.log2_chroma_w = s->out_desc->log2_chroma_w; + params.log2_chroma_h = s->out_desc->log2_chroma_h; } void *args[] = { ¶ms }; @@ -1265,11 +1368,12 @@ static int scalecuda_resize(AVFilterContext *ctx, const CUDAScaleFilter *filter = NULL; const CUDAScaleFilter *filter_uv = NULL; CUfunction func, func_uv; - /* The color matrix was computed at config time, so for yuv2rgb use the - * resolved color range to keep mpeg_range consistent with it. For plain + /* The color matrix was computed at config time, so for yuv2rgb/rgb2yuv use + * the resolved color range to keep mpeg_range consistent with it. For plain * scaling, follow the frame the texture was mapped from. */ - int mpeg_range = s->yuv2rgb ? (s->color_range != AVCOL_RANGE_JPEG) - : (in->color_range != AVCOL_RANGE_JPEG); + int mpeg_range = (s->yuv2rgb || s->rgb2yuv) ? + (s->color_range != AVCOL_RANGE_JPEG) : + (in->color_range != AVCOL_RANGE_JPEG); int ret; int out_planes = s->out_planes; @@ -1393,6 +1497,9 @@ static int cudascale_scale(AVFilterContext *ctx, AVFrame *out, AVFrame *in, if (s->yuv2rgb) { out->colorspace = AVCOL_SPC_RGB; out->color_range = AVCOL_RANGE_JPEG; + } else if (s->rgb2yuv) { + out->colorspace = s->colorspace; + out->color_range = s->color_range; } out->crop_top = 0; @@ -1497,6 +1604,14 @@ static const AVOption options[] = { { "increase", NULL, 0, AV_OPT_TYPE_CONST, {.i64 = SCALE_FORCE_OAR_INCREASE }, 0, 0, FLAGS, .unit = "force_oar" }, { "force_divisible_by", "enforce that the output resolution is divisible by a defined integer when force_original_aspect_ratio is used", OFFSET(force_divisible_by), AV_OPT_TYPE_INT, { .i64 = 1 }, 1, 256, FLAGS }, { "reset_sar", "reset SAR to 1 and scale to square pixels if scaling proportionally", OFFSET(reset_sar), AV_OPT_TYPE_BOOL, { .i64 = 0}, 0, 1, FLAGS }, + { "out_range", "Output color range", OFFSET(out_range), AV_OPT_TYPE_INT, { .i64 = AVCOL_RANGE_UNSPECIFIED }, AVCOL_RANGE_UNSPECIFIED, AVCOL_RANGE_JPEG, FLAGS, .unit = "range" }, + { "auto", NULL, 0, AV_OPT_TYPE_CONST, { .i64 = AVCOL_RANGE_UNSPECIFIED }, 0, 0, FLAGS, .unit = "range" }, + { "full", NULL, 0, AV_OPT_TYPE_CONST, { .i64 = AVCOL_RANGE_JPEG }, 0, 0, FLAGS, .unit = "range" }, + { "limited", NULL, 0, AV_OPT_TYPE_CONST, { .i64 = AVCOL_RANGE_MPEG }, 0, 0, FLAGS, .unit = "range" }, + { "jpeg", NULL, 0, AV_OPT_TYPE_CONST, { .i64 = AVCOL_RANGE_JPEG }, 0, 0, FLAGS, .unit = "range" }, + { "mpeg", NULL, 0, AV_OPT_TYPE_CONST, { .i64 = AVCOL_RANGE_MPEG }, 0, 0, FLAGS, .unit = "range" }, + { "tv", NULL, 0, AV_OPT_TYPE_CONST, { .i64 = AVCOL_RANGE_MPEG }, 0, 0, FLAGS, .unit = "range" }, + { "pc", NULL, 0, AV_OPT_TYPE_CONST, { .i64 = AVCOL_RANGE_JPEG }, 0, 0, FLAGS, .unit = "range" }, { NULL }, }; @@ -1538,7 +1653,7 @@ const FFFilter ff_vf_scale_cuda = { FILTER_INPUTS(cudascale_inputs), FILTER_OUTPUTS(cudascale_outputs), - FILTER_SINGLE_PIXFMT(AV_PIX_FMT_CUDA), + FILTER_QUERY_FUNC2(cudascale_query_formats), .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE, }; diff --git a/libavfilter/vf_scale_cuda.cu b/libavfilter/vf_scale_cuda.cu index 2a13ffd906..793decb8ee 100644 --- a/libavfilter/vf_scale_cuda.cu +++ b/libavfilter/vf_scale_cuda.cu @@ -1235,6 +1235,250 @@ __device__ void YuvToRgb_planar(CUDAScaleKernelParams params) dst[yo * dst_pitch + xo] = yuv8_to_rgba_generic(fy, fu, fv, params.color_matrix, channel_order); } +// RGB to YUV +// +// Y = m[0][0]*R + m[0][1]*G + m[0][2]*B (+ 16 for limited range) +// U = m[1][0]*R + m[1][1]*G + m[1][2]*B (+ 128) +// V = m[2][0]*R + m[2][1]*G + m[2][2]*B (+ 128) + +// Extract R,G,B from packed uchar4 pixel based on channel ordering. +// channel_order: 0 = BGRA (x=B,y=G,z=R), 1 = RGBA (x=R,y=G,z=B), +// 2 = BGR0 (x=B,y=G,z=R), 3 = RGB0 (x=R,y=G,z=B) +__device__ static inline void extract_rgb(uchar4 pixel, int channel_order, + float &r, float &g, float &b) +{ + switch (channel_order) { + case 0: // BGRA + case 2: // BGR0 + b = (float)pixel.x; g = (float)pixel.y; r = (float)pixel.z; + break; + case 1: // RGBA + case 3: // RGB0 + default: + r = (float)pixel.x; g = (float)pixel.y; b = (float)pixel.z; + break; + } +} + +__device__ static inline float rgb_to_y(float r, float g, float b, + const CUDAScaleColorMatrix &mat, int mpeg_range) +{ + float y = mat.m[0][0] * r + mat.m[0][1] * g + mat.m[0][2] * b; + if (mpeg_range) y += 16.0f; + return clamp_f(y, 0.0f, 255.0f); +} + +__device__ static inline float rgb_to_u(float r, float g, float b, + const CUDAScaleColorMatrix &mat) +{ + float u = mat.m[1][0] * r + mat.m[1][1] * g + mat.m[1][2] * b + 128.0f; + return clamp_f(u, 0.0f, 255.0f); +} + +__device__ static inline float rgb_to_v(float r, float g, float b, + const CUDAScaleColorMatrix &mat) +{ + float v = mat.m[2][0] * r + mat.m[2][1] * g + mat.m[2][2] * b + 128.0f; + return clamp_f(v, 0.0f, 255.0f); +} + +// RGB->YUV Y kernel (8-bit output) +template<typename out_y_T, int channel_order> +__device__ void RgbToYuv_Y(CUDAScaleKernelParams params) +{ + int xo = blockIdx.x * blockDim.x + threadIdx.x; + int yo = blockIdx.y * blockDim.y + threadIdx.y; + if (yo >= params.dst_height || xo >= params.dst_width) return; + + uchar4 pixel = tex2D<uchar4>( + params.src_tex[0], + xo + params.src_left + 0.5f, + yo + params.src_top + 0.5f); + + float r, g, b; + extract_rgb(pixel, channel_order, r, g, b); + float y = rgb_to_y(r, g, b, params.color_matrix, params.mpeg_range); + + out_y_T *dst_y = (out_y_T*)params.dst[0]; + int pitch = params.dst_pitch / sizeof(out_y_T); + dst_y[yo * pitch + xo] = (out_y_T)y; +} + +// RGB->YUV Y kernel for semiplanar10/16 output +template<int channel_order> +__device__ void RgbToYuv_Y_semiplanar16(CUDAScaleKernelParams params) +{ + int xo = blockIdx.x * blockDim.x + threadIdx.x; + int yo = blockIdx.y * blockDim.y + threadIdx.y; + if (yo >= params.dst_height || xo >= params.dst_width) return; + + uchar4 pixel = tex2D<uchar4>( + params.src_tex[0], + xo + params.src_left + 0.5f, + yo + params.src_top + 0.5f); + + float r, g, b; + extract_rgb(pixel, channel_order, r, g, b); + float y = rgb_to_y(r, g, b, params.color_matrix, params.mpeg_range); + + ushort *dst_y = (ushort*)params.dst[0]; + int pitch = params.dst_pitch / sizeof(ushort); + dst_y[yo * pitch + xo] = (ushort)clamp_f(y * 256.0f, 0.0f, 65535.0f); +} + +// RGB->YUV Y kernel for planar10 output: native 10-bit Y value +template<int channel_order> +__device__ void RgbToYuv_Y_planar10(CUDAScaleKernelParams params) +{ + int xo = blockIdx.x * blockDim.x + threadIdx.x; + int yo = blockIdx.y * blockDim.y + threadIdx.y; + if (yo >= params.dst_height || xo >= params.dst_width) return; + + uchar4 pixel = tex2D<uchar4>( + params.src_tex[0], + xo + params.src_left + 0.5f, + yo + params.src_top + 0.5f); + + float r, g, b; + extract_rgb(pixel, channel_order, r, g, b); + float y = rgb_to_y(r, g, b, params.color_matrix, params.mpeg_range); + + ushort *dst_y = (ushort*)params.dst[0]; + int pitch = params.dst_pitch / sizeof(ushort); + dst_y[yo * pitch + xo] = (ushort)clamp_f(y * 4.0f, 0.0f, 1023.0f); +} + +// RGB->YUV UV kernel for semiplanar output (NV12, NV16) +template<typename out_uv_T, int channel_order> +__device__ void RgbToYuv_semiplanar_UV(CUDAScaleKernelParams params) +{ + int xo = blockIdx.x * blockDim.x + threadIdx.x; + int yo = blockIdx.y * blockDim.y + threadIdx.y; + if (yo >= params.dst_height || xo >= params.dst_width) return; + + int chroma_w = 1 << params.log2_chroma_w; + int chroma_h = 1 << params.log2_chroma_h; + + float u_acc = 0.0f, v_acc = 0.0f; + int count = 0; + for (int dy = 0; dy < chroma_h; dy++) { + for (int dx = 0; dx < chroma_w; dx++) { + int lx = xo * chroma_w + dx; + int ly = yo * chroma_h + dy; + if (lx >= params.src_width || ly >= params.src_height) continue; + + uchar4 pixel = tex2D<uchar4>( + params.src_tex[0], + lx + params.src_left + 0.5f, + ly + params.src_top + 0.5f); + + float r, g, b; + extract_rgb(pixel, channel_order, r, g, b); + u_acc += rgb_to_u(r, g, b, params.color_matrix); + v_acc += rgb_to_v(r, g, b, params.color_matrix); + count++; + } + } + + float u_avg = u_acc / (float)count; + float v_avg = v_acc / (float)count; + + out_uv_T *dst_uv = (out_uv_T*)params.dst[1]; + int pitch = params.dst_pitch / sizeof(out_uv_T); + dst_uv[yo * pitch + xo] = make_uchar2((uchar)u_avg, (uchar)v_avg); +} + +// RGB->YUV UV kernel for semiplanar10/16 output +template<int channel_order> +__device__ void RgbToYuv_semiplanar16_UV(CUDAScaleKernelParams params) +{ + int xo = blockIdx.x * blockDim.x + threadIdx.x; + int yo = blockIdx.y * blockDim.y + threadIdx.y; + if (yo >= params.dst_height || xo >= params.dst_width) return; + + int chroma_w = 1 << params.log2_chroma_w; + int chroma_h = 1 << params.log2_chroma_h; + + float u_acc = 0.0f, v_acc = 0.0f; + int count = 0; + for (int dy = 0; dy < chroma_h; dy++) { + for (int dx = 0; dx < chroma_w; dx++) { + int lx = xo * chroma_w + dx; + int ly = yo * chroma_h + dy; + if (lx >= params.src_width || ly >= params.src_height) continue; + + uchar4 pixel = tex2D<uchar4>( + params.src_tex[0], + lx + params.src_left + 0.5f, + ly + params.src_top + 0.5f); + + float r, g, b; + extract_rgb(pixel, channel_order, r, g, b); + u_acc += rgb_to_u(r, g, b, params.color_matrix); + v_acc += rgb_to_v(r, g, b, params.color_matrix); + count++; + } + } + + float u_avg = u_acc / (float)count; + float v_avg = v_acc / (float)count; + + ushort2 *dst_uv = (ushort2*)params.dst[1]; + int pitch = params.dst_pitch / sizeof(ushort2); + ushort u_val = (ushort)clamp_f(u_avg * 256.0f, 0.0f, 65535.0f); + ushort v_val = (ushort)clamp_f(v_avg * 256.0f, 0.0f, 65535.0f); + dst_uv[yo * pitch + xo] = make_ushort2(u_val, v_val); +} + +// RGB->YUV UV kernel for planar output (YUV420P, YUV422P, YUV444P, YUV420P10) +template<typename out_uv_T, int channel_order> +__device__ void RgbToYuv_planar_UV(CUDAScaleKernelParams params) +{ + int xo = blockIdx.x * blockDim.x + threadIdx.x; + int yo = blockIdx.y * blockDim.y + threadIdx.y; + if (yo >= params.dst_height || xo >= params.dst_width) return; + + int chroma_w = 1 << params.log2_chroma_w; + int chroma_h = 1 << params.log2_chroma_h; + + float u_acc = 0.0f, v_acc = 0.0f; + int count = 0; + for (int dy = 0; dy < chroma_h; dy++) { + for (int dx = 0; dx < chroma_w; dx++) { + int lx = xo * chroma_w + dx; + int ly = yo * chroma_h + dy; + if (lx >= params.src_width || ly >= params.src_height) continue; + + uchar4 pixel = tex2D<uchar4>( + params.src_tex[0], + lx + params.src_left + 0.5f, + ly + params.src_top + 0.5f); + + float r, g, b; + extract_rgb(pixel, channel_order, r, g, b); + u_acc += rgb_to_u(r, g, b, params.color_matrix); + v_acc += rgb_to_v(r, g, b, params.color_matrix); + count++; + } + } + + float u_avg = u_acc / (float)count; + float v_avg = v_acc / (float)count; + + out_uv_T *dst_u = (out_uv_T*)params.dst[1]; + out_uv_T *dst_v = (out_uv_T*)params.dst[2]; + int pitch = params.dst_pitch / sizeof(out_uv_T); + + if (sizeof(out_uv_T) == 1) { + dst_u[yo * pitch + xo] = (out_uv_T)u_avg; + dst_v[yo * pitch + xo] = (out_uv_T)v_avg; + } else { + // 10-bit planar: native 10-bit value = 8-bit * 4 + dst_u[yo * pitch + xo] = (out_uv_T)(clamp_f(u_avg * 4.0f, 0.0f, 1023.0f)); + dst_v[yo * pitch + xo] = (out_uv_T)(clamp_f(v_avg * 4.0f, 0.0f, 1023.0f)); + } +} + template<typename T> __device__ static inline T Subsample_Bilinear(cudaTextureObject_t tex, int xo, int yo, @@ -1718,4 +1962,104 @@ __global__ void Subsample_Nearest_planar10_bgr0(CUDAScaleKernelParams p) { YuvToRgb_planar<ushort, ushort, 2>(p); } __global__ void Subsample_Nearest_planar10_rgb0(CUDAScaleKernelParams p) { YuvToRgb_planar<ushort, ushort, 3>(p); } + +// --- RGB to YUV kernel exports --- +// Two-pass: Y kernel writes Y plane, UV kernel writes chroma planes. +// No resize, color conversion only. + +// bgra -> semiplanar8 (NV12, NV16) +__global__ void Subsample_Nearest_bgra_semiplanar8(CUDAScaleKernelParams p) +{ RgbToYuv_Y<uchar, 0>(p); } +__global__ void Subsample_Nearest_bgra_semiplanar8_uv(CUDAScaleKernelParams p) +{ RgbToYuv_semiplanar_UV<uchar2, 0>(p); } + +// rgba -> semiplanar8 +__global__ void Subsample_Nearest_rgba_semiplanar8(CUDAScaleKernelParams p) +{ RgbToYuv_Y<uchar, 1>(p); } +__global__ void Subsample_Nearest_rgba_semiplanar8_uv(CUDAScaleKernelParams p) +{ RgbToYuv_semiplanar_UV<uchar2, 1>(p); } + +// bgr0 -> semiplanar8 +__global__ void Subsample_Nearest_bgr0_semiplanar8(CUDAScaleKernelParams p) +{ RgbToYuv_Y<uchar, 2>(p); } +__global__ void Subsample_Nearest_bgr0_semiplanar8_uv(CUDAScaleKernelParams p) +{ RgbToYuv_semiplanar_UV<uchar2, 2>(p); } + +// rgb0 -> semiplanar8 +__global__ void Subsample_Nearest_rgb0_semiplanar8(CUDAScaleKernelParams p) +{ RgbToYuv_Y<uchar, 3>(p); } +__global__ void Subsample_Nearest_rgb0_semiplanar8_uv(CUDAScaleKernelParams p) +{ RgbToYuv_semiplanar_UV<uchar2, 3>(p); } + +// bgra -> semiplanar10 (P010, P210) +__global__ void Subsample_Nearest_bgra_semiplanar10(CUDAScaleKernelParams p) +{ RgbToYuv_Y_semiplanar16<0>(p); } +__global__ void Subsample_Nearest_bgra_semiplanar10_uv(CUDAScaleKernelParams p) +{ RgbToYuv_semiplanar16_UV<0>(p); } + +// rgba -> semiplanar10 +__global__ void Subsample_Nearest_rgba_semiplanar10(CUDAScaleKernelParams p) +{ RgbToYuv_Y_semiplanar16<1>(p); } +__global__ void Subsample_Nearest_rgba_semiplanar10_uv(CUDAScaleKernelParams p) +{ RgbToYuv_semiplanar16_UV<1>(p); } + +// bgr0 -> semiplanar10 +__global__ void Subsample_Nearest_bgr0_semiplanar10(CUDAScaleKernelParams p) +{ RgbToYuv_Y_semiplanar16<2>(p); } +__global__ void Subsample_Nearest_bgr0_semiplanar10_uv(CUDAScaleKernelParams p) +{ RgbToYuv_semiplanar16_UV<2>(p); } + +// rgb0 -> semiplanar10 +__global__ void Subsample_Nearest_rgb0_semiplanar10(CUDAScaleKernelParams p) +{ RgbToYuv_Y_semiplanar16<3>(p); } +__global__ void Subsample_Nearest_rgb0_semiplanar10_uv(CUDAScaleKernelParams p) +{ RgbToYuv_semiplanar16_UV<3>(p); } + +// bgra -> planar8 (YUV420P, YUV422P, YUV444P) +__global__ void Subsample_Nearest_bgra_planar8(CUDAScaleKernelParams p) +{ RgbToYuv_Y<uchar, 0>(p); } +__global__ void Subsample_Nearest_bgra_planar8_uv(CUDAScaleKernelParams p) +{ RgbToYuv_planar_UV<uchar, 0>(p); } + +// rgba -> planar8 +__global__ void Subsample_Nearest_rgba_planar8(CUDAScaleKernelParams p) +{ RgbToYuv_Y<uchar, 1>(p); } +__global__ void Subsample_Nearest_rgba_planar8_uv(CUDAScaleKernelParams p) +{ RgbToYuv_planar_UV<uchar, 1>(p); } + +// bgr0 -> planar8 +__global__ void Subsample_Nearest_bgr0_planar8(CUDAScaleKernelParams p) +{ RgbToYuv_Y<uchar, 2>(p); } +__global__ void Subsample_Nearest_bgr0_planar8_uv(CUDAScaleKernelParams p) +{ RgbToYuv_planar_UV<uchar, 2>(p); } + +// rgb0 -> planar8 +__global__ void Subsample_Nearest_rgb0_planar8(CUDAScaleKernelParams p) +{ RgbToYuv_Y<uchar, 3>(p); } +__global__ void Subsample_Nearest_rgb0_planar8_uv(CUDAScaleKernelParams p) +{ RgbToYuv_planar_UV<uchar, 3>(p); } + +// bgra -> planar10 (YUV420P10, YUV422P10, YUV444P10) +__global__ void Subsample_Nearest_bgra_planar10(CUDAScaleKernelParams p) +{ RgbToYuv_Y_planar10<0>(p); } +__global__ void Subsample_Nearest_bgra_planar10_uv(CUDAScaleKernelParams p) +{ RgbToYuv_planar_UV<ushort, 0>(p); } + +// rgba -> planar10 +__global__ void Subsample_Nearest_rgba_planar10(CUDAScaleKernelParams p) +{ RgbToYuv_Y_planar10<1>(p); } +__global__ void Subsample_Nearest_rgba_planar10_uv(CUDAScaleKernelParams p) +{ RgbToYuv_planar_UV<ushort, 1>(p); } + +// bgr0 -> planar10 +__global__ void Subsample_Nearest_bgr0_planar10(CUDAScaleKernelParams p) +{ RgbToYuv_Y_planar10<2>(p); } +__global__ void Subsample_Nearest_bgr0_planar10_uv(CUDAScaleKernelParams p) +{ RgbToYuv_planar_UV<ushort, 2>(p); } + +// rgb0 -> planar10 +__global__ void Subsample_Nearest_rgb0_planar10(CUDAScaleKernelParams p) +{ RgbToYuv_Y_planar10<3>(p); } +__global__ void Subsample_Nearest_rgb0_planar10_uv(CUDAScaleKernelParams p) +{ RgbToYuv_planar_UV<ushort, 3>(p); } } -- 2.52.0 From d017367ebaa9d8afa38cd5756a5cf6e8bcfb6295 Mon Sep 17 00:00:00 2001 From: Zhao Zhili <[email protected]> Date: Mon, 2 Mar 2026 00:41:04 +0800 Subject: [PATCH 3/6] avfilter/vf_scale_cuda: add FATE tests for YUV/RGB conversion Add CUDA hardware-accelerated FATE tests for scale_cuda YUV<->RGB color conversion. Tests require a CUDA GPU and are gated behind HWACCEL=cuda: make fate-filter-scale_cuda HWACCEL=cuda --- tests/Makefile | 10 ++++ tests/fate/filter-video.mak | 54 +++++++++++++++++++ tests/ref/fate/filter-scale_cuda-bgra-nv12 | 10 ++++ tests/ref/fate/filter-scale_cuda-bgra-p010 | 10 ++++ tests/ref/fate/filter-scale_cuda-bgra-yuv420p | 10 ++++ tests/ref/fate/filter-scale_cuda-nv12-bgra | 10 ++++ .../fate/filter-scale_cuda-nv12-bgra-bt709 | 10 ++++ .../ref/fate/filter-scale_cuda-nv12-bgra-full | 10 ++++ .../fate/filter-scale_cuda-nv12-nv12-scale | 10 ++++ tests/ref/fate/filter-scale_cuda-nv12-rgba | 10 ++++ .../ref/fate/filter-scale_cuda-odd-bgra-nv12 | 10 ++++ .../fate/filter-scale_cuda-odd-bgra-yuv420p | 10 ++++ .../ref/fate/filter-scale_cuda-odd-nv12-bgra | 10 ++++ .../fate/filter-scale_cuda-odd-yuv420p-bgra | 10 ++++ tests/ref/fate/filter-scale_cuda-p010-bgra | 10 ++++ tests/ref/fate/filter-scale_cuda-rgba-nv12 | 10 ++++ .../ref/fate/filter-scale_cuda-roundtrip-nv12 | 10 ++++ .../fate/filter-scale_cuda-roundtrip-yuv444p | 10 ++++ tests/ref/fate/filter-scale_cuda-yuv420p-bgra | 10 ++++ tests/ref/fate/filter-scale_cuda-yuv444p-bgra | 10 ++++ 20 files changed, 244 insertions(+) create mode 100644 tests/ref/fate/filter-scale_cuda-bgra-nv12 create mode 100644 tests/ref/fate/filter-scale_cuda-bgra-p010 create mode 100644 tests/ref/fate/filter-scale_cuda-bgra-yuv420p create mode 100644 tests/ref/fate/filter-scale_cuda-nv12-bgra create mode 100644 tests/ref/fate/filter-scale_cuda-nv12-bgra-bt709 create mode 100644 tests/ref/fate/filter-scale_cuda-nv12-bgra-full create mode 100644 tests/ref/fate/filter-scale_cuda-nv12-nv12-scale create mode 100644 tests/ref/fate/filter-scale_cuda-nv12-rgba create mode 100644 tests/ref/fate/filter-scale_cuda-odd-bgra-nv12 create mode 100644 tests/ref/fate/filter-scale_cuda-odd-bgra-yuv420p create mode 100644 tests/ref/fate/filter-scale_cuda-odd-nv12-bgra create mode 100644 tests/ref/fate/filter-scale_cuda-odd-yuv420p-bgra create mode 100644 tests/ref/fate/filter-scale_cuda-p010-bgra create mode 100644 tests/ref/fate/filter-scale_cuda-rgba-nv12 create mode 100644 tests/ref/fate/filter-scale_cuda-roundtrip-nv12 create mode 100644 tests/ref/fate/filter-scale_cuda-roundtrip-yuv444p create mode 100644 tests/ref/fate/filter-scale_cuda-yuv420p-bgra create mode 100644 tests/ref/fate/filter-scale_cuda-yuv444p-bgra diff --git a/tests/Makefile b/tests/Makefile index bd83c8df91..33f318fd01 100644 --- a/tests/Makefile +++ b/tests/Makefile @@ -139,6 +139,16 @@ CONFIG_LARGE_TESTS:= !CONFIG_LARGE_TESTS:=yes endif +# Enable CUDA hardware-accelerated tests when HWACCEL=cuda is set on +# the make command line (e.g. make fate-filter-scale_cuda HWACCEL=cuda). +ifeq ($(HWACCEL), cuda) +CONFIG_HWACCEL_CUDA_TESTS:=yes +!CONFIG_HWACCEL_CUDA_TESTS:= +else +CONFIG_HWACCEL_CUDA_TESTS:= +!CONFIG_HWACCEL_CUDA_TESTS:=yes +endif + include $(SRC_PATH)/$(APITESTSDIR)/Makefile include $(SRC_PATH)/tests/fate/acodec.mak diff --git a/tests/fate/filter-video.mak b/tests/fate/filter-video.mak index ebbc0a7d1c..84afbcc2a2 100644 --- a/tests/fate/filter-video.mak +++ b/tests/fate/filter-video.mak @@ -1018,4 +1018,58 @@ FATE-$(CONFIG_AVFILTER) += $(FATE_LIBAVFILTER-yes) fate-vfilter: $(FATE_FILTER-yes) $(FATE_FILTER_SAMPLES-yes) $(FATE_FILTER_VSYNTH-yes) $(FATE_LIBAVFILTER-yes) +# scale_cuda filter tests (YUV<->RGB color space conversion) +# These require a CUDA GPU and are only run when HWACCEL=cuda is set: +# make fate-filter-scale_cuda HWACCEL=cuda +# +SCALE_CUDA_DEPS = SCALE_CUDA_FILTER HWUPLOAD_CUDA_FILTER HWDOWNLOAD_FILTER \ + FORMAT_FILTER TESTSRC2_FILTER HWACCEL_CUDA_TESTS + +SCALE_CUDA_ALLYES = $(call ALLYES, $(SCALE_CUDA_DEPS)) + +sc_yuv_rgb = format=$(1),hwupload_cuda,scale_cuda=format=$(2),hwdownload,format=$(2) +sc_rgb_yuv = format=$(1),hwupload_cuda,scale_cuda=format=$(2):out_range=limited,hwdownload,format=$(2) +sc_roundtrip = format=$(1),hwupload_cuda,scale_cuda=format=$(2),scale_cuda=format=$(3):out_range=limited,hwdownload,format=$(3) + +FATE_SCALE_CUDA_YUV_RGB := $(addprefix fate-filter-scale_cuda-, nv12-bgra yuv420p-bgra p010-bgra yuv444p-bgra nv12-rgba) +FATE_SCALE_CUDA_RGB_YUV := $(addprefix fate-filter-scale_cuda-, bgra-nv12 bgra-yuv420p bgra-p010 rgba-nv12) +FATE_SCALE_CUDA_SPECIAL := $(addprefix fate-filter-scale_cuda-, nv12-bgra-bt709 nv12-bgra-full roundtrip-nv12 roundtrip-yuv444p nv12-nv12-scale) +FATE_SCALE_CUDA_ODD := $(addprefix fate-filter-scale_cuda-, odd-bgra-nv12 odd-bgra-yuv420p odd-nv12-bgra odd-yuv420p-bgra) + +FATE_SCALE_CUDA-$(SCALE_CUDA_ALLYES) := $(FATE_SCALE_CUDA_YUV_RGB) $(FATE_SCALE_CUDA_RGB_YUV) \ + $(FATE_SCALE_CUDA_SPECIAL) $(FATE_SCALE_CUDA_ODD) + +$(FATE_SCALE_CUDA-yes): CMD = framecrc -lavfi "testsrc2=s=$(SIZE),$(LAVFI)" -frames:v 5 + +$(FATE_SCALE_CUDA_YUV_RGB) $(FATE_SCALE_CUDA_RGB_YUV) $(FATE_SCALE_CUDA_SPECIAL): SIZE = 160x120 +$(FATE_SCALE_CUDA_ODD): SIZE = 161x121 + +fate-filter-scale_cuda-nv12-bgra: LAVFI = $(call sc_yuv_rgb,nv12,bgra) +fate-filter-scale_cuda-yuv420p-bgra: LAVFI = $(call sc_yuv_rgb,yuv420p,bgra) +fate-filter-scale_cuda-p010-bgra: LAVFI = $(call sc_yuv_rgb,p010,bgra) +fate-filter-scale_cuda-yuv444p-bgra: LAVFI = $(call sc_yuv_rgb,yuv444p,bgra) +fate-filter-scale_cuda-nv12-rgba: LAVFI = $(call sc_yuv_rgb,nv12,rgba) + +fate-filter-scale_cuda-bgra-nv12: LAVFI = $(call sc_rgb_yuv,bgra,nv12) +fate-filter-scale_cuda-bgra-yuv420p: LAVFI = $(call sc_rgb_yuv,bgra,yuv420p) +fate-filter-scale_cuda-bgra-p010: LAVFI = $(call sc_rgb_yuv,bgra,p010) +fate-filter-scale_cuda-rgba-nv12: LAVFI = $(call sc_rgb_yuv,rgba,nv12) + +fate-filter-scale_cuda-nv12-bgra-bt709: LAVFI = format=nv12,setparams=colorspace=bt709:range=tv,hwupload_cuda,scale_cuda=format=bgra,hwdownload,format=bgra +fate-filter-scale_cuda-nv12-bgra-full: LAVFI = format=bgra,hwupload_cuda,scale_cuda=format=nv12:out_range=full,scale_cuda=format=bgra,hwdownload,format=bgra + +fate-filter-scale_cuda-roundtrip-nv12: LAVFI = $(call sc_roundtrip,nv12,bgra,nv12) +fate-filter-scale_cuda-roundtrip-yuv444p: LAVFI = $(call sc_roundtrip,bgra,yuv444p,bgra) +fate-filter-scale_cuda-nv12-nv12-scale: LAVFI = format=nv12,hwupload_cuda,scale_cuda=w=320:h=240:format=nv12,hwdownload,format=nv12 + +fate-filter-scale_cuda-odd-bgra-nv12: LAVFI = $(call sc_rgb_yuv,bgra,nv12) +fate-filter-scale_cuda-odd-bgra-yuv420p: LAVFI = $(call sc_rgb_yuv,bgra,yuv420p) +fate-filter-scale_cuda-odd-nv12-bgra: LAVFI = $(call sc_yuv_rgb,nv12,bgra) +fate-filter-scale_cuda-odd-yuv420p-bgra: LAVFI = $(call sc_yuv_rgb,yuv420p,bgra) + +FATE_HW-yes += $(FATE_SCALE_CUDA-yes) + +fate-filter-scale_cuda: $(FATE_SCALE_CUDA-yes) + + fate-filter: fate-afilter fate-vfilter $(FATE_METADATA_FILTER-yes) $(FATE_FILTER_FFPROBE-yes) diff --git a/tests/ref/fate/filter-scale_cuda-bgra-nv12 b/tests/ref/fate/filter-scale_cuda-bgra-nv12 new file mode 100644 index 0000000000..8a9e4d8c80 --- /dev/null +++ b/tests/ref/fate/filter-scale_cuda-bgra-nv12 @@ -0,0 +1,10 @@ +#tb 0: 1/25 +#media_type 0: video +#codec_id 0: rawvideo +#dimensions 0: 160x120 +#sar 0: 1/1 +0, 0, 0, 1, 28800, 0xea53c72b +0, 1, 1, 1, 28800, 0x77cbc057 +0, 2, 2, 1, 28800, 0x429ec8e8 +0, 3, 3, 1, 28800, 0x079bbc00 +0, 4, 4, 1, 28800, 0x3666bd3c diff --git a/tests/ref/fate/filter-scale_cuda-bgra-p010 b/tests/ref/fate/filter-scale_cuda-bgra-p010 new file mode 100644 index 0000000000..5417969495 --- /dev/null +++ b/tests/ref/fate/filter-scale_cuda-bgra-p010 @@ -0,0 +1,10 @@ +#tb 0: 1/25 +#media_type 0: video +#codec_id 0: rawvideo +#dimensions 0: 160x120 +#sar 0: 1/1 +0, 0, 0, 1, 57600, 0x732d73d1 +0, 1, 1, 1, 57600, 0xaa68a3d5 +0, 2, 2, 1, 57600, 0xe64492ed +0, 3, 3, 1, 57600, 0x815af146 +0, 4, 4, 1, 57600, 0xe3784548 diff --git a/tests/ref/fate/filter-scale_cuda-bgra-yuv420p b/tests/ref/fate/filter-scale_cuda-bgra-yuv420p new file mode 100644 index 0000000000..3ca7abbaa1 --- /dev/null +++ b/tests/ref/fate/filter-scale_cuda-bgra-yuv420p @@ -0,0 +1,10 @@ +#tb 0: 1/25 +#media_type 0: video +#codec_id 0: rawvideo +#dimensions 0: 160x120 +#sar 0: 1/1 +0, 0, 0, 1, 28800, 0xa927c72b +0, 1, 1, 1, 28800, 0x09bac057 +0, 2, 2, 1, 28800, 0x656cc8e8 +0, 3, 3, 1, 28800, 0x034bbc00 +0, 4, 4, 1, 28800, 0xa605bd3c diff --git a/tests/ref/fate/filter-scale_cuda-nv12-bgra b/tests/ref/fate/filter-scale_cuda-nv12-bgra new file mode 100644 index 0000000000..5c347f73e1 --- /dev/null +++ b/tests/ref/fate/filter-scale_cuda-nv12-bgra @@ -0,0 +1,10 @@ +#tb 0: 1/25 +#media_type 0: video +#codec_id 0: rawvideo +#dimensions 0: 160x120 +#sar 0: 1/1 +0, 0, 0, 1, 76800, 0x96543ee7 +0, 1, 1, 1, 76800, 0xbd823cd0 +0, 2, 2, 1, 76800, 0x95a663df +0, 3, 3, 1, 76800, 0x5d3e552d +0, 4, 4, 1, 76800, 0x86086305 diff --git a/tests/ref/fate/filter-scale_cuda-nv12-bgra-bt709 b/tests/ref/fate/filter-scale_cuda-nv12-bgra-bt709 new file mode 100644 index 0000000000..6b98ffa6ce --- /dev/null +++ b/tests/ref/fate/filter-scale_cuda-nv12-bgra-bt709 @@ -0,0 +1,10 @@ +#tb 0: 1/25 +#media_type 0: video +#codec_id 0: rawvideo +#dimensions 0: 160x120 +#sar 0: 1/1 +0, 0, 0, 1, 76800, 0x37efb894 +0, 1, 1, 1, 76800, 0x6adfb9f3 +0, 2, 2, 1, 76800, 0xa4b2e023 +0, 3, 3, 1, 76800, 0x6ec5d6fc +0, 4, 4, 1, 76800, 0x4cb3e85a diff --git a/tests/ref/fate/filter-scale_cuda-nv12-bgra-full b/tests/ref/fate/filter-scale_cuda-nv12-bgra-full new file mode 100644 index 0000000000..1474cdd34c --- /dev/null +++ b/tests/ref/fate/filter-scale_cuda-nv12-bgra-full @@ -0,0 +1,10 @@ +#tb 0: 1/25 +#media_type 0: video +#codec_id 0: rawvideo +#dimensions 0: 160x120 +#sar 0: 1/1 +0, 0, 0, 1, 76800, 0x87ed72f0 +0, 1, 1, 1, 76800, 0x3e5369e2 +0, 2, 2, 1, 76800, 0x3d61922a +0, 3, 3, 1, 76800, 0x491982ad +0, 4, 4, 1, 76800, 0x73459545 diff --git a/tests/ref/fate/filter-scale_cuda-nv12-nv12-scale b/tests/ref/fate/filter-scale_cuda-nv12-nv12-scale new file mode 100644 index 0000000000..939c2997d2 --- /dev/null +++ b/tests/ref/fate/filter-scale_cuda-nv12-nv12-scale @@ -0,0 +1,10 @@ +#tb 0: 1/25 +#media_type 0: video +#codec_id 0: rawvideo +#dimensions 0: 320x240 +#sar 0: 1/1 +0, 0, 0, 1, 115200, 0x59730662 +0, 1, 1, 1, 115200, 0x6ca8ee33 +0, 2, 2, 1, 115200, 0xa1fd11f9 +0, 3, 3, 1, 115200, 0x6e29e275 +0, 4, 4, 1, 115200, 0xf6e8e5ba diff --git a/tests/ref/fate/filter-scale_cuda-nv12-rgba b/tests/ref/fate/filter-scale_cuda-nv12-rgba new file mode 100644 index 0000000000..29d5b42693 --- /dev/null +++ b/tests/ref/fate/filter-scale_cuda-nv12-rgba @@ -0,0 +1,10 @@ +#tb 0: 1/25 +#media_type 0: video +#codec_id 0: rawvideo +#dimensions 0: 160x120 +#sar 0: 1/1 +0, 0, 0, 1, 76800, 0x376a3ee7 +0, 1, 1, 1, 76800, 0xed0f3cd0 +0, 2, 2, 1, 76800, 0xa9ef63df +0, 3, 3, 1, 76800, 0xd4a8552d +0, 4, 4, 1, 76800, 0xc7e26305 diff --git a/tests/ref/fate/filter-scale_cuda-odd-bgra-nv12 b/tests/ref/fate/filter-scale_cuda-odd-bgra-nv12 new file mode 100644 index 0000000000..6201113fbd --- /dev/null +++ b/tests/ref/fate/filter-scale_cuda-odd-bgra-nv12 @@ -0,0 +1,10 @@ +#tb 0: 1/25 +#media_type 0: video +#codec_id 0: rawvideo +#dimensions 0: 161x121 +#sar 0: 1/1 +0, 0, 0, 1, 29363, 0x46aee7f5 +0, 1, 1, 1, 29363, 0x775be080 +0, 2, 2, 1, 29363, 0xa995e922 +0, 3, 3, 1, 29363, 0x1a8edbe1 +0, 4, 4, 1, 29363, 0x5715dda0 diff --git a/tests/ref/fate/filter-scale_cuda-odd-bgra-yuv420p b/tests/ref/fate/filter-scale_cuda-odd-bgra-yuv420p new file mode 100644 index 0000000000..d8b02020e1 --- /dev/null +++ b/tests/ref/fate/filter-scale_cuda-odd-bgra-yuv420p @@ -0,0 +1,10 @@ +#tb 0: 1/25 +#media_type 0: video +#codec_id 0: rawvideo +#dimensions 0: 161x121 +#sar 0: 1/1 +0, 0, 0, 1, 29363, 0x876ae7f5 +0, 1, 1, 1, 29363, 0x5d79e080 +0, 2, 2, 1, 29363, 0xe379e922 +0, 3, 3, 1, 29363, 0xedfadbe1 +0, 4, 4, 1, 29363, 0xbcefdda0 diff --git a/tests/ref/fate/filter-scale_cuda-odd-nv12-bgra b/tests/ref/fate/filter-scale_cuda-odd-nv12-bgra new file mode 100644 index 0000000000..5c347f73e1 --- /dev/null +++ b/tests/ref/fate/filter-scale_cuda-odd-nv12-bgra @@ -0,0 +1,10 @@ +#tb 0: 1/25 +#media_type 0: video +#codec_id 0: rawvideo +#dimensions 0: 160x120 +#sar 0: 1/1 +0, 0, 0, 1, 76800, 0x96543ee7 +0, 1, 1, 1, 76800, 0xbd823cd0 +0, 2, 2, 1, 76800, 0x95a663df +0, 3, 3, 1, 76800, 0x5d3e552d +0, 4, 4, 1, 76800, 0x86086305 diff --git a/tests/ref/fate/filter-scale_cuda-odd-yuv420p-bgra b/tests/ref/fate/filter-scale_cuda-odd-yuv420p-bgra new file mode 100644 index 0000000000..5c347f73e1 --- /dev/null +++ b/tests/ref/fate/filter-scale_cuda-odd-yuv420p-bgra @@ -0,0 +1,10 @@ +#tb 0: 1/25 +#media_type 0: video +#codec_id 0: rawvideo +#dimensions 0: 160x120 +#sar 0: 1/1 +0, 0, 0, 1, 76800, 0x96543ee7 +0, 1, 1, 1, 76800, 0xbd823cd0 +0, 2, 2, 1, 76800, 0x95a663df +0, 3, 3, 1, 76800, 0x5d3e552d +0, 4, 4, 1, 76800, 0x86086305 diff --git a/tests/ref/fate/filter-scale_cuda-p010-bgra b/tests/ref/fate/filter-scale_cuda-p010-bgra new file mode 100644 index 0000000000..a301e41e14 --- /dev/null +++ b/tests/ref/fate/filter-scale_cuda-p010-bgra @@ -0,0 +1,10 @@ +#tb 0: 1/25 +#media_type 0: video +#codec_id 0: rawvideo +#dimensions 0: 160x120 +#sar 0: 1/1 +0, 0, 0, 1, 76800, 0xfbe39f76 +0, 1, 1, 1, 76800, 0x5e369de0 +0, 2, 2, 1, 76800, 0xdd48c539 +0, 3, 3, 1, 76800, 0x102fb716 +0, 4, 4, 1, 76800, 0xbce8c492 diff --git a/tests/ref/fate/filter-scale_cuda-rgba-nv12 b/tests/ref/fate/filter-scale_cuda-rgba-nv12 new file mode 100644 index 0000000000..8a9e4d8c80 --- /dev/null +++ b/tests/ref/fate/filter-scale_cuda-rgba-nv12 @@ -0,0 +1,10 @@ +#tb 0: 1/25 +#media_type 0: video +#codec_id 0: rawvideo +#dimensions 0: 160x120 +#sar 0: 1/1 +0, 0, 0, 1, 28800, 0xea53c72b +0, 1, 1, 1, 28800, 0x77cbc057 +0, 2, 2, 1, 28800, 0x429ec8e8 +0, 3, 3, 1, 28800, 0x079bbc00 +0, 4, 4, 1, 28800, 0x3666bd3c diff --git a/tests/ref/fate/filter-scale_cuda-roundtrip-nv12 b/tests/ref/fate/filter-scale_cuda-roundtrip-nv12 new file mode 100644 index 0000000000..c964aaad74 --- /dev/null +++ b/tests/ref/fate/filter-scale_cuda-roundtrip-nv12 @@ -0,0 +1,10 @@ +#tb 0: 1/25 +#media_type 0: video +#codec_id 0: rawvideo +#dimensions 0: 160x120 +#sar 0: 1/1 +0, 0, 0, 1, 28800, 0x38538295 +0, 1, 1, 1, 28800, 0x206c7cba +0, 2, 2, 1, 28800, 0x5fae8676 +0, 3, 3, 1, 28800, 0x8e547a90 +0, 4, 4, 1, 28800, 0x65597ad1 diff --git a/tests/ref/fate/filter-scale_cuda-roundtrip-yuv444p b/tests/ref/fate/filter-scale_cuda-roundtrip-yuv444p new file mode 100644 index 0000000000..455852b608 --- /dev/null +++ b/tests/ref/fate/filter-scale_cuda-roundtrip-yuv444p @@ -0,0 +1,10 @@ +#tb 0: 1/25 +#media_type 0: video +#codec_id 0: rawvideo +#dimensions 0: 160x120 +#sar 0: 1/1 +0, 0, 0, 1, 76800, 0x985449bc +0, 1, 1, 1, 76800, 0xda4b3f74 +0, 2, 2, 1, 76800, 0x416861cf +0, 3, 3, 1, 76800, 0x3dd24fea +0, 4, 4, 1, 76800, 0xab1c6151 diff --git a/tests/ref/fate/filter-scale_cuda-yuv420p-bgra b/tests/ref/fate/filter-scale_cuda-yuv420p-bgra new file mode 100644 index 0000000000..5c347f73e1 --- /dev/null +++ b/tests/ref/fate/filter-scale_cuda-yuv420p-bgra @@ -0,0 +1,10 @@ +#tb 0: 1/25 +#media_type 0: video +#codec_id 0: rawvideo +#dimensions 0: 160x120 +#sar 0: 1/1 +0, 0, 0, 1, 76800, 0x96543ee7 +0, 1, 1, 1, 76800, 0xbd823cd0 +0, 2, 2, 1, 76800, 0x95a663df +0, 3, 3, 1, 76800, 0x5d3e552d +0, 4, 4, 1, 76800, 0x86086305 diff --git a/tests/ref/fate/filter-scale_cuda-yuv444p-bgra b/tests/ref/fate/filter-scale_cuda-yuv444p-bgra new file mode 100644 index 0000000000..c42be9c3e2 --- /dev/null +++ b/tests/ref/fate/filter-scale_cuda-yuv444p-bgra @@ -0,0 +1,10 @@ +#tb 0: 1/25 +#media_type 0: video +#codec_id 0: rawvideo +#dimensions 0: 160x120 +#sar 0: 1/1 +0, 0, 0, 1, 76800, 0x7ddcbb72 +0, 1, 1, 1, 76800, 0x25ebb398 +0, 2, 2, 1, 76800, 0x1082d76f +0, 3, 3, 1, 76800, 0x26f4c738 +0, 4, 4, 1, 76800, 0x39cddc12 -- 2.52.0 From ffa3153fbc3182cf1c526f5f1a3c9c6dcc7c28d5 Mon Sep 17 00:00:00 2001 From: Timo Rothenpieler <[email protected]> Date: Sun, 27 Sep 2026 22:04:30 +0200 Subject: [PATCH 4/6] doc/filters: document scale_cuda enhancements --- doc/filters.texi | 19 +++++++++++++++++-- 1 file changed, 17 insertions(+), 2 deletions(-) diff --git a/doc/filters.texi b/doc/filters.texi index 04496d338f..5b234e443e 100644 --- a/doc/filters.texi +++ b/doc/filters.texi @@ -27889,8 +27889,6 @@ set with @option{param}; the default is 3. Controls the output pixel format. By default, or if none is specified, the input pixel format is used. -The filter does not support converting between YUV and RGB pixel formats. - @item passthrough If set to 0, every frame is processed, even if no conversion is necessary. This mode can be useful to use the filter as a buffer for a downstream @@ -27918,6 +27916,15 @@ Work the same as the identical @ref{scale} filter options. @item reset_sar Works the same as the identical @ref{scale} filter option. +@item out_range +Set the output YCbCr sample range used when converting from RGB to YUV. One of +@code{auto}/@code{unknown} (the default), @code{jpeg}/@code{full}/@code{pc}, or +@code{mpeg}/@code{limited}/@code{tv}. + +With @code{auto} the range is left to filter-graph negotiation; a range that is +still unspecified at that point is treated as limited. The option has no effect +when converting from YUV to RGB, which always outputs full range. + @end table @subsection Examples @@ -27942,6 +27949,14 @@ decoders frame pool. @example scale_cuda=passthrough=0 @end example + +@item +Convert to RGB for a filter that requires it, then back to NV12 with limited +range. Note the separate instances, as scaling and color space conversion cannot +be combined. +@example +scale_cuda=format=bgra,scale_cuda=format=nv12:out_range=limited +@end example @end itemize @section thumbnail_cuda -- 2.52.0 From f54e60eecdd3759d735ca4a5040b77140258dfc7 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Reinhard=20Ni=C3=9Fl?= <[email protected]> Date: Thu, 8 Oct 2026 00:34:19 +0200 Subject: [PATCH 5/6] avfilter/vf_scale_cuda: round YUV/RGB color stores MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit The 10-bit planar path already divides by 4 in float, so code 512 stays 128 and the low bits reach the matrix. The stores then truncate. Round half away from zero on the YUV to RGB cast and on the RGB to YUV stores. CUDA FATE checksums for these conversions are left unchanged; they need a GPU run after this rounding. Co-authored-by: Cursor <[email protected]> Signed-off-by: Reinhard Nißl <[email protected]> --- libavfilter/vf_scale_cuda.cu | 33 ++++++++++++++++++++------------- 1 file changed, 20 insertions(+), 13 deletions(-) diff --git a/libavfilter/vf_scale_cuda.cu b/libavfilter/vf_scale_cuda.cu index 793decb8ee..2e070ca992 100644 --- a/libavfilter/vf_scale_cuda.cu +++ b/libavfilter/vf_scale_cuda.cu @@ -1118,6 +1118,12 @@ __device__ static inline float clamp_f(float x, float lo, float hi) return x < lo ? lo : (x > hi ? hi : x); } +/* Half away from zero, then clamp. A bare cast truncates. */ +__device__ static inline float round_code(float x, float lo, float hi) +{ + return clamp_f(x + 0.5f, lo, hi); +} + // YUV to RGB CONVERSION // R = m[0][0]*(Y-Yoff) + m[0][1]*(U-128) + m[0][2]*(V-128) // G = m[1][0]*(Y-Yoff) + m[1][1]*(U-128) + m[1][2]*(V-128) @@ -1134,9 +1140,9 @@ __device__ static inline uchar4 yuv8_to_rgba_generic( float g = mat.m[1][0] * fy + mat.m[1][1] * fu + mat.m[1][2] * fv; float b = mat.m[2][0] * fy + mat.m[2][1] * fu + mat.m[2][2] * fv; - uchar rc = (uchar)clamp_f(r, 0.0f, 255.0f); - uchar gc = (uchar)clamp_f(g, 0.0f, 255.0f); - uchar bc = (uchar)clamp_f(b, 0.0f, 255.0f); + uchar rc = (uchar)round_code(r, 0.0f, 255.0f); + uchar gc = (uchar)round_code(g, 0.0f, 255.0f); + uchar bc = (uchar)round_code(b, 0.0f, 255.0f); switch (channel_order) { case 0: return make_uchar4(bc, gc, rc, 0xFF); // BGRA @@ -1301,7 +1307,7 @@ __device__ void RgbToYuv_Y(CUDAScaleKernelParams params) out_y_T *dst_y = (out_y_T*)params.dst[0]; int pitch = params.dst_pitch / sizeof(out_y_T); - dst_y[yo * pitch + xo] = (out_y_T)y; + dst_y[yo * pitch + xo] = (out_y_T)round_code(y, 0.0f, 255.0f); } // RGB->YUV Y kernel for semiplanar10/16 output @@ -1323,7 +1329,7 @@ __device__ void RgbToYuv_Y_semiplanar16(CUDAScaleKernelParams params) ushort *dst_y = (ushort*)params.dst[0]; int pitch = params.dst_pitch / sizeof(ushort); - dst_y[yo * pitch + xo] = (ushort)clamp_f(y * 256.0f, 0.0f, 65535.0f); + dst_y[yo * pitch + xo] = (ushort)round_code(y * 256.0f, 0.0f, 65535.0f); } // RGB->YUV Y kernel for planar10 output: native 10-bit Y value @@ -1345,7 +1351,7 @@ __device__ void RgbToYuv_Y_planar10(CUDAScaleKernelParams params) ushort *dst_y = (ushort*)params.dst[0]; int pitch = params.dst_pitch / sizeof(ushort); - dst_y[yo * pitch + xo] = (ushort)clamp_f(y * 4.0f, 0.0f, 1023.0f); + dst_y[yo * pitch + xo] = (ushort)round_code(y * 4.0f, 0.0f, 1023.0f); } // RGB->YUV UV kernel for semiplanar output (NV12, NV16) @@ -1385,7 +1391,8 @@ __device__ void RgbToYuv_semiplanar_UV(CUDAScaleKernelParams params) out_uv_T *dst_uv = (out_uv_T*)params.dst[1]; int pitch = params.dst_pitch / sizeof(out_uv_T); - dst_uv[yo * pitch + xo] = make_uchar2((uchar)u_avg, (uchar)v_avg); + dst_uv[yo * pitch + xo] = make_uchar2((uchar)round_code(u_avg, 0.0f, 255.0f), + (uchar)round_code(v_avg, 0.0f, 255.0f)); } // RGB->YUV UV kernel for semiplanar10/16 output @@ -1425,8 +1432,8 @@ __device__ void RgbToYuv_semiplanar16_UV(CUDAScaleKernelParams params) ushort2 *dst_uv = (ushort2*)params.dst[1]; int pitch = params.dst_pitch / sizeof(ushort2); - ushort u_val = (ushort)clamp_f(u_avg * 256.0f, 0.0f, 65535.0f); - ushort v_val = (ushort)clamp_f(v_avg * 256.0f, 0.0f, 65535.0f); + ushort u_val = (ushort)round_code(u_avg * 256.0f, 0.0f, 65535.0f); + ushort v_val = (ushort)round_code(v_avg * 256.0f, 0.0f, 65535.0f); dst_uv[yo * pitch + xo] = make_ushort2(u_val, v_val); } @@ -1470,12 +1477,12 @@ __device__ void RgbToYuv_planar_UV(CUDAScaleKernelParams params) int pitch = params.dst_pitch / sizeof(out_uv_T); if (sizeof(out_uv_T) == 1) { - dst_u[yo * pitch + xo] = (out_uv_T)u_avg; - dst_v[yo * pitch + xo] = (out_uv_T)v_avg; + dst_u[yo * pitch + xo] = (out_uv_T)round_code(u_avg, 0.0f, 255.0f); + dst_v[yo * pitch + xo] = (out_uv_T)round_code(v_avg, 0.0f, 255.0f); } else { // 10-bit planar: native 10-bit value = 8-bit * 4 - dst_u[yo * pitch + xo] = (out_uv_T)(clamp_f(u_avg * 4.0f, 0.0f, 1023.0f)); - dst_v[yo * pitch + xo] = (out_uv_T)(clamp_f(v_avg * 4.0f, 0.0f, 1023.0f)); + dst_u[yo * pitch + xo] = (out_uv_T)round_code(u_avg * 4.0f, 0.0f, 1023.0f); + dst_v[yo * pitch + xo] = (out_uv_T)round_code(v_avg * 4.0f, 0.0f, 1023.0f); } } -- 2.52.0 From b3337e205a40d79622a8f335ed968d5076f75bd6 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Reinhard=20Ni=C3=9Fl?= <[email protected]> Date: Wed, 14 Oct 2026 00:14:47 +0200 Subject: [PATCH 6/6] avfilter/vf_scale_cuda: describe the YUV matrix on the output link MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit A following scale_cuda reads the link colorspace. The conversion now writes the matrix and range it actually used, and a pure scale no longer advertises out_range. Full-range 10-bit uses the code maximum, so 1023 is white and 1020 is 254, and P010 stores that code in the high bits. BT.2020 constant luminance is rejected. Limited YCgCo uses the luma excursion for chroma. FATE checksums are from an nvcc run of the 18 conversion tests. Co-authored-by: Cursor <[email protected]> Signed-off-by: Reinhard Nißl <[email protected]> --- doc/filters.texi | 8 +- libavfilter/vf_scale_cuda.c | 51 +++++++-- libavfilter/vf_scale_cuda.cu | 105 ++++++++++++------ libavfilter/vf_scale_cuda.h | 2 + tests/fate/filter-video.mak | 6 +- tests/ref/fate/filter-scale_cuda-bgra-nv12 | 10 +- tests/ref/fate/filter-scale_cuda-bgra-p010 | 10 +- tests/ref/fate/filter-scale_cuda-bgra-yuv420p | 10 +- tests/ref/fate/filter-scale_cuda-nv12-bgra | 10 +- .../fate/filter-scale_cuda-nv12-bgra-bt709 | 10 +- .../ref/fate/filter-scale_cuda-nv12-bgra-full | 10 +- tests/ref/fate/filter-scale_cuda-nv12-rgba | 10 +- .../ref/fate/filter-scale_cuda-odd-bgra-nv12 | 10 +- .../fate/filter-scale_cuda-odd-bgra-yuv420p | 10 +- .../ref/fate/filter-scale_cuda-odd-nv12-bgra | 10 +- .../fate/filter-scale_cuda-odd-yuv420p-bgra | 10 +- tests/ref/fate/filter-scale_cuda-p010-bgra | 10 +- tests/ref/fate/filter-scale_cuda-rgba-nv12 | 10 +- .../ref/fate/filter-scale_cuda-roundtrip-nv12 | 10 +- .../fate/filter-scale_cuda-roundtrip-yuv444p | 10 +- tests/ref/fate/filter-scale_cuda-yuv420p-bgra | 10 +- tests/ref/fate/filter-scale_cuda-yuv444p-bgra | 10 +- 22 files changed, 208 insertions(+), 134 deletions(-) diff --git a/doc/filters.texi b/doc/filters.texi index 5b234e443e..c8bd21a3bb 100644 --- a/doc/filters.texi +++ b/doc/filters.texi @@ -27918,12 +27918,12 @@ Works the same as the identical @ref{scale} filter option. @item out_range Set the output YCbCr sample range used when converting from RGB to YUV. One of -@code{auto}/@code{unknown} (the default), @code{jpeg}/@code{full}/@code{pc}, or +@code{auto} (the default), @code{jpeg}/@code{full}/@code{pc}, or @code{mpeg}/@code{limited}/@code{tv}. -With @code{auto} the range is left to filter-graph negotiation; a range that is -still unspecified at that point is treated as limited. The option has no effect -when converting from YUV to RGB, which always outputs full range. +With @code{auto} the output range is limited. The option has no effect +when converting from YUV to RGB, which always outputs full range, or when +the filter is not converting between RGB and YUV. @end table diff --git a/libavfilter/vf_scale_cuda.c b/libavfilter/vf_scale_cuda.c index 046d396122..a570c1bbe1 100644 --- a/libavfilter/vf_scale_cuda.c +++ b/libavfilter/vf_scale_cuda.c @@ -468,6 +468,14 @@ static const char* get_format_name(enum AVPixelFormat fmt) return NULL; } +/* H.273 YCgCo uses the luma excursion for chroma as well. */ +static int chroma_uses_luma_span(enum AVColorSpace colorspace) +{ + return colorspace == AVCOL_SPC_YCGCO + || colorspace == AVCOL_SPC_YCGCO_RE + || colorspace == AVCOL_SPC_YCGCO_RO; +} + static int compute_yuv2rgb_matrix(void *log_ctx, CUDAScaleColorMatrix *mat, enum AVColorSpace colorspace, bool limited_range) @@ -477,7 +485,8 @@ static int compute_yuv2rgb_matrix(void *log_ctx, CUDAScaleColorMatrix *mat, float uv_scale = 1.0f; double rgb2yuv[3][3], yuv2rgb[3][3]; - if (!coeffs) { + /* Constant luminance is not the matrix returned for these coefficients. */ + if (!coeffs || colorspace == AVCOL_SPC_BT2020_CL) { av_log(log_ctx, AV_LOG_ERROR, "Unsupported colorspace %d for YUV->RGB conversion\n", colorspace); return AVERROR(EINVAL); @@ -488,7 +497,8 @@ static int compute_yuv2rgb_matrix(void *log_ctx, CUDAScaleColorMatrix *mat, if (limited_range) { y_scale = 255.0f / (235.0f - 16.0f); - uv_scale = 255.0f / (240.0f - 16.0f); + uv_scale = chroma_uses_luma_span(colorspace) ? y_scale + : 255.0f / (240.0f - 16.0f); } for (int i = 0; i < 3; i++) { @@ -509,7 +519,7 @@ static int compute_rgb2yuv_matrix(void *log_ctx, CUDAScaleColorMatrix *mat, float y_scale = 1.0f; float uv_scale = 1.0f; - if (!coeffs) { + if (!coeffs || colorspace == AVCOL_SPC_BT2020_CL) { av_log(log_ctx, AV_LOG_ERROR, "Unsupported colorspace %d for RGB->YUV conversion\n", colorspace); return AVERROR(EINVAL); @@ -519,7 +529,8 @@ static int compute_rgb2yuv_matrix(void *log_ctx, CUDAScaleColorMatrix *mat, if (limited_range) { y_scale = (235.0f - 16.0f) / 255.0f; - uv_scale = (240.0f - 16.0f) / 255.0f; + uv_scale = chroma_uses_luma_span(colorspace) ? y_scale + : (240.0f - 16.0f) / 255.0f; } for (int i = 0; i < 3; i++) { @@ -605,6 +616,7 @@ static av_cold int init_processing_chain(AVFilterContext *ctx, int in_width, int return AVERROR(ENOSYS); } + AVFilterLink *outlink = ctx->outputs[0]; enum AVColorSpace cs = inlink->colorspace; bool limited_range; @@ -622,6 +634,11 @@ static av_cold int init_processing_chain(AVFilterContext *ctx, int in_width, int s->colorspace = cs; s->color_range = limited_range ? AVCOL_RANGE_MPEG : AVCOL_RANGE_JPEG; + /* The next filter reads the link, not the frame tag. RGB output is + full range. The matrix that was applied stays on s->colorspace + only as the input description; the output link is RGB. */ + outlink->colorspace = AVCOL_SPC_RGB; + outlink->color_range = AVCOL_RANGE_JPEG; /* YUV->RGB only supports Nearest for now. Forcing it here also makes * the generic filter path get disabled below. */ @@ -658,14 +675,23 @@ static av_cold int init_processing_chain(AVFilterContext *ctx, int in_width, int if (cs == AVCOL_SPC_UNSPECIFIED || cs == AVCOL_SPC_RGB) cs = AVCOL_SPC_SMPTE170M; - limited_range = (outlink->color_range != AVCOL_RANGE_JPEG); + /* Unspecified out_range is limited. An explicit jpeg/full wins. + The output link has to carry that choice: a later scale_cuda + builds its matrix from the link, and a pure scale must not + inherit out_range from query_formats. */ + if (s->out_range == AVCOL_RANGE_JPEG) + limited_range = false; + else + limited_range = true; ret = compute_rgb2yuv_matrix(ctx, &s->color_matrix, cs, limited_range); if (ret < 0) return ret; s->colorspace = cs; - s->color_range = outlink->color_range; + s->color_range = limited_range ? AVCOL_RANGE_MPEG : AVCOL_RANGE_JPEG; + outlink->colorspace = cs; + outlink->color_range = s->color_range; /* RGB->YUV only supports Nearest for now. Forcing it here also makes * the generic filter path get disabled below. */ @@ -1126,9 +1152,10 @@ static int cudascale_query_formats(const AVFilterContext *ctx, if ((ret = ff_set_common_formats_from_list2(ctx, cfg_in, cfg_out, pix_fmts)) < 0) return ret; - formats = s->out_range != AVCOL_RANGE_UNSPECIFIED - ? ff_make_formats_list_singleton(s->out_range) - : ff_all_color_ranges(); + /* Range and matrix are written on the output link when a conversion + actually runs. Advertising out_range here also relabels a scale + that leaves the samples unchanged. */ + formats = ff_all_color_ranges(); if ((ret = ff_formats_ref(formats, &cfg_out[0]->color_ranges)) < 0) return ret; @@ -1337,6 +1364,8 @@ static int call_resize_kernel(AVFilterContext *ctx, CUfunction func, .src_storage_max = ((1U << s->in_desc->comp[0].depth) - 1) << s->in_desc->comp[0].shift, .color_matrix = s->color_matrix, + .dst_depth = s->out_desc->comp[0].depth, + .dst_shift = s->out_desc->comp[0].shift, }; if (filter) { @@ -1497,9 +1526,13 @@ static int cudascale_scale(AVFilterContext *ctx, AVFrame *out, AVFrame *in, if (s->yuv2rgb) { out->colorspace = AVCOL_SPC_RGB; out->color_range = AVCOL_RANGE_JPEG; + out->chroma_location = AVCHROMA_LOC_UNSPECIFIED; } else if (s->rgb2yuv) { out->colorspace = s->colorspace; out->color_range = s->color_range; + out->chroma_location = (s->out_desc->log2_chroma_w || s->out_desc->log2_chroma_h) + ? AVCHROMA_LOC_CENTER + : AVCHROMA_LOC_UNSPECIFIED; } out->crop_top = 0; diff --git a/libavfilter/vf_scale_cuda.cu b/libavfilter/vf_scale_cuda.cu index 2e070ca992..187d40f116 100644 --- a/libavfilter/vf_scale_cuda.cu +++ b/libavfilter/vf_scale_cuda.cu @@ -1124,6 +1124,58 @@ __device__ static inline float round_code(float x, float lo, float hi) return clamp_f(x + 0.5f, lo, hi); } +/* Stored samples enter the 8-bit matrix domain. + Limited codes are the 8-bit codes shifted by depth-8. Full-range + codes are scaled by the code maximum, so 1023 is 255 and the + chroma midpoint stays 128. P010 keeps that code in the high bits. */ +__device__ static inline void yuv_codes_to_u8(float raw_y, float raw_u, float raw_v, + int depth, int storage_max, int mpeg_range, + float &fy, float &fu, float &fv) +{ + if (depth <= 8) { + fy = raw_y; + fu = raw_u; + fv = raw_v; + } else if (mpeg_range) { + float code_div = (float)(1 << (depth - 8)) * + ((float)storage_max / (float)((1 << depth) - 1)); + fy = raw_y / code_div; + fu = raw_u / code_div; + fv = raw_v / code_div; + } else { + float peak = (float)((1 << depth) - 1); + float mid = (float)(1 << (depth - 1)); + float to_code = peak / (float)storage_max; + + fy = raw_y * 255.f / (float)storage_max; + fu = (raw_u * to_code - mid) * (255.f / peak) + 128.f; + fv = (raw_v * to_code - mid) * (255.f / peak) + 128.f; + } + if (mpeg_range) + fy -= 16.f; + fu -= 128.f; + fv -= 128.f; +} + +/* 8-bit matrix result back to a stored code. The shift clears padding. + Full-range luma maps 255 onto the code maximum. Full-range chroma + keeps the midpoint (128 -> 512 at 10-bit), which a plain multiply + by 1023/255 would move. */ +__device__ static inline ushort u8_to_stored(float code8, int depth, int shift, + int mpeg_range, int chroma) +{ + float peak = (float)((1 << depth) - 1); + float code; + + if (mpeg_range) + code = code8 * (float)(1 << (depth - 8)); + else if (chroma) + code = (code8 - 128.f) * peak / 255.f + (float)(1 << (depth - 1)); + else + code = code8 * peak / 255.f; + return (ushort)round_code(code, 0.f, peak) << shift; +} + // YUV to RGB CONVERSION // R = m[0][0]*(Y-Yoff) + m[0][1]*(U-128) + m[0][2]*(V-128) // G = m[1][0]*(Y-Yoff) + m[1][1]*(U-128) + m[1][2]*(V-128) @@ -1176,20 +1228,9 @@ __device__ void YuvToRgb_semiplanar(CUDAScaleKernelParams params) in_uv_T uv_raw = tex2D<in_uv_T>(params.src_tex[1], cx, cy); float fy, fu, fv; - if (sizeof(in_y_T) == 1) { - fy = (float)(int)y_raw; - fu = (float)(int)uv_raw.x; - fv = (float)(int)uv_raw.y; - } else { - // 10/16-bit semiplanar: normalize to 8-bit range by dividing by 256 - fy = (float)(int)y_raw / 256.0f; - fu = (float)(int)uv_raw.x / 256.0f; - fv = (float)(int)uv_raw.y / 256.0f; - } - - if (params.mpeg_range) fy -= 16.0f; - fu -= 128.0f; - fv -= 128.0f; + yuv_codes_to_u8((float)(int)y_raw, (float)(int)uv_raw.x, (float)(int)uv_raw.y, + params.src_depth, params.src_storage_max, params.mpeg_range, + fy, fu, fv); uchar4 *dst = (uchar4*)params.dst[0]; int dst_pitch = params.dst_pitch / sizeof(uchar4); @@ -1221,20 +1262,9 @@ __device__ void YuvToRgb_planar(CUDAScaleKernelParams params) in_uv_T v_raw = tex2D<in_uv_T>(params.src_tex[2], cx, cy); float fy, fu, fv; - if (sizeof(in_y_T) == 1) { - fy = (float)(int)y_raw; - fu = (float)(int)u_raw; - fv = (float)(int)v_raw; - } else { - // planar10: 10-bit native values, divide by 4 - fy = (float)(int)y_raw / 4.0f; - fu = (float)(int)u_raw / 4.0f; - fv = (float)(int)v_raw / 4.0f; - } - - if (params.mpeg_range) fy -= 16.0f; - fu -= 128.0f; - fv -= 128.0f; + yuv_codes_to_u8((float)(int)y_raw, (float)(int)u_raw, (float)(int)v_raw, + params.src_depth, params.src_storage_max, params.mpeg_range, + fy, fu, fv); uchar4 *dst = (uchar4*)params.dst[0]; int dst_pitch = params.dst_pitch / sizeof(uchar4); @@ -1329,7 +1359,8 @@ __device__ void RgbToYuv_Y_semiplanar16(CUDAScaleKernelParams params) ushort *dst_y = (ushort*)params.dst[0]; int pitch = params.dst_pitch / sizeof(ushort); - dst_y[yo * pitch + xo] = (ushort)round_code(y * 256.0f, 0.0f, 65535.0f); + dst_y[yo * pitch + xo] = u8_to_stored(y, params.dst_depth, params.dst_shift, + params.mpeg_range, 0); } // RGB->YUV Y kernel for planar10 output: native 10-bit Y value @@ -1351,7 +1382,8 @@ __device__ void RgbToYuv_Y_planar10(CUDAScaleKernelParams params) ushort *dst_y = (ushort*)params.dst[0]; int pitch = params.dst_pitch / sizeof(ushort); - dst_y[yo * pitch + xo] = (ushort)round_code(y * 4.0f, 0.0f, 1023.0f); + dst_y[yo * pitch + xo] = u8_to_stored(y, params.dst_depth, params.dst_shift, + params.mpeg_range, 0); } // RGB->YUV UV kernel for semiplanar output (NV12, NV16) @@ -1432,8 +1464,10 @@ __device__ void RgbToYuv_semiplanar16_UV(CUDAScaleKernelParams params) ushort2 *dst_uv = (ushort2*)params.dst[1]; int pitch = params.dst_pitch / sizeof(ushort2); - ushort u_val = (ushort)round_code(u_avg * 256.0f, 0.0f, 65535.0f); - ushort v_val = (ushort)round_code(v_avg * 256.0f, 0.0f, 65535.0f); + ushort u_val = u8_to_stored(u_avg, params.dst_depth, params.dst_shift, + params.mpeg_range, 1); + ushort v_val = u8_to_stored(v_avg, params.dst_depth, params.dst_shift, + params.mpeg_range, 1); dst_uv[yo * pitch + xo] = make_ushort2(u_val, v_val); } @@ -1480,9 +1514,10 @@ __device__ void RgbToYuv_planar_UV(CUDAScaleKernelParams params) dst_u[yo * pitch + xo] = (out_uv_T)round_code(u_avg, 0.0f, 255.0f); dst_v[yo * pitch + xo] = (out_uv_T)round_code(v_avg, 0.0f, 255.0f); } else { - // 10-bit planar: native 10-bit value = 8-bit * 4 - dst_u[yo * pitch + xo] = (out_uv_T)round_code(u_avg * 4.0f, 0.0f, 1023.0f); - dst_v[yo * pitch + xo] = (out_uv_T)round_code(v_avg * 4.0f, 0.0f, 1023.0f); + dst_u[yo * pitch + xo] = (out_uv_T)u8_to_stored(u_avg, params.dst_depth, + params.dst_shift, params.mpeg_range, 1); + dst_v[yo * pitch + xo] = (out_uv_T)u8_to_stored(v_avg, params.dst_depth, + params.dst_shift, params.mpeg_range, 1); } } diff --git a/libavfilter/vf_scale_cuda.h b/libavfilter/vf_scale_cuda.h index a2a2bcb9a5..b4c62844f6 100644 --- a/libavfilter/vf_scale_cuda.h +++ b/libavfilter/vf_scale_cuda.h @@ -59,6 +59,8 @@ typedef struct { int log2_chroma_w; int log2_chroma_h; CUDAScaleColorMatrix color_matrix; + int dst_depth; + int dst_shift; } CUDAScaleKernelParams; #endif diff --git a/tests/fate/filter-video.mak b/tests/fate/filter-video.mak index 84afbcc2a2..95d4396477 100644 --- a/tests/fate/filter-video.mak +++ b/tests/fate/filter-video.mak @@ -1023,7 +1023,9 @@ fate-vfilter: $(FATE_FILTER-yes) $(FATE_FILTER_SAMPLES-yes) $(FATE_FILTER_VSYNTH # make fate-filter-scale_cuda HWACCEL=cuda # SCALE_CUDA_DEPS = SCALE_CUDA_FILTER HWUPLOAD_CUDA_FILTER HWDOWNLOAD_FILTER \ - FORMAT_FILTER TESTSRC2_FILTER HWACCEL_CUDA_TESTS + FORMAT_FILTER TESTSRC2_FILTER HWACCEL_CUDA_TESTS \ + FFMPEG RAWVIDEO_ENCODER FRAMECRC_MUXER PIPE_PROTOCOL \ + SETPARAMS_FILTER SCALE_CUDA_ALLYES = $(call ALLYES, $(SCALE_CUDA_DEPS)) @@ -1069,6 +1071,8 @@ fate-filter-scale_cuda-odd-yuv420p-bgra: LAVFI = $(call sc_yuv_rgb,yuv420p,bgra) FATE_HW-yes += $(FATE_SCALE_CUDA-yes) +$(FATE_SCALE_CUDA-yes): $(FFMPEG) + fate-filter-scale_cuda: $(FATE_SCALE_CUDA-yes) diff --git a/tests/ref/fate/filter-scale_cuda-bgra-nv12 b/tests/ref/fate/filter-scale_cuda-bgra-nv12 index 8a9e4d8c80..356f15e1de 100644 --- a/tests/ref/fate/filter-scale_cuda-bgra-nv12 +++ b/tests/ref/fate/filter-scale_cuda-bgra-nv12 @@ -3,8 +3,8 @@ #codec_id 0: rawvideo #dimensions 0: 160x120 #sar 0: 1/1 -0, 0, 0, 1, 28800, 0xea53c72b -0, 1, 1, 1, 28800, 0x77cbc057 -0, 2, 2, 1, 28800, 0x429ec8e8 -0, 3, 3, 1, 28800, 0x079bbc00 -0, 4, 4, 1, 28800, 0x3666bd3c +0, 0, 0, 1, 28800, 0xe07afcc1 +0, 1, 1, 1, 28800, 0x3442f609 +0, 2, 2, 1, 28800, 0x67fefe95 +0, 3, 3, 1, 28800, 0x1865f1d5 +0, 4, 4, 1, 28800, 0x4381f390 diff --git a/tests/ref/fate/filter-scale_cuda-bgra-p010 b/tests/ref/fate/filter-scale_cuda-bgra-p010 index 5417969495..7d06390a98 100644 --- a/tests/ref/fate/filter-scale_cuda-bgra-p010 +++ b/tests/ref/fate/filter-scale_cuda-bgra-p010 @@ -3,8 +3,8 @@ #codec_id 0: rawvideo #dimensions 0: 160x120 #sar 0: 1/1 -0, 0, 0, 1, 57600, 0x732d73d1 -0, 1, 1, 1, 57600, 0xaa68a3d5 -0, 2, 2, 1, 57600, 0xe64492ed -0, 3, 3, 1, 57600, 0x815af146 -0, 4, 4, 1, 57600, 0xe3784548 +0, 0, 0, 1, 57600, 0x6fbb0d63 +0, 1, 1, 1, 57600, 0x60a0233d +0, 2, 2, 1, 57600, 0x32e1bcd3 +0, 3, 3, 1, 57600, 0xca98b8a8 +0, 4, 4, 1, 57600, 0xb164a9c3 diff --git a/tests/ref/fate/filter-scale_cuda-bgra-yuv420p b/tests/ref/fate/filter-scale_cuda-bgra-yuv420p index 3ca7abbaa1..0c04668d7d 100644 --- a/tests/ref/fate/filter-scale_cuda-bgra-yuv420p +++ b/tests/ref/fate/filter-scale_cuda-bgra-yuv420p @@ -3,8 +3,8 @@ #codec_id 0: rawvideo #dimensions 0: 160x120 #sar 0: 1/1 -0, 0, 0, 1, 28800, 0xa927c72b -0, 1, 1, 1, 28800, 0x09bac057 -0, 2, 2, 1, 28800, 0x656cc8e8 -0, 3, 3, 1, 28800, 0x034bbc00 -0, 4, 4, 1, 28800, 0xa605bd3c +0, 0, 0, 1, 28800, 0xac0efcc1 +0, 1, 1, 1, 28800, 0x4d54f609 +0, 2, 2, 1, 28800, 0x297ffe95 +0, 3, 3, 1, 28800, 0x2d2df1d5 +0, 4, 4, 1, 28800, 0xc5b8f390 diff --git a/tests/ref/fate/filter-scale_cuda-nv12-bgra b/tests/ref/fate/filter-scale_cuda-nv12-bgra index 5c347f73e1..d67d3fdb24 100644 --- a/tests/ref/fate/filter-scale_cuda-nv12-bgra +++ b/tests/ref/fate/filter-scale_cuda-nv12-bgra @@ -3,8 +3,8 @@ #codec_id 0: rawvideo #dimensions 0: 160x120 #sar 0: 1/1 -0, 0, 0, 1, 76800, 0x96543ee7 -0, 1, 1, 1, 76800, 0xbd823cd0 -0, 2, 2, 1, 76800, 0x95a663df -0, 3, 3, 1, 76800, 0x5d3e552d -0, 4, 4, 1, 76800, 0x86086305 +0, 0, 0, 1, 76800, 0x1bad7907 +0, 1, 1, 1, 76800, 0x9bd3774d +0, 2, 2, 1, 76800, 0xa1129eee +0, 3, 3, 1, 76800, 0x04458f95 +0, 4, 4, 1, 76800, 0xc1ca9de7 diff --git a/tests/ref/fate/filter-scale_cuda-nv12-bgra-bt709 b/tests/ref/fate/filter-scale_cuda-nv12-bgra-bt709 index 6b98ffa6ce..a8831d3ead 100644 --- a/tests/ref/fate/filter-scale_cuda-nv12-bgra-bt709 +++ b/tests/ref/fate/filter-scale_cuda-nv12-bgra-bt709 @@ -3,8 +3,8 @@ #codec_id 0: rawvideo #dimensions 0: 160x120 #sar 0: 1/1 -0, 0, 0, 1, 76800, 0x37efb894 -0, 1, 1, 1, 76800, 0x6adfb9f3 -0, 2, 2, 1, 76800, 0xa4b2e023 -0, 3, 3, 1, 76800, 0x6ec5d6fc -0, 4, 4, 1, 76800, 0x4cb3e85a +0, 0, 0, 1, 76800, 0x5a8fe7c7 +0, 1, 1, 1, 76800, 0xa05ee72f +0, 2, 2, 1, 76800, 0xbd650e8a +0, 3, 3, 1, 76800, 0xa9c0052e +0, 4, 4, 1, 76800, 0xb9c81713 diff --git a/tests/ref/fate/filter-scale_cuda-nv12-bgra-full b/tests/ref/fate/filter-scale_cuda-nv12-bgra-full index 1474cdd34c..a68d66f92f 100644 --- a/tests/ref/fate/filter-scale_cuda-nv12-bgra-full +++ b/tests/ref/fate/filter-scale_cuda-nv12-bgra-full @@ -3,8 +3,8 @@ #codec_id 0: rawvideo #dimensions 0: 160x120 #sar 0: 1/1 -0, 0, 0, 1, 76800, 0x87ed72f0 -0, 1, 1, 1, 76800, 0x3e5369e2 -0, 2, 2, 1, 76800, 0x3d61922a -0, 3, 3, 1, 76800, 0x491982ad -0, 4, 4, 1, 76800, 0x73459545 +0, 0, 0, 1, 76800, 0xa3b31af3 +0, 1, 1, 1, 76800, 0x170314e5 +0, 2, 2, 1, 76800, 0x092f3b3b +0, 3, 3, 1, 76800, 0x3b362d6c +0, 4, 4, 1, 76800, 0x69893cdb diff --git a/tests/ref/fate/filter-scale_cuda-nv12-rgba b/tests/ref/fate/filter-scale_cuda-nv12-rgba index 29d5b42693..c6ae2055db 100644 --- a/tests/ref/fate/filter-scale_cuda-nv12-rgba +++ b/tests/ref/fate/filter-scale_cuda-nv12-rgba @@ -3,8 +3,8 @@ #codec_id 0: rawvideo #dimensions 0: 160x120 #sar 0: 1/1 -0, 0, 0, 1, 76800, 0x376a3ee7 -0, 1, 1, 1, 76800, 0xed0f3cd0 -0, 2, 2, 1, 76800, 0xa9ef63df -0, 3, 3, 1, 76800, 0xd4a8552d -0, 4, 4, 1, 76800, 0xc7e26305 +0, 0, 0, 1, 76800, 0xe9ea7907 +0, 1, 1, 1, 76800, 0xf7fc774d +0, 2, 2, 1, 76800, 0xe2c39eee +0, 3, 3, 1, 76800, 0xa9598f95 +0, 4, 4, 1, 76800, 0x32359de7 diff --git a/tests/ref/fate/filter-scale_cuda-odd-bgra-nv12 b/tests/ref/fate/filter-scale_cuda-odd-bgra-nv12 index 6201113fbd..b93db9d358 100644 --- a/tests/ref/fate/filter-scale_cuda-odd-bgra-nv12 +++ b/tests/ref/fate/filter-scale_cuda-odd-bgra-nv12 @@ -3,8 +3,8 @@ #codec_id 0: rawvideo #dimensions 0: 161x121 #sar 0: 1/1 -0, 0, 0, 1, 29363, 0x46aee7f5 -0, 1, 1, 1, 29363, 0x775be080 -0, 2, 2, 1, 29363, 0xa995e922 -0, 3, 3, 1, 29363, 0x1a8edbe1 -0, 4, 4, 1, 29363, 0x5715dda0 +0, 0, 0, 1, 29363, 0x031a1e9b +0, 1, 1, 1, 29363, 0x145716f8 +0, 2, 2, 1, 29363, 0xf9952059 +0, 3, 3, 1, 29363, 0xe20b11e3 +0, 4, 4, 1, 29363, 0xb65e14bd diff --git a/tests/ref/fate/filter-scale_cuda-odd-bgra-yuv420p b/tests/ref/fate/filter-scale_cuda-odd-bgra-yuv420p index d8b02020e1..cfde71820b 100644 --- a/tests/ref/fate/filter-scale_cuda-odd-bgra-yuv420p +++ b/tests/ref/fate/filter-scale_cuda-odd-bgra-yuv420p @@ -3,8 +3,8 @@ #codec_id 0: rawvideo #dimensions 0: 161x121 #sar 0: 1/1 -0, 0, 0, 1, 29363, 0x876ae7f5 -0, 1, 1, 1, 29363, 0x5d79e080 -0, 2, 2, 1, 29363, 0xe379e922 -0, 3, 3, 1, 29363, 0xedfadbe1 -0, 4, 4, 1, 29363, 0xbcefdda0 +0, 0, 0, 1, 29363, 0x92691e9b +0, 1, 1, 1, 29363, 0x707a16f8 +0, 2, 2, 1, 29363, 0x881c2059 +0, 3, 3, 1, 29363, 0x759911e3 +0, 4, 4, 1, 29363, 0x6e5314bd diff --git a/tests/ref/fate/filter-scale_cuda-odd-nv12-bgra b/tests/ref/fate/filter-scale_cuda-odd-nv12-bgra index 5c347f73e1..d67d3fdb24 100644 --- a/tests/ref/fate/filter-scale_cuda-odd-nv12-bgra +++ b/tests/ref/fate/filter-scale_cuda-odd-nv12-bgra @@ -3,8 +3,8 @@ #codec_id 0: rawvideo #dimensions 0: 160x120 #sar 0: 1/1 -0, 0, 0, 1, 76800, 0x96543ee7 -0, 1, 1, 1, 76800, 0xbd823cd0 -0, 2, 2, 1, 76800, 0x95a663df -0, 3, 3, 1, 76800, 0x5d3e552d -0, 4, 4, 1, 76800, 0x86086305 +0, 0, 0, 1, 76800, 0x1bad7907 +0, 1, 1, 1, 76800, 0x9bd3774d +0, 2, 2, 1, 76800, 0xa1129eee +0, 3, 3, 1, 76800, 0x04458f95 +0, 4, 4, 1, 76800, 0xc1ca9de7 diff --git a/tests/ref/fate/filter-scale_cuda-odd-yuv420p-bgra b/tests/ref/fate/filter-scale_cuda-odd-yuv420p-bgra index 5c347f73e1..d67d3fdb24 100644 --- a/tests/ref/fate/filter-scale_cuda-odd-yuv420p-bgra +++ b/tests/ref/fate/filter-scale_cuda-odd-yuv420p-bgra @@ -3,8 +3,8 @@ #codec_id 0: rawvideo #dimensions 0: 160x120 #sar 0: 1/1 -0, 0, 0, 1, 76800, 0x96543ee7 -0, 1, 1, 1, 76800, 0xbd823cd0 -0, 2, 2, 1, 76800, 0x95a663df -0, 3, 3, 1, 76800, 0x5d3e552d -0, 4, 4, 1, 76800, 0x86086305 +0, 0, 0, 1, 76800, 0x1bad7907 +0, 1, 1, 1, 76800, 0x9bd3774d +0, 2, 2, 1, 76800, 0xa1129eee +0, 3, 3, 1, 76800, 0x04458f95 +0, 4, 4, 1, 76800, 0xc1ca9de7 diff --git a/tests/ref/fate/filter-scale_cuda-p010-bgra b/tests/ref/fate/filter-scale_cuda-p010-bgra index a301e41e14..e855f7d700 100644 --- a/tests/ref/fate/filter-scale_cuda-p010-bgra +++ b/tests/ref/fate/filter-scale_cuda-p010-bgra @@ -3,8 +3,8 @@ #codec_id 0: rawvideo #dimensions 0: 160x120 #sar 0: 1/1 -0, 0, 0, 1, 76800, 0xfbe39f76 -0, 1, 1, 1, 76800, 0x5e369de0 -0, 2, 2, 1, 76800, 0xdd48c539 -0, 3, 3, 1, 76800, 0x102fb716 -0, 4, 4, 1, 76800, 0xbce8c492 +0, 0, 0, 1, 76800, 0x3c5fd1a5 +0, 1, 1, 1, 76800, 0x42cad0ba +0, 2, 2, 1, 76800, 0x0084f803 +0, 3, 3, 1, 76800, 0x9db5eaac +0, 4, 4, 1, 76800, 0xadf0f8bd diff --git a/tests/ref/fate/filter-scale_cuda-rgba-nv12 b/tests/ref/fate/filter-scale_cuda-rgba-nv12 index 8a9e4d8c80..356f15e1de 100644 --- a/tests/ref/fate/filter-scale_cuda-rgba-nv12 +++ b/tests/ref/fate/filter-scale_cuda-rgba-nv12 @@ -3,8 +3,8 @@ #codec_id 0: rawvideo #dimensions 0: 160x120 #sar 0: 1/1 -0, 0, 0, 1, 28800, 0xea53c72b -0, 1, 1, 1, 28800, 0x77cbc057 -0, 2, 2, 1, 28800, 0x429ec8e8 -0, 3, 3, 1, 28800, 0x079bbc00 -0, 4, 4, 1, 28800, 0x3666bd3c +0, 0, 0, 1, 28800, 0xe07afcc1 +0, 1, 1, 1, 28800, 0x3442f609 +0, 2, 2, 1, 28800, 0x67fefe95 +0, 3, 3, 1, 28800, 0x1865f1d5 +0, 4, 4, 1, 28800, 0x4381f390 diff --git a/tests/ref/fate/filter-scale_cuda-roundtrip-nv12 b/tests/ref/fate/filter-scale_cuda-roundtrip-nv12 index c964aaad74..ed810379e2 100644 --- a/tests/ref/fate/filter-scale_cuda-roundtrip-nv12 +++ b/tests/ref/fate/filter-scale_cuda-roundtrip-nv12 @@ -3,8 +3,8 @@ #codec_id 0: rawvideo #dimensions 0: 160x120 #sar 0: 1/1 -0, 0, 0, 1, 28800, 0x38538295 -0, 1, 1, 1, 28800, 0x206c7cba -0, 2, 2, 1, 28800, 0x5fae8676 -0, 3, 3, 1, 28800, 0x8e547a90 -0, 4, 4, 1, 28800, 0x65597ad1 +0, 0, 0, 1, 28800, 0x5d58cb67 +0, 1, 1, 1, 28800, 0x9c84c57a +0, 2, 2, 1, 28800, 0xda77cf55 +0, 3, 3, 1, 28800, 0x07cfc371 +0, 4, 4, 1, 28800, 0xe1fcc3c9 diff --git a/tests/ref/fate/filter-scale_cuda-roundtrip-yuv444p b/tests/ref/fate/filter-scale_cuda-roundtrip-yuv444p index 455852b608..6a148c13d7 100644 --- a/tests/ref/fate/filter-scale_cuda-roundtrip-yuv444p +++ b/tests/ref/fate/filter-scale_cuda-roundtrip-yuv444p @@ -3,8 +3,8 @@ #codec_id 0: rawvideo #dimensions 0: 160x120 #sar 0: 1/1 -0, 0, 0, 1, 76800, 0x985449bc -0, 1, 1, 1, 76800, 0xda4b3f74 -0, 2, 2, 1, 76800, 0x416861cf -0, 3, 3, 1, 76800, 0x3dd24fea -0, 4, 4, 1, 76800, 0xab1c6151 +0, 0, 0, 1, 76800, 0x4e8ef7dd +0, 1, 1, 1, 76800, 0xf1dbefed +0, 2, 2, 1, 76800, 0xc70a13f4 +0, 3, 3, 1, 76800, 0xc6f40342 +0, 4, 4, 1, 76800, 0x0fcd185a diff --git a/tests/ref/fate/filter-scale_cuda-yuv420p-bgra b/tests/ref/fate/filter-scale_cuda-yuv420p-bgra index 5c347f73e1..d67d3fdb24 100644 --- a/tests/ref/fate/filter-scale_cuda-yuv420p-bgra +++ b/tests/ref/fate/filter-scale_cuda-yuv420p-bgra @@ -3,8 +3,8 @@ #codec_id 0: rawvideo #dimensions 0: 160x120 #sar 0: 1/1 -0, 0, 0, 1, 76800, 0x96543ee7 -0, 1, 1, 1, 76800, 0xbd823cd0 -0, 2, 2, 1, 76800, 0x95a663df -0, 3, 3, 1, 76800, 0x5d3e552d -0, 4, 4, 1, 76800, 0x86086305 +0, 0, 0, 1, 76800, 0x1bad7907 +0, 1, 1, 1, 76800, 0x9bd3774d +0, 2, 2, 1, 76800, 0xa1129eee +0, 3, 3, 1, 76800, 0x04458f95 +0, 4, 4, 1, 76800, 0xc1ca9de7 diff --git a/tests/ref/fate/filter-scale_cuda-yuv444p-bgra b/tests/ref/fate/filter-scale_cuda-yuv444p-bgra index c42be9c3e2..d63ba75de2 100644 --- a/tests/ref/fate/filter-scale_cuda-yuv444p-bgra +++ b/tests/ref/fate/filter-scale_cuda-yuv444p-bgra @@ -3,8 +3,8 @@ #codec_id 0: rawvideo #dimensions 0: 160x120 #sar 0: 1/1 -0, 0, 0, 1, 76800, 0x7ddcbb72 -0, 1, 1, 1, 76800, 0x25ebb398 -0, 2, 2, 1, 76800, 0x1082d76f -0, 3, 3, 1, 76800, 0x26f4c738 -0, 4, 4, 1, 76800, 0x39cddc12 +0, 0, 0, 1, 76800, 0x69e7f5cb +0, 1, 1, 1, 76800, 0x7ac4ee0e +0, 2, 2, 1, 76800, 0x3f87122e +0, 3, 3, 1, 76800, 0xf16001a9 +0, 4, 4, 1, 76800, 0xde7116af -- 2.52.0 _______________________________________________ ffmpeg-devel mailing list -- [email protected] To unsubscribe send an email to [email protected]
