2 * This file is part of FFmpeg.
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.
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.
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
20 #include "libavutil/avassert.h"
21 #include "libavutil/bprint.h"
22 #include "libavutil/common.h"
23 #include "libavutil/imgutils.h"
24 #include "libavutil/mem.h"
25 #include "libavutil/opt.h"
26 #include "libavutil/pixdesc.h"
31 #include "opencl_source.h"
33 #include "colorspace.h"
36 // - seperate peak-detection from tone-mapping kernel to solve
37 // one-frame-delay issue.
38 // - import colorspace matrix generation from vf_colorspace.c
39 // - more format support
41 #define DETECTION_FRAMES 63
43 enum TonemapAlgorithm {
54 typedef struct TonemapOpenCLContext {
55 OpenCLFilterContext ocf;
57 enum AVColorSpace colorspace, colorspace_in, colorspace_out;
58 enum AVColorTransferCharacteristic trc, trc_in, trc_out;
59 enum AVColorPrimaries primaries, primaries_in, primaries_out;
60 enum AVColorRange range, range_in, range_out;
61 enum AVChromaLocation chroma_loc;
63 enum TonemapAlgorithm tonemap;
64 enum AVPixelFormat format;
69 double scene_threshold;
72 cl_command_queue command_queue;
74 } TonemapOpenCLContext;
76 static const char *yuv_coff[AVCOL_SPC_NB] = {
77 [AVCOL_SPC_BT709] = "rgb2yuv_bt709",
78 [AVCOL_SPC_BT2020_NCL] = "rgb2yuv_bt2020",
81 static const char *rgb_coff[AVCOL_SPC_NB] = {
82 [AVCOL_SPC_BT709] = "yuv2rgb_bt709",
83 [AVCOL_SPC_BT2020_NCL] = "yuv2rgb_bt2020",
86 static const char *linearize_funcs[AVCOL_TRC_NB] = {
87 [AVCOL_TRC_SMPTE2084] = "eotf_st2084",
88 [AVCOL_TRC_ARIB_STD_B67] = "inverse_oetf_hlg",
91 static const char *delinearize_funcs[AVCOL_TRC_NB] = {
92 [AVCOL_TRC_BT709] = "inverse_eotf_bt1886",
93 [AVCOL_TRC_BT2020_10] = "inverse_eotf_bt1886",
96 static const struct LumaCoefficients luma_coefficients[AVCOL_SPC_NB] = {
97 [AVCOL_SPC_BT709] = { 0.2126, 0.7152, 0.0722 },
98 [AVCOL_SPC_BT2020_NCL] = { 0.2627, 0.6780, 0.0593 },
101 static struct PrimaryCoefficients primaries_table[AVCOL_PRI_NB] = {
102 [AVCOL_PRI_BT709] = { 0.640, 0.330, 0.300, 0.600, 0.150, 0.060 },
103 [AVCOL_PRI_BT2020] = { 0.708, 0.292, 0.170, 0.797, 0.131, 0.046 },
106 static struct WhitepointCoefficients whitepoint_table[AVCOL_PRI_NB] = {
107 [AVCOL_PRI_BT709] = { 0.3127, 0.3290 },
108 [AVCOL_PRI_BT2020] = { 0.3127, 0.3290 },
111 static const char *tonemap_func[TONEMAP_MAX] = {
112 [TONEMAP_NONE] = "direct",
113 [TONEMAP_LINEAR] = "linear",
114 [TONEMAP_GAMMA] = "gamma",
115 [TONEMAP_CLIP] = "clip",
116 [TONEMAP_REINHARD] = "reinhard",
117 [TONEMAP_HABLE] = "hable",
118 [TONEMAP_MOBIUS] = "mobius",
121 static void get_rgb2rgb_matrix(enum AVColorPrimaries in, enum AVColorPrimaries out,
122 double rgb2rgb[3][3]) {
123 double rgb2xyz[3][3], xyz2rgb[3][3];
125 ff_fill_rgb2xyz_table(&primaries_table[out], &whitepoint_table[out], rgb2xyz);
126 ff_matrix_invert_3x3(rgb2xyz, xyz2rgb);
127 ff_fill_rgb2xyz_table(&primaries_table[in], &whitepoint_table[in], rgb2xyz);
128 ff_matrix_mul_3x3(rgb2rgb, rgb2xyz, xyz2rgb);
131 #define OPENCL_SOURCE_NB 3
132 // Average light level for SDR signals. This is equal to a signal level of 0.5
133 // under a typical presentation gamma of about 2.0.
134 static const float sdr_avg = 0.25f;
136 static int tonemap_opencl_init(AVFilterContext *avctx)
138 TonemapOpenCLContext *ctx = avctx->priv;
139 int rgb2rgb_passthrough = 1;
140 double rgb2rgb[3][3];
141 struct LumaCoefficients luma_src, luma_dst;
145 const char *opencl_sources[OPENCL_SOURCE_NB];
147 av_bprint_init(&header, 1024, AV_BPRINT_SIZE_AUTOMATIC);
149 switch(ctx->tonemap) {
151 if (isnan(ctx->param))
154 case TONEMAP_REINHARD:
155 if (!isnan(ctx->param))
156 ctx->param = (1.0f - ctx->param) / ctx->param;
159 if (isnan(ctx->param))
164 if (isnan(ctx->param))
168 ctx->target_peak = 1.0f;
169 av_log(ctx, AV_LOG_DEBUG, "tone mapping transfer from %s to %s\n",
170 av_color_transfer_name(ctx->trc_in),
171 av_color_transfer_name(ctx->trc_out));
172 av_log(ctx, AV_LOG_DEBUG, "mapping colorspace from %s to %s\n",
173 av_color_space_name(ctx->colorspace_in),
174 av_color_space_name(ctx->colorspace_out));
175 av_log(ctx, AV_LOG_DEBUG, "mapping primaries from %s to %s\n",
176 av_color_primaries_name(ctx->primaries_in),
177 av_color_primaries_name(ctx->primaries_out));
178 av_log(ctx, AV_LOG_DEBUG, "mapping range from %s to %s\n",
179 av_color_range_name(ctx->range_in),
180 av_color_range_name(ctx->range_out));
181 // checking valid value just because of limited implementaion
182 // please remove when more functionalities are implemented
183 av_assert0(ctx->trc_out == AVCOL_TRC_BT709 ||
184 ctx->trc_out == AVCOL_TRC_BT2020_10);
185 av_assert0(ctx->trc_in == AVCOL_TRC_SMPTE2084||
186 ctx->trc_in == AVCOL_TRC_ARIB_STD_B67);
187 av_assert0(ctx->colorspace_in == AVCOL_SPC_BT2020_NCL ||
188 ctx->colorspace_in == AVCOL_SPC_BT709);
189 av_assert0(ctx->primaries_in == AVCOL_PRI_BT2020 ||
190 ctx->primaries_in == AVCOL_PRI_BT709);
192 av_bprintf(&header, "__constant const float tone_param = %.4ff;\n",
194 av_bprintf(&header, "__constant const float desat_param = %.4ff;\n",
196 av_bprintf(&header, "__constant const float target_peak = %.4ff;\n",
198 av_bprintf(&header, "__constant const float sdr_avg = %.4ff;\n", sdr_avg);
199 av_bprintf(&header, "__constant const float scene_threshold = %.4ff;\n",
200 ctx->scene_threshold);
201 av_bprintf(&header, "#define TONE_FUNC %s\n", tonemap_func[ctx->tonemap]);
202 av_bprintf(&header, "#define DETECTION_FRAMES %d\n", DETECTION_FRAMES);
204 if (ctx->primaries_out != ctx->primaries_in) {
205 get_rgb2rgb_matrix(ctx->primaries_in, ctx->primaries_out, rgb2rgb);
206 rgb2rgb_passthrough = 0;
208 if (ctx->range_in == AVCOL_RANGE_JPEG)
209 av_bprintf(&header, "#define FULL_RANGE_IN\n");
211 if (ctx->range_out == AVCOL_RANGE_JPEG)
212 av_bprintf(&header, "#define FULL_RANGE_OUT\n");
214 av_bprintf(&header, "#define chroma_loc %d\n", (int)ctx->chroma_loc);
216 if (rgb2rgb_passthrough)
217 av_bprintf(&header, "#define RGB2RGB_PASSTHROUGH\n");
219 av_bprintf(&header, "__constant float rgb2rgb[9] = {\n");
220 av_bprintf(&header, " %.4ff, %.4ff, %.4ff,\n",
221 rgb2rgb[0][0], rgb2rgb[0][1], rgb2rgb[0][2]);
222 av_bprintf(&header, " %.4ff, %.4ff, %.4ff,\n",
223 rgb2rgb[1][0], rgb2rgb[1][1], rgb2rgb[1][2]);
224 av_bprintf(&header, " %.4ff, %.4ff, %.4ff};\n",
225 rgb2rgb[2][0], rgb2rgb[2][1], rgb2rgb[2][2]);
228 av_bprintf(&header, "#define rgb_matrix %s\n",
229 rgb_coff[ctx->colorspace_in]);
230 av_bprintf(&header, "#define yuv_matrix %s\n",
231 yuv_coff[ctx->colorspace_out]);
233 luma_src = luma_coefficients[ctx->colorspace_in];
234 luma_dst = luma_coefficients[ctx->colorspace_out];
235 av_bprintf(&header, "constant float3 luma_src = {%.4ff, %.4ff, %.4ff};\n",
236 luma_src.cr, luma_src.cg, luma_src.cb);
237 av_bprintf(&header, "constant float3 luma_dst = {%.4ff, %.4ff, %.4ff};\n",
238 luma_dst.cr, luma_dst.cg, luma_dst.cb);
240 av_bprintf(&header, "#define linearize %s\n", linearize_funcs[ctx->trc_in]);
241 av_bprintf(&header, "#define delinearize %s\n",
242 delinearize_funcs[ctx->trc_out]);
244 if (ctx->trc_in == AVCOL_TRC_ARIB_STD_B67)
245 av_bprintf(&header, "#define ootf_impl ootf_hlg\n");
247 if (ctx->trc_out == AVCOL_TRC_ARIB_STD_B67)
248 av_bprintf(&header, "#define inverse_ootf_impl inverse_ootf_hlg\n");
250 av_log(avctx, AV_LOG_DEBUG, "Generated OpenCL header:\n%s\n", header.str);
251 opencl_sources[0] = header.str;
252 opencl_sources[1] = ff_opencl_source_tonemap;
253 opencl_sources[2] = ff_opencl_source_colorspace_common;
254 err = ff_opencl_filter_load_program(avctx, opencl_sources, OPENCL_SOURCE_NB);
256 av_bprint_finalize(&header, NULL);
260 ctx->command_queue = clCreateCommandQueue(ctx->ocf.hwctx->context,
261 ctx->ocf.hwctx->device_id,
263 CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create OpenCL "
264 "command queue %d.\n", cle);
266 ctx->kernel = clCreateKernel(ctx->ocf.program, "tonemap", &cle);
267 CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create kernel %d.\n", cle);
270 clCreateBuffer(ctx->ocf.hwctx->context, 0,
271 (2 * DETECTION_FRAMES + 7) * sizeof(unsigned),
273 CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create util buffer: %d.\n", cle);
275 ctx->initialised = 1;
280 clReleaseMemObject(ctx->util_mem);
281 if (ctx->command_queue)
282 clReleaseCommandQueue(ctx->command_queue);
284 clReleaseKernel(ctx->kernel);
288 static int tonemap_opencl_config_output(AVFilterLink *outlink)
290 AVFilterContext *avctx = outlink->src;
291 TonemapOpenCLContext *s = avctx->priv;
293 if (s->format == AV_PIX_FMT_NONE)
294 av_log(avctx, AV_LOG_WARNING, "format not set, use default format NV12\n");
296 if (s->format != AV_PIX_FMT_P010 &&
297 s->format != AV_PIX_FMT_NV12) {
298 av_log(avctx, AV_LOG_ERROR, "unsupported output format,"
299 "only p010/nv12 supported now\n");
300 return AVERROR(EINVAL);
304 s->ocf.output_format = s->format == AV_PIX_FMT_NONE ? AV_PIX_FMT_NV12 : s->format;
305 ret = ff_opencl_filter_config_output(outlink);
312 static int launch_kernel(AVFilterContext *avctx, cl_kernel kernel,
313 AVFrame *output, AVFrame *input, float peak) {
314 TonemapOpenCLContext *ctx = avctx->priv;
315 int err = AVERROR(ENOSYS);
316 size_t global_work[2];
317 size_t local_work[2];
320 CL_SET_KERNEL_ARG(kernel, 0, cl_mem, &output->data[0]);
321 CL_SET_KERNEL_ARG(kernel, 1, cl_mem, &input->data[0]);
322 CL_SET_KERNEL_ARG(kernel, 2, cl_mem, &output->data[1]);
323 CL_SET_KERNEL_ARG(kernel, 3, cl_mem, &input->data[1]);
324 CL_SET_KERNEL_ARG(kernel, 4, cl_mem, &ctx->util_mem);
325 CL_SET_KERNEL_ARG(kernel, 5, cl_float, &peak);
329 // Note the work size based on uv plane, as we process a 2x2 quad in one workitem
330 err = ff_opencl_filter_work_size_from_image(avctx, global_work, output,
335 cle = clEnqueueNDRangeKernel(ctx->command_queue, kernel, 2, NULL,
336 global_work, local_work,
338 CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to enqueue kernel: %d.\n", cle);
344 static int tonemap_opencl_filter_frame(AVFilterLink *inlink, AVFrame *input)
346 AVFilterContext *avctx = inlink->dst;
347 AVFilterLink *outlink = avctx->outputs[0];
348 TonemapOpenCLContext *ctx = avctx->priv;
349 AVFrame *output = NULL;
352 double peak = ctx->peak;
354 AVHWFramesContext *input_frames_ctx =
355 (AVHWFramesContext*)input->hw_frames_ctx->data;
357 av_log(ctx, AV_LOG_DEBUG, "Filter input: %s, %ux%u (%"PRId64").\n",
358 av_get_pix_fmt_name(input->format),
359 input->width, input->height, input->pts);
361 if (!input->hw_frames_ctx)
362 return AVERROR(EINVAL);
364 output = ff_get_video_buffer(outlink, outlink->w, outlink->h);
366 err = AVERROR(ENOMEM);
370 err = av_frame_copy_props(output, input);
375 peak = ff_determine_signal_peak(input);
378 output->color_trc = ctx->trc;
379 if (ctx->primaries != -1)
380 output->color_primaries = ctx->primaries;
381 if (ctx->colorspace != -1)
382 output->colorspace = ctx->colorspace;
383 if (ctx->range != -1)
384 output->color_range = ctx->range;
386 ctx->trc_in = input->color_trc;
387 ctx->trc_out = output->color_trc;
388 ctx->colorspace_in = input->colorspace;
389 ctx->colorspace_out = output->colorspace;
390 ctx->primaries_in = input->color_primaries;
391 ctx->primaries_out = output->color_primaries;
392 ctx->range_in = input->color_range;
393 ctx->range_out = output->color_range;
394 ctx->chroma_loc = output->chroma_location;
396 if (!ctx->initialised) {
397 if (!(input->color_trc == AVCOL_TRC_SMPTE2084 ||
398 input->color_trc == AVCOL_TRC_ARIB_STD_B67)) {
399 av_log(ctx, AV_LOG_ERROR, "unsupported transfer function characteristic.\n");
400 err = AVERROR(ENOSYS);
404 if (input_frames_ctx->sw_format != AV_PIX_FMT_P010) {
405 av_log(ctx, AV_LOG_ERROR, "unsupported format in tonemap_opencl.\n");
406 err = AVERROR(ENOSYS);
410 err = tonemap_opencl_init(avctx);
415 switch(input_frames_ctx->sw_format) {
416 case AV_PIX_FMT_P010:
417 err = launch_kernel(avctx, ctx->kernel, output, input, peak);
418 if (err < 0) goto fail;
421 err = AVERROR(ENOSYS);
425 cle = clFinish(ctx->command_queue);
426 CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to finish command queue: %d.\n", cle);
428 av_frame_free(&input);
430 ff_update_hdr_metadata(output, ctx->target_peak);
432 av_log(ctx, AV_LOG_DEBUG, "Tone-mapping output: %s, %ux%u (%"PRId64").\n",
433 av_get_pix_fmt_name(output->format),
434 output->width, output->height, output->pts);
437 uint32_t *ptr, *max_total_p, *avg_total_p, *frame_number_p;
438 float peak_detected, avg_detected;
439 unsigned map_size = (2 * DETECTION_FRAMES + 7) * sizeof(unsigned);
440 ptr = (void *)clEnqueueMapBuffer(ctx->command_queue, ctx->util_mem,
441 CL_TRUE, CL_MAP_READ, 0, map_size,
442 0, NULL, NULL, &cle);
443 // For the layout of the util buffer, refer tonemap.cl
445 max_total_p = ptr + 2 * (DETECTION_FRAMES + 1) + 1;
446 avg_total_p = max_total_p + 1;
447 frame_number_p = avg_total_p + 2;
448 peak_detected = (float)*max_total_p / (REFERENCE_WHITE * (*frame_number_p));
449 avg_detected = (float)*avg_total_p / (REFERENCE_WHITE * (*frame_number_p));
450 av_log(ctx, AV_LOG_DEBUG, "peak %f, avg %f will be used for next frame\n",
451 peak_detected, avg_detected);
452 clEnqueueUnmapMemObject(ctx->command_queue, ctx->util_mem, ptr, 0,
458 return ff_filter_frame(outlink, output);
461 clFinish(ctx->command_queue);
462 av_frame_free(&input);
463 av_frame_free(&output);
467 static av_cold void tonemap_opencl_uninit(AVFilterContext *avctx)
469 TonemapOpenCLContext *ctx = avctx->priv;
473 clReleaseMemObject(ctx->util_mem);
475 cle = clReleaseKernel(ctx->kernel);
476 if (cle != CL_SUCCESS)
477 av_log(avctx, AV_LOG_ERROR, "Failed to release "
478 "kernel: %d.\n", cle);
481 if (ctx->command_queue) {
482 cle = clReleaseCommandQueue(ctx->command_queue);
483 if (cle != CL_SUCCESS)
484 av_log(avctx, AV_LOG_ERROR, "Failed to release "
485 "command queue: %d.\n", cle);
488 ff_opencl_filter_uninit(avctx);
491 #define OFFSET(x) offsetof(TonemapOpenCLContext, x)
492 #define FLAGS (AV_OPT_FLAG_FILTERING_PARAM | AV_OPT_FLAG_VIDEO_PARAM)
493 static const AVOption tonemap_opencl_options[] = {
494 { "tonemap", "tonemap algorithm selection", OFFSET(tonemap), AV_OPT_TYPE_INT, {.i64 = TONEMAP_NONE}, TONEMAP_NONE, TONEMAP_MAX - 1, FLAGS, "tonemap" },
495 { "none", 0, 0, AV_OPT_TYPE_CONST, {.i64 = TONEMAP_NONE}, 0, 0, FLAGS, "tonemap" },
496 { "linear", 0, 0, AV_OPT_TYPE_CONST, {.i64 = TONEMAP_LINEAR}, 0, 0, FLAGS, "tonemap" },
497 { "gamma", 0, 0, AV_OPT_TYPE_CONST, {.i64 = TONEMAP_GAMMA}, 0, 0, FLAGS, "tonemap" },
498 { "clip", 0, 0, AV_OPT_TYPE_CONST, {.i64 = TONEMAP_CLIP}, 0, 0, FLAGS, "tonemap" },
499 { "reinhard", 0, 0, AV_OPT_TYPE_CONST, {.i64 = TONEMAP_REINHARD}, 0, 0, FLAGS, "tonemap" },
500 { "hable", 0, 0, AV_OPT_TYPE_CONST, {.i64 = TONEMAP_HABLE}, 0, 0, FLAGS, "tonemap" },
501 { "mobius", 0, 0, AV_OPT_TYPE_CONST, {.i64 = TONEMAP_MOBIUS}, 0, 0, FLAGS, "tonemap" },
502 { "transfer", "set transfer characteristic", OFFSET(trc), AV_OPT_TYPE_INT, {.i64 = AVCOL_TRC_BT709}, -1, INT_MAX, FLAGS, "transfer" },
503 { "t", "set transfer characteristic", OFFSET(trc), AV_OPT_TYPE_INT, {.i64 = AVCOL_TRC_BT709}, -1, INT_MAX, FLAGS, "transfer" },
504 { "bt709", 0, 0, AV_OPT_TYPE_CONST, {.i64 = AVCOL_TRC_BT709}, 0, 0, FLAGS, "transfer" },
505 { "bt2020", 0, 0, AV_OPT_TYPE_CONST, {.i64 = AVCOL_TRC_BT2020_10}, 0, 0, FLAGS, "transfer" },
506 { "matrix", "set colorspace matrix", OFFSET(colorspace), AV_OPT_TYPE_INT, {.i64 = -1}, -1, INT_MAX, FLAGS, "matrix" },
507 { "m", "set colorspace matrix", OFFSET(colorspace), AV_OPT_TYPE_INT, {.i64 = -1}, -1, INT_MAX, FLAGS, "matrix" },
508 { "bt709", 0, 0, AV_OPT_TYPE_CONST, {.i64 = AVCOL_SPC_BT709}, 0, 0, FLAGS, "matrix" },
509 { "bt2020", 0, 0, AV_OPT_TYPE_CONST, {.i64 = AVCOL_SPC_BT2020_NCL}, 0, 0, FLAGS, "matrix" },
510 { "primaries", "set color primaries", OFFSET(primaries), AV_OPT_TYPE_INT, {.i64 = -1}, -1, INT_MAX, FLAGS, "primaries" },
511 { "p", "set color primaries", OFFSET(primaries), AV_OPT_TYPE_INT, {.i64 = -1}, -1, INT_MAX, FLAGS, "primaries" },
512 { "bt709", 0, 0, AV_OPT_TYPE_CONST, {.i64 = AVCOL_PRI_BT709}, 0, 0, FLAGS, "primaries" },
513 { "bt2020", 0, 0, AV_OPT_TYPE_CONST, {.i64 = AVCOL_PRI_BT2020}, 0, 0, FLAGS, "primaries" },
514 { "range", "set color range", OFFSET(range), AV_OPT_TYPE_INT, {.i64 = -1}, -1, INT_MAX, FLAGS, "range" },
515 { "r", "set color range", OFFSET(range), AV_OPT_TYPE_INT, {.i64 = -1}, -1, INT_MAX, FLAGS, "range" },
516 { "tv", 0, 0, AV_OPT_TYPE_CONST, {.i64 = AVCOL_RANGE_MPEG}, 0, 0, FLAGS, "range" },
517 { "pc", 0, 0, AV_OPT_TYPE_CONST, {.i64 = AVCOL_RANGE_JPEG}, 0, 0, FLAGS, "range" },
518 { "limited", 0, 0, AV_OPT_TYPE_CONST, {.i64 = AVCOL_RANGE_MPEG}, 0, 0, FLAGS, "range" },
519 { "full", 0, 0, AV_OPT_TYPE_CONST, {.i64 = AVCOL_RANGE_JPEG}, 0, 0, FLAGS, "range" },
520 { "format", "output pixel format", OFFSET(format), AV_OPT_TYPE_PIXEL_FMT, {.i64 = AV_PIX_FMT_NONE}, AV_PIX_FMT_NONE, INT_MAX, FLAGS, "fmt" },
521 { "peak", "signal peak override", OFFSET(peak), AV_OPT_TYPE_DOUBLE, {.dbl = 0}, 0, DBL_MAX, FLAGS },
522 { "param", "tonemap parameter", OFFSET(param), AV_OPT_TYPE_DOUBLE, {.dbl = NAN}, DBL_MIN, DBL_MAX, FLAGS },
523 { "desat", "desaturation parameter", OFFSET(desat_param), AV_OPT_TYPE_DOUBLE, {.dbl = 0.5}, 0, DBL_MAX, FLAGS },
524 { "threshold", "scene detection threshold", OFFSET(scene_threshold), AV_OPT_TYPE_DOUBLE, {.dbl = 0.2}, 0, DBL_MAX, FLAGS },
528 AVFILTER_DEFINE_CLASS(tonemap_opencl);
530 static const AVFilterPad tonemap_opencl_inputs[] = {
533 .type = AVMEDIA_TYPE_VIDEO,
534 .filter_frame = &tonemap_opencl_filter_frame,
535 .config_props = &ff_opencl_filter_config_input,
540 static const AVFilterPad tonemap_opencl_outputs[] = {
543 .type = AVMEDIA_TYPE_VIDEO,
544 .config_props = &tonemap_opencl_config_output,
549 AVFilter ff_vf_tonemap_opencl = {
550 .name = "tonemap_opencl",
551 .description = NULL_IF_CONFIG_SMALL("perform HDR to SDR conversion with tonemapping"),
552 .priv_size = sizeof(TonemapOpenCLContext),
553 .priv_class = &tonemap_opencl_class,
554 .init = &ff_opencl_filter_init,
555 .uninit = &tonemap_opencl_uninit,
556 .query_formats = &ff_opencl_filter_query_formats,
557 .inputs = tonemap_opencl_inputs,
558 .outputs = tonemap_opencl_outputs,
559 .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE,