FFmpeg  4.0
vf_overlay_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/log.h"
20 #include "libavutil/mem.h"
21 #include "libavutil/opt.h"
22 #include "libavutil/pixdesc.h"
23 
24 #include "avfilter.h"
25 #include "framesync.h"
26 #include "internal.h"
27 #include "opencl.h"
28 #include "opencl_source.h"
29 #include "video.h"
30 
31 typedef struct OverlayOpenCLContext {
33 
35  cl_kernel kernel;
36  cl_command_queue command_queue;
37 
39 
40  int nb_planes;
44 
48 
50  enum AVPixelFormat main_format,
51  enum AVPixelFormat overlay_format)
52 {
53  OverlayOpenCLContext *ctx = avctx->priv;
54  cl_int cle;
55  const char *source = ff_opencl_source_overlay;
56  const char *kernel;
57  const AVPixFmtDescriptor *main_desc, *overlay_desc;
58  int err, i, main_planes, overlay_planes;
59 
60  main_desc = av_pix_fmt_desc_get(main_format);
61  overlay_desc = av_pix_fmt_desc_get(overlay_format);
62 
63  main_planes = overlay_planes = 0;
64  for (i = 0; i < main_desc->nb_components; i++)
65  main_planes = FFMAX(main_planes,
66  main_desc->comp[i].plane + 1);
67  for (i = 0; i < overlay_desc->nb_components; i++)
68  overlay_planes = FFMAX(overlay_planes,
69  overlay_desc->comp[i].plane + 1);
70 
71  ctx->nb_planes = main_planes;
72  ctx->x_subsample = 1 << main_desc->log2_chroma_w;
73  ctx->y_subsample = 1 << main_desc->log2_chroma_h;
74 
75  if (ctx->x_position % ctx->x_subsample ||
76  ctx->y_position % ctx->y_subsample) {
77  av_log(avctx, AV_LOG_WARNING, "Warning: overlay position (%d, %d) "
78  "does not match subsampling (%d, %d).\n",
79  ctx->x_position, ctx->y_position,
80  ctx->x_subsample, ctx->y_subsample);
81  }
82 
83  if (main_planes == overlay_planes) {
84  if (main_desc->nb_components == overlay_desc->nb_components)
85  kernel = "overlay_no_alpha";
86  else
87  kernel = "overlay_internal_alpha";
88  ctx->alpha_separate = 0;
89  } else {
90  kernel = "overlay_external_alpha";
91  ctx->alpha_separate = 1;
92  }
93 
94  av_log(avctx, AV_LOG_DEBUG, "Using kernel %s.\n", kernel);
95 
96  err = ff_opencl_filter_load_program(avctx, &source, 1);
97  if (err < 0)
98  goto fail;
99 
100  ctx->command_queue = clCreateCommandQueue(ctx->ocf.hwctx->context,
101  ctx->ocf.hwctx->device_id,
102  0, &cle);
103  if (!ctx->command_queue) {
104  av_log(avctx, AV_LOG_ERROR, "Failed to create OpenCL "
105  "command queue: %d.\n", cle);
106  err = AVERROR(EIO);
107  goto fail;
108  }
109 
110  ctx->kernel = clCreateKernel(ctx->ocf.program, kernel, &cle);
111  if (!ctx->kernel) {
112  av_log(avctx, AV_LOG_ERROR, "Failed to create kernel: %d.\n", cle);
113  err = AVERROR(EIO);
114  goto fail;
115  }
116 
117  ctx->initialised = 1;
118  return 0;
119 
120 fail:
121  if (ctx->command_queue)
122  clReleaseCommandQueue(ctx->command_queue);
123  if (ctx->kernel)
124  clReleaseKernel(ctx->kernel);
125  return err;
126 }
127 
129 {
130  AVFilterContext *avctx = fs->parent;
131  AVFilterLink *outlink = avctx->outputs[0];
132  OverlayOpenCLContext *ctx = avctx->priv;
133  AVFrame *input_main, *input_overlay;
134  AVFrame *output;
135  cl_mem mem;
136  cl_int cle, x, y;
137  size_t global_work[2];
138  int kernel_arg = 0;
139  int err, plane;
140 
141  err = ff_framesync_get_frame(fs, 0, &input_main, 0);
142  if (err < 0)
143  return err;
144  err = ff_framesync_get_frame(fs, 1, &input_overlay, 0);
145  if (err < 0)
146  return err;
147 
148  if (!ctx->initialised) {
149  AVHWFramesContext *main_fc =
150  (AVHWFramesContext*)input_main->hw_frames_ctx->data;
151  AVHWFramesContext *overlay_fc =
152  (AVHWFramesContext*)input_overlay->hw_frames_ctx->data;
153 
154  err = overlay_opencl_load(avctx, main_fc->sw_format,
155  overlay_fc->sw_format);
156  if (err < 0)
157  return err;
158  }
159 
160  output = ff_get_video_buffer(outlink, outlink->w, outlink->h);
161  if (!output) {
162  err = AVERROR(ENOMEM);
163  goto fail;
164  }
165 
166  for (plane = 0; plane < ctx->nb_planes; plane++) {
167  kernel_arg = 0;
168 
169  mem = (cl_mem)output->data[plane];
170  cle = clSetKernelArg(ctx->kernel, kernel_arg++, sizeof(cl_mem), &mem);
171  if (cle != CL_SUCCESS)
172  goto fail_kernel_arg;
173 
174  mem = (cl_mem)input_main->data[plane];
175  cle = clSetKernelArg(ctx->kernel, kernel_arg++, sizeof(cl_mem), &mem);
176  if (cle != CL_SUCCESS)
177  goto fail_kernel_arg;
178 
179  mem = (cl_mem)input_overlay->data[plane];
180  cle = clSetKernelArg(ctx->kernel, kernel_arg++, sizeof(cl_mem), &mem);
181  if (cle != CL_SUCCESS)
182  goto fail_kernel_arg;
183 
184  if (ctx->alpha_separate) {
185  mem = (cl_mem)input_overlay->data[ctx->nb_planes];
186  cle = clSetKernelArg(ctx->kernel, kernel_arg++, sizeof(cl_mem), &mem);
187  if (cle != CL_SUCCESS)
188  goto fail_kernel_arg;
189  }
190 
191  x = ctx->x_position / (plane == 0 ? 1 : ctx->x_subsample);
192  y = ctx->y_position / (plane == 0 ? 1 : ctx->y_subsample);
193 
194  cle = clSetKernelArg(ctx->kernel, kernel_arg++, sizeof(cl_int), &x);
195  if (cle != CL_SUCCESS)
196  goto fail_kernel_arg;
197  cle = clSetKernelArg(ctx->kernel, kernel_arg++, sizeof(cl_int), &y);
198  if (cle != CL_SUCCESS)
199  goto fail_kernel_arg;
200 
201  if (ctx->alpha_separate) {
202  cl_int alpha_adj_x = plane == 0 ? 1 : ctx->x_subsample;
203  cl_int alpha_adj_y = plane == 0 ? 1 : ctx->y_subsample;
204 
205  cle = clSetKernelArg(ctx->kernel, kernel_arg++, sizeof(cl_int), &alpha_adj_x);
206  if (cle != CL_SUCCESS)
207  goto fail_kernel_arg;
208  cle = clSetKernelArg(ctx->kernel, kernel_arg++, sizeof(cl_int), &alpha_adj_y);
209  if (cle != CL_SUCCESS)
210  goto fail_kernel_arg;
211  }
212 
213  err = ff_opencl_filter_work_size_from_image(avctx, global_work,
214  output, plane, 0);
215  if (err < 0)
216  goto fail;
217 
218  cle = clEnqueueNDRangeKernel(ctx->command_queue, ctx->kernel, 2, NULL,
219  global_work, NULL, 0, NULL, NULL);
220  if (cle != CL_SUCCESS) {
221  av_log(avctx, AV_LOG_ERROR, "Failed to enqueue "
222  "overlay kernel for plane %d: %d.\n", cle, plane);
223  err = AVERROR(EIO);
224  goto fail;
225  }
226  }
227 
228  cle = clFinish(ctx->command_queue);
229  if (cle != CL_SUCCESS) {
230  av_log(avctx, AV_LOG_ERROR, "Failed to finish "
231  "command queue: %d.\n", cle);
232  err = AVERROR(EIO);
233  goto fail;
234  }
235 
236  err = av_frame_copy_props(output, input_main);
237 
238  av_log(avctx, AV_LOG_DEBUG, "Filter output: %s, %ux%u (%"PRId64").\n",
239  av_get_pix_fmt_name(output->format),
240  output->width, output->height, output->pts);
241 
242  return ff_filter_frame(outlink, output);
243 
244 fail_kernel_arg:
245  av_log(avctx, AV_LOG_ERROR, "Failed to set kernel arg %d: %d.\n",
246  kernel_arg, cle);
247  err = AVERROR(EIO);
248 fail:
249  av_frame_free(&output);
250  return err;
251 }
252 
254 {
255  AVFilterContext *avctx = outlink->src;
256  OverlayOpenCLContext *ctx = avctx->priv;
257  int err;
258 
259  err = ff_opencl_filter_config_output(outlink);
260  if (err < 0)
261  return err;
262 
263  err = ff_framesync_init_dualinput(&ctx->fs, avctx);
264  if (err < 0)
265  return err;
266 
267  return ff_framesync_configure(&ctx->fs);
268 }
269 
271 {
272  OverlayOpenCLContext *ctx = avctx->priv;
273 
275 
276  return ff_opencl_filter_init(avctx);
277 }
278 
280 {
281  OverlayOpenCLContext *ctx = avctx->priv;
282 
283  return ff_framesync_activate(&ctx->fs);
284 }
285 
287 {
288  OverlayOpenCLContext *ctx = avctx->priv;
289  cl_int cle;
290 
291  if (ctx->kernel) {
292  cle = clReleaseKernel(ctx->kernel);
293  if (cle != CL_SUCCESS)
294  av_log(avctx, AV_LOG_ERROR, "Failed to release "
295  "kernel: %d.\n", cle);
296  }
297 
298  if (ctx->command_queue) {
299  cle = clReleaseCommandQueue(ctx->command_queue);
300  if (cle != CL_SUCCESS)
301  av_log(avctx, AV_LOG_ERROR, "Failed to release "
302  "command queue: %d.\n", cle);
303  }
304 
306 
307  ff_framesync_uninit(&ctx->fs);
308 }
309 
310 #define OFFSET(x) offsetof(OverlayOpenCLContext, x)
311 #define FLAGS (AV_OPT_FLAG_FILTERING_PARAM | AV_OPT_FLAG_VIDEO_PARAM)
313  { "x", "Overlay x position",
314  OFFSET(x_position), AV_OPT_TYPE_INT, { .i64 = 0 }, 0, INT_MAX, .flags = FLAGS },
315  { "y", "Overlay y position",
316  OFFSET(y_position), AV_OPT_TYPE_INT, { .i64 = 0 }, 0, INT_MAX, .flags = FLAGS },
317  { NULL },
318 };
319 
320 AVFILTER_DEFINE_CLASS(overlay_opencl);
321 
323  {
324  .name = "main",
325  .type = AVMEDIA_TYPE_VIDEO,
326  .config_props = &ff_opencl_filter_config_input,
327  },
328  {
329  .name = "overlay",
330  .type = AVMEDIA_TYPE_VIDEO,
331  .config_props = &ff_opencl_filter_config_input,
332  },
333  { NULL }
334 };
335 
337  {
338  .name = "default",
339  .type = AVMEDIA_TYPE_VIDEO,
340  .config_props = &overlay_opencl_config_output,
341  },
342  { NULL }
343 };
344 
346  .name = "overlay_opencl",
347  .description = NULL_IF_CONFIG_SMALL("Overlay one video on top of another"),
348  .priv_size = sizeof(OverlayOpenCLContext),
349  .priv_class = &overlay_opencl_class,
354  .inputs = overlay_opencl_inputs,
355  .outputs = overlay_opencl_outputs,
356  .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE,
357 };
int plane
Definition: avisynth_c.h:422
int plane
Which of the 4 planes contains the component.
Definition: pixdesc.h:35
#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
static const AVFilterPad overlay_opencl_outputs[]
const AVPixFmtDescriptor * av_pix_fmt_desc_get(enum AVPixelFormat pix_fmt)
Definition: pixdesc.c:2363
This structure describes decoded (raw) audio or video data.
Definition: frame.h:218
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:281
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:63
int ff_opencl_filter_query_formats(AVFilterContext *avctx)
Return that all inputs and outputs support only AV_PIX_FMT_OPENCL.
Definition: opencl.c:31
#define AV_LOG_WARNING
Something somehow does not look correct.
Definition: log.h:182
Main libavfilter public API header.
Memory handling functions.
static av_cold int init(AVCodecContext *avctx)
Definition: avrndec.c:35
int ff_framesync_configure(FFFrameSync *fs)
Configure a frame sync structure.
Definition: framesync.c:117
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
int ff_opencl_filter_init(AVFilterContext *avctx)
Initialise an OpenCL filter context.
Definition: opencl.c:150
uint8_t log2_chroma_w
Amount to shift the luma width right to find the chroma width.
Definition: pixdesc.h:92
AVFILTER_DEFINE_CLASS(overlay_opencl)
AVOpenCLDeviceContext * hwctx
Definition: opencl.h:40
AVBufferRef * hw_frames_ctx
For hwaccel-format frames, this should be a reference to the AVHWFramesContext describing the frame...
Definition: frame.h:556
static int overlay_opencl_load(AVFilterContext *avctx, enum AVPixelFormat main_format, enum AVPixelFormat overlay_format)
const char * name
Pad name.
Definition: internal.h:60
AVFilterContext * parent
Parent filter context.
Definition: framesync.h:152
static int overlay_opencl_activate(AVFilterContext *avctx)
int ff_filter_frame(AVFilterLink *link, AVFrame *frame)
Send a frame of data to the next filter.
Definition: avfilter.c:1080
static int activate(AVFilterContext *ctx)
Definition: af_amix.c:421
AVFilter ff_vf_overlay_opencl
AVComponentDescriptor comp[4]
Parameters that describe how pixels are packed.
Definition: pixdesc.h:117
int mem
Definition: avisynth_c.h:821
#define av_cold
Definition: attributes.h:82
static av_cold int uninit(AVCodecContext *avctx)
Definition: crystalhd.c:279
cl_command_queue command_queue
AVOptions.
int ff_framesync_init_dualinput(FFFrameSync *fs, AVFilterContext *parent)
Initialize a frame sync structure for dualinput.
Definition: framesync.c:361
static int overlay_opencl_blend(FFFrameSync *fs)
cl_device_id device_id
The primary device ID of the device.
int64_t pts
Presentation timestamp in time_base units (time when frame should be shown to user).
Definition: frame.h:311
static av_cold int overlay_opencl_init(AVFilterContext *avctx)
int ff_opencl_filter_config_output(AVFilterLink *outlink)
Create a suitable hardware frames context for the output.
Definition: opencl.c:99
const char * ff_opencl_source_overlay
#define av_log(a,...)
A filter pad used for either input or output.
Definition: internal.h:54
int width
Definition: frame.h:276
#define AV_LOG_ERROR
Something went wrong and cannot losslessly be recovered.
Definition: log.h:176
uint8_t log2_chroma_h
Amount to shift the luma height right to find the chroma height.
Definition: pixdesc.h:101
void ff_framesync_uninit(FFFrameSync *fs)
Free all memory currently allocated.
Definition: framesync.c:293
Frame sync structure.
Definition: framesync.h:146
#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:202
#define NULL_IF_CONFIG_SMALL(x)
Return NULL if CONFIG_SMALL is true, otherwise the argument without modification. ...
Definition: internal.h:186
void * priv
private data for use by the filter
Definition: avfilter.h:353
#define AV_LOG_DEBUG
Stuff which is only useful for libav* developers.
Definition: log.h:197
static const AVFilterPad overlay_opencl_inputs[]
int ff_framesync_activate(FFFrameSync *fs)
Examine the frames in the filter&#39;s input and try to produce output.
Definition: framesync.c:344
int(* on_event)(struct FFFrameSync *fs)
Callback called when a frame event is ready.
Definition: framesync.h:172
#define FFMAX(a, b)
Definition: common.h:94
#define fail()
Definition: checkasm.h:116
uint8_t nb_components
The number of components each pixel has, (1-4)
Definition: pixdesc.h:83
AVFormatContext * ctx
Definition: movenc.c:48
#define OFFSET(x)
static const AVFilterPad inputs[]
Definition: af_acontrast.c:193
static const AVFilterPad outputs[]
Definition: af_acontrast.c:203
int format
format of the frame, -1 if unknown or unset Values correspond to enum AVPixelFormat for video frames...
Definition: frame.h:291
Descriptor that unambiguously describes how the bits of a pixel are stored in the up to 4 data planes...
Definition: pixdesc.h:81
uint8_t * data
The data buffer.
Definition: buffer.h:89
static int overlay_opencl_config_output(AVFilterLink *outlink)
Filter definition.
Definition: avfilter.h:144
static av_cold void overlay_opencl_uninit(AVFilterContext *avctx)
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
AVFilterLink ** outputs
array of pointers to output links
Definition: avfilter.h:350
uint8_t * data[AV_NUM_DATA_POINTERS]
pointer to the picture/channel planes.
Definition: frame.h:232
static const AVOption overlay_opencl_options[]
static int query_formats(AVFilterContext *ctx)
Definition: aeval.c:244
void ff_opencl_filter_uninit(AVFilterContext *avctx)
Uninitialise an OpenCL filter context.
Definition: opencl.c:159
#define FLAGS
OpenCLFilterContext ocf
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
int height
Definition: frame.h:276
cl_program program
Definition: opencl.h:42
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:174
const char * av_get_pix_fmt_name(enum AVPixelFormat pix_fmt)
Return the short name for a pixel format, NULL in case pix_fmt is unknown.
Definition: pixdesc.c:2279
internal API functions
int ff_framesync_get_frame(FFFrameSync *fs, unsigned in, AVFrame **rframe, unsigned get)
Get the current frame in an input.
Definition: framesync.c:256
enum AVPixelFormat sw_format
The pixel format identifying the actual data layout of the hardware frames.
Definition: hwcontext.h:221
AVPixelFormat
Pixel format.
Definition: pixfmt.h:60
int av_frame_copy_props(AVFrame *dst, const AVFrame *src)
Copy only "metadata" fields from src to dst.
Definition: frame.c:652