Commit 6ba18c7b7e for ffmpeg
commit 6ba18c7b7ecf02032013f09674d1a73106d47e92
Author: Zhao Zhili <quinkblack@foxmail.com>
Date: Mon Mar 2 00:40:10 2026 +0800
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, with full range high bit depth input normalized by
the code maximum. BT.2020 constant luminance is rejected, as it is not
a plain matrix transform. Results are rounded to the nearest value.
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.
Hardware cropping is honored, by checking the visible frame size
against the output size per frame.
As the software format of CUDA frames is unknown during negotiation,
the input and output color properties are negotiated independently.
RGB output is always negotiated as full range RGB, so that following
filters see the properties of the converted frames, while other output
follows the input. Plain scaling rejects a matrix or range on the
output that differs from the input, as it cannot convert the samples.
Co-authored-by: Timo Rothenpieler <timo@rothenpieler.org>
Co-authored-by: Reinhard Nißl <rnissl@gmx.de>
Assisted-By: Claude Opus 5.5
diff --git a/libavfilter/vf_scale_cuda.c b/libavfilter/vf_scale_cuda.c
index 747cc624b3..dad827de42 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,7 +40,9 @@
#include "libswscale/filters.h"
#include "avfilter.h"
+#include "colorspace.h"
#include "filters.h"
+#include "formats.h"
#include "scale_eval.h"
#include "video.h"
@@ -191,6 +195,10 @@ typedef struct CUDAScaleContext {
int use_filters_opt; /* -1 for auto */
float param;
+
+ bool yuv2rgb;
+ enum AVColorRange color_range; ///< resolved color range for conversion
+ CUDAScaleColorMatrix color_matrix;
} CUDAScaleContext;
/*
@@ -457,6 +465,48 @@ 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)
+{
+ 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];
+
+ /* 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);
+ }
+
+ 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 = chroma_uses_luma_span(colorspace) ? y_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;
@@ -487,8 +537,10 @@ static av_cold int init_processing_chain(AVFilterContext *ctx, int in_width, int
int out_width, int out_height)
{
CUDAScaleContext *s = ctx->priv;
- FilterLink *inl = ff_filter_link(ctx->inputs[0]);
- FilterLink *outl = ff_filter_link(ctx->outputs[0]);
+ AVFilterLink *inlink = ctx->inputs[0];
+ AVFilterLink *outlink = ctx->outputs[0];
+ FilterLink *inl = ff_filter_link(inlink);
+ FilterLink *outl = ff_filter_link(outlink);
AVHWFramesContext *in_frames_ctx;
@@ -517,6 +569,70 @@ 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) {
+ /* Hardware cropping keeps the link dimensions, so the exact size
+ * check happens per frame. Only reject what can never match. */
+ if (out_width > in_width || out_height > in_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;
+
+ if (cs == AVCOL_SPC_RGB) {
+ av_log(ctx, AV_LOG_ERROR,
+ "YUV input is tagged with the RGB colorspace\n");
+ return AVERROR(EINVAL);
+ }
+
+ 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->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->out_desc->flags & AV_PIX_FMT_FLAG_RGB) &&
+ (inlink->colorspace != outlink->colorspace ||
+ inlink->color_range != outlink->color_range)) {
+ /* Without conversion, the samples keep their encoding. */
+ av_log(ctx, AV_LOG_ERROR,
+ "Cannot convert from %s %s range to %s %s range\n",
+ av_color_space_name(inlink->colorspace),
+ av_color_range_name(inlink->color_range),
+ av_color_space_name(outlink->colorspace),
+ av_color_range_name(outlink->color_range));
+ return AVERROR(ENOSYS);
+ }
+
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 +787,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;
@@ -945,6 +1065,40 @@ 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 };
+ const AVPixFmtDescriptor *out_desc = s->format != AV_PIX_FMT_NONE ?
+ av_pix_fmt_desc_get(s->format) : NULL;
+ const int out_rgb = out_desc && (out_desc->flags & AV_PIX_FMT_FLAG_RGB);
+ int ret;
+
+ if ((ret = ff_set_common_formats_from_list2(ctx, cfg_in, cfg_out, pix_fmts)) < 0)
+ return ret;
+
+ /* The software format of CUDA frames is unknown during negotiation, so
+ * whether a YUV->RGB conversion happens is only known at configuration
+ * time. Negotiate the input and output properties independently; when
+ * nothing constrains the output, it follows the input. Configuration
+ * rejects mismatches that the selected processing cannot produce. */
+ if ((ret = ff_formats_ref(ff_all_color_spaces(), &cfg_in[0]->color_spaces)) < 0 ||
+ (ret = ff_formats_ref(ff_all_color_ranges(), &cfg_in[0]->color_ranges)) < 0)
+ return ret;
+
+ if ((ret = ff_formats_ref(out_rgb ? ff_make_formats_list_singleton(AVCOL_SPC_RGB)
+ : ff_all_color_spaces(),
+ &cfg_out[0]->color_spaces)) < 0 ||
+ (ret = ff_formats_ref(out_rgb ? ff_make_formats_list_singleton(AVCOL_RANGE_JPEG)
+ : ff_all_color_ranges(),
+ &cfg_out[0]->color_ranges)) < 0)
+ return ret;
+
+ return 0;
+}
+
static av_cold int cudascale_config_props(AVFilterLink *outlink)
{
AVFilterContext *ctx = outlink->src;
@@ -1146,6 +1300,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 +1309,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 +1329,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 +1454,14 @@ static int cudascale_scale(AVFilterContext *ctx, AVFrame *out, AVFrame *in,
if (ret < 0)
goto fail;
+ if (s->yuv2rgb) {
+ /* The conversion matrix was built from the negotiated link
+ * properties, so describe the output frames the same way. */
+ out->colorspace = outlink->colorspace;
+ out->color_range = outlink->color_range;
+ out->chroma_location = AVCHROMA_LOC_UNSPECIFIED;
+ }
+
out->crop_top = 0;
out->crop_bottom = 0;
out->crop_left = 0;
@@ -1328,6 +1500,17 @@ static int cudascale_filter_frame(AVFilterLink *link, AVFrame *in)
if (ret < 0)
goto fail;
+ /* The conversion kernels read the visible area 1:1 without resizing. */
+ if (s->yuv2rgb &&
+ (in_width != outlink->w || in_height != outlink->h)) {
+ av_log(ctx, AV_LOG_ERROR,
+ "Visible input size %dx%d does not match output size %dx%d; "
+ "resizing is not supported during YUV->RGB conversion\n",
+ in_width, in_height, outlink->w, outlink->h);
+ ret = AVERROR(ENOSYS);
+ goto fail;
+ }
+
out = av_frame_alloc();
if (!out) {
ret = AVERROR(ENOMEM);
@@ -1433,7 +1616,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 7243777651..147ccab3a3 100644
--- a/libavfilter/vf_scale_cuda.cu
+++ b/libavfilter/vf_scale_cuda.cu
@@ -1113,6 +1113,145 @@ __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);
+}
+
+/* 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);
+}
+
+/* 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;
+}
+
+// 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)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
+ 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;
+ 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);
+ 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;
+ 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);
+ 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 +1692,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