2015-09-30 3:45 GMT+08:00 Timo Rothenpieler <t...@rothenpieler.org>:
> > Signed-off-by: Timo Rothenpieler <t...@rothenpieler.org>
> > ---
> > doc/filters.texi | 5 +
> > libavfilter/chromakey_opencl_kernel.h | 98 +++++++++++++++++++
> > libavfilter/opencl_allkernels.c | 2 +
> > libavfilter/vf_chromakey.c | 179
> +++++++++++++++++++++++++++++++++-
> > 4 files changed, 283 insertions(+), 1 deletion(-)
> > create mode 100644 libavfilter/chromakey_opencl_kernel.h
> >
> > diff --git a/doc/filters.texi b/doc/filters.texi
> > index 044876c..4faf4b9 100644
> > --- a/doc/filters.texi
> > +++ b/doc/filters.texi
> > @@ -3734,6 +3734,11 @@ Signals that the color passed is already in YUV
> instead of RGB.
> >
> > Litteral colors like "green" or "red" don't make sense with this
> enabled anymore.
> > This can be used to pass exact YUV values as hexadecimal numbers.
> > +
> > +@item opencl
> > +If set to 1, specify using OpenCL capabilities, only available if
> > +FFmpeg was configured with @code{--enable-opencl}. Default value is 0.
> > +
> > @end table
> >
> > @subsection Examples
> > diff --git a/libavfilter/chromakey_opencl_kernel.h
> b/libavfilter/chromakey_opencl_kernel.h
> > new file mode 100644
> > index 0000000..56bbc79
> > --- /dev/null
> > +++ b/libavfilter/chromakey_opencl_kernel.h
> > @@ -0,0 +1,98 @@
> > +/*
> > + * Copyright (c) 2015 Timo Rothenpieler <t...@rothenpieler.org>
> > + *
> > + * This file is part of FFmpeg.
> > + *
> > + * FFmpeg is free software; you can redistribute it and/or
> > + * modify it under the terms of the GNU Lesser General Public
> > + * License as published by the Free Software Foundation; either
> > + * version 2.1 of the License, or (at your option) any later version.
> > + *
> > + * FFmpeg is distributed in the hope that it will be useful,
> > + * but WITHOUT ANY WARRANTY; without even the implied warranty of
> > + * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU
> > + * Lesser General Public License for more details.
> > + *
> > + * You should have received a copy of the GNU Lesser General Public
> > + * License along with FFmpeg; if not, write to the Free Software
> > + * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA
> 02110-1301 USA
> > + */
> > +
> > +#ifndef AVFILTER_CHROMAKEY_OPENCL_KERNEL_H
> > +#define AVFILTER_CHROMAKEY_OPENCL_KERNEL_H
> > +
> > +#include "libavutil/opencl.h"
> > +
> > +const char *ff_kernel_chromakey_opencl = AV_OPENCL_KERNEL(
> > +
> > +inline unsigned char get_pixel(global unsigned char *src,
> > + int x,
> > + int y,
> > + int w,
> > + int h,
> > + int linesize,
> > + int hsub_log2,
> > + int vsub_log2,
> > + unsigned char def)
> > +{
> > + if (x < 0 || x >= w || y < 0 || x >= w)
> > + return def;
> > +
> > + x >>= hsub_log2;
> > + y >>= vsub_log2;
> > +
> > + return src[linesize * y + x];
> > +}
> > +
> > +kernel void chromakey(global unsigned char *src_u,
> > + global unsigned char *src_v,
> > + global unsigned char *dst,
> > + int linesize_u,
> > + int linesize_v,
> > + int linesize_a,
> > + int height,
> > + int width,
> > + int hsub_log2,
> > + int vsub_log2,
> > + unsigned char chromakey_u,
> > + unsigned char chromakey_v,
> > + float similarity,
> > + float blend
> > + )
> > +{
> > + int x = get_global_id(0);
> > + int y = get_global_id(1);
> > + unsigned char res;
> > +
> > + int xo, yo, du, dv;
> > + double diff = 0.0;
> > +
> > + if (x >= width || y >= height)
> > + return;
> > +
> > + for (yo = 0; yo < 3; yo++) {
> > + for (xo = 0; xo < 3; xo++) {
> > + du = get_pixel(src_u, x + xo - 1, y + yo - 1, width,
> height, linesize_u, hsub_log2, vsub_log2, chromakey_u);
> > + dv = get_pixel(src_v, x + xo - 1, y + yo - 1, width,
> height, linesize_v, hsub_log2, vsub_log2, chromakey_v);
> > +
> > + du -= chromakey_u;
> > + dv -= chromakey_v;
> > +
> > + diff += sqrt((du * du + dv * dv) / (double)(255.0 * 255.0));
> > + }
> > + }
> > +
> > + diff /= 9.0;
> > +
> > + if (blend > 0.0001) {
> > + res = clamp((diff - similarity) / blend, 0.0, 1.0) * 255.0;
> > + } else {
> > + res = (diff > similarity) ? 255 : 0;
> > + }
> > +
> > + dst[linesize_a * y + x] = res;
> > +}
> > +
> > +);
> > +
> > +#endif /* AVFILTER_CHROMAKEY_OPENCL_KERNEL_H */
> > diff --git a/libavfilter/opencl_allkernels.c
> b/libavfilter/opencl_allkernels.c
> > index 6d80fa8..fc05e66 100644
> > --- a/libavfilter/opencl_allkernels.c
> > +++ b/libavfilter/opencl_allkernels.c
> > @@ -23,6 +23,7 @@
> > #include "libavutil/opencl.h"
> > #include "deshake_opencl_kernel.h"
> > #include "unsharp_opencl_kernel.h"
> > +#include "chromakey_opencl_kernel.h"
> > #endif
> >
> > #define OPENCL_REGISTER_KERNEL_CODE(X, x)
> \
> > @@ -37,5 +38,6 @@ void ff_opencl_register_filter_kernel_code_all(void)
> > #if CONFIG_OPENCL
> > OPENCL_REGISTER_KERNEL_CODE(DESHAKE, deshake);
> > OPENCL_REGISTER_KERNEL_CODE(UNSHARP, unsharp);
> > + OPENCL_REGISTER_KERNEL_CODE(CHROMAKEY, chromakey);
> > #endif
> > }
> > diff --git a/libavfilter/vf_chromakey.c b/libavfilter/vf_chromakey.c
> > index 47fdea631..8f15f3e 100644
> > --- a/libavfilter/vf_chromakey.c
> > +++ b/libavfilter/vf_chromakey.c
> > @@ -25,6 +25,10 @@
> > #include "internal.h"
> > #include "video.h"
> >
> > +#if CONFIG_OPENCL
> > +#include "libavutil/opencl_internal.h"
> > +#endif
> > +
> > typedef struct ChromakeyContext {
> > const AVClass *class;
> >
> > @@ -35,8 +39,152 @@ typedef struct ChromakeyContext {
> > float blend;
> >
> > int is_yuv;
> > +
> > + int opencl;
> > +
> > +#if CONFIG_OPENCL
> > + cl_command_queue command_queue;
> > + cl_program program;
> > + cl_kernel kernel;
> > +
> > + cl_mem cl_inbuf_u;
> > + size_t cl_inbuf_u_size;
> > + cl_mem cl_inbuf_v;
> > + size_t cl_inbuf_v_size;
> > + cl_mem cl_outbuf;
> > + size_t cl_outbuf_size;
> > +#endif
> > } ChromakeyContext;
> >
> > +#if CONFIG_OPENCL
> > +#define ROUND_TO_16(a) (((((a) - 1)/16)+1)*16)
> > +
> > +static av_cold int opencl_chromakey_init(AVFilterContext *avctx)
> > +{
> > + int res = 0;
> > + ChromakeyContext *ctx = avctx->priv;
> > +
> > + if (res = av_opencl_init(NULL))
> > + return res;
> > +
> > + ctx->command_queue = av_opencl_get_command_queue();
> > + if (!ctx->command_queue) {
> > + av_log(ctx, AV_LOG_ERROR, "Unable to get OpenCL command queue
> in filter 'chromakey'\n");
> > + return AVERROR(EINVAL);
> > + }
> > +
> > + ctx->program = av_opencl_compile("chromakey", NULL);
> > + if (!ctx->program) {
> > + av_log(ctx, AV_LOG_ERROR, "OpenCL failed to compile program
> 'chromakey'\n");
> > + return AVERROR(EINVAL);
> > + }
> > +
> > + ctx->kernel = clCreateKernel(ctx->program, "chromakey", &res);
> > + if (res != CL_SUCCESS) {
> > + av_log(ctx, AV_LOG_ERROR, "OpenCL failed to create kernel
> 'chromakey'\n");
> > + return AVERROR(EINVAL);
> > + }
> > +
> > + return res;
> > +}
> > +
> > +static av_cold void opencl_chromakey_uninit(AVFilterContext *avctx)
> > +{
> > + ChromakeyContext *ctx = avctx->priv;
> > +
> > + if (ctx->cl_inbuf_u)
> > + av_opencl_buffer_release(&ctx->cl_inbuf_u);
> > + if (ctx->cl_inbuf_v)
> > + av_opencl_buffer_release(&ctx->cl_inbuf_v);
> > + if (ctx->cl_outbuf)
> > + av_opencl_buffer_release(&ctx->cl_outbuf);
> > + if (ctx->kernel)
> > + clReleaseKernel(ctx->kernel);
> > + if (ctx->program)
> > + clReleaseProgram(ctx->program);
> > +
> > + ctx->command_queue = NULL;
> > +
> > + av_opencl_uninit();
> > +}
> > +
> > +static int opencl_chromakey_frame(AVFilterContext *avctx, AVFrame
> *frame)
> > +{
> > + ChromakeyContext *ctx = avctx->priv;
> > + int res = 0;
> > + int hsub_log2 = 0, vsub_log2 = 0;
> > +
> > + size_t global_work_size[2] = { (size_t)ROUND_TO_16(frame->width),
> (size_t)ROUND_TO_16(frame->height) };
> > +
> > + FFOpenclParam param = { 0 };
> > + param.ctx = avctx;
> > + param.kernel = ctx->kernel;
> > +
> > + if (frame->format == AV_PIX_FMT_YUVA420P || frame->format ==
> AV_PIX_FMT_YUVA422P)
> > + hsub_log2 = 1;
> > +
> > + if (frame->format == AV_PIX_FMT_YUVA420P)
> > + vsub_log2 = 1;
> > +
> > + if (!ctx->cl_inbuf_u || !ctx->cl_inbuf_v || !ctx->cl_outbuf) {
> > + ctx->cl_inbuf_u_size = frame->linesize[1] * (frame->height >>
> vsub_log2);
> > + ctx->cl_inbuf_v_size = frame->linesize[2] * (frame->height >>
> vsub_log2);
> > + ctx->cl_outbuf_size = frame->linesize[3] * frame->height;
> > +
> > + res = av_opencl_buffer_create(&ctx->cl_inbuf_u,
> ctx->cl_inbuf_u_size, CL_MEM_READ_ONLY, NULL);
> > + if (res)
> > + return res;
> > +
> > + res = av_opencl_buffer_create(&ctx->cl_inbuf_v,
> ctx->cl_inbuf_v_size, CL_MEM_READ_ONLY, NULL);
> > + if (res)
> > + return res;
> > +
> > + res = av_opencl_buffer_create(&ctx->cl_outbuf,
> ctx->cl_outbuf_size, CL_MEM_READ_WRITE, NULL);
> > + if (res)
> > + return res;
> > + }
> > +
> > + res = av_opencl_buffer_write(ctx->cl_inbuf_u, frame->data[1],
> ctx->cl_inbuf_u_size);
> > + if (res)
> > + return res;
> > +
> > + res = av_opencl_buffer_write(ctx->cl_inbuf_v, frame->data[2],
> ctx->cl_inbuf_v_size);
> > + if (res)
> > + return res;
> > +
> > + res = avpriv_opencl_set_parameter(¶m,
> > +
> FF_OPENCL_PARAM_INFO(ctx->cl_inbuf_u),
> > +
> FF_OPENCL_PARAM_INFO(ctx->cl_inbuf_v),
> > +
> FF_OPENCL_PARAM_INFO(ctx->cl_outbuf),
> > +
> FF_OPENCL_PARAM_INFO(frame->linesize[1]),
> > +
> FF_OPENCL_PARAM_INFO(frame->linesize[2]),
> > +
> FF_OPENCL_PARAM_INFO(frame->linesize[3]),
> > +
> FF_OPENCL_PARAM_INFO(frame->height),
> > +
> FF_OPENCL_PARAM_INFO(frame->width),
> > + FF_OPENCL_PARAM_INFO(hsub_log2),
> > + FF_OPENCL_PARAM_INFO(vsub_log2),
> > +
> FF_OPENCL_PARAM_INFO(ctx->chromakey_uv[0]),
> > +
> FF_OPENCL_PARAM_INFO(ctx->chromakey_uv[1]),
> > +
> FF_OPENCL_PARAM_INFO(ctx->similarity),
> > + FF_OPENCL_PARAM_INFO(ctx->blend),
> > + NULL);
> > + if (res)
> > + return res;
> > +
> > + res = clEnqueueNDRangeKernel(ctx->command_queue, ctx->kernel, 2,
> NULL, global_work_size, NULL, 0, NULL, NULL);
> > + if (res != CL_SUCCESS) {
> > + av_log(ctx, AV_LOG_ERROR, "OpenCL run kernel error occurred:
> %s\n", av_opencl_errstr(res));
> > + return AVERROR_EXTERNAL;
> > + }
> > +
> > + res = av_opencl_buffer_read(frame->data[3], ctx->cl_outbuf,
> ctx->cl_outbuf_size);
> > + if (res)
> > + return res;
> > +
> > + return res;
> > +}
> > +#endif
> > +
> > static uint8_t do_chromakey_pixel(ChromakeyContext *ctx, uint8_t u[9],
> uint8_t v[9])
> > {
> > double diff = 0.0;
> > @@ -110,10 +258,18 @@ static int do_chromakey_slice(AVFilterContext
> *avctx, void *arg, int jobnr, int
> > static int filter_frame(AVFilterLink *link, AVFrame *frame)
> > {
> > AVFilterContext *avctx = link->dst;
> > + ChromakeyContext *ctx = avctx->priv;
> > int res;
> >
> > - if (res = avctx->internal->execute(avctx, do_chromakey_slice,
> frame, NULL, FFMIN(frame->height, avctx->graph->nb_threads)))
> > + if (CONFIG_OPENCL && ctx->opencl) {
> > +#if CONFIG_OPENCL
> > + if (res = opencl_chromakey_frame(avctx, frame)) {
> > + return res;
> > + }
> > +#endif
> > + } else if (res = avctx->internal->execute(avctx,
> do_chromakey_slice, frame, NULL, FFMIN(frame->height,
> avctx->graph->nb_threads))) {
> > return res;
> > + }
> >
> > return ff_filter_frame(avctx->outputs[0], frame);
> > }
> > @@ -134,9 +290,28 @@ static av_cold int
> initialize_chromakey(AVFilterContext *avctx)
> > ctx->chromakey_uv[1] = RGB_TO_V(ctx->chromakey_rgba);
> > }
> >
> > + if (ctx->opencl) {
> > +#if CONFIG_OPENCL
> > + return opencl_chromakey_init(avctx);
> > +#else
> > + av_log(ctx, AV_LOG_ERROR, "OpenCL support was not enabled in
> this build, cannot be selected\n");
> > + return AVERROR(EINVAL);
> > +#endif
> > + }
> > +
> > return 0;
> > }
> >
> > +static av_cold void uninitialize_chromakey(AVFilterContext *avctx)
> > +{
> > +#if CONFIG_OPENCL
> > + ChromakeyContext *ctx = avctx->priv;
> > +
> > + if (ctx->opencl)
> > + opencl_chromakey_uninit(avctx);
> > +#endif
> > +}
> > +
> > static av_cold int query_formats(AVFilterContext *avctx)
> > {
> > static const enum AVPixelFormat pixel_fmts[] = {
> > @@ -181,6 +356,7 @@ static const AVOption chromakey_options[] = {
> > { "similarity", "set the chromakey similarity value",
> OFFSET(similarity), AV_OPT_TYPE_FLOAT, { .dbl = 0.01 }, 0.01, 1.0, FLAGS },
> > { "blend", "set the chromakey key blend value", OFFSET(blend),
> AV_OPT_TYPE_FLOAT, { .dbl = 0.0 }, 0.0, 1.0, FLAGS },
> > { "yuv", "color parameter is in yuv instead of rgb",
> OFFSET(is_yuv), AV_OPT_TYPE_BOOL, { .i64 = 0 }, 0, 1, FLAGS },
> > + { "opencl", "use OpenCL filtering capabilities", OFFSET(opencl),
> AV_OPT_TYPE_BOOL, { .i64 = 0 }, 0, 1, FLAGS },
> > { NULL }
> > };
> >
> > @@ -192,6 +368,7 @@ AVFilter ff_vf_chromakey = {
> > .priv_size = sizeof(ChromakeyContext),
> > .priv_class = &chromakey_class,
> > .init = initialize_chromakey,
> > + .uninit = uninitialize_chromakey,
> > .query_formats = query_formats,
> > .inputs = chromakey_inputs,
> > .outputs = chromakey_outputs,
> >
>
> ping once again
>
> Hi
Could you describe how to verify it, and how can I test it?
Thanks
Best regards
>
> _______________________________________________
> ffmpeg-devel mailing list
> ffmpeg-devel@ffmpeg.org
> http://ffmpeg.org/mailman/listinfo/ffmpeg-devel
>
>
_______________________________________________
ffmpeg-devel mailing list
ffmpeg-devel@ffmpeg.org
http://ffmpeg.org/mailman/listinfo/ffmpeg-devel
