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 | 1 + configure | 1 + doc/filters.texi | 22 ++ 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, 686 insertions(+) create mode 100644 libavfilter/opencl/gopromax.cl create mode 100644 libavfilter/vf_gopromax_opencl.c diff --git a/Changelog b/Changelog index 2348c65173..ceaf75c649 100644 --- a/Changelog +++ b/Changelog @@ -20,6 +20,7 @@ version <next>: - mjpeg_nvjpeg NVIDIA JPEG encoder - yuvrgb_cuda filter - gopromax2 GoPro Max 2 CPU stitch filter +- gopromax_opencl GoPro Max 1 stitch filter version 9.0: - Extend AMF Color Converter (vf_vpp_amf) HDR capabilities diff --git a/configure b/configure index 83a4570b8c..81eb1c841a 100755 --- a/configure +++ b/configure @@ -4135,6 +4135,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 14efd1d1b2..b8872611b7 100644 --- a/doc/filters.texi +++ b/doc/filters.texi @@ -27648,6 +27648,28 @@ JPEG (full) range @section yuvrgb_cuda @anchor{gopromax2} @section gopromax2 +@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. Max 2 uses @ref{gopromax2}. + +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 + Convert the two GoPro Max 2 .360 EAC tracks to equirectangular projection on the CPU. diff --git a/libavfilter/Makefile b/libavfilter/Makefile index 4b65c28314..065723d687 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_GOPROMAX2_FILTER) += vf_gopromax2.o framesync.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 diff --git a/libavfilter/allfilters.c b/libavfilter/allfilters.c index 4ee16d9f51..0c26769757 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_gopromax2; +extern const FFFilter ff_vf_gopromax_opencl; extern const FFFilter ff_vf_gradfun; extern const FFFilter ff_vf_graphmonitor; extern const FFFilter ff_vf_grayworld; 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.43.0 _______________________________________________ ffmpeg-devel mailing list -- [email protected] To unsubscribe send an email to [email protected]
