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 "filters.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 
49 };
50 
51 typedef struct TonemapOpenCLContext {
53 
54  enum AVColorSpace colorspace, colorspace_in, colorspace_out;
57  enum AVColorRange range, range_in, range_out;
59 
60  /* enum TonemapAlgorithm */
61  int tonemap;
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[] = {
75  [AVCOL_TRC_SMPTE2084] = "eotf_st2084",
76  [AVCOL_TRC_ARIB_STD_B67] = "inverse_oetf_hlg",
77 };
78 
79 static const char *const delinearize_funcs[] = {
80  [AVCOL_TRC_BT709] = "inverse_eotf_bt1886",
81  [AVCOL_TRC_BT2020_10] = "inverse_eotf_bt1886",
82 };
83 
84 static const char *const tonemap_func[TONEMAP_MAX] = {
85  [TONEMAP_NONE] = "direct",
86  [TONEMAP_LINEAR] = "linear",
87  [TONEMAP_GAMMA] = "gamma",
88  [TONEMAP_CLIP] = "clip",
89  [TONEMAP_REINHARD] = "reinhard",
90  [TONEMAP_HABLE] = "hable",
91  [TONEMAP_MOBIUS] = "mobius",
92 };
93 
95  double rgb2rgb[3][3]) {
96  double rgb2xyz[3][3], xyz2rgb[3][3];
97 
98  const AVColorPrimariesDesc *in_primaries = av_csp_primaries_desc_from_id(in);
100 
101  if (!in_primaries || !out_primaries)
102  return AVERROR(EINVAL);
103 
104  ff_fill_rgb2xyz_table(&out_primaries->prim, &out_primaries->wp, rgb2xyz);
105  ff_matrix_invert_3x3(rgb2xyz, xyz2rgb);
106  ff_fill_rgb2xyz_table(&in_primaries->prim, &in_primaries->wp, rgb2xyz);
107  ff_matrix_mul_3x3(rgb2rgb, rgb2xyz, xyz2rgb);
108 
109  return 0;
110 }
111 
112 #define OPENCL_SOURCE_NB 3
113 // Average light level for SDR signals. This is equal to a signal level of 0.5
114 // under a typical presentation gamma of about 2.0.
115 static const float sdr_avg = 0.25f;
116 
118 {
119  TonemapOpenCLContext *ctx = avctx->priv;
120  int rgb2rgb_passthrough = 1;
121  double rgb2rgb[3][3], rgb2yuv[3][3], yuv2rgb[3][3];
122  const AVLumaCoefficients *luma_src, *luma_dst;
123  cl_int cle;
124  int err;
125  AVBPrint header;
126  const char *opencl_sources[OPENCL_SOURCE_NB];
127 
129 
130  switch(ctx->tonemap) {
131  case TONEMAP_GAMMA:
132  if (isnan(ctx->param))
133  ctx->param = 1.8f;
134  break;
135  case TONEMAP_REINHARD:
136  if (!isnan(ctx->param))
137  ctx->param = (1.0f - ctx->param) / ctx->param;
138  break;
139  case TONEMAP_MOBIUS:
140  if (isnan(ctx->param))
141  ctx->param = 0.3f;
142  break;
143  }
144 
145  if (isnan(ctx->param))
146  ctx->param = 1.0f;
147 
148  // SDR peak is 1.0f
149  ctx->target_peak = 1.0f;
150  av_log(ctx, AV_LOG_DEBUG, "tone mapping transfer from %s to %s\n",
151  av_color_transfer_name(ctx->trc_in),
152  av_color_transfer_name(ctx->trc_out));
153  av_log(ctx, AV_LOG_DEBUG, "mapping colorspace from %s to %s\n",
154  av_color_space_name(ctx->colorspace_in),
155  av_color_space_name(ctx->colorspace_out));
156  av_log(ctx, AV_LOG_DEBUG, "mapping primaries from %s to %s\n",
157  av_color_primaries_name(ctx->primaries_in),
158  av_color_primaries_name(ctx->primaries_out));
159  av_log(ctx, AV_LOG_DEBUG, "mapping range from %s to %s\n",
160  av_color_range_name(ctx->range_in),
161  av_color_range_name(ctx->range_out));
162  // checking valid value just because of limited implementation
163  // please remove when more functionalities are implemented
164  av_assert0(ctx->trc_out == AVCOL_TRC_BT709 ||
165  ctx->trc_out == AVCOL_TRC_BT2020_10);
166  av_assert0(ctx->trc_in == AVCOL_TRC_SMPTE2084||
167  ctx->trc_in == AVCOL_TRC_ARIB_STD_B67);
168  av_assert0(ctx->colorspace_in == AVCOL_SPC_BT2020_NCL ||
169  ctx->colorspace_in == AVCOL_SPC_BT709);
170  av_assert0(ctx->primaries_in == AVCOL_PRI_BT2020 ||
171  ctx->primaries_in == AVCOL_PRI_BT709);
172 
173  av_bprintf(&header, "__constant const float tone_param = %.4ff;\n",
174  ctx->param);
175  av_bprintf(&header, "__constant const float desat_param = %.4ff;\n",
176  ctx->desat_param);
177  av_bprintf(&header, "__constant const float target_peak = %.4ff;\n",
178  ctx->target_peak);
179  av_bprintf(&header, "__constant const float sdr_avg = %.4ff;\n", sdr_avg);
180  av_bprintf(&header, "__constant const float scene_threshold = %.4ff;\n",
181  ctx->scene_threshold);
182  av_bprintf(&header, "#define TONE_FUNC %s\n", tonemap_func[ctx->tonemap]);
183  av_bprintf(&header, "#define DETECTION_FRAMES %d\n", DETECTION_FRAMES);
184 
185  if (ctx->primaries_out != ctx->primaries_in) {
186  if ((err = get_rgb2rgb_matrix(ctx->primaries_in, ctx->primaries_out, rgb2rgb)) < 0)
187  goto fail;
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 = av_csp_luma_coeffs_from_avcsp(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 = av_csp_luma_coeffs_from_avcsp(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);
222 
223  ff_fill_rgb2yuv_table(luma_src, rgb2yuv);
226 
227  av_bprintf(&header, "constant float3 luma_src = {%.4ff, %.4ff, %.4ff};\n",
228  av_q2d(luma_src->cr), av_q2d(luma_src->cg), av_q2d(luma_src->cb));
229  av_bprintf(&header, "constant float3 luma_dst = {%.4ff, %.4ff, %.4ff};\n",
230  av_q2d(luma_dst->cr), av_q2d(luma_dst->cg), av_q2d(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_source_tonemap_cl;
245  opencl_sources[2] = ff_source_colorspace_common_cl;
246  err = ff_opencl_filter_load_program(avctx, opencl_sources, OPENCL_SOURCE_NB);
247 
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:
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 
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;
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 
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 
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  input_frames_ctx = (AVHWFramesContext*)input->hw_frames_ctx->data;
356 
357  output = ff_get_video_buffer(outlink, outlink->w, outlink->h);
358  if (!output) {
359  err = AVERROR(ENOMEM);
360  goto fail;
361  }
362 
364  if (err < 0)
365  goto fail;
366 
367  if (!peak)
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 
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);
457  return err;
458 }
459 
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 
482 }
483 
484 #define OFFSET(x) offsetof(TonemapOpenCLContext, x)
485 #define FLAGS (AV_OPT_FLAG_FILTERING_PARAM | AV_OPT_FLAG_VIDEO_PARAM)
487  { "tonemap", "tonemap algorithm selection", OFFSET(tonemap), AV_OPT_TYPE_INT, {.i64 = TONEMAP_NONE}, TONEMAP_NONE, TONEMAP_MAX - 1, FLAGS, .unit = "tonemap" },
488  { "none", 0, 0, AV_OPT_TYPE_CONST, {.i64 = TONEMAP_NONE}, 0, 0, FLAGS, .unit = "tonemap" },
489  { "linear", 0, 0, AV_OPT_TYPE_CONST, {.i64 = TONEMAP_LINEAR}, 0, 0, FLAGS, .unit = "tonemap" },
490  { "gamma", 0, 0, AV_OPT_TYPE_CONST, {.i64 = TONEMAP_GAMMA}, 0, 0, FLAGS, .unit = "tonemap" },
491  { "clip", 0, 0, AV_OPT_TYPE_CONST, {.i64 = TONEMAP_CLIP}, 0, 0, FLAGS, .unit = "tonemap" },
492  { "reinhard", 0, 0, AV_OPT_TYPE_CONST, {.i64 = TONEMAP_REINHARD}, 0, 0, FLAGS, .unit = "tonemap" },
493  { "hable", 0, 0, AV_OPT_TYPE_CONST, {.i64 = TONEMAP_HABLE}, 0, 0, FLAGS, .unit = "tonemap" },
494  { "mobius", 0, 0, AV_OPT_TYPE_CONST, {.i64 = TONEMAP_MOBIUS}, 0, 0, FLAGS, .unit = "tonemap" },
495  { "transfer", "set transfer characteristic", OFFSET(trc), AV_OPT_TYPE_INT, {.i64 = AVCOL_TRC_BT709}, -1, INT_MAX, FLAGS, .unit = "transfer" },
496  { "t", "set transfer characteristic", OFFSET(trc), AV_OPT_TYPE_INT, {.i64 = AVCOL_TRC_BT709}, -1, INT_MAX, FLAGS, .unit = "transfer" },
497  { "bt709", 0, 0, AV_OPT_TYPE_CONST, {.i64 = AVCOL_TRC_BT709}, 0, 0, FLAGS, .unit = "transfer" },
498  { "bt2020", 0, 0, AV_OPT_TYPE_CONST, {.i64 = AVCOL_TRC_BT2020_10}, 0, 0, FLAGS, .unit = "transfer" },
499  { "matrix", "set colorspace matrix", OFFSET(colorspace), AV_OPT_TYPE_INT, {.i64 = -1}, -1, INT_MAX, FLAGS, .unit = "matrix" },
500  { "m", "set colorspace matrix", OFFSET(colorspace), AV_OPT_TYPE_INT, {.i64 = -1}, -1, INT_MAX, FLAGS, .unit = "matrix" },
501  { "bt709", 0, 0, AV_OPT_TYPE_CONST, {.i64 = AVCOL_SPC_BT709}, 0, 0, FLAGS, .unit = "matrix" },
502  { "bt2020", 0, 0, AV_OPT_TYPE_CONST, {.i64 = AVCOL_SPC_BT2020_NCL}, 0, 0, FLAGS, .unit = "matrix" },
503  { "primaries", "set color primaries", OFFSET(primaries), AV_OPT_TYPE_INT, {.i64 = -1}, -1, INT_MAX, FLAGS, .unit = "primaries" },
504  { "p", "set color primaries", OFFSET(primaries), AV_OPT_TYPE_INT, {.i64 = -1}, -1, INT_MAX, FLAGS, .unit = "primaries" },
505  { "bt709", 0, 0, AV_OPT_TYPE_CONST, {.i64 = AVCOL_PRI_BT709}, 0, 0, FLAGS, .unit = "primaries" },
506  { "bt2020", 0, 0, AV_OPT_TYPE_CONST, {.i64 = AVCOL_PRI_BT2020}, 0, 0, FLAGS, .unit = "primaries" },
507  { "range", "set color range", OFFSET(range), AV_OPT_TYPE_INT, {.i64 = -1}, -1, INT_MAX, FLAGS, .unit = "range" },
508  { "r", "set color range", OFFSET(range), AV_OPT_TYPE_INT, {.i64 = -1}, -1, INT_MAX, FLAGS, .unit = "range" },
509  { "tv", 0, 0, AV_OPT_TYPE_CONST, {.i64 = AVCOL_RANGE_MPEG}, 0, 0, FLAGS, .unit = "range" },
510  { "pc", 0, 0, AV_OPT_TYPE_CONST, {.i64 = AVCOL_RANGE_JPEG}, 0, 0, FLAGS, .unit = "range" },
511  { "limited", 0, 0, AV_OPT_TYPE_CONST, {.i64 = AVCOL_RANGE_MPEG}, 0, 0, FLAGS, .unit = "range" },
512  { "full", 0, 0, AV_OPT_TYPE_CONST, {.i64 = AVCOL_RANGE_JPEG}, 0, 0, FLAGS, .unit = "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, .unit = "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 
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 
533  {
534  .name = "default",
535  .type = AVMEDIA_TYPE_VIDEO,
536  .config_props = &tonemap_opencl_config_output,
537  },
538 };
539 
541  .p.name = "tonemap_opencl",
542  .p.description = NULL_IF_CONFIG_SMALL("Perform HDR to SDR conversion with tonemapping."),
543  .p.priv_class = &tonemap_opencl_class,
544  .p.flags = AVFILTER_FLAG_HWDEVICE,
545  .priv_size = sizeof(TonemapOpenCLContext),
551  .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE,
552 };
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:118
AV_LOG_WARNING
#define AV_LOG_WARNING
Something somehow does not look correct.
Definition: log.h:216
AVPixelFormat
AVPixelFormat
Pixel format.
Definition: pixfmt.h:71
TONEMAP_REINHARD
@ TONEMAP_REINHARD
Definition: vf_tonemap_opencl.c:45
TonemapOpenCLContext::format
enum AVPixelFormat format
Definition: vf_tonemap_opencl.c:62
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:84
AVColorTransferCharacteristic
AVColorTransferCharacteristic
Color Transfer Characteristic.
Definition: pixfmt.h:666
CL_SET_KERNEL_ARG
#define CL_SET_KERNEL_ARG(kernel, arg_num, type, arg)
set argument to specific Kernel.
Definition: opencl.h:61
out
static FILE * out
Definition: movenc.c:55
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
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:1067
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:226
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: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:64
ff_vf_tonemap_opencl
const FFFilter ff_vf_tonemap_opencl
Definition: vf_tonemap_opencl.c:540
FILTER_INPUTS
#define FILTER_INPUTS(array)
Definition: filters.h:264
test::height
int height
Definition: vc1dsp.c:40
AVFrame
This structure describes decoded (raw) audio or video data.
Definition: frame.h:427
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:777
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:68
opencl.h
AVOption
AVOption.
Definition: opt.h:429
rgb2yuv
static const char rgb2yuv[]
Definition: vf_scale_vulkan.c:84
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:159
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:636
AVFilter::name
const char * name
Filter name.
Definition: avfilter.h:220
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
TonemapOpenCLContext::tonemap
int tonemap
Definition: vf_tonemap_opencl.c:61
primaries
enum AVColorPrimaries primaries
Definition: mediacodec_wrapper.c:2612
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:266
TonemapOpenCLContext::util_mem
cl_mem util_mem
Definition: vf_tonemap_opencl.c:71
av_color_space_name
const char * av_color_space_name(enum AVColorSpace space)
Definition: pixdesc.c:3856
TonemapOpenCLContext::peak
double peak
Definition: vf_tonemap_opencl.c:63
AVFilterContext::priv
void * priv
private data for use by the filter
Definition: avfilter.h:289
fail
#define fail()
Definition: checkasm.h:220
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:110
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:83
AVFilterPad
A filter pad used for either input or output.
Definition: filters.h:40
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
get_rgb2rgb_matrix
static int get_rgb2rgb_matrix(enum AVColorPrimaries in, enum AVColorPrimaries out, double rgb2rgb[3][3])
Definition: vf_tonemap_opencl.c:94
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:210
av_cold
#define av_cold
Definition: attributes.h:106
FFFilter
Definition: filters.h:267
float
float
Definition: af_crystalizer.c:122
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:95
TonemapOpenCLContext::scene_threshold
double scene_threshold
Definition: vf_tonemap_opencl.c:67
AV_OPT_TYPE_DOUBLE
@ AV_OPT_TYPE_DOUBLE
Underlying C type is double.
Definition: opt.h:267
av_q2d
static double av_q2d(AVRational a)
Convert an AVRational to a double.
Definition: rational.h:104
av_assert0
#define av_assert0(cond)
assert() equivalent, that is always enabled.
Definition: avassert.h:42
filters.h
AV_LOG_DEBUG
#define AV_LOG_DEBUG
Stuff which is only useful for libav* developers.
Definition: log.h:231
ctx
static AVFormatContext * ctx
Definition: movenc.c:49
AVLumaCoefficients::cg
AVRational cg
Definition: csp.h:49
FILTER_OUTPUTS
#define FILTER_OUTPUTS(array)
Definition: filters.h:265
NAN
#define NAN
Definition: mathematics.h:115
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:3772
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:213
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:599
format
New swscale design to change SwsGraph is what coordinates multiple passes These can include cascaded scaling error diffusion and so on Or we could have separate passes for the vertical and horizontal scaling In between each SwsPass lies a fully allocated image buffer Graph passes may have different levels of e g we can have a single threaded error diffusion pass following a multi threaded scaling pass SwsGraph is internally recreated whenever the image format
Definition: swscale-v2.txt:14
TonemapOpenCLContext::command_queue
cl_command_queue command_queue
Definition: vf_tonemap_opencl.c:70
isnan
#define isnan(x)
Definition: libm.h:342
AVCOL_PRI_BT709
@ AVCOL_PRI_BT709
also ITU-R BT1361 / IEC 61966-2-4 / SMPTE RP 177 Annex B
Definition: pixfmt.h:638
av_color_primaries_name
const char * av_color_primaries_name(enum AVColorPrimaries primaries)
Definition: pixdesc.c:3790
AVCOL_TRC_BT2020_10
@ AVCOL_TRC_BT2020_10
ITU-R BT2020 for 10-bit system.
Definition: pixfmt.h:681
AV_PIX_FMT_OPENCL
@ AV_PIX_FMT_OPENCL
Hardware surfaces for OpenCL.
Definition: pixfmt.h:358
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:647
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: filters.h:208
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:683
test::width
int width
Definition: vc1dsp.c:39
init
int(* init)(AVBSFContext *ctx)
Definition: dts2pts.c:550
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:94
av_bprint_finalize
int av_bprint_finalize(AVBPrint *buf, char **ret_str)
Finalize a print buffer.
Definition: bprint.c:235
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:115
AVFILTER_FLAG_HWDEVICE
#define AVFILTER_FLAG_HWDEVICE
The filter can create hardware frames using AVFilterContext.hw_device_ctx.
Definition: avfilter.h:188
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:1896
delinearize_funcs
static const char *const delinearize_funcs[]
Definition: vf_tonemap_opencl.c:79
range
enum AVColorRange range
Definition: mediacodec_wrapper.c:2594
header
static const uint8_t header[24]
Definition: sdr2.c:68
REFERENCE_WHITE
#define REFERENCE_WHITE
Definition: colorspace.h:27
OPENCL_SOURCE_NB
#define OPENCL_SOURCE_NB
Definition: vf_tonemap_opencl.c:112
linearize_funcs
static const char *const linearize_funcs[]
Definition: vf_tonemap_opencl.c:74
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:46
AVCOL_TRC_BT709
@ AVCOL_TRC_BT709
also ITU-R BT1361
Definition: pixfmt.h:668
AVChromaLocation
AVChromaLocation
Location of chroma samples.
Definition: pixfmt.h:796
AVFILTER_DEFINE_CLASS
AVFILTER_DEFINE_CLASS(tonemap_opencl)
TonemapOpenCLContext::trc_out
enum AVColorTransferCharacteristic trc trc_in trc_out
Definition: vf_tonemap_opencl.c:55
uninit
static void uninit(AVBSFContext *ctx)
Definition: pcm_rechunk.c:68
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:711
TonemapOpenCLContext::param
double param
Definition: vf_tonemap_opencl.c:64
AVColorSpace
AVColorSpace
YUV colorspace type.
Definition: pixfmt.h:700
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:117
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:329
AVFilterPad::name
const char * name
Pad name.
Definition: filters.h:46
AVCOL_RANGE_MPEG
@ AVCOL_RANGE_MPEG
Narrow or limited range content.
Definition: pixfmt.h:760
AVHWFramesContext
This struct describes a set or pool of "hardware" frames (i.e.
Definition: hwcontext.h:118
ff_opencl_filter_init
int ff_opencl_filter_init(AVFilterContext *avctx)
Initialise an OpenCL filter context.
Definition: opencl.c:135
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:96
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:122
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:687
AV_PIX_FMT_NONE
@ AV_PIX_FMT_NONE
Definition: pixfmt.h:72
AV_OPT_TYPE_INT
@ AV_OPT_TYPE_INT
Underlying C type is int.
Definition: opt.h:259
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:144
AV_OPT_TYPE_PIXEL_FMT
@ AV_OPT_TYPE_PIXEL_FMT
Underlying C type is enum AVPixelFormat.
Definition: opt.h:307
AVFilterContext
An instance of a filter.
Definition: avfilter.h:274
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:602
AVMEDIA_TYPE_VIDEO
@ AVMEDIA_TYPE_VIDEO
Definition: avutil.h:200
FFFilter::p
AVFilter p
The public AVFilter.
Definition: filters.h:271
TONEMAP_CLIP
@ TONEMAP_CLIP
Definition: vf_tonemap_opencl.c:44
AVColorPrimariesDesc::prim
AVPrimaryCoefficients prim
Definition: csp.h:80
imgutils.h
OFFSET
#define OFFSET(x)
Definition: vf_tonemap_opencl.c:484
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
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:702
AVColorRange
AVColorRange
Visual content value range.
Definition: pixfmt.h:742
AV_OPT_TYPE_CONST
@ AV_OPT_TYPE_CONST
Special option type for declaring named constants.
Definition: opt.h:299
FILTER_SINGLE_PIXFMT
#define FILTER_SINGLE_PIXFMT(pix_fmt_)
Definition: filters.h:254
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:263
av_color_transfer_name
const char * av_color_transfer_name(enum AVColorTransferCharacteristic transfer)
Definition: pixdesc.c:3823
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:3376
AVFilterContext::outputs
AVFilterLink ** outputs
array of pointers to output links
Definition: avfilter.h:286