[FFmpeg-devel] [PATCH] filters: Add an OpenCL filter for filtering GoPro Max native .360 files into standard equirectangular (default) or youtube equiangular cubemap (eac) projection
Ronan .
gopromax_opencl at outlook.fr
Sun May 9 09:54:09 EEST 2021
>From 47c39e21f81b6536c96b91aebfd755e3f8a39463 Mon Sep 17 00:00:00 2001
From: Ronan LE MEILLAT <gropromax_opencl at outlook.fr>
Date: Sun, 9 May 2021 08:42:58 +0200
Subject: [PATCH] filters: Add an OpenCL filter for filtering GoPro Max native .360 files into standard equirectangular (default) or youtube equiangular cubemap (eac) projection
Signed-off-by: Ronan LE MEILLAT <gopromax_opencl at outlook.fr>
---
doc/filters.texi | 26 +++
libavfilter/Makefile | 2 +
libavfilter/allfilters.c | 1 +
libavfilter/opencl/gopromax.cl | 305 ++++++++++++++++++++++++++++
libavfilter/opencl_source.h | 2 +-
libavfilter/vf_gopromax_opencl.c | 330 +++++++++++++++++++++++++++++++
6 files changed, 665 insertions(+), 1 deletion(-)
create mode 100644 libavfilter/opencl/gopromax.cl
create mode 100644 libavfilter/vf_gopromax_opencl.c
diff --git a/doc/filters.texi b/doc/filters.texi
index b405cc5dfb..0f92163f0e 100644
--- a/doc/filters.texi
+++ b/doc/filters.texi
@@ -23041,6 +23041,32 @@ Apply dilation filter with threshold0 set to 30, threshold1 set 40, threshold2 s
@end example
@end itemize
+ at 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.
+
+It accepts the following options:
+
+ at table @option
+
+ at item eac
+If @code{1} the ouptut is EAC. If @code{0} (default) the ouptut is equirectangular. 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 contains also 2x64 bits of overlapped area. The filter removes these two areas.
+If eac option is not specified, the output is equirectangular.
+
+ at end table
+
+ at subsection Example
+
+ at itemize
+ at item
+Convert .360 to equirectangular with OpenCL 0:1 device.
+ at example
+-i INPUT -hwaccel auto -hwaccel auto -init_hw_device opencl:0.1 -filter_hw_device opencl0 -v verbose -filter_complex '[0:0]format=yuv420p,hwupload[a] , [0:4]format=yuv420p,hwupload[b], [a][b]gopromax_opencl, hwdownload,format=yuv420p' OUTPUT
+ at end example
+ at end itemize
+
@section nlmeans_opencl
Non-local Means denoise filter through OpenCL, this filter accepts same options as @ref{nlmeans}.
diff --git a/libavfilter/Makefile b/libavfilter/Makefile
index 6c22d0404e..a3ca814c14 100644
--- a/libavfilter/Makefile
+++ b/libavfilter/Makefile
@@ -477,6 +477,8 @@ OBJS-$(CONFIG_VIGNETTE_FILTER) += vf_vignette.o
OBJS-$(CONFIG_VMAFMOTION_FILTER) += vf_vmafmotion.o framesync.o
OBJS-$(CONFIG_VPP_QSV_FILTER) += vf_vpp_qsv.o
OBJS-$(CONFIG_VSTACK_FILTER) += vf_stack.o framesync.o
+OBJS-$(CONFIG_GOPROMAX_OPENCL_FILTER) += vf_gopromax_opencl.o opencl.o \
+ opencl/gopromax.o framesync.o
OBJS-$(CONFIG_W3FDIF_FILTER) += vf_w3fdif.o
OBJS-$(CONFIG_WAVEFORM_FILTER) += vf_waveform.o
OBJS-$(CONFIG_WEAVE_FILTER) += vf_weave.o
diff --git a/libavfilter/allfilters.c b/libavfilter/allfilters.c
index 87c3661cf4..5a0d10fcd0 100644
--- a/libavfilter/allfilters.c
+++ b/libavfilter/allfilters.c
@@ -456,6 +456,7 @@ extern const AVFilter ff_vf_vignette;
extern const AVFilter ff_vf_vmafmotion;
extern const AVFilter ff_vf_vpp_qsv;
extern const AVFilter ff_vf_vstack;
+extern const AVFilter ff_vf_gopromax_opencl;
extern const AVFilter ff_vf_w3fdif;
extern const AVFilter ff_vf_waveform;
extern const AVFilter ff_vf_weave;
diff --git a/libavfilter/opencl/gopromax.cl b/libavfilter/opencl/gopromax.cl
new file mode 100644
index 0000000000..adea5d8677
--- /dev/null
+++ b/libavfilter/opencl/gopromax.cl
@@ -0,0 +1,305 @@
+/*
+ * Copyright (c) 2021 Ronan LE MEILLAT
+ *
+ * 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
+ */
+
+#define OVERLAP 64
+#define CUT 688
+#define BASESIZE 4096 //OVERLAP and CUT are based on this size
+
+
+#define FOV 360.0f
+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,
+};
+
+float2 rotate_cube_face(float2 uv, int rotation);
+int2 transpose_gopromax_overlap(int2 xy, int2 dim);
+float3 equirect_to_xyz(int2 xy,int2 size);
+float2 xyz_to_cube(float3 xyz, int *direction, int *face);
+float2 xyz_to_eac(float3 xyz, int2 size);
+
+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;
+}
+
+float3 equirect_to_xyz(int2 xy,int2 size)
+{
+ float3 xyz;
+ float phi = ((2.f * ((float)xy.x) + 0.5f) / ((float)size.x) - 1.f) * M_PI ;
+ float theta = ((2.f * ((float)xy.y) + 0.5f) / ((float)size.y) - 1.f) * M_PI_2;
+
+ xyz.x = cos(theta) * sin(phi);
+ xyz.y = sin(theta);
+ xyz.z = cos(theta) * cos(phi);
+
+ return xyz;
+}
+
+float2 xyz_to_cube(float3 xyz, int *direction, int *face)
+{
+ float phi = atan2(xyz.x, xyz.z);
+ float theta = asin(xyz.y);
+ float phi_norm, theta_threshold;
+ int face_rotation;
+ float2 uv;
+ //int direction;
+
+ if (phi >= -M_PI_4 && phi < M_PI_4) {
+ *direction = FRONT;
+ phi_norm = phi;
+ } else if (phi >= -(M_PI_2 + M_PI_4) && phi < -M_PI_4) {
+ *direction = LEFT;
+ phi_norm = phi + M_PI_2;
+ } else if (phi >= M_PI_4 && phi < M_PI_2 + M_PI_4) {
+ *direction = RIGHT;
+ phi_norm = phi - M_PI_2;
+ } else {
+ *direction = BACK;
+ phi_norm = phi + ((phi > 0.f) ? -M_PI : M_PI);
+ }
+
+ theta_threshold = atan(cos(phi_norm));
+ if (theta > theta_threshold) {
+ *direction = DOWN;
+ } else if (theta < -theta_threshold) {
+ *direction = UP;
+ }
+
+ 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;
+ face_rotation = ROT_0;
+ break;
+ case LEFT:
+ uv.x = -xyz.z / xyz.x;
+ uv.y = -xyz.y / xyz.x;
+ *face = TOP_LEFT;
+ face_rotation = ROT_0;
+ break;
+ case UP:
+ uv.x = -xyz.x / xyz.y;
+ uv.y = -xyz.z / xyz.y;
+ *face = BOTTOM_RIGHT;
+ face_rotation = ROT_270;
+ uv = rotate_cube_face(uv,face_rotation);
+ break;
+ case DOWN:
+ uv.x = xyz.x / xyz.y;
+ uv.y = -xyz.z / xyz.y;
+ *face = BOTTOM_LEFT;
+ face_rotation = ROT_270;
+ uv = rotate_cube_face(uv,face_rotation);
+ break;
+ case FRONT:
+ uv.x = xyz.x / xyz.z;
+ uv.y = xyz.y / xyz.z;
+ *face = TOP_MIDDLE;
+ face_rotation = ROT_0;
+ break;
+ case BACK:
+ uv.x = xyz.x / xyz.z;
+ uv.y = -xyz.y / xyz.z;
+ *face = BOTTOM_MIDDLE;
+ face_rotation = ROT_90;
+ uv = rotate_cube_face(uv,face_rotation);
+ break;
+ }
+
+ return uv;
+}
+
+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 direction, face;
+ int u_face, v_face;
+ float2 uv = xyz_to_cube(xyz,&direction,&face);
+
+ u_face = face % 3;
+ v_face = face / 3;
+ //eac expansion
+ uv.x = M_2_PI * atan(uv.x) + 0.5f;
+ uv.y = M_2_PI * atan(uv.y) + 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;
+
+ return uv;
+}
+
+const sampler_t sampler = (CLK_NORMALIZED_COORDS_FALSE |
+ CLK_ADDRESS_CLAMP_TO_EDGE |
+ CLK_FILTER_NEAREST);
+
+int2 transpose_gopromax_overlap(int2 xy, int2 dim)
+{
+ int2 ret;
+ int cut = dim.x*CUT/BASESIZE;
+ int overlap = dim.x*OVERLAP/BASESIZE;
+ if (xy.x<cut)
+ {
+ ret = xy;
+ }
+ else if ((xy.x>=cut) && (xy.x< (dim.x-cut)))
+ {
+ ret.x = xy.x+overlap;
+ ret.y = xy.y;
+ }
+ else
+ {
+ ret.x = xy.x+2*overlap;
+ ret.y = xy.y;
+ }
+ return ret;
+}
+__kernel void gopromax_equirectangular(__write_only image2d_t dst,
+ __read_only image2d_t gopromax_front,
+ __read_only image2d_t gopromax_rear)
+{
+
+ 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(gopromax_front);
+ int2 eac_size = (int2)(src_size.x-2*(src_size.x*OVERLAP/BASESIZE),dst_size.y);
+
+ int half_eight = src_size.y;
+
+ float3 xyz = equirect_to_xyz(loc,dst_size);
+
+ float2 uv = xyz_to_eac(xyz,eac_size);
+
+ int2 xy = convert_int2(floor(uv));
+
+ xy = transpose_gopromax_overlap(xy,eac_size);
+
+ if (xy.y<half_eight)
+ {
+ val = read_imagef(gopromax_front,sampler,xy);
+ }
+ else
+ {
+ val = read_imagef(gopromax_rear,sampler,(int2)(xy.x, (xy.y-half_eight)));
+ }
+
+ write_imagef(dst, loc, val);
+
+}
+
+__kernel void gopromax_stack(__write_only image2d_t dst,
+ __read_only image2d_t gopromax_front,
+ __read_only image2d_t gopromax_rear)
+{
+ const sampler_t sampler = (CLK_NORMALIZED_COORDS_FALSE |
+ CLK_ADDRESS_CLAMP_TO_EDGE |
+ CLK_FILTER_NEAREST);
+
+ float4 val;
+ int2 loc = (int2)(get_global_id(0), get_global_id(1));
+ int2 dst_size = get_image_dim(dst);
+ int half_height = dst_size.y / 2;
+ int cut0 = dst_size.x * CUT / (BASESIZE-2*OVERLAP);
+ int cut1 = dst_size.x - cut0;
+ int overlap = dst_size.x * OVERLAP / (BASESIZE-2*OVERLAP);
+
+ int x;
+ if (loc.x < (cut0-overlap))
+ {
+ x = loc.x;
+ }
+ else if ( (loc.x>=(cut0-overlap)) && ( loc.x < ( cut1 + overlap) ) )
+ {
+ x = loc.x + overlap;
+ }
+ else if ( loc.x >= ( cut1 - 2*overlap) )
+ {
+ x = loc.x + 2*overlap;
+ }
+
+ if (loc.y < half_height)
+ {
+ val = read_imagef(gopromax_front, sampler, (int2)(x, loc.y));
+ }
+ else
+ {
+ val = read_imagef(gopromax_rear, sampler, (int2)(x, loc.y-half_height));
+ }
+
+ write_imagef(dst, loc, val);
+}
diff --git a/libavfilter/opencl_source.h b/libavfilter/opencl_source.h
index 7e8133090e..66a205dcc7 100644
--- a/libavfilter/opencl_source.h
+++ b/libavfilter/opencl_source.h
@@ -32,5 +32,5 @@ extern const char *ff_opencl_source_tonemap;
extern const char *ff_opencl_source_transpose;
extern const char *ff_opencl_source_unsharp;
extern const char *ff_opencl_source_xfade;
-
+extern const char *ff_opencl_source_gopromax;
#endif /* AVFILTER_OPENCL_SOURCE_H */
diff --git a/libavfilter/vf_gopromax_opencl.c b/libavfilter/vf_gopromax_opencl.c
new file mode 100644
index 0000000000..e6394fb2fa
--- /dev/null
+++ b/libavfilter/vf_gopromax_opencl.c
@@ -0,0 +1,330 @@
+/*
+ * Copyright (c) 2021 Ronan LE MEILLAT
+ *
+ * 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 "framesync.h"
+#include "internal.h"
+#include "opencl.h"
+#include "opencl_source.h"
+#include "video.h"
+
+#define _WIDTH 5376
+#define _HEIGHT 2688
+#define OVERLAP 64
+#define CUT 688
+#define BASESIZE 4096 //OVERLAP and CUT are based on this size
+
+typedef struct GoProMaxOpenCLContext {
+ OpenCLFilterContext ocf;
+
+ int initialised;
+ cl_kernel kernel;
+ cl_command_queue command_queue;
+
+ FFFrameSync fs;
+
+ int nb_planes;
+ int x_subsample;
+ int y_subsample;
+ int alpha_separate;
+
+ int eac_output;
+} GoProMaxOpenCLContext;
+
+static int gopromax_opencl_load(AVFilterContext *avctx,
+ enum AVPixelFormat gopromax_front_format,
+ enum AVPixelFormat gopromax_rear_format)
+{
+ GoProMaxOpenCLContext *ctx = avctx->priv;
+ cl_int cle;
+ const char *source = ff_opencl_source_gopromax;
+ const char *kernel;
+ const AVPixFmtDescriptor *gopromax_front_desc, *gopromax_rear_desc;
+ int err, i, gopromax_front_planes, gopromax_rear_planes;
+
+ gopromax_front_desc = av_pix_fmt_desc_get(gopromax_front_format);
+ gopromax_rear_desc = av_pix_fmt_desc_get(gopromax_rear_format);
+ gopromax_front_planes = gopromax_rear_planes = 0;
+ for (i = 0; i < gopromax_front_desc->nb_components; i++)
+ gopromax_front_planes = FFMAX(gopromax_front_planes,
+ gopromax_front_desc->comp[i].plane + 1);
+ for (i = 0; i < gopromax_rear_desc->nb_components; i++)
+ gopromax_rear_planes = FFMAX(gopromax_rear_planes,
+ gopromax_rear_desc->comp[i].plane + 1);
+
+ ctx->nb_planes = gopromax_front_planes;
+ ctx->x_subsample = 1 << gopromax_front_desc->log2_chroma_w;
+ ctx->y_subsample = 1 << gopromax_front_desc->log2_chroma_h;
+
+
+ if (ctx->eac_output >0 )
+ {
+ kernel = "gopromax_stack";
+ }
+ else {
+ kernel = "gopromax_equirectangular";
+ }
+
+ av_log(avctx, AV_LOG_DEBUG, "Using kernel %s.\n", kernel);
+
+ err = ff_opencl_filter_load_program(avctx, &source, 1);
+ av_log(avctx, AV_LOG_VERBOSE,"OpenCL Kernel %s loaded err=%d\n",kernel,err);
+ if (err < 0)
+ goto fail;
+
+ ctx->command_queue = clCreateCommandQueue(ctx->ocf.hwctx->context,
+ ctx->ocf.hwctx->device_id,
+ 0, &cle);
+ av_log(avctx, AV_LOG_VERBOSE,"Leaving loading\n");
+ 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 *gopromax_front, *gotpromax_rear;
+ AVFrame *output;
+ cl_mem mem;
+ cl_int cle;//, x, y;
+ size_t global_work[2];
+ int kernel_arg = 0;
+ int err, plane;
+
+ err = ff_framesync_get_frame(fs, 0, &gopromax_front, 0);
+ if (err < 0)
+ return err;
+ err = ff_framesync_get_frame(fs, 1, &gotpromax_rear, 0);
+ if (err < 0)
+ return err;
+
+ if (!ctx->initialised) {
+ AVHWFramesContext *gopromax_front_fc =
+ (AVHWFramesContext*)gopromax_front->hw_frames_ctx->data;
+ AVHWFramesContext *gopromax_rear_fc =
+ (AVHWFramesContext*)gotpromax_rear->hw_frames_ctx->data;
+ err = gopromax_opencl_load(avctx, gopromax_front_fc->sw_format,
+ gopromax_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)gopromax_front->data[plane];
+ CL_SET_KERNEL_ARG(ctx->kernel, kernel_arg, cl_mem, &mem);
+ kernel_arg++;
+
+ mem = (cl_mem)gotpromax_rear->data[plane];
+ CL_SET_KERNEL_ARG(ctx->kernel, kernel_arg, cl_mem, &mem);
+ 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 %dx%d 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, gopromax_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;
+ av_log(avctx, AV_LOG_VERBOSE,"Setting output\n");
+ GoProMaxOpenCLContext *ctx = avctx->priv;
+ av_log(avctx, AV_LOG_VERBOSE,"Geting filtercontext\n");
+ AVFilterLink *inlink = avctx->inputs[0];
+ const AVPixFmtDescriptor *desc_in = av_pix_fmt_desc_get(inlink->format);
+
+ int height = avctx->inputs[0]->h;
+ int width = avctx->inputs[0]->w;
+ int err;
+
+ if (desc_in->log2_chroma_w != desc_in->log2_chroma_h) {
+ av_log(avctx, AV_LOG_ERROR, "Input format %s not supported.\n",
+ desc_in->name);
+ return AVERROR(EINVAL);
+ }
+
+ if (ctx->eac_output==0)
+ {
+ ctx->ocf.output_width = 4*height;
+ ctx->ocf.output_height = 2*height;
+ }
+ else
+ {
+ int overlap = width * OVERLAP / BASESIZE;
+ ctx->ocf.output_width = width - 2*overlap;
+ ctx->ocf.output_height = 2*height;
+ }
+
+ err = ff_opencl_filter_config_output(outlink);
+ av_log(avctx, AV_LOG_VERBOSE,"Output config ok w=%d h=%d err=%d\n",outlink->w, outlink->h, err);
+ if (err < 0)
+ return err;
+
+ err = ff_framesync_init_dualinput(&ctx->fs, avctx);
+ av_log(avctx, AV_LOG_VERBOSE,"Dualinput config ok err=%d\n",err);
+ 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[] = {
+ { "eac", "output Equiangular cubemap",
+ OFFSET(eac_output), AV_OPT_TYPE_INT, { .i64 = 0 }, 0, INT_MAX, .flags = FLAGS },
+ { NULL },
+};
+
+AVFILTER_DEFINE_CLASS(gopromax_opencl);
+
+static const AVFilterPad gopromax_opencl_inputs[] = {
+ {
+ .name = "gopromax_front",
+ .type = AVMEDIA_TYPE_VIDEO,
+ .config_props = &ff_opencl_filter_config_input,
+ },
+ {
+ .name = "gopromax_rear",
+ .type = AVMEDIA_TYPE_VIDEO,
+ .config_props = &ff_opencl_filter_config_input,
+ },
+ { NULL }
+};
+
+static const AVFilterPad gopromax_opencl_outputs[] = {
+ {
+ .name = "default",
+ .type = AVMEDIA_TYPE_VIDEO,
+ .config_props = &gopromax_opencl_config_output,
+ },
+ { NULL }
+};
+
+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,
+ .query_formats = &ff_opencl_filter_query_formats,
+ .activate = &gopromax_opencl_activate,
+ .inputs = gopromax_opencl_inputs,
+ .outputs = gopromax_opencl_outputs,
+ .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE,
+};
--
2.30.1 (Apple Git-130)
More information about the ffmpeg-devel
mailing list