From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: from ffbox0-bg.ffmpeg.org (ffbox0-bg.ffmpeg.org [79.124.17.100]) by master.gitmailbox.com (Postfix) with ESMTPS id 06E814F445 for ; Mon, 16 Jun 2025 23:40:01 +0000 (UTC) Received: from [127.0.1.1] (localhost [127.0.0.1]) by ffbox0-bg.ffmpeg.org (Postfix) with ESMTP id 5EF0368CAD6; Tue, 17 Jun 2025 02:39:58 +0300 (EEST) Received: from mail-pl1-f174.google.com (mail-pl1-f174.google.com [209.85.214.174]) by ffbox0-bg.ffmpeg.org (Postfix) with ESMTPS id 82BC168CA76 for ; Tue, 17 Jun 2025 02:39:52 +0300 (EEST) Received: by mail-pl1-f174.google.com with SMTP id d9443c01a7336-23602481460so47362195ad.0 for ; Mon, 16 Jun 2025 16:39:52 -0700 (PDT) DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=gmail.com; s=20230601; t=1750117190; x=1750721990; darn=ffmpeg.org; h=content-transfer-encoding:mime-version:message-id:date:subject:to :from:from:to:cc:subject:date:message-id:reply-to; bh=ymktlyhnrFyR66vnWvx/X0iGEucUXRJBAuBf1Kb6gHs=; b=SNjz97RuNPveC57/oy7nEjCBtL7ST57DHuc+TBke4NnbeIkmbshLK3vQIe9DT4oZx6 n3Nax+R1KaxvmZm+QkeYw/K1pH8CFuXjWZsMlfa/KIHnK0qzDCmM8JluED6c/gBXOKSs oxAVlFcyQaPytEwg9weO/+jYdkxZwgPRmfGKqtP0jPepGMZX2M0hac1gAy1Mu/lbLHVo RjH+3N5/VXRQnqEejbWg96yBp9rbU9bhsXVwlzVyWmfR2fXhMm9C/K+ht/itfJnimbjS qOr9BFzTsiSi1EBAVQXn/sHM8dj/oe08KhNkr7XQzCeMA6RsILJbS7gm8SuPvO6qf0+x ekXA== X-Google-DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=1e100.net; s=20230601; t=1750117190; x=1750721990; h=content-transfer-encoding:mime-version:message-id:date:subject:to :from:x-gm-message-state:from:to:cc:subject:date:message-id:reply-to; bh=ymktlyhnrFyR66vnWvx/X0iGEucUXRJBAuBf1Kb6gHs=; b=V+KLu5lpc8PmTWjxf6DmJfv91LKzRFbOxQswqa6ObGOlov0diENKszeINr5627G7/O RZfg+/yn8AvWIqMnddpBWmyXHMgdP4+xn2/pbpArf4CNByenx1B3bxAAPTRA462au70w GeVZXHzn8Fobi5xwzj355Ue6IolZzL+EepujmOZHim7UqrevxAdh9mh+WsQBBMXr0lBI rOZYzZtvtnm9ruPTwIJE4hm2n0bisFMTYiSE/DCO2OISQUI51mzQr/NxtiCbcyD28cB/ OR2vNYh5J9ZujiRVfEFB8wP0+48FSRt84ynKyEPamFFXl/sT8iYIHpiEGMOcaTHZETm3 LkXw== X-Gm-Message-State: AOJu0Yy1vZnn7ATPoDngVQg1al6jEVNFXLFZ//JwF293eU6QufD/+Q3G kdtV4FXIWsDqyUazwV3xCbfAtwhHgP2/6BgGDm6BJ4ezpdqZIMOwXcVptSOpgg== X-Gm-Gg: ASbGnctsj98q0YDOTe5NUzIlOnvKB/bf/4Zkam73suN94X7CzpodVsj/HTPoGAy41Sg w9KanJKLQbPSREfzKZUnCIpzwrJhZ2psDW8kvrNLSn7vx9iaN8UDDmwcCjXTvUBjgCrarTWyVwr Ytn1N6u7xzkVAdrhjxz5kDhRnW6403H1L/TVD0zvipa1WIKAoA/F7D7V6xbfNHtuI7VBPqVoh3J CbiD0GysNZEIsqItFrALzB0jdfNBNzO1fMn/cB3LnS4qZgWXe9/H6vzP8hCrHZQiXnyw/sdvIvy LtTT4J1D4KJsljOWJiy66cr6U3dC8Du7+n80TU5pfG6lfhXOOhVzxwHhUo66EjTbsiE55J5wdX6 km9I7 X-Google-Smtp-Source: AGHT+IF6JXDuo+SbneDEVn2BY5fmqkiFCheiom2FB47tAMTuw4zA8vT3ZEDDBw3qTMANE0HRAG9AwA== X-Received: by 2002:a17:90b:1dcd:b0:312:f263:954a with SMTP id 98e67ed59e1d1-313f1c7d8c9mr16034882a91.5.1750117189517; Mon, 16 Jun 2025 16:39:49 -0700 (PDT) Received: from JESTRADA-PC.localdomain ([98.97.30.252]) by smtp.gmail.com with ESMTPSA id 98e67ed59e1d1-313c19d1eafsm9340905a91.18.2025.06.16.16.39.48 for (version=TLS1_3 cipher=TLS_AES_256_GCM_SHA384 bits=256/256); Mon, 16 Jun 2025 16:39:49 -0700 (PDT) From: Jorge Estrada To: ffmpeg-devel@ffmpeg.org Date: Mon, 16 Jun 2025 16:39:22 -0700 Message-Id: <20250616233931.192291-1-jestrada.list@gmail.com> X-Mailer: git-send-email 2.34.1 MIME-Version: 1.0 Subject: [FFmpeg-devel] [PATCH] avfilter: add pad_cuda filter X-BeenThere: ffmpeg-devel@ffmpeg.org X-Mailman-Version: 2.1.29 Precedence: list List-Id: FFmpeg development discussions and patches List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , Reply-To: FFmpeg development discussions and patches Content-Type: text/plain; charset="us-ascii" Content-Transfer-Encoding: 7bit Errors-To: ffmpeg-devel-bounces@ffmpeg.org Sender: "ffmpeg-devel" Archived-At: List-Archive: List-Post: From: Jorge Estrada <--global> This patch adds the pad_cuda video filter. A filter similar to the existing pad filter but accelerated by CUDA. The filter shares the same options as the software pad filter. Example usage: ffmpeg -hwaccel cuda -hwaccel_output_format cuda -i input.mp4 -vf "pad_cuda=w=iw+100:h=ih+100:x=-1:y=-1:color=red" out.mp4 --- Changelog | 1 + configure | 2 + doc/filters.texi | 82 +++++ libavfilter/Makefile | 1 + libavfilter/allfilters.c | 1 + libavfilter/version.h | 2 +- libavfilter/vf_pad_cuda.c | 650 +++++++++++++++++++++++++++++++++++++ libavfilter/vf_pad_cuda.cu | 65 ++++ 8 files changed, 803 insertions(+), 1 deletion(-) create mode 100644 libavfilter/vf_pad_cuda.c create mode 100644 libavfilter/vf_pad_cuda.cu diff --git a/Changelog b/Changelog index 4217449438..ad7906acf9 100644 --- a/Changelog +++ b/Changelog @@ -18,6 +18,7 @@ version : - APV encoding support through a libopenapv wrapper - VVC decoder supports all content of SCC (Screen Content Coding): IBC (Inter Block Copy), Palette Mode and ACT (Adaptive Color Transform +- pad_cuda filter version 7.1: diff --git a/configure b/configure index 534b443f7d..e68d612867 100755 --- a/configure +++ b/configure @@ -3357,6 +3357,8 @@ thumbnail_cuda_filter_deps_any="cuda_nvcc cuda_llvm" transpose_npp_filter_deps="ffnvcodec libnpp" overlay_cuda_filter_deps="ffnvcodec" overlay_cuda_filter_deps_any="cuda_nvcc cuda_llvm" +pad_cuda_filter_deps="ffnvcodec" +pad_cuda_filter_deps_any="cuda_nvcc cuda_llvm" sharpen_npp_filter_deps="ffnvcodec libnpp" ddagrab_filter_deps="d3d11va IDXGIOutput1 DXGI_OUTDUPL_FRAME_INFO" diff --git a/doc/filters.texi b/doc/filters.texi index f32fc23c70..68be03c1ab 100644 --- a/doc/filters.texi +++ b/doc/filters.texi @@ -26682,6 +26682,88 @@ See @ref{framesync}. This filter also supports the @ref{framesync} options. + +@anchor{pad_cuda} +@subsection pad_cuda +Add paddings to an input video stream using CUDA. + +This filter is the CUDA-accelerated version of the @ref{pad} filter. It accepts CUDA frames as input and the following options: + +@table @option +@item width, w +@item height, h +Specify an expression for the size of the output image with the paddings. If the value for @var{width} or @var{height} is 0, the corresponding input size is used for the output. + +The default value for both @var{width} and @var{height} is 0. + +@item x +@item y +Specify the offsets to place the input image at within the padded area, with respect to the top/left border of the output image. + +If @var{x} or @var{y} evaluates to a negative number, the input image will be centered on the padded area. + +The default value for both @var{x} and @var{y} is 0. + +@item color +Specify the color of the padded area. For the syntax of this option, check the @ref{color syntax,,"Color" section in the ffmpeg-utils manual,ffmpeg-utils}. + +The default value of @var{color} is "black". + +@item eval +Specify when to evaluate the @var{width}, @var{height}, @var{x}, and @var{y} expressions. + +It accepts the following values: +@table @samp +@item init +Only evaluate expressions once during the filter initialization. (default) +@item frame +Evaluate expressions for each incoming frame. +@end table + +@item aspect +Pad to a specified aspect ratio instead of a fixed resolution. +@end table + +The value for the @var{width}, @var{height}, @var{x}, and @var{y} options are expressions containing the following constants: +@table @var +@item in_w, iw +The input video width. +@item in_h, ih +The input video height. +@item out_w, ow +The output width (the size of the padded area), as specified by the @var{width} expressions. +@item out_h, oh +The output height (the size of the padded area), as specified by the @var{height} expressions. +@item x +@item y +The x and y offsets as specified by the @var{x} and @var{y} expressions. +@item a +Same as @code{iw / ih}. +@item sar +The input sample aspect ratio. +@item dar +The input display aspect ratio, equivalent to @code{(iw / ih) * sar}. +@item hsub +@item vsub +The horizontal and vertical chroma subsample values. +@end table + +@subsubsection Examples + +@itemize +@item +Add a 200-pixel black border to all sides of a video frame: +@example +ffmpeg -hwaccel cuda -hwaccel_output_format cuda -i input.mp4 -vf "pad_cuda=w=iw+400:h=ih+400:x=200:y=200" -c:v h264_nvenc out.mp4 +@end example + +@item +Pad the input video to a 16:9 aspect ratio, filling with the color "blue": +@example +ffmpeg -hwaccel cuda -hwaccel_output_format cuda -i input.mp4 -vf "pad_cuda=w=ih*16/9/sar:h=ih:x=(ow-iw)/2:y=(oh-ih)/2:color=blue" -c:v h264_nvenc out.mp4 +@end example +@end itemize + @anchor{scale_cuda} @subsection scale_cuda diff --git a/libavfilter/Makefile b/libavfilter/Makefile index 97f8f17272..9e9153f5b0 100644 --- a/libavfilter/Makefile +++ b/libavfilter/Makefile @@ -422,6 +422,7 @@ OBJS-$(CONFIG_OVERLAY_VAAPI_FILTER) += vf_overlay_vaapi.o framesync.o v OBJS-$(CONFIG_OVERLAY_VULKAN_FILTER) += vf_overlay_vulkan.o vulkan.o vulkan_filter.o OBJS-$(CONFIG_OWDENOISE_FILTER) += vf_owdenoise.o OBJS-$(CONFIG_PAD_FILTER) += vf_pad.o +OBJS-$(CONFIG_PAD_CUDA_FILTER) += vf_pad_cuda.o vf_pad_cuda.ptx.o cuda/load_helper.o OBJS-$(CONFIG_PAD_OPENCL_FILTER) += vf_pad_opencl.o opencl.o opencl/pad.o OBJS-$(CONFIG_PALETTEGEN_FILTER) += vf_palettegen.o palette.o OBJS-$(CONFIG_PALETTEUSE_FILTER) += vf_paletteuse.o framesync.o palette.o diff --git a/libavfilter/allfilters.c b/libavfilter/allfilters.c index 3bc045b28f..409099bf1f 100644 --- a/libavfilter/allfilters.c +++ b/libavfilter/allfilters.c @@ -395,6 +395,7 @@ extern const FFFilter ff_vf_overlay_vulkan; extern const FFFilter ff_vf_overlay_cuda; extern const FFFilter ff_vf_owdenoise; extern const FFFilter ff_vf_pad; +extern const FFFilter ff_vf_pad_cuda; extern const FFFilter ff_vf_pad_opencl; extern const FFFilter ff_vf_palettegen; extern const FFFilter ff_vf_paletteuse; diff --git a/libavfilter/version.h b/libavfilter/version.h index d5a6bc143a..1e884d9b44 100644 --- a/libavfilter/version.h +++ b/libavfilter/version.h @@ -31,7 +31,7 @@ #include "version_major.h" -#define LIBAVFILTER_VERSION_MINOR 0 +#define LIBAVFILTER_VERSION_MINOR 1 #define LIBAVFILTER_VERSION_MICRO 100 diff --git a/libavfilter/vf_pad_cuda.c b/libavfilter/vf_pad_cuda.c new file mode 100644 index 0000000000..a9a9036d71 --- /dev/null +++ b/libavfilter/vf_pad_cuda.c @@ -0,0 +1,650 @@ +/* + * This file is part of FFmpeg. + * + * FFmpeg is free software; you can redistribute it and/or + * modify it under the terms of the GNU Lesser General Public + * License as published by the Free Software Foundation; either + * version 2.1 of the License, or (at your option) any later version. + * + * FFmpeg is distributed in the hope that it will be useful, + * but WITHOUT ANY WARRANTY; without even the implied warranty of + * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU + * Lesser General Public License for more details. + * + * You should have received a copy of the GNU Lesser General Public + * License along with FFmpeg; if not, write to the Free Software + * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA + */ + +/** + * @file + * CUDA video padding filter + */ + +#include +#include + +#include "filters.h" +#include "libavutil/avstring.h" +#include "libavutil/common.h" +#include "libavutil/cuda_check.h" +#include "libavutil/eval.h" +#include "libavutil/hwcontext.h" +#include "libavutil/hwcontext_cuda_internal.h" +#include "libavutil/imgutils.h" +#include "libavutil/internal.h" +#include "libavutil/opt.h" +#include "libavutil/pixdesc.h" + +#include "cuda/load_helper.h" + +#define CHECK_CU(x) FF_CUDA_CHECK_DL(ctx, device_hwctx->internal->cuda_dl, x) +#define DIV_UP(a, b) ( ((a) + (b) - 1) / (b) ) +#define BLOCK_X 32 +#define BLOCK_Y 16 + +static const enum AVPixelFormat supported_formats[] = { + AV_PIX_FMT_YUV420P, + AV_PIX_FMT_YUV444P, + AV_PIX_FMT_YUVA420P, + AV_PIX_FMT_YUVA444P, + AV_PIX_FMT_NV12, +}; + +typedef struct CUDAPadContext { + const AVClass *class; + + AVBufferRef *frames_ctx; + + int w, h; ///< output dimensions, a value of 0 will result in the input size + int x, y; ///< offsets of the input area with respect to the padded area + int in_w, in_h; ///< width and height for the padded input video + + char *w_expr; ///< width expression + char *h_expr; ///< height expression + char *x_expr; ///< x offset expression + char *y_expr; ///< y offset expression + + uint8_t rgba_color[4]; ///< color for the padding area + uint8_t parsed_color[4]; + AVRational aspect; + + int eval_mode; + + int last_out_w, last_out_h; ///< used to evaluate the prior output width and height with the incoming frame + + AVCUDADeviceContext *hwctx; + CUmodule cu_module; + CUfunction cu_func_planar; + CUfunction cu_func_uv; +} CUDAPadContext; + +static const char *const var_names[] = { + "in_w", "iw", + "in_h", "ih", + "out_w", "ow", + "out_h", "oh", + "x", + "y", + "a", + "sar", + "dar", + "hsub", + "vsub", + NULL +}; + +enum { + VAR_IN_W, + VAR_IW, + VAR_IN_H, + VAR_IH, + VAR_OUT_W, + VAR_OW, + VAR_OUT_H, + VAR_OH, + VAR_X, + VAR_Y, + VAR_A, + VAR_SAR, + VAR_DAR, + VAR_HSUB, + VAR_VSUB, + VARS_NB +}; + +enum EvalMode { + EVAL_MODE_INIT, + EVAL_MODE_FRAME, + EVAL_MODE_NB +}; + +static int eval_expr(AVFilterContext *ctx) +{ + CUDAPadContext *s = ctx->priv; + AVFilterLink *inlink = ctx->inputs[0]; + const AVPixFmtDescriptor *desc = av_pix_fmt_desc_get(inlink->format); + + double var_values[VARS_NB], res; + char *expr; + int ret; + + var_values[VAR_IN_W] = var_values[VAR_IW] = s->in_w; + var_values[VAR_IN_H] = var_values[VAR_IH] = s->in_h; + var_values[VAR_OUT_W] = var_values[VAR_OW] = NAN; + var_values[VAR_OUT_H] = var_values[VAR_OH] = NAN; + var_values[VAR_A] = (double)s->in_w / s->in_h; + var_values[VAR_SAR] = inlink->sample_aspect_ratio.num ? + (double)inlink->sample_aspect_ratio.num / + inlink->sample_aspect_ratio.den : 1; + var_values[VAR_DAR] = var_values[VAR_A] * var_values[VAR_SAR]; + var_values[VAR_HSUB] = 1 << desc->log2_chroma_w; + var_values[VAR_VSUB] = 1 << desc->log2_chroma_h; + + expr = s->w_expr; + ret = av_expr_parse_and_eval(&res, expr, var_names, var_values, NULL, NULL, NULL, NULL, NULL, 0, ctx); + if (ret < 0) + goto fail; + + s->w = res; + if (s->w < 0) { + av_log(ctx, AV_LOG_ERROR, "Width expression is negative.\n"); + ret = AVERROR(EINVAL); + goto fail; + } + + var_values[VAR_OUT_W] = var_values[VAR_OW] = s->w; + + expr = s->h_expr; + ret = av_expr_parse_and_eval(&res, expr, var_names, var_values, NULL, NULL, NULL, NULL, NULL, 0, ctx); + if (ret < 0) + goto fail; + + s->h = res; + if (s->h < 0) { + av_log(ctx, AV_LOG_ERROR, "Height expression is negative.\n"); + ret = AVERROR(EINVAL); + goto fail; + } + var_values[VAR_OUT_H] = var_values[VAR_OH] = s->h; + + if (!s->h) + s->h = s->in_h; + + var_values[VAR_OUT_H] = var_values[VAR_OH] = s->h; + + + expr = s->w_expr; + ret = av_expr_parse_and_eval(&res, expr, var_names, var_values, NULL, NULL, NULL, NULL, NULL, 0, ctx); + if (ret < 0) + goto fail; + + s->w = res; + if (s->w < 0) { + av_log(ctx, AV_LOG_ERROR, "Width expression is negative.\n"); + ret = AVERROR(EINVAL); + goto fail; + } + if (!s->w) + s->w = s->in_w; + + var_values[VAR_OUT_W] = var_values[VAR_OW] = s->w; + + + expr = s->x_expr; + ret = av_expr_parse_and_eval(&res, expr, var_names, var_values, NULL, NULL, NULL, NULL, NULL, 0, ctx); + if (ret < 0) + goto fail; + + s->x = res; + + + expr = s->y_expr; + ret = av_expr_parse_and_eval(&res, expr, var_names, var_values, NULL, NULL, NULL, NULL, NULL, 0, ctx); + if (ret < 0) + goto fail; + + s->y = res; + + if (s->x < 0 || s->x + s->in_w > s->w) { + s->x = (s->w - s->in_w) / 2; + av_log(ctx, AV_LOG_VERBOSE, "centering X offset.\n"); + } + + if (s->y < 0 || s->y + s->in_h > s->h) { + s->y = (s->h - s->in_h) / 2; + av_log(ctx, AV_LOG_VERBOSE, "centering Y offset.\n"); + } + + s->w = av_clip(s->w, 1, INT_MAX); + s->h = av_clip(s->h, 1, INT_MAX); + + if (s->w < s->in_w || s->h < s->in_h) { + av_log(ctx, AV_LOG_ERROR, "Padded size < input size.\n"); + return AVERROR(EINVAL); + } + + av_log(ctx, AV_LOG_DEBUG, + "w:%d h:%d -> w:%d h:%d x:%d y:%d color:0x%02X%02X%02X%02X\n", + inlink->w, inlink->h, s->w, s->h, s->x, s->y, s->rgba_color[0], + s->rgba_color[1], s->rgba_color[2], s->rgba_color[3]); + + return 0; + +fail: + av_log(ctx, AV_LOG_ERROR, "Error evaluating '%s'\n", expr); + return ret; +} + +static int cuda_pad_alloc_out_frames_ctx(AVFilterContext *ctx, AVBufferRef **out_frames_ctx, const int width, const int height) +{ + AVFilterLink *inlink = ctx->inputs[0]; + FilterLink *inl = ff_filter_link(inlink); + AVHWFramesContext *in_frames_ctx = (AVHWFramesContext *)inl->hw_frames_ctx->data; + int ret; + + *out_frames_ctx = av_hwframe_ctx_alloc(in_frames_ctx->device_ref); + if (!*out_frames_ctx) { + return AVERROR(ENOMEM); + } + + AVHWFramesContext *out_fc = (AVHWFramesContext *)(*out_frames_ctx)->data; + out_fc->format = AV_PIX_FMT_CUDA; + out_fc->sw_format = in_frames_ctx->sw_format; + + out_fc->width = FFALIGN(width, 32); + out_fc->height = FFALIGN(height, 32); + + ret = av_hwframe_ctx_init(*out_frames_ctx); + if (ret < 0) { + av_log(ctx, AV_LOG_ERROR, "Failed to init output ctx\n"); + av_buffer_unref(out_frames_ctx); + return ret; + } + + return 0; +} + +static av_cold int cuda_pad_init(AVFilterContext *ctx) +{ + CUDAPadContext *s = ctx->priv; + if (!s) { + av_log(ctx, AV_LOG_ERROR, "Failed to allocate CUDAPadContext.\n"); + return AVERROR(ENOMEM); + } + + s->last_out_w = -1; + s->last_out_h = -1; + + return 0; +} + +static av_cold void cuda_pad_uninit(AVFilterContext *ctx) +{ + CUDAPadContext *s = ctx->priv; + av_buffer_unref(&s->frames_ctx); + + if (s->cu_module) { + s->hwctx->internal->cuda_dl->cuModuleUnload(s->cu_module); + s->cu_module = NULL; + } +} + +static av_cold int cuda_pad_load_functions(AVFilterContext *ctx) +{ + CUDAPadContext *s = ctx->priv; + CudaFunctions *cu = s->hwctx->internal->cuda_dl; + CUcontext dummy_cu_ctx; + int ret; + + AVCUDADeviceContext *device_hwctx = s->hwctx; + + extern const unsigned char ff_vf_pad_cuda_ptx_data[]; + extern const unsigned int ff_vf_pad_cuda_ptx_len; + + ret = CHECK_CU(cu->cuCtxPushCurrent(device_hwctx->cuda_ctx)); + if (ret < 0) + return ret; + + ret = ff_cuda_load_module(ctx, device_hwctx, &s->cu_module, + ff_vf_pad_cuda_ptx_data, ff_vf_pad_cuda_ptx_len); + if (ret < 0) { + av_log(ctx, AV_LOG_ERROR, "Failed to load CUDA module\n"); + goto end; + } + + ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_planar, s->cu_module, "pad_planar_cuda")); + if (ret < 0) { + av_log(ctx, AV_LOG_ERROR, "Failed to load pad_planar_cuda\n"); + goto end; + } + + ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_uv, s->cu_module, "pad_uv_cuda")); + if (ret < 0) + av_log(ctx, AV_LOG_ERROR, "Failed to load pad_uv_cuda\n"); + +end: + CHECK_CU(cu->cuCtxPopCurrent(&dummy_cu_ctx)); + + return ret; +} + +static int cuda_pad_config_props(AVFilterLink *outlink) +{ + AVFilterContext *ctx = outlink->src; + CUDAPadContext *s = ctx->priv; + + AVFilterLink *inlink = ctx->inputs[0]; + FilterLink *inl = ff_filter_link(inlink); + + FilterLink *ol = ff_filter_link(outlink); + + AVHWFramesContext *in_frames_ctx; + int format_supported = 0; + int ret; + + s->in_w = inlink->w; + s->in_h = inlink->h; + ret = eval_expr(ctx); + if (ret < 0) + return ret; + + if (!inl->hw_frames_ctx) { + av_log(ctx, AV_LOG_ERROR, "No hw context provided on input\n"); + return AVERROR(EINVAL); + } + + in_frames_ctx = (AVHWFramesContext *)inl->hw_frames_ctx->data; + s->hwctx = in_frames_ctx->device_ctx->hwctx; + + for (int i = 0; i < FF_ARRAY_ELEMS(supported_formats); i++) { + if (in_frames_ctx->sw_format == supported_formats[i]) { + format_supported = 1; + break; + } + } + if (!format_supported) { + av_log(ctx, AV_LOG_ERROR, "Unsupported input format.\n"); + return AVERROR(EINVAL); + } + + uint8_t R = s->rgba_color[0]; + uint8_t G = s->rgba_color[1]; + uint8_t B = s->rgba_color[2]; + + int Y = (( 66 * R + 129 * G + 25 * B + 128) >> 8) + 16; + int U = ((-38 * R - 74 * G + 112 * B + 128) >> 8) + 128; + int V = ((112 * R - 94 * G - 18 * B + 128) >> 8) + 128; + s->parsed_color[0] = av_clip_uint8(Y); + s->parsed_color[1] = av_clip_uint8(U); + s->parsed_color[2] = av_clip_uint8(V); + s->parsed_color[3] = s->rgba_color[3]; + + ret = cuda_pad_alloc_out_frames_ctx(ctx, &s->frames_ctx, s->w, s->h); + if (ret < 0) + return ret; + + ol->hw_frames_ctx = av_buffer_ref(s->frames_ctx); + if (!ol->hw_frames_ctx) + return AVERROR(ENOMEM); + + outlink->w = s->w; + outlink->h = s->h; + outlink->time_base = inlink->time_base; + outlink->format = AV_PIX_FMT_CUDA; + + s->last_out_w = s->w; + s->last_out_h = s->h; + + ret = cuda_pad_load_functions(ctx); + if (ret < 0) + return ret; + + return 0; +} + +static int cuda_pad_pad(AVFilterContext *ctx, AVFrame *out, const AVFrame *in) +{ + CUDAPadContext *s = ctx->priv; + FilterLink *inl = ff_filter_link(ctx->inputs[0]); + + AVHWFramesContext *in_frames_ctx = + (AVHWFramesContext *)inl->hw_frames_ctx->data; + const AVPixFmtDescriptor *desc_in = + av_pix_fmt_desc_get(in_frames_ctx->sw_format); + + CudaFunctions *cu = s->hwctx->internal->cuda_dl; + AVCUDADeviceContext *device_hwctx = s->hwctx; + int ret; + + + const int nb_planes = av_pix_fmt_count_planes(in_frames_ctx->sw_format); + for (int plane = 0; plane < nb_planes; plane++) { + + int hsub = (plane == 1 || plane == 2) ? desc_in->log2_chroma_w : 0; + int vsub = (plane == 1 || plane == 2) ? desc_in->log2_chroma_h : 0; + + if (in_frames_ctx->sw_format == AV_PIX_FMT_NV12 && plane == 1) { + hsub = desc_in->log2_chroma_w; + vsub = desc_in->log2_chroma_h; + } + + int src_w = AV_CEIL_RSHIFT(s->in_w, hsub); + int src_h = AV_CEIL_RSHIFT(s->in_h, vsub); + + int dst_w = AV_CEIL_RSHIFT(s->w, hsub); + int dst_h = AV_CEIL_RSHIFT(s->h, vsub); + + int y_plane_offset = AV_CEIL_RSHIFT(s->y, vsub); + int x_plane_offset = AV_CEIL_RSHIFT(s->x, hsub); + + if (x_plane_offset + src_w > dst_w || y_plane_offset + src_h > dst_h) { + av_log(ctx, AV_LOG_ERROR, + "ROI out of bounds in plane %d: offset=(%d,%d) in=(%dx%d) " + "out=(%dx%d)\n", + plane, x_plane_offset, y_plane_offset, src_w, src_h, dst_w, dst_h); + return AVERROR(EINVAL); + } + + CUfunction cuda_func; + uchar2 fill_val_uv; + uint8_t fill_val_planar; + void *p_fill_val; + + if (in_frames_ctx->sw_format == AV_PIX_FMT_NV12 && plane == 1) { + cuda_func = s->cu_func_uv; + fill_val_uv.x = s->parsed_color[1]; + fill_val_uv.y = s->parsed_color[2]; + p_fill_val = &fill_val_uv; + } else { + cuda_func = s->cu_func_planar; + fill_val_planar = s->parsed_color[plane]; + p_fill_val = &fill_val_planar; + } + + CUdeviceptr d_dst = (CUdeviceptr)out->data[plane]; + CUdeviceptr d_src = (CUdeviceptr)in->data[plane]; + + int src_linesize = in->linesize[plane]; + + void *kernel_args[] = { + &d_dst, &out->linesize[plane], &dst_w, &dst_h, + &d_src, &src_linesize, &src_w, &src_h, + &x_plane_offset, &y_plane_offset, p_fill_val + }; + + unsigned int grid_x = DIV_UP(dst_w, BLOCK_X); + unsigned int grid_y = DIV_UP(dst_h, BLOCK_Y); + + ret = CHECK_CU(cu->cuLaunchKernel(cuda_func, grid_x, grid_y, 1, + BLOCK_X, BLOCK_Y, 1, + 0, s->hwctx->stream, kernel_args, NULL)); + + if (ret < 0) { + av_log(ctx, AV_LOG_ERROR, "Failed to launch kernel for plane %d\n", plane); + return ret; + } + } + + return 0; +} + +static int cuda_pad_filter_frame(AVFilterLink *inlink, AVFrame *in) +{ + AVFilterContext *ctx = inlink->dst; + CUDAPadContext *s = ctx->priv; + AVFilterLink *outlink = ctx->outputs[0]; + + FilterLink *outl = ff_filter_link(outlink); + + AVHWFramesContext *out_frames_ctx = + (AVHWFramesContext *)outl->hw_frames_ctx->data; + AVCUDADeviceContext *device_hwctx = out_frames_ctx->device_ctx->hwctx; + + int ret; + + if (s->eval_mode == EVAL_MODE_FRAME) { + s->in_w = in->width; + s->in_h = in->height; + s->aspect = in->sample_aspect_ratio; + + ret = eval_expr(ctx); + if (ret < 0) { + av_frame_free(&in); + return ret; + } + } + + + if (s->x == 0 && s->y == 0 && + s->w == in->width && s->h == in->height) { + av_log(ctx, AV_LOG_DEBUG, "No border. Passing the frame unmodified.\n"); + s->last_out_w = s->w; + s->last_out_h = s->h; + return ff_filter_frame(outlink, in); + } + + + if (s->w != s->last_out_w || s->h != s->last_out_h) { + + av_buffer_unref(&s->frames_ctx); + + ret = cuda_pad_alloc_out_frames_ctx(ctx, &s->frames_ctx, s->w, s->h); + if (ret < 0) + return ret; + + av_buffer_unref(&outl->hw_frames_ctx); + outl->hw_frames_ctx = av_buffer_ref(s->frames_ctx); + if (!outl->hw_frames_ctx) { + av_frame_free(&in); + av_log(ctx, AV_LOG_ERROR, "Failed to allocate output frame context.\n"); + return AVERROR(ENOMEM); + } + outlink->w = s->w; + outlink->h = s->h; + + s->last_out_w = s->w; + s->last_out_h = s->h; + } + + AVFrame *out = av_frame_alloc(); + if (!out) { + av_frame_free(&in); + av_log(ctx, AV_LOG_ERROR, "Failed to allocate output AVFrame.\n"); + return AVERROR(ENOMEM); + } + ret = av_hwframe_get_buffer(outl->hw_frames_ctx, out, 0); + if (ret < 0) { + av_log(ctx, AV_LOG_ERROR, "Unable to get output buffer: %s\n", + av_err2str(ret)); + av_frame_free(&out); + av_frame_free(&in); + return ret; + } + + CUcontext dummy; + ret = CHECK_CU(device_hwctx->internal->cuda_dl->cuCtxPushCurrent( + device_hwctx->cuda_ctx)); + if (ret < 0) { + av_frame_free(&out); + av_frame_free(&in); + return ret; + } + + ret = cuda_pad_pad(ctx, out, in); + + CHECK_CU(device_hwctx->internal->cuda_dl->cuCtxPopCurrent(&dummy)); + + if (ret < 0) { + av_frame_free(&out); + av_frame_free(&in); + return ret; + } + + av_frame_copy_props(out, in); + out->width = s->w; + out->height = s->h; + + + av_reduce(&out->sample_aspect_ratio.num, &out->sample_aspect_ratio.den, + (int64_t)in->sample_aspect_ratio.num * out->height * in->width, + (int64_t)in->sample_aspect_ratio.den * out->width * in->height, + INT_MAX); + + av_frame_free(&in); + return ff_filter_frame(outlink, out); +} + +#define OFFSET(x) offsetof(CUDAPadContext, x) +#define FLAGS (AV_OPT_FLAG_FILTERING_PARAM | AV_OPT_FLAG_VIDEO_PARAM) + +static const AVOption cuda_pad_options[] = { + { "width", "set the pad area width expression", OFFSET(w_expr), AV_OPT_TYPE_STRING, {.str = "iw"}, 0, 0, FLAGS }, + { "w", "set the pad area width expression", OFFSET(w_expr), AV_OPT_TYPE_STRING, {.str = "iw"}, 0, 0, FLAGS }, + { "height", "set the pad area height expression", OFFSET(h_expr), AV_OPT_TYPE_STRING, {.str = "ih"}, 0, 0, FLAGS }, + { "h", "set the pad area height expression", OFFSET(h_expr), AV_OPT_TYPE_STRING, {.str = "ih"}, 0, 0, FLAGS }, + { "x", "set the x offset expression for the input image position", OFFSET(x_expr), AV_OPT_TYPE_STRING, {.str = "0"}, 0, 0, FLAGS }, + { "y", "set the y offset expression for the input image position", OFFSET(y_expr), AV_OPT_TYPE_STRING, {.str = "0"}, 0, 0, FLAGS }, + { "color", "set the color of the padded area border", OFFSET(rgba_color), AV_OPT_TYPE_COLOR, {.str = "black"}, .flags = FLAGS }, + { "eval", "specify when to evaluate expressions", OFFSET(eval_mode), AV_OPT_TYPE_INT, {.i64 = EVAL_MODE_INIT}, 0, EVAL_MODE_NB-1, FLAGS, .unit = "eval" }, + { "init", "eval expressions once during initialization", 0, AV_OPT_TYPE_CONST, {.i64=EVAL_MODE_INIT}, .flags = FLAGS, .unit = "eval" }, + { "frame", "eval expressions during initialization and per-frame", 0, AV_OPT_TYPE_CONST, {.i64=EVAL_MODE_FRAME}, .flags = FLAGS, .unit = "eval" }, + { "aspect", "pad to fit an aspect instead of a resolution", OFFSET(aspect), AV_OPT_TYPE_RATIONAL, {.dbl = 0}, 0, DBL_MAX, FLAGS }, + { NULL } +}; + +static const AVClass cuda_pad_class = { + .class_name = "pad_cuda", + .item_name = av_default_item_name, + .option = cuda_pad_options, + .version = LIBAVUTIL_VERSION_INT, +}; + +static const AVFilterPad cuda_pad_inputs[] = {{ + .name = "default", + .type = AVMEDIA_TYPE_VIDEO, + .filter_frame = cuda_pad_filter_frame +}}; + +static const AVFilterPad cuda_pad_outputs[] = {{ + .name = "default", + .type = AVMEDIA_TYPE_VIDEO, + .config_props = cuda_pad_config_props, +}}; + +const FFFilter ff_vf_pad_cuda = { + .p.name = "pad_cuda", + .p.description = NULL_IF_CONFIG_SMALL("CUDA-based GPU padding filter"), + .init = cuda_pad_init, + .uninit = cuda_pad_uninit, + + .p.priv_class = &cuda_pad_class, + + FILTER_INPUTS(cuda_pad_inputs), + FILTER_OUTPUTS(cuda_pad_outputs), + + FILTER_SINGLE_PIXFMT(AV_PIX_FMT_CUDA), + + .priv_size = sizeof(CUDAPadContext), + .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE, +}; \ No newline at end of file diff --git a/libavfilter/vf_pad_cuda.cu b/libavfilter/vf_pad_cuda.cu new file mode 100644 index 0000000000..4524372e17 --- /dev/null +++ b/libavfilter/vf_pad_cuda.cu @@ -0,0 +1,65 @@ +/* + * This file is part of FFmpeg. + * + * FFmpeg is free software; you can redistribute it and/or + * modify it under the terms of the GNU Lesser General Public + * License as published by the Free Software Foundation; either + * version 2.1 of the License, or (at your option) any later version. + * + * FFmpeg is distributed in the hope that it will be useful, + * but WITHOUT ANY WARRANTY; without even the implied warranty of + * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU + * Lesser General Public License for more details. + * + * You should have received a copy of the GNU Lesser General Public + * License along with FFmpeg; if not, write to the Free Software + * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA + */ + + +extern "C" { + +__global__ void pad_planar_cuda(unsigned char* dst, int dst_pitch, int dst_w, int dst_h, + const unsigned char* src, int src_pitch, int src_w, int src_h, + int roi_x, int roi_y, unsigned char fill_val) +{ + const int x = blockIdx.x * blockDim.x + threadIdx.x; + const int y = blockIdx.y * blockDim.y + threadIdx.y; + + if (x >= dst_w || y >= dst_h) { + return; + } + + if (x >= roi_x && x < (roi_x + src_w) && y >= roi_y && y < (roi_y + src_h)) { + const int src_x = x - roi_x; + const int src_y = y - roi_y; + dst[y * dst_pitch + x] = src[src_y * src_pitch + src_x]; + } else { + dst[y * dst_pitch + x] = fill_val; + } +} + +__global__ void pad_uv_cuda(uchar2* dst, int dst_pitch, int dst_w, int dst_h, + const uchar2* src, int src_pitch, int src_w, int src_h, + int roi_x, int roi_y, uchar2 fill_val) +{ + const int x = blockIdx.x * blockDim.x + threadIdx.x; + const int y = blockIdx.y * blockDim.y + threadIdx.y; + + if (x >= dst_w || y >= dst_h) { + return; + } + + const int dst_pitch_uchar2 = dst_pitch / sizeof(uchar2); + const int src_pitch_uchar2 = src_pitch / sizeof(uchar2); + + if (x >= roi_x && x < (roi_x + src_w) && y >= roi_y && y < (roi_y + src_h)) { + const int src_x = x - roi_x; + const int src_y = y - roi_y; + dst[y * dst_pitch_uchar2 + x] = src[src_y * src_pitch_uchar2 + src_x]; + } else { + dst[y * dst_pitch_uchar2 + x] = fill_val; + } +} + +} \ No newline at end of file -- 2.34.1 _______________________________________________ ffmpeg-devel mailing list ffmpeg-devel@ffmpeg.org https://ffmpeg.org/mailman/listinfo/ffmpeg-devel To unsubscribe, visit link above, or email ffmpeg-devel-request@ffmpeg.org with subject "unsubscribe".