git.sesse.net Git - ffmpeg/blob - libavfilter/vf_colorkey_opencl.c

   1 /*
   2  * This file is part of FFmpeg.
   3  *
   4  * FFmpeg is free software; you can redistribute it and/or
   5  * modify it under the terms of the GNU Lesser General Public
   6  * License as published by the Free Software Foundation; either
   7  * version 2.1 of the License, or (at your option) any later version.
   8  *
   9  * FFmpeg is distributed in the hope that it will be useful,
  10  * but WITHOUT ANY WARRANTY; without even the implied warranty of
  11  * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU
  12  * Lesser General Public License for more details.
  13  *
  14  * You should have received a copy of the GNU Lesser General Public
  15  * License along with FFmpeg; if not, write to the Free Software
  16  * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA
  17  */
  18
  19 #include "libavutil/opt.h"
  20 #include "libavutil/imgutils.h"
  21 #include "avfilter.h"
  22 #include "formats.h"
  23 #include "internal.h"
  24 #include "opencl.h"
  25 #include "opencl_source.h"
  26 #include "video.h"
  27
  28 typedef struct ColorkeyOpenCLContext {
  29     OpenCLFilterContext ocf;
  30     // Whether or not the above `OpenCLFilterContext` has been initialized
  31     int initialized;
  32
  33     cl_command_queue command_queue;
  34     cl_kernel kernel_colorkey;
  35
  36     // The color we are supposed to replace with transparency
  37     uint8_t colorkey_rgba[4];
  38     // Stored as a normalized float for passing to the OpenCL kernel
  39     cl_float4 colorkey_rgba_float;
  40     // Similarity percentage compared to `colorkey_rgba`, ranging from `0.01` to `1.0`
  41     // where `0.01` matches only the key color and `1.0` matches all colors
  42     float similarity;
  43     // Blending percentage where `0.0` results in fully transparent pixels, `1.0` results
  44     // in fully opaque pixels, and numbers in between result in transparency that varies
  45     // based on the similarity to the key color
  46     float blend;
  47 } ColorkeyOpenCLContext;
  48
  49 static int colorkey_opencl_init(AVFilterContext *avctx)
  50 {
  51     ColorkeyOpenCLContext *ctx = avctx->priv;
  52     cl_int cle;
  53     int err;
  54
  55     err = ff_opencl_filter_load_program(avctx, &ff_opencl_source_colorkey, 1);
  56     if (err < 0)
  57         goto fail;
  58
  59     ctx->command_queue = clCreateCommandQueue(
  60         ctx->ocf.hwctx->context,
  61         ctx->ocf.hwctx->device_id,
  62         0,
  63         &cle
  64     );
  65
  66     CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create OpenCL command queue %d.\n", cle);
  67
  68     if (ctx->blend > 0.0001) {
  69         ctx->kernel_colorkey = clCreateKernel(ctx->ocf.program, "colorkey_blend", &cle);
  70         CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create colorkey_blend kernel: %d.\n", cle);
  71     } else {
  72         ctx->kernel_colorkey = clCreateKernel(ctx->ocf.program, "colorkey", &cle);
  73         CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create colorkey kernel: %d.\n", cle);
  74     }
  75
  76     for (int i = 0; i < 4; ++i) {
  77         ctx->colorkey_rgba_float.s[i] = (float)ctx->colorkey_rgba[i] / 255.0;
  78     }
  79
  80     ctx->initialized = 1;
  81     return 0;
  82
  83 fail:
  84     if (ctx->command_queue)
  85         clReleaseCommandQueue(ctx->command_queue);
  86     if (ctx->kernel_colorkey)
  87         clReleaseKernel(ctx->kernel_colorkey);
  88     return err;
  89 }
  90
  91 static int filter_frame(AVFilterLink *link, AVFrame *input_frame)
  92 {
  93     AVFilterContext *avctx = link->dst;
  94     AVFilterLink *outlink = avctx->outputs[0];
  95     ColorkeyOpenCLContext *colorkey_ctx = avctx->priv;
  96     AVFrame *output_frame = NULL;
  97     int err;
  98     cl_int cle;
  99     size_t global_work[2];
 100     cl_mem src, dst;
 101
 102     if (!input_frame->hw_frames_ctx)
 103         return AVERROR(EINVAL);
 104
 105     if (!colorkey_ctx->initialized) {
 106         AVHWFramesContext *input_frames_ctx =
 107             (AVHWFramesContext*)input_frame->hw_frames_ctx->data;
 108         int fmt = input_frames_ctx->sw_format;
 109
 110         // Make sure the input is a format we support
 111         if (fmt != AV_PIX_FMT_ARGB &&
 112             fmt != AV_PIX_FMT_RGBA &&
 113             fmt != AV_PIX_FMT_ABGR &&
 114             fmt != AV_PIX_FMT_BGRA
 115         ) {
 116             av_log(avctx, AV_LOG_ERROR, "unsupported (non-RGB) format in colorkey_opencl.\n");
 117             err = AVERROR(ENOSYS);
 118             goto fail;
 119         }
 120
 121         err = colorkey_opencl_init(avctx);
 122         if (err < 0)
 123             goto fail;
 124     }
 125
 126     // This filter only operates on RGB data and we know that will be on the first plane
 127     src = (cl_mem)input_frame->data[0];
 128     output_frame = ff_get_video_buffer(outlink, outlink->w, outlink->h);
 129     if (!output_frame) {
 130         err = AVERROR(ENOMEM);
 131         goto fail;
 132     }
 133     dst = (cl_mem)output_frame->data[0];
 134
 135     CL_SET_KERNEL_ARG(colorkey_ctx->kernel_colorkey, 0, cl_mem, &src);
 136     CL_SET_KERNEL_ARG(colorkey_ctx->kernel_colorkey, 1, cl_mem, &dst);
 137     CL_SET_KERNEL_ARG(colorkey_ctx->kernel_colorkey, 2, cl_float4, &colorkey_ctx->colorkey_rgba_float);
 138     CL_SET_KERNEL_ARG(colorkey_ctx->kernel_colorkey, 3, float, &colorkey_ctx->similarity);
 139     if (colorkey_ctx->blend > 0.0001) {
 140         CL_SET_KERNEL_ARG(colorkey_ctx->kernel_colorkey, 4, float, &colorkey_ctx->blend);
 141     }
 142
 143     err = ff_opencl_filter_work_size_from_image(avctx, global_work, input_frame, 0, 0);
 144     if (err < 0)
 145         goto fail;
 146
 147     cle = clEnqueueNDRangeKernel(
 148         colorkey_ctx->command_queue,
 149         colorkey_ctx->kernel_colorkey,
 150         2,
 151         NULL,
 152         global_work,
 153         NULL,
 154         0,
 155         NULL,
 156         NULL
 157     );
 158
 159     CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to enqueue colorkey kernel: %d.\n", cle);
 160
 161     // Run queued kernel
 162     cle = clFinish(colorkey_ctx->command_queue);
 163     CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to finish command queue: %d.\n", cle);
 164
 165     err = av_frame_copy_props(output_frame, input_frame);
 166     if (err < 0)
 167         goto fail;
 168
 169     av_frame_free(&input_frame);
 170
 171     return ff_filter_frame(outlink, output_frame);
 172
 173 fail:
 174     clFinish(colorkey_ctx->command_queue);
 175     av_frame_free(&input_frame);
 176     av_frame_free(&output_frame);
 177     return err;
 178 }
 179
 180 static av_cold void colorkey_opencl_uninit(AVFilterContext *avctx)
 181 {
 182     ColorkeyOpenCLContext *ctx = avctx->priv;
 183     cl_int cle;
 184
 185     if (ctx->kernel_colorkey) {
 186         cle = clReleaseKernel(ctx->kernel_colorkey);
 187         if (cle != CL_SUCCESS)
 188             av_log(avctx, AV_LOG_ERROR, "Failed to release "
 189                    "kernel: %d.\n", cle);
 190     }
 191
 192     if (ctx->command_queue) {
 193         cle = clReleaseCommandQueue(ctx->command_queue);
 194         if (cle != CL_SUCCESS)
 195             av_log(avctx, AV_LOG_ERROR, "Failed to release "
 196                    "command queue: %d.\n", cle);
 197     }
 198
 199     ff_opencl_filter_uninit(avctx);
 200 }
 201
 202 static const AVFilterPad colorkey_opencl_inputs[] = {
 203     {
 204         .name = "default",
 205         .type = AVMEDIA_TYPE_VIDEO,
 206         .filter_frame = filter_frame,
 207         .config_props = &ff_opencl_filter_config_input,
 208     },
 209     { NULL }
 210 };
 211
 212 static const AVFilterPad colorkey_opencl_outputs[] = {
 213     {
 214         .name = "default",
 215         .type = AVMEDIA_TYPE_VIDEO,
 216         .config_props = &ff_opencl_filter_config_output,
 217     },
 218     { NULL }
 219 };
 220
 221 #define OFFSET(x) offsetof(ColorkeyOpenCLContext, x)
 222 #define FLAGS AV_OPT_FLAG_FILTERING_PARAM|AV_OPT_FLAG_VIDEO_PARAM
 223
 224 static const AVOption colorkey_opencl_options[] = {
 225     { "color", "set the colorkey key color", OFFSET(colorkey_rgba), AV_OPT_TYPE_COLOR, { .str = "black" }, 0, 0, FLAGS },
 226     { "similarity", "set the colorkey similarity value", OFFSET(similarity), AV_OPT_TYPE_FLOAT, { .dbl = 0.01 }, 0.01, 1.0, FLAGS },
 227     { "blend", "set the colorkey key blend value", OFFSET(blend), AV_OPT_TYPE_FLOAT, { .dbl = 0.0 }, 0.0, 1.0, FLAGS },
 228     { NULL }
 229 };
 230
 231 AVFILTER_DEFINE_CLASS(colorkey_opencl);
 232
 233 const AVFilter ff_vf_colorkey_opencl = {
 234     .name           = "colorkey_opencl",
 235     .description    = NULL_IF_CONFIG_SMALL("Turns a certain color into transparency. Operates on RGB colors."),
 236     .priv_size      = sizeof(ColorkeyOpenCLContext),
 237     .priv_class     = &colorkey_opencl_class,
 238     .init           = &ff_opencl_filter_init,
 239     .uninit         = &colorkey_opencl_uninit,
 240     .query_formats  = &ff_opencl_filter_query_formats,
 241     .inputs         = colorkey_opencl_inputs,
 242     .outputs        = colorkey_opencl_outputs,
 243     .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE
 244 };