[FFmpeg-devel] [PATCH v2] avfilter: add CUDA-accelerated transpose filter

[Faeez Kadiri via ffmpeg-devel <ffmpeg-devel@ffmpeg.org>]

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 <f1k2faeez@gmail.com>
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 <next>:
 - 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<uchar>(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<ushort>(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<uchar4>(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<uchar>(src_tex_1, src_x + 0.5f, src_y + 0.5f);
+    dst_2[yo * pitch + xo] = tex2D<uchar>(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<ushort>(src_tex_1, src_x + 0.5f, src_y + 0.5f);
+    dst_2[yo * pitch + xo] = tex2D<ushort>(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<uchar2>(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<ushort2>(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