From patchwork Tue Jun 5 23:45:29 2018 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Danil Iashchenko X-Patchwork-Id: 9265 Delivered-To: ffmpegpatchwork@gmail.com Received: by 2002:a02:11c:0:0:0:0:0 with SMTP id c28-v6csp137428jad; Tue, 5 Jun 2018 16:53:16 -0700 (PDT) X-Google-Smtp-Source: ADUXVKKJ94hElu3Q6rUgkbCCv2t3zR4PwMXjCSgLMZtxCZSazF497ACjemHHmB56MaFOwBWf1F6G X-Received: by 2002:a1c:c46:: with SMTP id 67-v6mr180793wmm.79.1528242796798; Tue, 05 Jun 2018 16:53:16 -0700 (PDT) ARC-Seal: i=1; a=rsa-sha256; t=1528242796; cv=none; d=google.com; s=arc-20160816; b=x/4gobGvG83bHoi9pCdah/aL+ykXOaKOC0Ggy8cJkc0Nt5kMqXSjfiNRxIzuPb+4jJ pIlIwmZTajMhz/4Mu7E0Co3p/uHpvHaUbxop5YCn/LBEV6eiaqZE/FLWxiFGKyRrPbq5 FB7yCq3FvR9smRIYQ5FtY0IzXef8q19SLM9NePVkRZc97kdAQdohdfARcs1ukqMAXo6C VsjBJcRBFLFcONZPSFWU8E6B+6LeTWz1fN644OKROtpfiwtcV/9dCDDrUvCviIbQXWy9 tZbvMoCzo87qCwvNoLTX4a6S9p8D3Td1B5d9D+HHY7ysK+c/bpItQc8HThmC2OXJsfjw Yvqw== 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:references:in-reply-to:message-id:date :to:from:dkim-signature:delivered-to:arc-authentication-results; bh=ac91EaKgVr6fj5L69hY5RZrWavp6jBknVDTux+4ohWs=; b=fUNSDirpCk8cG5WN2pVzZ6qoHz5OFmb5OoB9bi8BC35kqatfw5T3pXKDmGEL9EA22o lhJGfvkiCOinNBYTxUn5oDADFajDSjIJxbHt5h236xN6KPjKdBWJc70flAYKLFe/Ig+f S9VKsY+yY5MO0lre/mylpZECYQ3iHZlS1iHMi7KLCTq6p9JSX6jYTu3SaCmsZ3ybkJwI QqxifX9uB6K7dbuhqxHkuNteETZaiyhJBwS8FsyTN4VJCw07xpd37gcEc4d12Mq+3RIe +glNFHUXy24y2VP9qhfnFGf0iyqTaLEO8Na1mEFCDCp8mX4ik11FOlz8Rj6xhnYrTS11 ILDg== ARC-Authentication-Results: i=1; mx.google.com; dkim=neutral (body hash did not verify) header.i=@gmail.com header.s=20161025 header.b=rMfzw3Xl; 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 q76-v6si23599940wrb.145.2018.06.05.16.53.15; Tue, 05 Jun 2018 16:53:16 -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=rMfzw3Xl; 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 80F8D68A84E; Wed, 6 Jun 2018 02:52:26 +0300 (EEST) X-Original-To: ffmpeg-devel@ffmpeg.org Delivered-To: ffmpeg-devel@ffmpeg.org Received: from mail-lf0-f68.google.com (mail-lf0-f68.google.com [209.85.215.68]) by ffbox0-bg.mplayerhq.hu (Postfix) with ESMTPS id 2095668A76B for ; Wed, 6 Jun 2018 02:52:20 +0300 (EEST) Received: by mail-lf0-f68.google.com with SMTP id d24-v6so6247349lfa.8 for ; Tue, 05 Jun 2018 16:53:07 -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:in-reply-to:references; bh=dOD2aVoldrT3JN2Dfg9iw3mb2+zDZl7oNakwBlqiAP8=; b=rMfzw3Xl8hrDVk9UhiqllqMVIHkYpt647dfNCtdKCFOTEzCkr1B06JalkD2Fab+S+M CcbT0qORPwEhWtUd6j11JjULdXrmLFokb/6jhKdsQg8gJCcrr4bdQjt/xGl/X/jogx1G ivx8FkFtRzs1OXtwvq3BrZ7ez4aOdEr74c4RPQG8pc+fCZnL6lcOuULLioUfDxif6ewX fpiCBuwF59ro1vYzUB1Ss31WhhI3NXx8OnFCoB4/0SGZ+SKAEDHBIt8OhlO6BxZTpqEP T2f4QyN/63x2vi0Ojltko8Losy4DrIuJ0QjRi/d6GXeOZ9rRljtlA6BNiD2UFrnVdgLN 4OYg== 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:in-reply-to :references; bh=dOD2aVoldrT3JN2Dfg9iw3mb2+zDZl7oNakwBlqiAP8=; b=ByhMlE2x8H/+wWiJVqOoeftFiNEBsZ1CAtsv2/2zY1HaOiDZkv10wzAbq/aPk7TvUq uVOkRG5ctk1uoG0Pw4TXCzW5fIIn4D/1ukZ282+eXy6/kNcI/WpPGaDgSnPCO7preP7f vSfvTlhZ9uPN2XTkAfl3x+0e8dnYO2Lh+M99scoOjbrMysskitQS52S6+1fldUzRoefx dyZl2k405CwvYRE+8fSf56kn5iiYfC3L9F1xLIFwECOKJok21zW6hcltzumModkdXM2R Zx2uCoWO8KIUfnHTmaRylzZfX7cwkqV2gxlYiRz7+VtCyJ+0jmCxGgarGZrIFdS1rYOE 47ZQ== X-Gm-Message-State: APt69E0SpT1K7mBsRgCqKW/qo/Xi3V0g+3K6NUVwIuB46h2zpN/+XtVH o7ir8XlPMCl2CpcuLqsUlHbHdK8= X-Received: by 2002:a19:7b0c:: with SMTP id w12-v6mr357649lfc.22.1528242334261; Tue, 05 Jun 2018 16:45:34 -0700 (PDT) Received: from dan-acer.lan (campus.ifmo.ru. [194.85.161.2]) by smtp.gmail.com with ESMTPSA id s4-v6sm2268228ljs.38.2018.06.05.16.45.33 (version=TLS1_2 cipher=ECDHE-RSA-AES128-SHA bits=128/128); Tue, 05 Jun 2018 16:45:33 -0700 (PDT) From: Danil Iashchenko To: ffmpeg-devel@ffmpeg.org Date: Wed, 6 Jun 2018 02:45:29 +0300 Message-Id: <1528242329-6121-1-git-send-email-danyaschenko@gmail.com> X-Mailer: git-send-email 2.7.4 In-Reply-To: <148B1B7A67D1C24B9EF0BE42EA49770684F99773@SHSMSX103.ccr.corp.intel.com> References: <148B1B7A67D1C24B9EF0BE42EA49770684F99773@SHSMSX103.ccr.corp.intel.com> Subject: [FFmpeg-devel] [PATCH] libavfilter/boxblur_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" Behaves like existing boxblur filter. --- Thanks! Fixed. libavfilter/Makefile | 2 + libavfilter/allfilters.c | 1 + libavfilter/vf_avgblur_opencl.c | 419 ++++++++++++++++++++++++++++++---------- 3 files changed, 324 insertions(+), 98 deletions(-) diff --git a/libavfilter/Makefile b/libavfilter/Makefile index c68ef05..6f00059 100644 --- a/libavfilter/Makefile +++ b/libavfilter/Makefile @@ -153,6 +153,8 @@ OBJS-$(CONFIG_BLACKDETECT_FILTER) += vf_blackdetect.o OBJS-$(CONFIG_BLACKFRAME_FILTER) += vf_blackframe.o OBJS-$(CONFIG_BLEND_FILTER) += vf_blend.o framesync.o OBJS-$(CONFIG_BOXBLUR_FILTER) += vf_boxblur.o +OBJS-$(CONFIG_BOXBLUR_OPENCL_FILTER) += vf_avgblur_opencl.o opencl.o \ + opencl/avgblur.o OBJS-$(CONFIG_BWDIF_FILTER) += vf_bwdif.o OBJS-$(CONFIG_CHROMAKEY_FILTER) += vf_chromakey.o OBJS-$(CONFIG_CIESCOPE_FILTER) += vf_ciescope.o diff --git a/libavfilter/allfilters.c b/libavfilter/allfilters.c index b44093d..97d92a0 100644 --- a/libavfilter/allfilters.c +++ b/libavfilter/allfilters.c @@ -146,6 +146,7 @@ extern AVFilter ff_vf_blackdetect; extern AVFilter ff_vf_blackframe; extern AVFilter ff_vf_blend; extern AVFilter ff_vf_boxblur; +extern AVFilter ff_vf_boxblur_opencl; extern AVFilter ff_vf_bwdif; extern AVFilter ff_vf_chromakey; extern AVFilter ff_vf_ciescope; diff --git a/libavfilter/vf_avgblur_opencl.c b/libavfilter/vf_avgblur_opencl.c index 48cebb5..d4759de 100644 --- a/libavfilter/vf_avgblur_opencl.c +++ b/libavfilter/vf_avgblur_opencl.c @@ -1,5 +1,6 @@ /* * Copyright (c) 2018 Dylan Fernando + * Copyright (c) 2018 Danil Iashchenko * * This file is part of FFmpeg. * @@ -23,6 +24,7 @@ #include "libavutil/mem.h" #include "libavutil/opt.h" #include "libavutil/pixdesc.h" +#include "libavutil/eval.h" #include "avfilter.h" #include "internal.h" @@ -31,6 +33,35 @@ #include "video.h" +static const char *const var_names[] = { + "w", + "h", + "cw", + "ch", + "hsub", + "vsub", + NULL +}; + + +enum var_name { + VAR_W, + VAR_H, + VAR_CW, + VAR_CH, + VAR_HSUB, + VAR_VSUB, + VARS_NB +}; + + +typedef struct FilterParam { + int radius; + int power; + char *radius_expr; +} FilterParam; + + typedef struct AverageBlurOpenCLContext { OpenCLFilterContext ocf; @@ -39,13 +70,26 @@ typedef struct AverageBlurOpenCLContext { cl_kernel kernel_vert; cl_command_queue command_queue; - int radius; + int radiusH; int radiusV; int planes; + FilterParam luma_param; + FilterParam chroma_param; + FilterParam alpha_param; + int hsub, vsub; + int radius[4]; + int power[4]; + } AverageBlurOpenCLContext; +#define Y 0 +#define U 1 +#define V 2 +#define A 3 + + static int avgblur_opencl_init(AVFilterContext *avctx) { AverageBlurOpenCLContext *ctx = avctx->priv; @@ -80,10 +124,6 @@ static int avgblur_opencl_init(AVFilterContext *avctx) goto fail; } - if (ctx->radiusV <= 0) { - ctx->radiusV = ctx->radius; - } - ctx->initialised = 1; return 0; @@ -97,6 +137,122 @@ fail: return err; } + +static int avgblur_opencl_make_filter_params(AVFilterLink *inlink) +{ + AVFilterContext *ctx = inlink->dst; + AverageBlurOpenCLContext *s = ctx->priv; + int i; + + if (s->radiusV <= 0) { + s->radiusV = s->radiusH; + } + + for (i = 0; i < 4; i++) { + s->power[i] = 1; + } + return 0; +} + + +static int boxblur_opencl_make_filter_params(AVFilterLink *inlink) +{ + const AVPixFmtDescriptor *desc = av_pix_fmt_desc_get(inlink->format); + AVFilterContext *ctx = inlink->dst; + AverageBlurOpenCLContext *s = ctx->priv; + int w = inlink->w, h = inlink->h; + int cw, ch; + double var_values[VARS_NB], res; + char *expr; + int ret, i; + + if (!s->luma_param.radius_expr) { + av_log(s, AV_LOG_ERROR, "Luma radius expression is not set.\n"); + return AVERROR(EINVAL); + } + + /* fill missing params */ + if (!s->chroma_param.radius_expr) { + s->chroma_param.radius_expr = av_strdup(s->luma_param.radius_expr); + if (!s->chroma_param.radius_expr) + return AVERROR(ENOMEM); + } + if (s->chroma_param.power < 0) + s->chroma_param.power = s->luma_param.power; + + if (!s->alpha_param.radius_expr) { + s->alpha_param.radius_expr = av_strdup(s->luma_param.radius_expr); + if (!s->alpha_param.radius_expr) + return AVERROR(ENOMEM); + } + if (s->alpha_param.power < 0) + s->alpha_param.power = s->luma_param.power; + + s->hsub = desc->log2_chroma_w; + s->vsub = desc->log2_chroma_h; + + var_values[VAR_W] = inlink->w; + var_values[VAR_H] = inlink->h; + var_values[VAR_CW] = cw = w>>s->hsub; + var_values[VAR_CH] = ch = h>>s->vsub; + var_values[VAR_HSUB] = 1<hsub; + var_values[VAR_VSUB] = 1<vsub; + +#define EVAL_RADIUS_EXPR(comp) \ + expr = s->comp##_param.radius_expr; \ + ret = av_expr_parse_and_eval(&res, expr, var_names, var_values, \ + NULL, NULL, NULL, NULL, NULL, 0, ctx); \ + s->comp##_param.radius = res; \ + if (ret < 0) { \ + av_log(NULL, AV_LOG_ERROR, \ + "Error when evaluating " #comp " radius expression '%s'\n", expr); \ + return ret; \ + } + EVAL_RADIUS_EXPR(luma); + EVAL_RADIUS_EXPR(chroma); + EVAL_RADIUS_EXPR(alpha); + + av_log(ctx, AV_LOG_VERBOSE, + "luma_radius:%d luma_power:%d " + "chroma_radius:%d chroma_power:%d " + "alpha_radius:%d alpha_power:%d " + "w:%d chroma_w:%d h:%d chroma_h:%d\n", + s->luma_param .radius, s->luma_param .power, + s->chroma_param.radius, s->chroma_param.power, + s->alpha_param .radius, s->alpha_param .power, + w, cw, h, ch); + +#define CHECK_RADIUS_VAL(w_, h_, comp) \ + if (s->comp##_param.radius < 0 || \ + 2*s->comp##_param.radius > FFMIN(w_, h_)) { \ + av_log(ctx, AV_LOG_ERROR, \ + "Invalid " #comp " radius value %d, must be >= 0 and <= %d\n", \ + s->comp##_param.radius, FFMIN(w_, h_)/2); \ + return AVERROR(EINVAL); \ + } + CHECK_RADIUS_VAL(w, h, luma); + CHECK_RADIUS_VAL(cw, ch, chroma); + CHECK_RADIUS_VAL(w, h, alpha); + + s->radius[Y] = s->luma_param.radius; + s->radius[U] = s->radius[V] = s->chroma_param.radius; + s->radius[A] = s->alpha_param.radius; + + s->power[Y] = s->luma_param.power; + s->power[U] = s->power[V] = s->chroma_param.power; + s->power[A] = s->alpha_param.power; + + for (i = 0; i < 4; i++) { + if (s->power[i] == 0) { + s->power[i] = 1; + s->radius[i] = 0; + } + } + + return 0; +} + + static int avgblur_opencl_filter_frame(AVFilterLink *inlink, AVFrame *input) { AVFilterContext *avctx = inlink->dst; @@ -107,7 +263,7 @@ static int avgblur_opencl_filter_frame(AVFilterLink *inlink, AVFrame *input) cl_int cle; size_t global_work[2]; cl_mem src, dst, inter; - int err, p, radius_x, radius_y; + int err, p, radius_x, radius_y, i; av_log(ctx, AV_LOG_DEBUG, "Filter input: %s, %ux%u (%"PRId64").\n", av_get_pix_fmt_name(input->format), @@ -121,6 +277,16 @@ static int avgblur_opencl_filter_frame(AVFilterLink *inlink, AVFrame *input) if (err < 0) goto fail; + if (!strcmp(avctx->filter->name, "avgblur_opencl")) { + err = avgblur_opencl_make_filter_params(inlink); + if (err < 0) + goto fail; + } else if (!strcmp(avctx->filter->name, "boxblur_opencl")) { + err = boxblur_opencl_make_filter_params(inlink); + if (err < 0) + goto fail; + } + } output = ff_get_video_buffer(outlink, outlink->w, outlink->h); @@ -128,7 +294,6 @@ static int avgblur_opencl_filter_frame(AVFilterLink *inlink, AVFrame *input) err = AVERROR(ENOMEM); goto fail; } - intermediate = ff_get_video_buffer(outlink, outlink->w, outlink->h); if (!intermediate) { err = AVERROR(ENOMEM); @@ -137,13 +302,13 @@ static int avgblur_opencl_filter_frame(AVFilterLink *inlink, AVFrame *input) for (p = 0; p < FF_ARRAY_ELEMS(output->data); p++) { src = (cl_mem) input->data[p]; - dst = (cl_mem)output->data[p]; - inter = (cl_mem) intermediate->data[p]; + dst = (cl_mem) output->data[p]; + inter = (cl_mem)intermediate->data[p]; if (!dst) break; - radius_x = ctx->radius; + radius_x = ctx->radiusH; radius_y = ctx->radiusV; if (!(ctx->planes & (1 << p))) { @@ -151,88 +316,98 @@ static int avgblur_opencl_filter_frame(AVFilterLink *inlink, AVFrame *input) radius_y = 0; } - cle = clSetKernelArg(ctx->kernel_horiz, 0, sizeof(cl_mem), &inter); - if (cle != CL_SUCCESS) { - av_log(avctx, AV_LOG_ERROR, "Failed to set kernel " - "destination image argument: %d.\n", cle); - err = AVERROR_UNKNOWN; - goto fail; - } - cle = clSetKernelArg(ctx->kernel_horiz, 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); - err = AVERROR_UNKNOWN; - goto fail; - } - cle = clSetKernelArg(ctx->kernel_horiz, 2, sizeof(cl_int), &radius_x); - if (cle != CL_SUCCESS) { - av_log(avctx, AV_LOG_ERROR, "Failed to set kernel " - "sizeX argument: %d.\n", cle); - err = AVERROR_UNKNOWN; - goto fail; - } - - err = ff_opencl_filter_work_size_from_image(avctx, global_work, - intermediate, p, 0); - if (err < 0) - goto fail; - 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_horiz, 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 = clSetKernelArg(ctx->kernel_vert, 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); - err = AVERROR_UNKNOWN; - goto fail; + for (i = 0; i < ctx->power[p]; i++) { + cle = clSetKernelArg(ctx->kernel_horiz, 0, sizeof(cl_mem), i == 0 ? &inter : &dst); + if (cle != CL_SUCCESS) { + av_log(avctx, AV_LOG_ERROR, "Failed to set kernel " + "destination image argument: %d.\n", cle); + err = AVERROR_UNKNOWN; + goto fail; + } + cle = clSetKernelArg(ctx->kernel_horiz, 1, sizeof(cl_mem), i == 0 ? &src : &inter); + if (cle != CL_SUCCESS) { + av_log(avctx, AV_LOG_ERROR, "Failed to set kernel " + "source image argument: %d.\n", cle); + err = AVERROR_UNKNOWN; + goto fail; + } + + if (!strcmp(avctx->filter->name, "avgblur_opencl")) { + cle = clSetKernelArg(ctx->kernel_horiz, 2, sizeof(cl_int), &radius_x); + } else if (!strcmp(avctx->filter->name, "boxblur_opencl")) { + cle = clSetKernelArg(ctx->kernel_horiz, 2, sizeof(cl_int), &ctx->radius[p]); + } + if (cle != CL_SUCCESS) { + av_log(avctx, AV_LOG_ERROR, "Failed to set kernel " + "radius argument: %d.\n", cle); + err = AVERROR_UNKNOWN; + goto fail; + } + err = ff_opencl_filter_work_size_from_image(avctx, global_work, + i == 0 ? intermediate : output, p, 0); + if (err < 0) + goto fail; + + cle = clEnqueueNDRangeKernel(ctx->command_queue, ctx->kernel_horiz, 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); + + err = ff_opencl_filter_work_size_from_image(avctx, global_work, + i == 0 ? output : intermediate, p, 0); + + + cle = clSetKernelArg(ctx->kernel_vert, 0, sizeof(cl_mem), i == 0 ? &dst : &inter); + + if (cle != CL_SUCCESS) { + av_log(avctx, AV_LOG_ERROR, "Failed to set kernel " + "destination image argument: %d.\n", cle); + err = AVERROR_UNKNOWN; + goto fail; + } + cle = clSetKernelArg(ctx->kernel_vert, 1, sizeof(cl_mem), i == 0 ? &inter : &dst); + if (cle != CL_SUCCESS) { + av_log(avctx, AV_LOG_ERROR, "Failed to set kernel " + "source image argument: %d.\n", cle); + err = AVERROR_UNKNOWN; + goto fail; + } + if (!strcmp(avctx->filter->name, "avgblur_opencl")) { + cle = clSetKernelArg(ctx->kernel_vert, 2, sizeof(cl_int), &radius_y); + } else if (!strcmp(avctx->filter->name, "boxblur_opencl")) { + cle = clSetKernelArg(ctx->kernel_vert, 2, sizeof(cl_int), &ctx->radius[p]); + } + if (cle != CL_SUCCESS) { + av_log(avctx, AV_LOG_ERROR, "Failed to set kernel " + "radius argument: %d.\n", cle); + err = AVERROR_UNKNOWN; + goto fail; + } + + cle = clEnqueueNDRangeKernel(ctx->command_queue, ctx->kernel_vert, 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 ((i == 0 && ctx->power[p] > 1) || (i && i == ctx->power[p] - 1)) { + FFSWAP(cl_mem, inter, dst); + } } - cle = clSetKernelArg(ctx->kernel_vert, 1, sizeof(cl_mem), &inter); - if (cle != CL_SUCCESS) { - av_log(avctx, AV_LOG_ERROR, "Failed to set kernel " - "source image argument: %d.\n", cle); - err = AVERROR_UNKNOWN; - goto fail; - } - cle = clSetKernelArg(ctx->kernel_vert, 2, sizeof(cl_int), &radius_y); - if (cle != CL_SUCCESS) { - av_log(avctx, AV_LOG_ERROR, "Failed to set kernel " - "sizeY argument: %d.\n", cle); - err = AVERROR_UNKNOWN; - goto fail; - } - - err = ff_opencl_filter_work_size_from_image(avctx, global_work, - output, p, 0); - if (err < 0) - goto fail; - - 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_vert, 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); @@ -264,12 +439,12 @@ fail: return err; } + static av_cold void avgblur_opencl_uninit(AVFilterContext *avctx) { AverageBlurOpenCLContext *ctx = avctx->priv; cl_int cle; - if (ctx->kernel_horiz) { cle = clReleaseKernel(ctx->kernel_horiz); if (cle != CL_SUCCESS) @@ -294,16 +469,6 @@ static av_cold void avgblur_opencl_uninit(AVFilterContext *avctx) ff_opencl_filter_uninit(avctx); } -#define OFFSET(x) offsetof(AverageBlurOpenCLContext, x) -#define FLAGS (AV_OPT_FLAG_FILTERING_PARAM | AV_OPT_FLAG_VIDEO_PARAM) -static const AVOption avgblur_opencl_options[] = { - { "sizeX", "set horizontal size", OFFSET(radius), AV_OPT_TYPE_INT, {.i64=1}, 1, 1024, FLAGS }, - { "planes", "set planes to filter", OFFSET(planes), AV_OPT_TYPE_INT, {.i64=0xF}, 0, 0xF, FLAGS }, - { "sizeY", "set vertical size", OFFSET(radiusV), AV_OPT_TYPE_INT, {.i64=0}, 0, 1024, FLAGS }, - { NULL } -}; - -AVFILTER_DEFINE_CLASS(avgblur_opencl); static const AVFilterPad avgblur_opencl_inputs[] = { { @@ -315,6 +480,7 @@ static const AVFilterPad avgblur_opencl_inputs[] = { { NULL } }; + static const AVFilterPad avgblur_opencl_outputs[] = { { .name = "default", @@ -324,6 +490,22 @@ static const AVFilterPad avgblur_opencl_outputs[] = { { NULL } }; + +#define OFFSET(x) offsetof(AverageBlurOpenCLContext, x) +#define FLAGS (AV_OPT_FLAG_FILTERING_PARAM | AV_OPT_FLAG_VIDEO_PARAM) + +static const AVOption avgblur_opencl_options[] = { + { "sizeX", "set horizontal size", OFFSET(radiusH), AV_OPT_TYPE_INT, {.i64=1}, 1, 1024, FLAGS }, + { "planes", "set planes to filter", OFFSET(planes), AV_OPT_TYPE_INT, {.i64=0xF}, 0, 0xF, FLAGS }, + { "sizeY", "set vertical size", OFFSET(radiusV), AV_OPT_TYPE_INT, {.i64=0}, 0, 1024, FLAGS }, + { NULL } +}; + +AVFILTER_DEFINE_CLASS(avgblur_opencl); + + +#if CONFIG_AVGBLUR_OPENCL_FILTER + AVFilter ff_vf_avgblur_opencl = { .name = "avgblur_opencl", .description = NULL_IF_CONFIG_SMALL("Apply average blur filter"), @@ -336,3 +518,44 @@ AVFilter ff_vf_avgblur_opencl = { .outputs = avgblur_opencl_outputs, .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE, }; + +#endif /* CONFIG_AVGBLUR_OPENCL_FILTER */ + + +#if CONFIG_BOXBLUR_OPENCL_FILTER + +static const AVOption boxblur_opencl_options[] = { + { "luma_radius", "Radius of the luma blurring box", OFFSET(luma_param.radius_expr), AV_OPT_TYPE_STRING, {.str="2"}, .flags = FLAGS }, + { "lr", "Radius of the luma blurring box", OFFSET(luma_param.radius_expr), AV_OPT_TYPE_STRING, {.str="2"}, .flags = FLAGS }, + { "luma_power", "How many times should the boxblur be applied to luma", OFFSET(luma_param.power), AV_OPT_TYPE_INT, {.i64=2}, 0, INT_MAX, .flags = FLAGS }, + { "lp", "How many times should the boxblur be applied to luma", OFFSET(luma_param.power), AV_OPT_TYPE_INT, {.i64=2}, 0, INT_MAX, .flags = FLAGS }, + + { "chroma_radius", "Radius of the chroma blurring box", OFFSET(chroma_param.radius_expr), AV_OPT_TYPE_STRING, {.str=NULL}, .flags = FLAGS }, + { "cr", "Radius of the chroma blurring box", OFFSET(chroma_param.radius_expr), AV_OPT_TYPE_STRING, {.str=NULL}, .flags = FLAGS }, + { "chroma_power", "How many times should the boxblur be applied to chroma", OFFSET(chroma_param.power), AV_OPT_TYPE_INT, {.i64=-1}, -1, INT_MAX, .flags = FLAGS }, + { "cp", "How many times should the boxblur be applied to chroma", OFFSET(chroma_param.power), AV_OPT_TYPE_INT, {.i64=-1}, -1, INT_MAX, .flags = FLAGS }, + + { "alpha_radius", "Radius of the alpha blurring box", OFFSET(alpha_param.radius_expr), AV_OPT_TYPE_STRING, {.str=NULL}, .flags = FLAGS }, + { "ar", "Radius of the alpha blurring box", OFFSET(alpha_param.radius_expr), AV_OPT_TYPE_STRING, {.str=NULL}, .flags = FLAGS }, + { "alpha_power", "How many times should the boxblur be applied to alpha", OFFSET(alpha_param.power), AV_OPT_TYPE_INT, {.i64=-1}, -1, INT_MAX, .flags = FLAGS }, + { "ap", "How many times should the boxblur be applied to alpha", OFFSET(alpha_param.power), AV_OPT_TYPE_INT, {.i64=-1}, -1, INT_MAX, .flags = FLAGS }, + + { NULL } +}; + +AVFILTER_DEFINE_CLASS(boxblur_opencl); + +AVFilter ff_vf_boxblur_opencl = { + .name = "boxblur_opencl", + .description = NULL_IF_CONFIG_SMALL("Apply boxblur filter to input video"), + .priv_size = sizeof(AverageBlurOpenCLContext), + .priv_class = &boxblur_opencl_class, + .init = &ff_opencl_filter_init, + .uninit = &avgblur_opencl_uninit, + .query_formats = &ff_opencl_filter_query_formats, + .inputs = avgblur_opencl_inputs, + .outputs = avgblur_opencl_outputs, + .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE, +}; + +#endif /* CONFIG_BOXBLUR_OPENCL_FILTER */