From patchwork Thu Apr 12 22:29:20 2018 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Gabriel Machado X-Patchwork-Id: 8432 Delivered-To: ffmpegpatchwork@gmail.com Received: by 10.2.1.70 with SMTP id c67csp100239jad; Thu, 12 Apr 2018 15:29:32 -0700 (PDT) X-Google-Smtp-Source: AIpwx4/fVx4ZiGvh07zqxQWeC6UIfy1pRh1NnKE6kcRTR2fHzEbHg3Vy1KeEXyqv2dv2bvvRH5kK X-Received: by 10.223.138.246 with SMTP id z51mr1869825wrz.150.1523572171957; Thu, 12 Apr 2018 15:29:31 -0700 (PDT) ARC-Seal: i=1; a=rsa-sha256; t=1523572171; cv=none; d=google.com; s=arc-20160816; b=Pwy5RLKdFOGonn6YmK/4+vMemimiahT1ZV3QL8v11dauCBGzibdKyOoeIryhH8atYR FWOwXvvDyvroml/chXPW7aq9VTugIpHt+R05o/vJ8w1294o6keMH1Zn/HlK4aDzji6js lN4+X/RjUmx2BjfOhjCdmKGzj7ysQMwhZbz/NCq2MniFGGLCp1RDRAyJrRHcVuk2+ML/ Ajcc/0h+ATBI9+KadvL/KFLcSMPxSAeCvMNwl5KUlDDCyAs3zL1Gavi0tbmbN+qm4MXr R09s4cbsW1wI4fnkNkFK+UBzjMEHGPbaArbbOiwThFJfKsLGYhJ9y4QVjrC9d+uibNH1 z3hw== 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: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=nraL5wMs8gHUWLLxOvY9MI5ZHOpSBdpsP/oatetd39M=; b=Nj3l24rSzQ6T7kT2QfKu7TomIgCI+vMLrZmtr7+WscCNnY6kbsWwn4iCcFJ/dyG/vY x/1fbdKyf9S810yMmKYYNZw48aG1MiKLQ6aSX+8+z+/c09RxPvGHoNdDX8Bym5QTj1A4 LEHshzfam1oepdNULTWctGUlhmJyYgNo6SqLq47fx1ZvPaNXReV4gLfzs35Ba1+zBKfJ uK14On6ytw7m5eiAbC0TRDj11XiEdkFXCN9bRIERvXnX5gpMovj5b8rlDVexAzqeKzFl f6X0RYKG1qPVHrS/FwEs8PbK/SjnEv9iTuNYQL4rvFDvoDd3rRVwN/UuNt8obC8EFtYe PCbA== ARC-Authentication-Results: i=1; mx.google.com; dkim=neutral (body hash did not verify) header.i=@live.com header.s=selector1 header.b=eUxPWoih; 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 n82si280290wma.43.2018.04.12.15.29.31; Thu, 12 Apr 2018 15:29:31 -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=eUxPWoih; 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 342AE689E69; Fri, 13 Apr 2018 01:29:05 +0300 (EEST) X-Original-To: ffmpeg-devel@ffmpeg.org Delivered-To: ffmpeg-devel@ffmpeg.org Received: from NAM02-BL2-obe.outbound.protection.outlook.com (mail-oln040092003056.outbound.protection.outlook.com [40.92.3.56]) by ffbox0-bg.mplayerhq.hu (Postfix) with ESMTPS id 707B6689812 for ; Fri, 13 Apr 2018 01:28:58 +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=eVd0QmOj/4wqLdRrRKYn+Jjvewdm01zP+R1yP3+RAgo=; b=eUxPWoihh6hTzmVTLiKXe/yYKAFTx5birGukGvHBAmApb7MB9PHLOaK1KtWp/0blKb7yG4F6YnAgPysRHYg46Nn9SYnaaGWtnqYyFaJlaI9Nj2eKr5LhVTadZeGzXTz0HaOarEcpoCyHZ9eEpP6LaqWGALBLcFpRIEpRRxtRrVTLMQsT23YxRwSJf8FEpsXx4Pqz5mz7lAkLUPd9ZzTv7r/8N/Xk3wZb+S8+mJ+/2udV/Jzk8lU00xwBihbIapRhwGcWMFkhpsCa6DQetk1BBPe83d79dFg2djPP2cKWSj7lfdP0/5qa1Rgs4nJXegC3xs9XYlTv/D3L0jq1Qd+3kg== Received: from BL2NAM02FT036.eop-nam02.prod.protection.outlook.com (10.152.76.55) by BL2NAM02HT233.eop-nam02.prod.protection.outlook.com (10.152.77.208) with Microsoft SMTP Server (version=TLS1_2, cipher=TLS_ECDHE_RSA_WITH_AES_256_CBC_SHA384_P384) id 15.20.653.8; Thu, 12 Apr 2018 22:29:21 +0000 Received: from RO1PR80MB0059.lamprd80.prod.outlook.com (10.152.76.55) by BL2NAM02FT036.mail.protection.outlook.com (10.152.77.154) with Microsoft SMTP Server (version=TLS1_2, cipher=TLS_ECDHE_RSA_WITH_AES_128_CBC_SHA256_P256) id 15.20.653.8 via Frontend Transport; Thu, 12 Apr 2018 22:29:20 +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.0675.012; Thu, 12 Apr 2018 22:29:20 +0000 From: Gabriel Machado To: "ffmpeg-devel@ffmpeg.org" Thread-Topic: [FFmpeg-devel][PATCH v6] avfilter: add OpenCL scale filter Thread-Index: AQHT0q2zc5NURrgi30+7KAEXGR448w== Date: Thu, 12 Apr 2018 22:29:20 +0000 Message-ID: References: <148B1B7A67D1C24B9EF0BE42EA49770684EC6631@SHSMSX103.ccr.corp.intel.com> In-Reply-To: <148B1B7A67D1C24B9EF0BE42EA49770684EC6631@SHSMSX103.ccr.corp.intel.com> Accept-Language: en-GB, pt-BR, en-US Content-Language: en-US X-MS-Has-Attach: X-MS-TNEF-Correlator: x-incomingtopheadermarker: OriginalChecksum:BF364B26DE1FB9C2D79AA6559B5EB67FF32EC3300A1B9F017734E14CBA3FBA3B; UpperCasedChecksum:0FD023EE9E97D79D67266A76CC5A36D1EB822E94E36679F1433396FC1762C155; SizeAsReceived:7302; Count:46 x-ms-exchange-messagesentrepresentingtype: 1 x-tmn: [mNT8sI/vEYZxibxi49Sepd5+gd654FZP] x-ms-publictraffictype: Email x-microsoft-exchange-diagnostics: 1; BL2NAM02HT233; 7:kBeBDahEc3f9ssDHI1X/x7YRhiXNiddPYA2h07YbbA5s8oL4gP2baQ1FPeArwDWhPjZk+mxdFXA2i+XMp9vgBLra4Ru5H4Tp7ZN2uOZKmfKF07LVzkor6b+pg/dsgb5mj74AgvyGYAtEXOkKEVJCiJ6d8YC2FAAAwW4gP21z71BxMZpGrrvJshJ3ZgB2ReIrIM3zExBMpbyLuyDOg8yU+2WL4+WPKOIwSH2esaJ8aylXD9rnSqjQMjtPFcYAHzC3 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:BL2NAM02HT233; x-ms-traffictypediagnostic: BL2NAM02HT233: x-exchange-antispam-report-cfa-test: BCL:0; PCL:0; RULEID:(444000031); SRVR:BL2NAM02HT233; BCL:0; PCL:0; RULEID:; SRVR:BL2NAM02HT233; x-forefront-prvs: 06400060E1 x-forefront-antispam-report: SFV:NSPM; SFS:(7070007)(98901004); DIR:OUT; SFP:1901; SCL:1; SRVR:BL2NAM02HT233; H:RO1PR80MB0059.lamprd80.prod.outlook.com; FPR:; SPF:None; LANG:; x-microsoft-antispam-message-info: DZxdoQiyptisqJq+xPro/s7jhKoYVoAA5INM9wgnt1sGLZKtWZBRl8egCm8KqXRlLXRdviLLEyg6Q+apVPTV1KV55OkUtn7AiaRgLBBfhNu2WR/Ds5TTPrVgtHSgIr7jPvCBj/Ltt0r93Ig3iSoWh1rZZoJl3jV9b51CIcTiccienAObkQr63w1pLihdwnco Content-ID: <1E6094F9B2BCF745897025452E5E5469@lamprd80.prod.outlook.com> MIME-Version: 1.0 X-MS-Office365-Filtering-Correlation-Id: cb372d15-ffd6-4524-3470-08d5a0c4d5b1 X-OriginatorOrg: live.com X-MS-Exchange-CrossTenant-RMS-PersistedConsumerOrg: d4d70346-2c10-4f39-8c00-e767963926d9 X-MS-Exchange-CrossTenant-Network-Message-Id: cb372d15-ffd6-4524-3470-08d5a0c4d5b1 X-MS-Exchange-CrossTenant-rms-persistedconsumerorg: d4d70346-2c10-4f39-8c00-e767963926d9 X-MS-Exchange-CrossTenant-originalarrivaltime: 12 Apr 2018 22:29:20.8550 (UTC) X-MS-Exchange-CrossTenant-fromentityheader: Internet X-MS-Exchange-CrossTenant-id: 84df9e7f-e9f6-40af-b435-aaaaaaaaaaaa X-MS-Exchange-Transport-CrossTenantHeadersStamped: BL2NAM02HT233 Subject: Re: [FFmpeg-devel] [PATCH v6] 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" On 4/8/18 11:37 PM Song, Ruiling wrote: > > + float2 src_coord = (convert_float2(dst_pos) + 0.5) / dst_size; > For the floating point constant, it is better to add suffix 'f' to the number to tell the compiler that this is a single floating point value (like 0.5f). > The reason is if no suffix provided, floating point constants will be treated as double precision float number. > OpenCL C language derive this feature from C language. And most GPUs are performance sensitive to single floating point/ double floating point operations. > If you intentionally need double precision, it is ok to leave it as it is. > There are also some other places like this in the OpenCL file. Changed to floating point constants, thanks. - Gabriel From b5cbdef9abc3727e1456915e50bbf1d9c87751d7 Mon Sep 17 00:00:00 2001 From: Gabriel Machado Date: Sat, 31 Mar 2018 23:03:18 -0300 Subject: [PATCH] avfilter: add OpenCL scale filter --- configure | 1 + libavfilter/Makefile | 1 + libavfilter/allfilters.c | 1 + libavfilter/opencl/scale.cl | 67 ++++++ libavfilter/opencl_source.h | 1 + libavfilter/vf_scale_opencl.c | 544 ++++++++++++++++++++++++++++++++++++++++++ 6 files changed, 615 insertions(+) create mode 100644 libavfilter/opencl/scale.cl create mode 100644 libavfilter/vf_scale_opencl.c diff --git a/configure b/configure index 1e192f8..7d0858c 100755 --- a/configure +++ b/configure @@ -2879,6 +2879,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 6eac828..3073881 100644 --- a/libavfilter/allfilters.c +++ b/libavfilter/allfilters.c @@ -293,6 +293,7 @@ extern AVFilter ff_vf_sab; extern AVFilter ff_vf_scale; extern AVFilter ff_vf_scale_cuda; extern AVFilter ff_vf_scale_npp; +extern AVFilter ff_vf_scale_opencl; extern AVFilter ff_vf_scale_qsv; extern AVFilter ff_vf_scale_vaapi; extern AVFilter ff_vf_scale2ref; diff --git a/libavfilter/opencl/scale.cl b/libavfilter/opencl/scale.cl new file mode 100644 index 0000000..6583bb5 --- /dev/null +++ b/libavfilter/opencl/scale.cl @@ -0,0 +1,67 @@ +/* + * 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 dst_pos = {get_global_id(0), get_global_id(1)}; + float2 dst_size = {get_global_size(0), get_global_size(1)}; + + float2 src_coord = (convert_float2(dst_pos) + 0.5f) / dst_size; + + float4 c = read_imagef(src, sampler, src_coord); + write_imagef(dst, dst_pos, c); +} + +__kernel void scale(__write_only image2d_t dst, + __read_only image2d_t src, + __constant float *cx, + __constant float *cy, + int2 flt_size) +{ + const sampler_t s_img = (CLK_NORMALIZED_COORDS_FALSE | + CLK_ADDRESS_CLAMP_TO_EDGE | + CLK_FILTER_NEAREST); + + int2 dst_pos = {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 src_coord = (convert_float2(dst_pos) + 0.5f) * src_size / dst_size; + + int2 src_pos = convert_int2(floor(src_coord - 0.5f)); + + float4 col = 0.0f; + for (int i = 0; i < flt_size.y; ++i) { + float4 s = 0.0f; + for (int j = 0; j < flt_size.x; ++j) { + float4 c = read_imagef(src, s_img, src_pos + (int2){flt_size.x/2 - j, flt_size.y/2 - i}); + s += c * cx[dst_pos.x * flt_size.x + j]; + } + col += s * cy[dst_pos.y * flt_size.y + i]; + } + + write_imagef(dst, dst_pos, 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..250edc4 --- /dev/null +++ b/libavfilter/vf_scale_opencl.c @@ -0,0 +1,544 @@ +/* + * 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" + +enum filters { + F_AREA, + F_BICUBIC, + F_BILINEAR, + F_GAUSSIAN, + F_LANCZOS, + F_NEIGHBOR, + F_SINC, + F_SPLINE, + F_EXPERIMENTAL +}; + +static int filter_radius[] = { + [F_AREA] = 1, + [F_BICUBIC] = 2, + [F_BILINEAR] = 1, + [F_GAUSSIAN] = 4, + [F_LANCZOS] = 3, + [F_NEIGHBOR] = -1, + [F_SINC] = 10, + [F_SPLINE] = 10, + [F_EXPERIMENTAL] = 4 +}; + +typedef struct ScaleOpenCLContext { + OpenCLFilterContext ocf; + + cl_command_queue command_queue; + cl_mem cx, cy; + cl_kernel kernel; + const char *kernel_name; + + char *w_expr, *h_expr; + int dst_w, dst_h; + int src_w, src_h; + int algorithm; + + cl_int2 flt_size; + int initialised; +} ScaleOpenCLContext; + +static float netravali(float t, float B, float C) +{ + if (t > 2) { + return 0; + } else { + float tt = t*t; + float ttt = t*tt; + if (t < 1) { + return ((12 - 9 * B - 6 * C) * ttt + + (-18 + 12 * B + 6 * C) * tt + + (6 - 2 * B)) / 6; + } else { + return ((-B - 6 * C) * ttt + + (6 * B + 30 * C) * tt + + (-12 * B - 48 * C) * t + + (8 * B + 24 * C)) / 6; + } + } +} + +static float sinc(float t) +{ + return (t == 0) ? 1.0 : sin(t * M_PI) / (t * M_PI); +} + +static float lanczos(float t, float a) +{ + return (t < a) ? sinc(t) * sinc(t / a) : 0; +} + +static double spline(double a, double b, double c, double d, double dist) +{ + if (dist <= 1.0) + return ((d * dist + c) * dist + b) * dist + a; + else + return spline(0.0, + b + 2.0 * c + 3.0 * d, + c + 3.0 * d, + -b - 3.0 * c - 6.0 * d, + dist - 1.0); +} + +static float calc_weight(int algorithm, float ratio, float t) +{ + t = fabs(t); + + switch (algorithm) { + case F_AREA: { + float t2 = t - 0.5; + if (t2 * ratio < -0.5) + return 1; + else if (t2 * ratio < 0.5) + return -t2 * ratio + 0.5; + else + return 0; + } + + case F_BICUBIC: { + const float B = 0, C = 0.6; + return netravali(t, B, C); + } + + case F_BILINEAR: + return t < 1 ? (1 - t) : 0; + + case F_EXPERIMENTAL: { + double A = 1.0; + double c; + + if (t < 1.0) + c = cos(t * M_PI); + else + c = -1.0; + if (c < 0.0) + c = -pow(-c, A); + else + c = pow(c, A); + return c * 0.5 + 0.5; + } + + case F_GAUSSIAN: { + const float p = 3.0; + return exp2(-p * t * t); + } + + case F_LANCZOS: { + return lanczos(t, filter_radius[algorithm]); + } + + case F_NEIGHBOR: + return 1; + + case F_SINC: + return sinc(t); + + case F_SPLINE: { + const double p = -2.196152422706632; + return spline(1.0, 0.0, p, -p - 1.0, t); + } + } + + return 0; +} + +static int scale_opencl_init(AVFilterContext *avctx) +{ + ScaleOpenCLContext *ctx = avctx->priv; + cl_int cle; + int err; + int i, j; + int filterw, filterh; + float scalex, scaley; + float *cx = NULL, *cy = NULL; + + if (ctx->algorithm == F_NEIGHBOR) { + ctx->kernel_name = "neighbor"; + } else { + ctx->kernel_name = "scale"; + + scalex = FFMAX((float) ctx->src_w / ctx->dst_w, 1); + scaley = FFMAX((float) ctx->src_h / ctx->dst_h, 1); + filterw = ceil(2 * filter_radius[ctx->algorithm] * scalex); + filterh = ceil(2 * filter_radius[ctx->algorithm] * scaley); + + filterw = FFMIN(filterw, ctx->src_w - 2); + filterw = FFMAX(filterw, 1); + filterh = FFMIN(filterh, ctx->src_h - 2); + filterh = FFMAX(filterh, 1); + + ctx->flt_size.s[0] = filterw; + ctx->flt_size.s[1] = filterh; + + av_log(avctx, AV_LOG_INFO, "filter size: %d %d.\n", filterw, filterh); + + cx = av_malloc_array(ctx->dst_w * filterw, sizeof(cl_float)); + cy = av_malloc_array(ctx->dst_h * filterh, sizeof(cl_float)); + + if (!cx || !cy) { + err = AVERROR(ENOMEM); + goto fail; + } + + for (i = 0; i < ctx->dst_w; ++i) { + float s_x = (i + 0.5) * ctx->src_w / ctx->dst_w - 0.5; + float t = s_x - floor(s_x); // fract + + float sum = 0; + for (j = 0; j < filterw; ++j) { + int x = filterw/2 - j; + sum += cx[i * filterw + j] = calc_weight(ctx->algorithm, + scalex, + (x - t) / scalex); + } + + for (j = 0; j < filterw; ++j) + cx[i * filterw + j] /= sum; + } + + for (i = 0; i < ctx->dst_h; ++i) { + float s_y = (i + 0.5) * ctx->src_h / ctx->dst_h - 0.5; + float t = s_y - floor(s_y); // fract + + float sum = 0; + for (j = 0; j < filterh; ++j) { + int y = filterh/2 - j; + sum += cy[i * filterh + j] = calc_weight(ctx->algorithm, + scaley, + (y - t) / scaley); + } + + for (j = 0; j < filterh; ++j) + cy[i * filterh + j] /= sum; + } + + ctx->cx = clCreateBuffer(ctx->ocf.hwctx->context, + CL_MEM_READ_ONLY | + CL_MEM_COPY_HOST_PTR | + CL_MEM_HOST_NO_ACCESS, + ctx->dst_w * filterw * sizeof(cl_float), + cx, + &cle); + + ctx->cy = clCreateBuffer(ctx->ocf.hwctx->context, + CL_MEM_READ_ONLY | + CL_MEM_COPY_HOST_PTR | + CL_MEM_HOST_NO_ACCESS, + ctx->dst_h * filterh * sizeof(cl_float), + cy, + &cle); + av_free(cx); + av_free(cy); + if (!ctx->cx || !ctx->cy) { + av_log(avctx, AV_LOG_ERROR, "Failed to create weights buffer: %d.\n", cle); + err = AVERROR(EIO); + goto fail; + } + } + + 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, ctx->kernel_name, &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); + if (ctx->cx) + clReleaseMemObject(ctx->cx); + if (ctx->cy) + clReleaseMemObject(ctx->cy); + if (cx) + av_free(cx); + if (cy) + av_free(cy); + return err; +} + +static int config_input_props(AVFilterLink *inlink) +{ + AVFilterContext *avctx = inlink->dst; + AVFilterLink *outlink = avctx->outputs[0]; + ScaleOpenCLContext *ctx = avctx->priv; + int ret; + + if ((ret = ff_scale_eval_dimensions(ctx, + ctx->w_expr, ctx->h_expr, + inlink, outlink, + &ctx->dst_w, &ctx->dst_h)) < 0) + return ret; + + if (((int64_t) ctx->dst_h * inlink->w) > INT_MAX || + ((int64_t) ctx->dst_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 = ctx->dst_w; + ctx->ocf.output_height = ctx->dst_h; + ctx->src_w = inlink->w; + ctx->src_h = inlink->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; + } + + if (ctx->algorithm != F_NEIGHBOR) { + cle = clSetKernelArg(ctx->kernel, 2, sizeof(cl_mem), &ctx->cx); + if (cle != CL_SUCCESS) { + av_log(avctx, AV_LOG_ERROR, "Failed to set kernel " + "cx argument: %d.\n", cle); + goto fail; + } + cle = clSetKernelArg(ctx->kernel, 3, sizeof(cl_mem), &ctx->cy); + if (cle != CL_SUCCESS) { + av_log(avctx, AV_LOG_ERROR, "Failed to set kernel " + "cy argument: %d.\n", cle); + goto fail; + } + cle = clSetKernelArg(ctx->kernel, 4, sizeof(cl_int2), &ctx->flt_size); + if (cle != CL_SUCCESS) { + av_log(avctx, AV_LOG_ERROR, "Failed to set kernel " + "filter size 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); + } + + if (ctx->cx) { + cle = clReleaseMemObject(ctx->cx); + if (cle != CL_SUCCESS) + av_log(avctx, AV_LOG_ERROR, "Failed to release " + "weights buffer: %d.\n", cle); + } + + if (ctx->cy) { + cle = clReleaseMemObject(ctx->cy); + if (cle != CL_SUCCESS) + av_log(avctx, AV_LOG_ERROR, "Failed to release " + "weights buffer: %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" }, + { "area", "area averaging", 0, AV_OPT_TYPE_CONST, { .i64 = F_AREA }, 0, 0, FLAGS, "algorithm" }, + { "bicubic", "bicubic", 0, AV_OPT_TYPE_CONST, { .i64 = F_BICUBIC }, 0, 0, FLAGS, "algorithm" }, + { "bilinear", "bilinear", 0, AV_OPT_TYPE_CONST, { .i64 = F_BILINEAR }, 0, 0, FLAGS, "algorithm" }, + { "experimental", "experimental", 0, AV_OPT_TYPE_CONST, { .i64 = F_EXPERIMENTAL }, 0, 0, FLAGS, "algorithm" }, + { "gauss", "Gaussian", 0, AV_OPT_TYPE_CONST, { .i64 = F_GAUSSIAN }, 0, 0, FLAGS, "algorithm" }, + { "lanczos", "Lanczos", 0, AV_OPT_TYPE_CONST, { .i64 = F_LANCZOS }, 0, 0, FLAGS, "algorithm" }, + { "neighbor", "nearest neighbor", 0, AV_OPT_TYPE_CONST, { .i64 = F_NEIGHBOR }, 0, 0, FLAGS, "algorithm" }, + { "sinc", "sinc", 0, AV_OPT_TYPE_CONST, { .i64 = F_SINC }, 0, 0, FLAGS, "algorithm" }, + { "spline", "bicubic spline", 0, AV_OPT_TYPE_CONST, { .i64 = F_SPLINE }, 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, +};