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 _______________________________________________ ffmpeg-devel mailing list ffmpeg-devel@ffmpeg.org http://ffmpeg.org/mailman/listinfo/ffmpeg-devel