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