[FFmpeg-devel] [PATCH] libavfilter/boxblur_opencl filter.
Mark Thompson
sw at jkqxz.net
Fri Jun 8 01:56:07 EEST 2018
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
More information about the ffmpeg-devel
mailing list