diff mbox

[FFmpeg-devel,v2] lavfi: add colorkey_opencl filter

Message ID 20190414042742.17782-1-cldfire3@gmail.com
State Superseded
Headers show

Commit Message

Jarek Samic April 14, 2019, 4:27 a.m. UTC
This is a direct port of the CPU filter.

Signed-off-by: Jarek Samic <cldfire3@gmail.com>
---
I've made the changes requested from the first patch. I also investigated splitting the kernel into two kernels in order to remove the blending if branch; I noticed negligible performance improvement (if any at all) with my test case and hardware, but I've left it split up as it's possible that it makes a difference with different hardware (and it's very little change in the code).

 configure                        |   1 +
 doc/filters.texi                 |  33 +++++
 libavfilter/Makefile             |   2 +
 libavfilter/allfilters.c         |   1 +
 libavfilter/opencl/colorkey.cl   |  53 +++++++
 libavfilter/opencl_source.h      |   1 +
 libavfilter/vf_colorkey_opencl.c | 243 +++++++++++++++++++++++++++++++
 7 files changed, 334 insertions(+)
 create mode 100644 libavfilter/opencl/colorkey.cl
 create mode 100644 libavfilter/vf_colorkey_opencl.c

Comments

Mark Thompson April 15, 2019, 11:06 p.m. UTC | #1
On 14/04/2019 05:27, Jarek Samic wrote:
> This is a direct port of the CPU filter.
> 
> Signed-off-by: Jarek Samic <cldfire3@gmail.com>
> ---
> I've made the changes requested from the first patch. I also investigated splitting the kernel into two kernels in order to remove the blending if branch; I noticed negligible performance improvement (if any at all) with my test case and hardware, but I've left it split up as it's possible that it makes a difference with different hardware (and it's very little change in the code).

Fair enough, that makes sense :)

