FFmpeg
vf_tonemap_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 #include <float.h>
19 
20 #include "libavutil/avassert.h"
21 #include "libavutil/common.h"
22 #include "libavutil/imgutils.h"
23 #include "libavutil/opt.h"
24 #include "libavutil/pixdesc.h"
25 
26 #include "avfilter.h"
27 #include "internal.h"
28 #include "opencl.h"
29 #include "opencl_source.h"
30 #include "video.h"
31 #include "colorspace.h"
32 
33 // TODO:
34 // - separate peak-detection from tone-mapping kernel to solve
35 // one-frame-delay issue.
36 // - more format support
37 
38 #define DETECTION_FRAMES 63
39 
40 enum TonemapAlgorithm {
41  TONEMAP_NONE,
42  TONEMAP_LINEAR,
43  TONEMAP_GAMMA,
44  TONEMAP_CLIP,
45  TONEMAP_REINHARD,
46  TONEMAP_HABLE,
47  TONEMAP_MOBIUS,
48  TONEMAP_MAX,
49 };
50 
51 typedef struct TonemapOpenCLContext {
52  OpenCLFilterContext ocf;
53 
54  enum AVColorSpace colorspace, colorspace_in, colorspace_out;
55  enum AVColorTransferCharacteristic trc, trc_in, trc_out;
56  enum AVColorPrimaries primaries, primaries_in, primaries_out;
57  enum AVColorRange range, range_in, range_out;
58  enum AVChromaLocation chroma_loc;
59 
60  enum TonemapAlgorithm tonemap;
61  enum AVPixelFormat format;
62  double peak;
63  double param;
64  double desat_param;
65  double target_peak;
66  double scene_threshold;
67  int initialised;
68  cl_kernel kernel;
69  cl_command_queue command_queue;
70  cl_mem util_mem;
71 } TonemapOpenCLContext;
72 
73 static const char *const linearize_funcs[AVCOL_TRC_NB] = {
74  [AVCOL_TRC_SMPTE2084] = "eotf_st2084",
75  [AVCOL_TRC_ARIB_STD_B67] = "inverse_oetf_hlg",
76 };
77 
78 static const char *const delinearize_funcs[AVCOL_TRC_NB] = {
79  [AVCOL_TRC_BT709] = "inverse_eotf_bt1886",
80  [AVCOL_TRC_BT2020_10] = "inverse_eotf_bt1886",
81 };
82 
83 static const char *const tonemap_func[TONEMAP_MAX] = {
84  [TONEMAP_NONE] = "direct",
85  [TONEMAP_LINEAR] = "linear",
86  [TONEMAP_GAMMA] = "gamma",
87  [TONEMAP_CLIP] = "clip",
88  [TONEMAP_REINHARD] = "reinhard",
89  [TONEMAP_HABLE] = "hable",
90  [TONEMAP_MOBIUS] = "mobius",
91 };
92 
93 static int get_rgb2rgb_matrix(enum AVColorPrimaries in, enum AVColorPrimaries out,
94  double rgb2rgb[3][3]) {
95  double rgb2xyz[3][3], xyz2rgb[3][3];
96 
97  const AVColorPrimariesDesc *in_primaries = av_csp_primaries_desc_from_id(in);
98  const AVColorPrimariesDesc *out_primaries = av_csp_primaries_desc_from_id(out);
99 
100  if (!in_primaries || !out_primaries)
101  return AVERROR(EINVAL);
102 
103  ff_fill_rgb2xyz_table(&out_primaries->prim, &out_primaries->wp, rgb2xyz);
104  ff_matrix_invert_3x3(rgb2xyz, xyz2rgb);
105  ff_fill_rgb2xyz_table(&in_primaries->prim, &in_primaries->wp, rgb2xyz);
106  ff_matrix_mul_3x3(rgb2rgb, rgb2xyz, xyz2rgb);
107 
108  return 0;
109 }
110 
111 #define OPENCL_SOURCE_NB 3
112 // Average light level for SDR signals. This is equal to a signal level of 0.5
113 // under a typical presentation gamma of about 2.0.
114 static const float sdr_avg = 0.25f;
115 
116 static int tonemap_opencl_init(AVFilterContext *avctx)
117 {
118  TonemapOpenCLContext *ctx = avctx->priv;
119  int rgb2rgb_passthrough = 1;
120  double rgb2rgb[3][3], rgb2yuv[3][3], yuv2rgb[3][3];
121  const AVLumaCoefficients *luma_src, *luma_dst;
122  cl_int cle;
123  int err;
124  AVBPrint header;
125  const char *opencl_sources[OPENCL_SOURCE_NB];
126 
127  av_bprint_init(&header, 1024, AV_BPRINT_SIZE_AUTOMATIC);
128 
129  switch(ctx->tonemap) {
130  case TONEMAP_GAMMA:
131  if (isnan(ctx->param))
132  ctx->param = 1.8f;
133  break;
134  case TONEMAP_REINHARD:
135  if (!isnan(ctx->param))
136  ctx->param = (1.0f - ctx->param) / ctx->param;
137  break;
138  case TONEMAP_MOBIUS:
139  if (isnan(ctx->param))
140  ctx->param = 0.3f;
141  break;
142  }
143 
144  if (isnan(ctx->param))
145  ctx->param = 1.0f;
146 
147  // SDR peak is 1.0f
148  ctx->target_peak = 1.0f;
149  av_log(ctx, AV_LOG_DEBUG, "tone mapping transfer from %s to %s\n",
150  av_color_transfer_name(ctx->trc_in),
151  av_color_transfer_name(ctx->trc_out));
152  av_log(ctx, AV_LOG_DEBUG, "mapping colorspace from %s to %s\n",
153  av_color_space_name(ctx->colorspace_in),
154  av_color_space_name(ctx->colorspace_out));
155  av_log(ctx, AV_LOG_DEBUG, "mapping primaries from %s to %s\n",
156  av_color_primaries_name(ctx->primaries_in),
157  av_color_primaries_name(ctx->primaries_out));
158  av_log(ctx, AV_LOG_DEBUG, "mapping range from %s to %s\n",
159  av_color_range_name(ctx->range_in),
160  av_color_range_name(ctx->range_out));
161  // checking valid value just because of limited implementaion
162  // please remove when more functionalities are implemented
163  av_assert0(ctx->trc_out == AVCOL_TRC_BT709 ||
164  ctx->trc_out == AVCOL_TRC_BT2020_10);
165  av_assert0(ctx->trc_in == AVCOL_TRC_SMPTE2084||
166  ctx->trc_in == AVCOL_TRC_ARIB_STD_B67);
167  av_assert0(ctx->colorspace_in == AVCOL_SPC_BT2020_NCL ||
168  ctx->colorspace_in == AVCOL_SPC_BT709);
169  av_assert0(ctx->primaries_in == AVCOL_PRI_BT2020 ||
170  ctx->primaries_in == AVCOL_PRI_BT709);
171 
172  av_bprintf(&header, "__constant const float tone_param = %.4ff;\n",
173  ctx->param);
174  av_bprintf(&header, "__constant const float desat_param = %.4ff;\n",
175  ctx->desat_param);
176  av_bprintf(&header, "__constant const float target_peak = %.4ff;\n",
177  ctx->target_peak);
178  av_bprintf(&header, "__constant const float sdr_avg = %.4ff;\n", sdr_avg);
179  av_bprintf(&header, "__constant const float scene_threshold = %.4ff;\n",
180  ctx->scene_threshold);
181  av_bprintf(&header, "#define TONE_FUNC %s\n", tonemap_func[ctx->tonemap]);
182  av_bprintf(&header, "#define DETECTION_FRAMES %d\n", DETECTION_FRAMES);
183 
184  if (ctx->primaries_out != ctx->primaries_in) {
185  if ((err = get_rgb2rgb_matrix(ctx->primaries_in, ctx->primaries_out, rgb2rgb)) < 0)
186  goto fail;
187  rgb2rgb_passthrough = 0;
188  }
189  if (ctx->range_in == AVCOL_RANGE_JPEG)
190  av_bprintf(&header, "#define FULL_RANGE_IN\n");
191 
192  if (ctx->range_out == AVCOL_RANGE_JPEG)
193  av_bprintf(&header, "#define FULL_RANGE_OUT\n");
194 
195  av_bprintf(&header, "#define chroma_loc %d\n", (int)ctx->chroma_loc);
196 
197  if (rgb2rgb_passthrough)
198  av_bprintf(&header, "#define RGB2RGB_PASSTHROUGH\n");
199  else
200  ff_opencl_print_const_matrix_3x3(&header, "rgb2rgb", rgb2rgb);
201 
202 
203  luma_src = av_csp_luma_coeffs_from_avcsp(ctx->colorspace_in);
204  if (!luma_src) {
205  err = AVERROR(EINVAL);
206  av_log(avctx, AV_LOG_ERROR, "unsupported input colorspace %d (%s)\n",
207  ctx->colorspace_in, av_color_space_name(ctx->colorspace_in));
208  goto fail;
209  }
210 
211  luma_dst = av_csp_luma_coeffs_from_avcsp(ctx->colorspace_out);
212  if (!luma_dst) {
213  err = AVERROR(EINVAL);
214  av_log(avctx, AV_LOG_ERROR, "unsupported output colorspace %d (%s)\n",
215  ctx->colorspace_out, av_color_space_name(ctx->colorspace_out));
216  goto fail;
217  }
218 
219  ff_fill_rgb2yuv_table(luma_dst, rgb2yuv);
220  ff_opencl_print_const_matrix_3x3(&header, "yuv_matrix", rgb2yuv);
221 
222  ff_fill_rgb2yuv_table(luma_src, rgb2yuv);
223  ff_matrix_invert_3x3(rgb2yuv, yuv2rgb);
224  ff_opencl_print_const_matrix_3x3(&header, "rgb_matrix", yuv2rgb);
225 
226  av_bprintf(&header, "constant float3 luma_src = {%.4ff, %.4ff, %.4ff};\n",
227  av_q2d(luma_src->cr), av_q2d(luma_src->cg), av_q2d(luma_src->cb));
228  av_bprintf(&header, "constant float3 luma_dst = {%.4ff, %.4ff, %.4ff};\n",
229  av_q2d(luma_dst->cr), av_q2d(luma_dst->cg), av_q2d(luma_dst->cb));
230 
231  av_bprintf(&header, "#define linearize %s\n", linearize_funcs[ctx->trc_in]);
232  av_bprintf(&header, "#define delinearize %s\n",
233  delinearize_funcs[ctx->trc_out]);
234 
235  if (ctx->trc_in == AVCOL_TRC_ARIB_STD_B67)
236  av_bprintf(&header, "#define ootf_impl ootf_hlg\n");
237 
238  if (ctx->trc_out == AVCOL_TRC_ARIB_STD_B67)
239  av_bprintf(&header, "#define inverse_ootf_impl inverse_ootf_hlg\n");
240 
241  av_log(avctx, AV_LOG_DEBUG, "Generated OpenCL header:\n%s\n", header.str);
242  opencl_sources[0] = header.str;
243  opencl_sources[1] = ff_source_tonemap_cl;
244  opencl_sources[2] = ff_source_colorspace_common_cl;
245  err = ff_opencl_filter_load_program(avctx, opencl_sources, OPENCL_SOURCE_NB);
246 
247  av_bprint_finalize(&header, NULL);
248  if (err < 0)
249  goto fail;
250 
251  ctx->command_queue = clCreateCommandQueue(ctx->ocf.hwctx->context,
252  ctx->ocf.hwctx->device_id,
253  0, &cle);
254  CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create OpenCL "
255  "command queue %d.\n", cle);
256 
257  ctx->kernel = clCreateKernel(ctx->ocf.program, "tonemap", &cle);
258  CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create kernel %d.\n", cle);
259 
260  ctx->util_mem =
261  clCreateBuffer(ctx->ocf.hwctx->context, 0,
262  (2 * DETECTION_FRAMES + 7) * sizeof(unsigned),
263  NULL, &cle);
264  CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create util buffer: %d.\n", cle);
265 
266  ctx->initialised = 1;
267  return 0;
268 
269 fail:
270  av_bprint_finalize(&header, NULL);
271  if (ctx->util_mem)
272  clReleaseMemObject(ctx->util_mem);
273  if (ctx->command_queue)
274  clReleaseCommandQueue(ctx->command_queue);
275  if (ctx->kernel)
276  clReleaseKernel(ctx->kernel);
277  return err;
278 }
279 
280 static int tonemap_opencl_config_output(AVFilterLink *outlink)
281 {
282  AVFilterContext *avctx = outlink->src;
283  TonemapOpenCLContext *s = avctx->priv;
284  int ret;
285  if (s->format == AV_PIX_FMT_NONE)
286  av_log(avctx, AV_LOG_WARNING, "format not set, use default format NV12\n");
287  else {
288  if (s->format != AV_PIX_FMT_P010 &&
289  s->format != AV_PIX_FMT_NV12) {
290  av_log(avctx, AV_LOG_ERROR, "unsupported output format,"
291  "only p010/nv12 supported now\n");
292  return AVERROR(EINVAL);
293  }
294  }
295 
296  s->ocf.output_format = s->format == AV_PIX_FMT_NONE ? AV_PIX_FMT_NV12 : s->format;
297  ret = ff_opencl_filter_config_output(outlink);
298  if (ret < 0)
299  return ret;
300 
301  return 0;
302 }
303 
304 static int launch_kernel(AVFilterContext *avctx, cl_kernel kernel,
305  AVFrame *output, AVFrame *input, float peak) {
306  TonemapOpenCLContext *ctx = avctx->priv;
307  int err = AVERROR(ENOSYS);
308  size_t global_work[2];
309  size_t local_work[2];
310  cl_int cle;
311 
312  CL_SET_KERNEL_ARG(kernel, 0, cl_mem, &output->data[0]);
313  CL_SET_KERNEL_ARG(kernel, 1, cl_mem, &input->data[0]);
314  CL_SET_KERNEL_ARG(kernel, 2, cl_mem, &output->data[1]);
315  CL_SET_KERNEL_ARG(kernel, 3, cl_mem, &input->data[1]);
316  CL_SET_KERNEL_ARG(kernel, 4, cl_mem, &ctx->util_mem);
317  CL_SET_KERNEL_ARG(kernel, 5, cl_float, &peak);
318 
319  local_work[0] = 16;
320  local_work[1] = 16;
321  // Note the work size based on uv plane, as we process a 2x2 quad in one workitem
322  err = ff_opencl_filter_work_size_from_image(avctx, global_work, output,
323  1, 16);
324  if (err < 0)
325  return err;
326 
327  cle = clEnqueueNDRangeKernel(ctx->command_queue, kernel, 2, NULL,
328  global_work, local_work,
329  0, NULL, NULL);
330  CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to enqueue kernel: %d.\n", cle);
331  return 0;
332 fail:
333  return err;
334 }
335 
336 static int tonemap_opencl_filter_frame(AVFilterLink *inlink, AVFrame *input)
337 {
338  AVFilterContext *avctx = inlink->dst;
339  AVFilterLink *outlink = avctx->outputs[0];
340  TonemapOpenCLContext *ctx = avctx->priv;
341  AVFrame *output = NULL;
342  cl_int cle;
343  int err;
344  double peak = ctx->peak;
345 
346  AVHWFramesContext *input_frames_ctx =
347  (AVHWFramesContext*)input->hw_frames_ctx->data;
348 
349  av_log(ctx, AV_LOG_DEBUG, "Filter input: %s, %ux%u (%"PRId64").\n",
350  av_get_pix_fmt_name(input->format),
351  input->width, input->height, input->pts);
352 
353  if (!input->hw_frames_ctx)
354  return AVERROR(EINVAL);
355 
356  output = ff_get_video_buffer(outlink, outlink->w, outlink->h);
357  if (!output) {
358  err = AVERROR(ENOMEM);
359  goto fail;
360  }
361 
362  err = av_frame_copy_props(output, input);
363  if (err < 0)
364  goto fail;
365 
366  if (!peak)
367  peak = ff_determine_signal_peak(input);
368 
369  if (ctx->trc != -1)
370  output->color_trc = ctx->trc;
371  if (ctx->primaries != -1)
372  output->color_primaries = ctx->primaries;
373  if (ctx->colorspace != -1)
374  output->colorspace = ctx->colorspace;
375  if (ctx->range != -1)
376  output->color_range = ctx->range;
377 
378  ctx->trc_in = input->color_trc;
379  ctx->trc_out = output->color_trc;
380  ctx->colorspace_in = input->colorspace;
381  ctx->colorspace_out = output->colorspace;
382  ctx->primaries_in = input->color_primaries;
383  ctx->primaries_out = output->color_primaries;
384  ctx->range_in = input->color_range;
385  ctx->range_out = output->color_range;
386  ctx->chroma_loc = output->chroma_location;
387 
388  if (!ctx->initialised) {
389  if (!(input->color_trc == AVCOL_TRC_SMPTE2084 ||
390  input->color_trc == AVCOL_TRC_ARIB_STD_B67)) {
391  av_log(ctx, AV_LOG_ERROR, "unsupported transfer function characteristic.\n");
392  err = AVERROR(ENOSYS);
393  goto fail;
394  }
395 
396  if (input_frames_ctx->sw_format != AV_PIX_FMT_P010) {
397  av_log(ctx, AV_LOG_ERROR, "unsupported format in tonemap_opencl.\n");
398  err = AVERROR(ENOSYS);
399  goto fail;
400  }
401 
402  err = tonemap_opencl_init(avctx);
403  if (err < 0)
404  goto fail;
405  }
406 
407  switch(input_frames_ctx->sw_format) {
408  case AV_PIX_FMT_P010:
409  err = launch_kernel(avctx, ctx->kernel, output, input, peak);
410  if (err < 0) goto fail;
411  break;
412  default:
413  err = AVERROR(ENOSYS);
414  goto fail;
415  }
416 
417  cle = clFinish(ctx->command_queue);
418  CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to finish command queue: %d.\n", cle);
419 
420  av_frame_free(&input);
421 
422  ff_update_hdr_metadata(output, ctx->target_peak);
423 
424  av_log(ctx, AV_LOG_DEBUG, "Tone-mapping output: %s, %ux%u (%"PRId64").\n",
425  av_get_pix_fmt_name(output->format),
426  output->width, output->height, output->pts);
427 #ifndef NDEBUG
428  {
429  uint32_t *ptr, *max_total_p, *avg_total_p, *frame_number_p;
430  float peak_detected, avg_detected;
431  unsigned map_size = (2 * DETECTION_FRAMES + 7) * sizeof(unsigned);
432  ptr = (void *)clEnqueueMapBuffer(ctx->command_queue, ctx->util_mem,
433  CL_TRUE, CL_MAP_READ, 0, map_size,
434  0, NULL, NULL, &cle);
435  // For the layout of the util buffer, refer tonemap.cl
436  if (ptr) {
437  max_total_p = ptr + 2 * (DETECTION_FRAMES + 1) + 1;
438  avg_total_p = max_total_p + 1;
439  frame_number_p = avg_total_p + 2;
440  peak_detected = (float)*max_total_p / (REFERENCE_WHITE * (*frame_number_p));
441  avg_detected = (float)*avg_total_p / (REFERENCE_WHITE * (*frame_number_p));
442  av_log(ctx, AV_LOG_DEBUG, "peak %f, avg %f will be used for next frame\n",
443  peak_detected, avg_detected);
444  clEnqueueUnmapMemObject(ctx->command_queue, ctx->util_mem, ptr, 0,
445  NULL, NULL);
446  }
447  }
448 #endif
449 
450  return ff_filter_frame(outlink, output);
451 
452 fail:
453  clFinish(ctx->command_queue);
454  av_frame_free(&input);
455  av_frame_free(&output);
456  return err;
457 }
458 
459 static av_cold void tonemap_opencl_uninit(AVFilterContext *avctx)
460 {
461  TonemapOpenCLContext *ctx = avctx->priv;
462  cl_int cle;
463 
464  if (ctx->util_mem)
465  clReleaseMemObject(ctx->util_mem);
466  if (ctx->kernel) {
467  cle = clReleaseKernel(ctx->kernel);
468  if (cle != CL_SUCCESS)
469  av_log(avctx, AV_LOG_ERROR, "Failed to release "
470  "kernel: %d.\n", cle);
471  }
472 
473  if (ctx->command_queue) {
474  cle = clReleaseCommandQueue(ctx->command_queue);
475  if (cle != CL_SUCCESS)
476  av_log(avctx, AV_LOG_ERROR, "Failed to release "
477  "command queue: %d.\n", cle);
478  }
479 
480  ff_opencl_filter_uninit(avctx);
481 }
482 
483 #define OFFSET(x) offsetof(TonemapOpenCLContext, x)
484 #define FLAGS (AV_OPT_FLAG_FILTERING_PARAM | AV_OPT_FLAG_VIDEO_PARAM)
485 static const AVOption tonemap_opencl_options[] = {
486  { "tonemap", "tonemap algorithm selection", OFFSET(tonemap), AV_OPT_TYPE_INT, {.i64 = TONEMAP_NONE}, TONEMAP_NONE, TONEMAP_MAX - 1, FLAGS, "tonemap" },
487  { "none", 0, 0, AV_OPT_TYPE_CONST, {.i64 = TONEMAP_NONE}, 0, 0, FLAGS, "tonemap" },
488  { "linear", 0, 0, AV_OPT_TYPE_CONST, {.i64 = TONEMAP_LINEAR}, 0, 0, FLAGS, "tonemap" },
489  { "gamma", 0, 0, AV_OPT_TYPE_CONST, {.i64 = TONEMAP_GAMMA}, 0, 0, FLAGS, "tonemap" },
490  { "clip", 0, 0, AV_OPT_TYPE_CONST, {.i64 = TONEMAP_CLIP}, 0, 0, FLAGS, "tonemap" },
491  { "reinhard", 0, 0, AV_OPT_TYPE_CONST, {.i64 = TONEMAP_REINHARD}, 0, 0, FLAGS, "tonemap" },
492  { "hable", 0, 0, AV_OPT_TYPE_CONST, {.i64 = TONEMAP_HABLE}, 0, 0, FLAGS, "tonemap" },
493  { "mobius", 0, 0, AV_OPT_TYPE_CONST, {.i64 = TONEMAP_MOBIUS}, 0, 0, FLAGS, "tonemap" },
494  { "transfer", "set transfer characteristic", OFFSET(trc), AV_OPT_TYPE_INT, {.i64 = AVCOL_TRC_BT709}, -1, INT_MAX, FLAGS, "transfer" },
495  { "t", "set transfer characteristic", OFFSET(trc), AV_OPT_TYPE_INT, {.i64 = AVCOL_TRC_BT709}, -1, INT_MAX, FLAGS, "transfer" },
496  { "bt709", 0, 0, AV_OPT_TYPE_CONST, {.i64 = AVCOL_TRC_BT709}, 0, 0, FLAGS, "transfer" },
497  { "bt2020", 0, 0, AV_OPT_TYPE_CONST, {.i64 = AVCOL_TRC_BT2020_10}, 0, 0, FLAGS, "transfer" },
498  { "matrix", "set colorspace matrix", OFFSET(colorspace), AV_OPT_TYPE_INT, {.i64 = -1}, -1, INT_MAX, FLAGS, "matrix" },
499  { "m", "set colorspace matrix", OFFSET(colorspace), AV_OPT_TYPE_INT, {.i64 = -1}, -1, INT_MAX, FLAGS, "matrix" },
500  { "bt709", 0, 0, AV_OPT_TYPE_CONST, {.i64 = AVCOL_SPC_BT709}, 0, 0, FLAGS, "matrix" },
501  { "bt2020", 0, 0, AV_OPT_TYPE_CONST, {.i64 = AVCOL_SPC_BT2020_NCL}, 0, 0, FLAGS, "matrix" },
502  { "primaries", "set color primaries", OFFSET(primaries), AV_OPT_TYPE_INT, {.i64 = -1}, -1, INT_MAX, FLAGS, "primaries" },
503  { "p", "set color primaries", OFFSET(primaries), AV_OPT_TYPE_INT, {.i64 = -1}, -1, INT_MAX, FLAGS, "primaries" },
504  { "bt709", 0, 0, AV_OPT_TYPE_CONST, {.i64 = AVCOL_PRI_BT709}, 0, 0, FLAGS, "primaries" },
505  { "bt2020", 0, 0, AV_OPT_TYPE_CONST, {.i64 = AVCOL_PRI_BT2020}, 0, 0, FLAGS, "primaries" },
506  { "range", "set color range", OFFSET(range), AV_OPT_TYPE_INT, {.i64 = -1}, -1, INT_MAX, FLAGS, "range" },
507  { "r", "set color range", OFFSET(range), AV_OPT_TYPE_INT, {.i64 = -1}, -1, INT_MAX, FLAGS, "range" },
508  { "tv", 0, 0, AV_OPT_TYPE_CONST, {.i64 = AVCOL_RANGE_MPEG}, 0, 0, FLAGS, "range" },
509  { "pc", 0, 0, AV_OPT_TYPE_CONST, {.i64 = AVCOL_RANGE_JPEG}, 0, 0, FLAGS, "range" },
510  { "limited", 0, 0, AV_OPT_TYPE_CONST, {.i64 = AVCOL_RANGE_MPEG}, 0, 0, FLAGS, "range" },
511  { "full", 0, 0, AV_OPT_TYPE_CONST, {.i64 = AVCOL_RANGE_JPEG}, 0, 0, FLAGS, "range" },
512  { "format", "output pixel format", OFFSET(format), AV_OPT_TYPE_PIXEL_FMT, {.i64 = AV_PIX_FMT_NONE}, AV_PIX_FMT_NONE, INT_MAX, FLAGS, "fmt" },
513  { "peak", "signal peak override", OFFSET(peak), AV_OPT_TYPE_DOUBLE, {.dbl = 0}, 0, DBL_MAX, FLAGS },
514  { "param", "tonemap parameter", OFFSET(param), AV_OPT_TYPE_DOUBLE, {.dbl = NAN}, DBL_MIN, DBL_MAX, FLAGS },
515  { "desat", "desaturation parameter", OFFSET(desat_param), AV_OPT_TYPE_DOUBLE, {.dbl = 0.5}, 0, DBL_MAX, FLAGS },
516  { "threshold", "scene detection threshold", OFFSET(scene_threshold), AV_OPT_TYPE_DOUBLE, {.dbl = 0.2}, 0, DBL_MAX, FLAGS },
517  { NULL }
518 };
519 
520 AVFILTER_DEFINE_CLASS(tonemap_opencl);
521 
522 static const AVFilterPad tonemap_opencl_inputs[] = {
523  {
524  .name = "default",
525  .type = AVMEDIA_TYPE_VIDEO,
526  .filter_frame = &tonemap_opencl_filter_frame,
527  .config_props = &ff_opencl_filter_config_input,
528  },
529 };
530 
531 static const AVFilterPad tonemap_opencl_outputs[] = {
532  {
533  .name = "default",
534  .type = AVMEDIA_TYPE_VIDEO,
535  .config_props = &tonemap_opencl_config_output,
536  },
537 };
538 
539 const AVFilter ff_vf_tonemap_opencl = {
540  .name = "tonemap_opencl",
541  .description = NULL_IF_CONFIG_SMALL("Perform HDR to SDR conversion with tonemapping."),
542  .priv_size = sizeof(TonemapOpenCLContext),
543  .priv_class = &tonemap_opencl_class,
544  .init = &ff_opencl_filter_init,
545  .uninit = &tonemap_opencl_uninit,
546  FILTER_INPUTS(tonemap_opencl_inputs),
547  FILTER_OUTPUTS(tonemap_opencl_outputs),
548  FILTER_SINGLE_PIXFMT(AV_PIX_FMT_OPENCL),
549  .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE,
550  .flags = AVFILTER_FLAG_HWDEVICE,
551 };
AVLumaCoefficients::cr
AVRational cr
Definition: csp.h:49
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:108
AV_LOG_WARNING
#define AV_LOG_WARNING
Something somehow does not look correct.
Definition: log.h:186
AVPixelFormat
AVPixelFormat
Pixel format.
Definition: pixfmt.h:64
TONEMAP_REINHARD
@ TONEMAP_REINHARD
Definition: vf_tonemap_opencl.c:45
TonemapOpenCLContext::format
enum AVPixelFormat format
Definition: vf_tonemap_opencl.c:61
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
tonemap_func
static const char *const tonemap_func[TONEMAP_MAX]
Definition: vf_tonemap_opencl.c:83
AVColorTransferCharacteristic
AVColorTransferCharacteristic
Color Transfer Characteristic.
Definition: pixfmt.h:566
CL_SET_KERNEL_ARG
#define CL_SET_KERNEL_ARG(kernel, arg_num, type, arg)
set argument to specific Kernel.
Definition: opencl.h:61
out
FILE * out
Definition: movenc.c:54
AVColorPrimariesDesc::wp
AVWhitepointCoefficients wp
Definition: csp.h:79
av_bprint_init
void av_bprint_init(AVBPrint *buf, unsigned size_init, unsigned size_max)
Definition: bprint.c:69
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:364
TONEMAP_MAX
@ TONEMAP_MAX
Definition: vf_tonemap_opencl.c:48
AVColorPrimariesDesc
Struct that contains both white point location and primaries location, providing the complete descrip...
Definition: csp.h:78
ff_filter_frame
int ff_filter_frame(AVFilterLink *link, AVFrame *frame)
Send a frame of data to the next filter.
Definition: avfilter.c:978
launch_kernel
static int launch_kernel(AVFilterContext *avctx, cl_kernel kernel, AVFrame *output, AVFrame *input, float peak)
Definition: vf_tonemap_opencl.c:304
ff_matrix_invert_3x3
void ff_matrix_invert_3x3(const double in[3][3], double out[3][3])
Definition: colorspace.c:27
tonemap_opencl_config_output
static int tonemap_opencl_config_output(AVFilterLink *outlink)
Definition: vf_tonemap_opencl.c:280
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
inlink
The exact code depends on how similar the blocks are and how related they are to the and needs to apply these operations to the correct inlink or outlink if there are several Macros are available to factor that when no extra processing is inlink
Definition: filter_design.txt:212
TonemapOpenCLContext::kernel
cl_kernel kernel
Definition: vf_tonemap_opencl.c:68
TONEMAP_HABLE
@ TONEMAP_HABLE
Definition: vf_tonemap_opencl.c:46
av_frame_free
void av_frame_free(AVFrame **frame)
Free the frame and any dynamically allocated objects in it, e.g.
Definition: frame.c:100
test::height
int height
Definition: vc1dsp.c:39
AVFrame
This structure describes decoded (raw) audio or video data.
Definition: frame.h:340
AVCOL_TRC_NB
@ AVCOL_TRC_NB
Not part of ABI.
Definition: pixfmt.h:588
pixdesc.h
TonemapOpenCLContext::target_peak
double target_peak
Definition: vf_tonemap_opencl.c:65
AVCOL_RANGE_JPEG
@ AVCOL_RANGE_JPEG
Full range content.
Definition: pixfmt.h:669
DETECTION_FRAMES
#define DETECTION_FRAMES
Definition: vf_tonemap_opencl.c:38
av_csp_luma_coeffs_from_avcsp
const struct AVLumaCoefficients * av_csp_luma_coeffs_from_avcsp(enum AVColorSpace csp)
Retrieves the Luma coefficients necessary to construct a conversion matrix from an enum constant desc...
Definition: csp.c:58
TonemapOpenCLContext::initialised
int initialised
Definition: vf_tonemap_opencl.c:67
opencl.h
AVOption
AVOption.
Definition: opt.h:251
rgb2yuv
static const char rgb2yuv[]
Definition: vf_scale_vulkan.c:70
ff_determine_signal_peak
double ff_determine_signal_peak(AVFrame *in)
Definition: colorspace.c:153
float.h
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
AVLumaCoefficients
Struct containing luma coefficients to be used for RGB to YUV/YCoCg, or similar calculations.
Definition: csp.h:48
AVColorPrimaries
AVColorPrimaries
Chromaticity coordinates of the source primaries.
Definition: pixfmt.h:541
AVFilter::name
const char * name
Filter name.
Definition: avfilter.h:170
video.h
TONEMAP_GAMMA
@ TONEMAP_GAMMA
Definition: vf_tonemap_opencl.c:43
TonemapOpenCLContext::colorspace_out
enum AVColorSpace colorspace colorspace_in colorspace_out
Definition: vf_tonemap_opencl.c:54
primaries
enum AVColorPrimaries primaries
Definition: mediacodec_wrapper.c:2664
delinearize_funcs
static const char *const delinearize_funcs[AVCOL_TRC_NB]
Definition: vf_tonemap_opencl.c:78
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
TonemapOpenCLContext::util_mem
cl_mem util_mem
Definition: vf_tonemap_opencl.c:70
av_color_space_name
const char * av_color_space_name(enum AVColorSpace space)
Definition: pixdesc.c:3312
TonemapOpenCLContext::peak
double peak
Definition: vf_tonemap_opencl.c:62
AVFilterContext::priv
void * priv
private data for use by the filter
Definition: avfilter.h:412
fail
#define fail()
Definition: checkasm.h:138
FLAGS
#define FLAGS
Definition: vf_tonemap_opencl.c:484
TonemapOpenCLContext::chroma_loc
enum AVChromaLocation chroma_loc
Definition: vf_tonemap_opencl.c:58
TonemapOpenCLContext
Definition: vf_tonemap_opencl.c:51
colorspace.h
AV_BPRINT_SIZE_AUTOMATIC
#define AV_BPRINT_SIZE_AUTOMATIC
tonemap
static void tonemap(TonemapContext *s, AVFrame *out, const AVFrame *in, const AVPixFmtDescriptor *desc, int x, int y, double peak)
Definition: vf_tonemap.c:108
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:47
TonemapOpenCLContext::primaries_out
enum AVColorPrimaries primaries primaries_in primaries_out
Definition: vf_tonemap_opencl.c:56
tonemap_opencl_filter_frame
static int tonemap_opencl_filter_frame(AVFilterLink *inlink, AVFrame *input)
Definition: vf_tonemap_opencl.c:336
get_rgb2rgb_matrix
static int get_rgb2rgb_matrix(enum AVColorPrimaries in, enum AVColorPrimaries out, double rgb2rgb[3][3])
Definition: vf_tonemap_opencl.c:93
TonemapOpenCLContext::ocf
OpenCLFilterContext ocf
Definition: vf_tonemap_opencl.c:52
avassert.h
AV_LOG_ERROR
#define AV_LOG_ERROR
Something went wrong and cannot losslessly be recovered.
Definition: log.h:180
av_cold
#define av_cold
Definition: attributes.h:90
float
float
Definition: af_crystalizer.c:121
s
#define s(width, name)
Definition: cbs_vp9.c:198
av_csp_primaries_desc_from_id
const AVColorPrimariesDesc * av_csp_primaries_desc_from_id(enum AVColorPrimaries prm)
Retrieves a complete gamut description from an enum constant describing the color primaries.
Definition: csp.c:90
format
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 format(the sample packing is implied by the sample format) and sample rate. The lists are not just lists
TonemapOpenCLContext::scene_threshold
double scene_threshold
Definition: vf_tonemap_opencl.c:66
AV_OPT_TYPE_DOUBLE
@ AV_OPT_TYPE_DOUBLE
Definition: opt.h:227
av_q2d
static double av_q2d(AVRational a)
Convert an AVRational to a double.
Definition: rational.h:104
init
int(* init)(AVBSFContext *ctx)
Definition: dts2pts_bsf.c:365
av_assert0
#define av_assert0(cond)
assert() equivalent, that is always enabled.
Definition: avassert.h:40
AV_LOG_DEBUG
#define AV_LOG_DEBUG
Stuff which is only useful for libav* developers.
Definition: log.h:201
ctx
AVFormatContext * ctx
Definition: movenc.c:48
AVLumaCoefficients::cg
AVRational cg
Definition: csp.h:49
NAN
#define NAN
Definition: mathematics.h:115
FILTER_INPUTS
#define FILTER_INPUTS(array)
Definition: internal.h:192
tonemap_opencl_outputs
static const AVFilterPad tonemap_opencl_outputs[]
Definition: vf_tonemap_opencl.c:531
if
if(ret)
Definition: filter_design.txt:179
tonemap_opencl_uninit
static av_cold void tonemap_opencl_uninit(AVFilterContext *avctx)
Definition: vf_tonemap_opencl.c:459
av_color_range_name
const char * av_color_range_name(enum AVColorRange range)
Definition: pixdesc.c:3252
ff_matrix_mul_3x3
void ff_matrix_mul_3x3(double dst[3][3], const double src1[3][3], const double src2[3][3])
Definition: colorspace.c:54
TONEMAP_NONE
@ TONEMAP_NONE
Definition: vf_tonemap_opencl.c:41
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:736
TonemapOpenCLContext::command_queue
cl_command_queue command_queue
Definition: vf_tonemap_opencl.c:69
isnan
#define isnan(x)
Definition: libm.h:340
AVCOL_PRI_BT709
@ AVCOL_PRI_BT709
also ITU-R BT1361 / IEC 61966-2-4 / SMPTE RP 177 Annex B
Definition: pixfmt.h:543
av_color_primaries_name
const char * av_color_primaries_name(enum AVColorPrimaries primaries)
Definition: pixdesc.c:3270
AVCOL_TRC_BT2020_10
@ AVCOL_TRC_BT2020_10
ITU-R BT2020 for 10-bit system.
Definition: pixfmt.h:581
AV_PIX_FMT_OPENCL
@ AV_PIX_FMT_OPENCL
Hardware surfaces for OpenCL.
Definition: pixfmt.h:355
TonemapOpenCLContext::range_out
enum AVColorRange range range_in range_out
Definition: vf_tonemap_opencl.c:57
AVCOL_PRI_BT2020
@ AVCOL_PRI_BT2020
ITU-R BT2020.
Definition: pixfmt.h:552
TonemapAlgorithm
TonemapAlgorithm
Definition: vf_tonemap.c:41
AVCOL_TRC_SMPTE2084
@ AVCOL_TRC_SMPTE2084
SMPTE ST 2084 for 10-, 12-, 14- and 16-bit systems.
Definition: pixfmt.h:583
test::width
int width
Definition: vc1dsp.c:38
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_bprint_finalize
int av_bprint_finalize(AVBPrint *buf, char **ret_str)
Finalize a print buffer.
Definition: bprint.c:240
tonemap_opencl_inputs
static const AVFilterPad tonemap_opencl_inputs[]
Definition: vf_tonemap_opencl.c:522
sdr_avg
static const float sdr_avg
Definition: vf_tonemap_opencl.c:114
AVFILTER_FLAG_HWDEVICE
#define AVFILTER_FLAG_HWDEVICE
The filter can create hardware frames using AVFilterContext.hw_device_ctx.
Definition: avfilter.h:138
ff_source_tonemap_cl
const char * ff_source_tonemap_cl
ff_update_hdr_metadata
void ff_update_hdr_metadata(AVFrame *in, double peak)
Definition: colorspace.c:178
xyz2rgb
static const float xyz2rgb[3][3]
Definition: tiff.c:1873
range
enum AVColorRange range
Definition: mediacodec_wrapper.c:2646
header
static const uint8_t header[24]
Definition: sdr2.c:67
REFERENCE_WHITE
#define REFERENCE_WHITE
Definition: colorspace.h:27
OPENCL_SOURCE_NB
#define OPENCL_SOURCE_NB
Definition: vf_tonemap_opencl.c:111
opencl_source.h
input
and forward the test the status of outputs and forward it to the corresponding return FFERROR_NOT_READY If the filters stores internally one or a few frame for some input
Definition: filter_design.txt:172
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
AVCOL_TRC_BT709
@ AVCOL_TRC_BT709
also ITU-R BT1361
Definition: pixfmt.h:568
internal.h
AVChromaLocation
AVChromaLocation
Location of chroma samples.
Definition: pixfmt.h:688
AVFILTER_DEFINE_CLASS
AVFILTER_DEFINE_CLASS(tonemap_opencl)
FILTER_SINGLE_PIXFMT
#define FILTER_SINGLE_PIXFMT(pix_fmt_)
Definition: internal.h:182
TonemapOpenCLContext::trc_out
enum AVColorTransferCharacteristic trc trc_in trc_out
Definition: vf_tonemap_opencl.c:55
ff_fill_rgb2yuv_table
void ff_fill_rgb2yuv_table(const AVLumaCoefficients *coeffs, double rgb2yuv[3][3])
Definition: colorspace.c:125
AVCOL_SPC_BT2020_NCL
@ AVCOL_SPC_BT2020_NCL
ITU-R BT2020 non-constant luminance system.
Definition: pixfmt.h:606
TonemapOpenCLContext::param
double param
Definition: vf_tonemap_opencl.c:63
AVColorSpace
AVColorSpace
YUV colorspace type.
Definition: pixfmt.h:595
common.h
TonemapOpenCLContext::desat_param
double desat_param
Definition: vf_tonemap_opencl.c:64
tonemap_opencl_init
static int tonemap_opencl_init(AVFilterContext *avctx)
Definition: vf_tonemap_opencl.c:116
ff_source_colorspace_common_cl
const char * ff_source_colorspace_common_cl
ff_opencl_print_const_matrix_3x3
void ff_opencl_print_const_matrix_3x3(AVBPrint *buf, const char *name_str, double mat[3][3])
Print a 3x3 matrix into a buffer as __constant array, which could be included in an OpenCL program.
Definition: opencl.c:326
AVFilterPad::name
const char * name
Pad name.
Definition: internal.h:53
TonemapOpenCLContext::tonemap
enum TonemapAlgorithm tonemap
Definition: vf_tonemap_opencl.c:60
AVCOL_RANGE_MPEG
@ AVCOL_RANGE_MPEG
Narrow or limited range content.
Definition: pixfmt.h:652
ff_vf_tonemap_opencl
const AVFilter ff_vf_tonemap_opencl
Definition: vf_tonemap_opencl.c:539
AVFilter
Filter definition.
Definition: avfilter.h:166
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:132
ret
ret
Definition: filter_design.txt:187
AV_PIX_FMT_NV12
@ AV_PIX_FMT_NV12
planar YUV 4:2:0, 12bpp, 1 plane for Y and 1 plane for the UV components, which are interleaved (firs...
Definition: pixfmt.h:89
ff_fill_rgb2xyz_table
void ff_fill_rgb2xyz_table(const AVPrimaryCoefficients *coeffs, const AVWhitepointCoefficients *wp, double rgb2xyz[3][3])
Definition: colorspace.c:79
av_bprintf
void av_bprintf(AVBPrint *buf, const char *fmt,...)
Definition: bprint.c:99
TONEMAP_MOBIUS
@ TONEMAP_MOBIUS
Definition: vf_tonemap_opencl.c:47
AVCOL_TRC_ARIB_STD_B67
@ AVCOL_TRC_ARIB_STD_B67
ARIB STD-B67, known as "Hybrid log-gamma".
Definition: pixfmt.h:587
linearize_funcs
static const char *const linearize_funcs[AVCOL_TRC_NB]
Definition: vf_tonemap_opencl.c:73
AV_PIX_FMT_NONE
@ AV_PIX_FMT_NONE
Definition: pixfmt.h:65
AV_OPT_TYPE_INT
@ AV_OPT_TYPE_INT
Definition: opt.h:225
avfilter.h
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
AV_OPT_TYPE_PIXEL_FMT
@ AV_OPT_TYPE_PIXEL_FMT
Definition: opt.h:236
AVFilterContext
An instance of a filter.
Definition: avfilter.h:397
tonemap_opencl_options
static const AVOption tonemap_opencl_options[]
Definition: vf_tonemap_opencl.c:485
AV_PIX_FMT_P010
#define AV_PIX_FMT_P010
Definition: pixfmt.h:514
AVMEDIA_TYPE_VIDEO
@ AVMEDIA_TYPE_VIDEO
Definition: avutil.h:201
TONEMAP_CLIP
@ TONEMAP_CLIP
Definition: vf_tonemap_opencl.c:44
FILTER_OUTPUTS
#define FILTER_OUTPUTS(array)
Definition: internal.h:193
AVColorPrimariesDesc::prim
AVPrimaryCoefficients prim
Definition: csp.h:80
imgutils.h
OFFSET
#define OFFSET(x)
Definition: vf_tonemap_opencl.c:483
av_log
#define av_log(a,...)
Definition: tableprint_vlc.h:27
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
AVLumaCoefficients::cb
AVRational cb
Definition: csp.h:49
AVCOL_SPC_BT709
@ AVCOL_SPC_BT709
also ITU-R BT1361 / IEC 61966-2-4 xvYCC709 / derived in SMPTE RP 177 Annex B
Definition: pixfmt.h:597
AVColorRange
AVColorRange
Visual content value range.
Definition: pixfmt.h:634
AV_OPT_TYPE_CONST
@ AV_OPT_TYPE_CONST
Definition: opt.h:234
TONEMAP_LINEAR
@ TONEMAP_LINEAR
Definition: vf_tonemap_opencl.c:42
yuv2rgb
static void yuv2rgb(uint8_t *out, int ridx, int Y, int U, int V)
Definition: g2meet.c:262
av_color_transfer_name
const char * av_color_transfer_name(enum AVColorTransferCharacteristic transfer)
Definition: pixdesc.c:3291
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:2856
AVFilterContext::outputs
AVFilterLink ** outputs
array of pointers to output links
Definition: avfilter.h:409
Generated on Sat Sep 23 2023 19:22:14 for FFmpeg by   doxygen 1.8.17