From patchwork Tue Mar 27 04:48:22 2018 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Gabriel Machado X-Patchwork-Id: 8186 Delivered-To: ffmpegpatchwork@gmail.com Received: by 10.2.1.70 with SMTP id c67csp4321164jad; Mon, 26 Mar 2018 21:48:38 -0700 (PDT) X-Google-Smtp-Source: AG47ELsL58OWljumHPhxNNHGAke9VnM0XtHi06Hfuqgn8vB5qF4mZ5NnH5AQzrsPn0hY9ULoRPBE X-Received: by 10.223.135.83 with SMTP id 19mr18107528wrz.150.1522126118017; Mon, 26 Mar 2018 21:48:38 -0700 (PDT) ARC-Seal: i=1; a=rsa-sha256; t=1522126117; cv=none; d=google.com; s=arc-20160816; b=X/BnyNDF2pRPNeb/1XkIP0WJvQaPwnFtaTSFX5bZZOJ/7ufMx0pCVYNe6VxSVdiyUh r4iBJw3+ZYl9JuT2obgappukftIxtqCF4kIxjOpjIppyKNY2RrZqP4r+qR4UKBraPB+b deU+SmPgvGnj1U2ZPw+KoXpzFluHt9jUN0EwUgsPsffcQ2EdZilweTpD9akxa7QbePhn W3iO2m0o3gt/hE0mnYOlvDH7cBvGQx3H28U7ICHrT3MfejXnzoAnk1eABa2EnRwieL54 dCUzKgDV8QKxOH0gHZafd6NepuBEJa0HMQEsc3sVckAOV20QTiJw6TPaoaVCRPGFWUs4 OtTQ== ARC-Message-Signature: i=1; a=rsa-sha256; c=relaxed/relaxed; d=google.com; s=arc-20160816; h=sender:errors-to:content-transfer-encoding:cc:reply-to :list-subscribe:list-help:list-post:list-archive:list-unsubscribe :list-id:precedence:subject:mime-version:spamdiagnosticmetadata :spamdiagnosticoutput:content-language:accept-language:message-id :date:thread-index:thread-topic:to:from:dkim-signature:delivered-to :arc-authentication-results; bh=P9TnsDEyHGgk5Zkl8pluWqEiiWTbaa+hPLcUUpMWVks=; b=rtZ3vbYwL+SJrBP7MwFLnkPkCv6rFMea3BdG1mO4+Hy9Jspv9uU/B+7E/325FTO0pZ Gkbuu7a3caRJEPobsh/VzIewMz7X2xXB1f942yhqT02AUW/9WlU5QzySKc9hqMS7n7sm nido4cZg1bp2+IxItv4vPRClucbhM5rFSw/3yd3R5V2pNXCS8xg5Gf/Bm9/GQvAJLpJ1 fJ9szdQGSY5AM/CojQQJfRzAGPepL6C8UpTtSER8FxtJ2+TRzmZGYKILNwSEos9tAnzF 7AX1nxTlkV6+ncsmMDNIbjjK0mte1WBNQzEjYDn17XZFVsm3BkJcLTatijUlPwPpBzhs h6Lg== ARC-Authentication-Results: i=1; mx.google.com; dkim=neutral (body hash did not verify) header.i=@live.com header.s=selector1 header.b=XJbDqVWS; 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 71si432783wmg.60.2018.03.26.21.48.37; Mon, 26 Mar 2018 21:48:37 -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=XJbDqVWS; 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 D441C689A34; Tue, 27 Mar 2018 07:48:17 +0300 (EEST) X-Original-To: ffmpeg-devel@ffmpeg.org Delivered-To: ffmpeg-devel@ffmpeg.org Received: from NAM04-CO1-obe.outbound.protection.outlook.com (mail-oln040092010076.outbound.protection.outlook.com [40.92.10.76]) by ffbox0-bg.mplayerhq.hu (Postfix) with ESMTPS id 1A3ED680754 for ; Tue, 27 Mar 2018 07:48:11 +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=0e3QeT2vh22r3xt1LDrxR6KUGpVl8X68n9YyQBG9oYc=; b=XJbDqVWS1le3BNcTXfdNtpsk1+/QWQxF05M6nLQNrMN3by7DX6Nu1zWt6vG6J0chb6NIJx6F7VRPssZbP8x2ctm0/oDU7omIHz3fi9FBGJ+W/yJPC0eG11rV+LT2rf1gkIqhgNBJIgVVmEx+Lu2IQRZEQ2LGrjiMnMJbffXgp8qsfwxpo+yt1cQgY6ak6Pdq8oIFoyUrecV+VJZrHw/64RCmZrPZb/2c7WAW++1FSB3zjbrDEnDxR9pzHjJ0AUufKGZ+Z5C+ckJeolCOElEIufaBc3A76iyQxzkbYkIeIrGuLD762ttRqKi7GchE16lHg0+6cTPtxTh+KVSup+ykgg== Received: from CO1NAM04FT044.eop-NAM04.prod.protection.outlook.com (10.152.90.52) by CO1NAM04HT248.eop-NAM04.prod.protection.outlook.com (10.152.90.251) with Microsoft SMTP Server (version=TLS1_2, cipher=TLS_ECDHE_RSA_WITH_AES_256_CBC_SHA384_P256) id 15.20.527.25; Tue, 27 Mar 2018 04:48:24 +0000 Received: from RO1PR80MB0059.lamprd80.prod.outlook.com (10.152.90.57) by CO1NAM04FT044.mail.protection.outlook.com (10.152.91.95) with Microsoft SMTP Server (version=TLS1_2, cipher=TLS_ECDHE_RSA_WITH_AES_256_CBC_SHA384_P256) id 15.20.527.25 via Frontend Transport; Tue, 27 Mar 2018 04:48:24 +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 04:48:23 +0000 From: Gabriel Machado To: "ffmpeg-devel@ffmpeg.org" Thread-Topic: [PATCH] avfilter: add OpenCL scale filter Thread-Index: AQHTxYbVm8mET90mLEiydsb7kqalWA== Date: Tue, 27 Mar 2018 04:48:22 +0000 Message-ID: Accept-Language: en-GB, pt-BR, en-US Content-Language: en-US X-MS-Has-Attach: X-MS-TNEF-Correlator: x-incomingtopheadermarker: OriginalChecksum:11FD0F0123B14401946E9B9DDF0841C05E7186E8D86649FF0D7A4F2A1C4E3F07; UpperCasedChecksum:69BB759BF61B445B2AD9DD75D9095E2B6E54B257772A17A5CFA02ACD1EF7FF88; SizeAsReceived:6991; Count:45 x-ms-exchange-messagesentrepresentingtype: 1 x-tmn: [Smbf81uDSaZN4bRsnCZLjnyuG1jol9lh] x-ms-publictraffictype: Email x-microsoft-exchange-diagnostics: 1; CO1NAM04HT248; 6:Z+XBkZrUsgADxTf4WhaiKbPmpbA+unqPS4eE2HPBmabEJVn1WCfMskeUmxnuUvL8PdNFo1QMWHGI0OUJwxmFfM39HuwiwjQLxZAAxRpZpnf+C2fNB9S/EPSDc2mJwGE7IgwfwQqAFDNkJ6qLJGR8yX7jGmM6BLDezbkQLliwakREUQwmxz4670d8s0VBRrMDRWTijOlwK4A8tF0fXr1pvRCO/YwZqSYyWF74AbCz4VsMacoJESJrvW5q8F/emVybblO061zRG5P4o67svUioG1pMecXiN2zxbTskCQ+BIZTk3jN1ldDRbPHeKo9re8G+jzZxIDNQGPJq0rWeFO+BdfPjQUrTJbx+6lYSxTvxXCE=; 5:th6diyfNlyTi0w9qOKsnhY4DrN/JIoIuNbi+WeSQ2tjfy2nLt+lXeAvf/we8YtmsoFu5ZlB1y+OL3hcjOWYQRo9yGYq2qsgBTcLWBmeB15WrLQaY8dBnCZhV5OERbCclHy3Pb/lR6dUzajD9KfXeXV7OG/73LW1pyPLG403BKhc=; 24:smh0clcooFP9PPsgbTE05EvWOy0ENl/AFd81tuMtfkgJtjM/AKtEBNUIW+7VzU8OOVOxwaIzXowkXm0VuZFwF/GXqTJcGgdh6FRv1iBL+b0=; 7:3bNO5xhPvPH/yBs6I8lv3eB34rmDy6cckuWrAvYS2XJgFOXiAcSFJ5+vFIOUO5zn3JNmLawcOJ1d3sddXAapRusTUQkY0Dbh/aFMc+PuqZJCL7SlnLQD72NHSWblmVOfxC4NINVCjgupAJKOlo9Y0o/PlNoMYP5urXjLVDkmWkVbBoGHZZJm0hG8al45mV1Dkl/TGkV2xUnI81m63d+sR5rj4TdrcYAzqmD8BWhkNtUMYaMjHgdMRkv5sUtRmrKp x-incomingheadercount: 45 x-eopattributedmessage: 0 x-microsoft-antispam: UriScan:; BCL:0; PCL:0; RULEID:(7020095)(201702061078)(5061506573)(5061507331)(1603103135)(2017031320274)(2017031324274)(2017031323274)(201702181274)(2017031322404)(1601125374)(1603101448)(1701031045); SRVR:CO1NAM04HT248; x-ms-exchange-slblob-mailprops: UPd532NOaID5XcXkg4/ZbAL4jRvaiFZAI92SkwhP4BCTJChWvHl6JGgdiIbNXKFfJOHeAJ4PmY4DJnPoOBNTCAVPXzB1D9UQviSY5xEvGmL29U0r0wtAmRP1YxY/8Dg85JM7YUSnqLEdt+VNMPDoyTcY2xssRb9RMmaWZyEge53VwezzelS4EPW+n4ucBUhPZqfqHeSwKwWXD3KzJPNAYYUPRYInd7IbllK8nAqkDcQ3Dt/Ncl7ylYbywxVf5OWBGykAl0ZbSHgMEZkE/e/f1Lg8+4fWr9uLnl8yGQl2W/mKxfnu82xmvfBPJ56p8VlSThL4F29Q4NLCshJ4xfaaHpRSIcubF3uDlNzwUD+EDfOLhcZVYOFvt5vanGG5hX60FH2MYu0JalciMD9BXXu2OdGVirFGNvj/6/AfGXQoVTARj1oHOgKLWFvJJ/e4MlIvXKNDl2CyKtDD9Ke54PyMZp0WNb6KujYRHcR+6feVRVps8lxv+Yk8pzh5mP8U0SbFoLBlwtCNiMvhiPgaZDQb2Z3Vx21M57HDxw5k0W73/8eds9nJ5MZgoSGnRq/6w7PClhkqvyDlV9uvicTE4TZ3x5VTdHSf52QrTD6eNwG4gz5THcuEcLg5XbdvHqBui4esj6VwTbHusICwzgzANuAmrogq6Ai0GbzhZ12FGu38nL9j7AIrr453/jJTVMUf7f5hZ/fVKU7YHvBUoYbndngAj8+2Q/lWioOhJh9ls6ywXfSHeZCWvH3NOwbGDFIbXWVWtfyd54zbHsQbaB3YTirfIjCSHOshf8YCsm5nJ2Wfu7iVRfQSA55IqWhHXls8iXj1Jm8mrNN4tggGM+ILDqFn1oIDti9FHarUmdCdoCY/YNAuRxBev5vqUw== x-ms-traffictypediagnostic: CO1NAM04HT248: x-ms-office365-filtering-correlation-id: 4b42009a-2338-4ccb-9149-08d5939df7ca x-exchange-antispam-report-cfa-test: BCL:0; PCL:0; RULEID:(444000031); SRVR:CO1NAM04HT248; BCL:0; PCL:0; RULEID:; SRVR:CO1NAM04HT248; x-forefront-prvs: 0624A2429E x-forefront-antispam-report: SFV:NSPM; SFS:(7070007)(98901004); DIR:OUT; SFP:1901; SCL:1; SRVR:CO1NAM04HT248; H:RO1PR80MB0059.lamprd80.prod.outlook.com; FPR:; SPF:None; LANG:; x-microsoft-antispam-message-info: qUS35t2mH9uyLCKm4h4AUt/YU/hHnl620slpaHsJvU8vxa4uoDJ7+lrmlX1QRHiQwDl0MEbY3AZ1V6dSQ4FGnVJTyDy8D18Qpvw2+elTlF85+6BZM0JbhxXh9kUf30WbMNL3ZZeF/ZtkyJx7iDQCEWbuOXyzqiCET7enGB46h7XYs9AbKFJw5ysPSalt71Qh spamdiagnosticoutput: 1:99 spamdiagnosticmetadata: NSPM MIME-Version: 1.0 X-OriginatorOrg: live.com X-MS-Exchange-CrossTenant-Network-Message-Id: 4b42009a-2338-4ccb-9149-08d5939df7ca X-MS-Exchange-CrossTenant-originalarrivaltime: 27 Mar 2018 04:48:23.0128 (UTC) X-MS-Exchange-CrossTenant-fromentityheader: Internet X-MS-Exchange-CrossTenant-id: 84df9e7f-e9f6-40af-b435-aaaaaaaaaaaa X-MS-Exchange-Transport-CrossTenantHeadersStamped: CO1NAM04HT248 Subject: [FFmpeg-devel] [PATCH] 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 Cc: Gabriel Machado Errors-To: ffmpeg-devel-bounces@ffmpeg.org Sender: "ffmpeg-devel" From: Gabriel Machado Some scaling filters implemented as OpenCL kernels. Can be used as: scale_opencl=::flags= where can be `neighbor', `bilinear', `bicubic' or `fast_bicubic' This is an initial draft, there's still a long way to go in terms of completeness, configurability and performance. --- configure | 1 + libavfilter/Makefile | 1 + libavfilter/allfilters.c | 1 + libavfilter/opencl/scale.cl | 165 ++++++++++++++++++++++++ libavfilter/opencl_source.h | 1 + libavfilter/vf_scale_opencl.c | 289 ++++++++++++++++++++++++++++++++++++++++++ 6 files changed, 458 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..b0e6cb2 --- /dev/null +++ b/libavfilter/opencl/scale.cl @@ -0,0 +1,165 @@ +/* + * 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); +} + +// https://www.shadertoy.com/view/4df3Dn +// 4x4 bicubic filter using 4 bilinear texture lookups +// cubic B-spline basis functions +float w0(float a) { return (1.0/6.0)*(a*(a*(-a + 3.0) - 3.0) + 1.0); } +float w1(float a) { return (1.0/6.0)*(a*a*(3.0*a - 6.0) + 4.0); } +float w2(float a) { return (1.0/6.0)*(a*(a*(-3.0*a + 3.0) + 3.0) + 1.0); } +float w3(float a) { return (1.0/6.0)*(a*a*a); } + +float g0(float a) { return w0(a) + w1(a); } +float g1(float a) { return w2(a) + w3(a); } + +float h0(float a) { return -1.0 + w1(a) / g0(a); } +float h1(float a) { return 1.0 + w3(a) / g1(a); } + +__kernel void fast_bicubic(__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 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 * src_size + 0.5; + + float2 src_coordf; + float2 frac = fract(src_pos, &src_coordf); + int2 src_coord = convert_int2(src_coordf); + + float g0x = g0(frac.x); + float g1x = g1(frac.x); + float h0x = h0(frac.x); + float h1x = h1(frac.x); + float h0y = h0(frac.y); + float h1y = h1(frac.y); + + float2 p0 = ((float2){src_coord.x + h0x, src_coord.y + h0y} - 0.5) / src_size; + float2 p1 = ((float2){src_coord.x + h1x, src_coord.y + h0y} - 0.5) / src_size; + float2 p2 = ((float2){src_coord.x + h0x, src_coord.y + h1y} - 0.5) / src_size; + float2 p3 = ((float2){src_coord.x + h1x, src_coord.y + h1y} - 0.5) / src_size; + + float4 c = g0(frac.y) * (g0x * read_imagef(src, sampler, p0) + + g1x * read_imagef(src, sampler, p1)) + + g1(frac.y) * (g0x * read_imagef(src, sampler, p2) + + g1x * read_imagef(src, sampler, p3)); + + write_imagef(dst, dst_coord, c); +} 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..14479c7 --- /dev/null +++ b/libavfilter/vf_scale_opencl.c @@ -0,0 +1,289 @@ +/* + * 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 +#define F_FAST_BICUBIC 3 + +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", + "fast_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" }, + { "fast_bicubic", "fast bicubic", 0, AV_OPT_TYPE_CONST, { .i64 = F_FAST_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, +};