From patchwork Tue Nov 14 19:47:27 2017 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Mark Thompson X-Patchwork-Id: 6064 Delivered-To: ffmpegpatchwork@gmail.com Received: by 10.2.161.94 with SMTP id m30csp3722013jah; Tue, 14 Nov 2017 11:49:18 -0800 (PST) X-Google-Smtp-Source: AGs4zMY4RvVcfsfRXFPvNP69dVJ0Fvszrb8vgFbqDZFjC4C649qqD23uuOn08kNY78Z0aciS8BNV X-Received: by 10.223.198.18 with SMTP id n18mr10413757wrg.96.1510688958842; Tue, 14 Nov 2017 11:49:18 -0800 (PST) ARC-Seal: i=1; a=rsa-sha256; t=1510688958; cv=none; d=google.com; s=arc-20160816; b=q3YUz2zjajreXpNk+p4ucYdLgBkWFC0g0HL5sQTsenm17o5anhTDTedK5/K326twsP Z0/NnAQVZQBOGUqPJjKnVxVx7y9gn2B68KhkK5vDPtyazYSH6fIO0vtQTjqdTFEkmDky IIjkqnQKifsH5RGVL30HCp89CddwldRd4dTLQWf9xLppJvrtaiL6UzLLK4JeXSQ1ct1e dR8nTKSVTjWQb0/LR0KVUvgr4NAX/lp/lUH7DtcRvWIAEgHaSNnrdW7+MUAh5lYMIbQZ +nXoozMQFwM/dNlLnw8lYQLchPRbws5YqDUvT6B9jUqMgNKH4pzOgmJimSCWpu2WUISQ 0YiQ== ARC-Message-Signature: i=1; a=rsa-sha256; c=relaxed/relaxed; d=google.com; s=arc-20160816; h=sender:errors-to:content-transfer-encoding:mime-version:reply-to :list-subscribe:list-help:list-post:list-archive:list-unsubscribe :list-id:precedence:subject:references:in-reply-to:message-id:date :to:from:dkim-signature:delivered-to:arc-authentication-results; bh=9/GwUR6ynfHgfXDG3wlFhLmNzJEy2lU+EhQeh35rIKQ=; b=kdx4l+ItmsIktNib4ua+9wqDH7ZjEAsNQLqJv8YNVfjFr7ZNJP9OSw5ZL+aUVgFjzm zvWDGWT8/HHjXANlaWypvb8G92N467uN9sem9GIdT8sy8WapUKIh6H7dpgksv2lIo6Fd HSQeeqwSbUEBtjCnXWnNNCsfPMpSA9V82vEN5IV94RQmb0HovlROWGs104kg71XmJMoB 078VfD3TvdyCTp5f6w6La/1FJDtXQFIKNgsGa+vUxucbi8fiMK6TjhvT3phOYHtf71pw PSCijPd9M+E/IrgxR1JWMZIrhR5KelKdP+GLKapA+WcmtBQTPhsuvjM7/B9efxR0vtuQ ZhrQ== ARC-Authentication-Results: i=1; mx.google.com; dkim=neutral (body hash did not verify) header.i=@jkqxz-net.20150623.gappssmtp.com header.s=20150623 header.b=WP3J+HP1; 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 Return-Path: Received: from ffbox0-bg.mplayerhq.hu (ffbox0-bg.ffmpeg.org. [79.124.17.100]) by mx.google.com with ESMTP id o22si17479974wrc.101.2017.11.14.11.49.18; Tue, 14 Nov 2017 11:49:18 -0800 (PST) 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=@jkqxz-net.20150623.gappssmtp.com header.s=20150623 header.b=WP3J+HP1; 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 Received: from [127.0.1.1] (localhost [127.0.0.1]) by ffbox0-bg.mplayerhq.hu (Postfix) with ESMTP id D023568A1B4; Tue, 14 Nov 2017 21:47:35 +0200 (EET) X-Original-To: ffmpeg-devel@ffmpeg.org Delivered-To: ffmpeg-devel@ffmpeg.org Received: from mail-wr0-f176.google.com (mail-wr0-f176.google.com [209.85.128.176]) by ffbox0-bg.mplayerhq.hu (Postfix) with ESMTPS id 9181068043B for ; Tue, 14 Nov 2017 21:47:32 +0200 (EET) Received: by mail-wr0-f176.google.com with SMTP id l22so18487884wrc.11 for ; Tue, 14 Nov 2017 11:47:47 -0800 (PST) DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=jkqxz-net.20150623.gappssmtp.com; s=20150623; h=from:to:subject:date:message-id:in-reply-to:references; bh=k4d4IGGq6ktVVAjD9Y1+rWnmm/q3YLSZwyjdPkJiGws=; b=WP3J+HP1GSaXyr+gqa/WXJsKleCh3JMTZxJIOv4DiA6IXrJsOu5ZZYlbqTS7BfTF20 GkNw/qtUibKg2adk2TCLJCVqjyWFWqhsx1BMxmD7ch5ymNyUnfBOiVLaDpEMLsWpu62i WkPzpb//bSud8XwGF0jPZeLTqRnu0euRHIunKc24sN+HU8y/DKWZBdj871gYQPWxoHgH NucLg20c8X08iccsyVMxG9UDMBrGh3w/XZhN0Wkl+LF/KgFswokxxRGEJbt3edHJRwem Frt19soV9MFA68xdaEVkjFEvbaT6prxsVeQjZ5ZDhWaG4VpvQJuPiH3FroMdZw4N44b6 c6RQ== X-Google-DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=1e100.net; s=20161025; h=x-gm-message-state:from:to:subject:date:message-id:in-reply-to :references; bh=k4d4IGGq6ktVVAjD9Y1+rWnmm/q3YLSZwyjdPkJiGws=; b=hJ/GZMtOzm7uWEJIov+mheGs0YM4nTcwiuaanANAfPhG8WneA0+cKtPHJ35X9v1+ov sqmTpcua/JdGf+xE8otzPzUkUeDSelcsRlIMoGY8yMCHdQguuYUX0k65gZ8uDGQ+afv9 dWSidk3YVrqRCaTYFv19H08+EGzLZH1elZ0vuQ0tphh9tqBA34NU3Irz7ZBhGcgu/vnc QWsDxk3VPRMPD69QgTNIAPyPtKVlsFOxz9UXJKd3+Jf7Gbma87ATqagU6w9KJiW0S8yJ jRCND2dUpiADeTSfS5FNQvHalSCkATet1YeUsph2gVf772nLC7A8IZWSVh1zygZLmHXV mi9A== X-Gm-Message-State: AJaThX4fovjDIW9ALCpiEeD9jebzuq0kcO28eG4Ut7b6PH1TxO7CAX7t Y7muwN76WORp6of0CYehM5fspHl7 X-Received: by 10.223.148.69 with SMTP id 63mr11876474wrq.89.1510688866479; Tue, 14 Nov 2017 11:47:46 -0800 (PST) Received: from rywe.jkqxz.net (cpc91242-cmbg18-2-0-cust650.5-4.cable.virginm.net. [82.8.130.139]) by smtp.gmail.com with ESMTPSA id v35sm38938226wrc.13.2017.11.14.11.47.45 for (version=TLS1_2 cipher=ECDHE-RSA-AES128-GCM-SHA256 bits=128/128); Tue, 14 Nov 2017 11:47:45 -0800 (PST) From: Mark Thompson To: ffmpeg-devel@ffmpeg.org Date: Tue, 14 Nov 2017 19:47:27 +0000 Message-Id: <20171114194730.11052-13-sw@jkqxz.net> X-Mailer: git-send-email 2.11.0 In-Reply-To: <20171114194730.11052-1-sw@jkqxz.net> References: <20171114194730.11052-1-sw@jkqxz.net> Subject: [FFmpeg-devel] [PATCH 12/15] lavfi: Add OpenCL unsharp mask 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 MIME-Version: 1.0 Errors-To: ffmpeg-devel-bounces@ffmpeg.org Sender: "ffmpeg-devel" Intended to replace existing opencl mode of the unsharp filter. Supports many more pixel formats and works without immediate upload and download of frame data. The options are compatible with the existing filter. --- The most obvious missing optimisation is to split the kernels into multiple versions for different pixel formats (float4 everywhere is quite wasteful), but that would make the code a lot bigger. configure | 1 + libavfilter/Makefile | 2 + libavfilter/allfilters.c | 1 + libavfilter/opencl/unsharp.cl | 99 +++++++++ libavfilter/opencl_source.h | 1 + libavfilter/vf_unsharp_opencl.c | 482 ++++++++++++++++++++++++++++++++++++++++ 6 files changed, 586 insertions(+) create mode 100644 libavfilter/opencl/unsharp.cl create mode 100644 libavfilter/vf_unsharp_opencl.c diff --git a/configure b/configure index d718cf9f4c..a38b2cc4dd 100755 --- a/configure +++ b/configure @@ -3270,6 +3270,7 @@ tinterlace_filter_deps="gpl" tinterlace_merge_test_deps="tinterlace_filter" tinterlace_pad_test_deps="tinterlace_filter" tonemap_filter_deps="const_nan" +unsharp_opencl_filter_deps="opencl" uspp_filter_deps="gpl avcodec" unsharp_filter_suggest="opencl" vaguedenoiser_filter_deps="gpl" diff --git a/libavfilter/Makefile b/libavfilter/Makefile index 8f1c5ad22c..a17bcdd749 100644 --- a/libavfilter/Makefile +++ b/libavfilter/Makefile @@ -330,6 +330,8 @@ OBJS-$(CONFIG_TRANSPOSE_FILTER) += vf_transpose.o OBJS-$(CONFIG_TRIM_FILTER) += trim.o OBJS-$(CONFIG_UNPREMULTIPLY_FILTER) += vf_premultiply.o framesync.o OBJS-$(CONFIG_UNSHARP_FILTER) += vf_unsharp.o +OBJS-$(CONFIG_UNSHARP_OPENCL_FILTER) += vf_unsharp_opencl.o opencl.o \ + opencl/unsharp.o OBJS-$(CONFIG_USPP_FILTER) += vf_uspp.o OBJS-$(CONFIG_VAGUEDENOISER_FILTER) += vf_vaguedenoiser.o OBJS-$(CONFIG_VECTORSCOPE_FILTER) += vf_vectorscope.o diff --git a/libavfilter/allfilters.c b/libavfilter/allfilters.c index 156ec019dd..1ae1681a45 100644 --- a/libavfilter/allfilters.c +++ b/libavfilter/allfilters.c @@ -338,6 +338,7 @@ static void register_all(void) REGISTER_FILTER(TRIM, trim, vf); REGISTER_FILTER(UNPREMULTIPLY, unpremultiply, vf); REGISTER_FILTER(UNSHARP, unsharp, vf); + REGISTER_FILTER(UNSHARP_OPENCL, unsharp_opencl, vf); REGISTER_FILTER(USPP, uspp, vf); REGISTER_FILTER(VAGUEDENOISER, vaguedenoiser, vf); REGISTER_FILTER(VECTORSCOPE, vectorscope, vf); diff --git a/libavfilter/opencl/unsharp.cl b/libavfilter/opencl/unsharp.cl new file mode 100644 index 0000000000..e629834e50 --- /dev/null +++ b/libavfilter/opencl/unsharp.cl @@ -0,0 +1,99 @@ +/* + * 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 unsharp_global(__write_only image2d_t dst, + __read_only image2d_t src, + int size_x, + int size_y, + float amount, + __constant float *coef_matrix) +{ + const sampler_t sampler = (CLK_NORMALIZED_COORDS_FALSE | + CLK_FILTER_NEAREST); + int2 loc = (int2)(get_global_id(0), get_global_id(1)); + int2 centre = (int2)(size_x / 2, size_y / 2); + + float4 val = read_imagef(src, sampler, loc); + float4 sum = 0.0f; + int x, y; + + for (y = 0; y < size_y; y++) { + for (x = 0; x < size_x; x++) { + int2 pos = loc + (int2)(x, y) - centre; + sum += coef_matrix[y * size_x + x] * + read_imagef(src, sampler, pos); + } + } + + write_imagef(dst, loc, val + (val - sum) * amount); +} + +__kernel void unsharp_local(__write_only image2d_t dst, + __read_only image2d_t src, + int size_x, + int size_y, + float amount, + __constant float *coef_x, + __constant float *coef_y) +{ + const sampler_t sampler = (CLK_NORMALIZED_COORDS_FALSE | + CLK_ADDRESS_CLAMP_TO_EDGE | + CLK_FILTER_NEAREST); + int2 block = (int2)(get_group_id(0), get_group_id(1)) * 16; + int2 pos = (int2)(get_local_id(0), get_local_id(1)); + + __local float4 tmp[32][32]; + + int rad_x = size_x / 2; + int rad_y = size_y / 2; + int x, y; + + for (y = 0; y <= 1; y++) { + for (x = 0; x <= 1; x++) { + tmp[pos.y + 16 * y][pos.x + 16 * x] = + read_imagef(src, sampler, block + pos + (int2)(16 * x - 8, 16 * y - 8)); + } + } + + barrier(CLK_LOCAL_MEM_FENCE); + + float4 val = tmp[pos.y + 8][pos.x + 8]; + + float4 horiz[2]; + for (y = 0; y <= 1; y++) { + horiz[y] = 0.0f; + for (x = 0; x < size_x; x++) + horiz[y] += coef_x[x] * tmp[pos.y + y * 16][pos.x + 8 + x - rad_x]; + } + + barrier(CLK_LOCAL_MEM_FENCE); + + for (y = 0; y <= 1; y++) { + tmp[pos.y + y * 16][pos.x + 8] = horiz[y]; + } + + barrier(CLK_LOCAL_MEM_FENCE); + + float4 sum = 0.0f; + for (y = 0; y < size_y; y++) + sum += coef_y[y] * tmp[pos.y + 8 + y - rad_y][pos.x + 8]; + + if (block.x + pos.x < get_image_width(dst) && + block.y + pos.y < get_image_height(dst)) + write_imagef(dst, block + pos, val + (val - sum) * amount); +} diff --git a/libavfilter/opencl_source.h b/libavfilter/opencl_source.h index e7af58bcfa..23cdfc6ac9 100644 --- a/libavfilter/opencl_source.h +++ b/libavfilter/opencl_source.h @@ -20,5 +20,6 @@ #define AVFILTER_OPENCL_SOURCE_H extern const char *ff_opencl_source_overlay; +extern const char *ff_opencl_source_unsharp; #endif /* AVFILTER_OPENCL_SOURCE_H */ diff --git a/libavfilter/vf_unsharp_opencl.c b/libavfilter/vf_unsharp_opencl.c new file mode 100644 index 0000000000..c2ebf70ad9 --- /dev/null +++ b/libavfilter/vf_unsharp_opencl.c @@ -0,0 +1,482 @@ +/* + * 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 "video.h" + +#define MAX_DIAMETER 23 + +typedef struct UnsharpOpenCLContext { + OpenCLFilterContext ocf; + + int initialised; + cl_kernel kernel; + cl_command_queue command_queue; + + float luma_size_x; + float luma_size_y; + float luma_amount; + float chroma_size_x; + float chroma_size_y; + float chroma_amount; + + int global; + + int nb_planes; + struct { + float blur_x[MAX_DIAMETER]; + float blur_y[MAX_DIAMETER]; + + cl_mem matrix; + cl_mem coef_x; + cl_mem coef_y; + + cl_int size_x; + cl_int size_y; + cl_float amount; + cl_float threshold; + } plane[4]; +} UnsharpOpenCLContext; + + +static int unsharp_opencl_init(AVFilterContext *avctx) +{ + UnsharpOpenCLContext *ctx = avctx->priv; + cl_int cle; + int err; + + err = ff_opencl_filter_load_program(avctx, &ff_opencl_source_unsharp, 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; + } + + // Use global kernel if mask size will be too big for the local store.. + ctx->global = (ctx->luma_size_x > 17.0f || + ctx->luma_size_y > 17.0f || + ctx->chroma_size_x > 17.0f || + ctx->chroma_size_y > 17.0f); + + ctx->kernel = clCreateKernel(ctx->ocf.program, + ctx->global ? "unsharp_global" + : "unsharp_local", &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 unsharp_opencl_make_filter_params(AVFilterContext *avctx) +{ + UnsharpOpenCLContext *ctx = avctx->priv; + const AVPixFmtDescriptor *desc; + float *matrix; + double val, sum; + cl_int cle; + cl_mem buffer; + size_t matrix_bytes; + float diam_x, diam_y, amount; + int err, p, x, y, size_x, size_y; + + desc = av_pix_fmt_desc_get(ctx->ocf.output_format); + + ctx->nb_planes = 0; + for (p = 0; p < desc->nb_components; p++) + ctx->nb_planes = FFMAX(ctx->nb_planes, desc->comp[p].plane + 1); + + for (p = 0; p < ctx->nb_planes; p++) { + if (p == 0 || (desc->flags & AV_PIX_FMT_FLAG_RGB)) { + diam_x = ctx->luma_size_x; + diam_y = ctx->luma_size_y; + amount = ctx->luma_amount; + } else { + diam_x = ctx->chroma_size_x; + diam_y = ctx->chroma_size_y; + amount = ctx->chroma_amount; + } + size_x = (int)ceil(diam_x) | 1; + size_y = (int)ceil(diam_y) | 1; + matrix_bytes = size_x * size_y * sizeof(float); + + matrix = av_malloc(matrix_bytes); + if (!matrix) { + err = AVERROR(ENOMEM); + goto fail; + } + + sum = 0.0; + for (x = 0; x < size_x; x++) { + double dx = (double)(x - size_x / 2) / diam_x; + sum += ctx->plane[p].blur_x[x] = exp(-16.0 * (dx * dx)); + } + for (x = 0; x < size_x; x++) + ctx->plane[p].blur_x[x] /= sum; + + sum = 0.0; + for (y = 0; y < size_y; y++) { + double dy = (double)(y - size_y / 2) / diam_y; + sum += ctx->plane[p].blur_y[y] = exp(-16.0 * (dy * dy)); + } + for (y = 0; y < size_y; y++) + ctx->plane[p].blur_y[y] /= sum; + + for (y = 0; y < size_y; y++) { + for (x = 0; x < size_x; x++) { + val = ctx->plane[p].blur_x[x] * ctx->plane[p].blur_y[y]; + matrix[y * size_x + x] = val; + } + } + + if (ctx->global) { + buffer = clCreateBuffer(ctx->ocf.hwctx->context, + CL_MEM_READ_ONLY | + CL_MEM_COPY_HOST_PTR | + CL_MEM_HOST_NO_ACCESS, + matrix_bytes, matrix, &cle); + if (!buffer) { + av_log(avctx, AV_LOG_ERROR, "Failed to create matrix buffer: " + "%d.\n", cle); + err = AVERROR(EIO); + goto fail; + } + ctx->plane[p].matrix = buffer; + } else { + buffer = clCreateBuffer(ctx->ocf.hwctx->context, + CL_MEM_READ_ONLY | + CL_MEM_COPY_HOST_PTR | + CL_MEM_HOST_NO_ACCESS, + sizeof(ctx->plane[p].blur_x), + ctx->plane[p].blur_x, &cle); + if (!buffer) { + av_log(avctx, AV_LOG_ERROR, "Failed to create x-coef buffer: " + "%d.\n", cle); + err = AVERROR(EIO); + goto fail; + } + ctx->plane[p].coef_x = buffer; + + buffer = clCreateBuffer(ctx->ocf.hwctx->context, + CL_MEM_READ_ONLY | + CL_MEM_COPY_HOST_PTR | + CL_MEM_HOST_NO_ACCESS, + sizeof(ctx->plane[p].blur_y), + ctx->plane[p].blur_y, &cle); + if (!buffer) { + av_log(avctx, AV_LOG_ERROR, "Failed to create y-coef buffer: " + "%d.\n", cle); + err = AVERROR(EIO); + goto fail; + } + ctx->plane[p].coef_y = buffer; + } + + av_freep(&matrix); + + ctx->plane[p].size_x = size_x; + ctx->plane[p].size_y = size_y; + ctx->plane[p].amount = amount; + } + + err = 0; +fail: + av_freep(&matrix); + return err; +} + +static int unsharp_opencl_filter_frame(AVFilterLink *inlink, AVFrame *input) +{ + AVFilterContext *avctx = inlink->dst; + AVFilterLink *outlink = avctx->outputs[0]; + UnsharpOpenCLContext *ctx = avctx->priv; + AVFrame *output = NULL; + cl_int cle; + size_t global_work[2]; + size_t local_work[2]; + cl_mem src, dst; + int err, p; + + 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 = unsharp_opencl_init(avctx); + if (err < 0) + goto fail; + + err = unsharp_opencl_make_filter_params(avctx); + if (err < 0) + goto fail; + } + + output = ff_get_video_buffer(outlink, outlink->w, outlink->h); + if (!output) { + err = AVERROR(ENOMEM); + goto fail; + } + + 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; + } + cle = clSetKernelArg(ctx->kernel, 2, sizeof(cl_int), &ctx->plane[p].size_x); + if (cle != CL_SUCCESS) { + av_log(avctx, AV_LOG_ERROR, "Failed to set kernel " + "matrix size argument: %d.\n", cle); + goto fail; + } + cle = clSetKernelArg(ctx->kernel, 3, sizeof(cl_int), &ctx->plane[p].size_y); + if (cle != CL_SUCCESS) { + av_log(avctx, AV_LOG_ERROR, "Failed to set kernel " + "matrix size argument: %d.\n", cle); + goto fail; + } + cle = clSetKernelArg(ctx->kernel, 4, sizeof(cl_float), &ctx->plane[p].amount); + if (cle != CL_SUCCESS) { + av_log(avctx, AV_LOG_ERROR, "Failed to set kernel " + "amount argument: %d.\n", cle); + goto fail; + } + if (ctx->global) { + cle = clSetKernelArg(ctx->kernel, 5, sizeof(cl_mem), &ctx->plane[p].matrix); + if (cle != CL_SUCCESS) { + av_log(avctx, AV_LOG_ERROR, "Failed to set kernel " + "matrix argument: %d.\n", cle); + goto fail; + } + } else { + cle = clSetKernelArg(ctx->kernel, 5, sizeof(cl_mem), &ctx->plane[p].coef_x); + if (cle != CL_SUCCESS) { + av_log(avctx, AV_LOG_ERROR, "Failed to set kernel " + "x-coef argument: %d.\n", cle); + goto fail; + } + cle = clSetKernelArg(ctx->kernel, 6, sizeof(cl_mem), &ctx->plane[p].coef_y); + if (cle != CL_SUCCESS) { + av_log(avctx, AV_LOG_ERROR, "Failed to set kernel " + "y-coef argument: %d.\n", cle); + goto fail; + } + } + + if (ctx->global) { + global_work[0] = output->width; + global_work[1] = output->height; + } else { + global_work[0] = FFALIGN(output->width, 16); + global_work[1] = FFALIGN(output->height, 16); + local_work[0] = 16; + local_work[1] = 16; + } + + av_log(avctx, AV_LOG_DEBUG, "Run kernel on plane %d " + "(%zux%zu).\n", p, global_work[0], global_work[1]); + + cle = clEnqueueNDRangeKernel(ctx->command_queue, ctx->kernel, 2, NULL, + global_work, ctx->global ? NULL : local_work, + 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; + } + + err = av_frame_copy_props(output, input); + if (err < 0) + 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 unsharp_opencl_uninit(AVFilterContext *avctx) +{ + UnsharpOpenCLContext *ctx = avctx->priv; + cl_int cle; + int i; + + for (i = 0; i < ctx->nb_planes; i++) { + if (ctx->plane[i].matrix) + clReleaseMemObject(ctx->plane[i].matrix); + if (ctx->plane[i].coef_x) + clReleaseMemObject(ctx->plane[i].coef_x); + if (ctx->plane[i].coef_y) + clReleaseMemObject(ctx->plane[i].coef_y); + } + + 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(UnsharpOpenCLContext, x) +#define FLAGS (AV_OPT_FLAG_FILTERING_PARAM | AV_OPT_FLAG_VIDEO_PARAM) +static const AVOption unsharp_opencl_options[] = { + { "luma_msize_x", "Set luma mask horizontal diameter (pixels)", + OFFSET(luma_size_x), AV_OPT_TYPE_FLOAT, + { .dbl = 5.0 }, 1, MAX_DIAMETER, FLAGS }, + { "lx", "Set luma mask horizontal diameter (pixels)", + OFFSET(luma_size_x), AV_OPT_TYPE_FLOAT, + { .dbl = 5.0 }, 1, MAX_DIAMETER, FLAGS }, + { "luma_msize_y", "Set luma mask vertical diameter (pixels)", + OFFSET(luma_size_y), AV_OPT_TYPE_FLOAT, + { .dbl = 5.0 }, 1, MAX_DIAMETER, FLAGS }, + { "ly", "Set luma mask vertical diameter (pixels)", + OFFSET(luma_size_y), AV_OPT_TYPE_FLOAT, + { .dbl = 5.0 }, 1, MAX_DIAMETER, FLAGS }, + { "luma_amount", "Set luma amount (multiplier)", + OFFSET(luma_amount), AV_OPT_TYPE_FLOAT, + { .dbl = 1.0 }, -10, 10, FLAGS }, + { "la", "Set luma amount (multiplier)", + OFFSET(luma_amount), AV_OPT_TYPE_FLOAT, + { .dbl = 1.0 }, -10, 10, FLAGS }, + + { "chroma_msize_x", "Set chroma mask horizontal diameter (pixels after subsampling)", + OFFSET(chroma_size_x), AV_OPT_TYPE_FLOAT, + { .dbl = 5.0 }, 1, MAX_DIAMETER, FLAGS }, + { "cx", "Set chroma mask horizontal diameter (pixels after subsampling)", + OFFSET(chroma_size_x), AV_OPT_TYPE_FLOAT, + { .dbl = 5.0 }, 1, MAX_DIAMETER, FLAGS }, + { "chroma_msize_y", "Set chroma mask vertical diameter (pixels after subsampling)", + OFFSET(chroma_size_y), AV_OPT_TYPE_FLOAT, + { .dbl = 5.0 }, 1, MAX_DIAMETER, FLAGS }, + { "cy", "Set chroma mask vertical diameter (pixels after subsampling)", + OFFSET(chroma_size_y), AV_OPT_TYPE_FLOAT, + { .dbl = 5.0 }, 1, MAX_DIAMETER, FLAGS }, + { "chroma_amount", "Set chroma amount (multiplier)", + OFFSET(chroma_amount), AV_OPT_TYPE_FLOAT, + { .dbl = 0.0 }, -10, 10, FLAGS }, + { "ca", "Set chroma amount (multiplier)", + OFFSET(chroma_amount), AV_OPT_TYPE_FLOAT, + { .dbl = 0.0 }, -10, 10, FLAGS }, + + { NULL } +}; + +AVFILTER_DEFINE_CLASS(unsharp_opencl); + +static const AVFilterPad unsharp_opencl_inputs[] = { + { + .name = "default", + .type = AVMEDIA_TYPE_VIDEO, + .filter_frame = &unsharp_opencl_filter_frame, + .config_props = &ff_opencl_filter_config_input, + }, + { NULL } +}; + +static const AVFilterPad unsharp_opencl_outputs[] = { + { + .name = "default", + .type = AVMEDIA_TYPE_VIDEO, + .config_props = &ff_opencl_filter_config_output, + }, + { NULL } +}; + +AVFilter ff_vf_unsharp_opencl = { + .name = "unsharp_opencl", + .description = NULL_IF_CONFIG_SMALL("Apply unsharp mask to input video"), + .priv_size = sizeof(UnsharpOpenCLContext), + .priv_class = &unsharp_opencl_class, + .init = &ff_opencl_filter_init, + .uninit = &unsharp_opencl_uninit, + .query_formats = &ff_opencl_filter_query_formats, + .inputs = unsharp_opencl_inputs, + .outputs = unsharp_opencl_outputs, + .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE, +};