FFmpeg  4.3.6
vf_colorkey_opencl.c
Go to the documentation of this file.
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 {
30  // Whether or not the above `OpenCLFilterContext` has been initialized
32 
33  cl_command_queue command_queue;
34  cl_kernel kernel_colorkey;
35 
36  // The color we are supposed to replace with transparency
38  // Stored as a normalized float for passing to the OpenCL kernel
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;
48 
50 {
51  ColorkeyOpenCLContext *ctx = avctx->priv;
52  cl_int cle;
53  int err;
54 
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;
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 
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 
200 }
201 
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 
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 
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 
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,
241  .inputs = colorkey_opencl_inputs,
242  .outputs = colorkey_opencl_outputs,
243  .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE
244 };
#define NULL
Definition: coverity.c:32
#define FF_FILTER_FLAG_HWFRAME_AWARE
The filter is aware of hardware frames, and any hardware frame context should not be automatically pr...
Definition: internal.h:365
const char * ff_opencl_source_colorkey
Definition: colorkey.c:2
This structure describes decoded (raw) audio or video data.
Definition: frame.h:300
AVOption.
Definition: opt.h:246
int ff_opencl_filter_work_size_from_image(AVFilterContext *avctx, size_t *work_size, AVFrame *frame, int plane, int block_alignment)
Find the work size needed needed for a given plane of an image.
Definition: opencl.c:278
int ff_opencl_filter_config_input(AVFilterLink *inlink)
Check that the input link contains a suitable hardware frames context and extract the device from it...
Definition: opencl.c:60
int ff_opencl_filter_query_formats(AVFilterContext *avctx)
Return that all inputs and outputs support only AV_PIX_FMT_OPENCL.
Definition: opencl.c:28
cl_command_queue command_queue
misc image utilities
Main libavfilter public API header.
static av_cold int init(AVCodecContext *avctx)
Definition: avrndec.c:35
AVFrame * ff_get_video_buffer(AVFilterLink *link, int w, int h)
Request a picture buffer with a specific set of permissions.
Definition: video.c:104
static const AVFilterPad colorkey_opencl_inputs[]
int ff_opencl_filter_init(AVFilterContext *avctx)
Initialise an OpenCL filter context.
Definition: opencl.c:147
AVOpenCLDeviceContext * hwctx
Definition: opencl.h:41
AVBufferRef * hw_frames_ctx
For hwaccel-format frames, this should be a reference to the AVHWFramesContext describing the frame...
Definition: frame.h:639
const char * name
Pad name.
Definition: internal.h:60
static const AVFilterPad colorkey_opencl_outputs[]
int ff_filter_frame(AVFilterLink *link, AVFrame *frame)
Send a frame of data to the next filter.
Definition: avfilter.c:1075
uint8_t
#define av_cold
Definition: attributes.h:88
static av_cold void colorkey_opencl_uninit(AVFilterContext *avctx)
static av_cold int uninit(AVCodecContext *avctx)
Definition: crystalhd.c:279
AVOptions.
cl_device_id device_id
The primary device ID of the device.
packed ABGR 8:8:8:8, 32bpp, ABGRABGR...
Definition: pixfmt.h:94
int ff_opencl_filter_config_output(AVFilterLink *outlink)
Create a suitable hardware frames context for the output.
Definition: opencl.c:96
#define av_log(a,...)
A filter pad used for either input or output.
Definition: internal.h:54
#define src
Definition: vp8dsp.c:254
#define i(width, name, range_min, range_max)
Definition: cbs_h2645.c:269
#define AV_LOG_ERROR
Something went wrong and cannot losslessly be recovered.
Definition: log.h:176
#define AVERROR(e)
Definition: error.h:43
void av_frame_free(AVFrame **frame)
Free the frame and any dynamically allocated objects in it, e.g.
Definition: frame.c:203
#define NULL_IF_CONFIG_SMALL(x)
Return NULL if CONFIG_SMALL is true, otherwise the argument without modification. ...
Definition: internal.h:188
packed BGRA 8:8:8:8, 32bpp, BGRABGRA...
Definition: pixfmt.h:95
void * priv
private data for use by the filter
Definition: avfilter.h:353
packed ARGB 8:8:8:8, 32bpp, ARGBARGB...
Definition: pixfmt.h:92
#define fail()
Definition: checkasm.h:123
packed RGBA 8:8:8:8, 32bpp, RGBARGBA...
Definition: pixfmt.h:93
AVFilter ff_vf_colorkey_opencl
AVFormatContext * ctx
Definition: movenc.c:48
static int filter_frame(AVFilterLink *link, AVFrame *input_frame)
static const AVFilterPad inputs[]
Definition: af_acontrast.c:193
static const AVFilterPad outputs[]
Definition: af_acontrast.c:203
#define CL_SET_KERNEL_ARG(kernel, arg_num, type, arg)
set argument to specific Kernel.
Definition: opencl.h:61
uint8_t * data
The data buffer.
Definition: buffer.h:89
static int output_frame(H264Context *h, AVFrame *dst, H264Picture *srcp)
Definition: h264dec.c:853
Filter definition.
Definition: avfilter.h:144
This struct describes a set or pool of "hardware" frames (i.e.
Definition: hwcontext.h:124
const char * name
Filter name.
Definition: avfilter.h:148
AVFilterLink ** outputs
array of pointers to output links
Definition: avfilter.h:350
AVFILTER_DEFINE_CLASS(colorkey_opencl)
static const AVOption colorkey_opencl_options[]
#define CL_FAIL_ON_ERROR(errcode,...)
A helper macro to handle OpenCL errors.
Definition: opencl.h:74
uint8_t * data[AV_NUM_DATA_POINTERS]
pointer to the picture/channel planes.
Definition: frame.h:314
static int query_formats(AVFilterContext *ctx)
Definition: aeval.c:244
if(ret< 0)
Definition: vf_mcdeint.c:279
#define FLAGS
void ff_opencl_filter_uninit(AVFilterContext *avctx)
Uninitialise an OpenCL filter context.
Definition: opencl.c:156
cl_context context
The OpenCL context which will contain all operations and frames on this device.
An instance of a filter.
Definition: avfilter.h:338
static int colorkey_opencl_init(AVFilterContext *avctx)
#define OFFSET(x)
cl_program program
Definition: opencl.h:43
int ff_opencl_filter_load_program(AVFilterContext *avctx, const char **program_source_array, int nb_strings)
Load a new OpenCL program from strings in memory.
Definition: opencl.c:171
internal API functions
OpenCLFilterContext ocf
enum AVPixelFormat sw_format
The pixel format identifying the actual data layout of the hardware frames.
Definition: hwcontext.h:222
int av_frame_copy_props(AVFrame *dst, const AVFrame *src)
Copy only "metadata" fields from src to dst.
Definition: frame.c:659