| Message ID | 1528242329-6121-1-git-send-email-danyaschenko@gmail.com |
|---|---|
| State | Superseded |
| Headers | show |
On 06/06/18 00:45, Danil Iashchenko wrote: > 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 ^ There's a tab here. > 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 > ... > + > +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<<s->hsub; > + var_values[VAR_VSUB] = 1<<s->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; Most of this function is duplicating code from vf_boxblur.c. Can you move it into another file (boxblur.c?) and then both filters can call it from there? > +} > + > + > 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); I don't think you should need to clFinish() after each step? Also the return value of clFinish() should be checked. > + if ((i == 0 && ctx->power[p] > 1) || (i && i == ctx->power[p] - 1)) { > + FFSWAP(cl_mem, inter, dst); > + } So the first step does src -- horizontal -> inter inter -- vertical -> dst and every step thereafter does: inter -- horizontal -> dst dst -- vertical -> inter but dst and inter got swapped after the first step? After some thought I think that does the right thing, but it could be clearer. Possibly I am not getting this right, but I think something like: i == 0 ? src : dst -- horizontal -> inter inter -- vertical -> dst would do the right thing without any swapping? > } > - 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); The options and class definition should be inside the #if. > + > + > +#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 */ > Doing some testing with this it all looks good. Thanks, - Mark
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<<s->hsub; + var_values[VAR_VSUB] = 1<<s->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 */