PR #24724 opened by rnissl URL: https://code.ffmpeg.org/FFmpeg/FFmpeg/pulls/24724 Patch URL: https://code.ffmpeg.org/FFmpeg/FFmpeg/pulls/24724.patch
This is 7/9 in the GoPro Max review order. It applies on current master by itself. OpenCL stitch for GoPro Max 2 .360, side/center/side EAC to equirectangular. Based on gopromax_opencl (Max 1) by Ronan LE MEILLAT and TADANO Tokumei. 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 3cd42f7a107ceb5777b3294470a67a41b5231af5 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Reinhard=20Ni=C3=9Fl?= <[email protected]> Date: Sat, 26 Sep 2026 21:38:39 +0200 Subject: [PATCH] avfilter: add gopromax2_opencl GoPro Max 2 EAC stitch MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit OpenCL stitch for GoPro Max 2 .360, side/center/side EAC to equirectangular. Based on gopromax_opencl (Max 1) by Ronan LE MEILLAT and TADANO Tokumei. Co-authored-by: Cursor <[email protected]> Signed-off-by: Reinhard Nißl <[email protected]> --- Changelog | 2 + configure | 1 + doc/filters.texi | 30 +++ libavfilter/Makefile | 2 + libavfilter/allfilters.c | 1 + libavfilter/opencl/gopromax2.cl | 178 ++++++++++++++ libavfilter/opencl_source.h | 1 + libavfilter/vf_gopromax2_opencl.c | 386 ++++++++++++++++++++++++++++++ 8 files changed, 601 insertions(+) create mode 100644 libavfilter/opencl/gopromax2.cl create mode 100644 libavfilter/vf_gopromax2_opencl.c diff --git a/Changelog b/Changelog index 1cfa9db1b6..169866e83d 100644 --- a/Changelog +++ b/Changelog @@ -16,6 +16,8 @@ version <next>: - DSD (dsd_msbf) encoder +- gopromax2_opencl GoPro Max 2 OpenCL 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..1942fb9a66 100755 --- a/configure +++ b/configure @@ -4128,6 +4128,7 @@ avgblur_opencl_filter_deps="opencl" avgblur_vulkan_filter_deps="vulkan spirv_compiler" azmq_filter_deps="libzmq" blackdetect_vulkan_filter_deps="vulkan spirv_compiler" +gopromax2_opencl_filter_deps="opencl" blackframe_filter_deps="gpl" blend_vulkan_filter_deps="vulkan spirv_compiler" boxblur_filter_deps="gpl" diff --git a/doc/filters.texi b/doc/filters.texi index 1a9bbeff7e..90a2c5eb85 100644 --- a/doc/filters.texi +++ b/doc/filters.texi @@ -28432,6 +28432,36 @@ The inputs have same memory layout for color channels , the overlay has addition @end itemize +@anchor{gopromax2_opencl} +@section gopromax2_opencl + +Convert the two GoPro Max 2 .360 EAC tracks to equirectangular projection +with OpenCL. + +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 + +@example +ffmpeg -init_hw_device opencl=ocl -filter_hw_device ocl -i INPUT.360 -filter_complex "[0:0]hwupload[a];[0:4]hwupload[b];[a][b]gopromax2_opencl=yaw=-90,hwdownload,format=nv12" out.mp4 +@end example + @section pad_opencl Add paddings to the input image, and place the original input at the diff --git a/libavfilter/Makefile b/libavfilter/Makefile index ceea328e67..ab96bf25fc 100644 --- a/libavfilter/Makefile +++ b/libavfilter/Makefile @@ -338,6 +338,8 @@ OBJS-$(CONFIG_GBLUR_VULKAN_FILTER) += vf_gblur_vulkan.o vulkan.o vulka OBJS-$(CONFIG_GEQ_FILTER) += vf_geq.o OBJS-$(CONFIG_GRADFUN_FILTER) += vf_gradfun.o OBJS-$(CONFIG_GRAPHMONITOR_FILTER) += f_graphmonitor.o +OBJS-$(CONFIG_GOPROMAX2_OPENCL_FILTER) += vf_gopromax2_opencl.o opencl.o \ + opencl/gopromax2.o framesync.o OBJS-$(CONFIG_GRAYWORLD_FILTER) += vf_grayworld.o OBJS-$(CONFIG_GREYEDGE_FILTER) += vf_colorconstancy.o OBJS-$(CONFIG_GUIDED_FILTER) += vf_guided.o framesync.o diff --git a/libavfilter/allfilters.c b/libavfilter/allfilters.c index 3b93fe12f6..408a16b0b2 100644 --- a/libavfilter/allfilters.c +++ b/libavfilter/allfilters.c @@ -312,6 +312,7 @@ extern const FFFilter ff_vf_gblur_vulkan; extern const FFFilter ff_vf_geq; extern const FFFilter ff_vf_gradfun; extern const FFFilter ff_vf_graphmonitor; +extern const FFFilter ff_vf_gopromax2_opencl; extern const FFFilter ff_vf_grayworld; extern const FFFilter ff_vf_greyedge; extern const FFFilter ff_vf_guided; diff --git a/libavfilter/opencl/gopromax2.cl b/libavfilter/opencl/gopromax2.cl new file mode 100644 index 0000000000..5f96892412 --- /dev/null +++ b/libavfilter/opencl/gopromax2.cl @@ -0,0 +1,178 @@ +/* + * GoPro Max 2 EAC -> equirect (Paul Bourke max2sphere layout). + * + * Input tracks are 5952x1920: side 2016 / center 1920 / side 2016, + * overlap blend in the side faces (default 48). Rear DOWN/BACK/TOP + * get a 90° UV rotate then V-invert, matching native Max 2 .360. + * + * 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 + */ + +enum { + FACE_LEFT, + FACE_RIGHT, + FACE_FRONT, + FACE_BACK, + FACE_TOP, + FACE_DOWN +}; + +const sampler_t sampler_linear = (CLK_NORMALIZED_COORDS_FALSE | + CLK_ADDRESS_CLAMP_TO_EDGE | + CLK_FILTER_LINEAR); + +static float4 sample_img(__read_only image2d_t im, float x, float y) +{ + int2 dim = get_image_dim(im); + x = clamp(x, 0.f, (float)(dim.x - 1)); + y = clamp(y, 0.f, (float)(dim.y - 1)); + return read_imagef(im, sampler_linear, (float2)(x + 0.5f, y + 0.5f)); +} + +static int find_face(float lon, float lat, float *u, float *v) +{ + float cl = cos(lat); + float px = cl * sin(lon); + float py = cl * cos(lon); + float pz = sin(lat); + float ax = fabs(px), ay = fabs(py), az = fabs(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 / M_PI_F; + float qx = atan(px / dom) * k; + float qy = atan(py / dom) * k; + float qz = atan(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 = clamp(*u, 0.f, 1.f - 1e-6f); + *v = clamp(*v, 0.f, 1.f - 1e-6f); + return face; +} + +static float4 sample_src(int use_front, float x, float y, + __read_only image2d_t front, + __read_only image2d_t rear) +{ + if (use_front) + return sample_img(front, x, y); + return sample_img(rear, x, y); +} + +static float4 get_colour(int face, float u, float v, + __read_only image2d_t front, + __read_only image2d_t rear, + int side, int center, int blendw) +{ + int2 dim = get_image_dim(front); + int use_front = (face == FACE_FRONT || face == FACE_LEFT || face == FACE_RIGHT); + int x0, w; + + /* RotateUV90 for rear-strip faces, then V-invert (real MAX files). */ + 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 = clamp(u, 0.f, 1.f - 1e-6f); + v = clamp(v, 0.f, 1.f - 1e-6f); + + if (face == FACE_FRONT || face == FACE_BACK) { + x0 = side; + w = center; + return sample_src(use_front, x0 + u * (float)w, v * (float)dim.y, + front, rear); + } + + 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)dim.y; + if (uL <= 0.5f - 2.f * duv) { + return sample_src(use_front, x0 + uL * (float)w, y, front, rear); + } else if (uR >= 0.5f + 2.f * duv) { + return sample_src(use_front, x0 + uR * (float)w, y, front, rear); + } else { + float4 c1 = sample_src(use_front, x0 + uL * (float)w, y, front, rear); + float4 c2 = sample_src(use_front, x0 + uR * (float)w, y, front, rear); + return mix(c1, c2, (uL - 0.5f + 2.f * duv) / (2.f * duv)); + } + } +} + +__kernel void gopromax2_equirect(__write_only image2d_t dst, + __read_only image2d_t front, + __read_only image2d_t rear, + int side, int center, int blend, + float yaw) +{ + int2 loc = (int2)(get_global_id(0), get_global_id(1)); + int2 dst_size = get_image_dim(dst); + + if (loc.x >= dst_size.x || loc.y >= dst_size.y) + return; + + float lat = M_PI_F * 0.5f - (loc.y + 0.5f) / (float)dst_size.y * M_PI_F; + float lon = (loc.x + 0.5f) / (float)dst_size.x * 2.f * M_PI_F - M_PI_F + yaw; + if (lon > M_PI_F) + lon -= 2.f * M_PI_F; + if (lon < -M_PI_F) + lon += 2.f * M_PI_F; + + float u, v; + int face = find_face(lon, lat, &u, &v); + float4 val = get_colour(face, u, v, front, rear, side, center, blend); + write_imagef(dst, loc, val); +} diff --git a/libavfilter/opencl_source.h b/libavfilter/opencl_source.h index b6930fb686..c229b10400 100644 --- a/libavfilter/opencl_source.h +++ b/libavfilter/opencl_source.h @@ -25,6 +25,7 @@ extern const char *ff_source_colorspace_common_cl; extern const char *ff_source_convolution_cl; extern const char *ff_source_deshake_cl; extern const char *ff_source_neighbor_cl; +extern const char *ff_source_gopromax2_cl; extern const char *ff_source_nlmeans_cl; extern const char *ff_source_overlay_cl; extern const char *ff_source_pad_cl; diff --git a/libavfilter/vf_gopromax2_opencl.c b/libavfilter/vf_gopromax2_opencl.c new file mode 100644 index 0000000000..57a7c9cffc --- /dev/null +++ b/libavfilter/vf_gopromax2_opencl.c @@ -0,0 +1,386 @@ +/* + * GoPro Max 2 .360 to equirectangular (OpenCL). + * Layout matches _stitch_max2: side/center/side with overlap in the sides. + * + * 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 "libavutil/log.h" +#include "libavutil/mem.h" +#include "libavutil/opt.h" +#include "libavutil/pixdesc.h" + +#ifndef M_PI +#define M_PI 3.14159265358979323846 +#endif + +#include "avfilter.h" +#include "filters.h" +#include "framesync.h" +#include "opencl.h" +#include "opencl_source.h" +#include "video.h" + +typedef struct GoProMax2OpenCLContext { + OpenCLFilterContext ocf; + + int initialised; + cl_kernel kernel; + cl_command_queue command_queue; + + FFFrameSync fs; + + int nb_planes; + + int width, height; + int side, center, blend; + float yaw; + + int in_w, in_h; + int layout_side, layout_center, layout_blend; +} GoProMax2OpenCLContext; + +static int gopromax2_opencl_load(AVFilterContext *avctx, + enum AVPixelFormat front_format, + enum AVPixelFormat rear_format) +{ + GoProMax2OpenCLContext *ctx = avctx->priv; + cl_int cle; + const char *source = ff_source_gopromax2_cl; + const AVPixFmtDescriptor *front_desc, *rear_desc; + int err, i, front_planes, rear_planes; + + front_desc = av_pix_fmt_desc_get(front_format); + rear_desc = av_pix_fmt_desc_get(rear_format); + front_planes = rear_planes = 0; + for (i = 0; i < front_desc->nb_components; i++) + front_planes = FFMAX(front_planes, + front_desc->comp[i].plane + 1); + for (i = 0; i < rear_desc->nb_components; i++) + rear_planes = FFMAX(rear_planes, + rear_desc->comp[i].plane + 1); + + if (front_planes != rear_planes) { + av_log(avctx, AV_LOG_ERROR, "Front and rear formats must have " + "the same number of planes.\n"); + return AVERROR(EINVAL); + } + + ctx->nb_planes = front_planes; + + err = ff_opencl_filter_load_program(avctx, &source, 1); + if (err < 0) + goto fail; + + ctx->command_queue = clCreateCommandQueue(ctx->ocf.hwctx->context, + ctx->ocf.hwctx->device_id, + 0, &cle); + CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create OpenCL " + "command queue %d.\n", cle); + + ctx->kernel = clCreateKernel(ctx->ocf.program, "gopromax2_equirect", &cle); + CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create kernel %d.\n", cle); + + ctx->initialised = 1; + return 0; + +fail: + if (ctx->command_queue) + clReleaseCommandQueue(ctx->command_queue); + if (ctx->kernel) + clReleaseKernel(ctx->kernel); + return err; +} + +static void gopromax2_plane_layout(GoProMax2OpenCLContext *ctx, + const AVPixFmtDescriptor *desc, + int plane, int *side, int *center, int *blend) +{ + int pw = ctx->in_w; + int ph = ctx->in_h; + + if (plane == 1 || plane == 2) { + pw = AV_CEIL_RSHIFT(ctx->in_w, desc->log2_chroma_w); + ph = AV_CEIL_RSHIFT(ctx->in_h, desc->log2_chroma_h); + } + + *side = ctx->layout_side * pw / ctx->in_w; + *center = ctx->layout_center * ph / ctx->in_h; + *blend = FFMAX(1, ctx->layout_blend * pw / ctx->in_w); +} + +static int gopromax2_opencl_filter(FFFrameSync *fs) +{ + AVFilterContext *avctx = fs->parent; + AVFilterLink *outlink = avctx->outputs[0]; + GoProMax2OpenCLContext *ctx = avctx->priv; + AVFrame *input_front, *input_rear; + AVFrame *output; + AVHWFramesContext *front_fc; + const AVPixFmtDescriptor *desc; + cl_mem mem; + cl_int cle; + float yaw; + size_t global_work[2]; + int kernel_arg = 0; + int err, plane; + int side, center, blend; + + err = ff_framesync_get_frame(fs, 0, &input_front, 0); + if (err < 0) + return err; + err = ff_framesync_get_frame(fs, 1, &input_rear, 0); + if (err < 0) + return err; + + front_fc = (AVHWFramesContext *)input_front->hw_frames_ctx->data; + desc = av_pix_fmt_desc_get(front_fc->sw_format); + + if (!ctx->initialised) { + AVHWFramesContext *rear_fc = + (AVHWFramesContext *)input_rear->hw_frames_ctx->data; + + err = gopromax2_opencl_load(avctx, front_fc->sw_format, + rear_fc->sw_format); + if (err < 0) + return err; + } + + output = ff_get_video_buffer(outlink, outlink->w, outlink->h); + if (!output) { + err = AVERROR(ENOMEM); + goto fail; + } + + yaw = ctx->yaw * (float)M_PI / 180.f; + + for (plane = 0; plane < ctx->nb_planes; plane++) { + kernel_arg = 0; + gopromax2_plane_layout(ctx, desc, plane, &side, ¢er, &blend); + + mem = (cl_mem)output->data[plane]; + CL_SET_KERNEL_ARG(ctx->kernel, kernel_arg, cl_mem, &mem); + kernel_arg++; + + mem = (cl_mem)input_front->data[plane]; + CL_SET_KERNEL_ARG(ctx->kernel, kernel_arg, cl_mem, &mem); + kernel_arg++; + + mem = (cl_mem)input_rear->data[plane]; + CL_SET_KERNEL_ARG(ctx->kernel, kernel_arg, cl_mem, &mem); + kernel_arg++; + + CL_SET_KERNEL_ARG(ctx->kernel, kernel_arg, cl_int, &side); + kernel_arg++; + CL_SET_KERNEL_ARG(ctx->kernel, kernel_arg, cl_int, ¢er); + kernel_arg++; + CL_SET_KERNEL_ARG(ctx->kernel, kernel_arg, cl_int, &blend); + kernel_arg++; + CL_SET_KERNEL_ARG(ctx->kernel, kernel_arg, cl_float, &yaw); + kernel_arg++; + + err = ff_opencl_filter_work_size_from_image(avctx, global_work, + output, plane, 0); + if (err < 0) + goto fail; + + av_log(avctx, AV_LOG_DEBUG, + "gopromax2_opencl plane %d work %ux%u side %d center %d blend %d\n", + plane, (unsigned)global_work[0], (unsigned)global_work[1], + side, center, blend); + + cle = clEnqueueNDRangeKernel(ctx->command_queue, ctx->kernel, 2, NULL, + global_work, NULL, 0, NULL, NULL); + CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to enqueue gopromax2 kernel " + "for plane %d: %d.\n", plane, cle); + } + + cle = clFinish(ctx->command_queue); + CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to finish command queue: %d.\n", cle); + + err = av_frame_copy_props(output, input_front); + if (err < 0) + goto fail; + + return ff_filter_frame(outlink, output); + +fail: + av_frame_free(&output); + return err; +} + +static int gopromax2_opencl_config_output(AVFilterLink *outlink) +{ + AVFilterContext *avctx = outlink->src; + GoProMax2OpenCLContext *ctx = avctx->priv; + int in_w = avctx->inputs[0]->w; + int in_h = avctx->inputs[0]->h; + int err; + + if (in_w <= 0 || in_h <= 0) + return AVERROR(EINVAL); + + ctx->in_w = in_w; + ctx->in_h = in_h; + ctx->layout_center = ctx->center > 0 ? ctx->center : in_h; + ctx->layout_side = ctx->side > 0 ? ctx->side : (in_w - ctx->layout_center) / 2; + ctx->layout_blend = ctx->blend > 0 ? ctx->blend : 48; + + if (ctx->layout_side <= 0 || ctx->layout_center <= 0 || + 2 * ctx->layout_side + ctx->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, ctx->layout_side, ctx->layout_center); + return AVERROR(EINVAL); + } + + if (ctx->width > 0 && ctx->height > 0) { + if (ctx->width != ctx->height * 2) { + av_log(avctx, AV_LOG_ERROR, + "Specified size (%dx%d) is not 2:1 equirect.\n", + ctx->width, ctx->height); + return AVERROR(EINVAL); + } + ctx->ocf.output_width = ctx->width; + ctx->ocf.output_height = ctx->height; + } else if (ctx->width > 0 || ctx->height > 0) { + av_log(avctx, AV_LOG_ERROR, + "Both width and height values should be specified.\n"); + return AVERROR(EINVAL); + } else { + ctx->ocf.output_width = 4 * in_h; + ctx->ocf.output_height = 2 * in_h; + } + + av_log(avctx, AV_LOG_INFO, + "gopromax2_opencl in %dx%d side %d center %d blend %d yaw %.1f out %dx%d\n", + in_w, in_h, ctx->layout_side, ctx->layout_center, ctx->layout_blend, + ctx->yaw, ctx->ocf.output_width, ctx->ocf.output_height); + + err = ff_opencl_filter_config_output(outlink); + if (err < 0) + return err; + + err = ff_framesync_init_dualinput(&ctx->fs, avctx); + if (err < 0) + return err; + + return ff_framesync_configure(&ctx->fs); +} + +static av_cold int gopromax2_opencl_init(AVFilterContext *avctx) +{ + GoProMax2OpenCLContext *ctx = avctx->priv; + + ctx->fs.on_event = &gopromax2_opencl_filter; + + return ff_opencl_filter_init(avctx); +} + +static int gopromax2_opencl_activate(AVFilterContext *avctx) +{ + GoProMax2OpenCLContext *ctx = avctx->priv; + + return ff_framesync_activate(&ctx->fs); +} + +static av_cold void gopromax2_opencl_uninit(AVFilterContext *avctx) +{ + GoProMax2OpenCLContext *ctx = avctx->priv; + cl_int cle; + + if (ctx->kernel) { + cle = clReleaseKernel(ctx->kernel); + if (cle != CL_SUCCESS) + av_log(avctx, AV_LOG_ERROR, "Failed to release " + "kernel: %d.\n", cle); + } + + if (ctx->command_queue) { + cle = clReleaseCommandQueue(ctx->command_queue); + if (cle != CL_SUCCESS) + av_log(avctx, AV_LOG_ERROR, "Failed to release " + "command queue: %d.\n", cle); + } + + ff_opencl_filter_uninit(avctx); + + ff_framesync_uninit(&ctx->fs); +} + +#define OFFSET(x) offsetof(GoProMax2OpenCLContext, x) +#define FLAGS (AV_OPT_FLAG_FILTERING_PARAM | AV_OPT_FLAG_VIDEO_PARAM) +static const AVOption gopromax2_opencl_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_opencl); + +static const AVFilterPad gopromax2_opencl_inputs[] = { + { + .name = "front", + .type = AVMEDIA_TYPE_VIDEO, + .config_props = &ff_opencl_filter_config_input, + }, + { + .name = "rear", + .type = AVMEDIA_TYPE_VIDEO, + .config_props = &ff_opencl_filter_config_input, + }, +}; + +static const AVFilterPad gopromax2_opencl_outputs[] = { + { + .name = "default", + .type = AVMEDIA_TYPE_VIDEO, + .config_props = &gopromax2_opencl_config_output, + }, +}; + +const FFFilter ff_vf_gopromax2_opencl = { + .p.name = "gopromax2_opencl", + .p.description = NULL_IF_CONFIG_SMALL("GoPro Max 2 .360 to equirectangular projection"), + .p.priv_class = &gopromax2_opencl_class, + .p.flags = AVFILTER_FLAG_HWDEVICE, + .priv_size = sizeof(GoProMax2OpenCLContext), + .init = &gopromax2_opencl_init, + .uninit = &gopromax2_opencl_uninit, + .activate = &gopromax2_opencl_activate, + FILTER_INPUTS(gopromax2_opencl_inputs), + FILTER_OUTPUTS(gopromax2_opencl_outputs), + FILTER_SINGLE_PIXFMT(AV_PIX_FMT_OPENCL), + .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE, +}; -- 2.52.0 _______________________________________________ ffmpeg-devel mailing list -- [email protected] To unsubscribe send an email to [email protected]
