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 585F24F561 for ; Fri, 27 Feb 2026 18:36:15 +0000 (UTC) Authentication-Results: ffbox; dkim=fail (body hash mismatch (got b'Z/8BR3GHrYc2HTyp/ihkgFCejNGqXgxwJCs6VMn0Wr4=', expected b'VgLrP19/WlNpsKpDXD7Ot5FPAJ3RPLDr3/rsSY0O+Ys=')) header.d=gmail.com header.a=rsa-sha256 DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/simple; d=ffmpeg.org; i=@ffmpeg.org; q=dns/txt; s=mail; t=1772217346; h=to : date : message-id : in-reply-to : references : mime-version : reply-to : subject : list-id : list-archive : list-archive : list-help : list-owner : list-post : list-subscribe : list-unsubscribe : from : cc : content-type : content-transfer-encoding : from; bh=IcQgaUavTWwBp4xlYHt9XviKS9nmDaTSUbcbXtBvLXU=; b=mN9VA/XpDS6gA4s20zkEmFhoMnZwxDzZLSZ+N6CK5NzHqthtIL1VsGLRYvSljVHdbhqDE GA51jEV2ug5M1oSCXwS2PuOnkDU+yGpKa48Jo0r5B2JdtAQgYKVBJyqwsz5ZzDMgn93Vr4z Jiw2AoptojbhBJFlq4T64EwVWnh50bbf8IuKBmJOC2wUPzYMRKkKPE20qJvzcbjx4z9NanU 9gqwU7wejmjegDpz59owO58vuGX9WED4jq6pqVWS1Fpaadlfff3PYVblp4D4bkGhYd9IOtT D60PndhM8vX3iywFGGKE2qXXTbR3ZqnmlrI4nEMemtmc7NigGtRjM8YSmReg== Received: from [172.18.0.3] (unknown [172.18.0.3]) by ffbox0-bg.ffmpeg.org (Postfix) with ESMTP id 1B9DF6913C7; Fri, 27 Feb 2026 20:35:46 +0200 (EET) ARC-Seal: i=1; cv=none; a=rsa-sha256; d=ffmpeg.org; s=arc; t=1772217328; b=IKGyI+aKFsJf20kNHBZjwuajj+soBtyPFbEvv59xAO6cGYPVidzd3cSI3a9ZePvaNGXu4 q2tJZDzg0lMJgWg3QssvutbuKM+okZxz/DvQA9yoAP8Rg4ng3VE4tj5Nr/6pKyF0dJvP60y ol4bkOL157CO3veNS/vnTio+uTtgb0LpgwH+F08XjeaehPuAMe2XSvHSlgeLQURDiiBarhP S/oxElLWa2k+7dS8dhklJvILLy8zCyLwkiO11/pvNy29GpBx7DvWLzVYpB1vJkhpidHTcPX a78GlfO603LOEZPAC6KapAAjg/PVvV5rTm0aWPQwUyfIcn6dhxLdz3LhK9Ew== ARC-Message-Signature: i=1; a=rsa-sha256; c=relaxed/relaxed; d=ffmpeg.org; s=arc; t=1772217328; h=from : sender : reply-to : subject : date : message-id : to : cc : mime-version : content-type : content-transfer-encoding : content-id : content-description : resent-date : resent-from : resent-sender : resent-to : resent-cc : resent-message-id : in-reply-to : references : list-id : list-help : list-unsubscribe : list-subscribe : list-post : list-owner : list-archive; bh=Z/8BR3GHrYc2HTyp/ihkgFCejNGqXgxwJCs6VMn0Wr4=; b=abXjZCDZHYjvG6ftAdocVvlDSkdPiG2imEQdwpLgfD/QLF3BBSe2xrZEqIrTxP/UXnikS 9YcsMhv3QcJUk3yXUtaNlua0WxcDGmMsXJSELaCNoHxd3KKQwhLbmd2zx43AxFTe+fPJN1D EaJNskaRiU/mQKBHhk46f6vbEnXtQSumUPQLe3in4bCQD+zg8xRnCLRUM+u63dNsYXVEzGN LTd2cHZ5eFbo+fh0SYxEWfJM/0gAVOXQP/5wTB6DWzx41r+4sdU/6ERK+J5CVKxkvIP+gZu sJrlQu3o4fZXKwEiU4V6ifYI52+nomjUhJGvEjgExX/iHd0CYC+vyogAeo/A== ARC-Authentication-Results: i=1; ffmpeg.org; dkim=pass header.d=gmail.com; arc=none; dmarc=pass header.from=gmail.com policy.dmarc=quarantine Authentication-Results: ffmpeg.org; dkim=pass header.d=gmail.com; arc=none (Message is not ARC signed); dmarc=pass (Used From Domain Record) header.from=gmail.com policy.dmarc=quarantine Received: from mail-pl1-f169.google.com (mail-pl1-f169.google.com [209.85.214.169]) by ffbox0-bg.ffmpeg.org (Postfix) with ESMTPS id 19D3869120F for ; Fri, 27 Feb 2026 20:35:15 +0200 (EET) Received: by mail-pl1-f169.google.com with SMTP id d9443c01a7336-2aaf9191da3so14491505ad.2 for ; Fri, 27 Feb 2026 10:35:14 -0800 (PST) DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=gmail.com; s=20230601; t=1772217313; x=1772822113; darn=ffmpeg.org; h=content-transfer-encoding:mime-version:references:in-reply-to :message-id:date:subject:cc:to:from:from:to:cc:subject:date :message-id:reply-to; bh=VgLrP19/WlNpsKpDXD7Ot5FPAJ3RPLDr3/rsSY0O+Ys=; b=YrBzH5mq77W9g/D93qO0fpMHo3EI6+3UKCC63ttFWBefYQQjvXP0PZ8DWrl5Taqs7O slmXUt52O7IQdLJJmuis1oOsypOK+8AkdtWpeLP67x45E2aNbKWKsutNHZp7lMmEJOUZ achrHFiGj8yNijkOzphqzyaIU67atFC9ED7vsoP3hnkwWWnPiJXYJc4lCZa2NBLFSvDn i05iyF0Aj/5V8u7nTtaSgpcTvQbFiQvjPekS6CVeA8Dfc+7g+BwnwVTHlm5pjd7joiCK 7OGhcRPQjT1Aw2yKtaecz3QpZpP6rQvpyjr2yF+qOH41xZfQPuXR8LiI2jJYPCroJ775 4JvA== X-Google-DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=1e100.net; s=20230601; t=1772217313; x=1772822113; h=content-transfer-encoding:mime-version:references:in-reply-to :message-id:date:subject:cc:to:from:x-gm-gg:x-gm-message-state:from :to:cc:subject:date:message-id:reply-to; bh=VgLrP19/WlNpsKpDXD7Ot5FPAJ3RPLDr3/rsSY0O+Ys=; b=wGJLoMZz0+5rPWU3vx0shVq3E14JUMy9Qb6lvg+aSiIeUaa284pkUQublKCOoRfbNx ekiAkZh6/7pWnR3brTHh3RsVlwPqYWAoC3HWHQZ5Yw8EdWAQ7yN6AEBh+F00W2U0fKBL ogAu3SRw8d4iHS/EI2lBqRWjtqY9xn8RL6jbhk8fPQP0TTSwOEfFZhhlcuxiXvxchq61 mzfPGsZLqFC5ZW2mCPnAVmgG0Ak1GhflAJsZbQ24Ko5faDsq/R/CV9Ic6emLUZp4MZ95 b1ozJq6YCVj/RL6JEr3DYy4Du16TcT4CwhqdEveuKFeBAhDz7scUCIwn+mtvBt1SWHWu U0RQ== X-Gm-Message-State: AOJu0YxIUFRHOjXBRqDf/QpDUq+G4h5ZM9hg1jd7GkXDkdIbiWsZ0A1R haKxoE0k/et/DVt5mDisEK1IpizTXuSpl83mCBi1k2A3SXjYsoxKtYFgTpkNYdMw X-Gm-Gg: ATEYQzxGRFCxHNpLmVY7T+ukAslap9VAit22dfDr28Bedy7dDKvsruOFCPlSGLe1lub V+u53CaAX9b/YOhzhySVvduI5iG57tGXyRNM5PSLUfQiKWDwNcrQOJMwho99tBF2hVVoDc+RmeV IrN4XkaJ+BxI8x3wxnF2REHL1b/vtUTXEPuoTrYJOB3Q3MntPgfh69IgFhPRp5guhmnmDFFdceX QIVx+jymP9K5HBChmSi7VWiwiMotdBctSgzrG5Hb5hMsr4jhpWgBzJ+ZYwTQT16j2sckF0AScQl FWrgVqePuY4liQddL4ajR7efJqcorC6XYxm40YvJucBeYdDtoBH/Zmj+draaQC0KWOGeLFsZ1ZA mx23JG13YpwhN8nuNDwo+Pq8TOLuMm5wVN/ZtyGwqjKvZ0f10sV1LWX07tLr0OkAhfzA5qNlwn5 r4oNHXMnXnGvjADZvzVHcXNv3wDb7ApAbH774GCFk= X-Received: by 2002:a17:902:e744:b0:2aa:ea3d:a37b with SMTP id d9443c01a7336-2ae2e251a6bmr37555425ad.2.1772217312498; Fri, 27 Feb 2026 10:35:12 -0800 (PST) Received: from velotio-ThinkPad-E14.. ([2405:201:1008:b04e:dc77:fc69:d4ab:6766]) by smtp.gmail.com with ESMTPSA id 41be03b00d2f7-c70fa62147esm5268606a12.12.2026.02.27.10.35.10 (version=TLS1_3 cipher=TLS_AES_256_GCM_SHA384 bits=256/256); Fri, 27 Feb 2026 10:35:11 -0800 (PST) To: ffmpeg-devel@ffmpeg.org Date: Sat, 28 Feb 2026 00:05:08 +0530 Message-Id: <20260227183508.667671-1-f1k2faeez@gmail.com> X-Mailer: git-send-email 2.34.1 In-Reply-To: <20250605110938.686643-1-f1k2faeez@gmail.com> References: <20250605110938.686643-1-f1k2faeez@gmail.com> MIME-Version: 1.0 Message-ID-Hash: QK4QTQOI7UJNEF26AQWSC5RI4KSE6RUE X-Message-ID-Hash: QK4QTQOI7UJNEF26AQWSC5RI4KSE6RUE X-MailFrom: SRS0=OUq/=A7=gmail.com=f1k2faeez@ffmpeg.org X-Mailman-Rule-Misses: dmarc-mitigation; no-senders; approved; loop; banned-address; header-match-ffmpeg-devel.ffmpeg.org-0; header-match-ffmpeg-devel.ffmpeg.org-1; header-match-ffmpeg-devel.ffmpeg.org-2; header-match-ffmpeg-devel.ffmpeg.org-3; emergency; member-moderation; nonmember-moderation; administrivia; implicit-dest; max-recipients; max-size; news-moderation; no-subject; digests; suspicious-header X-Mailman-Version: 3.3.10 Precedence: list Reply-To: FFmpeg development discussions and patches Subject: [FFmpeg-devel] [PATCH v2] avfilter: add CUDA-accelerated transpose filter List-Id: FFmpeg development discussions and patches Archived-At: Archived-At: List-Archive: List-Archive: List-Help: List-Owner: List-Post: List-Subscribe: List-Unsubscribe: From: Faeez Kadiri via ffmpeg-devel Cc: Faeez Kadiri Content-Type: text/plain; charset="us-ascii" Content-Transfer-Encoding: 7bit Archived-At: List-Archive: List-Post: Add a new CUDA-accelerated transpose filter (transpose_cuda) that provides hardware-accelerated video transposition on NVIDIA GPUs. Supported operations: - 90 degree clockwise/counter-clockwise rotation (with optional flip) - 180 degree rotation - Horizontal and vertical flip Supported pixel formats: - YUV420P, YUV444P, YUV420P10, YUV444P10, YUV444P16 (planar) - NV12, P010, P016 (semi-planar) - RGB32, BGR32, 0RGB32, 0BGR32 (packed) The implementation uses element-size based CUDA kernels that are selected dynamically from pixel format descriptors, allowing format support to be extended without adding new kernels. The filter is API-compatible with the existing transpose filter and includes passthrough mode support. Signed-off-by: Faeez Kadiri Made-with: Cursor --- Changelog | 2 +- configure | 2 + doc/filters.texi | 53 ++++ libavfilter/Makefile | 1 + libavfilter/allfilters.c | 1 + libavfilter/vf_transpose_cuda.c | 505 +++++++++++++++++++++++++++++++ libavfilter/vf_transpose_cuda.cu | 138 +++++++++ 7 files changed, 701 insertions(+), 1 deletion(-) create mode 100644 libavfilter/vf_transpose_cuda.c create mode 100644 libavfilter/vf_transpose_cuda.cu diff --git a/Changelog b/Changelog index 26416cb1d6..9f0a4d8ede 100644 --- a/Changelog +++ b/Changelog @@ -24,6 +24,7 @@ version : - Remove the old HLS protocol handler - Vulkan compute codec optimizations - swscale Vulkan support +- Transpose CUDA filter (transpose_cuda) version 8.0: @@ -62,7 +63,6 @@ version 8.0: - ffprobe -codec option - HDR10+ metadata passthrough when decoding/encoding with libaom-av1 - version 7.1: - Raw Captions with Time (RCWT) closed caption demuxer - LC3/LC3plus decoding/encoding using external library liblc3 diff --git a/configure b/configure index 87a9c02686..b87daee557 100755 --- a/configure +++ b/configure @@ -3509,6 +3509,8 @@ scale_cuda_filter_deps="ffnvcodec" scale_cuda_filter_deps_any="cuda_nvcc cuda_llvm" thumbnail_cuda_filter_deps="ffnvcodec" thumbnail_cuda_filter_deps_any="cuda_nvcc cuda_llvm" +transpose_cuda_filter_deps="ffnvcodec" +transpose_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" diff --git a/doc/filters.texi b/doc/filters.texi index e49dd9ef0d..d4de7f14be 100644 --- a/doc/filters.texi +++ b/doc/filters.texi @@ -27291,6 +27291,59 @@ Thumbnails are extracted from every @var{n}=150-frame batch, selecting one per b @end itemize +@anchor{transpose_cuda} +@subsection transpose_cuda + +Transpose rows with columns in the input video and optionally flip it using +CUDA. This is the CUDA variant of the @ref{transpose} filter, which shares +mostly the same options. In addition to the standard transpose directions, this +filter also supports 180 degree rotation, horizontal flip, and vertical flip. + +It accepts the following parameters: + +@table @option + +@item dir +Specify the transposition direction. + +Can assume the following values: +@table @samp +@item cclock_flip +Rotate by 90 degrees counterclockwise and vertically flip. (default) + +@item clock +Rotate by 90 degrees clockwise. + +@item cclock +Rotate by 90 degrees counterclockwise. + +@item clock_flip +Rotate by 90 degrees clockwise and vertically flip. + +@item reversal +Rotate by 180 degrees. + +@item hflip +Flip horizontally. + +@item vflip +Flip vertically. +@end table + +@item passthrough +Do not apply the transposition if the input geometry matches the one +specified by the specified value. It accepts the following values: +@table @samp +@item none +Always apply transposition. (default) +@item portrait +Preserve portrait geometry (when @var{height} >= @var{width}). +@item landscape +Preserve landscape geometry (when @var{width} >= @var{height}). +@end table + +@end table + @subsection yadif_cuda Deinterlace the input video using the @ref{yadif} algorithm, but implemented diff --git a/libavfilter/Makefile b/libavfilter/Makefile index 6ecacc346b..4462ac2970 100644 --- a/libavfilter/Makefile +++ b/libavfilter/Makefile @@ -545,6 +545,7 @@ OBJS-$(CONFIG_TONEMAP_OPENCL_FILTER) += vf_tonemap_opencl.o opencl.o \ OBJS-$(CONFIG_TONEMAP_VAAPI_FILTER) += vf_tonemap_vaapi.o vaapi_vpp.o OBJS-$(CONFIG_TPAD_FILTER) += vf_tpad.o OBJS-$(CONFIG_TRANSPOSE_FILTER) += vf_transpose.o +OBJS-$(CONFIG_TRANSPOSE_CUDA_FILTER) += vf_transpose_cuda.o vf_transpose_cuda.ptx.o cuda/load_helper.o OBJS-$(CONFIG_TRANSPOSE_NPP_FILTER) += vf_transpose_npp.o OBJS-$(CONFIG_TRANSPOSE_OPENCL_FILTER) += vf_transpose_opencl.o opencl.o opencl/transpose.o OBJS-$(CONFIG_TRANSPOSE_VAAPI_FILTER) += vf_transpose_vaapi.o vaapi_vpp.o diff --git a/libavfilter/allfilters.c b/libavfilter/allfilters.c index 458f8c5373..f94a64aabc 100644 --- a/libavfilter/allfilters.c +++ b/libavfilter/allfilters.c @@ -512,6 +512,7 @@ extern const FFFilter ff_vf_tonemap_opencl; extern const FFFilter ff_vf_tonemap_vaapi; extern const FFFilter ff_vf_tpad; extern const FFFilter ff_vf_transpose; +extern const FFFilter ff_vf_transpose_cuda; extern const FFFilter ff_vf_transpose_npp; extern const FFFilter ff_vf_transpose_opencl; extern const FFFilter ff_vf_transpose_vaapi; diff --git a/libavfilter/vf_transpose_cuda.c b/libavfilter/vf_transpose_cuda.c new file mode 100644 index 0000000000..925a82894c --- /dev/null +++ b/libavfilter/vf_transpose_cuda.c @@ -0,0 +1,505 @@ +/* + * Copyright (c) 2025, Faeez Kadiri < f1k2faeez at gmail dot com> + * + * This file is part of FFmpeg. + * + * FFmpeg is free software; you can redistribute it and/or + * modify it under the terms of the GNU Lesser General Public + * License as published by the Free Software Foundation; either + * version 2.1 of the License, or (at your option) any later version. + * + * FFmpeg is distributed in the hope that it will be useful, + * but WITHOUT ANY WARRANTY; without even the implied warranty of + * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU + * Lesser General Public License for more details. + * + * You should have received a copy of the GNU Lesser General Public + * License along with FFmpeg; if not, write to the Free Software + * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA + */ + +#include "libavutil/common.h" +#include "libavutil/hwcontext.h" +#include "libavutil/hwcontext_cuda_internal.h" +#include "libavutil/cuda_check.h" +#include "libavutil/internal.h" +#include "libavutil/opt.h" +#include "libavutil/pixdesc.h" + +#include "avfilter.h" +#include "filters.h" +#include "video.h" +#include "transpose.h" + +#include "cuda/load_helper.h" + +static const enum AVPixelFormat supported_formats[] = { + AV_PIX_FMT_YUV420P, + AV_PIX_FMT_YUV444P, + AV_PIX_FMT_YUV420P10, + AV_PIX_FMT_YUV444P10, + AV_PIX_FMT_YUV444P16, + AV_PIX_FMT_NV12, + AV_PIX_FMT_P010, + AV_PIX_FMT_P016, + AV_PIX_FMT_0RGB32, + AV_PIX_FMT_0BGR32, + AV_PIX_FMT_RGB32, + AV_PIX_FMT_BGR32, +}; + +#define DIV_UP(a, b) ( ((a) + (b) - 1) / (b) ) +#define BLOCKX 32 +#define BLOCKY 16 + +#define CHECK_CU(x) FF_CUDA_CHECK_DL(ctx, s->hwctx->internal->cuda_dl, x) + +typedef struct TransposeCUDAContext { + const AVClass *class; + + AVCUDADeviceContext *hwctx; + AVBufferRef *device_ref; + + const AVPixFmtDescriptor *in_desc; + int in_planes; + int in_plane_depths[4]; + int in_plane_channels[4]; + + CUmodule cu_module; + CUstream cu_stream; + + CUfunction cu_func; + CUfunction cu_func_uv; + + int passthrough_mode; + int dir; +} TransposeCUDAContext; + +static int format_is_supported(enum AVPixelFormat fmt) +{ + for (int i = 0; i < FF_ARRAY_ELEMS(supported_formats); i++) + if (supported_formats[i] == fmt) + return 1; + return 0; +} + +static av_cold void set_format_info(AVFilterContext *ctx, enum AVPixelFormat format) +{ + TransposeCUDAContext *s = ctx->priv; + + s->in_desc = av_pix_fmt_desc_get(format); + s->in_planes = av_pix_fmt_count_planes(format); + + for (int i = 0; i < s->in_desc->nb_components; i++) { + int d = (s->in_desc->comp[i].depth + 7) / 8; + int p = s->in_desc->comp[i].plane; + s->in_plane_channels[p] = FFMAX(s->in_plane_channels[p], + s->in_desc->comp[i].step / d); + s->in_plane_depths[p] = s->in_desc->comp[i].depth; + } +} + +static const char *get_func_name(int depth, int channels) +{ + if (channels == 4 && depth <= 8) + return "Transpose_uchar4"; + if (channels == 2 && depth <= 8) + return "Transpose_uchar2"; + if (channels == 2 && depth > 8) + return "Transpose_ushort2"; + if (depth > 8) + return "Transpose_ushort"; + return "Transpose_uchar"; +} + +static const char *get_uv_func_name(int depth, int channels) +{ + if (channels >= 2 && depth <= 8) + return "Transpose_uchar2"; + if (channels >= 2 && depth > 8) + return "Transpose_ushort2"; + if (depth > 8) + return "Transpose_ushort_uv"; + return "Transpose_uchar_uv"; +} + +static av_cold int transpose_cuda_load_functions(AVFilterContext *ctx) +{ + TransposeCUDAContext *s = ctx->priv; + CUcontext dummy, cuda_ctx = s->hwctx->cuda_ctx; + CudaFunctions *cu = s->hwctx->internal->cuda_dl; + const char *func_name; + int ret; + + extern const unsigned char ff_vf_transpose_cuda_ptx_data[]; + extern const unsigned int ff_vf_transpose_cuda_ptx_len; + + ret = CHECK_CU(cu->cuCtxPushCurrent(cuda_ctx)); + if (ret < 0) + return ret; + + ret = ff_cuda_load_module(ctx, s->hwctx, &s->cu_module, + ff_vf_transpose_cuda_ptx_data, + ff_vf_transpose_cuda_ptx_len); + if (ret < 0) + goto fail; + + func_name = get_func_name(s->in_plane_depths[0], s->in_plane_channels[0]); + ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_func, s->cu_module, func_name)); + if (ret < 0) { + av_log(ctx, AV_LOG_FATAL, "Failed loading %s\n", func_name); + goto fail; + } + + if (s->in_planes > 1) { + func_name = get_uv_func_name(s->in_plane_depths[1], s->in_plane_channels[1]); + ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_uv, s->cu_module, func_name)); + if (ret < 0) { + av_log(ctx, AV_LOG_FATAL, "Failed loading %s\n", func_name); + goto fail; + } + } + +fail: + CHECK_CU(cu->cuCtxPopCurrent(&dummy)); + return ret; +} + +static av_cold int init_processing_chain(AVFilterContext *ctx, + int out_width, int out_height) +{ + TransposeCUDAContext *s = ctx->priv; + FilterLink *inl = ff_filter_link(ctx->inputs[0]); + FilterLink *outl = ff_filter_link(ctx->outputs[0]); + AVHWFramesContext *in_frames_ctx; + AVBufferRef *hw_frames_ctx; + AVHWFramesContext *out_frames_ctx; + int 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; + + if (!format_is_supported(in_frames_ctx->sw_format)) { + av_log(ctx, AV_LOG_ERROR, "Unsupported format: %s\n", + av_get_pix_fmt_name(in_frames_ctx->sw_format)); + return AVERROR(ENOSYS); + } + + set_format_info(ctx, in_frames_ctx->sw_format); + + s->device_ref = av_buffer_ref(in_frames_ctx->device_ref); + if (!s->device_ref) + return AVERROR(ENOMEM); + + s->hwctx = in_frames_ctx->device_ctx->hwctx; + s->cu_stream = s->hwctx->stream; + + hw_frames_ctx = av_hwframe_ctx_alloc(s->device_ref); + if (!hw_frames_ctx) + return AVERROR(ENOMEM); + + out_frames_ctx = (AVHWFramesContext *)hw_frames_ctx->data; + out_frames_ctx->format = AV_PIX_FMT_CUDA; + out_frames_ctx->sw_format = in_frames_ctx->sw_format; + out_frames_ctx->width = out_width; + out_frames_ctx->height = out_height; + + ret = av_hwframe_ctx_init(hw_frames_ctx); + if (ret < 0) { + av_buffer_unref(&hw_frames_ctx); + return ret; + } + + av_buffer_unref(&outl->hw_frames_ctx); + outl->hw_frames_ctx = hw_frames_ctx; + + return 0; +} + +static av_cold void transpose_cuda_uninit(AVFilterContext *ctx) +{ + TransposeCUDAContext *s = ctx->priv; + + if (s->hwctx && s->cu_module) { + CudaFunctions *cu = s->hwctx->internal->cuda_dl; + CUcontext dummy; + + CHECK_CU(cu->cuCtxPushCurrent(s->hwctx->cuda_ctx)); + CHECK_CU(cu->cuModuleUnload(s->cu_module)); + s->cu_module = NULL; + CHECK_CU(cu->cuCtxPopCurrent(&dummy)); + } + + av_buffer_unref(&s->device_ref); + s->hwctx = NULL; +} + +static int transpose_cuda_call_kernel(AVFilterContext *ctx, CUfunction func, + CUtexObject src_tex[4], + AVFrame *out_frame, + int width, int height, + int dst_width, int dst_height, + int dst_pitch, + int src_width, int src_height, + int dir) +{ + TransposeCUDAContext *s = ctx->priv; + CudaFunctions *cu = s->hwctx->internal->cuda_dl; + + CUdeviceptr dst_devptr[4] = { + (CUdeviceptr)out_frame->data[0], (CUdeviceptr)out_frame->data[1], + (CUdeviceptr)out_frame->data[2], (CUdeviceptr)out_frame->data[3] + }; + + void *args[] = { + &src_tex[0], &src_tex[1], &src_tex[2], &src_tex[3], + &dst_devptr[0], &dst_devptr[1], &dst_devptr[2], &dst_devptr[3], + &width, &height, &dst_pitch, + &dst_width, &dst_height, + &src_width, &src_height, + &dir + }; + + return CHECK_CU(cu->cuLaunchKernel(func, + DIV_UP(width, BLOCKX), DIV_UP(height, BLOCKY), 1, + BLOCKX, BLOCKY, 1, + 0, s->cu_stream, args, NULL)); +} + +static int transpose_cuda_filter_frame(AVFilterLink *inlink, AVFrame *in) +{ + AVFilterContext *ctx = inlink->dst; + AVFilterLink *outlink = ctx->outputs[0]; + TransposeCUDAContext *s = ctx->priv; + AVFrame *out = NULL; + CUtexObject tex[4] = { 0, 0, 0, 0 }; + CUcontext dummy; + CudaFunctions *cu; + int ret, i; + + if (s->passthrough_mode) + return ff_filter_frame(outlink, in); + + cu = s->hwctx->internal->cuda_dl; + + ret = CHECK_CU(cu->cuCtxPushCurrent(s->hwctx->cuda_ctx)); + if (ret < 0) + return ret; + + out = ff_get_video_buffer(outlink, outlink->w, outlink->h); + if (!out) { + ret = AVERROR(ENOMEM); + goto fail; + } + + ret = av_frame_copy_props(out, in); + if (ret < 0) + goto fail; + + for (i = 0; i < s->in_planes; i++) { + CUDA_TEXTURE_DESC tex_desc = { + .filterMode = CU_TR_FILTER_MODE_POINT, + .flags = CU_TRSF_READ_AS_INTEGER, + }; + + CUDA_RESOURCE_DESC res_desc = { + .resType = CU_RESOURCE_TYPE_PITCH2D, + .res.pitch2D.format = s->in_plane_depths[i] <= 8 ? + CU_AD_FORMAT_UNSIGNED_INT8 : + CU_AD_FORMAT_UNSIGNED_INT16, + .res.pitch2D.numChannels = s->in_plane_channels[i], + .res.pitch2D.pitchInBytes = in->linesize[i], + .res.pitch2D.devPtr = (CUdeviceptr)in->data[i], + }; + + if (i == 1 || i == 2) { + res_desc.res.pitch2D.width = AV_CEIL_RSHIFT(in->width, s->in_desc->log2_chroma_w); + res_desc.res.pitch2D.height = AV_CEIL_RSHIFT(in->height, s->in_desc->log2_chroma_h); + } else { + res_desc.res.pitch2D.width = in->width; + res_desc.res.pitch2D.height = in->height; + } + + ret = CHECK_CU(cu->cuTexObjectCreate(&tex[i], &res_desc, &tex_desc, NULL)); + if (ret < 0) + goto fail; + } + + ret = transpose_cuda_call_kernel(ctx, s->cu_func, tex, out, + out->width, out->height, + out->width, out->height, + out->linesize[0], + in->width, in->height, s->dir); + if (ret < 0) + goto fail; + + if (s->in_planes > 1) { + ret = transpose_cuda_call_kernel(ctx, s->cu_func_uv, tex, out, + AV_CEIL_RSHIFT(out->width, s->in_desc->log2_chroma_w), + AV_CEIL_RSHIFT(out->height, s->in_desc->log2_chroma_h), + out->width, out->height, + out->linesize[1], + AV_CEIL_RSHIFT(in->width, s->in_desc->log2_chroma_w), + AV_CEIL_RSHIFT(in->height, s->in_desc->log2_chroma_h), + s->dir); + if (ret < 0) + goto fail; + } + + switch (s->dir) { + case TRANSPOSE_CCLOCK_FLIP: + case TRANSPOSE_CLOCK: + case TRANSPOSE_CCLOCK: + case TRANSPOSE_CLOCK_FLIP: + if (in->sample_aspect_ratio.num == 0) { + out->sample_aspect_ratio = in->sample_aspect_ratio; + } else { + out->sample_aspect_ratio.num = in->sample_aspect_ratio.den; + out->sample_aspect_ratio.den = in->sample_aspect_ratio.num; + } + break; + default: + out->sample_aspect_ratio = in->sample_aspect_ratio; + break; + } + + for (i = 0; i < FF_ARRAY_ELEMS(tex); i++) + if (tex[i]) + CHECK_CU(cu->cuTexObjectDestroy(tex[i])); + + CHECK_CU(cu->cuCtxPopCurrent(&dummy)); + + av_frame_free(&in); + return ff_filter_frame(outlink, out); + +fail: + for (i = 0; i < FF_ARRAY_ELEMS(tex); i++) + if (tex[i]) + CHECK_CU(cu->cuTexObjectDestroy(tex[i])); + + av_frame_free(&in); + av_frame_free(&out); + CHECK_CU(cu->cuCtxPopCurrent(&dummy)); + return ret; +} + +static av_cold int transpose_cuda_config_output(AVFilterLink *outlink) +{ + AVFilterContext *ctx = outlink->src; + TransposeCUDAContext *s = ctx->priv; + AVFilterLink *inlink = ctx->inputs[0]; + int out_w, out_h; + int ret; + + if ((inlink->w >= inlink->h && s->passthrough_mode == TRANSPOSE_PT_TYPE_LANDSCAPE) || + (inlink->w <= inlink->h && s->passthrough_mode == TRANSPOSE_PT_TYPE_PORTRAIT)) { + FilterLink *inl = ff_filter_link(inlink); + FilterLink *outl = ff_filter_link(outlink); + outlink->w = inlink->w; + outlink->h = inlink->h; + if (inl->hw_frames_ctx) + outl->hw_frames_ctx = av_buffer_ref(inl->hw_frames_ctx); + av_log(ctx, AV_LOG_VERBOSE, + "w:%d h:%d -> w:%d h:%d (passthrough mode)\n", + inlink->w, inlink->h, inlink->w, inlink->h); + return 0; + } + s->passthrough_mode = TRANSPOSE_PT_TYPE_NONE; + + switch (s->dir) { + case TRANSPOSE_CCLOCK_FLIP: + case TRANSPOSE_CCLOCK: + case TRANSPOSE_CLOCK: + case TRANSPOSE_CLOCK_FLIP: + out_w = inlink->h; + out_h = inlink->w; + break; + default: + out_w = inlink->w; + out_h = inlink->h; + break; + } + + outlink->w = out_w; + outlink->h = out_h; + + ret = init_processing_chain(ctx, out_w, out_h); + if (ret < 0) + return ret; + + ret = transpose_cuda_load_functions(ctx); + if (ret < 0) + return ret; + + av_log(ctx, AV_LOG_VERBOSE, + "w:%d h:%d dir:%d -> w:%d h:%d rotation:%s vflip:%d\n", + inlink->w, inlink->h, s->dir, out_w, out_h, + s->dir == 1 || s->dir == 3 ? "clockwise" : "counterclockwise", + s->dir == 0 || s->dir == 3); + + return 0; +} + +static AVFrame *get_video_buffer(AVFilterLink *inlink, int w, int h) +{ + TransposeCUDAContext *s = inlink->dst->priv; + + return s->passthrough_mode ? + ff_null_get_video_buffer(inlink, w, h) : + ff_default_get_video_buffer(inlink, w, h); +} + +#define OFFSET(x) offsetof(TransposeCUDAContext, x) +#define FLAGS (AV_OPT_FLAG_VIDEO_PARAM | AV_OPT_FLAG_FILTERING_PARAM) +static const AVOption transpose_cuda_options[] = { + { "dir", "set transpose direction", OFFSET(dir), AV_OPT_TYPE_INT, { .i64 = TRANSPOSE_CCLOCK_FLIP }, 0, 6, FLAGS, .unit = "dir" }, + { "cclock_flip", "rotate counter-clockwise with vertical flip", 0, AV_OPT_TYPE_CONST, { .i64 = TRANSPOSE_CCLOCK_FLIP }, .flags=FLAGS, .unit = "dir" }, + { "clock", "rotate clockwise", 0, AV_OPT_TYPE_CONST, { .i64 = TRANSPOSE_CLOCK }, .flags=FLAGS, .unit = "dir" }, + { "cclock", "rotate counter-clockwise", 0, AV_OPT_TYPE_CONST, { .i64 = TRANSPOSE_CCLOCK }, .flags=FLAGS, .unit = "dir" }, + { "clock_flip", "rotate clockwise with vertical flip", 0, AV_OPT_TYPE_CONST, { .i64 = TRANSPOSE_CLOCK_FLIP }, .flags=FLAGS, .unit = "dir" }, + { "reversal", "rotate by half-turn", 0, AV_OPT_TYPE_CONST, { .i64 = TRANSPOSE_REVERSAL }, .flags=FLAGS, .unit = "dir" }, + { "hflip", "flip horizontally", 0, AV_OPT_TYPE_CONST, { .i64 = TRANSPOSE_HFLIP }, .flags=FLAGS, .unit = "dir" }, + { "vflip", "flip vertically", 0, AV_OPT_TYPE_CONST, { .i64 = TRANSPOSE_VFLIP }, .flags=FLAGS, .unit = "dir" }, + + { "passthrough", "do not apply transposition if the input matches the specified geometry", + OFFSET(passthrough_mode), AV_OPT_TYPE_INT, {.i64=TRANSPOSE_PT_TYPE_NONE}, 0, INT_MAX, FLAGS, .unit = "passthrough" }, + { "none", "always apply transposition", 0, AV_OPT_TYPE_CONST, {.i64=TRANSPOSE_PT_TYPE_NONE}, INT_MIN, INT_MAX, FLAGS, .unit = "passthrough" }, + { "portrait", "preserve portrait geometry", 0, AV_OPT_TYPE_CONST, {.i64=TRANSPOSE_PT_TYPE_PORTRAIT}, INT_MIN, INT_MAX, FLAGS, .unit = "passthrough" }, + { "landscape", "preserve landscape geometry", 0, AV_OPT_TYPE_CONST, {.i64=TRANSPOSE_PT_TYPE_LANDSCAPE}, INT_MIN, INT_MAX, FLAGS, .unit = "passthrough" }, + + { NULL } +}; + +AVFILTER_DEFINE_CLASS(transpose_cuda); + +static const AVFilterPad transpose_cuda_inputs[] = { + { + .name = "default", + .type = AVMEDIA_TYPE_VIDEO, + .filter_frame = transpose_cuda_filter_frame, + .get_buffer.video = get_video_buffer, + }, +}; + +static const AVFilterPad transpose_cuda_outputs[] = { + { + .name = "default", + .type = AVMEDIA_TYPE_VIDEO, + .config_props = transpose_cuda_config_output, + }, +}; + +const FFFilter ff_vf_transpose_cuda = { + .p.name = "transpose_cuda", + .p.description = NULL_IF_CONFIG_SMALL("CUDA accelerated video transpose"), + .p.priv_class = &transpose_cuda_class, + .priv_size = sizeof(TransposeCUDAContext), + .uninit = transpose_cuda_uninit, + FILTER_INPUTS(transpose_cuda_inputs), + FILTER_OUTPUTS(transpose_cuda_outputs), + FILTER_SINGLE_PIXFMT(AV_PIX_FMT_CUDA), + .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE, +}; diff --git a/libavfilter/vf_transpose_cuda.cu b/libavfilter/vf_transpose_cuda.cu new file mode 100644 index 0000000000..fd37fd2191 --- /dev/null +++ b/libavfilter/vf_transpose_cuda.cu @@ -0,0 +1,138 @@ +/* + * Copyright (c) 2025, Faeez Kadiri < f1k2faeez at gmail dot com> + * + * This file is part of FFmpeg. + * + * FFmpeg is free software; you can redistribute it and/or + * modify it under the terms of the GNU Lesser General Public + * License as published by the Free Software Foundation; either + * version 2.1 of the License, or (at your option) any later version. + * + * FFmpeg is distributed in the hope that it will be useful, + * but WITHOUT ANY WARRANTY; without even the implied warranty of + * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU + * Lesser General Public License for more details. + * + * You should have received a copy of the GNU Lesser General Public + * License along with FFmpeg; if not, write to the Free Software + * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA + */ + +#include "cuda/vector_helpers.cuh" +#include "transpose.h" + +#define FIXED_PITCH(T) \ + (dst_pitch / sizeof(T)) + +__device__ static inline void get_transpose_coords(int src_x, int src_y, + int src_width, int src_height, + int *dst_x, int *dst_y, + int dst_width, int dst_height, + int dir) +{ + switch (dir) { + case TRANSPOSE_CCLOCK_FLIP: + *dst_x = src_y; + *dst_y = src_x; + break; + case TRANSPOSE_CLOCK: + *dst_x = src_y; + *dst_y = src_width - 1 - src_x; + break; + case TRANSPOSE_CCLOCK: + *dst_x = src_height - 1 - src_y; + *dst_y = src_x; + break; + case TRANSPOSE_CLOCK_FLIP: + *dst_x = src_height - 1 - src_y; + *dst_y = src_width - 1 - src_x; + break; + case TRANSPOSE_REVERSAL: + *dst_x = src_width - 1 - src_x; + *dst_y = src_height - 1 - src_y; + break; + case TRANSPOSE_HFLIP: + *dst_x = src_width - 1 - src_x; + *dst_y = src_y; + break; + case TRANSPOSE_VFLIP: + *dst_x = src_x; + *dst_y = src_height - 1 - src_y; + break; + default: + *dst_x = src_x; + *dst_y = src_y; + break; + } +} + +#define TRANSPOSE_KERNEL_ARGS(T) \ + cudaTextureObject_t src_tex_0, cudaTextureObject_t src_tex_1, \ + cudaTextureObject_t src_tex_2, cudaTextureObject_t src_tex_3, \ + T *dst_0, T *dst_1, T *dst_2, T *dst_3, \ + int width, int height, int dst_pitch, \ + int dst_width, int dst_height, \ + int src_width, int src_height, int dir + +#define KERNEL_PREAMBLE(T) \ + int xo = blockIdx.x * blockDim.x + threadIdx.x; \ + int yo = blockIdx.y * blockDim.y + threadIdx.y; \ + if (xo >= width || yo >= height) return; \ + int src_x, src_y; \ + get_transpose_coords(xo, yo, width, height, \ + &src_x, &src_y, src_width, src_height, dir); + +extern "C" { + +__global__ void Transpose_uchar(TRANSPOSE_KERNEL_ARGS(uchar)) +{ + KERNEL_PREAMBLE(uchar) + dst_0[yo * FIXED_PITCH(uchar) + xo] = + tex2D(src_tex_0, src_x + 0.5f, src_y + 0.5f); +} + +__global__ void Transpose_ushort(TRANSPOSE_KERNEL_ARGS(ushort)) +{ + KERNEL_PREAMBLE(ushort) + dst_0[yo * FIXED_PITCH(ushort) + xo] = + tex2D(src_tex_0, src_x + 0.5f, src_y + 0.5f); +} + +__global__ void Transpose_uchar4(TRANSPOSE_KERNEL_ARGS(uchar4)) +{ + KERNEL_PREAMBLE(uchar4) + dst_0[yo * FIXED_PITCH(uchar4) + xo] = + tex2D(src_tex_0, src_x + 0.5f, src_y + 0.5f); +} + +__global__ void Transpose_uchar_uv(TRANSPOSE_KERNEL_ARGS(uchar)) +{ + KERNEL_PREAMBLE(uchar) + int pitch = FIXED_PITCH(uchar); + dst_1[yo * pitch + xo] = tex2D(src_tex_1, src_x + 0.5f, src_y + 0.5f); + dst_2[yo * pitch + xo] = tex2D(src_tex_2, src_x + 0.5f, src_y + 0.5f); +} + +__global__ void Transpose_ushort_uv(TRANSPOSE_KERNEL_ARGS(ushort)) +{ + KERNEL_PREAMBLE(ushort) + int pitch = FIXED_PITCH(ushort); + dst_1[yo * pitch + xo] = tex2D(src_tex_1, src_x + 0.5f, src_y + 0.5f); + dst_2[yo * pitch + xo] = tex2D(src_tex_2, src_x + 0.5f, src_y + 0.5f); +} + +__global__ void Transpose_uchar2(TRANSPOSE_KERNEL_ARGS(uchar2)) +{ + KERNEL_PREAMBLE(uchar2) + dst_1[yo * FIXED_PITCH(uchar2) + xo] = + tex2D(src_tex_1, src_x + 0.5f, src_y + 0.5f); +} + +__global__ void Transpose_ushort2(TRANSPOSE_KERNEL_ARGS(ushort2)) +{ + KERNEL_PREAMBLE(ushort2) + dst_1[yo * FIXED_PITCH(ushort2) + xo] = + tex2D(src_tex_1, src_x + 0.5f, src_y + 0.5f); +} + +} -- 2.34.1 _______________________________________________ ffmpeg-devel mailing list -- ffmpeg-devel@ffmpeg.org To unsubscribe send an email to ffmpeg-devel-leave@ffmpeg.org