diff --git a/doc/filters.texi b/doc/filters.texi index 328339d1e2..7ab0cfaab7 100644 --- a/doc/filters.texi +++ b/doc/filters.texi @@ -27892,7 +27892,11 @@ pixel format is used. Converting between @code{yuv420p}, @code{yuv422p}, @code{yuv444p}, @code{yuv420p10}, @code{yuv422p10}, @code{yuv444p10}, @code{nv12}, @code{nv16}, @code{p010} or @code{p210} and @code{bgra}, @code{rgba}, @code{bgr0} or -@code{rgb0} is supported, but cannot be combined with resizing. +@code{rgb0} is supported, also together with resizing. Resizing is then done +on the RGB side, after converting from YUV or before converting to YUV, and +always uses the generic filters selected with @option{interp_algo}. Resizing +the YUV side in a separate instance instead needs less intermediate memory and +may be faster, but filters chroma differently. YUV to RGB conversion uses the YCbCr matrix and range of the input, RGB to YUV conversion those selected for the output (see @option{out_color_matrix} and @option{out_range}). An unspecified matrix is treated as BT.601 and an @@ -27913,7 +27917,8 @@ If set to 1, filter with a generic weight LUT instead of using fixed-function shader kernels. May be faster or slower depending on the hardware. A value of @code{auto} (the default) enables this automatically when required for correct anti-aliasing when downscaling, and for high-quality Lanczos scaling. -Set this to 0 to use the legacy fixed 4-tap Lanczos kernel. +Set this to 0 to use the legacy fixed 4-tap Lanczos kernel. Resizing during +YUV to RGB or RGB to YUV conversion always uses the generic filters. @item param Algorithm-specific parameter. For @var{bicubic}, this affects the filter @@ -27976,11 +27981,16 @@ scale_cuda=passthrough=0 @item Convert BT.709 input to RGB for a filter that requires it, then back to -BT.709 NV12 with limited range. Note the separate instances, as scaling and -color space conversion cannot be combined. +BT.709 NV12 with limited range. @example scale_cuda=format=bgra,scale_cuda=format=nv12:out_color_matrix=bt709:out_range=limited @end example + +@item +Scale to 720p while converting to RGB. +@example +scale_cuda=-2:720:format=bgra +@end example @end itemize @section thumbnail_cuda diff --git a/libavfilter/vf_scale_cuda.c b/libavfilter/vf_scale_cuda.c index 80d26f650d..55839acfc4 100644 --- a/libavfilter/vf_scale_cuda.c +++ b/libavfilter/vf_scale_cuda.c @@ -141,11 +141,12 @@ typedef struct CUDATex { enum { CUDA_SCALE_BUF_IN, CUDA_SCALE_BUF_TMP, + CUDA_SCALE_BUF_TMP2, CUDA_SCALE_BUF_OUT, CUDA_SCALE_BUF_NB, }; -#define CUDA_SCALE_MAX_PASSES 2 +#define CUDA_SCALE_MAX_PASSES 3 /* * One kernel pass. Everything the kernels need beyond the geometry of the @@ -166,7 +167,9 @@ typedef struct CUDAScalePass { typedef struct CUDAScaleFilterSet { CUDAScalePass passes[CUDA_SCALE_MAX_PASSES]; int nb_passes; - CUDATex tmp; ///< CUDA_SCALE_BUF_TMP + CUDATex tmp[2]; ///< CUDA_SCALE_BUF_TMP, CUDA_SCALE_BUF_TMP2 + int in_use_linear; ///< texture filtering of the input + int in_as_integer; ///< integer texture reads of the input int in_width, in_height; int normalize_crop; } CUDAScaleFilterSet; @@ -208,6 +211,9 @@ typedef struct CUDAScaleContext { CUfunction cu_func_out_uv[2]; CUfunction cu_func_tmp; CUfunction cu_func_tmp_uv; + CUfunction cu_func_tmp_v; ///< vertical float pass before RGB->YUV + CUfunction cu_func_conv; ///< conversion with a float RGB side + CUfunction cu_func_conv_uv; CUstream cu_stream; int interp_algo; @@ -335,7 +341,8 @@ static void cuda_tex_uninit(CudaFunctions *cu, CUDATex *t) static void filter_set_uninit(CudaFunctions *cu, CUDAScaleFilterSet *set) { - cuda_tex_uninit(cu, &set->tmp); + for (int i = 0; i < FF_ARRAY_ELEMS(set->tmp); i++) + cuda_tex_uninit(cu, &set->tmp[i]); for (int i = 0; i < FF_ARRAY_ELEMS(set->passes); i++) { filter_uninit(cu, &set->passes[i].filter); filter_uninit(cu, &set->passes[i].filter_uv); @@ -406,7 +413,10 @@ fail: } static int inter_buf_init(AVFilterContext *ctx, CUDATex *tex, - int out_width, int in_height, + int width, int height, + int log2_chroma_w, int log2_chroma_h, + int nb_planes, const int depths[4], + const int channels[4], int use_float, unsigned int planes) { CUDAScaleContext *s = ctx->priv; @@ -414,25 +424,24 @@ static int inter_buf_init(AVFilterContext *ctx, CUDATex *tex, int ret = 0; *tex = (CUDATex) { - .width = out_width, - .height = in_height, - .crop_width = out_width, - .crop_height = in_height, - .log2_chroma_w = s->out_desc->log2_chroma_w, - .log2_chroma_h = s->in_desc->log2_chroma_h, + .width = width, + .height = height, + .crop_width = width, + .crop_height = height, + .log2_chroma_w = log2_chroma_w, + .log2_chroma_h = log2_chroma_h, }; - for (int i = 0; i < s->in_planes; i++) { + for (int i = 0; i < nb_planes; i++) { const int is_chroma = i == 1 || i == 2; const unsigned int plane = is_chroma ? CUDA_SCALE_PLANE_CHROMA : CUDA_SCALE_PLANE_PRIMARY; const int sub_x = is_chroma ? tex->log2_chroma_w : 0; const int sub_y = is_chroma ? tex->log2_chroma_h : 0; - const int plane_w = AV_CEIL_RSHIFT(out_width, sub_x); - const int plane_h = AV_CEIL_RSHIFT(in_height, sub_y); + const int plane_w = AV_CEIL_RSHIFT(width, sub_x); + const int plane_h = AV_CEIL_RSHIFT(height, sub_y); const int sizeof_pixel = (use_float ? sizeof(float) : - s->in_plane_depths[i] <= 8 ? 1 : 2) * - s->in_plane_channels[i]; + depths[i] <= 8 ? 1 : 2) * channels[i]; if (!(planes & plane)) continue; @@ -453,10 +462,10 @@ static int inter_buf_init(AVFilterContext *ctx, CUDATex *tex, CUDA_RESOURCE_DESC res_desc = { .resType = CU_RESOURCE_TYPE_PITCH2D, .res.pitch2D.format = use_float ? CU_AD_FORMAT_FLOAT : - s->in_plane_depths[i] <= 8 ? + depths[i] <= 8 ? CU_AD_FORMAT_UNSIGNED_INT8 : CU_AD_FORMAT_UNSIGNED_INT16, - .res.pitch2D.numChannels = s->in_plane_channels[i], + .res.pitch2D.numChannels = channels[i], .res.pitch2D.devPtr = tex->data[i], .res.pitch2D.pitchInBytes = pitch, .res.pitch2D.width = plane_w, @@ -639,14 +648,6 @@ static av_cold int init_processing_chain(AVFilterContext *ctx, int in_width, int 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; @@ -671,15 +672,6 @@ static av_cold int init_processing_chain(AVFilterContext *ctx, int in_width, int 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), @@ -689,13 +681,6 @@ static av_cold int init_processing_chain(AVFilterContext *ctx, int in_width, int s->rgb2yuv = (s->in_desc->flags & AV_PIX_FMT_FLAG_RGB) && !(s->out_desc->flags & AV_PIX_FMT_FLAG_RGB); if (s->rgb2yuv) { - if (out_width > in_width || out_height > in_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); - } - /* The matrix and range are those negotiated for the output, which * the output frames are tagged with. When unspecified, use BT.601 * (SMPTE 170M) limited range to match swscale behavior. */ @@ -713,15 +698,6 @@ static av_cold int init_processing_chain(AVFilterContext *ctx, int in_width, int s->color_range = limited_range ? AVCOL_RANGE_MPEG : AVCOL_RANGE_JPEG; - /* 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), @@ -799,7 +775,15 @@ static av_cold int init_processing_chain(AVFilterContext *ctx, int in_width, int s->use_filters_opt != 1) s->interp_algo = INTERP_ALGO_NEAREST; - if (s->interp_algo == INTERP_ALGO_NEAREST) { + if (s->yuv2rgb || s->rgb2yuv) { + /* Conversions plan their own passes, see + * cudascale_conv_set_init(). */ + if (s->use_filters_opt == 0 && + (in_width != out_width || in_height != out_height)) + av_log(ctx, AV_LOG_WARNING, "Scaling during YUV<->RGB " + "conversion always uses the generic filters\n"); + s->use_filters = 0; + } else if (s->interp_algo == INTERP_ALGO_NEAREST) { s->use_filters = 0; } else if (s->use_filters_opt >= 0) { s->use_filters = s->use_filters_opt; @@ -849,7 +833,9 @@ static av_cold int cudascale_load_functions(AVFilterContext *ctx) extern const unsigned char ff_vf_scale_cuda_ptx_data[]; extern const unsigned int ff_vf_scale_cuda_ptx_len; - switch (s->interp_algo) { + /* The direct conversion kernels sample 1:1; scaling during conversion + * uses the generic filters with interp_algo instead. */ + switch (s->yuv2rgb || s->rgb2yuv ? INTERP_ALGO_NEAREST : s->interp_algo) { case INTERP_ALGO_NEAREST: fixed_infix = "Nearest"; fixed_use_linear = 0; @@ -947,6 +933,43 @@ static av_cold int cudascale_load_functions(AVFilterContext *ctx) } } + if (s->yuv2rgb || s->rgb2yuv) { + const char *rgb_name = s->yuv2rgb ? out_fmt_name : in_fmt_name; + + snprintf(buf, sizeof(buf), "Conv_float_%s_%s", in_fmt_name, out_fmt_name); + ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_conv, s->cu_module, buf)); + if (ret < 0) + goto unsupported; + + snprintf(buf, sizeof(buf), "Subsample_Generic_float_h_%s_%s", rgb_name, rgb_name); + ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_tmp, s->cu_module, buf)); + if (ret < 0) + goto unsupported; + + if (s->yuv2rgb) { + static const char *const infix[] = { "Generic_h", "Generic_v" }; + + for (int dir = CUDA_SCALE_DIR_X; dir <= CUDA_SCALE_DIR_Y; dir++) { + snprintf(buf, sizeof(buf), "Subsample_%s_%s_%s", + infix[dir], rgb_name, rgb_name); + ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_out[dir], + s->cu_module, buf)); + if (ret < 0) + goto unsupported; + } + } else { + snprintf(buf, sizeof(buf), "Conv_float_%s_%s_uv", in_fmt_name, out_fmt_name); + ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_conv_uv, s->cu_module, buf)); + if (ret < 0) + goto unsupported; + + snprintf(buf, sizeof(buf), "Subsample_Generic_float_v_%s_%s", rgb_name, rgb_name); + ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_tmp_v, s->cu_module, buf)); + if (ret < 0) + goto unsupported; + } + } + goto fail; unsupported: @@ -980,7 +1003,7 @@ static int cudascale_filter_init(AVFilterContext *ctx, params.virtual_size = 0.0; } else { switch (s->interp_algo) { - case INTERP_ALGO_NEAREST: return 0; /* no weights needed */ + case INTERP_ALGO_NEAREST: params.scaler = SWS_SCALE_POINT; break; case INTERP_ALGO_BILINEAR: params.scaler = SWS_SCALE_BILINEAR; break; case INTERP_ALGO_LANCZOS: params.scaler = SWS_SCALE_LANCZOS; @@ -1046,6 +1069,103 @@ fail: return ret; } +/* + * Converting together with scaling scales on the RGB side, so that every + * scaling pass works on a single packed plane and reuses the generic RGB + * kernels: YUV->RGB converts at the visible input size before scaling, + * RGB->YUV after scaling at the output size. The RGB intermediates are + * normalized floats, so the scaler rounds and clamps only once. + */ +static int cudascale_conv_set_init(AVFilterContext *ctx, + CUDAScaleFilterSet *set, + int in_width, int in_height) +{ + CUDAScaleContext *s = ctx->priv; + AVFilterLink *outlink = ctx->outputs[0]; + const AVPixFmtDescriptor *rgb_desc = s->yuv2rgb ? s->out_desc : s->in_desc; + static const int rgb_depths[4] = { 8 }; + static const int rgb_channels[4] = { 4 }; + const int scale_x = in_width != outlink->w; + const int scale_y = in_height != outlink->h; + int src = CUDA_SCALE_BUF_IN, tmp = CUDA_SCALE_BUF_TMP; + CUDAScalePass *pass; + int ret; + + /* Scaling reads RGB input as normalized floats. */ + set->in_use_linear = 0; + set->in_as_integer = s->yuv2rgb; + + if (s->yuv2rgb) { + ret = inter_buf_init(ctx, &set->tmp[tmp - CUDA_SCALE_BUF_TMP], + in_width, in_height, 0, 0, 1, + rgb_depths, rgb_channels, 1, + CUDA_SCALE_PLANE_PRIMARY); + if (ret < 0) + return ret; + + set->passes[set->nb_passes++] = (CUDAScalePass) { + .func = s->cu_func_conv, + .planes = CUDA_SCALE_PLANE_PRIMARY, + .src = src, + .dst = tmp, + .src_desc = s->in_desc, + .dst_desc = rgb_desc, + }; + src = tmp++; + } + + for (int dir = CUDA_SCALE_DIR_X; dir <= CUDA_SCALE_DIR_Y; dir++) { + const int last = dir == CUDA_SCALE_DIR_Y || !scale_y; + const int to_out = last && s->yuv2rgb; + + if (!(dir == CUDA_SCALE_DIR_X ? scale_x : scale_y)) + continue; + + if (!to_out) { + ret = inter_buf_init(ctx, &set->tmp[tmp - CUDA_SCALE_BUF_TMP], + outlink->w, + dir == CUDA_SCALE_DIR_Y ? outlink->h : in_height, + 0, 0, 1, rgb_depths, rgb_channels, 1, + CUDA_SCALE_PLANE_PRIMARY); + if (ret < 0) + return ret; + } + + pass = &set->passes[set->nb_passes++]; + *pass = (CUDAScalePass) { + .func = to_out ? s->cu_func_out[dir] : + dir == CUDA_SCALE_DIR_X ? s->cu_func_tmp : + s->cu_func_tmp_v, + .planes = CUDA_SCALE_PLANE_PRIMARY, + .src = src, + .dst = to_out ? CUDA_SCALE_BUF_OUT : tmp, + .src_desc = rgb_desc, + .dst_desc = rgb_desc, + }; + ret = dir == CUDA_SCALE_DIR_X ? + cudascale_filter_init(ctx, &pass->filter, in_width, outlink->w, 0.0, 1) : + cudascale_filter_init(ctx, &pass->filter, in_height, outlink->h, 0.0, 1); + if (ret < 0) + return ret; + if (!to_out) + src = tmp++; + } + + if (s->rgb2yuv) { + set->passes[set->nb_passes++] = (CUDAScalePass) { + .func = s->cu_func_conv, + .func_uv = s->cu_func_conv_uv, + .planes = CUDA_SCALE_PLANE_PRIMARY | CUDA_SCALE_PLANE_CHROMA, + .src = src, + .dst = CUDA_SCALE_BUF_OUT, + .src_desc = rgb_desc, + .dst_desc = s->out_desc, + }; + } + + return 0; +} + static int cudascale_filter_set_init(AVFilterContext *ctx, CUDAScaleFilterSet *set, int in_width, int in_height, @@ -1071,9 +1191,16 @@ static int cudascale_filter_set_init(AVFilterContext *ctx, set->in_width = in_width; set->in_height = in_height; set->normalize_crop = normalize_crop; + set->in_use_linear = s->interp_use_linear; + set->in_as_integer = s->interp_as_integer; + + if ((s->yuv2rgb || s->rgb2yuv) && + (in_width != outlink->w || in_height != outlink->h)) + return cudascale_conv_set_init(ctx, set, in_width, in_height); if (!s->use_filters) { - /* The fixed function kernels scale and convert in a single pass. */ + /* The fixed function kernels scale and convert in a single pass, + * as do the conversion kernels without scaling. */ set->passes[set->nb_passes++] = (CUDAScalePass) { .func = s->cu_func_fixed, .func_uv = s->cu_func_fixed_uv, @@ -1175,8 +1302,10 @@ static int cudascale_filter_set_init(AVFilterContext *ctx, } if (pass_x && pass_x->dst == CUDA_SCALE_BUF_TMP) { - ret = inter_buf_init(ctx, &set->tmp, + ret = inter_buf_init(ctx, &set->tmp[0], outlink->w, in_height, + out_sub_x, in_sub_y, s->in_planes, + s->in_plane_depths, s->in_plane_channels, s->interp_algo == INTERP_ALGO_LANCZOS, pass_x->planes); if (ret < 0) @@ -1196,9 +1325,10 @@ static int cudascale_prepare_filter_set(AVFilterContext *ctx, CUDAScaleFilterSet old; int ret; - /* Only the generic filters depend on the visible input size. */ + /* Only the generic filters and conversions depend on the visible + * input size. */ if (s->filter_set.nb_passes && - (!s->use_filters || + ((!s->use_filters && !s->yuv2rgb && !s->rgb2yuv) || (s->filter_set.in_width == in_width && s->filter_set.in_height == in_height && s->filter_set.normalize_crop == normalize_crop))) @@ -1670,7 +1800,8 @@ static int cudascale_scale(AVFilterContext *ctx, AVFrame *out, AVFrame *in, CUDATex in_tex = {0}, out_tex = {0}; ret = cuda_tex_map_frame(ctx, in, s->in_plane_depths, s->in_plane_channels, &in_tex, - s->interp_use_linear, s->interp_as_integer); + s->filter_set.in_use_linear, + s->filter_set.in_as_integer); if (ret < 0) goto fail; @@ -1681,7 +1812,8 @@ static int cudascale_scale(AVFilterContext *ctx, AVFrame *out, AVFrame *in, /* Local views of the buffers; intermediates borrow textures below. */ CUDATex bufs[CUDA_SCALE_BUF_NB] = { [CUDA_SCALE_BUF_IN] = in_tex, - [CUDA_SCALE_BUF_TMP] = s->filter_set.tmp, + [CUDA_SCALE_BUF_TMP] = s->filter_set.tmp[0], + [CUDA_SCALE_BUF_TMP2] = s->filter_set.tmp[1], [CUDA_SCALE_BUF_OUT] = out_tex, }; @@ -1761,17 +1893,6 @@ 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 || s->rgb2yuv) && - (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); diff --git a/libavfilter/vf_scale_cuda.cu b/libavfilter/vf_scale_cuda.cu index 80aecde9ed..14bc355c47 100644 --- a/libavfilter/vf_scale_cuda.cu +++ b/libavfilter/vf_scale_cuda.cu @@ -1183,8 +1183,40 @@ __device__ static inline ushort u8_to_stored(float code8, int depth, int shift, // 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, +__device__ static inline void store_rgb(uchar4 *dst, float r, float g, float b, + int channel_order) +{ + 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: *dst = make_uchar4(bc, gc, rc, 0xFF); break; // BGRA + case 1: *dst = make_uchar4(rc, gc, bc, 0xFF); break; // RGBA + case 2: *dst = make_uchar4(bc, gc, rc, 0x00); break; // BGR0 + case 3: *dst = make_uchar4(rc, gc, bc, 0x00); break; // RGB0 + default: *dst = make_uchar4(bc, gc, rc, 0xFF); break; + } +} + +/* Float intermediate for scaling after conversion: normalized like the + * texture reads of the generic scaling kernels, and left unclamped so that + * the scaler rounds and clamps only once. */ +__device__ static inline void store_rgb(float4 *dst, float r, float g, float b, + int channel_order) +{ + const float n = 1.0f / 255.0f; + const float a = channel_order < 2 ? 1.0f : 0.0f; + + switch (channel_order) { + case 0: case 2: *dst = make_float4(b * n, g * n, r * n, a); break; // BGRA, BGR0 + default: *dst = make_float4(r * n, g * n, b * n, a); break; // RGBA, RGB0 + } +} + +template +__device__ static inline void yuv8_to_rgba_generic( + out_T *dst, float fy, float fu, float fv, const CUDAScaleColorMatrix &mat, int channel_order) { @@ -1192,22 +1224,12 @@ __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)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); - } + store_rgb(dst, r, g, b, channel_order); } // YUV to RGB kernel for semiplanar input (NV12, P010, NV16, P210) // No resize: src and dst have identical dimensions. -template +template __device__ void YuvToRgb_semiplanar(CUDAScaleKernelParams params) { int xo = blockIdx.x * blockDim.x + threadIdx.x; @@ -1232,14 +1254,15 @@ __device__ void YuvToRgb_semiplanar(CUDAScaleKernelParams params) 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); + out_T *dst = (out_T*)params.dst[0]; + int dst_pitch = params.dst_pitch / sizeof(out_T); + yuv8_to_rgba_generic(&dst[yo * dst_pitch + xo], 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 +template __device__ void YuvToRgb_planar(CUDAScaleKernelParams params) { int xo = blockIdx.x * blockDim.x + threadIdx.x; @@ -1266,9 +1289,10 @@ __device__ void YuvToRgb_planar(CUDAScaleKernelParams params) 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); + out_T *dst = (out_T*)params.dst[0]; + int dst_pitch = params.dst_pitch / sizeof(out_T); + yuv8_to_rgba_generic(&dst[yo * dst_pitch + xo], fy, fu, fv, + params.color_matrix, channel_order); } // RGB to YUV @@ -1277,35 +1301,51 @@ __device__ void YuvToRgb_planar(CUDAScaleKernelParams params) // 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. +// Extract R,G,B from a packed 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, +__device__ static inline void extract_rgb(float x, float y, float z, + 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; + b = x; g = y; r = z; break; case 1: // RGBA case 3: // RGB0 default: - r = (float)pixel.x; g = (float)pixel.y; b = (float)pixel.z; + r = x; g = y; b = z; break; } } +__device__ static inline void extract_rgb(uchar4 pixel, int channel_order, + float &r, float &g, float &b) +{ + extract_rgb((float)pixel.x, (float)pixel.y, (float)pixel.z, + channel_order, r, g, b); +} + +/* Normalized float intermediate produced by scaling before conversion. */ +__device__ static inline void extract_rgb(float4 pixel, int channel_order, + float &r, float &g, float &b) +{ + extract_rgb(pixel.x * 255.0f, pixel.y * 255.0f, pixel.z * 255.0f, + channel_order, r, g, b); +} + __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); + return y; } -/* Chroma is left unclamped so that subsampling averages the exact values; - * the stores clamp at destination precision. */ +/* Like luma, chroma is left unclamped, so that subsampling averages the + * exact values; the stores clamp at destination precision. */ __device__ static inline float rgb_to_u(float r, float g, float b, const CUDAScaleColorMatrix &mat) { @@ -1319,14 +1359,14 @@ __device__ static inline float rgb_to_v(float r, float g, float b, } // RGB->YUV Y kernel (8-bit output) -template +template __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( + in_T pixel = tex2D( params.src_tex[0], xo + params.src_left + 0.5f, yo + params.src_top + 0.5f); @@ -1341,14 +1381,14 @@ __device__ void RgbToYuv_Y(CUDAScaleKernelParams params) } // RGB->YUV Y kernel for semiplanar10/16 output -template +template __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( + in_T pixel = tex2D( params.src_tex[0], xo + params.src_left + 0.5f, yo + params.src_top + 0.5f); @@ -1363,31 +1403,16 @@ __device__ void RgbToYuv_Y_semiplanar16(CUDAScaleKernelParams params) params.mpeg_range, 0); } -// RGB->YUV Y kernel for planar10 output: native 10-bit Y value -template +// RGB->YUV Y kernel for planar10 output: the luma plane is stored like the +// semiplanar one, with dst_depth and dst_shift describing the code. +template __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( - 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] = u8_to_stored(y, params.dst_depth, params.dst_shift, - params.mpeg_range, 0); + RgbToYuv_Y_semiplanar16(params); } // RGB->YUV UV kernel for semiplanar output (NV12, NV16) -template +template __device__ void RgbToYuv_semiplanar_UV(CUDAScaleKernelParams params) { int xo = blockIdx.x * blockDim.x + threadIdx.x; @@ -1405,7 +1430,7 @@ __device__ void RgbToYuv_semiplanar_UV(CUDAScaleKernelParams params) int ly = yo * chroma_h + dy; if (lx >= params.src_width || ly >= params.src_height) continue; - uchar4 pixel = tex2D( + in_T pixel = tex2D( params.src_tex[0], lx + params.src_left + 0.5f, ly + params.src_top + 0.5f); @@ -1428,7 +1453,7 @@ __device__ void RgbToYuv_semiplanar_UV(CUDAScaleKernelParams params) } // RGB->YUV UV kernel for semiplanar10/16 output -template +template __device__ void RgbToYuv_semiplanar16_UV(CUDAScaleKernelParams params) { int xo = blockIdx.x * blockDim.x + threadIdx.x; @@ -1446,7 +1471,7 @@ __device__ void RgbToYuv_semiplanar16_UV(CUDAScaleKernelParams params) int ly = yo * chroma_h + dy; if (lx >= params.src_width || ly >= params.src_height) continue; - uchar4 pixel = tex2D( + in_T pixel = tex2D( params.src_tex[0], lx + params.src_left + 0.5f, ly + params.src_top + 0.5f); @@ -1472,7 +1497,7 @@ __device__ void RgbToYuv_semiplanar16_UV(CUDAScaleKernelParams params) } // RGB->YUV UV kernel for planar output (YUV420P, YUV422P, YUV444P, YUV420P10) -template +template __device__ void RgbToYuv_planar_UV(CUDAScaleKernelParams params) { int xo = blockIdx.x * blockDim.x + threadIdx.x; @@ -1490,7 +1515,7 @@ __device__ void RgbToYuv_planar_UV(CUDAScaleKernelParams params) int ly = yo * chroma_h + dy; if (lx >= params.src_width || ly >= params.src_height) continue; - uchar4 pixel = tex2D( + in_T pixel = tex2D( params.src_tex[0], lx + params.src_left + 0.5f, ly + params.src_top + 0.5f); @@ -1640,7 +1665,7 @@ __device__ static inline T Subsample_Generic(cudaTextureObject_t tex, ); } -template +template __device__ static inline floatT Subsample_GenericFloat(cudaTextureObject_t tex, int xo, int yo, int src_left, int src_top, @@ -1648,14 +1673,22 @@ __device__ static inline floatT Subsample_GenericFloat(cudaTextureObject_t tex, const int *offsets, int filter_size) { - const float *row = &weights[xo * filter_size]; - const float x = 0.5f + src_left + offsets[xo]; - const float y = 0.5f + src_top + yo; floatT sum; vec_set_scalar(sum, 0.0f); - for (int i = 0; i < filter_size; i++) - sum += tex2D(tex, x + i, y) * row[i]; + if (dir == SCALE_DIR_X) { + const float *row = &weights[xo * filter_size]; + const float x = 0.5f + src_left + offsets[xo]; + const float y = 0.5f + src_top + yo; + for (int i = 0; i < filter_size; i++) + sum += tex2D(tex, x + i, y) * row[i]; + } else { + const float *col = &weights[yo * filter_size]; + const float x = 0.5f + src_left + xo; + const float y = 0.5f + src_top + offsets[yo]; + for (int i = 0; i < filter_size; i++) + sum += tex2D(tex, x, y + i) * col[i]; + } return sum; } @@ -1961,6 +1994,27 @@ GENERIC_FLOAT_KERNEL_Y(bgr0, uchar4, float4) GENERIC_FLOAT_KERNEL_Y(rgba, uchar4, float4) GENERIC_FLOAT_KERNEL_Y(bgra, uchar4, float4) +/* Vertical float pass, used to scale RGB before converting it to YUV. */ +#define GENERIC_FLOAT_KERNEL_V(C, T, F) \ + __global__ void Subsample_Generic_float_v_##C##_##C( \ + 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; \ + F *dst = (F *)((char *)params.dst[0] + \ + yo * params.dst_pitch); \ + dst[xo] = Subsample_GenericFloat( \ + params.src_tex[0], xo, yo, params.src_left, params.src_top, \ + (const float *)params.weights, (const int *)params.offsets, \ + params.filter_size); \ + } + +GENERIC_FLOAT_KERNEL_V(rgb0, uchar4, float4) +GENERIC_FLOAT_KERNEL_V(bgr0, uchar4, float4) +GENERIC_FLOAT_KERNEL_V(rgba, uchar4, float4) +GENERIC_FLOAT_KERNEL_V(bgra, uchar4, float4) + // --- YUV to RGB kernel exports --- // Single-pass kernels (no resize, color conversion only) @@ -2104,4 +2158,46 @@ __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(p); } + +// --- Conversion kernels for a float RGB intermediate --- +// Converting together with scaling scales on the RGB side: YUV->RGB writes +// the intermediate that is then scaled, RGB->YUV reads the scaled one. + +#define CONV_FLOAT_TO_RGB(RGB, O) \ +__global__ void Conv_float_semiplanar8_##RGB(CUDAScaleKernelParams p) \ +{ YuvToRgb_semiplanar(p); } \ +__global__ void Conv_float_semiplanar10_##RGB(CUDAScaleKernelParams p) \ +{ YuvToRgb_semiplanar(p); } \ +__global__ void Conv_float_planar8_##RGB(CUDAScaleKernelParams p) \ +{ YuvToRgb_planar(p); } \ +__global__ void Conv_float_planar10_##RGB(CUDAScaleKernelParams p) \ +{ YuvToRgb_planar(p); } + +#define CONV_FLOAT_FROM_RGB(RGB, O) \ +__global__ void Conv_float_##RGB##_semiplanar8(CUDAScaleKernelParams p) \ +{ RgbToYuv_Y(p); } \ +__global__ void Conv_float_##RGB##_semiplanar8_uv(CUDAScaleKernelParams p) \ +{ RgbToYuv_semiplanar_UV(p); } \ +__global__ void Conv_float_##RGB##_semiplanar10(CUDAScaleKernelParams p) \ +{ RgbToYuv_Y_semiplanar16(p); } \ +__global__ void Conv_float_##RGB##_semiplanar10_uv(CUDAScaleKernelParams p) \ +{ RgbToYuv_semiplanar16_UV(p); } \ +__global__ void Conv_float_##RGB##_planar8(CUDAScaleKernelParams p) \ +{ RgbToYuv_Y(p); } \ +__global__ void Conv_float_##RGB##_planar8_uv(CUDAScaleKernelParams p) \ +{ RgbToYuv_planar_UV(p); } \ +__global__ void Conv_float_##RGB##_planar10(CUDAScaleKernelParams p) \ +{ RgbToYuv_Y_planar10(p); } \ +__global__ void Conv_float_##RGB##_planar10_uv(CUDAScaleKernelParams p) \ +{ RgbToYuv_planar_UV(p); } + +CONV_FLOAT_TO_RGB(bgra, 0) +CONV_FLOAT_TO_RGB(rgba, 1) +CONV_FLOAT_TO_RGB(bgr0, 2) +CONV_FLOAT_TO_RGB(rgb0, 3) + +CONV_FLOAT_FROM_RGB(bgra, 0) +CONV_FLOAT_FROM_RGB(rgba, 1) +CONV_FLOAT_FROM_RGB(bgr0, 2) +CONV_FLOAT_FROM_RGB(rgb0, 3) } diff --git a/tests/fate/filter-video.mak b/tests/fate/filter-video.mak index dcb3c0da92..4a065ec032 100644 --- a/tests/fate/filter-video.mak +++ b/tests/fate/filter-video.mak @@ -1032,15 +1032,17 @@ FATE_SCALE_CUDA_SPECIAL := $(addprefix fate-filter-scale_cuda-, nv12-bgra-bt709 crop-nv12-bgra crop-bgra-nv12 nv12-bgra-full-gbr roundtrip-nv12-bt709) FATE_SCALE_CUDA_ODD_RGB := $(addprefix fate-filter-scale_cuda-, odd-bgra-nv12 odd-bgra-yuv420p) FATE_SCALE_CUDA_ODD_YUV := $(addprefix fate-filter-scale_cuda-, odd-nv12-bgra odd-yuv420p-bgra) +FATE_SCALE_CUDA_RESIZE := $(addprefix fate-filter-scale_cuda-resize-, nv12-bgra-up p010-bgra-down yuv420p-rgba-w \ + bgra-nv12-down bgra-p010-up bgra-yuv444p-h crop-nv12-bgra) FATE_SCALE_CUDA := $(FATE_SCALE_CUDA_YUV_RGB) $(FATE_SCALE_CUDA_RGB_YUV) $(FATE_SCALE_CUDA_SPECIAL) \ - $(FATE_SCALE_CUDA_ODD_RGB) $(FATE_SCALE_CUDA_ODD_YUV) + $(FATE_SCALE_CUDA_ODD_RGB) $(FATE_SCALE_CUDA_ODD_YUV) $(FATE_SCALE_CUDA_RESIZE) FATE_SCALE_CUDA-$(call FILTERFRAMECRC, $(SCALE_CUDA_FILTERS), FFMPEG) := $(FATE_SCALE_CUDA) $(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_YUV_RGB) $(FATE_SCALE_CUDA_RGB_YUV) $(FATE_SCALE_CUDA_SPECIAL) $(FATE_SCALE_CUDA_RESIZE): SIZE = 160x120 $(FATE_SCALE_CUDA_ODD_RGB): SIZE = 161x121 # testsrc2 rounds subsampled YUV sizes down, so crop exactly to 161x121 instead. $(FATE_SCALE_CUDA_ODD_YUV): SIZE = 162x122 @@ -1077,6 +1079,20 @@ fate-filter-scale_cuda-roundtrip-nv12-bt709: LAVFI = format=nv12,setparams=color fate-filter-scale_cuda-crop-nv12-bgra: LAVFI = format=nv12,hwupload_cuda,crop=96:64:16:8,scale_cuda=96:64:format=bgra,hwdownload,format=bgra fate-filter-scale_cuda-crop-bgra-nv12: LAVFI = format=bgra,hwupload_cuda,crop=96:64:16:8,scale_cuda=96:64:format=nv12:out_range=limited,hwdownload,format=nv12 +# Resizing during conversion, in one or both directions. +fate-filter-scale_cuda-resize-nv12-bgra-up: LAVFI = format=nv12,hwupload_cuda,scale_cuda=240:180:format=bgra,hwdownload,format=bgra +fate-filter-scale_cuda-resize-p010-bgra-down: LAVFI = format=p010,hwupload_cuda,scale_cuda=80:60:format=bgra:interp_algo=lanczos,hwdownload,format=bgra +fate-filter-scale_cuda-resize-yuv420p-rgba-w: LAVFI = format=yuv420p,hwupload_cuda,scale_cuda=240:120:format=rgba:interp_algo=bilinear,hwdownload,format=rgba +fate-filter-scale_cuda-resize-bgra-nv12-down: LAVFI = format=bgra,hwupload_cuda,scale_cuda=80:60:format=nv12:out_range=limited,hwdownload,format=nv12 +fate-filter-scale_cuda-resize-bgra-p010-up: LAVFI = format=bgra,hwupload_cuda,scale_cuda=240:180:format=p010:out_range=limited,hwdownload,format=p010 +fate-filter-scale_cuda-resize-bgra-yuv444p-h: LAVFI = format=bgra,hwupload_cuda,scale_cuda=160:60:format=yuv444p:out_range=limited:interp_algo=nearest,hwdownload,format=yuv444p +fate-filter-scale_cuda-resize-crop-nv12-bgra: LAVFI = format=nv12,hwupload_cuda,crop=96:64:16:8,scale_cuda=192:128:format=bgra,hwdownload,format=bgra + +# A sharp step resized with Lanczos overshoots; limited range high bit depth luma keeps it up to the code maximum. +FATE_SCALE_CUDA_STEP-$(call FILTERFRAMECRC, $(SCALE_CUDA_FILTERS) COLOR HSTACK, FFMPEG) += fate-filter-scale_cuda-resize-step-p010 +fate-filter-scale_cuda-resize-step-p010: CMD = framecrc -lavfi "color=black:s=3x4,format=bgra[a];color=white:s=5x4,format=bgra[b];[a][b]hstack,hwupload_cuda,scale_cuda=32:4:format=p010:interp_algo=lanczos:out_range=limited,hwdownload,format=p010" -frames:v 1 +FATE_SCALE_CUDA-yes += $(FATE_SCALE_CUDA_STEP-yes) + 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 = format=nv12,crop=161:121:0:0:exact=1,hwupload_cuda,scale_cuda=format=bgra,hwdownload,format=bgra diff --git a/tests/ref/fate/filter-scale_cuda-resize-bgra-nv12-down b/tests/ref/fate/filter-scale_cuda-resize-bgra-nv12-down new file mode 100644 index 0000000000..c161296aaf --- /dev/null +++ b/tests/ref/fate/filter-scale_cuda-resize-bgra-nv12-down @@ -0,0 +1,10 @@ +#tb 0: 1/25 +#media_type 0: video +#codec_id 0: rawvideo +#dimensions 0: 80x60 +#sar 0: 1/1 +0, 0, 0, 1, 7200, 0x6a5afed1 +0, 1, 1, 1, 7200, 0x0014fd36 +0, 2, 2, 1, 7200, 0xd310ff42 +0, 3, 3, 1, 7200, 0xaaa2fc0d +0, 4, 4, 1, 7200, 0x6ab1fc8a diff --git a/tests/ref/fate/filter-scale_cuda-resize-bgra-p010-up b/tests/ref/fate/filter-scale_cuda-resize-bgra-p010-up new file mode 100644 index 0000000000..4349c323a6 --- /dev/null +++ b/tests/ref/fate/filter-scale_cuda-resize-bgra-p010-up @@ -0,0 +1,10 @@ +#tb 0: 1/25 +#media_type 0: video +#codec_id 0: rawvideo +#dimensions 0: 240x180 +#sar 0: 1/1 +0, 0, 0, 1, 129600, 0x0358b3d2 +0, 1, 1, 1, 129600, 0x92658173 +0, 2, 2, 1, 129600, 0x9b231dc4 +0, 3, 3, 1, 129600, 0x39b7a620 +0, 4, 4, 1, 129600, 0x756c2ef3 diff --git a/tests/ref/fate/filter-scale_cuda-resize-bgra-yuv444p-h b/tests/ref/fate/filter-scale_cuda-resize-bgra-yuv444p-h new file mode 100644 index 0000000000..5517a6d35c --- /dev/null +++ b/tests/ref/fate/filter-scale_cuda-resize-bgra-yuv444p-h @@ -0,0 +1,10 @@ +#tb 0: 1/25 +#media_type 0: video +#codec_id 0: rawvideo +#dimensions 0: 160x60 +#sar 0: 1/2 +0, 0, 0, 1, 28800, 0x28ee7eac +0, 1, 1, 1, 28800, 0x594582d8 +0, 2, 2, 1, 28800, 0x20f08796 +0, 3, 3, 1, 28800, 0xfa668f1f +0, 4, 4, 1, 28800, 0x60119873 diff --git a/tests/ref/fate/filter-scale_cuda-resize-crop-nv12-bgra b/tests/ref/fate/filter-scale_cuda-resize-crop-nv12-bgra new file mode 100644 index 0000000000..ac811555b1 --- /dev/null +++ b/tests/ref/fate/filter-scale_cuda-resize-crop-nv12-bgra @@ -0,0 +1,10 @@ +#tb 0: 1/25 +#media_type 0: video +#codec_id 0: rawvideo +#dimensions 0: 192x128 +#sar 0: 1/1 +0, 0, 0, 1, 98304, 0xd4a1e982 +0, 1, 1, 1, 98304, 0x53828346 +0, 2, 2, 1, 98304, 0xbb0ba085 +0, 3, 3, 1, 98304, 0xa7da1006 +0, 4, 4, 1, 98304, 0x0b600560 diff --git a/tests/ref/fate/filter-scale_cuda-resize-nv12-bgra-up b/tests/ref/fate/filter-scale_cuda-resize-nv12-bgra-up new file mode 100644 index 0000000000..d73d9dc517 --- /dev/null +++ b/tests/ref/fate/filter-scale_cuda-resize-nv12-bgra-up @@ -0,0 +1,10 @@ +#tb 0: 1/25 +#media_type 0: video +#codec_id 0: rawvideo +#dimensions 0: 240x180 +#sar 0: 1/1 +0, 0, 0, 1, 172800, 0xd2396d96 +0, 1, 1, 1, 172800, 0x2cbb6993 +0, 2, 2, 1, 172800, 0x7be2bd82 +0, 3, 3, 1, 172800, 0x50ea9b74 +0, 4, 4, 1, 172800, 0x1f9dbf30 diff --git a/tests/ref/fate/filter-scale_cuda-resize-p010-bgra-down b/tests/ref/fate/filter-scale_cuda-resize-p010-bgra-down new file mode 100644 index 0000000000..797c95756e --- /dev/null +++ b/tests/ref/fate/filter-scale_cuda-resize-p010-bgra-down @@ -0,0 +1,10 @@ +#tb 0: 1/25 +#media_type 0: video +#codec_id 0: rawvideo +#dimensions 0: 80x60 +#sar 0: 1/1 +0, 0, 0, 1, 19200, 0xd593b268 +0, 1, 1, 1, 19200, 0x5348b21c +0, 2, 2, 1, 19200, 0x32e5bb51 +0, 3, 3, 1, 19200, 0x8e77b7dd +0, 4, 4, 1, 19200, 0xffbdbc12 diff --git a/tests/ref/fate/filter-scale_cuda-resize-step-p010 b/tests/ref/fate/filter-scale_cuda-resize-step-p010 new file mode 100644 index 0000000000..c37767d63a --- /dev/null +++ b/tests/ref/fate/filter-scale_cuda-resize-step-p010 @@ -0,0 +1,6 @@ +#tb 0: 1/25 +#media_type 0: video +#codec_id 0: rawvideo +#dimensions 0: 32x4 +#sar 0: 1/4 +0, 0, 0, 1, 384, 0xcb318a70 diff --git a/tests/ref/fate/filter-scale_cuda-resize-yuv420p-rgba-w b/tests/ref/fate/filter-scale_cuda-resize-yuv420p-rgba-w new file mode 100644 index 0000000000..1b0b351cf8 --- /dev/null +++ b/tests/ref/fate/filter-scale_cuda-resize-yuv420p-rgba-w @@ -0,0 +1,10 @@ +#tb 0: 1/25 +#media_type 0: video +#codec_id 0: rawvideo +#dimensions 0: 240x120 +#sar 0: 2/3 +0, 0, 0, 1, 115200, 0x5a4aa6ee +0, 1, 1, 1, 115200, 0xac5fa2fe +0, 2, 2, 1, 115200, 0xd354dcd3 +0, 3, 3, 1, 115200, 0x90c6c54f +0, 4, 4, 1, 115200, 0x190ddce4