From patchwork Wed Jul 25 12:13:36 2018 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Danil Iashchenko X-Patchwork-Id: 9789 Delivered-To: ffmpegpatchwork@gmail.com Received: by 2002:a02:104:0:0:0:0:0 with SMTP id c4-v6csp738590jad; Wed, 25 Jul 2018 05:21:38 -0700 (PDT) X-Google-Smtp-Source: AAOMgpebOs3PC3mCNT2U4xoOtcBYKHfoiEyW8aXPkyKjC9mTGrNL9GlYSVraXg7yDT4QcN7Cg0Ux X-Received: by 2002:a5d:4c4c:: with SMTP id n12-v6mr13850619wrt.71.1532521298769; Wed, 25 Jul 2018 05:21:38 -0700 (PDT) ARC-Seal: i=1; a=rsa-sha256; t=1532521298; cv=none; d=google.com; s=arc-20160816; b=kPE0D3ucuqGEuBghACkcXFJjddQlBEbOjhyQ+FT51QFdfxkLWXzDLBVC0zUjOO+kyQ n4wnBTu+0+GLIaGE9FcClulRX6BTENzdZpk6jxhuIt/5b/UQFfZuRx2ByxTb6qgnGRlH BboKIpp3fApiJsVh1DESXNjKfJDUJKBhZ3yX2B+9Yx2S9IrkMtIfbBfM5x4XTmOvbVTS MtT+kP8yY3rdzkFCjtjCGdQMjk5Mf30/CUAjQzXh1k4ZZkvWs19xv3ueheRCDifwi84I 5Xj3K+Td8EzuwJlup7xE+gKeJQZNhHbFQaiYKiI562uw9DAV/LNEZVtAKSMgrFl9PD9r T6/g== 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:cc:reply-to :list-subscribe:list-help:list-post:list-archive:list-unsubscribe :list-id:precedence:subject:message-id:date:to:from:dkim-signature :delivered-to:arc-authentication-results; bh=X3Ugrf+V9jXM/LC6fybN8y/AUkfq7VAhmnaB6Kna5ns=; b=ExBTGtqLdbYdPDwpxC5eatfULwYlnDZMDlJeEiUGfW0afWTXRUbM3W4hVGuJ64CWx2 NaWZLCTYMQfVh292OKt8xJ8xtIcuDnlQkY9sE3fBtdgV2zT0LJ3LtOQMnyzfDKzjWYHg ntnJu1gRi6hXGaLSePQYsZ8KTdMSTFQhwPx+3zJz+5Y9uv7QCLaNTAPHIYACH+YbCCmD uJLyhcSDhMc5JTJTxcm2WTsEfUCeB/xn82MxRC0EhFDbrQbo0ONBy60+Ws5C+++OZ6VL HUHBhI8SguZoISr/4BTneleuNmjEy7945gWubc3C+CORrdvBqAS29JTvgQy+ulLuPyhG YkAQ== ARC-Authentication-Results: i=1; mx.google.com; dkim=neutral (body hash did not verify) header.i=@gmail.com header.s=20161025 header.b=Ur4utM53; 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=gmail.com Return-Path: Received: from ffbox0-bg.mplayerhq.hu (ffbox0-bg.ffmpeg.org. [79.124.17.100]) by mx.google.com with ESMTP id o5-v6si12388599wre.145.2018.07.25.05.21.37; Wed, 25 Jul 2018 05:21:38 -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=@gmail.com header.s=20161025 header.b=Ur4utM53; 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=gmail.com Received: from [127.0.1.1] (localhost [127.0.0.1]) by ffbox0-bg.mplayerhq.hu (Postfix) with ESMTP id EADB468A380; Wed, 25 Jul 2018 15:21:21 +0300 (EEST) X-Original-To: ffmpeg-devel@ffmpeg.org Delivered-To: ffmpeg-devel@ffmpeg.org Received: from mail-lf1-f50.google.com (mail-lf1-f50.google.com [209.85.167.50]) by ffbox0-bg.mplayerhq.hu (Postfix) with ESMTPS id 8596D6899A1 for ; Wed, 25 Jul 2018 15:21:15 +0300 (EEST) Received: by mail-lf1-f50.google.com with SMTP id a4-v6so5279891lff.5 for ; Wed, 25 Jul 2018 05:21:29 -0700 (PDT) DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=gmail.com; s=20161025; h=from:to:cc:subject:date:message-id; bh=uqwCVNkbgXSg4MeFBzLcXn9gCeqjPJ0QdH43UHcavmU=; b=Ur4utM53pgdCq0st8Rct4MoivUWZgus1euiGBsN9vkrxBsYIBl1vV9X70lZw8AGRw+ fIP4AfLU0qDeMqtS0PWLEdMV6k2yRrqcLScQnZUt8H9CVHQ4raTtBQamicA0rsgUKKeg q/S95+nTzLk49RO1Ix7DyYYl8gv2Trr95b+CBdbBC+d3NdbkQmhrWcjk+TI3BDWFg21G OZhfxYYwFMKd9f4YDdrzfCPVH1ly3+3CpO1FvotKmQebPtf3FZz0czz5AXw094TxqFXz RRkw7Pf2M9eSsp+a2nn+b7t61BRqbUqSn6pCuZ9bwWSeZzJqkfFtGikngh8YcD6P10gP NZww== X-Google-DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=1e100.net; s=20161025; h=x-gm-message-state:from:to:cc:subject:date:message-id; bh=uqwCVNkbgXSg4MeFBzLcXn9gCeqjPJ0QdH43UHcavmU=; b=IlmZMxcnjaGtU2z2GOel4fkWhYFH7zl6jy655Ne/H/PnpSdxPA+7VGR83BJBTo9vjm 5Ug16U6vqu+b1OKl2J9SNBJ9+FX0B1u+XWdhXC68OtRI5/vedfZMlw34AdTJxe7AhdOO 25y6T/bxeL2/KsmXO0EOcknyW4Wuc+l8/NQdttC4Ue2vjA9d0CVrj3cp27FUrSwtH1dx 1pZfq3WGpD/W5HrBB4Y/ucvebQqKm7JcOYjT7TaGAKU4MRFSocE13lE1rN9I5JMt/MN9 6JDnUQd06S4u7QLWG40TZ3Ftq+JCBUnpkRThoquxowy4kStM1YXhF4Fw/2PL/HZzfotu cdSA== X-Gm-Message-State: AOUpUlG6lMHbU725MK9nZs5mznTj9zU4584BiHGCHqG/ul7fjL9By732 WVyxH/sAhbkVnzr8NZG6uaYgt5I= X-Received: by 2002:a19:be54:: with SMTP id o81-v6mr12366832lff.31.1532520847775; Wed, 25 Jul 2018 05:14:07 -0700 (PDT) Received: from dan-acer.lan (campus.ifmo.ru. [194.85.161.2]) by smtp.gmail.com with ESMTPSA id z10-v6sm2764272ljh.57.2018.07.25.05.14.06 (version=TLS1_2 cipher=ECDHE-RSA-AES128-SHA bits=128/128); Wed, 25 Jul 2018 05:14:07 -0700 (PDT) From: Danil Iashchenko To: ffmpeg-devel@ffmpeg.org Date: Wed, 25 Jul 2018 15:13:36 +0300 Message-Id: <1532520816-4990-1-git-send-email-danyaschenko@gmail.com> X-Mailer: git-send-email 2.7.4 Subject: [FFmpeg-devel] [PATCH] lavfi: add lumakey_opencl 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: Danil Iashchenko MIME-Version: 1.0 Errors-To: ffmpeg-devel-bounces@ffmpeg.org Sender: "ffmpeg-devel" Add lumakey_opencl filter. Behaves like existing lumakey filter. --- configure | 1 + libavfilter/Makefile | 2 + libavfilter/allfilters.c | 1 + libavfilter/opencl/lumakey.cl | 43 +++++++ libavfilter/opencl_source.h | 1 + libavfilter/vf_lumakey_opencl.c | 243 ++++++++++++++++++++++++++++++++++++++++ 6 files changed, 291 insertions(+) create mode 100644 libavfilter/opencl/lumakey.cl create mode 100644 libavfilter/vf_lumakey_opencl.c diff --git a/configure b/configure index 5783407..9816ebb 100755 --- a/configure +++ b/configure @@ -3356,6 +3356,7 @@ interlace_filter_deps="gpl" kerndeint_filter_deps="gpl" ladspa_filter_deps="ladspa libdl" lensfun_filter_deps="liblensfun version3" +lumakey_opencl_filter_deps="opencl" lv2_filter_deps="lv2" mcdeint_filter_deps="avcodec gpl" movie_filter_deps="avcodec avformat" diff --git a/libavfilter/Makefile b/libavfilter/Makefile index 5d4549e..2a01bf3 100644 --- a/libavfilter/Makefile +++ b/libavfilter/Makefile @@ -253,6 +253,8 @@ OBJS-$(CONFIG_LIBVMAF_FILTER) += vf_libvmaf.o framesync.o OBJS-$(CONFIG_LIMITER_FILTER) += vf_limiter.o OBJS-$(CONFIG_LOOP_FILTER) += f_loop.o OBJS-$(CONFIG_LUMAKEY_FILTER) += vf_lumakey.o +OBJS-$(CONFIG_LUMAKEY_OPENCL_FILTER) += vf_lumakey_opencl.o opencl.o \ + opencl/lumakey.o OBJS-$(CONFIG_LUT_FILTER) += vf_lut.o OBJS-$(CONFIG_LUT2_FILTER) += vf_lut2.o framesync.o OBJS-$(CONFIG_LUT3D_FILTER) += vf_lut3d.o diff --git a/libavfilter/allfilters.c b/libavfilter/allfilters.c index 521bc53..065ad9f 100644 --- a/libavfilter/allfilters.c +++ b/libavfilter/allfilters.c @@ -242,6 +242,7 @@ extern AVFilter ff_vf_libvmaf; extern AVFilter ff_vf_limiter; extern AVFilter ff_vf_loop; extern AVFilter ff_vf_lumakey; +extern AVFilter ff_vf_lumakey_opencl; extern AVFilter ff_vf_lut; extern AVFilter ff_vf_lut2; extern AVFilter ff_vf_lut3d; diff --git a/libavfilter/opencl/lumakey.cl b/libavfilter/opencl/lumakey.cl new file mode 100644 index 0000000..dbee63e --- /dev/null +++ b/libavfilter/opencl/lumakey.cl @@ -0,0 +1,43 @@ +/* + * Copyright (c) 2018 Danil Iashchenko + * + * 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 lumakey_global(__write_only image2d_t dstAlpha, + __read_only image2d_t srcLuma, + float w, + float b, + int so) +{ + const sampler_t sampler = (CLK_NORMALIZED_COORDS_FALSE | + CLK_ADDRESS_CLAMP_TO_EDGE | + CLK_FILTER_NEAREST); + + int2 loc = (int2)(get_global_id(0), get_global_id(1)); + float4 lumaPix = read_imagef(srcLuma, sampler, loc) * 255; + + if (lumaPix.x >= b && lumaPix.x <= w) { + write_imagef(dstAlpha, loc, (float4)(0.0f)); + } else if (lumaPix.x > b - so && lumaPix.x < w + so) { + if (lumaPix.x < b) { + write_imagef(dstAlpha, loc, (float4)((1 - (lumaPix.x - b + so) / so))); + } else { + write_imagef(dstAlpha, loc, (float4)(((lumaPix.x - w) / so))); + } + } +} diff --git a/libavfilter/opencl_source.h b/libavfilter/opencl_source.h index de4e66e..dba701e 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_colorspace_common; extern const char *ff_opencl_source_convolution; +extern const char *ff_opencl_source_lumakey; extern const char *ff_opencl_source_overlay; extern const char *ff_opencl_source_tonemap; extern const char *ff_opencl_source_unsharp; diff --git a/libavfilter/vf_lumakey_opencl.c b/libavfilter/vf_lumakey_opencl.c new file mode 100644 index 0000000..8879a88 --- /dev/null +++ b/libavfilter/vf_lumakey_opencl.c @@ -0,0 +1,243 @@ +/* + * Copyright (c) 2018 Danil Iashchenko + * + * 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/opt.h" + + +#include "avfilter.h" +#include "internal.h" +#include "opencl.h" +#include "opencl_source.h" +#include "video.h" + +typedef struct LumakeyOpenCLContext { + OpenCLFilterContext ocf; + + int initialised; + cl_kernel kernel; + cl_command_queue command_queue; + + cl_int threshold; + cl_int tolerance; + cl_int softness; + + cl_float black; + cl_float white; + +} LumakeyOpenCLContext; + +static int lumakey_opencl_init(AVFilterContext *avctx) +{ + LumakeyOpenCLContext *ctx = avctx->priv; + cl_int cle; + int err; + + err = ff_opencl_filter_load_program(avctx, &ff_opencl_source_lumakey, 1); + if (err < 0) + goto fail; + + ctx->command_queue = clCreateCommandQueue(ctx->ocf.hwctx->context, + ctx->ocf.hwctx->device_id, + 0, &cle); + CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create OpenCL " + "command queue %d.\n", cle); + + ctx->kernel = clCreateKernel(ctx->ocf.program, "lumakey_global", &cle); + CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create " + "kernel %d.\n", cle); + + ctx->initialised = 1; + return 0; + +fail: + if (ctx->command_queue) + clReleaseCommandQueue(ctx->command_queue); + if (ctx->kernel) + clReleaseKernel(ctx->kernel); + return err; +} + + +static int lumakey_opencl_filter_frame(AVFilterLink *inlink, AVFrame *input) +{ + AVFilterContext *avctx = inlink->dst; + AVFilterLink *outlink = avctx->outputs[0]; + LumakeyOpenCLContext *ctx = avctx->priv; + AVFrame *output = NULL; + cl_int cle; + size_t global_work[2]; + cl_mem src, dst; + int err, i; + size_t origin[3] = {0, 0, 0}; + size_t region[3] = {0, 0, 1}; + + 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 = lumakey_opencl_init(avctx); + if (err < 0) + goto fail; + + ctx->white = ctx->threshold + ctx->tolerance; + ctx->black = ctx->threshold - ctx->tolerance; + } + + output = ff_get_video_buffer(outlink, outlink->w, outlink->h); + if (!output) { + err = AVERROR(ENOMEM); + goto fail; + } + + for (i = 0; i < FF_ARRAY_ELEMS(output->data) - 1; i++) { + src = (cl_mem) input->data[i]; + dst = (cl_mem)output->data[i]; + + if (!dst) + break; + + err = ff_opencl_filter_work_size_from_image(avctx, region, output, i, 0); + if (err < 0) + goto fail; + + cle = clEnqueueCopyImage(ctx->command_queue, src, dst, + origin, origin, region, 0, NULL, NULL); + CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to copy plane %d: %d.\n", + i, cle); + } + + src = (cl_mem) input->data[0]; + dst = (cl_mem)output->data[3]; + + CL_SET_KERNEL_ARG(ctx->kernel, 0, cl_mem, &dst); + CL_SET_KERNEL_ARG(ctx->kernel, 1, cl_mem, &src); + CL_SET_KERNEL_ARG(ctx->kernel, 2, cl_float, &ctx->white); + CL_SET_KERNEL_ARG(ctx->kernel, 3, cl_float, &ctx->black); + CL_SET_KERNEL_ARG(ctx->kernel, 4, cl_int, &ctx->softness); + + err = ff_opencl_filter_work_size_from_image(avctx, global_work, output, 0, 0); + if (err < 0) + goto fail; + + av_log(avctx, AV_LOG_DEBUG, "Run kernel on plane %d" + "(%"SIZE_SPECIFIER"x%"SIZE_SPECIFIER").\n", + 0, global_work[0], global_work[1]); + + cle = clEnqueueNDRangeKernel(ctx->command_queue, ctx->kernel, 2, NULL, + global_work, NULL, + 0, NULL, NULL); + CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to enqueue " + "kernel: %d.\n", cle); + + + cle = clFinish(ctx->command_queue); + CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to finish command queue: %d.\n", cle); + + err = av_frame_copy_props(output, input); + 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 lumakey_opencl_uninit(AVFilterContext *avctx) +{ + LumakeyOpenCLContext *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); +} + +static const AVFilterPad lumakey_opencl_inputs[] = { + { + .name = "default", + .type = AVMEDIA_TYPE_VIDEO, + .filter_frame = &lumakey_opencl_filter_frame, + .config_props = &ff_opencl_filter_config_input, + }, + { NULL } +}; + +static const AVFilterPad lumakey_opencl_outputs[] = { + { + .name = "default", + .type = AVMEDIA_TYPE_VIDEO, + .config_props = &ff_opencl_filter_config_output, + }, + { NULL } +}; + +#define OFFSET(x) offsetof(LumakeyOpenCLContext, x) +#define FLAGS (AV_OPT_FLAG_FILTERING_PARAM | AV_OPT_FLAG_VIDEO_PARAM) + +static const AVOption lumakey_opencl_options[] = { + { "threshold", "set the threshold value", OFFSET(threshold), AV_OPT_TYPE_INT, {.i64=0}, 0, UINT16_MAX, FLAGS }, + { "tolerance", "set the tolerance value", OFFSET(tolerance), AV_OPT_TYPE_INT, {.i64=1}, 0, UINT16_MAX, FLAGS }, + { "softness", "set the softness value", OFFSET(softness), AV_OPT_TYPE_INT, {.i64=0}, 0, UINT16_MAX, FLAGS }, + { NULL } +}; + +AVFILTER_DEFINE_CLASS(lumakey_opencl); + +AVFilter ff_vf_lumakey_opencl = { + .name = "lumakey_opencl", + .description = NULL_IF_CONFIG_SMALL("Apply convolution mask to input video"), + .priv_size = sizeof(LumakeyOpenCLContext), + .priv_class = &lumakey_opencl_class, + .init = &ff_opencl_filter_init, + .uninit = &lumakey_opencl_uninit, + .query_formats = &ff_opencl_filter_query_formats, + .inputs = lumakey_opencl_inputs, + .outputs = lumakey_opencl_outputs, + .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE, +};