From patchwork Tue Mar 27 15:09:08 2018 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Gabriel Machado X-Patchwork-Id: 8194 Delivered-To: ffmpegpatchwork@gmail.com Received: by 10.2.1.70 with SMTP id c67csp4882084jad; Tue, 27 Mar 2018 08:09:28 -0700 (PDT) X-Google-Smtp-Source: AG47ELsqx7O0TwwTsXuEHcR2NCY3Ox+KvPvRDX+8NbpnZQsMQADf1e7ota/whLKdYiv9/yxCEQTs X-Received: by 10.28.9.81 with SMTP id 78mr17803845wmj.23.1522163368711; Tue, 27 Mar 2018 08:09:28 -0700 (PDT) ARC-Seal: i=1; a=rsa-sha256; t=1522163368; cv=none; d=google.com; s=arc-20160816; b=kVpl3TqRytXuWdzdLt9f55Wh2ZMd4jIvSbZOTZSEUPvLeKvyq7ORZ5gjS1rbOl6pAO EQaiY/lU0po90UF7ocVLXMNaxD0B+RuYKozX1Be18IB8cFdCxqkGw4rWPNABKWk+gU3N ax23crpStsEsuCd6ugo1CkeIqqrJLaaUIGZWVrrpauy4Dc1IGQmJXxwQuk1kj9Lzp207 K6Q9QoBIRJlG6R1qBKRfbtzLaTWfk4KjPyaxLFGAbQFiCUg1TkFL54TRjru5icSVylss Nynw7LxIjkj/SPxliHZDcwXHpaNh2v9/leSGUF08BZos5Z+L08RtATCz1YETQJB65wq3 LxpA== ARC-Message-Signature: i=1; a=rsa-sha256; c=relaxed/relaxed; d=google.com; s=arc-20160816; h=sender:errors-to:content-transfer-encoding:reply-to:list-subscribe :list-help:list-post:list-archive:list-unsubscribe:list-id :precedence:subject:mime-version:content-id:spamdiagnosticmetadata :spamdiagnosticoutput:content-language:accept-language:in-reply-to :references:message-id:date:thread-index:thread-topic:to:from :dkim-signature:delivered-to:arc-authentication-results; bh=0oiwCeW8joxHQQyFqQHZoG7dXWgRrMY6XcqRj5a0Tts=; b=kSclWo5WWsFit0ugzeeJ1Jx2YkZGq+ymamqZgTvsinYcrIHdkJubbbVKyYl/h4O7az QSP6PSo7h1UzmVGABo/xwMCYQslf+c0T0gE8TdcpZVySE97eZY1O++lh8q8VYH/Wz7OG 0yVBTwPvzkfFySnI2Cmmia2lkOC1fefd5cqr+4zB2p/R9tMlYuqOUAnpXG1GUOU4ua99 vIskFkgSxFgdp4SvYmsU5+kqtcLYwLNRUA9M/zPBgO9yzQBzEpLyD3qIXebVT/KJosHw A9ShuiAEpBQ6f0wzRTYNj2YHV8gsmFlqAB1KKVU4CVv2wRHiLwDLC5z4LUYFaMDCVNw6 Wi/w== ARC-Authentication-Results: i=1; mx.google.com; dkim=neutral (body hash did not verify) header.i=@live.com header.s=selector1 header.b=WdbMC/Ml; 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=QUARANTINE dis=NONE) header.from=live.com Return-Path: Received: from ffbox0-bg.mplayerhq.hu (ffbox0-bg.ffmpeg.org. [79.124.17.100]) by mx.google.com with ESMTP id c136si1251997wmh.44.2018.03.27.08.09.15; Tue, 27 Mar 2018 08:09:28 -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=@live.com header.s=selector1 header.b=WdbMC/Ml; 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=QUARANTINE dis=NONE) header.from=live.com Received: from [127.0.1.1] (localhost [127.0.0.1]) by ffbox0-bg.mplayerhq.hu (Postfix) with ESMTP id BF94C689BCE; Tue, 27 Mar 2018 18:08:56 +0300 (EEST) X-Original-To: ffmpeg-devel@ffmpeg.org Delivered-To: ffmpeg-devel@ffmpeg.org Received: from NAM03-CO1-obe.outbound.protection.outlook.com (mail-oln040092007102.outbound.protection.outlook.com [40.92.7.102]) by ffbox0-bg.mplayerhq.hu (Postfix) with ESMTPS id 2A1FA68038A for ; Tue, 27 Mar 2018 18:08:54 +0300 (EEST) DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=live.com; s=selector1; h=From:Date:Subject:Message-ID:Content-Type:MIME-Version; bh=03xqBoCJXxIpQWCrcoo974/82vd9SeTBgFf+cqdd47M=; b=WdbMC/Mlioxc83IkGyKNSDeYfxNmZOuNF81CTatg8Lw82JyRn6AEprfWrUK0/C28KOgWaKk6dt+VT6zIh1MteyBHd9ktaXccoWQ2itoBUmq5Sv9/FXUCq4tUeGhf9vIaBGmvW6I0tJBVBvXAmQWK7K67O90vlUeR9qdr/yTdj60lGzXaW1h+IircUjQp4ZVb+419rA+vHSQ6nDO5Yv3FqT2zsRHvfQlgNJ2YAQrNhb/NkxwY4mtql/kl9tlxKF8rcGtZ8EPKJZe+qGKfduN5V6ApGegUtsuCMAM5uU2mw6wAfhnSj2/5RRmHAwbZU9idBxJIWGWq15kYMFYdDFOBmA== Received: from DM3NAM03FT042.eop-NAM03.prod.protection.outlook.com (10.152.82.60) by DM3NAM03HT069.eop-NAM03.prod.protection.outlook.com (10.152.83.119) with Microsoft SMTP Server (version=TLS1_2, cipher=TLS_ECDHE_RSA_WITH_AES_256_GCM_SHA384) id 15.20.631.7; Tue, 27 Mar 2018 15:09:09 +0000 Received: from RO1PR80MB0059.lamprd80.prod.outlook.com (10.152.82.56) by DM3NAM03FT042.mail.protection.outlook.com (10.152.83.174) with Microsoft SMTP Server (version=TLS1_2, cipher=TLS_ECDHE_RSA_WITH_AES_256_CBC_SHA384) id 15.20.631.7 via Frontend Transport; Tue, 27 Mar 2018 15:09:08 +0000 Received: from RO1PR80MB0059.lamprd80.prod.outlook.com ([fe80::4c33:a09c:e30f:e77c]) by RO1PR80MB0059.lamprd80.prod.outlook.com ([fe80::4c33:a09c:e30f:e77c%18]) with mapi id 15.20.0609.012; Tue, 27 Mar 2018 15:09:08 +0000 From: Gabriel Machado To: "ffmpeg-devel@ffmpeg.org" Thread-Topic: [FFmpeg-devel][PATCH v2 RESEND] avfilter: add OpenCL scale filter Thread-Index: AQHTxd2OnqFs0aMTWEi1zA7iokpIvQ== Date: Tue, 27 Mar 2018 15:09:08 +0000 Message-ID: References: In-Reply-To: Accept-Language: en-GB, pt-BR, en-US Content-Language: en-US X-MS-Has-Attach: X-MS-TNEF-Correlator: x-incomingtopheadermarker: OriginalChecksum:E34CE3512C859010CF8AFAC10B82A08FCD01675D11E2FD464E1EE2EF10B6CC64; UpperCasedChecksum:B5994589BB1475A65714FFF34220C7B73578CD4DF3DD658DE996E31064168AB2; SizeAsReceived:7213; Count:46 x-ms-exchange-messagesentrepresentingtype: 1 x-tmn: [Yj5f6wIC5uSDueSUybFm9kaaWp9g5oFp] x-ms-publictraffictype: Email x-microsoft-exchange-diagnostics: 1; DM3NAM03HT069; 7:0Qg/bofEH5dJAfDY8DAnsCjqhfyUdalYylRn7r3qTvEArjAmWCL5VCzKAHXDTAi2ljulwy9V7l4jY3fIT2R3hhbcU6ld4kPWxAkzjDFpXOLviuxqk19BQ3aWl1Hh1J4MQL0fx4U8g01A22YueCiBdq/raYID82j2hyptAUEbN94BVKkbFRukuM0bV1MX1ewKCM8nCdtJcMjlD7PWN7x0rfBudfVHiHUHW197yorJ837pBEReu8N3rjayp2cp6i9a x-incomingheadercount: 46 x-eopattributedmessage: 0 x-microsoft-antispam: UriScan:; BCL:0; PCL:0; RULEID:(7020095)(201702061078)(5061506573)(5061507331)(1603103135)(2017031320274)(2017031324274)(2017031323274)(2017031322404)(1601125374)(1603101448)(1701031045); SRVR:DM3NAM03HT069; x-ms-traffictypediagnostic: DM3NAM03HT069: x-ms-office365-filtering-correlation-id: fec50100-9908-49c1-e85d-08d593f4b043 x-exchange-antispam-report-cfa-test: BCL:0; PCL:0; RULEID:(444000031); SRVR:DM3NAM03HT069; BCL:0; PCL:0; RULEID:; SRVR:DM3NAM03HT069; x-forefront-prvs: 0624A2429E x-forefront-antispam-report: SFV:NSPM; SFS:(7070007)(98901004); DIR:OUT; SFP:1901; SCL:1; SRVR:DM3NAM03HT069; H:RO1PR80MB0059.lamprd80.prod.outlook.com; FPR:; SPF:None; LANG:; x-microsoft-antispam-message-info: XqG/267w0k1Z5EIn0xZKdbFsZYkDZDnkRDUTxfFGOSNP9TfbmmL/dsGqxpKHJcobFLRG+pxurI/BizzkzRkDvop2XDx2eCjJMf5osZ9ovHHhh/5z6WMacjp8Qhyg8Qta1ApHZ9moJc6Ww1g2Aeq5tMC65GxT5OYrFQbnRv9vurU0KHxYQDvlHkEanPrkWr6W spamdiagnosticoutput: 1:99 spamdiagnosticmetadata: NSPM Content-ID: <753ACCE091CFC443A01093CC251D35D9@lamprd80.prod.outlook.com> MIME-Version: 1.0 X-OriginatorOrg: live.com X-MS-Exchange-CrossTenant-Network-Message-Id: fec50100-9908-49c1-e85d-08d593f4b043 X-MS-Exchange-CrossTenant-originalarrivaltime: 27 Mar 2018 15:09:08.7118 (UTC) X-MS-Exchange-CrossTenant-fromentityheader: Internet X-MS-Exchange-CrossTenant-id: 84df9e7f-e9f6-40af-b435-aaaaaaaaaaaa X-MS-Exchange-Transport-CrossTenantHeadersStamped: DM3NAM03HT069 Subject: [FFmpeg-devel] [PATCH v2 RESEND] avfilter: add OpenCL scale filter X-BeenThere: ffmpeg-devel@ffmpeg.org X-Mailman-Version: 2.1.20 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 Errors-To: ffmpeg-devel-bounces@ffmpeg.org Sender: "ffmpeg-devel" Sorry, my mail client corrupted the patch. --- configure | 1 + libavfilter/Makefile | 1 + libavfilter/allfilters.c | 1 + libavfilter/opencl/scale.cl | 111 +++++++++++++++++ libavfilter/opencl_source.h | 1 + libavfilter/vf_scale_opencl.c | 284 ++++++++++++++++++++++++++++++++++++++++++ 6 files changed, 399 insertions(+) create mode 100644 libavfilter/opencl/scale.cl create mode 100644 libavfilter/vf_scale_opencl.c diff --git a/configure b/configure index 5ccf3ce..4007ee8 100755 --- a/configure +++ b/configure @@ -2821,6 +2821,7 @@ v4l2_m2m_deps_any="linux_videodev2_h" hwupload_cuda_filter_deps="ffnvcodec" scale_npp_filter_deps="ffnvcodec libnpp" +scale_opencl_filter_deps="opencl" scale_cuda_filter_deps="cuda_sdk" thumbnail_cuda_filter_deps="cuda_sdk" diff --git a/libavfilter/Makefile b/libavfilter/Makefile index a90ca30..6303cbd 100644 --- a/libavfilter/Makefile +++ b/libavfilter/Makefile @@ -302,6 +302,7 @@ OBJS-$(CONFIG_SAB_FILTER) += vf_sab.o OBJS-$(CONFIG_SCALE_FILTER) += vf_scale.o scale.o OBJS-$(CONFIG_SCALE_CUDA_FILTER) += vf_scale_cuda.o vf_scale_cuda.ptx.o OBJS-$(CONFIG_SCALE_NPP_FILTER) += vf_scale_npp.o scale.o +OBJS-$(CONFIG_SCALE_OPENCL_FILTER) += vf_scale_opencl.o opencl.o opencl/scale.o OBJS-$(CONFIG_SCALE_QSV_FILTER) += vf_scale_qsv.o OBJS-$(CONFIG_SCALE_VAAPI_FILTER) += vf_scale_vaapi.o scale.o vaapi_vpp.o OBJS-$(CONFIG_SCALE2REF_FILTER) += vf_scale.o scale.o diff --git a/libavfilter/allfilters.c b/libavfilter/allfilters.c index 1cf1340..3185b17 100644 --- a/libavfilter/allfilters.c +++ b/libavfilter/allfilters.c @@ -309,6 +309,7 @@ static void register_all(void) REGISTER_FILTER(SCALE, scale, vf); REGISTER_FILTER(SCALE_CUDA, scale_cuda, vf); REGISTER_FILTER(SCALE_NPP, scale_npp, vf); + REGISTER_FILTER(SCALE_OPENCL, scale_opencl, vf); REGISTER_FILTER(SCALE_QSV, scale_qsv, vf); REGISTER_FILTER(SCALE_VAAPI, scale_vaapi, vf); REGISTER_FILTER(SCALE2REF, scale2ref, vf); diff --git a/libavfilter/opencl/scale.cl b/libavfilter/opencl/scale.cl new file mode 100644 index 0000000..3436694 --- /dev/null +++ b/libavfilter/opencl/scale.cl @@ -0,0 +1,111 @@ +/* + * Copyright (c) 2018 Gabriel Machado + * + * 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 + */ + +__kernel void neighbor(__write_only image2d_t dst, + __read_only image2d_t src) +{ + const sampler_t sampler = (CLK_NORMALIZED_COORDS_TRUE | + CLK_ADDRESS_CLAMP_TO_EDGE | + CLK_FILTER_NEAREST); + + int2 coord = {get_global_id(0), get_global_id(1)}; + int2 size = {get_global_size(0), get_global_size(1)}; + + float2 pos = (convert_float2(coord) + 0.5) / convert_float2(size); + + float4 c = read_imagef(src, sampler, pos); + write_imagef(dst, coord, c); +} + +__kernel void bilinear(__write_only image2d_t dst, + __read_only image2d_t src) +{ + const sampler_t sampler = (CLK_NORMALIZED_COORDS_TRUE | + CLK_ADDRESS_CLAMP_TO_EDGE | + CLK_FILTER_LINEAR); + + int2 coord = {get_global_id(0), get_global_id(1)}; + int2 size = {get_global_size(0), get_global_size(1)}; + + float2 pos = (convert_float2(coord) + 0.5) / convert_float2(size); + + float4 c = read_imagef(src, sampler, pos); + write_imagef(dst, coord, c); +} + +// https://developer.nvidia.com/gpugems/GPUGems/gpugems_ch24.html +float MitchellNetravali(float x, float B, float C) +{ + float t = fabs(x); + float tt = t*t; + float ttt = tt*t; + + if (t < 1) { + return ((12 - 9 * B - 6 * C) * ttt + + (-18 + 12 * B + 6 * C) * tt + (6 - 2 * B)) / 6; + } else if ((t >= 1) && (t < 2)) { + return ((-B - 6 * C) * ttt + + (6 * B + 30 * C) * tt + (-12 * B - 48 * C) * + t + (8 * B + 24 * C)) / 6; + } else { + return 0; + } +} + +float4 cubic(float4 c0, float4 c1, float4 c2, float4 c3, float t) +{ + float B = 0, C = 0.6; // libswscale default + float a = MitchellNetravali(t + 1, B, C); + float b = MitchellNetravali(t, B, C); + float c = MitchellNetravali(1 - t, B, C); + float d = MitchellNetravali(2 - t, B, C); + return a*c0 + b*c1 + c*c2 + d*c3; +} + +__kernel void bicubic(__write_only image2d_t dst, + __read_only image2d_t src) +{ + const sampler_t sampler = (CLK_NORMALIZED_COORDS_FALSE | + CLK_ADDRESS_CLAMP_TO_EDGE | + CLK_FILTER_NEAREST); + + int2 dst_coord = {get_global_id(0), get_global_id(1)}; + + float2 dst_size = {get_global_size(0), get_global_size(1)}; + float2 src_size = convert_float2(get_image_dim(src)); + + float2 uv = convert_float2(dst_coord) / dst_size; + + float2 src_pos = uv * convert_float2(src_size) - 0.5; + + float2 src_coordf; + float2 t = fract(src_pos, &src_coordf); + int2 src_coord = convert_int2(src_coordf); + +#define TEX(x,y) read_imagef(src, sampler, src_coord + (int2){x,y}) + float4 col = cubic(cubic(TEX(-1,-1), TEX(0,-1), TEX(1,-1), TEX(2,-1), t.x), + cubic(TEX(-1, 0), TEX(0, 0), TEX(1, 0), TEX(2, 0), t.x), + cubic(TEX(-1, 1), TEX(0, 1), TEX(1, 1), TEX(2, 1), t.x), + cubic(TEX(-1, 2), TEX(0, 2), TEX(1, 2), TEX(2, 2), t.x), + t.y); +#undef TEX + + write_imagef(dst, dst_coord, col); +} diff --git a/libavfilter/opencl_source.h b/libavfilter/opencl_source.h index 4bb9969..e3bb887 100644 --- a/libavfilter/opencl_source.h +++ b/libavfilter/opencl_source.h @@ -22,6 +22,7 @@ extern const char *ff_opencl_source_avgblur; extern const char *ff_opencl_source_convolution; extern const char *ff_opencl_source_overlay; +extern const char *ff_opencl_source_scale; extern const char *ff_opencl_source_unsharp; #endif /* AVFILTER_OPENCL_SOURCE_H */ diff --git a/libavfilter/vf_scale_opencl.c b/libavfilter/vf_scale_opencl.c new file mode 100644 index 0000000..6482162 --- /dev/null +++ b/libavfilter/vf_scale_opencl.c @@ -0,0 +1,284 @@ +/* + * Copyright (c) 2018 Gabriel Machado + * + * 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/common.h" +#include "libavutil/imgutils.h" +#include "libavutil/mem.h" +#include "libavutil/opt.h" +#include "libavutil/pixdesc.h" + +#include "avfilter.h" +#include "internal.h" +#include "opencl.h" +#include "opencl_source.h" +#include "scale.h" +#include "video.h" + +#define F_NEIGHBOR 0 +#define F_BILINEAR 1 +#define F_BICUBIC 2 + +typedef struct ScaleOpenCLContext { + OpenCLFilterContext ocf; + + int initialised; + cl_kernel kernel; + cl_command_queue command_queue; + + char *w_expr; + char *h_expr; + int algorithm; +} ScaleOpenCLContext; + +static const char *kernel_name[] = { + "neighbor", + "bilinear", + "bicubic" +}; + +static int scale_opencl_init(AVFilterContext *avctx) +{ + ScaleOpenCLContext *ctx = avctx->priv; + cl_int cle; + int err; + + err = ff_opencl_filter_load_program(avctx, &ff_opencl_source_scale, 1); + if (err < 0) + goto fail; + + ctx->command_queue = clCreateCommandQueue(ctx->ocf.hwctx->context, + ctx->ocf.hwctx->device_id, + 0, &cle); + if (!ctx->command_queue) { + av_log(avctx, AV_LOG_ERROR, "Failed to create OpenCL " + "command queue: %d.\n", cle); + err = AVERROR(EIO); + goto fail; + } + + ctx->kernel = clCreateKernel(ctx->ocf.program, kernel_name[ctx->algorithm], &cle); + if (!ctx->kernel) { + av_log(avctx, AV_LOG_ERROR, "Failed to create kernel: %d.\n", cle); + err = AVERROR(EIO); + goto fail; + } + + ctx->initialised = 1; + + return 0; + +fail: + if (ctx->command_queue) + clReleaseCommandQueue(ctx->command_queue); + if (ctx->kernel) + clReleaseKernel(ctx->kernel); + return err; +} + +static int config_input_props(AVFilterLink *inlink) +{ + AVFilterContext *avctx = inlink->dst; + AVFilterLink *outlink = avctx->outputs[0]; + ScaleOpenCLContext *ctx = avctx->priv; + int w, h; + int ret; + + if ((ret = ff_scale_eval_dimensions(ctx, + ctx->w_expr, ctx->h_expr, + inlink, outlink, + &w, &h)) < 0) + return ret; + + if (((int64_t)h * inlink->w) > INT_MAX || + ((int64_t)w * inlink->h) > INT_MAX) + av_log(ctx, AV_LOG_ERROR, "Rescaled value for width or height is too big.\n"); + + ctx->ocf.output_width = w; + ctx->ocf.output_height = h; + + return ff_opencl_filter_config_input(inlink); +} + +static int scale_opencl_filter_frame(AVFilterLink *inlink, AVFrame *input) +{ + AVFilterContext *avctx = inlink->dst; + AVFilterLink *outlink = avctx->outputs[0]; + ScaleOpenCLContext *ctx = avctx->priv; + AVHWFramesContext *main_fc = (AVHWFramesContext *) inlink->hw_frames_ctx->data; + const AVPixFmtDescriptor *pix_desc = av_pix_fmt_desc_get(main_fc->sw_format); + AVFrame *output = NULL; + + cl_int cle; + size_t global_work[2]; + cl_mem src, dst; + int err, p; + int x_subsample = 1 << pix_desc->log2_chroma_w; + int y_subsample = 1 << pix_desc->log2_chroma_h; + + av_log(ctx, AV_LOG_DEBUG, "Filter input: %s, %ux%u (%"PRId64").\n", + av_get_pix_fmt_name(input->format), + input->width, input->height, input->pts); + + if (!input->hw_frames_ctx) + return AVERROR(EINVAL); + + if (!ctx->initialised) { + err = scale_opencl_init(avctx); + if (err < 0) + goto fail; + } + + output = ff_get_video_buffer(outlink, outlink->w, outlink->h); + if (!output) { + err = AVERROR(ENOMEM); + goto fail; + } + + err = av_frame_copy_props(output, input); + if (err < 0) + goto fail; + output->width = outlink->w; + output->height = outlink->h; + + for (p = 0; p < FF_ARRAY_ELEMS(output->data); p++) { + src = (cl_mem) input->data[p]; + dst = (cl_mem)output->data[p]; + + if (!dst) + break; + + cle = clSetKernelArg(ctx->kernel, 0, sizeof(cl_mem), &dst); + if (cle != CL_SUCCESS) { + av_log(avctx, AV_LOG_ERROR, "Failed to set kernel " + "destination image argument: %d.\n", cle); + goto fail; + } + cle = clSetKernelArg(ctx->kernel, 1, sizeof(cl_mem), &src); + if (cle != CL_SUCCESS) { + av_log(avctx, AV_LOG_ERROR, "Failed to set kernel " + "source image argument: %d.\n", cle); + goto fail; + } + + global_work[0] = output->width / (p ? x_subsample : 1); + global_work[1] = output->height / (p ? y_subsample : 1); + + av_log(avctx, AV_LOG_DEBUG, "Run kernel on plane %d " + "(%"SIZE_SPECIFIER"x%"SIZE_SPECIFIER").\n", + p, global_work[0], global_work[1]); + + cle = clEnqueueNDRangeKernel(ctx->command_queue, ctx->kernel, 2, NULL, + global_work, NULL, 0, NULL, NULL); + if (cle != CL_SUCCESS) { + av_log(avctx, AV_LOG_ERROR, "Failed to enqueue kernel: %d.\n", cle); + err = AVERROR(EIO); + goto fail; + } + } + + cle = clFinish(ctx->command_queue); + if (cle != CL_SUCCESS) { + av_log(avctx, AV_LOG_ERROR, "Failed to finish command queue: %d.\n", cle); + err = AVERROR(EIO); + goto fail; + } + + av_frame_free(&input); + + av_log(ctx, 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: + clFinish(ctx->command_queue); + av_frame_free(&input); + av_frame_free(&output); + return err; +} + +static av_cold void scale_opencl_uninit(AVFilterContext *avctx) +{ + ScaleOpenCLContext *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); +} + +#define OFFSET(x) offsetof(ScaleOpenCLContext, x) +#define FLAGS (AV_OPT_FLAG_FILTERING_PARAM | AV_OPT_FLAG_VIDEO_PARAM) +static const AVOption scale_opencl_options[] = { + { "w", "Output video width", OFFSET(w_expr), AV_OPT_TYPE_STRING, { .str = "iw" }, .flags = FLAGS }, + { "h", "Output video height", OFFSET(h_expr), AV_OPT_TYPE_STRING, { .str = "ih" }, .flags = FLAGS }, + { "algorithm", "Scaling algorithm", OFFSET(algorithm), AV_OPT_TYPE_INT, { .i64 = F_BICUBIC }, INT_MIN, INT_MAX, FLAGS, "algorithm" }, + { "bilinear", "bilinear", 0, AV_OPT_TYPE_CONST, { .i64 = F_BILINEAR }, 0, 0, FLAGS, "algorithm" }, + { "bicubic", "bicubic", 0, AV_OPT_TYPE_CONST, { .i64 = F_BICUBIC }, 0, 0, FLAGS, "algorithm" }, + { "neighbor", "nearest neighbor", 0, AV_OPT_TYPE_CONST, { .i64 = F_NEIGHBOR }, 0, 0, FLAGS, "algorithm" }, + { NULL } +}; + +AVFILTER_DEFINE_CLASS(scale_opencl); + +static const AVFilterPad scale_opencl_inputs[] = { + { + .name = "default", + .type = AVMEDIA_TYPE_VIDEO, + .filter_frame = &scale_opencl_filter_frame, + .config_props = &config_input_props, + }, + { NULL } +}; + +static const AVFilterPad scale_opencl_outputs[] = { + { + .name = "default", + .type = AVMEDIA_TYPE_VIDEO, + .config_props = &ff_opencl_filter_config_output, + }, + { NULL } +}; + +AVFilter ff_vf_scale_opencl = { + .name = "scale_opencl", + .description = NULL_IF_CONFIG_SMALL("Scale the input video size."), + .priv_size = sizeof(ScaleOpenCLContext), + .priv_class = &scale_opencl_class, + .init = &ff_opencl_filter_init, + .uninit = &scale_opencl_uninit, + .query_formats = &ff_opencl_filter_query_formats, + .inputs = scale_opencl_inputs, + .outputs = scale_opencl_outputs, + .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE, +};