PR #24723 opened by rnissl URL: https://code.ffmpeg.org/FFmpeg/FFmpeg/pulls/24723 Patch URL: https://code.ffmpeg.org/FFmpeg/FFmpeg/pulls/24723.patch
This is 8/9 in the GoPro Max review order. It applies on current master by itself. OpenCL stitch for GoPro Max 1 .360. Port of gopromax_opencl from gmat/goproMax-ffmpeg-v5 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 73930e00c2a6b8477296c5f701686ccc7eaf3644 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Reinhard=20Ni=C3=9Fl?= <[email protected]> Date: Sat, 26 Sep 2026 21:35:03 +0200 Subject: [PATCH] avfilter: add gopromax_opencl GoPro Max 1 stitch MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit OpenCL stitch for GoPro Max 1 .360. Port of gopromax_opencl from gmat/goproMax-ffmpeg-v5 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 | 23 ++ libavfilter/Makefile | 2 + libavfilter/allfilters.c | 1 + libavfilter/opencl/gopromax.cl | 288 ++++++++++++++++++++++++ libavfilter/opencl_source.h | 1 + libavfilter/vf_gopromax_opencl.c | 370 +++++++++++++++++++++++++++++++ 8 files changed, 688 insertions(+) create mode 100644 libavfilter/opencl/gopromax.cl create mode 100644 libavfilter/vf_gopromax_opencl.c diff --git a/Changelog b/Changelog index 1cfa9db1b6..cf06271a86 100644 --- a/Changelog +++ b/Changelog @@ -16,6 +16,8 @@ version <next>: - DSD (dsd_msbf) encoder +- gopromax_opencl GoPro Max 1 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..07d5896781 100755 --- a/configure +++ b/configure @@ -4126,6 +4126,7 @@ asr_filter_deps="pocketsphinx" ass_filter_deps="libass" avgblur_opencl_filter_deps="opencl" avgblur_vulkan_filter_deps="vulkan spirv_compiler" +gopromax_opencl_filter_deps="opencl" azmq_filter_deps="libzmq" blackdetect_vulkan_filter_deps="vulkan spirv_compiler" blackframe_filter_deps="gpl" diff --git a/doc/filters.texi b/doc/filters.texi index 1a9bbeff7e..431bf11d7c 100644 --- a/doc/filters.texi +++ b/doc/filters.texi @@ -28395,6 +28395,29 @@ Apply dilation filter with threshold0 set to 30, threshold1 set 40, threshold2 s Non-local Means denoise filter through OpenCL, this filter accepts same options as @ref{nlmeans}. +@anchor{gopromax_opencl} +@section gopromax_opencl + +OpenCL stitch for GoPro Max 1 .360 dual-EAC. Port of gopromax_opencl from +gmat by Ronan LE MEILLAT and TADANO Tokumei. GoPro Max 2 uses a different +face layout. + +Typical streams are @code{0:0} and @code{0:5}; check with ffprobe. + +@table @option +@item output +@samp{equirect} (default) or @samp{eac}. +@item w +@item h +Output size. +@item overlap +Overlapped pixels. @code{0} derives @code{width - 3*height}. +@end table + +@example +ffmpeg -i INPUT.360 -filter_complex "[0:0]hwupload[a];[0:5]hwupload[b];[a][b]gopromax_opencl,hwdownload,format=yuvj420p" out.mp4 +@end example + @section overlay_opencl Overlay one video on top of another. diff --git a/libavfilter/Makefile b/libavfilter/Makefile index ceea328e67..af0359f39c 100644 --- a/libavfilter/Makefile +++ b/libavfilter/Makefile @@ -337,6 +337,8 @@ OBJS-$(CONFIG_GBLUR_FILTER) += vf_gblur.o OBJS-$(CONFIG_GBLUR_VULKAN_FILTER) += vf_gblur_vulkan.o vulkan.o vulkan_filter.o OBJS-$(CONFIG_GEQ_FILTER) += vf_geq.o OBJS-$(CONFIG_GRADFUN_FILTER) += vf_gradfun.o +OBJS-$(CONFIG_GOPROMAX_OPENCL_FILTER) += vf_gopromax_opencl.o opencl.o \ + opencl/gopromax.o framesync.o OBJS-$(CONFIG_GRAPHMONITOR_FILTER) += f_graphmonitor.o OBJS-$(CONFIG_GRAYWORLD_FILTER) += vf_grayworld.o OBJS-$(CONFIG_GREYEDGE_FILTER) += vf_colorconstancy.o diff --git a/libavfilter/allfilters.c b/libavfilter/allfilters.c index 3b93fe12f6..3c04a51c63 100644 --- a/libavfilter/allfilters.c +++ b/libavfilter/allfilters.c @@ -311,6 +311,7 @@ extern const FFFilter ff_vf_gblur; extern const FFFilter ff_vf_gblur_vulkan; extern const FFFilter ff_vf_geq; extern const FFFilter ff_vf_gradfun; +extern const FFFilter ff_vf_gopromax_opencl; extern const FFFilter ff_vf_graphmonitor; extern const FFFilter ff_vf_grayworld; extern const FFFilter ff_vf_greyedge; diff --git a/libavfilter/opencl/gopromax.cl b/libavfilter/opencl/gopromax.cl new file mode 100644 index 0000000000..1ea420c594 --- /dev/null +++ b/libavfilter/opencl/gopromax.cl @@ -0,0 +1,288 @@ +/* + * Copyright (c) 2021 Ronan LE MEILLAT + * Copyright (c) 2024 TADANO Tokumei + * Copyright (c) 2026 Reinhard Nißl + * + * Port of gopromax_opencl from gmat/goproMax-ffmpeg-v5. + * + * 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 Faces { + TOP_LEFT, + TOP_MIDDLE, + TOP_RIGHT, + BOTTOM_LEFT, + BOTTOM_MIDDLE, + BOTTOM_RIGHT, + NB_FACES, +}; + +enum Direction { + RIGHT, + LEFT, + UP, + DOWN, + FRONT, + BACK, + NB_DIRECTIONS, +}; + +enum Rotation { + ROT_0, + ROT_90, + ROT_180, + ROT_270, + NB_ROTATIONS, +}; + +static float2 rotate_cube_face(float2 uv, int rotation) +{ + float2 ret_uv; + + switch (rotation) { + case ROT_0: + ret_uv = uv; + break; + case ROT_90: + ret_uv.x = -uv.y; + ret_uv.y = uv.x; + break; + case ROT_180: + ret_uv.x = -uv.x; + ret_uv.y = -uv.y; + break; + case ROT_270: + ret_uv.x = uv.y; + ret_uv.y = -uv.x; + break; + } + + return ret_uv; +} + +static float3 equirect_to_xyz(int2 xy, int2 size) +{ + float3 xyz; + float phi = ((2.f * ((float)xy.x) + 1.f) / ((float)size.x) - 1.f) * M_PI_F; + float theta = ((2.f * ((float)xy.y) + 1.f) / ((float)size.y) - 1.f) * M_PI_2_F; + + xyz.x = cos(theta) * sin(phi); + xyz.y = sin(theta); + xyz.z = cos(theta) * cos(phi); + + return xyz; +} + +static float2 xyz_to_cube(float3 xyz, int *face) +{ + float phi = atan2(xyz.x, xyz.z); + float theta = asin(xyz.y); + float phi_norm, theta_threshold; + float2 uv; + int direction; + + if (phi >= -M_PI_4_F && phi < M_PI_4_F) { + direction = FRONT; + phi_norm = phi; + } else if (phi >= -(M_PI_2_F + M_PI_4_F) && phi < -M_PI_4_F) { + direction = LEFT; + phi_norm = phi + M_PI_2_F; + } else if (phi >= M_PI_4_F && phi < M_PI_2_F + M_PI_4_F) { + direction = RIGHT; + phi_norm = phi - M_PI_2_F; + } else { + direction = BACK; + phi_norm = phi + ((phi > 0.f) ? -M_PI_F : M_PI_F); + } + + theta_threshold = atan(cos(phi_norm)); + if (theta > theta_threshold) { + direction = DOWN; + } else if (theta < -theta_threshold) { + direction = UP; + } + + switch (direction) { + case RIGHT: + uv.x = -xyz.z / xyz.x; + uv.y = xyz.y / xyz.x; + *face = TOP_RIGHT; + break; + case LEFT: + uv.x = -xyz.z / xyz.x; + uv.y = -xyz.y / xyz.x; + *face = TOP_LEFT; + break; + case UP: + uv.x = -xyz.x / xyz.y; + uv.y = -xyz.z / xyz.y; + *face = BOTTOM_RIGHT; + uv = rotate_cube_face(uv, ROT_270); + break; + case DOWN: + uv.x = xyz.x / xyz.y; + uv.y = -xyz.z / xyz.y; + *face = BOTTOM_LEFT; + uv = rotate_cube_face(uv, ROT_270); + break; + case FRONT: + uv.x = xyz.x / xyz.z; + uv.y = xyz.y / xyz.z; + *face = TOP_MIDDLE; + break; + case BACK: + uv.x = xyz.x / xyz.z; + uv.y = -xyz.y / xyz.z; + *face = BOTTOM_MIDDLE; + uv = rotate_cube_face(uv, ROT_90); + break; + } + + return uv; +} + +static float2 xyz_to_eac(float3 xyz, int2 size) +{ + float pixel_pad = 2; + float u_pad = pixel_pad / size.x; + float v_pad = pixel_pad / size.y; + + int face; + int u_face, v_face; + float2 uv = xyz_to_cube(xyz, &face); + + u_face = face % 3; + v_face = face / 3; + uv = M_2_PI_F * atan(uv) + 0.5f; + + uv.x = (uv.x + u_face) * (1.f - 2.f * u_pad) / 3.f + u_pad; + uv.y = uv.y * (0.5f - 2.f * v_pad) + v_pad + 0.5f * v_face; + + uv.x *= size.x; + uv.y *= size.y; + + uv -= 0.5f; + + return uv; +} + +const sampler_t sampler_nearest = (CLK_NORMALIZED_COORDS_FALSE | + CLK_ADDRESS_CLAMP_TO_EDGE | + CLK_FILTER_NEAREST); + +const sampler_t sampler_linear = (CLK_NORMALIZED_COORDS_FALSE | + CLK_ADDRESS_CLAMP_TO_EDGE | + CLK_FILTER_LINEAR); + +static float4 gopromax_to_eac(float2 uv, int overlap, __read_only image2d_t src) +{ + int2 dim = get_image_dim(src); + int cube_size = dim.y; + /* Max 1 overlap: width - 3*height (4096x1344 → 64, 2272x736 → 64). Also chroma. */ + int ov = overlap > 0 ? overlap : (dim.x - cube_size * 3); + if (ov < 0) + ov = 0; + int gap = (cube_size * 3 + ov * 2 - dim.x) / 2; + float2 uv2 = uv; + float a = 0.f; + float4 val; + bool is_aligned; + + if (uv.x < cube_size || uv.x > cube_size * 2) { + int dx = 0; + int cs = cube_size - gap; + float cx = fmod(uv.x, cube_size) * cs / cube_size; + if (uv.x >= cube_size * 2) + dx = cube_size * 2 + ov - gap; + if (cx >= (cs + ov) / 2) + dx += ov; + uv2.x = cx + dx; + if (ov > 0 && cx > (cs - ov) / 2 && cx < (cs + ov) / 2) + a = (cx - (cs - ov) / 2) / ov; + } else { + uv2.x += ov - gap; + } + + { + int2 d = convert_int2(ceil(uv2) - floor(uv2)); + is_aligned = (d.x == 0 && d.y == 0); + } + if (is_aligned) + val = read_imagef(src, sampler_nearest, uv2); + else + val = read_imagef(src, sampler_linear, uv2); + if (a > 0.f) { + float4 val2; + uv2.x += ov; + if (is_aligned) + val2 = read_imagef(src, sampler_nearest, uv2); + else + val2 = read_imagef(src, sampler_linear, uv2); + val = mix(val, val2, a); + } + + return val; +} + +__kernel void gopromax_equirectangular(__write_only image2d_t dst, + __read_only image2d_t front, + __read_only image2d_t rear, + int overlap) +{ + float4 val; + int2 loc = (int2)(get_global_id(0), get_global_id(1)); + + int2 dst_size = get_image_dim(dst); + int2 src_size = get_image_dim(front); + int2 eac_size = (int2)(src_size.y * 3, src_size.y * 2); + + float3 xyz = equirect_to_xyz(loc, dst_size); + float2 uv = xyz_to_eac(xyz, eac_size); + + if (uv.y >= src_size.y) { + uv.y -= src_size.y; + val = gopromax_to_eac(uv, overlap, rear); + } else { + val = gopromax_to_eac(uv, overlap, front); + } + + write_imagef(dst, loc, val); +} + +__kernel void gopromax_stack(__write_only image2d_t dst, + __read_only image2d_t front, + __read_only image2d_t rear, + int overlap) +{ + float4 val; + int2 loc = (int2)(get_global_id(0), get_global_id(1)); + int2 dst_size = get_image_dim(dst); + int2 src_size = get_image_dim(front); + float2 uv = convert_float2(loc); + + uv *= (float)src_size.y * 2 / dst_size.y; + + if (uv.y >= src_size.y) { + uv.y -= src_size.y; + val = gopromax_to_eac(uv, overlap, rear); + } else { + val = gopromax_to_eac(uv, overlap, front); + } + + write_imagef(dst, loc, val); +} diff --git a/libavfilter/opencl_source.h b/libavfilter/opencl_source.h index b6930fb686..92135c6a7d 100644 --- a/libavfilter/opencl_source.h +++ b/libavfilter/opencl_source.h @@ -24,6 +24,7 @@ extern const char *ff_source_colorkey_cl; 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_gopromax_cl; extern const char *ff_source_neighbor_cl; extern const char *ff_source_nlmeans_cl; extern const char *ff_source_overlay_cl; diff --git a/libavfilter/vf_gopromax_opencl.c b/libavfilter/vf_gopromax_opencl.c new file mode 100644 index 0000000000..18c1bd9990 --- /dev/null +++ b/libavfilter/vf_gopromax_opencl.c @@ -0,0 +1,370 @@ +/* + * Copyright (c) 2021 Ronan LE MEILLAT + * Copyright (c) 2024 TADANO Tokumei + * Copyright (c) 2026 Reinhard Nißl + * + * Port of gopromax_opencl from gmat/goproMax-ffmpeg-v5. + * + * 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/log.h" +#include "libavutil/mem.h" +#include "libavutil/opt.h" +#include "libavutil/pixdesc.h" + +#include "avfilter.h" +#include "filters.h" +#include "framesync.h" +#include "opencl.h" +#include "opencl_source.h" +#include "v360.h" +#include "video.h" + +typedef struct GoProMaxOpenCLContext { + OpenCLFilterContext ocf; + + int initialised; + cl_kernel kernel; + cl_command_queue command_queue; + + FFFrameSync fs; + + int nb_planes; + + int out; + int width, height; + int overlap; +} GoProMaxOpenCLContext; + +static int gopromax_opencl_load(AVFilterContext *avctx, + enum AVPixelFormat front_format, + enum AVPixelFormat rear_format) +{ + GoProMaxOpenCLContext *ctx = avctx->priv; + cl_int cle; + const char *source = ff_source_gopromax_cl; + const char *kernel; + 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; + + switch (ctx->out) { + case EQUIRECTANGULAR: + kernel = "gopromax_equirectangular"; + break; + case EQUIANGULAR: + kernel = "gopromax_stack"; + break; + default: + av_log(avctx, AV_LOG_ERROR, "Specified output format is not handled.\n"); + return AVERROR_BUG; + } + + av_log(avctx, AV_LOG_DEBUG, "Using kernel %s.\n", kernel); + + 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, kernel, &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 int gopromax_opencl_stack(FFFrameSync *fs) +{ + AVFilterContext *avctx = fs->parent; + AVFilterLink *outlink = avctx->outputs[0]; + GoProMaxOpenCLContext *ctx = avctx->priv; + AVFrame *input_front, *input_rear; + AVFrame *output; + cl_mem mem; + cl_int cle, overlap; + size_t global_work[2]; + int kernel_arg = 0; + int err, plane; + + 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; + + if (!ctx->initialised) { + AVHWFramesContext *front_fc = + (AVHWFramesContext*)input_front->hw_frames_ctx->data; + AVHWFramesContext *rear_fc = + (AVHWFramesContext*)input_rear->hw_frames_ctx->data; + + err = gopromax_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; + } + + for (plane = 0; plane < ctx->nb_planes; plane++) { + kernel_arg = 0; + + 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++; + + overlap = ctx->overlap; + /* 0 lets the kernel derive overlap as width - 3*height (correct per plane). */ + CL_SET_KERNEL_ARG(ctx->kernel, kernel_arg, cl_int, &overlap); + 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, + "gopromax_opencl plane %d work size %ux%u frame size %dx%d\n", + plane, (unsigned)global_work[0], (unsigned)global_work[1], + outlink->w, outlink->h); + + cle = clEnqueueNDRangeKernel(ctx->command_queue, ctx->kernel, 2, NULL, + global_work, NULL, 0, NULL, NULL); + CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to enqueue gopromax 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; + + av_log(avctx, AV_LOG_DEBUG, "Filter output: %s, %ux%u (%"PRId64").\n", + av_get_pix_fmt_name(output->format), + output->width, output->height, output->pts); + + return ff_filter_frame(outlink, output); + +fail: + av_frame_free(&output); + return err; +} + +static int gopromax_opencl_config_output(AVFilterLink *outlink) +{ + AVFilterContext *avctx = outlink->src; + GoProMaxOpenCLContext *ctx = avctx->priv; + int height = avctx->inputs[0]->h; + int err; + + switch (ctx->out) { + case EQUIRECTANGULAR: + 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 suitable.\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 * height; + ctx->ocf.output_height = 2 * height; + } + break; + case EQUIANGULAR: + if (ctx->width > 0 && ctx->height > 0) { + if (ctx->width * 2 != ctx->height * 3) { + av_log(avctx, AV_LOG_ERROR, + "Specified size (%dx%d) is not suitable.\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 = 3 * height; + ctx->ocf.output_height = 2 * height; + } + break; + default: + av_log(avctx, AV_LOG_ERROR, "Specified output format is not supported.\n"); + return AVERROR(EINVAL); + } + + if (!ctx->overlap) { + int auto_overlap = avctx->inputs[0]->w - 3 * avctx->inputs[0]->h; + av_log(avctx, AV_LOG_INFO, "gopromax_opencl overlap auto %d from %dx%d\n", + FFMAX(auto_overlap, 0), avctx->inputs[0]->w, avctx->inputs[0]->h); + } + + 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 gopromax_opencl_init(AVFilterContext *avctx) +{ + GoProMaxOpenCLContext *ctx = avctx->priv; + + ctx->fs.on_event = &gopromax_opencl_stack; + + return ff_opencl_filter_init(avctx); +} + +static int gopromax_opencl_activate(AVFilterContext *avctx) +{ + GoProMaxOpenCLContext *ctx = avctx->priv; + + return ff_framesync_activate(&ctx->fs); +} + +static av_cold void gopromax_opencl_uninit(AVFilterContext *avctx) +{ + GoProMaxOpenCLContext *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(GoProMaxOpenCLContext, x) +#define FLAGS (AV_OPT_FLAG_FILTERING_PARAM | AV_OPT_FLAG_VIDEO_PARAM) +static const AVOption gopromax_opencl_options[] = { + { "output", "set output projection", OFFSET(out), AV_OPT_TYPE_INT, {.i64=EQUIRECTANGULAR}, 0, NB_PROJECTIONS-1, FLAGS, .unit = "out" }, + { "e", "equirectangular", 0, AV_OPT_TYPE_CONST, {.i64=EQUIRECTANGULAR}, 0, 0, FLAGS, .unit = "out" }, + { "equirect", "equirectangular", 0, AV_OPT_TYPE_CONST, {.i64=EQUIRECTANGULAR}, 0, 0, FLAGS, .unit = "out" }, + { "eac", "equi-angular cubemap", 0, AV_OPT_TYPE_CONST, {.i64=EQUIANGULAR}, 0, 0, FLAGS, .unit = "out" }, + { "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 }, + { "overlap", "set overlapped pixels (0 = auto from frame size)", OFFSET(overlap), AV_OPT_TYPE_INT, {.i64=0}, 0, 1024, FLAGS }, + { NULL }, +}; + +AVFILTER_DEFINE_CLASS(gopromax_opencl); + +static const AVFilterPad gopromax_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 gopromax_opencl_outputs[] = { + { + .name = "default", + .type = AVMEDIA_TYPE_VIDEO, + .config_props = &gopromax_opencl_config_output, + }, +}; + +const FFFilter ff_vf_gopromax_opencl = { + .p.name = "gopromax_opencl", + .p.description = NULL_IF_CONFIG_SMALL("GoPro Max (Max 1) .360 to equirectangular projection"), + .p.priv_class = &gopromax_opencl_class, + .p.flags = AVFILTER_FLAG_HWDEVICE, + .priv_size = sizeof(GoProMaxOpenCLContext), + .init = &gopromax_opencl_init, + .uninit = &gopromax_opencl_uninit, + .activate = &gopromax_opencl_activate, + FILTER_INPUTS(gopromax_opencl_inputs), + FILTER_OUTPUTS(gopromax_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]
