FFmpeg
vf_remap_opencl.c
Go to the documentation of this file.
1 /*
2  * Copyright (c) 2022 Paul B Mahol
3  *
4  * This file is part of FFmpeg.
5  *
6  * FFmpeg is free software; you can redistribute it and/or
7  * modify it under the terms of the GNU Lesser General Public
8  * License as published by the Free Software Foundation; either
9  * version 2.1 of the License, or (at your option) any later version.
10  *
11  * FFmpeg is distributed in the hope that it will be useful,
12  * but WITHOUT ANY WARRANTY; without even the implied warranty of
13  * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU
14  * Lesser General Public License for more details.
15  *
16  * You should have received a copy of the GNU Lesser General Public
17  * License along with FFmpeg; if not, write to the Free Software
18  * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA
19  */
20 
21 #include "libavutil/colorspace.h"
22 #include "libavutil/pixdesc.h"
23 #include "libavutil/opt.h"
24 #include "avfilter.h"
25 #include "drawutils.h"
26 #include "filters.h"
27 #include "framesync.h"
28 #include "opencl.h"
29 #include "opencl_source.h"
30 #include "video.h"
31 
32 typedef struct RemapOpenCLContext {
34 
35  int nb_planes;
36  int interp;
37  uint8_t fill_rgba[4];
38  cl_float4 cl_fill_color;
39 
41  cl_kernel kernel;
42  cl_command_queue command_queue;
43 
46 
47 #define OFFSET(x) offsetof(RemapOpenCLContext, x)
48 #define FLAGS AV_OPT_FLAG_FILTERING_PARAM|AV_OPT_FLAG_VIDEO_PARAM
49 
50 static const AVOption remap_opencl_options[] = {
51  { "interp", "set interpolation method", OFFSET(interp), AV_OPT_TYPE_INT, {.i64=1}, 0, 1, FLAGS, .unit = "interp" },
52  { "near", NULL, 0, AV_OPT_TYPE_CONST, {.i64=0}, 0, 0, FLAGS, .unit = "interp" },
53  { "linear", NULL, 0, AV_OPT_TYPE_CONST, {.i64=1}, 0, 0, FLAGS, .unit = "interp" },
54  { "fill", "set the color of the unmapped pixels", OFFSET(fill_rgba), AV_OPT_TYPE_COLOR, {.str="black"}, .flags = FLAGS },
55  { NULL }
56 };
57 
58 AVFILTER_DEFINE_CLASS(remap_opencl);
59 
61 {
62  return ff_opencl_filter_init(avctx);
63 }
64 
65 static const char *kernels[] = { "remap_near", "remap_linear" };
66 
68  enum AVPixelFormat main_format,
69  enum AVPixelFormat xmap_format,
70  enum AVPixelFormat ymap_format)
71 {
72  RemapOpenCLContext *ctx = avctx->priv;
73  cl_int cle;
74  const char *source = ff_source_remap_cl;
75  const char *kernel = kernels[ctx->interp];
76  const AVPixFmtDescriptor *main_desc;
77  int err, main_planes;
78  const AVPixFmtDescriptor *desc = av_pix_fmt_desc_get(main_format);
79  int is_rgb = !!(desc->flags & AV_PIX_FMT_FLAG_RGB);
80  const float scale = 1.f / 255.f;
81  uint8_t rgba_map[4];
82 
83  ff_fill_rgba_map(rgba_map, main_format);
84 
85  if (is_rgb) {
86  ctx->cl_fill_color.s[rgba_map[0]] = ctx->fill_rgba[0] * scale;
87  ctx->cl_fill_color.s[rgba_map[1]] = ctx->fill_rgba[1] * scale;
88  ctx->cl_fill_color.s[rgba_map[2]] = ctx->fill_rgba[2] * scale;
89  ctx->cl_fill_color.s[rgba_map[3]] = ctx->fill_rgba[3] * scale;
90  } else {
91  ctx->cl_fill_color.s[0] = RGB_TO_Y_BT709(ctx->fill_rgba[0], ctx->fill_rgba[1], ctx->fill_rgba[2]) * scale;
92  ctx->cl_fill_color.s[1] = RGB_TO_U_BT709(ctx->fill_rgba[0], ctx->fill_rgba[1], ctx->fill_rgba[2], 0) * scale;
93  ctx->cl_fill_color.s[2] = RGB_TO_V_BT709(ctx->fill_rgba[0], ctx->fill_rgba[1], ctx->fill_rgba[2], 0) * scale;
94  ctx->cl_fill_color.s[3] = ctx->fill_rgba[3] * scale;
95  }
96 
97  main_desc = av_pix_fmt_desc_get(main_format);
98 
99  main_planes = 0;
100  for (int i = 0; i < main_desc->nb_components; i++)
101  main_planes = FFMAX(main_planes,
102  main_desc->comp[i].plane + 1);
103 
104  ctx->nb_planes = main_planes;
105 
106  err = ff_opencl_filter_load_program(avctx, &source, 1);
107  if (err < 0)
108  goto fail;
109 
110  ctx->command_queue = clCreateCommandQueue(ctx->ocf.hwctx->context,
111  ctx->ocf.hwctx->device_id,
112  0, &cle);
113  CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create OpenCL "
114  "command queue %d.\n", cle);
115 
116  ctx->kernel = clCreateKernel(ctx->ocf.program, kernel, &cle);
117  CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create kernel %d.\n", cle);
118 
119  ctx->initialised = 1;
120  return 0;
121 
122 fail:
123  if (ctx->command_queue)
124  clReleaseCommandQueue(ctx->command_queue);
125  if (ctx->kernel)
126  clReleaseKernel(ctx->kernel);
127  return err;
128 }
129 
131 {
132  AVFilterContext *avctx = fs->parent;
133  AVFilterLink *outlink = avctx->outputs[0];
134  RemapOpenCLContext *ctx = avctx->priv;
135  AVFrame *input_main, *input_xmap, *input_ymap;
136  AVFrame *output;
137  cl_mem mem;
138  cl_int cle;
139  size_t global_work[2];
140  int kernel_arg = 0;
141  int err, plane;
142 
143  err = ff_framesync_get_frame(fs, 0, &input_main, 0);
144  if (err < 0)
145  return err;
146  err = ff_framesync_get_frame(fs, 1, &input_xmap, 0);
147  if (err < 0)
148  return err;
149  err = ff_framesync_get_frame(fs, 2, &input_ymap, 0);
150  if (err < 0)
151  return err;
152 
153  if (!ctx->initialised) {
154  AVHWFramesContext *main_fc =
155  (AVHWFramesContext*)input_main->hw_frames_ctx->data;
156  AVHWFramesContext *xmap_fc =
157  (AVHWFramesContext*)input_xmap->hw_frames_ctx->data;
158  AVHWFramesContext *ymap_fc =
159  (AVHWFramesContext*)input_ymap->hw_frames_ctx->data;
160 
161  err = remap_opencl_load(avctx, main_fc->sw_format,
162  xmap_fc->sw_format,
163  ymap_fc->sw_format);
164  if (err < 0)
165  return err;
166  }
167 
168  output = ff_get_video_buffer(outlink, outlink->w, outlink->h);
169  if (!output) {
170  err = AVERROR(ENOMEM);
171  goto fail;
172  }
173 
174  for (plane = 0; plane < ctx->nb_planes; plane++) {
175  cl_float4 cl_fill_color;
176  kernel_arg = 0;
177 
178  if (ctx->nb_planes == 1)
179  cl_fill_color = ctx->cl_fill_color;
180  else
181  cl_fill_color.s[0] = ctx->cl_fill_color.s[plane];
182 
183  mem = (cl_mem)output->data[plane];
184  CL_SET_KERNEL_ARG(ctx->kernel, kernel_arg, cl_mem, &mem);
185  kernel_arg++;
186 
187  mem = (cl_mem)input_main->data[plane];
188  CL_SET_KERNEL_ARG(ctx->kernel, kernel_arg, cl_mem, &mem);
189  kernel_arg++;
190 
191  mem = (cl_mem)input_xmap->data[0];
192  CL_SET_KERNEL_ARG(ctx->kernel, kernel_arg, cl_mem, &mem);
193  kernel_arg++;
194 
195  mem = (cl_mem)input_ymap->data[0];
196  CL_SET_KERNEL_ARG(ctx->kernel, kernel_arg, cl_mem, &mem);
197  kernel_arg++;
198 
199  CL_SET_KERNEL_ARG(ctx->kernel, kernel_arg, cl_float4, &cl_fill_color);
200  kernel_arg++;
201 
202  err = ff_opencl_filter_work_size_from_image(avctx, global_work,
203  output, plane, 0);
204  if (err < 0)
205  goto fail;
206 
207  cle = clEnqueueNDRangeKernel(ctx->command_queue, ctx->kernel, 2, NULL,
208  global_work, NULL, 0, NULL, NULL);
209  CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to enqueue remap kernel "
210  "for plane %d: %d.\n", plane, cle);
211  }
212 
213  cle = clFinish(ctx->command_queue);
214  CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to finish command queue: %d.\n", cle);
215 
216  err = av_frame_copy_props(output, input_main);
217 
218  av_log(avctx, AV_LOG_DEBUG, "Filter output: %s, %ux%u (%"PRId64").\n",
219  av_get_pix_fmt_name(output->format),
220  output->width, output->height, output->pts);
221 
222  return ff_filter_frame(outlink, output);
223 
224 fail:
226  return err;
227 }
228 
229 static int config_output(AVFilterLink *outlink)
230 {
231  AVFilterContext *ctx = outlink->src;
232  RemapOpenCLContext *s = ctx->priv;
233  AVFilterLink *srclink = ctx->inputs[0];
234  AVFilterLink *xlink = ctx->inputs[1];
235  AVFilterLink *ylink = ctx->inputs[2];
236  FilterLink *il = ff_filter_link(srclink);
237  FilterLink *ol = ff_filter_link(outlink);
238  FFFrameSyncIn *in;
239  int ret;
240 
241  if (xlink->w != ylink->w || xlink->h != ylink->h) {
242  av_log(ctx, AV_LOG_ERROR, "Second input link %s parameters "
243  "(size %dx%d) do not match the corresponding "
244  "third input link %s parameters (%dx%d)\n",
245  ctx->input_pads[1].name, xlink->w, xlink->h,
246  ctx->input_pads[2].name, ylink->w, ylink->h);
247  return AVERROR(EINVAL);
248  }
249 
250  outlink->w = xlink->w;
251  outlink->h = xlink->h;
252  outlink->sample_aspect_ratio = srclink->sample_aspect_ratio;
253  ol->frame_rate = il->frame_rate;
254 
255  ret = ff_framesync_init(&s->fs, ctx, 3);
256  if (ret < 0)
257  return ret;
258 
259  in = s->fs.in;
260  in[0].time_base = srclink->time_base;
261  in[1].time_base = xlink->time_base;
262  in[2].time_base = ylink->time_base;
263  in[0].sync = 2;
264  in[0].before = EXT_STOP;
265  in[0].after = EXT_STOP;
266  in[1].sync = 1;
267  in[1].before = EXT_NULL;
268  in[1].after = EXT_INFINITY;
269  in[2].sync = 1;
270  in[2].before = EXT_NULL;
271  in[2].after = EXT_INFINITY;
272  s->fs.opaque = s;
273  s->fs.on_event = remap_opencl_process_frame;
274 
275  ret = ff_framesync_configure(&s->fs);
276  outlink->time_base = s->fs.time_base;
277  if (ret < 0)
278  return ret;
279 
280  s->ocf.output_width = outlink->w;
281  s->ocf.output_height = outlink->h;
282 
283  return ff_opencl_filter_config_output(outlink);
284 }
285 
287 {
288  RemapOpenCLContext *s = ctx->priv;
289  return ff_framesync_activate(&s->fs);
290 }
291 
293 {
294  RemapOpenCLContext *ctx = avctx->priv;
295  cl_int cle;
296 
297  if (ctx->kernel) {
298  cle = clReleaseKernel(ctx->kernel);
299  if (cle != CL_SUCCESS)
300  av_log(avctx, AV_LOG_ERROR, "Failed to release "
301  "kernel: %d.\n", cle);
302  }
303 
304  if (ctx->command_queue) {
305  cle = clReleaseCommandQueue(ctx->command_queue);
306  if (cle != CL_SUCCESS)
307  av_log(avctx, AV_LOG_ERROR, "Failed to release "
308  "command queue: %d.\n", cle);
309  }
310 
312 
313  ff_framesync_uninit(&ctx->fs);
314 }
315 
317  {
318  .name = "source",
319  .type = AVMEDIA_TYPE_VIDEO,
320  .config_props = &ff_opencl_filter_config_input,
321  },
322  {
323  .name = "xmap",
324  .type = AVMEDIA_TYPE_VIDEO,
325  .config_props = &ff_opencl_filter_config_input,
326  },
327  {
328  .name = "ymap",
329  .type = AVMEDIA_TYPE_VIDEO,
330  .config_props = &ff_opencl_filter_config_input,
331  },
332 };
333 
335  {
336  .name = "default",
337  .type = AVMEDIA_TYPE_VIDEO,
338  .config_props = config_output,
339  },
340 };
341 
343  .name = "remap_opencl",
344  .description = NULL_IF_CONFIG_SMALL("Remap pixels using OpenCL."),
345  .priv_size = sizeof(RemapOpenCLContext),
348  .activate = activate,
352  .priv_class = &remap_opencl_class,
353  .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE,
354  .flags = AVFILTER_FLAG_HWDEVICE,
355 };
remap_opencl_load
static int remap_opencl_load(AVFilterContext *avctx, enum AVPixelFormat main_format, enum AVPixelFormat xmap_format, enum AVPixelFormat ymap_format)
Definition: vf_remap_opencl.c:67
ff_get_video_buffer
AVFrame * ff_get_video_buffer(AVFilterLink *link, int w, int h)
Request a picture buffer with a specific set of permissions.
Definition: video.c:116
FFFrameSyncIn::time_base
AVRational time_base
Time base for the incoming frames.
Definition: framesync.h:117
ff_framesync_configure
int ff_framesync_configure(FFFrameSync *fs)
Configure a frame sync structure.
Definition: framesync.c:137
AVPixelFormat
AVPixelFormat
Pixel format.
Definition: pixfmt.h:71
AVERROR
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
opt.h
CL_SET_KERNEL_ARG
#define CL_SET_KERNEL_ARG(kernel, arg_num, type, arg)
set argument to specific Kernel.
Definition: opencl.h:61
ff_framesync_uninit
void ff_framesync_uninit(FFFrameSync *fs)
Free all memory currently allocated.
Definition: framesync.c:301
ff_filter_frame
int ff_filter_frame(AVFilterLink *link, AVFrame *frame)
Send a frame of data to the next filter.
Definition: avfilter.c:1061
av_pix_fmt_desc_get
const AVPixFmtDescriptor * av_pix_fmt_desc_get(enum AVPixelFormat pix_fmt)
Definition: pixdesc.c:3170
AVBufferRef::data
uint8_t * data
The data buffer.
Definition: buffer.h:90
ff_framesync_get_frame
int ff_framesync_get_frame(FFFrameSync *fs, unsigned in, AVFrame **rframe, unsigned get)
Get the current frame in an input.
Definition: framesync.c:269
output
filter_frame For filters that do not use the this method is called when a frame is pushed to the filter s input It can be called at any time except in a reentrant way If the input frame is enough to produce output
Definition: filter_design.txt:225
av_frame_free
void av_frame_free(AVFrame **frame)
Free the frame and any dynamically allocated objects in it, e.g.
Definition: frame.c:162
RemapOpenCLContext::cl_fill_color
cl_float4 cl_fill_color
Definition: vf_remap_opencl.c:38
FILTER_INPUTS
#define FILTER_INPUTS(array)
Definition: filters.h:262
AVFrame
This structure describes decoded (raw) audio or video data.
Definition: frame.h:389
pixdesc.h
opencl.h
AVOption
AVOption.
Definition: opt.h:429
RemapOpenCLContext::kernel
cl_kernel kernel
Definition: vf_remap_opencl.c:41
activate
static int activate(AVFilterContext *ctx)
Definition: vf_remap_opencl.c:286
ff_opencl_filter_load_program
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:159
FFMAX
#define FFMAX(a, b)
Definition: macros.h:47
config_output
static int config_output(AVFilterLink *outlink)
Definition: vf_remap_opencl.c:229
AVFilter::name
const char * name
Filter name.
Definition: avfilter.h:205
FFFrameSync
Frame sync structure.
Definition: framesync.h:168
EXT_INFINITY
@ EXT_INFINITY
Extend the frame to infinity.
Definition: framesync.h:75
video.h
ff_source_remap_cl
const char * ff_source_remap_cl
AVFrame::data
uint8_t * data[AV_NUM_DATA_POINTERS]
pointer to the picture/channel planes.
Definition: frame.h:410
EXT_STOP
@ EXT_STOP
Completely stop all streams with this one.
Definition: framesync.h:65
ff_opencl_filter_work_size_from_image
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:266
FLAGS
#define FLAGS
Definition: vf_remap_opencl.c:48
AVFilterContext::priv
void * priv
private data for use by the filter
Definition: avfilter.h:472
RemapOpenCLContext::interp
int interp
Definition: vf_remap_opencl.c:36
interp
interp
Definition: vf_curves.c:62
fail
#define fail()
Definition: checkasm.h:193
FFFrameSyncIn
Input stream structure.
Definition: framesync.h:102
EXT_NULL
@ EXT_NULL
Ignore this stream and continue processing the other ones.
Definition: framesync.h:70
RemapOpenCLContext
Definition: vf_remap_opencl.c:32
FFFrameSyncIn::sync
unsigned sync
Synchronization level: frames on input at the highest sync level will generate output frame events.
Definition: framesync.h:160
remap_opencl_options
static const AVOption remap_opencl_options[]
Definition: vf_remap_opencl.c:50
ff_opencl_filter_config_output
int ff_opencl_filter_config_output(AVFilterLink *outlink)
Create a suitable hardware frames context for the output.
Definition: opencl.c:83
AVFilterPad
A filter pad used for either input or output.
Definition: filters.h:38
colorspace.h
AV_LOG_ERROR
#define AV_LOG_ERROR
Something went wrong and cannot losslessly be recovered.
Definition: log.h:209
OFFSET
#define OFFSET(x)
Definition: vf_remap_opencl.c:47
av_cold
#define av_cold
Definition: attributes.h:90
remap_opencl_uninit
static av_cold void remap_opencl_uninit(AVFilterContext *avctx)
Definition: vf_remap_opencl.c:292
RGB_TO_Y_BT709
#define RGB_TO_Y_BT709(r, g, b)
Definition: vf_pseudocolor.c:674
s
#define s(width, name)
Definition: cbs_vp9.c:198
RGB_TO_U_BT709
#define RGB_TO_U_BT709(r1, g1, b1, max)
Definition: vf_pseudocolor.c:678
filters.h
AV_LOG_DEBUG
#define AV_LOG_DEBUG
Stuff which is only useful for libav* developers.
Definition: log.h:230
kernels
static const char * kernels[]
Definition: vf_remap_opencl.c:65
ctx
AVFormatContext * ctx
Definition: movenc.c:49
FILTER_OUTPUTS
#define FILTER_OUTPUTS(array)
Definition: filters.h:263
NULL
#define NULL
Definition: coverity.c:32
AVHWFramesContext::sw_format
enum AVPixelFormat sw_format
The pixel format identifying the actual data layout of the hardware frames.
Definition: hwcontext.h:210
av_frame_copy_props
int av_frame_copy_props(AVFrame *dst, const AVFrame *src)
Copy only "metadata" fields from src to dst.
Definition: frame.c:725
fs
#define fs(width, name, subs,...)
Definition: cbs_vp9.c:200
AVPixFmtDescriptor::nb_components
uint8_t nb_components
The number of components each pixel has, (1-4)
Definition: pixdesc.h:71
AV_OPT_TYPE_COLOR
@ AV_OPT_TYPE_COLOR
Underlying C type is uint8_t[4].
Definition: opt.h:323
AVComponentDescriptor::plane
int plane
Which of the 4 planes contains the component.
Definition: pixdesc.h:34
AV_PIX_FMT_OPENCL
@ AV_PIX_FMT_OPENCL
Hardware surfaces for OpenCL.
Definition: pixfmt.h:358
ff_vf_remap_opencl
const AVFilter ff_vf_remap_opencl
Definition: vf_remap_opencl.c:342
ff_filter_link
static FilterLink * ff_filter_link(AVFilterLink *link)
Definition: filters.h:197
FF_FILTER_FLAG_HWFRAME_AWARE
#define FF_FILTER_FLAG_HWFRAME_AWARE
The filter is aware of hardware frames, and any hardware frame context should not be automatically pr...
Definition: filters.h:206
source
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 the filter must be ready for frames arriving randomly on any input any filter with several inputs will most likely require some kind of queuing mechanism It is perfectly acceptable to have a limited queue and to drop frames when the inputs are too unbalanced request_frame For filters that do not use the this method is called when a frame is wanted on an output For a source
Definition: filter_design.txt:255
remap_opencl_process_frame
static int remap_opencl_process_frame(FFFrameSync *fs)
Definition: vf_remap_opencl.c:130
RGB_TO_V_BT709
#define RGB_TO_V_BT709(r1, g1, b1, max)
Definition: vf_pseudocolor.c:682
init
int(* init)(AVBSFContext *ctx)
Definition: dts2pts.c:368
NULL_IF_CONFIG_SMALL
#define NULL_IF_CONFIG_SMALL(x)
Return NULL if CONFIG_SMALL is true, otherwise the argument without modification.
Definition: internal.h:94
AV_PIX_FMT_FLAG_RGB
#define AV_PIX_FMT_FLAG_RGB
The pixel format contains RGB-like data (as opposed to YUV/grayscale).
Definition: pixdesc.h:136
AVFILTER_FLAG_HWDEVICE
#define AVFILTER_FLAG_HWDEVICE
The filter can create hardware frames using AVFilterContext.hw_device_ctx.
Definition: avfilter.h:173
RemapOpenCLContext::nb_planes
int nb_planes
Definition: vf_remap_opencl.c:35
opencl_source.h
RemapOpenCLContext::initialised
int initialised
Definition: vf_remap_opencl.c:40
ff_opencl_filter_config_input
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:46
uninit
static void uninit(AVBSFContext *ctx)
Definition: pcm_rechunk.c:68
RemapOpenCLContext::fill_rgba
uint8_t fill_rgba[4]
Definition: vf_remap_opencl.c:37
i
#define i(width, name, range_min, range_max)
Definition: cbs_h2645.c:256
RemapOpenCLContext::ocf
OpenCLFilterContext ocf
Definition: vf_remap_opencl.c:33
AVFilterPad::name
const char * name
Pad name.
Definition: filters.h:44
AVFilter
Filter definition.
Definition: avfilter.h:201
AVHWFramesContext
This struct describes a set or pool of "hardware" frames (i.e.
Definition: hwcontext.h:115
ff_opencl_filter_init
int ff_opencl_filter_init(AVFilterContext *avctx)
Initialise an OpenCL filter context.
Definition: opencl.c:135
ret
ret
Definition: filter_design.txt:187
ff_framesync_init
int ff_framesync_init(FFFrameSync *fs, AVFilterContext *parent, unsigned nb_in)
Initialize a frame sync structure.
Definition: framesync.c:86
RemapOpenCLContext::fs
FFFrameSync fs
Definition: vf_remap_opencl.c:44
AVFrame::hw_frames_ctx
AVBufferRef * hw_frames_ctx
For hwaccel-format frames, this should be a reference to the AVHWFramesContext describing the frame.
Definition: frame.h:740
remap_opencl_inputs
static const AVFilterPad remap_opencl_inputs[]
Definition: vf_remap_opencl.c:316
FFFrameSyncIn::before
enum FFFrameSyncExtMode before
Extrapolation mode for timestamps before the first frame.
Definition: framesync.h:107
framesync.h
remap_opencl_outputs
static const AVFilterPad remap_opencl_outputs[]
Definition: vf_remap_opencl.c:334
AV_OPT_TYPE_INT
@ AV_OPT_TYPE_INT
Underlying C type is int.
Definition: opt.h:259
avfilter.h
remap_opencl_init
static av_cold int remap_opencl_init(AVFilterContext *avctx)
Definition: vf_remap_opencl.c:60
AVPixFmtDescriptor::comp
AVComponentDescriptor comp[4]
Parameters that describe how pixels are packed.
Definition: pixdesc.h:105
OpenCLFilterContext
Definition: opencl.h:36
ff_opencl_filter_uninit
void ff_opencl_filter_uninit(AVFilterContext *avctx)
Uninitialise an OpenCL filter context.
Definition: opencl.c:144
AVFilterContext
An instance of a filter.
Definition: avfilter.h:457
desc
const char * desc
Definition: libsvtav1.c:79
AVMEDIA_TYPE_VIDEO
@ AVMEDIA_TYPE_VIDEO
Definition: avutil.h:201
AVFILTER_DEFINE_CLASS
AVFILTER_DEFINE_CLASS(remap_opencl)
AVPixFmtDescriptor
Descriptor that unambiguously describes how the bits of a pixel are stored in the up to 4 data planes...
Definition: pixdesc.h:69
scale
static void scale(int *out, const int *in, const int w, const int h, const int shift)
Definition: intra.c:291
ff_fill_rgba_map
int ff_fill_rgba_map(uint8_t *rgba_map, enum AVPixelFormat pix_fmt)
Definition: drawutils.c:79
av_log
#define av_log(a,...)
Definition: tableprint_vlc.h:27
FFFrameSyncIn::after
enum FFFrameSyncExtMode after
Extrapolation mode for timestamps after the last frame.
Definition: framesync.h:112
CL_FAIL_ON_ERROR
#define CL_FAIL_ON_ERROR(errcode,...)
A helper macro to handle OpenCL errors.
Definition: opencl.h:74
ff_framesync_activate
int ff_framesync_activate(FFFrameSync *fs)
Examine the frames in the filter's input and try to produce output.
Definition: framesync.c:352
drawutils.h
RemapOpenCLContext::command_queue
cl_command_queue command_queue
Definition: vf_remap_opencl.c:42
AV_OPT_TYPE_CONST
@ AV_OPT_TYPE_CONST
Special option type for declaring named constants.
Definition: opt.h:299
FILTER_SINGLE_PIXFMT
#define FILTER_SINGLE_PIXFMT(pix_fmt_)
Definition: filters.h:252
av_get_pix_fmt_name
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:3090
AVFilterContext::outputs
AVFilterLink ** outputs
array of pointers to output links
Definition: avfilter.h:469