PR #24728 opened by rnissl URL: https://code.ffmpeg.org/FFmpeg/FFmpeg/pulls/24728 Patch URL: https://code.ffmpeg.org/FFmpeg/FFmpeg/pulls/24728.patch
This is 3/9 in the GoPro Max review order. It applies on current master by itself. Two commits. The first adds ~1 ulp sinf/cosf/atanf polynomials so a cuda-llvm build matches nvcc at the cube-map seam. The second adds the filter: stitch the two GoPro Max 2 .360 EAC tracks to equirectangular on CUDA frames. Depends on ffnvcodec and cuda-nvcc or cuda-llvm. A stitched frame from this filter: https://www.mapillary.com/app/user/rnissl?lat=49.34297320000002&lng=12.365471200000002&z=17&x=0.44596925515485036&y=0.5035377617780364&zoom=0&pKey=1267299538827145&focus=photo Preferred review order, each request on its own: 1/9 avutil/log: flush stderr and match Windows warning yellow 2/9 avfilter: add yuvrgb_cuda 3/9 avfilter: add gopromax2_cuda GoPro Max 2 EAC stitch 4/9 avcodec: add mjpeg_nvjpeg NVIDIA JPEG encoder 5/9 avcodec/bsf: add jpegapp to inject JPEG APP segments 6/9 avfilter: add gopromax2 GoPro Max 2 EAC stitch 7/9 avfilter: add gopromax2_opencl GoPro Max 2 EAC stitch 8/9 avfilter: add gopromax_opencl GoPro Max 1 stitch 9/9 avcodec/mjpegenc: encode packed RGB via YUV conversion After 1-5, mapillary_tools processes GoPro Max 2 files GPU-only in my setup: scale_cuda=yuv444p10, yuvrgb_cuda=rgb0, gopromax2_cuda, mjpeg_nvjpeg. jpegapp inserts the APP1 segment while encoding, so the JPEG ffmpeg writes needs no postprocessing. After 6, the Max 2 CPU fallback works. After 7, Max 2 via OpenCL works. After 8, original GoPro Max (Max 1) via OpenCL works. 9/9 is needed when 4/9 cannot be used, for example because no nvjpeg library is present, and encoding falls back to the software mjpeg encoder. That encoder then receives packed RGB (rgb24). Upstream mjpegenc does not accept it; this change converts it to YUV first. The six patches I sent to ffmpeg-devel on 20 September are obsolete. Once the filter requests are merged, I will send one short documentation request that only adds the cross-references between these sections. From 143040699c1a201444a65a455b884bc5d9fbd045 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Reinhard=20Ni=C3=9Fl?= <[email protected]> Date: Fri, 25 Sep 2026 23:46:08 +0200 Subject: [PATCH 1/2] compat/cuda: polynomial sinf/cosf/atanf for cuda-llvm MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit clang -nocudalib has no libdevice. Aliasing sinf/cosf to the nvvm approximations moves the GoPro Max 2 cube-map seam by more than 1 LSB versus nvcc. Provide ~1 ulp polynomials instead. Co-authored-by: Cursor <[email protected]> Signed-off-by: Reinhard Nißl <[email protected]> --- compat/cuda/cuda_runtime.h | 91 ++++++++++++++++++++++++++++++++++++++ 1 file changed, 91 insertions(+) diff --git a/compat/cuda/cuda_runtime.h b/compat/cuda/cuda_runtime.h index d08511f72a..b717ed31c9 100644 --- a/compat/cuda/cuda_runtime.h +++ b/compat/cuda/cuda_runtime.h @@ -2,6 +2,7 @@ * Minimum CUDA compatibility definitions header * * Copyright (c) 2019 rcombs + * Copyright (c) 2026 Reinhard Nißl * * This file is part of FFmpeg. * @@ -190,6 +191,96 @@ static inline __device__ float __cosf(float a) { return __nvvm_cos_approx_f(a); static inline __device__ float __expf(float a) { return __nvvm_ex2_approx_f(a * (float)__builtin_log2(__builtin_exp(1))); } static inline __device__ float __powf(float a, float b) { return __nvvm_ex2_approx_f(__nvvm_lg2_approx_f(a) * b); } +static inline __device__ float fminf(float a, float b) { return min(a, b); } +static inline __device__ float fmaxf(float a, float b) { return max(a, b); } + +/* clang -nocudalib has no libdevice. Do not alias sinf/cosf to __sinf/__cosf + * (nvvm approx): cube-map seams in gopromax2_cuda need ~1 ulp. */ + +static inline __device__ float cuda_sin_poly(float x) +{ + float z = x * x; + return x * (1.0f + z * (-1.6666667163e-1f + + z * (8.3333473562e-3f + + z * (-1.9842602985e-4f + + z * 2.7600128038e-6f)))); +} + +static inline __device__ float cuda_cos_poly(float x) +{ + float z = x * x; + return 1.0f + z * (-0.5f + + z * (4.1666641831e-2f + + z * (-1.3888330439e-3f + + z * 2.4374417526e-5f))); +} + +static inline __device__ void cuda_sincos_reduce(float x, float *y, int *q) +{ + const float inv2pi = 0.63661977236758134308f; + const float pio2_1 = 1.57079631090164184570f; + const float pio2_2 = 1.58932547735281906091e-8f; + float n = rintf(x * inv2pi); + *q = (int)n; + *y = (x - n * pio2_1) - n * pio2_2; +} + +static inline __device__ float sinf(float x) +{ + float y; + int q; + cuda_sincos_reduce(x, &y, &q); + q &= 3; + if (q == 0) + return cuda_sin_poly(y); + if (q == 1) + return cuda_cos_poly(y); + if (q == 2) + return -cuda_sin_poly(y); + return -cuda_cos_poly(y); +} + +static inline __device__ float cosf(float x) +{ + float y; + int q; + cuda_sincos_reduce(x, &y, &q); + q &= 3; + if (q == 0) + return cuda_cos_poly(y); + if (q == 1) + return -cuda_sin_poly(y); + if (q == 2) + return -cuda_cos_poly(y); + return cuda_sin_poly(y); +} + +static inline __device__ float atanf(float x) +{ + const float pio2 = 1.57079632679489661923f; + const float pio4 = 0.78539816339744830962f; + float ax = fabsf(x); + int big = ax > 1.0f; + float t = big ? (1.0f / ax) : ax; + int hi = t > 0.41421356237f; + if (hi) + t = (t - 1.0f) / (t + 1.0f); + { + float z = t * t; + t = t * (1.0f + z * (-3.333329916e-1f + + z * (1.999868423e-1f + + z * (-1.425129980e-1f + + z * (1.064911537e-1f + + z * (-7.508206323e-2f + + z * 4.285914557e-2f)))))); + } + if (hi) + t += pio4; + if (big) + t = pio2 - t; + return (x < 0.0f) ? -t : t; +} + // Misc helper functions extern "C" __device__ int printf(const char*, ...); -- 2.52.0 From a979cd537706e79a6d2154d4ab0bf5ce35bb95af Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Reinhard=20Ni=C3=9Fl?= <[email protected]> Date: Sat, 26 Sep 2026 21:35:08 +0200 Subject: [PATCH 2/2] avfilter: add gopromax2_cuda GoPro Max 2 EAC stitch MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Stitch the two GoPro Max 2 .360 EAC tracks to equirectangular on CUDA frames. Depends on ffnvcodec and cuda-nvcc or cuda-llvm. The llvm build uses the sinf/cosf/atanf stubs from the previous commit. A stitched frame from this filter: https://www.mapillary.com/app/user/rnissl?lat=49.34297320000002&lng=12.365471200000002&z=17&x=0.44596925515485036&y=0.5035377617780364&zoom=0&pKey=1267299538827145&focus=photo Co-authored-by: Cursor <[email protected]> Signed-off-by: Reinhard Nißl <[email protected]> --- Changelog | 2 + configure | 2 + doc/filters.texi | 40 +++ libavfilter/Makefile | 2 + libavfilter/allfilters.c | 1 + libavfilter/vf_gopromax2_cuda.c | 511 +++++++++++++++++++++++++++++++ libavfilter/vf_gopromax2_cuda.cu | 216 +++++++++++++ 7 files changed, 774 insertions(+) create mode 100644 libavfilter/vf_gopromax2_cuda.c create mode 100644 libavfilter/vf_gopromax2_cuda.cu diff --git a/Changelog b/Changelog index 1cfa9db1b6..588e1b2045 100644 --- a/Changelog +++ b/Changelog @@ -16,6 +16,8 @@ version <next>: - DSD (dsd_msbf) encoder +- gopromax2_cuda GoPro Max 2 CUDA stitch filter + version 9.0: - Extend AMF Color Converter (vf_vpp_amf) HDR capabilities - LCEVC track muxing support in MP4 muxer diff --git a/configure b/configure index f686b727cf..afe307fa09 100755 --- a/configure +++ b/configure @@ -3552,6 +3552,8 @@ colorspace_cuda_filter_deps_any="cuda_nvcc cuda_llvm" hwupload_cuda_filter_deps="ffnvcodec" scale_cuda_filter_deps="ffnvcodec" scale_cuda_filter_deps_any="cuda_nvcc cuda_llvm" +gopromax2_cuda_filter_deps="ffnvcodec" +gopromax2_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" diff --git a/doc/filters.texi b/doc/filters.texi index 1a9bbeff7e..87f80af725 100644 --- a/doc/filters.texi +++ b/doc/filters.texi @@ -27877,6 +27877,46 @@ Thumbnails are extracted from every @var{n}=150-frame batch, selecting one per b @end itemize +@anchor{gopromax2_cuda} +@section gopromax2_cuda + +Convert the two GoPro Max 2 .360 EAC tracks to equirectangular projection +on CUDA frames. + +Each input track is @code{side} + @code{center} + @code{side} (GoPro Max 2: +5952x1920 = 2016 + 1920 + 2016). The duplicated overlap lives in the side +faces and is blended over @code{blend} pixels (default 48). Rear faces are +rotated 90 degrees in UV and V-inverted. Typical mapping is front @code{0:0} +and rear @code{0:4}. Default output size is 7680x3840. + +@table @option +@item w +@item h +Output equirectangular size. Must be 2:1. +@item side +Width of each side face in luma pixels. @code{0} is @code{(iw-center)/2}. +@item center +Width of the center face. @code{0} is the input height. +@item blend +Overlap blend width in luma pixels. Default 48. +@item yaw +Yaw rotation of the output equirect in degrees. Default 0. +@end table + +Inputs must already be CUDA frames. Accepted software formats are +@code{yuv444p}, @code{yuv444p10}, @code{yuv444p12}, @code{yuv444p16} and +packed RGB (@code{rgb0}, @code{bgr0}, @code{rgba}, @code{bgra}). +@ref{scale_cuda} can produce @code{yuv444p10}. It has no planar 10-bit to +@code{rgb0} path. + +@example +[0:0]scale_cuda=format=yuv444p10[a]; +[0:4]scale_cuda=format=yuv444p10[b]; +[a][b]gopromax2_cuda=yaw=-90 +@end example + +Encode on the GPU with @code{-c:v mjpeg_nvjpeg -quality 83 -chroma 420}. + @section transpose_cuda Transpose rows with columns in the input video and optionally flip it. diff --git a/libavfilter/Makefile b/libavfilter/Makefile index ceea328e67..52986afd43 100644 --- a/libavfilter/Makefile +++ b/libavfilter/Makefile @@ -530,6 +530,8 @@ OBJS-$(CONFIG_TELECINE_FILTER) += vf_telecine.o OBJS-$(CONFIG_THISTOGRAM_FILTER) += vf_histogram.o OBJS-$(CONFIG_THRESHOLD_FILTER) += vf_threshold.o framesync.o OBJS-$(CONFIG_THUMBNAIL_FILTER) += vf_thumbnail.o +OBJS-$(CONFIG_GOPROMAX2_CUDA_FILTER) += vf_gopromax2_cuda.o vf_gopromax2_cuda.ptx.o \ + framesync.o cuda/load_helper.o OBJS-$(CONFIG_THUMBNAIL_CUDA_FILTER) += vf_thumbnail_cuda.o vf_thumbnail_cuda.ptx.o \ cuda/load_helper.o OBJS-$(CONFIG_TILE_FILTER) += vf_tile.o diff --git a/libavfilter/allfilters.c b/libavfilter/allfilters.c index 3b93fe12f6..2da270f279 100644 --- a/libavfilter/allfilters.c +++ b/libavfilter/allfilters.c @@ -501,6 +501,7 @@ extern const FFFilter ff_vf_telecine; extern const FFFilter ff_vf_thistogram; extern const FFFilter ff_vf_threshold; extern const FFFilter ff_vf_thumbnail; +extern const FFFilter ff_vf_gopromax2_cuda; extern const FFFilter ff_vf_thumbnail_cuda; extern const FFFilter ff_vf_tile; extern const FFFilter ff_vf_tiltandshift; diff --git a/libavfilter/vf_gopromax2_cuda.c b/libavfilter/vf_gopromax2_cuda.c new file mode 100644 index 0000000000..b9ad74e6ad --- /dev/null +++ b/libavfilter/vf_gopromax2_cuda.c @@ -0,0 +1,511 @@ +/* + * GoPro Max 2 .360 to equirectangular (CUDA). + * Dual-EAC side/center/side stitch to 2:1 equirectangular. + * + * Copyright (c) 2021 Ronan LE MEILLAT + * Copyright (c) 2024 TADANO Tokumei + * Copyright (c) 2026 Reinhard Nißl + * + * Based on gopromax_opencl (Max 1) by Ronan LE MEILLAT and TADANO Tokumei. + * + * 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 <math.h> +#include <string.h> + +#include "libavutil/cuda_check.h" +#include "libavutil/hwcontext.h" +#include "libavutil/hwcontext_cuda_internal.h" +#include "libavutil/opt.h" +#include "libavutil/pixdesc.h" + +#ifndef M_PI +#define M_PI 3.14159265358979323846 +#endif + +#include "avfilter.h" +#include "cuda/load_helper.h" +#include "filters.h" +#include "framesync.h" +#include "video.h" + +#define CHECK_CU(x) FF_CUDA_CHECK_DL(avctx, s->hwctx->internal->cuda_dl, x) +#define DIV_UP(a, b) (((a) + (b) - 1) / (b)) +#define BLOCKX 32 +#define BLOCKY 16 + +static const enum AVPixelFormat supported_formats[] = { + AV_PIX_FMT_YUV444P, + AV_PIX_FMT_YUV444P10, + AV_PIX_FMT_YUV444P12, + AV_PIX_FMT_YUV444P16, + AV_PIX_FMT_RGB0, + AV_PIX_FMT_BGR0, + AV_PIX_FMT_RGBA, + AV_PIX_FMT_BGRA, + AV_PIX_FMT_NONE, +}; + +typedef struct GoProMax2CUDAContext { + const AVClass *class; + + AVCUDADeviceContext *hwctx; + AVBufferRef *hw_frames_ctx; + + CUmodule cu_module; + CUfunction cu_plane_u8; + CUfunction cu_plane_u16; + CUfunction cu_packed_u8; + CUstream cu_stream; + + FFFrameSync fs; + + int width, height; + int side, center, blend; + float yaw; + + int in_w, in_h; + int layout_side, layout_center, layout_blend; + int packed; + int bytes_ps; + int nb_planes; + enum AVPixelFormat sw_format; +} GoProMax2CUDAContext; + +static int format_is_supported(enum AVPixelFormat fmt) +{ + int i; + for (i = 0; supported_formats[i] != AV_PIX_FMT_NONE; i++) + if (supported_formats[i] == fmt) + return 1; + return 0; +} + +static av_cold void gopromax2_cuda_uninit(AVFilterContext *avctx) +{ + GoProMax2CUDAContext *s = avctx->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->hw_frames_ctx); + ff_framesync_uninit(&s->fs); +} + +static int gopromax2_cuda_load(AVFilterContext *avctx) +{ + GoProMax2CUDAContext *s = avctx->priv; + CudaFunctions *cu = s->hwctx->internal->cuda_dl; + CUcontext dummy, cuda_ctx = s->hwctx->cuda_ctx; + int ret; + + extern const unsigned char ff_vf_gopromax2_cuda_ptx_data[]; + extern const unsigned int ff_vf_gopromax2_cuda_ptx_len; + + ret = CHECK_CU(cu->cuCtxPushCurrent(cuda_ctx)); + if (ret < 0) + return ret; + + ret = ff_cuda_load_module(avctx, s->hwctx, &s->cu_module, + ff_vf_gopromax2_cuda_ptx_data, + ff_vf_gopromax2_cuda_ptx_len); + if (ret < 0) + goto fail; + + ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_plane_u8, s->cu_module, + "gopromax2_plane_u8")); + if (ret < 0) + goto fail; + ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_plane_u16, s->cu_module, + "gopromax2_plane_u16")); + if (ret < 0) + goto fail; + ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_packed_u8, s->cu_module, + "gopromax2_packed_u8")); + if (ret < 0) + goto fail; + +fail: + CHECK_CU(cu->cuCtxPopCurrent(&dummy)); + return ret; +} + +static int create_tex(AVFilterContext *avctx, CUtexObject *tex, + CUdeviceptr ptr, int linesize, int width, int height, + CUarray_format fmt, int channels) +{ + GoProMax2CUDAContext *s = avctx->priv; + CudaFunctions *cu = s->hwctx->internal->cuda_dl; + CUDA_RESOURCE_DESC res_desc; + CUDA_TEXTURE_DESC tex_desc; + + memset(&res_desc, 0, sizeof(res_desc)); + memset(&tex_desc, 0, sizeof(tex_desc)); + + res_desc.resType = CU_RESOURCE_TYPE_PITCH2D; + res_desc.res.pitch2D.format = fmt; + res_desc.res.pitch2D.numChannels = channels; + res_desc.res.pitch2D.width = width; + res_desc.res.pitch2D.height = height; + res_desc.res.pitch2D.pitchInBytes = linesize; + res_desc.res.pitch2D.devPtr = ptr; + + tex_desc.addressMode[0] = CU_TR_ADDRESS_MODE_CLAMP; + tex_desc.addressMode[1] = CU_TR_ADDRESS_MODE_CLAMP; + tex_desc.addressMode[2] = CU_TR_ADDRESS_MODE_CLAMP; + tex_desc.filterMode = CU_TR_FILTER_MODE_LINEAR; + + return CHECK_CU(cu->cuTexObjectCreate(tex, &res_desc, &tex_desc, NULL)); +} + +static int launch_kernel(AVFilterContext *avctx, CUfunction func, + CUtexObject front, CUtexObject rear, + CUdeviceptr dst, int dst_pitch, + int dst_w, int dst_h, int src_w, int src_h, + int side, int center, int blend, float yaw) +{ + GoProMax2CUDAContext *s = avctx->priv; + CudaFunctions *cu = s->hwctx->internal->cuda_dl; + + void *args[] = { + &front, &rear, &dst, &dst_pitch, + &dst_w, &dst_h, &src_w, &src_h, + &side, ¢er, &blend, &yaw + }; + + return CHECK_CU(cu->cuLaunchKernel(func, + DIV_UP(dst_w, BLOCKX), DIV_UP(dst_h, BLOCKY), 1, + BLOCKX, BLOCKY, 1, 0, s->cu_stream, args, NULL)); +} + +static int gopromax2_cuda_filter(FFFrameSync *fs) +{ + AVFilterContext *avctx = fs->parent; + AVFilterLink *outlink = avctx->outputs[0]; + GoProMax2CUDAContext *s = avctx->priv; + CudaFunctions *cu; + AVFrame *front, *rear, *output = NULL; + CUcontext dummy; + CUtexObject tex_f = 0, tex_r = 0; + float yaw; + int ret, plane, pushed = 0; + + ret = ff_framesync_get_frame(fs, 0, &front, 0); + if (ret < 0) + return ret; + ret = ff_framesync_get_frame(fs, 1, &rear, 0); + if (ret < 0) + return ret; + + cu = s->hwctx->internal->cuda_dl; + yaw = s->yaw * (float)M_PI / 180.f; + + output = ff_get_video_buffer(outlink, outlink->w, outlink->h); + if (!output) + return AVERROR(ENOMEM); + + ret = CHECK_CU(cu->cuCtxPushCurrent(s->hwctx->cuda_ctx)); + if (ret < 0) + goto fail; + pushed = 1; + + if (s->packed) { + ret = create_tex(avctx, &tex_f, (CUdeviceptr)front->data[0], + front->linesize[0], s->in_w, s->in_h, + CU_AD_FORMAT_UNSIGNED_INT8, 4); + if (ret < 0) + goto fail; + ret = create_tex(avctx, &tex_r, (CUdeviceptr)rear->data[0], + rear->linesize[0], s->in_w, s->in_h, + CU_AD_FORMAT_UNSIGNED_INT8, 4); + if (ret < 0) + goto fail; + + ret = launch_kernel(avctx, s->cu_packed_u8, tex_f, tex_r, + (CUdeviceptr)output->data[0], output->linesize[0], + outlink->w, outlink->h, s->in_w, s->in_h, + s->layout_side, s->layout_center, s->layout_blend, yaw); + if (ret < 0) + goto fail; + } else { + CUarray_format fmt = s->bytes_ps == 2 ? CU_AD_FORMAT_UNSIGNED_INT16 + : CU_AD_FORMAT_UNSIGNED_INT8; + CUfunction func = s->bytes_ps == 2 ? s->cu_plane_u16 : s->cu_plane_u8; + + for (plane = 0; plane < s->nb_planes; plane++) { + int src_w = s->in_w; + int src_h = s->in_h; + int dst_w = outlink->w; + int dst_h = outlink->h; + int dst_pitch = output->linesize[plane] / s->bytes_ps; + + if (tex_f) { + CHECK_CU(cu->cuTexObjectDestroy(tex_f)); + tex_f = 0; + } + if (tex_r) { + CHECK_CU(cu->cuTexObjectDestroy(tex_r)); + tex_r = 0; + } + + ret = create_tex(avctx, &tex_f, (CUdeviceptr)front->data[plane], + front->linesize[plane], src_w, src_h, fmt, 1); + if (ret < 0) + goto fail; + ret = create_tex(avctx, &tex_r, (CUdeviceptr)rear->data[plane], + rear->linesize[plane], src_w, src_h, fmt, 1); + if (ret < 0) + goto fail; + + ret = launch_kernel(avctx, func, tex_f, tex_r, + (CUdeviceptr)output->data[plane], dst_pitch, + dst_w, dst_h, src_w, src_h, + s->layout_side, s->layout_center, + s->layout_blend, yaw); + if (ret < 0) + goto fail; + } + } + + if (tex_f) + CHECK_CU(cu->cuTexObjectDestroy(tex_f)); + tex_f = 0; + if (tex_r) + CHECK_CU(cu->cuTexObjectDestroy(tex_r)); + tex_r = 0; + + ret = CHECK_CU(cu->cuCtxPopCurrent(&dummy)); + if (ret < 0) + goto fail; + pushed = 0; + + ret = av_frame_copy_props(output, front); + if (ret < 0) + goto fail; + + return ff_filter_frame(outlink, output); + +fail: + if (tex_f) + CHECK_CU(cu->cuTexObjectDestroy(tex_f)); + if (tex_r) + CHECK_CU(cu->cuTexObjectDestroy(tex_r)); + if (pushed) + CHECK_CU(cu->cuCtxPopCurrent(&dummy)); + av_frame_free(&output); + return ret; +} + +static int init_hwframe_ctx(AVFilterContext *avctx, AVBufferRef *device_ref, + int width, int height) +{ + GoProMax2CUDAContext *s = avctx->priv; + AVBufferRef *out_ref; + AVHWFramesContext *out_ctx; + int ret; + + out_ref = av_hwframe_ctx_alloc(device_ref); + if (!out_ref) + return AVERROR(ENOMEM); + out_ctx = (AVHWFramesContext *)out_ref->data; + + out_ctx->format = AV_PIX_FMT_CUDA; + out_ctx->sw_format = s->sw_format; + out_ctx->width = FFALIGN(width, 32); + out_ctx->height = FFALIGN(height, 32); + + ret = av_hwframe_ctx_init(out_ref); + if (ret < 0) { + av_buffer_unref(&out_ref); + return ret; + } + + av_buffer_unref(&s->hw_frames_ctx); + s->hw_frames_ctx = out_ref; + return 0; +} + +static int gopromax2_cuda_config_output(AVFilterLink *outlink) +{ + AVFilterContext *avctx = outlink->src; + GoProMax2CUDAContext *s = avctx->priv; + FilterLink *inl = ff_filter_link(avctx->inputs[0]); + FilterLink *inl1 = ff_filter_link(avctx->inputs[1]); + FilterLink *outl = ff_filter_link(outlink); + AVHWFramesContext *in_fc, *in_fc1; + const AVPixFmtDescriptor *desc; + int in_w, in_h, out_w, out_h, ret, i; + + if (!inl->hw_frames_ctx || !inl1->hw_frames_ctx) { + av_log(avctx, AV_LOG_ERROR, "CUDA frames context required on both inputs\n"); + return AVERROR(EINVAL); + } + + in_fc = (AVHWFramesContext *)inl->hw_frames_ctx->data; + in_fc1 = (AVHWFramesContext *)inl1->hw_frames_ctx->data; + + if (in_fc->sw_format != in_fc1->sw_format || + avctx->inputs[0]->w != avctx->inputs[1]->w || + avctx->inputs[0]->h != avctx->inputs[1]->h) { + av_log(avctx, AV_LOG_ERROR, "Front and rear must match in size and format\n"); + return AVERROR(EINVAL); + } + + if (!format_is_supported(in_fc->sw_format)) { + av_log(avctx, AV_LOG_ERROR, + "Unsupported CUDA sw_format %s. Convert on GPU first, e.g. " + "scale_cuda=format=yuv444p10 or scale_cuda=format=rgb0.\n", + av_get_pix_fmt_name(in_fc->sw_format)); + return AVERROR(ENOSYS); + } + + in_w = avctx->inputs[0]->w; + in_h = avctx->inputs[0]->h; + s->in_w = in_w; + s->in_h = in_h; + s->sw_format = in_fc->sw_format; + s->layout_center = s->center > 0 ? s->center : in_h; + s->layout_side = s->side > 0 ? s->side : (in_w - s->layout_center) / 2; + s->layout_blend = s->blend > 0 ? s->blend : 48; + + if (s->layout_side <= 0 || s->layout_center <= 0 || + 2 * s->layout_side + s->layout_center != in_w) { + av_log(avctx, AV_LOG_ERROR, + "Max 2 layout mismatch: input %dx%d, side %d center %d " + "(need 2*side+center = width).\n", + in_w, in_h, s->layout_side, s->layout_center); + return AVERROR(EINVAL); + } + + if (s->width > 0 && s->height > 0) { + if (s->width != s->height * 2) { + av_log(avctx, AV_LOG_ERROR, "Specified size (%dx%d) is not 2:1 equirect.\n", + s->width, s->height); + return AVERROR(EINVAL); + } + out_w = s->width; + out_h = s->height; + } else if (s->width > 0 || s->height > 0) { + av_log(avctx, AV_LOG_ERROR, "Both width and height values should be specified.\n"); + return AVERROR(EINVAL); + } else { + out_w = 4 * in_h; + out_h = 2 * in_h; + } + + desc = av_pix_fmt_desc_get(s->sw_format); + s->nb_planes = av_pix_fmt_count_planes(s->sw_format); + s->packed = (s->nb_planes == 1 && desc->nb_components >= 3); + s->bytes_ps = 1; + for (i = 0; i < desc->nb_components; i++) + s->bytes_ps = FFMAX(s->bytes_ps, (desc->comp[i].depth + 7) / 8); + + s->hwctx = ((AVHWDeviceContext *)in_fc->device_ref->data)->hwctx; + s->cu_stream = s->hwctx->stream; + + ret = init_hwframe_ctx(avctx, in_fc->device_ref, out_w, out_h); + if (ret < 0) + return ret; + + outlink->w = out_w; + outlink->h = out_h; + outlink->sample_aspect_ratio = (AVRational){ 1, 1 }; + outl->hw_frames_ctx = av_buffer_ref(s->hw_frames_ctx); + if (!outl->hw_frames_ctx) + return AVERROR(ENOMEM); + + av_log(avctx, AV_LOG_INFO, + "gopromax2_cuda %s in %dx%d side %d center %d blend %d yaw %.1f out %dx%d\n", + av_get_pix_fmt_name(s->sw_format), in_w, in_h, + s->layout_side, s->layout_center, s->layout_blend, s->yaw, out_w, out_h); + + ret = gopromax2_cuda_load(avctx); + if (ret < 0) + return ret; + + ret = ff_framesync_init_dualinput(&s->fs, avctx); + if (ret < 0) + return ret; + + return ff_framesync_configure(&s->fs); +} + +static av_cold int gopromax2_cuda_init(AVFilterContext *avctx) +{ + GoProMax2CUDAContext *s = avctx->priv; + s->fs.on_event = &gopromax2_cuda_filter; + return 0; +} + +static int gopromax2_cuda_activate(AVFilterContext *avctx) +{ + GoProMax2CUDAContext *s = avctx->priv; + return ff_framesync_activate(&s->fs); +} + +#define OFFSET(x) offsetof(GoProMax2CUDAContext, x) +#define FLAGS (AV_OPT_FLAG_FILTERING_PARAM | AV_OPT_FLAG_VIDEO_PARAM) +static const AVOption gopromax2_cuda_options[] = { + { "w", "output width", OFFSET(width), AV_OPT_TYPE_INT, {.i64=0}, 0, INT16_MAX, FLAGS }, + { "h", "output height", OFFSET(height), AV_OPT_TYPE_INT, {.i64=0}, 0, INT16_MAX, FLAGS }, + { "side", "side face width in luma pixels (0 = (in_w-center)/2)", + OFFSET(side), AV_OPT_TYPE_INT, {.i64=0}, 0, INT16_MAX, FLAGS }, + { "center", "center face width in luma pixels (0 = in_h)", + OFFSET(center), AV_OPT_TYPE_INT, {.i64=0}, 0, INT16_MAX, FLAGS }, + { "blend", "overlap blend width in luma pixels", + OFFSET(blend), AV_OPT_TYPE_INT, {.i64=48}, 1, 1024, FLAGS }, + { "yaw", "yaw rotation of the equirect in degrees", + OFFSET(yaw), AV_OPT_TYPE_FLOAT, {.dbl=0}, -360, 360, FLAGS }, + { NULL }, +}; + +AVFILTER_DEFINE_CLASS(gopromax2_cuda); + +static const AVFilterPad gopromax2_cuda_inputs[] = { + { .name = "front", .type = AVMEDIA_TYPE_VIDEO }, + { .name = "rear", .type = AVMEDIA_TYPE_VIDEO }, +}; + +static const AVFilterPad gopromax2_cuda_outputs[] = { + { + .name = "default", + .type = AVMEDIA_TYPE_VIDEO, + .config_props = &gopromax2_cuda_config_output, + }, +}; + +const FFFilter ff_vf_gopromax2_cuda = { + .p.name = "gopromax2_cuda", + .p.description = NULL_IF_CONFIG_SMALL("GoPro Max 2 .360 to equirectangular projection (CUDA)"), + .p.priv_class = &gopromax2_cuda_class, + .priv_size = sizeof(GoProMax2CUDAContext), + .init = &gopromax2_cuda_init, + .uninit = &gopromax2_cuda_uninit, + .activate = &gopromax2_cuda_activate, + FILTER_INPUTS(gopromax2_cuda_inputs), + FILTER_OUTPUTS(gopromax2_cuda_outputs), + FILTER_SINGLE_PIXFMT(AV_PIX_FMT_CUDA), + .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE, +}; diff --git a/libavfilter/vf_gopromax2_cuda.cu b/libavfilter/vf_gopromax2_cuda.cu new file mode 100644 index 0000000000..0c757c785c --- /dev/null +++ b/libavfilter/vf_gopromax2_cuda.cu @@ -0,0 +1,216 @@ +/* + * GoPro Max 2 EAC -> equirect (CUDA). + * + * Copyright (c) 2021 Ronan LE MEILLAT + * Copyright (c) 2024 TADANO Tokumei + * Copyright (c) 2026 Reinhard Nißl + * + * Based on gopromax_opencl (Max 1) by Ronan LE MEILLAT and TADANO Tokumei. + * + * 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" + +enum { + FACE_LEFT, + FACE_RIGHT, + FACE_FRONT, + FACE_BACK, + FACE_TOP, + FACE_DOWN +}; + +#define GOPROMAX2_PI 3.14159265358979323846f +#define GOPROMAX2_PI2 1.57079632679489661923f + +template<typename TexT> +__device__ static inline TexT sample_img(cudaTextureObject_t im, float x, float y, + int w, int h) +{ + x = fminf(fmaxf(x, 0.f), (float)(w - 1)); + y = fminf(fmaxf(y, 0.f), (float)(h - 1)); + return tex2D<TexT>(im, x + 0.5f, y + 0.5f); +} + +__device__ static inline int find_face(float lon, float lat, float *u, float *v) +{ + float cl = cosf(lat); + float px = cl * sinf(lon); + float py = cl * cosf(lon); + float pz = sinf(lat); + float ax = fabsf(px), ay = fabsf(py), az = fabsf(pz); + int face = FACE_FRONT; + + if (ax >= ay && ax >= az) + face = px < 0.f ? FACE_LEFT : FACE_RIGHT; + else if (ay > ax && ay >= az) + face = py >= 0.f ? FACE_FRONT : FACE_BACK; + else if (az > ax && az > ay) + face = pz >= 0.f ? FACE_TOP : FACE_DOWN; + + float dom = ax; + if (face == FACE_FRONT || face == FACE_BACK) + dom = ay; + else if (face == FACE_TOP || face == FACE_DOWN) + dom = az; + if (dom < 1e-12f) + dom = 1e-12f; + + float k = 4.f / GOPROMAX2_PI; + float qx = atanf(px / dom) * k; + float qy = atanf(py / dom) * k; + float qz = atanf(pz / dom) * k; + + switch (face) { + case FACE_LEFT: *u = (qy + 1.f) * 0.5f; *v = (qz + 1.f) * 0.5f; break; + case FACE_RIGHT: *u = (1.f - qy) * 0.5f; *v = (qz + 1.f) * 0.5f; break; + case FACE_FRONT: *u = (qx + 1.f) * 0.5f; *v = (qz + 1.f) * 0.5f; break; + case FACE_BACK: *u = (1.f - qx) * 0.5f; *v = (qz + 1.f) * 0.5f; break; + case FACE_TOP: *u = (1.f - qx) * 0.5f; *v = (qy + 1.f) * 0.5f; break; + default: *u = (1.f - qx) * 0.5f; *v = (1.f - qy) * 0.5f; break; + } + *u = fminf(fmaxf(*u, 0.f), 1.f - 1e-6f); + *v = fminf(fmaxf(*v, 0.f), 1.f - 1e-6f); + return face; +} + +template<typename TexT> +__device__ static inline TexT get_colour(int face, float u, float v, + cudaTextureObject_t front, + cudaTextureObject_t rear, + int src_w, int src_h, + int side, int center, int blendw) +{ + int use_front = (face == FACE_FRONT || face == FACE_LEFT || face == FACE_RIGHT); + cudaTextureObject_t src = use_front ? front : rear; + int x0, w; + + if (face == FACE_DOWN || face == FACE_BACK || face == FACE_TOP) { + float nu = v; + float nv = 1.f - 1e-6f - u; + u = nu; + v = nv; + } + v = 1.f - 1e-6f - v; + u = fminf(fmaxf(u, 0.f), 1.f - 1e-6f); + v = fminf(fmaxf(v, 0.f), 1.f - 1e-6f); + + if (face == FACE_FRONT || face == FACE_BACK) { + x0 = side; + w = center; + return sample_img<TexT>(src, x0 + u * (float)w, v * (float)src_h, src_w, src_h); + } + + if (face == FACE_LEFT || face == FACE_DOWN) { + x0 = 0; + w = side; + } else { + x0 = side + center; + w = side; + } + + { + float duv = (float)blendw / (float)w; + float uL = 2.f * (0.5f - duv) * u; + float uR = 2.f * (0.5f - duv) * (u - 0.5f) + 0.5f + duv; + float y = v * (float)src_h; + if (uL <= 0.5f - 2.f * duv) { + return sample_img<TexT>(src, x0 + uL * (float)w, y, src_w, src_h); + } else if (uR >= 0.5f + 2.f * duv) { + return sample_img<TexT>(src, x0 + uR * (float)w, y, src_w, src_h); + } else { + TexT c1 = sample_img<TexT>(src, x0 + uL * (float)w, y, src_w, src_h); + TexT c2 = sample_img<TexT>(src, x0 + uR * (float)w, y, src_w, src_h); + return lerp_scalar(c1, c2, (uL - 0.5f + 2.f * duv) / (2.f * duv)); + } + } +} + +__device__ static inline void equirect_uv(int x, int y, int dst_w, int dst_h, + float yaw, int *face, float *u, float *v) +{ + float lat = GOPROMAX2_PI2 - (y + 0.5f) / (float)dst_h * GOPROMAX2_PI; + float lon = (x + 0.5f) / (float)dst_w * 2.f * GOPROMAX2_PI - GOPROMAX2_PI + yaw; + if (lon > GOPROMAX2_PI) + lon -= 2.f * GOPROMAX2_PI; + if (lon < -GOPROMAX2_PI) + lon += 2.f * GOPROMAX2_PI; + *face = find_face(lon, lat, u, v); +} + +extern "C" { + +__global__ void gopromax2_plane_u8(cudaTextureObject_t front, cudaTextureObject_t rear, + unsigned char *dst, int dst_pitch, + int dst_w, int dst_h, int src_w, int src_h, + int side, int center, int blend, float yaw) +{ + int x = blockIdx.x * blockDim.x + threadIdx.x; + int y = blockIdx.y * blockDim.y + threadIdx.y; + if (x >= dst_w || y >= dst_h) + return; + + int face; + float u, v; + equirect_uv(x, y, dst_w, dst_h, yaw, &face, &u, &v); + float val = get_colour<float>(face, u, v, front, rear, src_w, src_h, + side, center, blend); + dst[y * dst_pitch + x] = (unsigned char)saturate_rintf(val, 255.f); +} + +__global__ void gopromax2_plane_u16(cudaTextureObject_t front, cudaTextureObject_t rear, + unsigned short *dst, int dst_pitch, + int dst_w, int dst_h, int src_w, int src_h, + int side, int center, int blend, float yaw) +{ + int x = blockIdx.x * blockDim.x + threadIdx.x; + int y = blockIdx.y * blockDim.y + threadIdx.y; + if (x >= dst_w || y >= dst_h) + return; + + int face; + float u, v; + equirect_uv(x, y, dst_w, dst_h, yaw, &face, &u, &v); + float val = get_colour<float>(face, u, v, front, rear, src_w, src_h, + side, center, blend); + dst[y * dst_pitch + x] = (unsigned short)saturate_rintf(val, 65535.f); +} + +__global__ void gopromax2_packed_u8(cudaTextureObject_t front, cudaTextureObject_t rear, + unsigned char *dst, int dst_pitch, + int dst_w, int dst_h, int src_w, int src_h, + int side, int center, int blend, float yaw) +{ + int x = blockIdx.x * blockDim.x + threadIdx.x; + int y = blockIdx.y * blockDim.y + threadIdx.y; + if (x >= dst_w || y >= dst_h) + return; + + int face; + float u, v; + equirect_uv(x, y, dst_w, dst_h, yaw, &face, &u, &v); + float4 val = get_colour<float4>(face, u, v, front, rear, src_w, src_h, + side, center, blend); + unsigned char *p = dst + y * dst_pitch + x * 4; + p[0] = (unsigned char)saturate_rintf(val.x, 255.f); + p[1] = (unsigned char)saturate_rintf(val.y, 255.f); + p[2] = (unsigned char)saturate_rintf(val.z, 255.f); + p[3] = (unsigned char)saturate_rintf(val.w, 255.f); +} + +} -- 2.52.0 _______________________________________________ ffmpeg-devel mailing list -- [email protected] To unsubscribe send an email to [email protected]
