| Message ID | 20240727043054.213199-2-aimingoff@pc.nifty.jp |
|---|---|
| State | New |
| Headers | show |
| Series | lavfi/vf_gopromax_opencl: add GoPor Max 360 video filter | expand |
| Context | Check | Description |
|---|---|---|
| andriy/make_x86 | fail | Make failed |
On 2024/07/27 13:30, TADANO Tokumei wrote: > Add an OpenCL filter for filtering GoPro Max native .360 files > into standard equirectangular or youtube equiangular cubemap (eac) > projection. > > The .360 file contains separated two video streams. > This filter combine two streams into single stream with standard > format. > --- > doc/filters.texi | 78 +++++++ > libavfilter/Makefile | 2 + > libavfilter/allfilters.c | 1 + > libavfilter/opencl/gopromax.cl | 280 ++++++++++++++++++++++++ > libavfilter/opencl_source.h | 1 + > libavfilter/vf_gopromax_opencl.c | 351 +++++++++++++++++++++++++++++++ > 6 files changed, 713 insertions(+) > create mode 100644 libavfilter/opencl/gopromax.cl > create mode 100644 libavfilter/vf_gopromax_opencl.c The patchwork failed, but it was caused by opencl.c (not by this patch): In file included from ./libavutil/common.h:48:0, from ./libavutil/avutil.h:301, from ./libavutil/opt.h:31, from libavdevice/sdl2.c:31: ./config.h:335:0: warning: 'HAVE_PTHREAD_SETNAME_NP' redefined #define HAVE_PTHREAD_SETNAME_NP 0 In file included from /usr/include/SDL2/SDL_stdinc.h:31:0, from /usr/include/SDL2/SDL_main.h:25, from /usr/include/SDL2/SDL.h:32, from libavdevice/sdl2.c:26: /usr/include/SDL2/SDL_config.h:186:0: note: this is the location of the previous definition #define HAVE_PTHREAD_SETNAME_NP 1 In file included from libavfilter/opencl.h:31:0, from libavfilter/opencl.c:26: ./libavutil/hwcontext_opencl.h:25:10: fatal error: CL/cl.h: No such file or directory #include ^~~~~~~~~ compilation terminated. make: *** [libavfilter/opencl.o] Error 1 > diff --git a/doc/filters.texi b/doc/filters.texi > index 2585d818ff..dd9ab47686 100644 > --- a/doc/filters.texi > +++ b/doc/filters.texi > @@ -27101,6 +27101,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 > > diff --git a/libavfilter/Makefile b/libavfilter/Makefile > index 63088e9286..9b5fa78920 100644 > --- a/libavfilter/Makefile > +++ b/libavfilter/Makefile > @@ -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 > diff --git a/libavfilter/allfilters.c b/libavfilter/allfilters.c > index 63600e9b58..5a517d8586 100644 > --- a/libavfilter/allfilters.c > +++ b/libavfilter/allfilters.c > @@ -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; > diff --git a/libavfilter/opencl/gopromax.cl b/libavfilter/opencl/gopromax.cl > new file mode 100644 > index 0000000000..64d2afe31c > --- /dev/null > +++ b/libavfilter/opencl/gopromax.cl > @@ -0,0 +1,280 @@ > +/* > + * 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; > + > + 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); > +} > 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..6867ca5cbb > --- /dev/null > +++ b/libavfilter/vf_gopromax_opencl.c > @@ -0,0 +1,351 @@ > +/* > + * 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 "internal.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, > +};
On Sun, Jul 28, 2024 at 01:42:09AM +0900, TADANO Tokumei wrote: > > On 2024/07/27 13:30, TADANO Tokumei wrote: > > Add an OpenCL filter for filtering GoPro Max native .360 files > > into standard equirectangular or youtube equiangular cubemap (eac) > > projection. > > > > The .360 file contains separated two video streams. > > This filter combine two streams into single stream with standard > > format. > > --- > > doc/filters.texi | 78 +++++++ > > libavfilter/Makefile | 2 + > > libavfilter/allfilters.c | 1 + > > libavfilter/opencl/gopromax.cl | 280 ++++++++++++++++++++++++ > > libavfilter/opencl_source.h | 1 + > > libavfilter/vf_gopromax_opencl.c | 351 +++++++++++++++++++++++++++++++ > > 6 files changed, 713 insertions(+) > > create mode 100644 libavfilter/opencl/gopromax.cl > > create mode 100644 libavfilter/vf_gopromax_opencl.c > > The patchwork failed, but it was caused by opencl.c (not by this patch): > > In file included from ./libavutil/common.h:48:0, > from ./libavutil/avutil.h:301, > from ./libavutil/opt.h:31, > from libavdevice/sdl2.c:31: > ./config.h:335:0: warning: 'HAVE_PTHREAD_SETNAME_NP' redefined > #define HAVE_PTHREAD_SETNAME_NP 0 > In file included from /usr/include/SDL2/SDL_stdinc.h:31:0, > from /usr/include/SDL2/SDL_main.h:25, > from /usr/include/SDL2/SDL.h:32, > from libavdevice/sdl2.c:26: > /usr/include/SDL2/SDL_config.h:186:0: note: this is the location of the previous definition > #define HAVE_PTHREAD_SETNAME_NP 1 > In file included from libavfilter/opencl.h:31:0, > from libavfilter/opencl.c:26: > ./libavutil/hwcontext_opencl.h:25:10: fatal error: CL/cl.h: No such file or directory > #include > ^~~~~~~~~ > compilation terminated. > make: *** [libavfilter/opencl.o] Error 1 with this patch it fails here on ubuntu: /usr/bin/ld: libavfilter/libavfilter.a(opencl.o): undefined reference to symbol 'clBuildProgram@@OPENCL_1.0' /usr/bin/ld: /usr/local/cuda/targets/x86_64-linux/lib/libOpenCL.so.1: error adding symbols: DSO missing from command line thx [...]
On 2024/07/28 18:26, Michael Niedermayer wrote: > On Sun, Jul 28, 2024 at 01:42:09AM +0900, TADANO Tokumei wrote: >> >> On 2024/07/27 13:30, TADANO Tokumei wrote: >>> Add an OpenCL filter for filtering GoPro Max native .360 files >>> into standard equirectangular or youtube equiangular cubemap (eac) >>> projection. >>> >>> The .360 file contains separated two video streams. >>> This filter combine two streams into single stream with standard >>> format. >>> --- >>> doc/filters.texi | 78 +++++++ >>> libavfilter/Makefile | 2 + >>> libavfilter/allfilters.c | 1 + >>> libavfilter/opencl/gopromax.cl | 280 ++++++++++++++++++++++++ >>> libavfilter/opencl_source.h | 1 + >>> libavfilter/vf_gopromax_opencl.c | 351 +++++++++++++++++++++++++++++++ >>> 6 files changed, 713 insertions(+) >>> create mode 100644 libavfilter/opencl/gopromax.cl >>> create mode 100644 libavfilter/vf_gopromax_opencl.c >> >> The patchwork failed, but it was caused by opencl.c (not by this patch): >> >> In file included from ./libavutil/common.h:48:0, >> from ./libavutil/avutil.h:301, >> from ./libavutil/opt.h:31, >> from libavdevice/sdl2.c:31: >> ./config.h:335:0: warning: 'HAVE_PTHREAD_SETNAME_NP' redefined >> #define HAVE_PTHREAD_SETNAME_NP 0 >> In file included from /usr/include/SDL2/SDL_stdinc.h:31:0, >> from /usr/include/SDL2/SDL_main.h:25, >> from /usr/include/SDL2/SDL.h:32, >> from libavdevice/sdl2.c:26: >> /usr/include/SDL2/SDL_config.h:186:0: note: this is the location of the previous definition >> #define HAVE_PTHREAD_SETNAME_NP 1 >> In file included from libavfilter/opencl.h:31:0, >> from libavfilter/opencl.c:26: >> ./libavutil/hwcontext_opencl.h:25:10: fatal error: CL/cl.h: No such file or directory >> #include >> ^~~~~~~~~ >> compilation terminated. >> make: *** [libavfilter/opencl.o] Error 1 > > with this patch it fails here on ubuntu: > /usr/bin/ld: libavfilter/libavfilter.a(opencl.o): undefined reference to symbol 'clBuildProgram@@OPENCL_1.0' > /usr/bin/ld: /usr/local/cuda/targets/x86_64-linux/lib/libOpenCL.so.1: error adding symbols: DSO missing from command line > > thx I tested on Ubuntu 22.04, and it works fine. As message shows, it seems the error is not related to this patch but OpenCL library. > [...] > > > _______________________________________________ > ffmpeg-devel mailing list > ffmpeg-devel@ffmpeg.org > https://ffmpeg.org/mailman/listinfo/ffmpeg-devel > > To unsubscribe, visit link above, or email > ffmpeg-devel-request@ffmpeg.org with subject "unsubscribe".
On 2024/07/29 1:30, TADANO Tokumei wrote: > On 2024/07/28 18:26, Michael Niedermayer wrote: >> On Sun, Jul 28, 2024 at 01:42:09AM +0900, TADANO Tokumei wrote: >>> >>> On 2024/07/27 13:30, TADANO Tokumei wrote: >>>> Add an OpenCL filter for filtering GoPro Max native .360 files >>>> into standard equirectangular or youtube equiangular cubemap (eac) >>>> projection. >>>> >>>> The .360 file contains separated two video streams. >>>> This filter combine two streams into single stream with standard >>>> format. >>>> --- >>>> doc/filters.texi | 78 +++++++ >>>> libavfilter/Makefile | 2 + >>>> libavfilter/allfilters.c | 1 + >>>> libavfilter/opencl/gopromax.cl | 280 ++++++++++++++++++++++++ >>>> libavfilter/opencl_source.h | 1 + >>>> libavfilter/vf_gopromax_opencl.c | 351 +++++++++++++++++++++++++++++++ >>>> 6 files changed, 713 insertions(+) >>>> create mode 100644 libavfilter/opencl/gopromax.cl >>>> create mode 100644 libavfilter/vf_gopromax_opencl.c >>> >>> The patchwork failed, but it was caused by opencl.c (not by this patch): >>> >>> In file included from ./libavutil/common.h:48:0, >>> from ./libavutil/avutil.h:301, >>> from ./libavutil/opt.h:31, >>> from libavdevice/sdl2.c:31: >>> ./config.h:335:0: warning: 'HAVE_PTHREAD_SETNAME_NP' redefined >>> #define HAVE_PTHREAD_SETNAME_NP 0 >>> In file included from /usr/include/SDL2/SDL_stdinc.h:31:0, >>> from /usr/include/SDL2/SDL_main.h:25, >>> from /usr/include/SDL2/SDL.h:32, >>> from libavdevice/sdl2.c:26: >>> /usr/include/SDL2/SDL_config.h:186:0: note: this is the location of the previous definition >>> #define HAVE_PTHREAD_SETNAME_NP 1 >>> In file included from libavfilter/opencl.h:31:0, >>> from libavfilter/opencl.c:26: >>> ./libavutil/hwcontext_opencl.h:25:10: fatal error: CL/cl.h: No such file or directory >>> #include >>> ^~~~~~~~~ >>> compilation terminated. >>> make: *** [libavfilter/opencl.o] Error 1 >> >> with this patch it fails here on ubuntu: >> /usr/bin/ld: libavfilter/libavfilter.a(opencl.o): undefined reference to symbol 'clBuildProgram@@OPENCL_1.0' >> /usr/bin/ld: /usr/local/cuda/targets/x86_64-linux/lib/libOpenCL.so.1: error adding symbols: DSO missing from command line >> >> thx > > I tested on Ubuntu 22.04, and it works fine. > As message shows, it seems the error is not related to this patch but OpenCL library. I found I have to add `gopromax_opencl_filter_deps="opencl"` in configure file. I'll amend the patch later. >> [...] >> >> >> _______________________________________________ >> ffmpeg-devel mailing list >> ffmpeg-devel@ffmpeg.org >> https://ffmpeg.org/mailman/listinfo/ffmpeg-devel >> >> To unsubscribe, visit link above, or email >> ffmpeg-devel-request@ffmpeg.org with subject "unsubscribe". > _______________________________________________ > ffmpeg-devel mailing list > ffmpeg-devel@ffmpeg.org > https://ffmpeg.org/mailman/listinfo/ffmpeg-devel > > To unsubscribe, visit link above, or email > ffmpeg-devel-request@ffmpeg.org with subject "unsubscribe".
On Mon, Jul 29, 2024 at 01:30:34AM +0900, TADANO Tokumei wrote: > On 2024/07/28 18:26, Michael Niedermayer wrote: > > On Sun, Jul 28, 2024 at 01:42:09AM +0900, TADANO Tokumei wrote: > > > > > > On 2024/07/27 13:30, TADANO Tokumei wrote: > > > > Add an OpenCL filter for filtering GoPro Max native .360 files > > > > into standard equirectangular or youtube equiangular cubemap (eac) > > > > projection. > > > > > > > > The .360 file contains separated two video streams. > > > > This filter combine two streams into single stream with standard > > > > format. > > > > --- > > > > doc/filters.texi | 78 +++++++ > > > > libavfilter/Makefile | 2 + > > > > libavfilter/allfilters.c | 1 + > > > > libavfilter/opencl/gopromax.cl | 280 ++++++++++++++++++++++++ > > > > libavfilter/opencl_source.h | 1 + > > > > libavfilter/vf_gopromax_opencl.c | 351 +++++++++++++++++++++++++++++++ > > > > 6 files changed, 713 insertions(+) > > > > create mode 100644 libavfilter/opencl/gopromax.cl > > > > create mode 100644 libavfilter/vf_gopromax_opencl.c > > > > > > The patchwork failed, but it was caused by opencl.c (not by this patch): > > > > > > In file included from ./libavutil/common.h:48:0, > > > from ./libavutil/avutil.h:301, > > > from ./libavutil/opt.h:31, > > > from libavdevice/sdl2.c:31: > > > ./config.h:335:0: warning: 'HAVE_PTHREAD_SETNAME_NP' redefined > > > #define HAVE_PTHREAD_SETNAME_NP 0 > > > In file included from /usr/include/SDL2/SDL_stdinc.h:31:0, > > > from /usr/include/SDL2/SDL_main.h:25, > > > from /usr/include/SDL2/SDL.h:32, > > > from libavdevice/sdl2.c:26: > > > /usr/include/SDL2/SDL_config.h:186:0: note: this is the location of the previous definition > > > #define HAVE_PTHREAD_SETNAME_NP 1 > > > In file included from libavfilter/opencl.h:31:0, > > > from libavfilter/opencl.c:26: > > > ./libavutil/hwcontext_opencl.h:25:10: fatal error: CL/cl.h: No such file or directory > > > #include > > > ^~~~~~~~~ > > > compilation terminated. > > > make: *** [libavfilter/opencl.o] Error 1 > > > > with this patch it fails here on ubuntu: > > /usr/bin/ld: libavfilter/libavfilter.a(opencl.o): undefined reference to symbol 'clBuildProgram@@OPENCL_1.0' > > /usr/bin/ld: /usr/local/cuda/targets/x86_64-linux/lib/libOpenCL.so.1: error adding symbols: DSO missing from command line > > > > thx > > I tested on Ubuntu 22.04, and it works fine. My failure was on a 20.04 thx [...]
On 2024/07/31 5:44, Michael Niedermayer wrote: > On Mon, Jul 29, 2024 at 01:30:34AM +0900, TADANO Tokumei wrote: >> On 2024/07/28 18:26, Michael Niedermayer wrote: >>> On Sun, Jul 28, 2024 at 01:42:09AM +0900, TADANO Tokumei wrote: >>>> >>>> On 2024/07/27 13:30, TADANO Tokumei wrote: >>>>> Add an OpenCL filter for filtering GoPro Max native .360 files >>>>> into standard equirectangular or youtube equiangular cubemap (eac) >>>>> projection. >>>>> >>>>> The .360 file contains separated two video streams. >>>>> This filter combine two streams into single stream with standard >>>>> format. >>>>> --- >>>>> doc/filters.texi | 78 +++++++ >>>>> libavfilter/Makefile | 2 + >>>>> libavfilter/allfilters.c | 1 + >>>>> libavfilter/opencl/gopromax.cl | 280 ++++++++++++++++++++++++ >>>>> libavfilter/opencl_source.h | 1 + >>>>> libavfilter/vf_gopromax_opencl.c | 351 +++++++++++++++++++++++++++++++ >>>>> 6 files changed, 713 insertions(+) >>>>> create mode 100644 libavfilter/opencl/gopromax.cl >>>>> create mode 100644 libavfilter/vf_gopromax_opencl.c >>>> >>>> The patchwork failed, but it was caused by opencl.c (not by this patch): >>>> >>>> In file included from ./libavutil/common.h:48:0, >>>> from ./libavutil/avutil.h:301, >>>> from ./libavutil/opt.h:31, >>>> from libavdevice/sdl2.c:31: >>>> ./config.h:335:0: warning: 'HAVE_PTHREAD_SETNAME_NP' redefined >>>> #define HAVE_PTHREAD_SETNAME_NP 0 >>>> In file included from /usr/include/SDL2/SDL_stdinc.h:31:0, >>>> from /usr/include/SDL2/SDL_main.h:25, >>>> from /usr/include/SDL2/SDL.h:32, >>>> from libavdevice/sdl2.c:26: >>>> /usr/include/SDL2/SDL_config.h:186:0: note: this is the location of the previous definition >>>> #define HAVE_PTHREAD_SETNAME_NP 1 >>>> In file included from libavfilter/opencl.h:31:0, >>>> from libavfilter/opencl.c:26: >>>> ./libavutil/hwcontext_opencl.h:25:10: fatal error: CL/cl.h: No such file or directory >>>> #include >>>> ^~~~~~~~~ >>>> compilation terminated. >>>> make: *** [libavfilter/opencl.o] Error 1 >>> >>> with this patch it fails here on ubuntu: >>> /usr/bin/ld: libavfilter/libavfilter.a(opencl.o): undefined reference to symbol 'clBuildProgram@@OPENCL_1.0' >>> /usr/bin/ld: /usr/local/cuda/targets/x86_64-linux/lib/libOpenCL.so.1: error adding symbols: DSO missing from command line >>> >>> thx >> >> I tested on Ubuntu 22.04, and it works fine. > > My failure was on a 20.04 > > thx Did you try v3 patch? Anyway, your OpenCL library seems to be Nvidia's proprietary one. I'm not sure, but I guess the library has some incompatibility like: https://github.com/OpenKinect/libfreenect2/issues/804#issuecomment-286515571 > [...] > > > _______________________________________________ > ffmpeg-devel mailing list > ffmpeg-devel@ffmpeg.org > https://ffmpeg.org/mailman/listinfo/ffmpeg-devel > > To unsubscribe, visit link above, or email > ffmpeg-devel-request@ffmpeg.org with subject "unsubscribe".
On Wed, Jul 31, 2024 at 11:45:29PM +0900, TADANO Tokumei wrote: > On 2024/07/31 5:44, Michael Niedermayer wrote: > > On Mon, Jul 29, 2024 at 01:30:34AM +0900, TADANO Tokumei wrote: > > > On 2024/07/28 18:26, Michael Niedermayer wrote: > > > > On Sun, Jul 28, 2024 at 01:42:09AM +0900, TADANO Tokumei wrote: > > > > > > > > > > On 2024/07/27 13:30, TADANO Tokumei wrote: > > > > > > Add an OpenCL filter for filtering GoPro Max native .360 files > > > > > > into standard equirectangular or youtube equiangular cubemap (eac) > > > > > > projection. > > > > > > > > > > > > The .360 file contains separated two video streams. > > > > > > This filter combine two streams into single stream with standard > > > > > > format. > > > > > > --- > > > > > > doc/filters.texi | 78 +++++++ > > > > > > libavfilter/Makefile | 2 + > > > > > > libavfilter/allfilters.c | 1 + > > > > > > libavfilter/opencl/gopromax.cl | 280 ++++++++++++++++++++++++ > > > > > > libavfilter/opencl_source.h | 1 + > > > > > > libavfilter/vf_gopromax_opencl.c | 351 +++++++++++++++++++++++++++++++ > > > > > > 6 files changed, 713 insertions(+) > > > > > > create mode 100644 libavfilter/opencl/gopromax.cl > > > > > > create mode 100644 libavfilter/vf_gopromax_opencl.c > > > > > > > > > > The patchwork failed, but it was caused by opencl.c (not by this patch): > > > > > > > > > > In file included from ./libavutil/common.h:48:0, > > > > > from ./libavutil/avutil.h:301, > > > > > from ./libavutil/opt.h:31, > > > > > from libavdevice/sdl2.c:31: > > > > > ./config.h:335:0: warning: 'HAVE_PTHREAD_SETNAME_NP' redefined > > > > > #define HAVE_PTHREAD_SETNAME_NP 0 > > > > > In file included from /usr/include/SDL2/SDL_stdinc.h:31:0, > > > > > from /usr/include/SDL2/SDL_main.h:25, > > > > > from /usr/include/SDL2/SDL.h:32, > > > > > from libavdevice/sdl2.c:26: > > > > > /usr/include/SDL2/SDL_config.h:186:0: note: this is the location of the previous definition > > > > > #define HAVE_PTHREAD_SETNAME_NP 1 > > > > > In file included from libavfilter/opencl.h:31:0, > > > > > from libavfilter/opencl.c:26: > > > > > ./libavutil/hwcontext_opencl.h:25:10: fatal error: CL/cl.h: No such file or directory > > > > > #include > > > > > ^~~~~~~~~ > > > > > compilation terminated. > > > > > make: *** [libavfilter/opencl.o] Error 1 > > > > > > > > with this patch it fails here on ubuntu: > > > > /usr/bin/ld: libavfilter/libavfilter.a(opencl.o): undefined reference to symbol 'clBuildProgram@@OPENCL_1.0' > > > > /usr/bin/ld: /usr/local/cuda/targets/x86_64-linux/lib/libOpenCL.so.1: error adding symbols: DSO missing from command line > > > > > > > > thx > > > > > > I tested on Ubuntu 22.04, and it works fine. > > > > My failure was on a 20.04 > > > > thx > > Did you try v3 patch? i only saw v3 after i tested v2 v3 no longer fails thx [...]
diff --git a/doc/filters.texi b/doc/filters.texi index 2585d818ff..dd9ab47686 100644 --- a/doc/filters.texi +++ b/doc/filters.texi @@ -27101,6 +27101,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 diff --git a/libavfilter/Makefile b/libavfilter/Makefile index 63088e9286..9b5fa78920 100644 --- a/libavfilter/Makefile +++ b/libavfilter/Makefile @@ -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 diff --git a/libavfilter/allfilters.c b/libavfilter/allfilters.c index 63600e9b58..5a517d8586 100644 --- a/libavfilter/allfilters.c +++ b/libavfilter/allfilters.c @@ -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; diff --git a/libavfilter/opencl/gopromax.cl b/libavfilter/opencl/gopromax.cl new file mode 100644 index 0000000000..64d2afe31c --- /dev/null +++ b/libavfilter/opencl/gopromax.cl @@ -0,0 +1,280 @@ +/* + * 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; + + 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); +} 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..6867ca5cbb --- /dev/null +++ b/libavfilter/vf_gopromax_opencl.c @@ -0,0 +1,351 @@ +/* + * 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 "internal.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, +};