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/imgutils.h"
23 #include "libavutil/pixdesc.h"
24 #include "libavutil/opt.h"
25 #include "avfilter.h"
26 #include "drawutils.h"
27 #include "formats.h"
28 #include "framesync.h"
29 #include "internal.h"
30 #include "opencl.h"
31 #include "opencl_source.h"
32 #include "video.h"
33 
34 typedef struct RemapOpenCLContext {
35  OpenCLFilterContext ocf;
36 
37  int nb_planes;
38  int interp;
39  uint8_t fill_rgba[4];
40  cl_float4 cl_fill_color;
41 
42  int initialised;
43  cl_kernel kernel;
44  cl_command_queue command_queue;
45 
46  FFFrameSync fs;
47 } RemapOpenCLContext;
48 
49 #define OFFSET(x) offsetof(RemapOpenCLContext, x)
50 #define FLAGS AV_OPT_FLAG_FILTERING_PARAM|AV_OPT_FLAG_VIDEO_PARAM
51 
52 static const AVOption remap_opencl_options[] = {
53  { "interp", "set interpolation method", OFFSET(interp), AV_OPT_TYPE_INT, {.i64=1}, 0, 1, FLAGS, "interp" },
54  { "near", NULL, 0, AV_OPT_TYPE_CONST, {.i64=0}, 0, 0, FLAGS, "interp" },
55  { "linear", NULL, 0, AV_OPT_TYPE_CONST, {.i64=1}, 0, 0, FLAGS, "interp" },
56  { "fill", "set the color of the unmapped pixels", OFFSET(fill_rgba), AV_OPT_TYPE_COLOR, {.str="black"}, .flags = FLAGS },
57  { NULL }
58 };
59 
60 AVFILTER_DEFINE_CLASS(remap_opencl);
61 
62 static av_cold int remap_opencl_init(AVFilterContext *avctx)
63 {
64  return ff_opencl_filter_init(avctx);
65 }
66 
67 static const char *kernels[] = { "remap_near", "remap_linear" };
68 
69 static int remap_opencl_load(AVFilterContext *avctx,
70  enum AVPixelFormat main_format,
71  enum AVPixelFormat xmap_format,
72  enum AVPixelFormat ymap_format)
73 {
74  RemapOpenCLContext *ctx = avctx->priv;
75  cl_int cle;
76  const char *source = ff_opencl_source_remap;
77  const char *kernel = kernels[ctx->interp];
78  const AVPixFmtDescriptor *main_desc;
79  int err, main_planes;
80  const AVPixFmtDescriptor *desc = av_pix_fmt_desc_get(main_format);
81  int is_rgb = !!(desc->flags & AV_PIX_FMT_FLAG_RGB);
82  const float scale = 1.f / 255.f;
83  uint8_t rgba_map[4];
84 
85  ff_fill_rgba_map(rgba_map, main_format);
86 
87  if (is_rgb) {
88  ctx->cl_fill_color.s[rgba_map[0]] = ctx->fill_rgba[0] * scale;
89  ctx->cl_fill_color.s[rgba_map[1]] = ctx->fill_rgba[1] * scale;
90  ctx->cl_fill_color.s[rgba_map[2]] = ctx->fill_rgba[2] * scale;
91  ctx->cl_fill_color.s[rgba_map[3]] = ctx->fill_rgba[3] * scale;
92  } else {
93  ctx->cl_fill_color.s[0] = RGB_TO_Y_BT709(ctx->fill_rgba[0], ctx->fill_rgba[1], ctx->fill_rgba[2]) * scale;
94  ctx->cl_fill_color.s[1] = RGB_TO_U_BT709(ctx->fill_rgba[0], ctx->fill_rgba[1], ctx->fill_rgba[2], 0) * scale;
95  ctx->cl_fill_color.s[2] = RGB_TO_V_BT709(ctx->fill_rgba[0], ctx->fill_rgba[1], ctx->fill_rgba[2], 0) * scale;
96  ctx->cl_fill_color.s[3] = ctx->fill_rgba[3] * scale;
97  }
98 
99  main_desc = av_pix_fmt_desc_get(main_format);
100 
101  main_planes = 0;
102  for (int i = 0; i < main_desc->nb_components; i++)
103  main_planes = FFMAX(main_planes,
104  main_desc->comp[i].plane + 1);
105 
106  ctx->nb_planes = main_planes;
107 
108  err = ff_opencl_filter_load_program(avctx, &source, 1);
109  if (err < 0)
110  goto fail;
111 
112  ctx->command_queue = clCreateCommandQueue(ctx->ocf.hwctx->context,
113  ctx->ocf.hwctx->device_id,
114  0, &cle);
115  CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create OpenCL "
116  "command queue %d.\n", cle);
117 
118  ctx->kernel = clCreateKernel(ctx->ocf.program, kernel, &cle);
119  CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create kernel %d.\n", cle);
120 
121  ctx->initialised = 1;
122  return 0;
123 
124 fail:
125  if (ctx->command_queue)
126  clReleaseCommandQueue(ctx->command_queue);
127  if (ctx->kernel)
128  clReleaseKernel(ctx->kernel);
129  return err;
130 }
131 
132 static int remap_opencl_process_frame(FFFrameSync *fs)
133 {
134  AVFilterContext *avctx = fs->parent;
135  AVFilterLink *outlink = avctx->outputs[0];
136  RemapOpenCLContext *ctx = avctx->priv;
137  AVFrame *input_main, *input_xmap, *input_ymap;
138  AVFrame *output;
139  cl_mem mem;
140  cl_int cle;
141  size_t global_work[2];
142  int kernel_arg = 0;
143  int err, plane;
144 
145  err = ff_framesync_get_frame(fs, 0, &input_main, 0);
146  if (err < 0)
147  return err;
148  err = ff_framesync_get_frame(fs, 1, &input_xmap, 0);
149  if (err < 0)
150  return err;
151  err = ff_framesync_get_frame(fs, 2, &input_ymap, 0);
152  if (err < 0)
153  return err;
154 
155  if (!ctx->initialised) {
156  AVHWFramesContext *main_fc =
157  (AVHWFramesContext*)input_main->hw_frames_ctx->data;
158  AVHWFramesContext *xmap_fc =
159  (AVHWFramesContext*)input_xmap->hw_frames_ctx->data;
160  AVHWFramesContext *ymap_fc =
161  (AVHWFramesContext*)input_ymap->hw_frames_ctx->data;
162 
163  err = remap_opencl_load(avctx, main_fc->sw_format,
164  xmap_fc->sw_format,
165  ymap_fc->sw_format);
166  if (err < 0)
167  return err;
168  }
169 
170  output = ff_get_video_buffer(outlink, outlink->w, outlink->h);
171  if (!output) {
172  err = AVERROR(ENOMEM);
173  goto fail;
174  }
175 
176  for (plane = 0; plane < ctx->nb_planes; plane++) {
177  cl_float4 cl_fill_color;
178  kernel_arg = 0;
179 
180  if (ctx->nb_planes == 1)
181  cl_fill_color = ctx->cl_fill_color;
182  else
183  cl_fill_color.s[0] = ctx->cl_fill_color.s[plane];
184 
185  mem = (cl_mem)output->data[plane];
186  CL_SET_KERNEL_ARG(ctx->kernel, kernel_arg, cl_mem, &mem);
187  kernel_arg++;
188 
189  mem = (cl_mem)input_main->data[plane];
190  CL_SET_KERNEL_ARG(ctx->kernel, kernel_arg, cl_mem, &mem);
191  kernel_arg++;
192 
193  mem = (cl_mem)input_xmap->data[0];
194  CL_SET_KERNEL_ARG(ctx->kernel, kernel_arg, cl_mem, &mem);
195  kernel_arg++;
196 
197  mem = (cl_mem)input_ymap->data[0];
198  CL_SET_KERNEL_ARG(ctx->kernel, kernel_arg, cl_mem, &mem);
199  kernel_arg++;
200 
201  CL_SET_KERNEL_ARG(ctx->kernel, kernel_arg, cl_float4, &cl_fill_color);
202  kernel_arg++;
203 
204  err = ff_opencl_filter_work_size_from_image(avctx, global_work,
205  output, plane, 0);
206  if (err < 0)
207  goto fail;
208 
209  cle = clEnqueueNDRangeKernel(ctx->command_queue, ctx->kernel, 2, NULL,
210  global_work, NULL, 0, NULL, NULL);
211  CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to enqueue remap kernel "
212  "for plane %d: %d.\n", plane, cle);
213  }
214 
215  cle = clFinish(ctx->command_queue);
216  CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to finish command queue: %d.\n", cle);
217 
218  err = av_frame_copy_props(output, input_main);
219 
220  av_log(avctx, AV_LOG_DEBUG, "Filter output: %s, %ux%u (%"PRId64").\n",
221  av_get_pix_fmt_name(output->format),
222  output->width, output->height, output->pts);
223 
224  return ff_filter_frame(outlink, output);
225 
226 fail:
227  av_frame_free(&output);
228  return err;
229 }
230 
231 static int config_output(AVFilterLink *outlink)
232 {
233  AVFilterContext *ctx = outlink->src;
234  RemapOpenCLContext *s = ctx->priv;
235  AVFilterLink *srclink = ctx->inputs[0];
236  AVFilterLink *xlink = ctx->inputs[1];
237  AVFilterLink *ylink = ctx->inputs[2];
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  outlink->frame_rate = srclink->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 
286 static int activate(AVFilterContext *ctx)
287 {
288  RemapOpenCLContext *s = ctx->priv;
289  return ff_framesync_activate(&s->fs);
290 }
291 
292 static av_cold void remap_opencl_uninit(AVFilterContext *avctx)
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 
311  ff_opencl_filter_uninit(avctx);
312 
313  ff_framesync_uninit(&ctx->fs);
314 }
315 
316 static const AVFilterPad remap_opencl_inputs[] = {
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 
334 static const AVFilterPad remap_opencl_outputs[] = {
335  {
336  .name = "default",
337  .type = AVMEDIA_TYPE_VIDEO,
338  .config_props = config_output,
339  },
340 };
341 
342 const AVFilter ff_vf_remap_opencl = {
343  .name = "remap_opencl",
344  .description = NULL_IF_CONFIG_SMALL("Remap pixels using OpenCL."),
345  .priv_size = sizeof(RemapOpenCLContext),
346  .init = remap_opencl_init,
347  .uninit = remap_opencl_uninit,
348  .activate = activate,
349  FILTER_INPUTS(remap_opencl_inputs),
350  FILTER_OUTPUTS(remap_opencl_outputs),
351  FILTER_SINGLE_PIXFMT(AV_PIX_FMT_OPENCL),
352  .priv_class = &remap_opencl_class,
353  .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE,
354 };
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:69
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:101
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:134
AVPixelFormat
AVPixelFormat
Pixel format.
Definition: pixfmt.h:64
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:304
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: internal.h:374
ff_filter_frame
int ff_filter_frame(AVFilterLink *link, AVFrame *frame)
Send a frame of data to the next filter.
Definition: avfilter.c:969
av_pix_fmt_desc_get
const AVPixFmtDescriptor * av_pix_fmt_desc_get(enum AVPixelFormat pix_fmt)
Definition: pixdesc.c:2888
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:267
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:99
RemapOpenCLContext::cl_fill_color
cl_float4 cl_fill_color
Definition: vf_remap_opencl.c:40
AVFrame
This structure describes decoded (raw) audio or video data.
Definition: frame.h:330
pixdesc.h
opencl.h
AVOption
AVOption.
Definition: opt.h:251
RemapOpenCLContext::kernel
cl_kernel kernel
Definition: vf_remap_opencl.c:43
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:157
FFMAX
#define FFMAX(a, b)
Definition: macros.h:47
config_output
static int config_output(AVFilterLink *outlink)
Definition: vf_remap_opencl.c:231
AVFilter::name
const char * name
Filter name.
Definition: avfilter.h:165
FFFrameSync
Frame sync structure.
Definition: framesync.h:168
EXT_INFINITY
@ EXT_INFINITY
Extend the frame to infinity.
Definition: framesync.h:75
video.h
AVFrame::data
uint8_t * data[AV_NUM_DATA_POINTERS]
pointer to the picture/channel planes.
Definition: frame.h:351
formats.h
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:264
FLAGS
#define FLAGS
Definition: vf_remap_opencl.c:50
AVFilterContext::priv
void * priv
private data for use by the filter
Definition: avfilter.h:407
RemapOpenCLContext::interp
int interp
Definition: vf_remap_opencl.c:38
interp
interp
Definition: vf_curves.c:61
fail
#define fail()
Definition: checkasm.h:134
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
scale
static av_always_inline float scale(float x, float s)
Definition: vf_v360.c:1389
RemapOpenCLContext
Definition: vf_remap_opencl.c:34
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:52
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:82
AVFilterPad
A filter pad used for either input or output.
Definition: internal.h:49
colorspace.h
AV_LOG_ERROR
#define AV_LOG_ERROR
Something went wrong and cannot losslessly be recovered.
Definition: log.h:180
OFFSET
#define OFFSET(x)
Definition: vf_remap_opencl.c:49
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:562
s
#define s(width, name)
Definition: cbs_vp9.c:256
RGB_TO_U_BT709
#define RGB_TO_U_BT709(r1, g1, b1, max)
Definition: vf_pseudocolor.c:566
init
int(* init)(AVBSFContext *ctx)
Definition: dts2pts_bsf.c:365
AV_LOG_DEBUG
#define AV_LOG_DEBUG
Stuff which is only useful for libav* developers.
Definition: log.h:201
kernels
static const char * kernels[]
Definition: vf_remap_opencl.c:67
ctx
AVFormatContext * ctx
Definition: movenc.c:48
FILTER_INPUTS
#define FILTER_INPUTS(array)
Definition: internal.h:194
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:222
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:594
fs
#define fs(width, name, subs,...)
Definition: cbs_vp9.c:258
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
Definition: opt.h:240
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:355
ff_vf_remap_opencl
const AVFilter ff_vf_remap_opencl
Definition: vf_remap_opencl.c:342
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:132
RGB_TO_V_BT709
#define RGB_TO_V_BT709(r1, g1, b1, max)
Definition: vf_pseudocolor.c:570
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:115
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
RemapOpenCLContext::nb_planes
int nb_planes
Definition: vf_remap_opencl.c:37
opencl_source.h
RemapOpenCLContext::initialised
int initialised
Definition: vf_remap_opencl.c:42
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
internal.h
FILTER_SINGLE_PIXFMT
#define FILTER_SINGLE_PIXFMT(pix_fmt_)
Definition: internal.h:184
RemapOpenCLContext::fill_rgba
uint8_t fill_rgba[4]
Definition: vf_remap_opencl.c:39
i
#define i(width, name, range_min, range_max)
Definition: cbs_h2645.c:269
RemapOpenCLContext::ocf
OpenCLFilterContext ocf
Definition: vf_remap_opencl.c:35
AVFilterPad::name
const char * name
Pad name.
Definition: internal.h:55
AVFilter
Filter definition.
Definition: avfilter.h:161
AVHWFramesContext
This struct describes a set or pool of "hardware" frames (i.e.
Definition: hwcontext.h:124
ff_opencl_filter_init
int ff_opencl_filter_init(AVFilterContext *avctx)
Initialise an OpenCL filter context.
Definition: opencl.c:133
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:46
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:678
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
Definition: opt.h:225
avfilter.h
remap_opencl_init
static av_cold int remap_opencl_init(AVFilterContext *avctx)
Definition: vf_remap_opencl.c:62
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:142
AVFilterContext
An instance of a filter.
Definition: avfilter.h:392
desc
const char * desc
Definition: libsvtav1.c:83
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
FILTER_OUTPUTS
#define FILTER_OUTPUTS(array)
Definition: internal.h:195
ff_fill_rgba_map
int ff_fill_rgba_map(uint8_t *rgba_map, enum AVPixelFormat pix_fmt)
Definition: drawutils.c:35
imgutils.h
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
uninit
static av_cold int uninit(AVCodecContext *avctx)
Definition: crystalhd.c:285
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:355
drawutils.h
RemapOpenCLContext::command_queue
cl_command_queue command_queue
Definition: vf_remap_opencl.c:44
AV_OPT_TYPE_CONST
@ AV_OPT_TYPE_CONST
Definition: opt.h:234
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:2808
ff_opencl_source_remap
const char * ff_opencl_source_remap
AVFilterContext::outputs
AVFilterLink ** outputs
array of pointers to output links
Definition: avfilter.h:404
Generated on Tue Feb 28 2023 21:33:58 for FFmpeg by   doxygen 1.8.17