[FFmpeg-devel] [PATCH v2] lavfi: add colorkey_opencl filter
Mark Thompson
sw at jkqxz.net
Tue Apr 16 02:06:23 EEST 2019
On 14/04/2019 05:27, Jarek Samic wrote:
> This is a direct port of the CPU filter.
>
> Signed-off-by: Jarek Samic <cldfire3 at 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
More information about the ffmpeg-devel
mailing list