>  configure                        |   1 +
>  doc/filters.texi                 |  33 +++++
>  libavfilter/Makefile             |   2 +
>  libavfilter/allfilters.c         |   1 +
>  libavfilter/opencl/colorkey.cl   |  53 +++++++
>  libavfilter/opencl_source.h      |   1 +
>  libavfilter/vf_colorkey_opencl.c | 243 +++++++++++++++++++++++++++++++
>  7 files changed, 334 insertions(+)
>  create mode 100644 libavfilter/opencl/colorkey.cl
>  create mode 100644 libavfilter/vf_colorkey_opencl.c
> 
> ...
> diff --git a/libavfilter/opencl/colorkey.cl b/libavfilter/opencl/colorkey.cl
> new file mode 100644
> index 0000000000..82ab5c8832
> --- /dev/null
> +++ b/libavfilter/opencl/colorkey.cl
> @@ -0,0 +1,53 @@
> +/*
> + * 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
> + */
> +
> +float4 get_pixel(image2d_t src, int2 loc) {
> +    const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE |
> +                                CLK_FILTER_NEAREST;

The Mali driver doesn't like this:

"""
[Parsed_colorkey_opencl_2 @ 0x83a040c0] Failed to build program: -11.
[Parsed_colorkey_opencl_2 @ 0x83a040c0] Build log:
<source>:21:21: error: declaring sampler variable in this context is not allowed
    const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE |
                    ^

error: Compiler frontend failed (error code 59)
"""

From the standard:

"""
The sampler type (sampler_t) can only be used as the type of a function argument or a
variable declared in the program scope or the outermost scope of a kernel function. The
behavior of a sampler variable declared in a non-outermost scope of a kernel function is
implementation-defined. A sampler argument or variable cannot be modified.
"""

I think move it into the program scope (and then inline the get_pixel function, since it no longer does very much).

> +
> +    return read_imagef(src, sampler, loc);
> +}
> +
> +__kernel void colorkey_blend(
> +    __read_only  image2d_t src,
> +    __write_only image2d_t dst,
> +    float4 colorkey_rgba,
> +    float similarity,
> +    float blend
> +) {
> +    int2 loc = (int2)(get_global_id(0), get_global_id(1));
> +    float4 pixel = get_pixel(src, loc);
> +    float diff = distance(pixel.xyz, colorkey_rgba.xyz);
> +
> +    pixel.s3 = clamp((diff - similarity) / blend, 0.0f, 1.0f);
> +    write_imagef(dst, loc, pixel);
> +}
> +
> +__kernel void colorkey(
> +    __read_only  image2d_t src,
> +    __write_only image2d_t dst,
> +    float4 colorkey_rgba,
> +    float similarity
> +) {
> +    int2 loc = (int2)(get_global_id(0), get_global_id(1));
> +    float4 pixel = get_pixel(src, loc);
> +    float diff = distance(pixel.xyz, colorkey_rgba.xyz);
> +
> +    pixel.s3 = (diff > similarity) ? 1.0 : 0.0;

1.0f, 0.0f.  (The compiler probably optimises this away, but it's sensible to get into the habit of always avoiding doubles.)

> +    write_imagef(dst, loc, pixel);
> +}
> diff --git a/libavfilter/opencl_source.h b/libavfilter/opencl_source.h
> index 4118138c30..51f7178cf2 100644
> --- a/libavfilter/opencl_source.h
> +++ b/libavfilter/opencl_source.h
> @@ -20,6 +20,7 @@
>  #define AVFILTER_OPENCL_SOURCE_H
>  
>  extern const char *ff_opencl_source_avgblur;
> +extern const char *ff_opencl_source_colorkey;
>  extern const char *ff_opencl_source_colorspace_common;
>  extern const char *ff_opencl_source_convolution;
>  extern const char *ff_opencl_source_neighbor;
> diff --git a/libavfilter/vf_colorkey_opencl.c b/libavfilter/vf_colorkey_opencl.c
> new file mode 100644
> index 0000000000..2790a01cae
> --- /dev/null
> +++ b/libavfilter/vf_colorkey_opencl.c
> @@ -0,0 +1,243 @@
> +/*
> + * 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
> + */
> +
> +#include "libavutil/opt.h"
> +#include "libavutil/imgutils.h"
> +#include "avfilter.h"
> +#include "formats.h"
> +#include "internal.h"
> +#include "opencl.h"
> +#include "opencl_source.h"
> +#include "video.h"
> +
> +typedef struct ColorkeyOpenCLContext {
> +    OpenCLFilterContext ocf;
> +    // Whether or not the above `OpenCLFilterContext` has been initialized
> +    int initialized;
> +
> +    cl_command_queue command_queue;
> +    cl_kernel kernel_colorkey;
> +
> +    // The color we are supposed to replace with transparency
> +    uint8_t colorkey_rgba[4];
> +    // Stored as a normalized float for passing to the OpenCL kernel
> +    cl_float4 colorkey_rgba_float;
> +    // Similarity percentage compared to `colorkey_rgba`, ranging from `0.01` to `1.0`
> +    // where `0.01` matches only the key color and `1.0` matches all colors
> +    float similarity;
> +    // Blending percentage where `0.0` results in fully transparent pixels, `1.0` results
> +    // in fully opaque pixels, and numbers in between result in transparency that varies
> +    // based on the similarity to the key color
> +    float blend;
> +} ColorkeyOpenCLContext;
> +
> +static int colorkey_opencl_init(AVFilterContext* avctx)

"AVFilterContext *avctx"

(* is part of the declarator, not the declaration-specifiers.  Consider the meaning of "struct foo* a, b;" to see why this matters.)

Also in more declarations below.

> +{
> +    ColorkeyOpenCLContext *ctx = avctx->priv;
> +    cl_int cle;
> +    int err;
> +
> +    err = ff_opencl_filter_load_program(avctx, &ff_opencl_source_colorkey, 1);
> +    if (err < 0)
> +        goto fail;
> +
> +    ctx->command_queue = clCreateCommandQueue(
> +        ctx->ocf.hwctx->context,
> +        ctx->ocf.hwctx->device_id,
> +        0, &cle
> +    );
> +
> +    CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create OpenCL command queue %d.\n", cle);
> +
> +    if (ctx->blend > 0.0001) {
> +        ctx->kernel_colorkey = clCreateKernel(ctx->ocf.program, "colorkey_blend", &cle);
> +        CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create colorkey_blend kernel: %d.\n", cle);
> +    } else {
> +        ctx->kernel_colorkey = clCreateKernel(ctx->ocf.program, "colorkey", &cle);
> +        CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create colorkey kernel: %d.\n", cle);
> +    }
> +
> +    for (int i = 0; i < 4; ++i) {
> +        ctx->colorkey_rgba_float.s[i] = (float)ctx->colorkey_rgba[i] / 255.0;
> +    }
> +
> +    ctx->initialized = 1;
> +    return 0;
> +
> +fail:
> +    if (ctx->command_queue)
> +        clReleaseCommandQueue(ctx->command_queue);
> +    if (ctx->kernel_colorkey)
> +        clReleaseKernel(ctx->kernel_colorkey);
> +    return err;
> +}
> ...

Thanks,

- Mark
Timo Rothenpieler April 16, 2019, 10:06 a.m. UTC | #2
More a general question regarding OpenCL filters than related to this 
specific one.
We are not integrating OpenCL acceleration into the relevant native 
filter anymore, like the very old original OpenCL infra worked, are we?
So that vf_colorkey would just use OpenCL acceleration on its own 
whenever available and fallback to the software path otherwise.



Timo
diff mbox

Patch

diff --git a/configure b/configure
index c2580b34c3..6515cdc149 100755
--- a/configure
+++ b/configure
@@ -3412,6 +3412,7 @@  boxblur_filter_deps="gpl"
 boxblur_opencl_filter_deps="opencl gpl"
 bs2b_filter_deps="libbs2b"
 colormatrix_filter_deps="gpl"
+colorkey_opencl_filter_deps="opencl"
 convolution_opencl_filter_deps="opencl"
 convolve_filter_deps="avcodec"
 convolve_filter_select="fft"
diff --git a/doc/filters.texi b/doc/filters.texi
index 867607d870..390c8b97cf 100644
--- a/doc/filters.texi
+++ b/doc/filters.texi
@@ -19030,6 +19030,39 @@  Apply erosion filter with threshold0 set to 30, threshold1 set 40, threshold2 se
 @end example
 @end itemize
 
+@section colorkey_opencl
+RGB colorspace color keying.
+
+The filter accepts the following options:
+
+@table @option
+@item color
+The color which will be replaced with transparency.
+
+@item similarity
+Similarity percentage with the key color.
+
+0.01 matches only the exact key color, while 1.0 matches everything.
+
+@item blend
+Blend percentage.
+
+0.0 makes pixels either fully transparent, or not transparent at all.
+
+Higher values result in semi-transparent pixels, with a higher transparency
+the more similar the pixels color is to the key color.
+@end table
+
+@subsection Examples
+
+@itemize
+@item
+Make every semi-green pixel in the input transparent with some slight blending:
+@example
+-i INPUT -vf "hwupload, colorkey_opencl=green:0.3:0.1, hwdownload" OUTPUT
+@end example
+@end itemize
+
 @section overlay_opencl
 
 Overlay one video on top of another.
diff --git a/libavfilter/Makefile b/libavfilter/Makefile
index fef6ec5c55..9589dd8747 100644
--- a/libavfilter/Makefile
+++ b/libavfilter/Makefile
@@ -176,6 +176,8 @@  OBJS-$(CONFIG_CODECVIEW_FILTER)              += vf_codecview.o
 OBJS-$(CONFIG_COLORBALANCE_FILTER)           += vf_colorbalance.o
 OBJS-$(CONFIG_COLORCHANNELMIXER_FILTER)      += vf_colorchannelmixer.o
 OBJS-$(CONFIG_COLORKEY_FILTER)               += vf_colorkey.o
+OBJS-$(CONFIG_COLORKEY_OPENCL_FILTER)        += vf_colorkey_opencl.o opencl.o \
+                                                opencl/colorkey.o
 OBJS-$(CONFIG_COLORLEVELS_FILTER)            += vf_colorlevels.o
 OBJS-$(CONFIG_COLORMATRIX_FILTER)            += vf_colormatrix.o
 OBJS-$(CONFIG_COLORSPACE_FILTER)             += vf_colorspace.o colorspace.o colorspacedsp.o
diff --git a/libavfilter/allfilters.c b/libavfilter/allfilters.c
index c51ae0f3c7..ff4eb5bf6b 100644
--- a/libavfilter/allfilters.c
+++ b/libavfilter/allfilters.c
@@ -165,6 +165,7 @@  extern AVFilter ff_vf_codecview;
 extern AVFilter ff_vf_colorbalance;
 extern AVFilter ff_vf_colorchannelmixer;
 extern AVFilter ff_vf_colorkey;
+extern AVFilter ff_vf_colorkey_opencl;
 extern AVFilter ff_vf_colorlevels;
 extern AVFilter ff_vf_colormatrix;
 extern AVFilter ff_vf_colorspace;
diff --git a/libavfilter/opencl/colorkey.cl b/libavfilter/opencl/colorkey.cl
new file mode 100644
index 0000000000..82ab5c8832
--- /dev/null
+++ b/libavfilter/opencl/colorkey.cl
@@ -0,0 +1,53 @@ 
+/*
+ * 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
+ */
+
+float4 get_pixel(image2d_t src, int2 loc) {
+    const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE |
+                                CLK_FILTER_NEAREST;
+
+    return read_imagef(src, sampler, loc);
+}
+
+__kernel void colorkey_blend(
+    __read_only  image2d_t src,
+    __write_only image2d_t dst,
+    float4 colorkey_rgba,
+    float similarity,
+    float blend
+) {
+    int2 loc = (int2)(get_global_id(0), get_global_id(1));
+    float4 pixel = get_pixel(src, loc);
+    float diff = distance(pixel.xyz, colorkey_rgba.xyz);
+
+    pixel.s3 = clamp((diff - similarity) / blend, 0.0f, 1.0f);
+    write_imagef(dst, loc, pixel);
+}
+
+__kernel void colorkey(
+    __read_only  image2d_t src,
+    __write_only image2d_t dst,
+    float4 colorkey_rgba,
+    float similarity
+) {
+    int2 loc = (int2)(get_global_id(0), get_global_id(1));
+    float4 pixel = get_pixel(src, loc);
+    float diff = distance(pixel.xyz, colorkey_rgba.xyz);
+
+    pixel.s3 = (diff > similarity) ? 1.0 : 0.0;
+    write_imagef(dst, loc, pixel);
+}
diff --git a/libavfilter/opencl_source.h b/libavfilter/opencl_source.h
index 4118138c30..51f7178cf2 100644
--- a/libavfilter/opencl_source.h
+++ b/libavfilter/opencl_source.h
@@ -20,6 +20,7 @@ 
 #define AVFILTER_OPENCL_SOURCE_H
 
 extern const char *ff_opencl_source_avgblur;
+extern const char *ff_opencl_source_colorkey;
 extern const char *ff_opencl_source_colorspace_common;
 extern const char *ff_opencl_source_convolution;
 extern const char *ff_opencl_source_neighbor;
diff --git a/libavfilter/vf_colorkey_opencl.c b/libavfilter/vf_colorkey_opencl.c
new file mode 100644
index 0000000000..2790a01cae
--- /dev/null
+++ b/libavfilter/vf_colorkey_opencl.c
@@ -0,0 +1,243 @@ 
+/*
+ * 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
+ */
+
+#include "libavutil/opt.h"
+#include "libavutil/imgutils.h"
+#include "avfilter.h"
+#include "formats.h"
+#include "internal.h"
+#include "opencl.h"
+#include "opencl_source.h"
+#include "video.h"
+
+typedef struct ColorkeyOpenCLContext {
+    OpenCLFilterContext ocf;
+    // Whether or not the above `OpenCLFilterContext` has been initialized
+    int initialized;
+
+    cl_command_queue command_queue;
+    cl_kernel kernel_colorkey;
+
+    // The color we are supposed to replace with transparency
+    uint8_t colorkey_rgba[4];
+    // Stored as a normalized float for passing to the OpenCL kernel
+    cl_float4 colorkey_rgba_float;
+    // Similarity percentage compared to `colorkey_rgba`, ranging from `0.01` to `1.0`
+    // where `0.01` matches only the key color and `1.0` matches all colors
+    float similarity;
+    // Blending percentage where `0.0` results in fully transparent pixels, `1.0` results
+    // in fully opaque pixels, and numbers in between result in transparency that varies
+    // based on the similarity to the key color
+    float blend;
+} ColorkeyOpenCLContext;
+
+static int colorkey_opencl_init(AVFilterContext* avctx)
+{
+    ColorkeyOpenCLContext *ctx = avctx->priv;
+    cl_int cle;
+    int err;
+
+    err = ff_opencl_filter_load_program(avctx, &ff_opencl_source_colorkey, 1);
+    if (err < 0)
+        goto fail;
+
+    ctx->command_queue = clCreateCommandQueue(
+        ctx->ocf.hwctx->context,
+        ctx->ocf.hwctx->device_id,
+        0, &cle
+    );
+
+    CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create OpenCL command queue %d.\n", cle);
+
+    if (ctx->blend > 0.0001) {
+        ctx->kernel_colorkey = clCreateKernel(ctx->ocf.program, "colorkey_blend", &cle);
+        CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create colorkey_blend kernel: %d.\n", cle);
+    } else {
+        ctx->kernel_colorkey = clCreateKernel(ctx->ocf.program, "colorkey", &cle);
+        CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create colorkey kernel: %d.\n", cle);
+    }
+
+    for (int i = 0; i < 4; ++i) {
+        ctx->colorkey_rgba_float.s[i] = (float)ctx->colorkey_rgba[i] / 255.0;
+    }
+
+    ctx->initialized = 1;
+    return 0;
+
+fail:
+    if (ctx->command_queue)
+        clReleaseCommandQueue(ctx->command_queue);
+    if (ctx->kernel_colorkey)
+        clReleaseKernel(ctx->kernel_colorkey);
+    return err;
+}
+
+static int filter_frame(AVFilterLink* link, AVFrame* input_frame)
+{
+    AVFilterContext* avctx = link->dst;
+    AVFilterLink* outlink = avctx->outputs[0];
+    ColorkeyOpenCLContext* colorkey_ctx = avctx->priv;
+    AVFrame* output_frame = NULL;
+    int err;
+    cl_int cle;
+    size_t global_work[2];
+    cl_mem src, dst;
+
+    if (!input_frame->hw_frames_ctx)
+        return AVERROR(EINVAL);
+
+    if (!colorkey_ctx->initialized) {
+        AVHWFramesContext *input_frames_ctx =
+            (AVHWFramesContext*)input_frame->hw_frames_ctx->data;
+        int fmt = input_frames_ctx->sw_format;
+
+        // Make sure the input is a format we support
+        if (fmt != AV_PIX_FMT_ARGB &&
+            fmt != AV_PIX_FMT_RGBA &&
+            fmt != AV_PIX_FMT_ABGR &&
+            fmt != AV_PIX_FMT_BGRA
+        ) {
+            av_log(avctx, AV_LOG_ERROR, "unsupported (non-RGB) format in colorkey_opencl.\n");
+            err = AVERROR(ENOSYS);
+            goto fail;
+        }
+
+        err = colorkey_opencl_init(avctx);
+        if (err < 0)
+            goto fail;
+    }
+
+    // This filter only operates on RGB data and we know that will be on the first plane
+    src = (cl_mem)input_frame->data[0];
+    output_frame = ff_get_video_buffer(outlink, outlink->w, outlink->h);
+    if (!output_frame) {
+        err = AVERROR(ENOMEM);
+        goto fail;
+    }
+    dst = (cl_mem)output_frame->data[0];
+
+    CL_SET_KERNEL_ARG(colorkey_ctx->kernel_colorkey, 0, cl_mem, &src);
+    CL_SET_KERNEL_ARG(colorkey_ctx->kernel_colorkey, 1, cl_mem, &dst);
+    CL_SET_KERNEL_ARG(colorkey_ctx->kernel_colorkey, 2, cl_float4, &colorkey_ctx->colorkey_rgba_float);
+    CL_SET_KERNEL_ARG(colorkey_ctx->kernel_colorkey, 3, float, &colorkey_ctx->similarity);
+    if (colorkey_ctx->blend > 0.0001) {
+        CL_SET_KERNEL_ARG(colorkey_ctx->kernel_colorkey, 4, float, &colorkey_ctx->blend);
+    }
+
+    err = ff_opencl_filter_work_size_from_image(avctx, global_work, input_frame, 0, 0);
+    if (err < 0)
+        goto fail;
+
+    cle = clEnqueueNDRangeKernel(
+        colorkey_ctx->command_queue,
+        colorkey_ctx->kernel_colorkey,
+        2,
+        NULL,
+        global_work,
+        NULL,
+        0,
+        NULL,
+        NULL
+    );
+
+    CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to enqueue colorkey kernel: %d.\n", cle);
+
+    // Run queued kernel
+    cle = clFinish(colorkey_ctx->command_queue);
+    CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to finish command queue: %d.\n", cle);
+
+    err = av_frame_copy_props(output_frame, input_frame);
+    if (err < 0)
+        goto fail;
+
+    av_frame_free(&input_frame);
+
+    return ff_filter_frame(outlink, output_frame);
+
+fail:
+    clFinish(colorkey_ctx->command_queue);
+    av_frame_free(&input_frame);
+    av_frame_free(&output_frame);
+    return err;
+}
+
+static av_cold void colorkey_opencl_uninit(AVFilterContext* avctx)
+{
+    ColorkeyOpenCLContext *ctx = avctx->priv;
+    cl_int cle;
+
+    if (ctx->kernel_colorkey) {
+        cle = clReleaseKernel(ctx->kernel_colorkey);
+        if (cle != CL_SUCCESS)
+            av_log(avctx, AV_LOG_ERROR, "Failed to release "
+                   "kernel: %d.\n", cle);
+    }
+
+    if (ctx->command_queue) {
+        cle = clReleaseCommandQueue(ctx->command_queue);
+        if (cle != CL_SUCCESS)
+            av_log(avctx, AV_LOG_ERROR, "Failed to release "
+                   "command queue: %d.\n", cle);
+    }
+
+    ff_opencl_filter_uninit(avctx);
+}
+
+static const AVFilterPad colorkey_opencl_inputs[] = {
+    {
+        .name = "default",
+        .type = AVMEDIA_TYPE_VIDEO,
+        .filter_frame = filter_frame,
+        .config_props = &ff_opencl_filter_config_input,
+    },
+    { NULL }
+};
+
+static const AVFilterPad colorkey_opencl_outputs[] = {
+    {
+        .name = "default",
+        .type = AVMEDIA_TYPE_VIDEO,
+        .config_props = &ff_opencl_filter_config_output,
+    },
+    { NULL }
+};
+
+#define OFFSET(x) offsetof(ColorkeyOpenCLContext, x)
+#define FLAGS AV_OPT_FLAG_FILTERING_PARAM|AV_OPT_FLAG_VIDEO_PARAM
+
+static const AVOption colorkey_opencl_options[] = {
+    { "color", "set the colorkey key color", OFFSET(colorkey_rgba), AV_OPT_TYPE_COLOR, { .str = "black" }, CHAR_MIN, CHAR_MAX, FLAGS },
+    { "similarity", "set the colorkey similarity value", OFFSET(similarity), AV_OPT_TYPE_FLOAT, { .dbl = 0.01 }, 0.01, 1.0, FLAGS },
+    { "blend", "set the colorkey key blend value", OFFSET(blend), AV_OPT_TYPE_FLOAT, { .dbl = 0.0 }, 0.0, 1.0, FLAGS },
+    { NULL }
+};
+
+AVFILTER_DEFINE_CLASS(colorkey_opencl);
+
+AVFilter ff_vf_colorkey_opencl = {
+    .name           = "colorkey_opencl",
+    .description    = NULL_IF_CONFIG_SMALL("Turns a certain color into transparency. Operates on RGB colors."),
+    .priv_size      = sizeof(ColorkeyOpenCLContext),
+    .priv_class     = &colorkey_opencl_class,
+    .init           = &ff_opencl_filter_init,
+    .uninit         = &colorkey_opencl_uninit,
+    .query_formats  = &ff_opencl_filter_query_formats,
+    .inputs         = colorkey_opencl_inputs,
+    .outputs        = colorkey_opencl_outputs,
+    .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE
+};