From patchwork Thu Sep 19 14:21:16 2024 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: TADANO Tokumei X-Patchwork-Id: 51650 Delivered-To: ffmpegpatchwork2@gmail.com Received: by 2002:a59:d154:0:b0:48e:c0f8:d0de with SMTP id bt20csp345354vqb; Thu, 19 Sep 2024 07:21:55 -0700 (PDT) X-Forwarded-Encrypted: i=2; AJvYcCXwFadhKHkhwRfvwcjgRoOYCnSlnCbvxZYDa6MhPfM3F8c4BvjAmisduiz70BHXIBDF3l6Yq4Rlt863gf50x2np@gmail.com X-Google-Smtp-Source: AGHT+IFs4iohSvpCCGvhnagqRxQPn4UZl77/nrapAq0fTxHq0mMlbtGDyfTBtMedAkvzukWUXihU X-Received: by 2002:a17:907:1c2a:b0:a8a:811e:3fd5 with SMTP id a640c23a62f3a-a9029448c6amr1252347266b.4.1726755715215; Thu, 19 Sep 2024 07:21:55 -0700 (PDT) ARC-Seal: i=1; a=rsa-sha256; t=1726755715; cv=none; d=google.com; s=arc-20240605; b=RSsQrXF9CUKmuEVZzWZG+C9QRPgd+SGIKGvVwJYm3F7eaYbxf3C3HKFSCq1qhnAPxl YgsCA4aiD6sl/GBAufnXYyt9+WTIWUbRyw5m5brGm4b2A2UzxEqqE6errnG7y9CH+QuV uVpbQhHPZcv7lie/Acfqfpmg4IDRmxjiBlBNNolRnmmisIgtWF2OEsHhvkPivhYg7XLp sjLmDpEi6zJ8La31x/4JxuDnQ+8DAm/nTrD6ALc5eu64dUGOPJYEWnvTMC5cMHipMkIY d+vNOm4gRq4gl/N1rvaoy89bIvSmq2lcbCl5nzfW/gxhXEKyowEJ4QrMvWZepHxs3nBB QQAA== ARC-Message-Signature: i=1; a=rsa-sha256; c=relaxed/relaxed; d=google.com; s=arc-20240605; h=sender:errors-to:content-transfer-encoding:cc:reply-to :list-subscribe:list-help:list-post:list-archive:list-unsubscribe :list-id:precedence:subject:dkim-signature:mime-version:references :in-reply-to:message-id:date:to:from:delivered-to; bh=cn0OQEaPJYbLySCZzd+9FiSrfRuUfvqbV83BMWy9i4Q=; fh=ssXC39eKeah6YcPRlt8jBK+nQIpHohGaLsoJm2kZFpQ=; b=jZQTMmT0aSPN4EEYfl07mf7pwj5FF4KKG6qp3L4Y4+2b/qhd1G5vpdsVIVx8Lnskdt Hi+DUaYqvqa0qvODtSu/RkVRMNvgsgquGeb0LaBvmOEIQ7l4CpBAe7iP42yjwRKyGjl5 FZAKTCC14Qv/gr6Yd85tNeonYKKTlCr5jprdO+LRTmRkkV+GyLi3Y8CfoQI8b7jbOuEl VAyMU0zhpBk5AxrIxtAi2nOC3llPvwUReFsmoU3KJIixu3fHDdqcvilGI1Bbcejeak3E wP5rXl6cyrd2EL6y4IX1zr+rKh6oEhqiI5T1YO4IK+/TjZ4bu+jnUZ0857427hMNRVqY XVIQ==; dara=google.com ARC-Authentication-Results: i=1; mx.google.com; dkim=neutral (body hash did not verify) header.i=@pc.nifty.jp header.s=default-1th84yt82rvi header.b=TzlV6xHi; spf=pass (google.com: domain of ffmpeg-devel-bounces@ffmpeg.org designates 79.124.17.100 as permitted sender) smtp.mailfrom=ffmpeg-devel-bounces@ffmpeg.org; dmarc=fail (p=NONE sp=NONE dis=NONE) header.from=pc.nifty.jp Return-Path: Received: from ffbox0-bg.mplayerhq.hu (ffbox0-bg.ffmpeg.org. [79.124.17.100]) by mx.google.com with ESMTP id a640c23a62f3a-a90612e782esi792825566b.499.2024.09.19.07.21.54; Thu, 19 Sep 2024 07:21:55 -0700 (PDT) Received-SPF: pass (google.com: domain of ffmpeg-devel-bounces@ffmpeg.org designates 79.124.17.100 as permitted sender) client-ip=79.124.17.100; Authentication-Results: mx.google.com; dkim=neutral (body hash did not verify) header.i=@pc.nifty.jp header.s=default-1th84yt82rvi header.b=TzlV6xHi; spf=pass (google.com: domain of ffmpeg-devel-bounces@ffmpeg.org designates 79.124.17.100 as permitted sender) smtp.mailfrom=ffmpeg-devel-bounces@ffmpeg.org; dmarc=fail (p=NONE sp=NONE dis=NONE) header.from=pc.nifty.jp Received: from [127.0.1.1] (localhost [127.0.0.1]) by ffbox0-bg.mplayerhq.hu (Postfix) with ESMTP id DAAEC68DBB8; Thu, 19 Sep 2024 17:21:47 +0300 (EEST) X-Original-To: ffmpeg-devel@ffmpeg.org Delivered-To: ffmpeg-devel@ffmpeg.org Received: from mta-snd-e02.mail.nifty.com (mta-snd-e02.mail.nifty.com [106.153.227.114]) by ffbox0-bg.mplayerhq.hu (Postfix) with ESMTPS id 8B2A968DB8A for ; Thu, 19 Sep 2024 17:21:39 +0300 (EEST) Received: from localhost.localdomain by mta-snd-e02.mail.nifty.com with ESMTP id <20240919142136014.BNQV.44461.localhost.localdomain@nifty.com>; Thu, 19 Sep 2024 23:21:36 +0900 From: TADANO Tokumei To: ffmpeg-devel@ffmpeg.org Date: Thu, 19 Sep 2024 23:21:16 +0900 Message-Id: <20240919142116.239725-2-aimingoff@pc.nifty.jp> X-Mailer: git-send-email 2.39.5 In-Reply-To: <20240919142116.239725-1-aimingoff@pc.nifty.jp> References: <20240919142116.239725-1-aimingoff@pc.nifty.jp> MIME-Version: 1.0 DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=pc.nifty.jp; s=default-1th84yt82rvi; t=1726755696; bh=yAARVzBQV0tAAfAImzy+khVOb11oh+/SE6SfkNkfwWE=; h=From:To:Cc:Subject:Date:In-Reply-To:References; b=TzlV6xHiSY8NIN/vO9OM8H2HOJ7gzo//aDVPB4uk7YHPBjRv1ryvpgy0YD9qkgJHiqqU+djc dEVG8o6V9vqGhQWdjLRYOFy+Guvc122KHsnRTFiNMOom1u+HZFiC8bN78+Q/wMR/hdIjrxx4q5 W1QWl9TCKuurOFt3bTDeh2D0AKyVQbvLyb2WiCzkp59A8DZgAUovjtq2sFco0jGC5iFZcdBqXp 0lR5TV27tmrcAIfA/rfKkoJCN99ExTP79RFc+Nl4xEcZRmHxB7yEbqCmx/ubjbwd5DPQVYc7Uz ZAdg7S/Ly4Nx4xO7WaqAOm0IWzcnjvthX8qwWjZiRT4FtHGg== Subject: [FFmpeg-devel] [PATCH v5 1/1] lavfi/vf_gopromax_opencl: add GoPro Max 360 video filter X-BeenThere: ffmpeg-devel@ffmpeg.org X-Mailman-Version: 2.1.29 Precedence: list List-Id: FFmpeg development discussions and patches List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , Reply-To: FFmpeg development discussions and patches Cc: TADANO Tokumei Errors-To: ffmpeg-devel-bounces@ffmpeg.org Sender: "ffmpeg-devel" X-TUID: 7T3+CaIlDkY6 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. --- configure | 1 + doc/filters.texi | 78 +++++++ libavfilter/Makefile | 2 + libavfilter/allfilters.c | 1 + libavfilter/opencl/gopromax.cl | 282 +++++++++++++++++++++++++ libavfilter/opencl_source.h | 1 + libavfilter/vf_gopromax_opencl.c | 350 +++++++++++++++++++++++++++++++ 7 files changed, 715 insertions(+) create mode 100644 libavfilter/opencl/gopromax.cl create mode 100644 libavfilter/vf_gopromax_opencl.c diff --git a/configure b/configure index d872213af7..83afcea5c6 100755 --- a/configure +++ b/configure @@ -3886,6 +3886,7 @@ frei0r_src_filter_deps="frei0r" fspp_filter_deps="gpl" fsync_filter_deps="avformat" gblur_vulkan_filter_deps="vulkan spirv_compiler" +gopromax_opencl_filter_deps="opencl" hflip_vulkan_filter_deps="vulkan spirv_compiler" histeq_filter_deps="gpl" hqdn3d_filter_deps="gpl" diff --git a/doc/filters.texi b/doc/filters.texi index db2f4b7ea7..02d4ffa753 100644 --- a/doc/filters.texi +++ b/doc/filters.texi @@ -27201,6 +27201,84 @@ Apply dilation filter with threshold0 set to 30, threshold1 set 40, threshold2 s @end example @end itemize +@anchor{gopromax_opencl} +@section gopromax_opencl + +Apply transformation of the two GoPro Max video streams to equirectangular or equiangular-cubemap projection. + +This filter is designed to use directly GoPro .360 files. +Native .360 files are sort of EAC files, in fact the front and rear lenses streams are the top and the bottom of the EAC projection. + +The .360 file contains two video streams. +Most of cases, one is stream #0:0, and the other is stream #0:5. +Please check actual stream number with @code{ffprobe} command. +This filter combine two streams to single stream. + +The .360 contains also 2x64 bits of overlapped area. +The filter blends overlapped images in these two areas. + +The filter accepts the following options: + +@table @option + +@item output +Set format of the output video. + +Available formats: + +@table @samp + +@item e +@item equirect +Equirectangular projection. + +@item eac +Equi-Angular Cubemap. + +@end table + +Default is @code{equirect}. + +@item w +@item h +Set the output video resolution. + +Default resolution depends on formats. + +@item overlap +Set number of overlapped pixels on input .360 video. + +No need to specify this option for native .360 video file. +This option is for rescaled video or future video format change. + +Default is @code{64}. + +@end table + +@subsection Example + +@itemize +@item +Convert .360 to Equirectangular projection. +@example +-i INPUT -filter_complex '[0:0]hwupload[a], [0:5]hwupload[b], [a][b]gopromax_opencl=w=4096:h=2048, hwdownload, format=yuvj420p' -map 0:a:0 -c:a copy OUTPUT +@end example + +Two video streams (#0:0 and #0:5) are combined and converted to default equirectangular projection with specified resolution. +First audio stream (GoPro AAC) is copied with the video stream. + +@item +Convert .360 to Equi-Angular Cubemap projection. +@example +-i INPUT -filter_complex '[0:0]hwupload[a], [0:5]hwupload[b], [a][b]gopromax_opencl=eac, hwdownload, format=yuvj420p, v360=eac:c3x2:w=1344:h=896' -map 0:1 -map 0:3 -c:a copy -c:u copy OUTPUT +@end example + +Two video streams (#0:0 and #0:5) are combined and converted to equi-angular cubemap projection, +then it is converted to c3x2 cubemap projection and shrunk by v360 filter. +Stream #0:1 (GoPro AAC) and stream #0:3 (GoPro MET) are copied with the video stream. + +@end itemize + @anchor{nlmeans_opencl} @section nlmeans_opencl diff --git a/libavfilter/Makefile b/libavfilter/Makefile index 91487afb21..bbbb234c1b 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 9819f0f95b..dc89f104d2 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..440469e2b8 --- /dev/null +++ b/libavfilter/opencl/gopromax.cl @@ -0,0 +1,282 @@ +/* + * Copyright (c) 2021 Ronan LE MEILLAT + * Copyright (c) 2024 TADANO Tokumei + * + * This file is part of FFmpeg. + * + * FFmpeg is free software; you can redistribute it and/or + * modify it under the terms of the GNU Lesser General Public + * License as published by the Free Software Foundation; either + * version 2.1 of the License, or (at your option) any later version. + * + * FFmpeg is distributed in the hope that it will be useful, + * but WITHOUT ANY WARRANTY; without even the implied warranty of + * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU + * Lesser General Public License for more details. + * + * You should have received a copy of the GNU Lesser General Public + * License along with FFmpeg; if not, write to the Free Software + * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA + */ + +enum Faces { + TOP_LEFT, + TOP_MIDDLE, + TOP_RIGHT, + BOTTOM_LEFT, + BOTTOM_MIDDLE, + BOTTOM_RIGHT, + NB_FACES, +}; + +enum Direction { + RIGHT, + LEFT, + UP, + DOWN, + FRONT, + BACK, + NB_DIRECTIONS, +}; + +enum Rotation { + ROT_0, + ROT_90, + ROT_180, + ROT_270, + NB_ROTATIONS, +}; + +static float2 rotate_cube_face(float2 uv, int rotation) +{ + float2 ret_uv; + + switch (rotation) { + case ROT_0: + ret_uv = uv; + break; + case ROT_90: + ret_uv.x = -uv.y; + ret_uv.y = uv.x; + break; + case ROT_180: + ret_uv.x = -uv.x; + ret_uv.y = -uv.y; + break; + case ROT_270: + ret_uv.x = uv.y; + ret_uv.y = -uv.x; + break; + } + + return ret_uv; +} + +static float3 equirect_to_xyz(int2 xy, int2 size) +{ + float3 xyz; + float phi = ((2.f * ((float)xy.x) + 1.f) / ((float)size.x) - 1.f) * M_PI_F ; + float theta = ((2.f * ((float)xy.y) + 1.f) / ((float)size.y) - 1.f) * M_PI_2_F; + + xyz.x = cos(theta) * sin(phi); + xyz.y = sin(theta); + xyz.z = cos(theta) * cos(phi); + + return xyz; +} + +static float2 xyz_to_cube(float3 xyz, int *face) +{ + float phi = atan2(xyz.x, xyz.z); + float theta = asin(xyz.y); + float phi_norm, theta_threshold; + float2 uv; + int direction; + + if (phi >= -M_PI_4_F && phi < M_PI_4_F) { + direction = FRONT; + phi_norm = phi; + } else if (phi >= -(M_PI_2_F + M_PI_4_F) && phi < -M_PI_4_F) { + direction = LEFT; + phi_norm = phi + M_PI_2_F; + } else if (phi >= M_PI_4_F && phi < M_PI_2_F + M_PI_4_F) { + direction = RIGHT; + phi_norm = phi - M_PI_2_F; + } else { + direction = BACK; + phi_norm = phi + ((phi > 0.f) ? -M_PI_F : M_PI_F); + } + + theta_threshold = atan(cos(phi_norm)); + if (theta > theta_threshold) { + direction = DOWN; + } else if (theta < -theta_threshold) { + direction = UP; + } + + switch (direction) { + case RIGHT: + uv.x = -xyz.z / xyz.x; + uv.y = xyz.y / xyz.x; + *face = TOP_RIGHT; + break; + case LEFT: + uv.x = -xyz.z / xyz.x; + uv.y = -xyz.y / xyz.x; + *face = TOP_LEFT; + break; + case UP: + uv.x = -xyz.x / xyz.y; + uv.y = -xyz.z / xyz.y; + *face = BOTTOM_RIGHT; + uv = rotate_cube_face(uv, ROT_270); + break; + case DOWN: + uv.x = xyz.x / xyz.y; + uv.y = -xyz.z / xyz.y; + *face = BOTTOM_LEFT; + uv = rotate_cube_face(uv, ROT_270); + break; + case FRONT: + uv.x = xyz.x / xyz.z; + uv.y = xyz.y / xyz.z; + *face = TOP_MIDDLE; + break; + case BACK: + uv.x = xyz.x / xyz.z; + uv.y = -xyz.y / xyz.z; + *face = BOTTOM_MIDDLE; + uv = rotate_cube_face(uv, ROT_90); + break; + } + + return uv; +} + +static float2 xyz_to_eac(float3 xyz, int2 size) +{ + float pixel_pad = 2; + float u_pad = pixel_pad / size.x; + float v_pad = pixel_pad / size.y; + + int face; + int u_face, v_face; + float2 uv = xyz_to_cube(xyz, &face); + + u_face = face % 3; + v_face = face / 3; + //eac expansion + uv = M_2_PI_F * atan(uv) + 0.5f; + + uv.x = (uv.x + u_face) * (1.f - 2.f * u_pad) / 3.f + u_pad; + uv.y = uv.y * (0.5f - 2.f * v_pad) + v_pad + 0.5f * v_face; + + uv.x *= size.x; + uv.y *= size.y; + + uv -= 0.5f; + + return uv; +} + +const sampler_t sampler_nearest = (CLK_NORMALIZED_COORDS_FALSE | + CLK_ADDRESS_CLAMP_TO_EDGE | + CLK_FILTER_NEAREST); + +const sampler_t sampler_linear = (CLK_NORMALIZED_COORDS_FALSE | + CLK_ADDRESS_CLAMP_TO_EDGE | + CLK_FILTER_LINEAR); + +static float4 gopromax_to_eac(float2 uv, int overlap, __read_only image2d_t src) +{ + int2 dim = get_image_dim(src); + int cube_size = dim.y; + int gap = (cube_size * 3 + overlap * 2 - dim.x) / 2; + float2 uv2 = uv; + float a = 0.f; + float4 val; + bool is_aligned; + + if (uv.x < cube_size || uv.x > cube_size * 2) { + int dx = 0; + int cs = cube_size - gap; + float cx = fmod(uv.x, cube_size) * cs / cube_size; + if (uv.x >= cube_size * 2) + dx = cube_size * 2 + overlap - gap; + if (cx >= (cs + overlap) / 2) + dx += overlap; + uv2.x = cx + dx; + if (cx > (cs - overlap) / 2 && cx < (cs + overlap) / 2) + a = (cx - (cs - overlap) / 2) / overlap; + } else { + uv2.x += overlap - gap; + } + + { + int2 d = convert_int2(ceil(uv2) - floor(uv2)); + is_aligned = (d.x == 0 && d.y == 0); + } + if (is_aligned) + val = read_imagef(src, sampler_nearest, uv2); + else + val = read_imagef(src, sampler_linear, uv2); + if (a > 0.f) { + float4 val2; + uv2.x += overlap; + if (is_aligned) + val2 = read_imagef(src, sampler_nearest, uv2); + else + val2 = read_imagef(src, sampler_linear, uv2); + val = mix(val, val2, a); + } + + return val; +} + +__kernel void gopromax_equirectangular(__write_only image2d_t dst, + __read_only image2d_t front, + __read_only image2d_t rear, + int overlap) +{ + float4 val; + int2 loc = (int2)(get_global_id(0), get_global_id(1)); + + int2 dst_size = get_image_dim(dst); + int2 src_size = get_image_dim(front); + int2 eac_size = (int2)(src_size.y * 3, src_size.y * 2); + + float3 xyz = equirect_to_xyz(loc, dst_size); + float2 uv = xyz_to_eac(xyz, eac_size); + + if (uv.y >= src_size.y) { + uv.y -= src_size.y; + val = gopromax_to_eac(uv, overlap, rear); + } else { + val = gopromax_to_eac(uv, overlap, front); + } + + write_imagef(dst, loc, val); +} + +__kernel void gopromax_stack(__write_only image2d_t dst, + __read_only image2d_t front, + __read_only image2d_t rear, + int overlap) +{ + float4 val; + int2 loc = (int2)(get_global_id(0), get_global_id(1)); + int2 dst_size = get_image_dim(dst); + int2 src_size = get_image_dim(front); + float2 uv = convert_float2(loc); + + uv *= (float)src_size.y * 2 / dst_size.y; + + if (uv.y >= src_size.y) { + uv.y -= src_size.y; + val = gopromax_to_eac(uv, overlap, rear); + } else { + val = gopromax_to_eac(uv, overlap, front); + } + + write_imagef(dst, loc, val); +} 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..3edfbf937b --- /dev/null +++ b/libavfilter/vf_gopromax_opencl.c @@ -0,0 +1,350 @@ +/* + * Copyright (c) 2021 Ronan LE MEILLAT + * Copyright (c) 2024 TADANO Tokumei + * + * This file is part of FFmpeg. + * + * FFmpeg is free software; you can redistribute it and/or + * modify it under the terms of the GNU Lesser General Public + * License as published by the Free Software Foundation; either + * version 2.1 of the License, or (at your option) any later version. + * + * FFmpeg is distributed in the hope that it will be useful, + * but WITHOUT ANY WARRANTY; without even the implied warranty of + * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU + * Lesser General Public License for more details. + * + * You should have received a copy of the GNU Lesser General Public + * License along with FFmpeg; if not, write to the Free Software + * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA + */ + +#include "libavutil/log.h" +#include "libavutil/mem.h" +#include "libavutil/pixdesc.h" +#include "libavutil/opt.h" + +#include "avfilter.h" +#include "framesync.h" +#include "opencl.h" +#include "opencl_source.h" +#include "video.h" +#include "v360.h" + +typedef struct GoProMaxOpenCLContext { + OpenCLFilterContext ocf; + + int initialised; + cl_kernel kernel; + cl_command_queue command_queue; + + FFFrameSync fs; + + int nb_planes; + + int out; + int width, height; + int overlap; +} GoProMaxOpenCLContext; + +static int gopromax_opencl_load(AVFilterContext *avctx, + enum AVPixelFormat front_format, + enum AVPixelFormat rear_format) +{ + GoProMaxOpenCLContext *ctx = avctx->priv; + cl_int cle; + const char *source = ff_source_gopromax_cl; + const char *kernel; + const AVPixFmtDescriptor *front_desc, *rear_desc; + int err, i, front_planes, rear_planes; + + front_desc = av_pix_fmt_desc_get(front_format); + rear_desc = av_pix_fmt_desc_get(rear_format); + front_planes = rear_planes = 0; + for (i = 0; i < front_desc->nb_components; i++) + front_planes = FFMAX(front_planes, + front_desc->comp[i].plane + 1); + for (i = 0; i < rear_desc->nb_components; i++) + rear_planes = FFMAX(rear_planes, + rear_desc->comp[i].plane + 1); + + ctx->nb_planes = front_planes; + + switch (ctx->out) { + case EQUIRECTANGULAR: + kernel = "gopromax_equirectangular"; + break; + case EQUIANGULAR: + kernel = "gopromax_stack"; + break; + default: + av_log(ctx, AV_LOG_ERROR, "Specified output format is not handled.\n"); + return AVERROR_BUG; + } + + av_log(avctx, AV_LOG_DEBUG, "Using kernel %s.\n", kernel); + + err = ff_opencl_filter_load_program(avctx, &source, 1); + if (err < 0) + goto fail; + + ctx->command_queue = clCreateCommandQueue(ctx->ocf.hwctx->context, + ctx->ocf.hwctx->device_id, + 0, &cle); + CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create OpenCL " + "command queue %d.\n", cle); + + ctx->kernel = clCreateKernel(ctx->ocf.program, kernel, &cle); + CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create kernel %d.\n", cle); + + ctx->initialised = 1; + return 0; + +fail: + if (ctx->command_queue) + clReleaseCommandQueue(ctx->command_queue); + if (ctx->kernel) + clReleaseKernel(ctx->kernel); + return err; +} + +static int gopromax_opencl_stack(FFFrameSync *fs) +{ + AVFilterContext *avctx = fs->parent; + AVFilterLink *outlink = avctx->outputs[0]; + GoProMaxOpenCLContext *ctx = avctx->priv; + AVFrame *input_front, *input_rear; + AVFrame *output; + cl_mem mem; + cl_int cle, overlap; + size_t global_work[2]; + int kernel_arg = 0; + int err, plane; + + err = ff_framesync_get_frame(fs, 0, &input_front, 0); + if (err < 0) + return err; + err = ff_framesync_get_frame(fs, 1, &input_rear, 0); + if (err < 0) + return err; + + if (!ctx->initialised) { + AVHWFramesContext *front_fc = + (AVHWFramesContext*)input_front->hw_frames_ctx->data; + AVHWFramesContext *rear_fc = + (AVHWFramesContext*)input_rear->hw_frames_ctx->data; + + err = gopromax_opencl_load(avctx, front_fc->sw_format, + rear_fc->sw_format); + if (err < 0) + return err; + } + + output = ff_get_video_buffer(outlink, outlink->w, outlink->h); + if (!output) { + err = AVERROR(ENOMEM); + goto fail; + } + + for (plane = 0; plane < ctx->nb_planes; plane++) { + kernel_arg = 0; + + mem = (cl_mem)output->data[plane]; + CL_SET_KERNEL_ARG(ctx->kernel, kernel_arg, cl_mem, &mem); + kernel_arg++; + + mem = (cl_mem)input_front->data[plane]; + CL_SET_KERNEL_ARG(ctx->kernel, kernel_arg, cl_mem, &mem); + kernel_arg++; + + mem = (cl_mem)input_rear->data[plane]; + CL_SET_KERNEL_ARG(ctx->kernel, kernel_arg, cl_mem, &mem); + kernel_arg++; + + overlap = ctx->overlap; + CL_SET_KERNEL_ARG(ctx->kernel, kernel_arg, cl_int, &overlap); + kernel_arg++; + + err = ff_opencl_filter_work_size_from_image(avctx, global_work, + output, plane, 0); + if (err < 0) + goto fail; + + av_log(avctx, AV_LOG_VERBOSE, + "In gopromax_opencl_stack for plane:%d %lux%lu frame size %dx%d\n", + plane, global_work[0], global_work[1], outlink->w, outlink->h); + + cle = clEnqueueNDRangeKernel(ctx->command_queue, ctx->kernel, 2, NULL, + global_work, NULL, 0, NULL, NULL); + CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to enqueue gopromax kernel " + "for plane %d: %d.\n", plane, cle); + } + + cle = clFinish(ctx->command_queue); + CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to finish command queue: %d.\n", cle); + + err = av_frame_copy_props(output, input_front); + + av_log(avctx, AV_LOG_DEBUG, "Filter output: %s, %ux%u (%"PRId64").\n", + av_get_pix_fmt_name(output->format), + output->width, output->height, output->pts); + + return ff_filter_frame(outlink, output); + +fail: + av_frame_free(&output); + return err; +} + +static int gopromax_opencl_config_output(AVFilterLink *outlink) +{ + AVFilterContext *avctx = outlink->src; + GoProMaxOpenCLContext *ctx = avctx->priv; + int height = avctx->inputs[0]->h; + int err; + + switch (ctx->out) { + case EQUIRECTANGULAR: + if (ctx->width > 0 && ctx->height > 0) { + if (ctx->width != ctx->height * 2) { + av_log(ctx, AV_LOG_ERROR, + "Specified size (%dx%d) is not suitable.\n", + ctx->width, ctx->height); + return AVERROR(EINVAL); + } + ctx->ocf.output_width = ctx->width; + ctx->ocf.output_height = ctx->height; + } else if (ctx->width > 0 || ctx->height > 0) { + av_log(ctx, AV_LOG_ERROR, + "Both width and height values should be specified.\n"); + return AVERROR(EINVAL); + } else { + ctx->ocf.output_width = 4 * height; + ctx->ocf.output_height = 2 * height; + } + break; + case EQUIANGULAR: + if (ctx->width > 0 && ctx->height > 0) { + if (ctx->width * 2 != ctx->height * 3) { + av_log(ctx, AV_LOG_ERROR, + "Specified size (%dx%d) is not suitable.\n", + ctx->width, ctx->height); + return AVERROR(EINVAL); + } + ctx->ocf.output_width = ctx->width; + ctx->ocf.output_height = ctx->height; + } else if (ctx->width > 0 || ctx->height > 0) { + av_log(ctx, AV_LOG_ERROR, + "Both width and height values should be specified.\n"); + return AVERROR(EINVAL); + } else { + ctx->ocf.output_width = 3 * height; + ctx->ocf.output_height = 2 * height; + } + break; + default: + av_log(ctx, AV_LOG_ERROR, "Specified output format is not supported.\n"); + return AVERROR(EINVAL); + } + + err = ff_opencl_filter_config_output(outlink); + if (err < 0) + return err; + + err = ff_framesync_init_dualinput(&ctx->fs, avctx); + if (err < 0) + return err; + + return ff_framesync_configure(&ctx->fs); +} + +static av_cold int gopromax_opencl_init(AVFilterContext *avctx) +{ + GoProMaxOpenCLContext *ctx = avctx->priv; + + ctx->fs.on_event = &gopromax_opencl_stack; + + return ff_opencl_filter_init(avctx); +} + +static int gopromax_opencl_activate(AVFilterContext *avctx) +{ + GoProMaxOpenCLContext *ctx = avctx->priv; + + return ff_framesync_activate(&ctx->fs); +} + +static av_cold void gopromax_opencl_uninit(AVFilterContext *avctx) +{ + GoProMaxOpenCLContext *ctx = avctx->priv; + cl_int cle; + + if (ctx->kernel) { + cle = clReleaseKernel(ctx->kernel); + if (cle != CL_SUCCESS) + av_log(avctx, AV_LOG_ERROR, "Failed to release " + "kernel: %d.\n", cle); + } + + if (ctx->command_queue) { + cle = clReleaseCommandQueue(ctx->command_queue); + if (cle != CL_SUCCESS) + av_log(avctx, AV_LOG_ERROR, "Failed to release " + "command queue: %d.\n", cle); + } + + ff_opencl_filter_uninit(avctx); + + ff_framesync_uninit(&ctx->fs); +} + +#define OFFSET(x) offsetof(GoProMaxOpenCLContext, x) +#define FLAGS (AV_OPT_FLAG_FILTERING_PARAM | AV_OPT_FLAG_VIDEO_PARAM) +static const AVOption gopromax_opencl_options[] = { + { "output", "set output projection", OFFSET(out), AV_OPT_TYPE_INT, {.i64=EQUIRECTANGULAR}, 0, NB_PROJECTIONS-1, FLAGS, .unit = "out" }, + { "e", "equirectangular", 0, AV_OPT_TYPE_CONST, {.i64=EQUIRECTANGULAR}, 0, 0, FLAGS, .unit = "out" }, + { "equirect", "equirectangular", 0, AV_OPT_TYPE_CONST, {.i64=EQUIRECTANGULAR}, 0, 0, FLAGS, .unit = "out" }, + { "eac", "equi-angular cubemap", 0, AV_OPT_TYPE_CONST, {.i64=EQUIANGULAR}, 0, 0, FLAGS, .unit = "out" }, + { "w", "output width", OFFSET(width), AV_OPT_TYPE_INT, {.i64=0}, 0, INT16_MAX, FLAGS, .unit = "w"}, + { "h", "output height", OFFSET(height), AV_OPT_TYPE_INT, {.i64=0}, 0, INT16_MAX, FLAGS, .unit = "h"}, + { "overlap", "set overlapped pixels", OFFSET(overlap), AV_OPT_TYPE_INT, {.i64=64}, 0, 128, FLAGS, .unit = "overlap"}, + { NULL }, +}; + +AVFILTER_DEFINE_CLASS(gopromax_opencl); + +static const AVFilterPad gopromax_opencl_inputs[] = { + { + .name = "front", + .type = AVMEDIA_TYPE_VIDEO, + .config_props = &ff_opencl_filter_config_input, + }, + { + .name = "rear", + .type = AVMEDIA_TYPE_VIDEO, + .config_props = &ff_opencl_filter_config_input, + }, +}; + +static const AVFilterPad gopromax_opencl_outputs[] = { + { + .name = "default", + .type = AVMEDIA_TYPE_VIDEO, + .config_props = &gopromax_opencl_config_output, + }, +}; + +const AVFilter ff_vf_gopromax_opencl = { + .name = "gopromax_opencl", + .description = NULL_IF_CONFIG_SMALL("GoProMax .360 to equirectangular projection"), + .priv_size = sizeof(GoProMaxOpenCLContext), + .priv_class = &gopromax_opencl_class, + .init = &gopromax_opencl_init, + .uninit = &gopromax_opencl_uninit, + .activate = &gopromax_opencl_activate, + FILTER_INPUTS(gopromax_opencl_inputs), + FILTER_OUTPUTS(gopromax_opencl_outputs), + FILTER_SINGLE_PIXFMT(AV_PIX_FMT_OPENCL), + .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE, + .flags = AVFILTER_FLAG_HWDEVICE, +};