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 "framesync.h"
27 #include "internal.h"
28 #include "opencl.h"
29 #include "opencl_source.h"
30 #include "video.h"
31 
32 typedef struct RemapOpenCLContext {
33  OpenCLFilterContext ocf;
34 
35  int nb_planes;
36  int interp;
37  uint8_t fill_rgba[4];
38  cl_float4 cl_fill_color;
39 
40  int initialised;
41  cl_kernel kernel;
42  cl_command_queue command_queue;
43 
44  FFFrameSync fs;
45 } RemapOpenCLContext;
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 
60 static av_cold int remap_opencl_init(AVFilterContext *avctx)
61 {
62  return ff_opencl_filter_init(avctx);
63 }
64 
65 static const char *kernels[] = { "remap_near", "remap_linear" };
66 
67 static int remap_opencl_load(AVFilterContext *avctx,
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 
130 static int remap_opencl_process_frame(FFFrameSync *fs)
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:
225  av_frame_free(&output);
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  FFFrameSyncIn *in;
237  int ret;
238 
239  if (xlink->w != ylink->w || xlink->h != ylink->h) {
240  av_log(ctx, AV_LOG_ERROR, "Second input link %s parameters "
241  "(size %dx%d) do not match the corresponding "
242  "third input link %s parameters (%dx%d)\n",
243  ctx->input_pads[1].name, xlink->w, xlink->h,
244  ctx->input_pads[2].name, ylink->w, ylink->h);
245  return AVERROR(EINVAL);
246  }
247 
248  outlink->w = xlink->w;
249  outlink->h = xlink->h;
250  outlink->sample_aspect_ratio = srclink->sample_aspect_ratio;
251  outlink->frame_rate = srclink->frame_rate;
252 
253  ret = ff_framesync_init(&s->fs, ctx, 3);
254  if (ret < 0)
255  return ret;
256 
257  in = s->fs.in;
258  in[0].time_base = srclink->time_base;
259  in[1].time_base = xlink->time_base;
260  in[2].time_base = ylink->time_base;
261  in[0].sync = 2;
262  in[0].before = EXT_STOP;
263  in[0].after = EXT_STOP;
264  in[1].sync = 1;
265  in[1].before = EXT_NULL;
266  in[1].after = EXT_INFINITY;
267  in[2].sync = 1;
268  in[2].before = EXT_NULL;
269  in[2].after = EXT_INFINITY;
270  s->fs.opaque = s;
271  s->fs.on_event = remap_opencl_process_frame;
272 
273  ret = ff_framesync_configure(&s->fs);
274  outlink->time_base = s->fs.time_base;
275  if (ret < 0)
276  return ret;
277 
278  s->ocf.output_width = outlink->w;
279  s->ocf.output_height = outlink->h;
280 
281  return ff_opencl_filter_config_output(outlink);
282 }
283 
284 static int activate(AVFilterContext *ctx)
285 {
286  RemapOpenCLContext *s = ctx->priv;
287  return ff_framesync_activate(&s->fs);
288 }
289 
290 static av_cold void remap_opencl_uninit(AVFilterContext *avctx)
291 {
292  RemapOpenCLContext *ctx = avctx->priv;
293  cl_int cle;
294 
295  if (ctx->kernel) {
296  cle = clReleaseKernel(ctx->kernel);
297  if (cle != CL_SUCCESS)
298  av_log(avctx, AV_LOG_ERROR, "Failed to release "
299  "kernel: %d.\n", cle);
300  }
301 
302  if (ctx->command_queue) {
303  cle = clReleaseCommandQueue(ctx->command_queue);
304  if (cle != CL_SUCCESS)
305  av_log(avctx, AV_LOG_ERROR, "Failed to release "
306  "command queue: %d.\n", cle);
307  }
308 
309  ff_opencl_filter_uninit(avctx);
310 
311  ff_framesync_uninit(&ctx->fs);
312 }
313 
314 static const AVFilterPad remap_opencl_inputs[] = {
315  {
316  .name = "source",
317  .type = AVMEDIA_TYPE_VIDEO,
318  .config_props = &ff_opencl_filter_config_input,
319  },
320  {
321  .name = "xmap",
322  .type = AVMEDIA_TYPE_VIDEO,
323  .config_props = &ff_opencl_filter_config_input,
324  },
325  {
326  .name = "ymap",
327  .type = AVMEDIA_TYPE_VIDEO,
328  .config_props = &ff_opencl_filter_config_input,
329  },
330 };
331 
332 static const AVFilterPad remap_opencl_outputs[] = {
333  {
334  .name = "default",
335  .type = AVMEDIA_TYPE_VIDEO,
336  .config_props = config_output,
337  },
338 };
339 
340 const AVFilter ff_vf_remap_opencl = {
341  .name = "remap_opencl",
342  .description = NULL_IF_CONFIG_SMALL("Remap pixels using OpenCL."),
343  .priv_size = sizeof(RemapOpenCLContext),
344  .init = remap_opencl_init,
345  .uninit = remap_opencl_uninit,
346  .activate = activate,
347  FILTER_INPUTS(remap_opencl_inputs),
348  FILTER_OUTPUTS(remap_opencl_outputs),
349  FILTER_SINGLE_PIXFMT(AV_PIX_FMT_OPENCL),
350  .priv_class = &remap_opencl_class,
351  .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE,
352  .flags = AVFILTER_FLAG_HWDEVICE,
353 };
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:112
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: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: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:351
ff_filter_frame
int ff_filter_frame(AVFilterLink *link, AVFrame *frame)
Send a frame of data to the next filter.
Definition: avfilter.c:1018
av_pix_fmt_desc_get
const AVPixFmtDescriptor * av_pix_fmt_desc_get(enum AVPixelFormat pix_fmt)
Definition: pixdesc.c:2962
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:130
RemapOpenCLContext::cl_fill_color
cl_float4 cl_fill_color
Definition: vf_remap_opencl.c:38
AVFrame
This structure describes decoded (raw) audio or video data.
Definition: frame.h:344
pixdesc.h
opencl.h
AVOption
AVOption.
Definition: opt.h:346
RemapOpenCLContext::kernel
cl_kernel kernel
Definition: vf_remap_opencl.c:41
activate
static int activate(AVFilterContext *ctx)
Definition: vf_remap_opencl.c:284
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:156
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:170
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:365
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:263
FLAGS
#define FLAGS
Definition: vf_remap_opencl.c:48
AVFilterContext::priv
void * priv
private data for use by the filter
Definition: avfilter.h:422
RemapOpenCLContext::interp
int interp
Definition: vf_remap_opencl.c:36
interp
interp
Definition: vf_curves.c:61
fail
#define fail()
Definition: checkasm.h:179
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:81
AVFilterPad
A filter pad used for either input or output.
Definition: internal.h:33
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: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:290
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
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:65
ctx
AVFormatContext * ctx
Definition: movenc.c:48
FILTER_INPUTS
#define FILTER_INPUTS(array)
Definition: internal.h:182
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:679
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
Definition: opt.h:250
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:340
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:365
scale
static void scale(int *out, const int *in, const int w, const int h, const int shift)
Definition: vvc_intra.c:291
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:106
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:138
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:45
internal.h
FILTER_SINGLE_PIXFMT
#define FILTER_SINGLE_PIXFMT(pix_fmt_)
Definition: internal.h:172
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:255
RemapOpenCLContext::ocf
OpenCLFilterContext ocf
Definition: vf_remap_opencl.c:33
AVFilterPad::name
const char * name
Pad name.
Definition: internal.h:39
AVFilter
Filter definition.
Definition: avfilter.h:166
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:132
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:695
remap_opencl_inputs
static const AVFilterPad remap_opencl_inputs[]
Definition: vf_remap_opencl.c:314
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:332
AV_OPT_TYPE_INT
@ AV_OPT_TYPE_INT
Definition: opt.h:235
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:141
AVFilterContext
An instance of a filter.
Definition: avfilter.h:407
desc
const char * desc
Definition: libsvtav1.c:75
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:183
ff_fill_rgba_map
int ff_fill_rgba_map(uint8_t *rgba_map, enum AVPixelFormat pix_fmt)
Definition: drawutils.c:35
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:355
drawutils.h
RemapOpenCLContext::command_queue
cl_command_queue command_queue
Definition: vf_remap_opencl.c:42
AV_OPT_TYPE_CONST
@ AV_OPT_TYPE_CONST
Definition: opt.h:244
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:2882
AVFilterContext::outputs
AVFilterLink ** outputs
array of pointers to output links
Definition: avfilter.h:419
Generated on Thu Sep 26 2024 23:15:39 for FFmpeg by   doxygen 1.8.17