avfilter/vf_chromakey: Add OpenCL acceleration

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

libavfilter/chromakey_opencl_kernel.h

@@ -0,0 +1,95 @@
+/*
+ * Copyright (c) 2015 Timo Rothenpieler <timo@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;
+
+    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 */

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
 }

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,150 @@ 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
+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)frame->width, (size_t)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(&param,
+                                      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 +256,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 +288,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 +354,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 +366,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,