@@ -3886,6 +3886,7 @@ frei0r_src_filter_deps="frei0r"
fspp_filter_deps="gpl"
fsync_filter_deps="avformat"
gblur_vulkan_filter_deps="vulkan spirv_compiler"
+gopromax_opencl_filter_deps="opencl"
hflip_vulkan_filter_deps="vulkan spirv_compiler"
histeq_filter_deps="gpl"
hqdn3d_filter_deps="gpl"
@@ -27201,6 +27201,84 @@ Apply dilation filter with threshold0 set to 30, threshold1 set 40, threshold2 s
@end example
@end itemize
+@anchor{gopromax_opencl}
+@section gopromax_opencl
+
+Apply transformation of the two GoPro Max video streams to equirectangular or equiangular-cubemap projection.
+
+This filter is designed to use directly GoPro .360 files.
+Native .360 files are sort of EAC files, in fact the front and rear lenses streams are the top and the bottom of the EAC projection.
+
+The .360 file contains two video streams.
+Most of cases, one is stream #0:0, and the other is stream #0:5.
+Please check actual stream number with @code{ffprobe} command.
+This filter combine two streams to single stream.
+
+The .360 contains also 2x64 bits of overlapped area.
+The filter blends overlapped images in these two areas.
+
+The filter accepts the following options:
+
+@table @option
+
+@item output
+Set format of the output video.
+
+Available formats:
+
+@table @samp
+
+@item e
+@item equirect
+Equirectangular projection.
+
+@item eac
+Equi-Angular Cubemap.
+
+@end table
+
+Default is @code{equirect}.
+
+@item w
+@item h
+Set the output video resolution.
+
+Default resolution depends on formats.
+
+@item overlap
+Set number of overlapped pixels on input .360 video.
+
+No need to specify this option for native .360 video file.
+This option is for rescaled video or future video format change.
+
+Default is @code{64}.
+
+@end table
+
+@subsection Example
+
+@itemize
+@item
+Convert .360 to Equirectangular projection.
+@example
+-i INPUT -filter_complex '[0:0]hwupload[a], [0:5]hwupload[b], [a][b]gopromax_opencl=w=4096:h=2048, hwdownload, format=yuvj420p' -map 0:a:0 -c:a copy OUTPUT
+@end example
+
+Two video streams (#0:0 and #0:5) are combined and converted to default equirectangular projection with specified resolution.
+First audio stream (GoPro AAC) is copied with the video stream.
+
+@item
+Convert .360 to Equi-Angular Cubemap projection.
+@example
+-i INPUT -filter_complex '[0:0]hwupload[a], [0:5]hwupload[b], [a][b]gopromax_opencl=eac, hwdownload, format=yuvj420p, v360=eac:c3x2:w=1344:h=896' -map 0:1 -map 0:3 -c:a copy -c:u copy OUTPUT
+@end example
+
+Two video streams (#0:0 and #0:5) are combined and converted to equi-angular cubemap projection,
+then it is converted to c3x2 cubemap projection and shrunk by v360 filter.
+Stream #0:1 (GoPro AAC) and stream #0:3 (GoPro MET) are copied with the video stream.
+
+@end itemize
+
@anchor{nlmeans_opencl}
@section nlmeans_opencl
@@ -327,6 +327,8 @@ OBJS-$(CONFIG_FSYNC_FILTER) += vf_fsync.o
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_GOPROMAX_OPENCL_FILTER) += vf_gopromax_opencl.o opencl.o \
+ opencl/gopromax.o framesync.o
OBJS-$(CONFIG_GRADFUN_FILTER) += vf_gradfun.o
OBJS-$(CONFIG_GRAPHMONITOR_FILTER) += f_graphmonitor.o
OBJS-$(CONFIG_GRAYWORLD_FILTER) += vf_grayworld.o
@@ -303,6 +303,7 @@ extern const AVFilter ff_vf_fsync;
extern const AVFilter ff_vf_gblur;
extern const AVFilter ff_vf_gblur_vulkan;
extern const AVFilter ff_vf_geq;
+extern const AVFilter ff_vf_gopromax_opencl;
extern const AVFilter ff_vf_gradfun;
extern const AVFilter ff_vf_graphmonitor;
extern const AVFilter ff_vf_grayworld;
new file mode 100644
@@ -0,0 +1,282 @@
+/*
+ * Copyright (c) 2021 Ronan LE MEILLAT
+ * Copyright (c) 2024 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 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;
+ //eac expansion
+ 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;
+ int gap = (cube_size * 3 + overlap * 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 + overlap - gap;
+ if (cx >= (cs + overlap) / 2)
+ dx += overlap;
+ uv2.x = cx + dx;
+ if (cx > (cs - overlap) / 2 && cx < (cs + overlap) / 2)
+ a = (cx - (cs - overlap) / 2) / overlap;
+ } else {
+ uv2.x += overlap - 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 += overlap;
+ 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);
+}
@@ -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;
new file mode 100644
@@ -0,0 +1,350 @@
+/*
+ * Copyright (c) 2021 Ronan LE MEILLAT
+ * Copyright (c) 2024 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 "libavutil/log.h"
+#include "libavutil/mem.h"
+#include "libavutil/pixdesc.h"
+#include "libavutil/opt.h"
+
+#include "avfilter.h"
+#include "framesync.h"
+#include "opencl.h"
+#include "opencl_source.h"
+#include "video.h"
+#include "v360.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);
+
+ ctx->nb_planes = front_planes;
+
+ switch (ctx->out) {
+ case EQUIRECTANGULAR:
+ kernel = "gopromax_equirectangular";
+ break;
+ case EQUIANGULAR:
+ kernel = "gopromax_stack";
+ break;
+ default:
+ av_log(ctx, 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;
+ 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_VERBOSE,
+ "In gopromax_opencl_stack for plane:%d %lux%lu frame size %dx%d\n",
+ plane, global_work[0], 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);
+
+ 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(ctx, 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(ctx, 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(ctx, 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(ctx, 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(ctx, AV_LOG_ERROR, "Specified output format is not supported.\n");
+ return AVERROR(EINVAL);
+ }
+
+ 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, .unit = "w"},
+ { "h", "output height", OFFSET(height), AV_OPT_TYPE_INT, {.i64=0}, 0, INT16_MAX, FLAGS, .unit = "h"},
+ { "overlap", "set overlapped pixels", OFFSET(overlap), AV_OPT_TYPE_INT, {.i64=64}, 0, 128, FLAGS, .unit = "overlap"},
+ { 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 AVFilter ff_vf_gopromax_opencl = {
+ .name = "gopromax_opencl",
+ .description = NULL_IF_CONFIG_SMALL("GoProMax .360 to equirectangular projection"),
+ .priv_size = sizeof(GoProMaxOpenCLContext),
+ .priv_class = &gopromax_opencl_class,
+ .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,
+ .flags = AVFILTER_FLAG_HWDEVICE,
+};