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 struct PrimaryCoefficients primaries_table[AVCOL_PRI_NB] = {
84  [AVCOL_PRI_BT709] = { 0.640, 0.330, 0.300, 0.600, 0.150, 0.060 },
85  [AVCOL_PRI_BT2020] = { 0.708, 0.292, 0.170, 0.797, 0.131, 0.046 },
86 };
87 
88 static const struct WhitepointCoefficients whitepoint_table[AVCOL_PRI_NB] = {
89  [AVCOL_PRI_BT709] = { 0.3127, 0.3290 },
90  [AVCOL_PRI_BT2020] = { 0.3127, 0.3290 },
91 };
92 
93 static const char *const tonemap_func[TONEMAP_MAX] = {
94  [TONEMAP_NONE] = "direct",
95  [TONEMAP_LINEAR] = "linear",
96  [TONEMAP_GAMMA] = "gamma",
97  [TONEMAP_CLIP] = "clip",
98  [TONEMAP_REINHARD] = "reinhard",
99  [TONEMAP_HABLE] = "hable",
100  [TONEMAP_MOBIUS] = "mobius",
101 };
102 
103 static void get_rgb2rgb_matrix(enum AVColorPrimaries in, enum AVColorPrimaries out,
104  double rgb2rgb[3][3]) {
105  double rgb2xyz[3][3], xyz2rgb[3][3];
106 
107  ff_fill_rgb2xyz_table(&primaries_table[out], &whitepoint_table[out], rgb2xyz);
108  ff_matrix_invert_3x3(rgb2xyz, xyz2rgb);
109  ff_fill_rgb2xyz_table(&primaries_table[in], &whitepoint_table[in], rgb2xyz);
110  ff_matrix_mul_3x3(rgb2rgb, rgb2xyz, xyz2rgb);
111 }
112 
113 #define OPENCL_SOURCE_NB 3
114 // Average light level for SDR signals. This is equal to a signal level of 0.5
115 // under a typical presentation gamma of about 2.0.
116 static const float sdr_avg = 0.25f;
117 
118 static int tonemap_opencl_init(AVFilterContext *avctx)
119 {
120  TonemapOpenCLContext *ctx = avctx->priv;
121  int rgb2rgb_passthrough = 1;
122  double rgb2rgb[3][3], rgb2yuv[3][3], yuv2rgb[3][3];
123  const struct LumaCoefficients *luma_src, *luma_dst;
124  cl_int cle;
125  int err;
126  AVBPrint header;
127  const char *opencl_sources[OPENCL_SOURCE_NB];
128 
129  av_bprint_init(&header, 1024, AV_BPRINT_SIZE_AUTOMATIC);
130 
131  switch(ctx->tonemap) {
132  case TONEMAP_GAMMA:
133  if (isnan(ctx->param))
134  ctx->param = 1.8f;
135  break;
136  case TONEMAP_REINHARD:
137  if (!isnan(ctx->param))
138  ctx->param = (1.0f - ctx->param) / ctx->param;
139  break;
140  case TONEMAP_MOBIUS:
141  if (isnan(ctx->param))
142  ctx->param = 0.3f;
143  break;
144  }
145 
146  if (isnan(ctx->param))
147  ctx->param = 1.0f;
148 
149  // SDR peak is 1.0f
150  ctx->target_peak = 1.0f;
151  av_log(ctx, AV_LOG_DEBUG, "tone mapping transfer from %s to %s\n",
152  av_color_transfer_name(ctx->trc_in),
153  av_color_transfer_name(ctx->trc_out));
154  av_log(ctx, AV_LOG_DEBUG, "mapping colorspace from %s to %s\n",
155  av_color_space_name(ctx->colorspace_in),
156  av_color_space_name(ctx->colorspace_out));
157  av_log(ctx, AV_LOG_DEBUG, "mapping primaries from %s to %s\n",
158  av_color_primaries_name(ctx->primaries_in),
159  av_color_primaries_name(ctx->primaries_out));
160  av_log(ctx, AV_LOG_DEBUG, "mapping range from %s to %s\n",
161  av_color_range_name(ctx->range_in),
162  av_color_range_name(ctx->range_out));
163  // checking valid value just because of limited implementaion
164  // please remove when more functionalities are implemented
165  av_assert0(ctx->trc_out == AVCOL_TRC_BT709 ||
166  ctx->trc_out == AVCOL_TRC_BT2020_10);
167  av_assert0(ctx->trc_in == AVCOL_TRC_SMPTE2084||
168  ctx->trc_in == AVCOL_TRC_ARIB_STD_B67);
169  av_assert0(ctx->colorspace_in == AVCOL_SPC_BT2020_NCL ||
170  ctx->colorspace_in == AVCOL_SPC_BT709);
171  av_assert0(ctx->primaries_in == AVCOL_PRI_BT2020 ||
172  ctx->primaries_in == AVCOL_PRI_BT709);
173 
174  av_bprintf(&header, "__constant const float tone_param = %.4ff;\n",
175  ctx->param);
176  av_bprintf(&header, "__constant const float desat_param = %.4ff;\n",
177  ctx->desat_param);
178  av_bprintf(&header, "__constant const float target_peak = %.4ff;\n",
179  ctx->target_peak);
180  av_bprintf(&header, "__constant const float sdr_avg = %.4ff;\n", sdr_avg);
181  av_bprintf(&header, "__constant const float scene_threshold = %.4ff;\n",
182  ctx->scene_threshold);
183  av_bprintf(&header, "#define TONE_FUNC %s\n", tonemap_func[ctx->tonemap]);
184  av_bprintf(&header, "#define DETECTION_FRAMES %d\n", DETECTION_FRAMES);
185 
186  if (ctx->primaries_out != ctx->primaries_in) {
187  get_rgb2rgb_matrix(ctx->primaries_in, ctx->primaries_out, rgb2rgb);
188  rgb2rgb_passthrough = 0;
189  }
190  if (ctx->range_in == AVCOL_RANGE_JPEG)
191  av_bprintf(&header, "#define FULL_RANGE_IN\n");
192 
193  if (ctx->range_out == AVCOL_RANGE_JPEG)
194  av_bprintf(&header, "#define FULL_RANGE_OUT\n");
195 
196  av_bprintf(&header, "#define chroma_loc %d\n", (int)ctx->chroma_loc);
197 
198  if (rgb2rgb_passthrough)
199  av_bprintf(&header, "#define RGB2RGB_PASSTHROUGH\n");
200  else
201  ff_opencl_print_const_matrix_3x3(&header, "rgb2rgb", rgb2rgb);
202 
203 
204  luma_src = ff_get_luma_coefficients(ctx->colorspace_in);
205  if (!luma_src) {
206  err = AVERROR(EINVAL);
207  av_log(avctx, AV_LOG_ERROR, "unsupported input colorspace %d (%s)\n",
208  ctx->colorspace_in, av_color_space_name(ctx->colorspace_in));
209  goto fail;
210  }
211 
212  luma_dst = ff_get_luma_coefficients(ctx->colorspace_out);
213  if (!luma_dst) {
214  err = AVERROR(EINVAL);
215  av_log(avctx, AV_LOG_ERROR, "unsupported output colorspace %d (%s)\n",
216  ctx->colorspace_out, av_color_space_name(ctx->colorspace_out));
217  goto fail;
218  }
219 
220  ff_fill_rgb2yuv_table(luma_dst, rgb2yuv);
221  ff_opencl_print_const_matrix_3x3(&header, "yuv_matrix", rgb2yuv);
222 
223  ff_fill_rgb2yuv_table(luma_src, rgb2yuv);
224  ff_matrix_invert_3x3(rgb2yuv, yuv2rgb);
225  ff_opencl_print_const_matrix_3x3(&header, "rgb_matrix", yuv2rgb);
226 
227  av_bprintf(&header, "constant float3 luma_src = {%.4ff, %.4ff, %.4ff};\n",
228  luma_src->cr, luma_src->cg, luma_src->cb);
229  av_bprintf(&header, "constant float3 luma_dst = {%.4ff, %.4ff, %.4ff};\n",
230  luma_dst->cr, luma_dst->cg, luma_dst->cb);
231 
232  av_bprintf(&header, "#define linearize %s\n", linearize_funcs[ctx->trc_in]);
233  av_bprintf(&header, "#define delinearize %s\n",
234  delinearize_funcs[ctx->trc_out]);
235 
236  if (ctx->trc_in == AVCOL_TRC_ARIB_STD_B67)
237  av_bprintf(&header, "#define ootf_impl ootf_hlg\n");
238 
239  if (ctx->trc_out == AVCOL_TRC_ARIB_STD_B67)
240  av_bprintf(&header, "#define inverse_ootf_impl inverse_ootf_hlg\n");
241 
242  av_log(avctx, AV_LOG_DEBUG, "Generated OpenCL header:\n%s\n", header.str);
243  opencl_sources[0] = header.str;
244  opencl_sources[1] = ff_opencl_source_tonemap;
245  opencl_sources[2] = ff_opencl_source_colorspace_common;
246  err = ff_opencl_filter_load_program(avctx, opencl_sources, OPENCL_SOURCE_NB);
247 
248  av_bprint_finalize(&header, NULL);
249  if (err < 0)
250  goto fail;
251 
252  ctx->command_queue = clCreateCommandQueue(ctx->ocf.hwctx->context,
253  ctx->ocf.hwctx->device_id,
254  0, &cle);
255  CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create OpenCL "
256  "command queue %d.\n", cle);
257 
258  ctx->kernel = clCreateKernel(ctx->ocf.program, "tonemap", &cle);
259  CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create kernel %d.\n", cle);
260 
261  ctx->util_mem =
262  clCreateBuffer(ctx->ocf.hwctx->context, 0,
263  (2 * DETECTION_FRAMES + 7) * sizeof(unsigned),
264  NULL, &cle);
265  CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create util buffer: %d.\n", cle);
266 
267  ctx->initialised = 1;
268  return 0;
269 
270 fail:
271  av_bprint_finalize(&header, NULL);
272  if (ctx->util_mem)
273  clReleaseMemObject(ctx->util_mem);
274  if (ctx->command_queue)
275  clReleaseCommandQueue(ctx->command_queue);
276  if (ctx->kernel)
277  clReleaseKernel(ctx->kernel);
278  return err;
279 }
280 
281 static int tonemap_opencl_config_output(AVFilterLink *outlink)
282 {
283  AVFilterContext *avctx = outlink->src;
284  TonemapOpenCLContext *s = avctx->priv;
285  int ret;
286  if (s->format == AV_PIX_FMT_NONE)
287  av_log(avctx, AV_LOG_WARNING, "format not set, use default format NV12\n");
288  else {
289  if (s->format != AV_PIX_FMT_P010 &&
290  s->format != AV_PIX_FMT_NV12) {
291  av_log(avctx, AV_LOG_ERROR, "unsupported output format,"
292  "only p010/nv12 supported now\n");
293  return AVERROR(EINVAL);
294  }
295  }
296 
297  s->ocf.output_format = s->format == AV_PIX_FMT_NONE ? AV_PIX_FMT_NV12 : s->format;
298  ret = ff_opencl_filter_config_output(outlink);
299  if (ret < 0)
300  return ret;
301 
302  return 0;
303 }
304 
305 static int launch_kernel(AVFilterContext *avctx, cl_kernel kernel,
306  AVFrame *output, AVFrame *input, float peak) {
307  TonemapOpenCLContext *ctx = avctx->priv;
308  int err = AVERROR(ENOSYS);
309  size_t global_work[2];
310  size_t local_work[2];
311  cl_int cle;
312 
313  CL_SET_KERNEL_ARG(kernel, 0, cl_mem, &output->data[0]);
314  CL_SET_KERNEL_ARG(kernel, 1, cl_mem, &input->data[0]);
315  CL_SET_KERNEL_ARG(kernel, 2, cl_mem, &output->data[1]);
316  CL_SET_KERNEL_ARG(kernel, 3, cl_mem, &input->data[1]);
317  CL_SET_KERNEL_ARG(kernel, 4, cl_mem, &ctx->util_mem);
318  CL_SET_KERNEL_ARG(kernel, 5, cl_float, &peak);
319 
320  local_work[0] = 16;
321  local_work[1] = 16;
322  // Note the work size based on uv plane, as we process a 2x2 quad in one workitem
323  err = ff_opencl_filter_work_size_from_image(avctx, global_work, output,
324  1, 16);
325  if (err < 0)
326  return err;
327 
328  cle = clEnqueueNDRangeKernel(ctx->command_queue, kernel, 2, NULL,
329  global_work, local_work,
330  0, NULL, NULL);
331  CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to enqueue kernel: %d.\n", cle);
332  return 0;
333 fail:
334  return err;
335 }
336 
337 static int tonemap_opencl_filter_frame(AVFilterLink *inlink, AVFrame *input)
338 {
339  AVFilterContext *avctx = inlink->dst;
340  AVFilterLink *outlink = avctx->outputs[0];
341  TonemapOpenCLContext *ctx = avctx->priv;
342  AVFrame *output = NULL;
343  cl_int cle;
344  int err;
345  double peak = ctx->peak;
346 
347  AVHWFramesContext *input_frames_ctx =
348  (AVHWFramesContext*)input->hw_frames_ctx->data;
349 
350  av_log(ctx, AV_LOG_DEBUG, "Filter input: %s, %ux%u (%"PRId64").\n",
351  av_get_pix_fmt_name(input->format),
352  input->width, input->height, input->pts);
353 
354  if (!input->hw_frames_ctx)
355  return AVERROR(EINVAL);
356 
357  output = ff_get_video_buffer(outlink, outlink->w, outlink->h);
358  if (!output) {
359  err = AVERROR(ENOMEM);
360  goto fail;
361  }
362 
363  err = av_frame_copy_props(output, input);
364  if (err < 0)
365  goto fail;
366 
367  if (!peak)
368  peak = ff_determine_signal_peak(input);
369 
370  if (ctx->trc != -1)
371  output->color_trc = ctx->trc;
372  if (ctx->primaries != -1)
373  output->color_primaries = ctx->primaries;
374  if (ctx->colorspace != -1)
375  output->colorspace = ctx->colorspace;
376  if (ctx->range != -1)
377  output->color_range = ctx->range;
378 
379  ctx->trc_in = input->color_trc;
380  ctx->trc_out = output->color_trc;
381  ctx->colorspace_in = input->colorspace;
382  ctx->colorspace_out = output->colorspace;
383  ctx->primaries_in = input->color_primaries;
384  ctx->primaries_out = output->color_primaries;
385  ctx->range_in = input->color_range;
386  ctx->range_out = output->color_range;
387  ctx->chroma_loc = output->chroma_location;
388 
389  if (!ctx->initialised) {
390  if (!(input->color_trc == AVCOL_TRC_SMPTE2084 ||
391  input->color_trc == AVCOL_TRC_ARIB_STD_B67)) {
392  av_log(ctx, AV_LOG_ERROR, "unsupported transfer function characteristic.\n");
393  err = AVERROR(ENOSYS);
394  goto fail;
395  }
396 
397  if (input_frames_ctx->sw_format != AV_PIX_FMT_P010) {
398  av_log(ctx, AV_LOG_ERROR, "unsupported format in tonemap_opencl.\n");
399  err = AVERROR(ENOSYS);
400  goto fail;
401  }
402 
403  err = tonemap_opencl_init(avctx);
404  if (err < 0)
405  goto fail;
406  }
407 
408  switch(input_frames_ctx->sw_format) {
409  case AV_PIX_FMT_P010:
410  err = launch_kernel(avctx, ctx->kernel, output, input, peak);
411  if (err < 0) goto fail;
412  break;
413  default:
414  err = AVERROR(ENOSYS);
415  goto fail;
416  }
417 
418  cle = clFinish(ctx->command_queue);
419  CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to finish command queue: %d.\n", cle);
420 
421  av_frame_free(&input);
422 
423  ff_update_hdr_metadata(output, ctx->target_peak);
424 
425  av_log(ctx, AV_LOG_DEBUG, "Tone-mapping output: %s, %ux%u (%"PRId64").\n",
426  av_get_pix_fmt_name(output->format),
427  output->width, output->height, output->pts);
428 #ifndef NDEBUG
429  {
430  uint32_t *ptr, *max_total_p, *avg_total_p, *frame_number_p;
431  float peak_detected, avg_detected;
432  unsigned map_size = (2 * DETECTION_FRAMES + 7) * sizeof(unsigned);
433  ptr = (void *)clEnqueueMapBuffer(ctx->command_queue, ctx->util_mem,
434  CL_TRUE, CL_MAP_READ, 0, map_size,
435  0, NULL, NULL, &cle);
436  // For the layout of the util buffer, refer tonemap.cl
437  if (ptr) {
438  max_total_p = ptr + 2 * (DETECTION_FRAMES + 1) + 1;
439  avg_total_p = max_total_p + 1;
440  frame_number_p = avg_total_p + 2;
441  peak_detected = (float)*max_total_p / (REFERENCE_WHITE * (*frame_number_p));
442  avg_detected = (float)*avg_total_p / (REFERENCE_WHITE * (*frame_number_p));
443  av_log(ctx, AV_LOG_DEBUG, "peak %f, avg %f will be used for next frame\n",
444  peak_detected, avg_detected);
445  clEnqueueUnmapMemObject(ctx->command_queue, ctx->util_mem, ptr, 0,
446  NULL, NULL);
447  }
448  }
449 #endif
450 
451  return ff_filter_frame(outlink, output);
452 
453 fail:
454  clFinish(ctx->command_queue);
455  av_frame_free(&input);
456  av_frame_free(&output);
457  return err;
458 }
459 
460 static av_cold void tonemap_opencl_uninit(AVFilterContext *avctx)
461 {
462  TonemapOpenCLContext *ctx = avctx->priv;
463  cl_int cle;
464 
465  if (ctx->util_mem)
466  clReleaseMemObject(ctx->util_mem);
467  if (ctx->kernel) {
468  cle = clReleaseKernel(ctx->kernel);
469  if (cle != CL_SUCCESS)
470  av_log(avctx, AV_LOG_ERROR, "Failed to release "
471  "kernel: %d.\n", cle);
472  }
473 
474  if (ctx->command_queue) {
475  cle = clReleaseCommandQueue(ctx->command_queue);
476  if (cle != CL_SUCCESS)
477  av_log(avctx, AV_LOG_ERROR, "Failed to release "
478  "command queue: %d.\n", cle);
479  }
480 
481  ff_opencl_filter_uninit(avctx);
482 }
483 
484 #define OFFSET(x) offsetof(TonemapOpenCLContext, x)
485 #define FLAGS (AV_OPT_FLAG_FILTERING_PARAM | AV_OPT_FLAG_VIDEO_PARAM)
486 static const AVOption tonemap_opencl_options[] = {
487  { "tonemap", "tonemap algorithm selection", OFFSET(tonemap), AV_OPT_TYPE_INT, {.i64 = TONEMAP_NONE}, TONEMAP_NONE, TONEMAP_MAX - 1, FLAGS, "tonemap" },
488  { "none", 0, 0, AV_OPT_TYPE_CONST, {.i64 = TONEMAP_NONE}, 0, 0, FLAGS, "tonemap" },
489  { "linear", 0, 0, AV_OPT_TYPE_CONST, {.i64 = TONEMAP_LINEAR}, 0, 0, FLAGS, "tonemap" },
490  { "gamma", 0, 0, AV_OPT_TYPE_CONST, {.i64 = TONEMAP_GAMMA}, 0, 0, FLAGS, "tonemap" },
491  { "clip", 0, 0, AV_OPT_TYPE_CONST, {.i64 = TONEMAP_CLIP}, 0, 0, FLAGS, "tonemap" },
492  { "reinhard", 0, 0, AV_OPT_TYPE_CONST, {.i64 = TONEMAP_REINHARD}, 0, 0, FLAGS, "tonemap" },
493  { "hable", 0, 0, AV_OPT_TYPE_CONST, {.i64 = TONEMAP_HABLE}, 0, 0, FLAGS, "tonemap" },
494  { "mobius", 0, 0, AV_OPT_TYPE_CONST, {.i64 = TONEMAP_MOBIUS}, 0, 0, FLAGS, "tonemap" },
495  { "transfer", "set transfer characteristic", OFFSET(trc), AV_OPT_TYPE_INT, {.i64 = AVCOL_TRC_BT709}, -1, INT_MAX, FLAGS, "transfer" },
496  { "t", "set transfer characteristic", OFFSET(trc), AV_OPT_TYPE_INT, {.i64 = AVCOL_TRC_BT709}, -1, INT_MAX, FLAGS, "transfer" },
497  { "bt709", 0, 0, AV_OPT_TYPE_CONST, {.i64 = AVCOL_TRC_BT709}, 0, 0, FLAGS, "transfer" },
498  { "bt2020", 0, 0, AV_OPT_TYPE_CONST, {.i64 = AVCOL_TRC_BT2020_10}, 0, 0, FLAGS, "transfer" },
499  { "matrix", "set colorspace matrix", OFFSET(colorspace), AV_OPT_TYPE_INT, {.i64 = -1}, -1, INT_MAX, FLAGS, "matrix" },
500  { "m", "set colorspace matrix", OFFSET(colorspace), AV_OPT_TYPE_INT, {.i64 = -1}, -1, INT_MAX, FLAGS, "matrix" },
501  { "bt709", 0, 0, AV_OPT_TYPE_CONST, {.i64 = AVCOL_SPC_BT709}, 0, 0, FLAGS, "matrix" },
502  { "bt2020", 0, 0, AV_OPT_TYPE_CONST, {.i64 = AVCOL_SPC_BT2020_NCL}, 0, 0, FLAGS, "matrix" },
503  { "primaries", "set color primaries", OFFSET(primaries), AV_OPT_TYPE_INT, {.i64 = -1}, -1, INT_MAX, FLAGS, "primaries" },
504  { "p", "set color primaries", OFFSET(primaries), AV_OPT_TYPE_INT, {.i64 = -1}, -1, INT_MAX, FLAGS, "primaries" },
505  { "bt709", 0, 0, AV_OPT_TYPE_CONST, {.i64 = AVCOL_PRI_BT709}, 0, 0, FLAGS, "primaries" },
506  { "bt2020", 0, 0, AV_OPT_TYPE_CONST, {.i64 = AVCOL_PRI_BT2020}, 0, 0, FLAGS, "primaries" },
507  { "range", "set color range", OFFSET(range), AV_OPT_TYPE_INT, {.i64 = -1}, -1, INT_MAX, FLAGS, "range" },
508  { "r", "set color range", OFFSET(range), AV_OPT_TYPE_INT, {.i64 = -1}, -1, INT_MAX, FLAGS, "range" },
509  { "tv", 0, 0, AV_OPT_TYPE_CONST, {.i64 = AVCOL_RANGE_MPEG}, 0, 0, FLAGS, "range" },
510  { "pc", 0, 0, AV_OPT_TYPE_CONST, {.i64 = AVCOL_RANGE_JPEG}, 0, 0, FLAGS, "range" },
511  { "limited", 0, 0, AV_OPT_TYPE_CONST, {.i64 = AVCOL_RANGE_MPEG}, 0, 0, FLAGS, "range" },
512  { "full", 0, 0, AV_OPT_TYPE_CONST, {.i64 = AVCOL_RANGE_JPEG}, 0, 0, FLAGS, "range" },
513  { "format", "output pixel format", OFFSET(format), AV_OPT_TYPE_PIXEL_FMT, {.i64 = AV_PIX_FMT_NONE}, AV_PIX_FMT_NONE, INT_MAX, FLAGS, "fmt" },
514  { "peak", "signal peak override", OFFSET(peak), AV_OPT_TYPE_DOUBLE, {.dbl = 0}, 0, DBL_MAX, FLAGS },
515  { "param", "tonemap parameter", OFFSET(param), AV_OPT_TYPE_DOUBLE, {.dbl = NAN}, DBL_MIN, DBL_MAX, FLAGS },
516  { "desat", "desaturation parameter", OFFSET(desat_param), AV_OPT_TYPE_DOUBLE, {.dbl = 0.5}, 0, DBL_MAX, FLAGS },
517  { "threshold", "scene detection threshold", OFFSET(scene_threshold), AV_OPT_TYPE_DOUBLE, {.dbl = 0.2}, 0, DBL_MAX, FLAGS },
518  { NULL }
519 };
520 
521 AVFILTER_DEFINE_CLASS(tonemap_opencl);
522 
523 static const AVFilterPad tonemap_opencl_inputs[] = {
524  {
525  .name = "default",
526  .type = AVMEDIA_TYPE_VIDEO,
527  .filter_frame = &tonemap_opencl_filter_frame,
528  .config_props = &ff_opencl_filter_config_input,
529  },
530 };
531 
532 static const AVFilterPad tonemap_opencl_outputs[] = {
533  {
534  .name = "default",
535  .type = AVMEDIA_TYPE_VIDEO,
536  .config_props = &tonemap_opencl_config_output,
537  },
538 };
539 
540 const AVFilter ff_vf_tonemap_opencl = {
541  .name = "tonemap_opencl",
542  .description = NULL_IF_CONFIG_SMALL("Perform HDR to SDR conversion with tonemapping."),
543  .priv_size = sizeof(TonemapOpenCLContext),
544  .priv_class = &tonemap_opencl_class,
545  .init = &ff_opencl_filter_init,
546  .uninit = &tonemap_opencl_uninit,
547  FILTER_INPUTS(tonemap_opencl_inputs),
548  FILTER_OUTPUTS(tonemap_opencl_outputs),
549  FILTER_SINGLE_PIXFMT(AV_PIX_FMT_OPENCL),
550  .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE,
551 };
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:98
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
init
static av_cold int init(AVCodecContext *avctx)
Definition: avrndec.c:31
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:93
AVColorTransferCharacteristic
AVColorTransferCharacteristic
Color Transfer Characteristic.
Definition: pixfmt.h:476
CL_SET_KERNEL_ARG
#define CL_SET_KERNEL_ARG(kernel, arg_num, type, arg)
set argument to specific Kernel.
Definition: opencl.h:61
av_bprint_finalize
int av_bprint_finalize(AVBPrint *buf, char **ret_str)
Finalize a print buffer.
Definition: bprint.c:234
out
FILE * out
Definition: movenc.c:54
av_bprint_init
void av_bprint_init(AVBPrint *buf, unsigned size_init, unsigned size_max)
Definition: bprint.c:68
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:371
TONEMAP_MAX
@ TONEMAP_MAX
Definition: vf_tonemap_opencl.c:48
ff_filter_frame
int ff_filter_frame(AVFilterLink *link, AVFrame *frame)
Send a frame of data to the next filter.
Definition: avfilter.c:1018
launch_kernel
static int launch_kernel(AVFilterContext *avctx, cl_kernel kernel, AVFrame *output, AVFrame *input, float peak)
Definition: vf_tonemap_opencl.c:305
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:281
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:112
ff_opencl_source_tonemap
const char * ff_opencl_source_tonemap
AVFrame
This structure describes decoded (raw) audio or video data.
Definition: frame.h:303
AVCOL_TRC_NB
@ AVCOL_TRC_NB
Not part of ABI.
Definition: pixfmt.h:498
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:579
DETECTION_FRAMES
#define DETECTION_FRAMES
Definition: vf_tonemap_opencl.c:38
LumaCoefficients::cb
double cb
Definition: colorspace.h:29
TonemapOpenCLContext::initialised
int initialised
Definition: vf_tonemap_opencl.c:67
opencl.h
AVOption
AVOption.
Definition: opt.h:247
LumaCoefficients
Definition: colorspace.h:28
rgb2yuv
static const char rgb2yuv[]
Definition: vf_scale_vulkan.c:65
get_rgb2rgb_matrix
static void get_rgb2rgb_matrix(enum AVColorPrimaries in, enum AVColorPrimaries out, double rgb2rgb[3][3])
Definition: vf_tonemap_opencl.c:103
ff_determine_signal_peak
double ff_determine_signal_peak(AVFrame *in)
Definition: colorspace.c:168
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
AVColorPrimaries
AVColorPrimaries
Chromaticity coordinates of the source primaries.
Definition: pixfmt.h:451
AVFilter::name
const char * name
Filter name.
Definition: avfilter.h:153
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
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
LumaCoefficients::cg
double cg
Definition: colorspace.h:29
av_color_space_name
const char * av_color_space_name(enum AVColorSpace space)
Definition: pixdesc.c:2952
TonemapOpenCLContext::peak
double peak
Definition: vf_tonemap_opencl.c:62
AVFilterContext::priv
void * priv
private data for use by the filter
Definition: avfilter.h:401
fail
#define fail()
Definition: checkasm.h:127
FLAGS
#define FLAGS
Definition: vf_tonemap_opencl.c:485
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:119
PrimaryCoefficients
Definition: colorspace.h:32
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:50
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:337
ff_get_luma_coefficients
const struct LumaCoefficients * ff_get_luma_coefficients(enum AVColorSpace csp)
Definition: colorspace.c:128
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
s
#define s(width, name)
Definition: cbs_vp9.c:257
AVCOL_PRI_NB
@ AVCOL_PRI_NB
Not part of ABI.
Definition: pixfmt.h:469
TonemapOpenCLContext::scene_threshold
double scene_threshold
Definition: vf_tonemap_opencl.c:66
AV_OPT_TYPE_DOUBLE
@ AV_OPT_TYPE_DOUBLE
Definition: opt.h:226
ff_fill_rgb2yuv_table
void ff_fill_rgb2yuv_table(const struct LumaCoefficients *coeffs, double rgb2yuv[3][3])
Definition: colorspace.c:141
av_assert0
#define av_assert0(cond)
assert() equivalent, that is always enabled.
Definition: avassert.h:37
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
ff_opencl_source_colorspace_common
const char * ff_opencl_source_colorspace_common
NAN
#define NAN
Definition: mathematics.h:64
FILTER_INPUTS
#define FILTER_INPUTS(array)
Definition: internal.h:191
tonemap_opencl_outputs
static const AVFilterPad tonemap_opencl_outputs[]
Definition: vf_tonemap_opencl.c:532
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:460
av_color_range_name
const char * av_color_range_name(enum AVColorRange range)
Definition: pixdesc.c:2892
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:537
ff_fill_rgb2xyz_table
void ff_fill_rgb2xyz_table(const struct PrimaryCoefficients *coeffs, const struct WhitepointCoefficients *wp, double rgb2xyz[3][3])
Definition: colorspace.c:68
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:453
av_color_primaries_name
const char * av_color_primaries_name(enum AVColorPrimaries primaries)
Definition: pixdesc.c:2910
AVCOL_TRC_BT2020_10
@ AVCOL_TRC_BT2020_10
ITU-R BT2020 for 10-bit system.
Definition: pixfmt.h:491
AV_PIX_FMT_OPENCL
@ AV_PIX_FMT_OPENCL
Hardware surfaces for OpenCL.
Definition: pixfmt.h:325
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:462
TonemapAlgorithm
TonemapAlgorithm
Definition: vf_tonemap.c:42
AVCOL_TRC_SMPTE2084
@ AVCOL_TRC_SMPTE2084
SMPTE ST 2084 for 10-, 12-, 14- and 16-bit systems.
Definition: pixfmt.h:493
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:117
tonemap_opencl_inputs
static const AVFilterPad tonemap_opencl_inputs[]
Definition: vf_tonemap_opencl.c:523
sdr_avg
static const float sdr_avg
Definition: vf_tonemap_opencl.c:116
primaries_table
static const struct PrimaryCoefficients primaries_table[AVCOL_PRI_NB]
Definition: vf_tonemap_opencl.c:83
ff_update_hdr_metadata
void ff_update_hdr_metadata(AVFrame *in, double peak)
Definition: colorspace.c:193
header
static const uint8_t header[24]
Definition: sdr2.c:67
whitepoint_table
static const struct WhitepointCoefficients whitepoint_table[AVCOL_PRI_NB]
Definition: vf_tonemap_opencl.c:88
REFERENCE_WHITE
#define REFERENCE_WHITE
Definition: colorspace.h:26
OPENCL_SOURCE_NB
#define OPENCL_SOURCE_NB
Definition: vf_tonemap_opencl.c:113
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:478
internal.h
AVChromaLocation
AVChromaLocation
Location of chroma samples.
Definition: pixfmt.h:598
AVFILTER_DEFINE_CLASS
AVFILTER_DEFINE_CLASS(tonemap_opencl)
WhitepointCoefficients
Definition: colorspace.h:36
FILTER_SINGLE_PIXFMT
#define FILTER_SINGLE_PIXFMT(pix_fmt_)
Definition: internal.h:181
TonemapOpenCLContext::trc_out
enum AVColorTransferCharacteristic trc trc_in trc_out
Definition: vf_tonemap_opencl.c:55
AVCOL_SPC_BT2020_NCL
@ AVCOL_SPC_BT2020_NCL
ITU-R BT2020 non-constant luminance system.
Definition: pixfmt.h:516
TonemapOpenCLContext::param
double param
Definition: vf_tonemap_opencl.c:63
AVColorSpace
AVColorSpace
YUV colorspace type.
Definition: pixfmt.h:505
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:118
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:56
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:562
ff_vf_tonemap_opencl
const AVFilter ff_vf_tonemap_opencl
Definition: vf_tonemap_opencl.c:540
AVFilter
Filter definition.
Definition: avfilter.h:149
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
av_bprintf
void av_bprintf(AVBPrint *buf, const char *fmt,...)
Definition: bprint.c:93
TONEMAP_MOBIUS
@ TONEMAP_MOBIUS
Definition: vf_tonemap_opencl.c:47
LumaCoefficients::cr
double cr
Definition: colorspace.h:29
AVCOL_TRC_ARIB_STD_B67
@ AVCOL_TRC_ARIB_STD_B67
ARIB STD-B67, known as "Hybrid log-gamma".
Definition: pixfmt.h:497
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:224
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:235
AVFilterContext
An instance of a filter.
Definition: avfilter.h:386
tonemap_opencl_options
static const AVOption tonemap_opencl_options[]
Definition: vf_tonemap_opencl.c:486
AV_PIX_FMT_P010
#define AV_PIX_FMT_P010
Definition: pixfmt.h:440
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:192
format
fg outputs[0] format
Definition: ffmpeg_filter.c:175
imgutils.h
OFFSET
#define OFFSET(x)
Definition: vf_tonemap_opencl.c:484
av_log
#define av_log(a,...)
Definition: tableprint_vlc.h:28
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:279
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:507
AVColorRange
AVColorRange
Visual content value range.
Definition: pixfmt.h:544
AV_OPT_TYPE_CONST
@ AV_OPT_TYPE_CONST
Definition: opt.h:233
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:261
av_color_transfer_name
const char * av_color_transfer_name(enum AVColorTransferCharacteristic transfer)
Definition: pixdesc.c:2931
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:2484
AVFilterContext::outputs
AVFilterLink ** outputs
array of pointers to output links
Definition: avfilter.h:398
Generated on Wed Oct 27 2021 19:21:58 for FFmpeg by   doxygen 1.8.17