lavfi: Add boxblur_opencl filter

Behaves like existing boxblur filter.

1 files changed, 182 insertions, 101 deletions

diff --git a/libavfilter/vf_avgblur_opencl.c b/libavfilter/vf_avgblur_opencl.c
index 48cebb5887..d1d3eb1983 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.
  *
@@ -20,16 +21,14 @@
 
 #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"
-
+#include "boxblur.h"
 
 typedef struct AverageBlurOpenCLContext {
     OpenCLFilterContext ocf;
@@ -39,10 +38,16 @@ 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 radius[4];
+    int power[4];
+
 } AverageBlurOpenCLContext;
 
 
@@ -80,10 +85,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 +98,60 @@ 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)
+{
+    AVFilterContext    *ctx = inlink->dst;
+    AverageBlurOpenCLContext *s = ctx->priv;
+    int err, i;
+
+    err = ff_boxblur_eval_filter_params(inlink,
+                                        &s->luma_param,
+                                        &s->chroma_param,
+                                        &s->alpha_param);
+
+    if (err != 0) {
+        av_log(ctx, AV_LOG_ERROR, "Failed to evaluate "
+               "filter params: %d.\n", err);
+        return err;
+    }
+
+    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 +162,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 +176,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 +193,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 +201,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 +215,57 @@ 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;
-        }
-        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;
+        for (i = 0; i < ctx->power[p]; i++) {
+            CL_SET_KERNEL_ARG(ctx->kernel_horiz, 0, cl_mem, &inter);
+            CL_SET_KERNEL_ARG(ctx->kernel_horiz, 1, cl_mem, i == 0 ? &src : &dst);
+            if (!strcmp(avctx->filter->name, "avgblur_opencl")) {
+                CL_SET_KERNEL_ARG(ctx->kernel_horiz, 2, cl_int, &radius_x);
+            } else if (!strcmp(avctx->filter->name, "boxblur_opencl")) {
+                CL_SET_KERNEL_ARG(ctx->kernel_horiz, 2, cl_int, &ctx->radius[p]);
+            }
+
+            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);
+
+            CL_SET_KERNEL_ARG(ctx->kernel_vert, 0, cl_mem, &dst);
+            CL_SET_KERNEL_ARG(ctx->kernel_vert, 1, cl_mem, &inter);
+
+            if (!strcmp(avctx->filter->name, "avgblur_opencl")) {
+                CL_SET_KERNEL_ARG(ctx->kernel_vert, 2, cl_int, &radius_y);
+            } else if (!strcmp(avctx->filter->name, "boxblur_opencl")) {
+                CL_SET_KERNEL_ARG(ctx->kernel_vert, 2, cl_int, &ctx->radius[p]);
+            }
+
+            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 +297,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 +327,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 +338,7 @@ static const AVFilterPad avgblur_opencl_inputs[] = {
     { NULL }
 };
 
+
 static const AVFilterPad avgblur_opencl_outputs[] = {
     {
         .name         = "default",
@@ -324,6 +348,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)
+
+#if CONFIG_AVGBLUR_OPENCL_FILTER
+
+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);
+
+
 AVFilter ff_vf_avgblur_opencl = {
     .name           = "avgblur_opencl",
     .description    = NULL_IF_CONFIG_SMALL("Apply average blur filter"),
@@ -336,3 +376,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 */