FFmpeg
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 &&
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" }, CHAR_MIN, CHAR_MAX, 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:385
This structure describes decoded (raw) audio or video data.
Definition: frame.h:268
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
const char * fmt
Definition: avisynth_c.h:861
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:99
static const AVFilterPad colorkey_opencl_inputs[]
#define src
Definition: vp8dsp.c:254
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:606
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:1080
uint8_t
#define av_cold
Definition: attributes.h:82
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 i(width, name, range_min, range_max)
Definition: cbs_h2645.c:260
const char * ff_opencl_source_colorkey
#define AV_LOG_ERROR
Something went wrong and cannot losslessly be recovered.
Definition: log.h:176
void av_frame_free(AVFrame **frame)
Free the frame and any dynamically allocated objects in it, e.g.
Definition: frame.c:202
#define NULL_IF_CONFIG_SMALL(x)
Return NULL if CONFIG_SMALL is true, otherwise the argument without modification. ...
Definition: internal.h:186
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:120
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 outputs[]
Definition: af_acontrast.c:203
if(ret)
#define CL_SET_KERNEL_ARG(kernel, arg_num, type, arg)
set argument to specific Kernel.
Definition: opencl.h:56
uint8_t * data
The data buffer.
Definition: buffer.h:89
these buffered frames must be flushed immediately if a new input produces new the filter must not call request_frame to get more It must just process the frame or queue it The task of requesting more frames is left to the filter s request_frame method or the application If a filter has several inputs
static int output_frame(H264Context *h, AVFrame *dst, H264Picture *srcp)
Definition: h264dec.c:832
Filter definition.
Definition: avfilter.h:144
This struct describes a set or pool of "hardware" frames (i.e.
Definition: hwcontext.h:123
const char * name
Filter name.
Definition: avfilter.h:148
Filter the word “frame” indicates either a video frame or a group of audio as stored in an AVFrame structure Format for each input and each output the list of supported formats For video that means pixel format For audio that means channel sample they are references to shared objects When the negotiation mechanism computes the intersection of the formats supported at each end of a link
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:69
uint8_t * data[AV_NUM_DATA_POINTERS]
pointer to the picture/channel planes.
Definition: frame.h:282
static int query_formats(AVFilterContext *ctx)
Definition: aeval.c:244
#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
Filter the word “frame” indicates either a video frame or a group of audio as stored in an AVFrame structure Format for each input and each output the list of supported formats For video that means pixel format For audio that means channel sample they are references to shared objects When the negotiation mechanism computes the intersection of the formats supported at each end of a all references to both lists are replaced with a reference to the intersection And when a single format is eventually chosen for a link amongst the remaining all references to the list are updated That means that if a filter requires that its input and output have the same format amongst a supported all it has to do is use a reference to the same list of formats query_formats can leave some formats unset and return AVERROR(EAGAIN) to cause the negotiation mechanism toagain later.That can be used by filters with complex requirements to use the format negotiated on one link to set the formats supported on another.Frame references ownership and permissions
OpenCLFilterContext ocf
enum AVPixelFormat sw_format
The pixel format identifying the actual data layout of the hardware frames.
Definition: hwcontext.h:221
int av_frame_copy_props(AVFrame *dst, const AVFrame *src)
Copy only "metadata" fields from src to dst.
Definition: frame.c:654