avfilter/af_afir: do not over allocate fft buffer
[ffmpeg.git] / libavfilter / vf_tonemap_opencl.c
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/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"
27
28 #include "avfilter.h"
29 #include "internal.h"
30 #include "opencl.h"
31 #include "opencl_source.h"
32 #include "video.h"
33 #include "colorspace.h"
34
35 // TODO:
36 // - separate 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
40
41 #define DETECTION_FRAMES 63
42
43 enum TonemapAlgorithm {
44     TONEMAP_NONE,
45     TONEMAP_LINEAR,
46     TONEMAP_GAMMA,
47     TONEMAP_CLIP,
48     TONEMAP_REINHARD,
49     TONEMAP_HABLE,
50     TONEMAP_MOBIUS,
51     TONEMAP_MAX,
52 };
53
54 typedef struct TonemapOpenCLContext {
55     OpenCLFilterContext ocf;
56
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;
62
63     enum TonemapAlgorithm tonemap;
64     enum AVPixelFormat    format;
65     double                peak;
66     double                param;
67     double                desat_param;
68     double                target_peak;
69     double                scene_threshold;
70     int                   initialised;
71     cl_kernel             kernel;
72     cl_command_queue      command_queue;
73     cl_mem                util_mem;
74 } TonemapOpenCLContext;
75
76 static const char *yuv_coff[AVCOL_SPC_NB] = {
77     [AVCOL_SPC_BT709] = "rgb2yuv_bt709",
78     [AVCOL_SPC_BT2020_NCL] = "rgb2yuv_bt2020",
79 };
80
81 static const char *rgb_coff[AVCOL_SPC_NB] = {
82     [AVCOL_SPC_BT709] = "yuv2rgb_bt709",
83     [AVCOL_SPC_BT2020_NCL] = "yuv2rgb_bt2020",
84 };
85
86 static const char *linearize_funcs[AVCOL_TRC_NB] = {
87     [AVCOL_TRC_SMPTE2084] = "eotf_st2084",
88     [AVCOL_TRC_ARIB_STD_B67] = "inverse_oetf_hlg",
89 };
90
91 static const char *delinearize_funcs[AVCOL_TRC_NB] = {
92     [AVCOL_TRC_BT709]     = "inverse_eotf_bt1886",
93     [AVCOL_TRC_BT2020_10] = "inverse_eotf_bt1886",
94 };
95
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 },
99 };
100
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 },
104 };
105
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 },
109 };
110
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",
119 };
120
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];
124
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);
129 }
130
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;
135
136 static int tonemap_opencl_init(AVFilterContext *avctx)
137 {
138     TonemapOpenCLContext *ctx = avctx->priv;
139     int rgb2rgb_passthrough = 1;
140     double rgb2rgb[3][3];
141     struct LumaCoefficients luma_src, luma_dst;
142     cl_int cle;
143     int err;
144     AVBPrint header;
145     const char *opencl_sources[OPENCL_SOURCE_NB];
146
147     av_bprint_init(&header, 1024, AV_BPRINT_SIZE_AUTOMATIC);
148
149     switch(ctx->tonemap) {
150     case TONEMAP_GAMMA:
151         if (isnan(ctx->param))
152             ctx->param = 1.8f;
153         break;
154     case TONEMAP_REINHARD:
155         if (!isnan(ctx->param))
156             ctx->param = (1.0f - ctx->param) / ctx->param;
157         break;
158     case TONEMAP_MOBIUS:
159         if (isnan(ctx->param))
160             ctx->param = 0.3f;
161         break;
162     }
163
164     if (isnan(ctx->param))
165         ctx->param = 1.0f;
166
167     // SDR peak is 1.0f
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);
191
192     av_bprintf(&header, "__constant const float tone_param = %.4ff;\n",
193                ctx->param);
194     av_bprintf(&header, "__constant const float desat_param = %.4ff;\n",
195                ctx->desat_param);
196     av_bprintf(&header, "__constant const float target_peak = %.4ff;\n",
197                ctx->target_peak);
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);
203
204     if (ctx->primaries_out != ctx->primaries_in) {
205         get_rgb2rgb_matrix(ctx->primaries_in, ctx->primaries_out, rgb2rgb);
206         rgb2rgb_passthrough = 0;
207     }
208     if (ctx->range_in == AVCOL_RANGE_JPEG)
209         av_bprintf(&header, "#define FULL_RANGE_IN\n");
210
211     if (ctx->range_out == AVCOL_RANGE_JPEG)
212         av_bprintf(&header, "#define FULL_RANGE_OUT\n");
213
214     av_bprintf(&header, "#define chroma_loc %d\n", (int)ctx->chroma_loc);
215
216     if (rgb2rgb_passthrough)
217         av_bprintf(&header, "#define RGB2RGB_PASSTHROUGH\n");
218     else {
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]);
226     }
227
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]);
232
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);
239
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]);
243
244     if (ctx->trc_in == AVCOL_TRC_ARIB_STD_B67)
245         av_bprintf(&header, "#define ootf_impl ootf_hlg\n");
246
247     if (ctx->trc_out == AVCOL_TRC_ARIB_STD_B67)
248         av_bprintf(&header, "#define inverse_ootf_impl inverse_ootf_hlg\n");
249
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);
255
256     av_bprint_finalize(&header, NULL);
257     if (err < 0)
258         goto fail;
259
260     ctx->command_queue = clCreateCommandQueue(ctx->ocf.hwctx->context,
261                                               ctx->ocf.hwctx->device_id,
262                                               0, &cle);
263     CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create OpenCL "
264                      "command queue %d.\n", cle);
265
266     ctx->kernel = clCreateKernel(ctx->ocf.program, "tonemap", &cle);
267     CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create kernel %d.\n", cle);
268
269     ctx->util_mem =
270         clCreateBuffer(ctx->ocf.hwctx->context, 0,
271                        (2 * DETECTION_FRAMES + 7) * sizeof(unsigned),
272                        NULL, &cle);
273     CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create util buffer: %d.\n", cle);
274
275     ctx->initialised = 1;
276     return 0;
277
278 fail:
279     if (ctx->util_mem)
280         clReleaseMemObject(ctx->util_mem);
281     if (ctx->command_queue)
282         clReleaseCommandQueue(ctx->command_queue);
283     if (ctx->kernel)
284         clReleaseKernel(ctx->kernel);
285     return err;
286 }
287
288 static int tonemap_opencl_config_output(AVFilterLink *outlink)
289 {
290     AVFilterContext *avctx = outlink->src;
291     TonemapOpenCLContext *s = avctx->priv;
292     int ret;
293     if (s->format == AV_PIX_FMT_NONE)
294         av_log(avctx, AV_LOG_WARNING, "format not set, use default format NV12\n");
295     else {
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);
301       }
302     }
303
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);
306     if (ret < 0)
307         return ret;
308
309     return 0;
310 }
311
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];
318     cl_int cle;
319
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);
326
327     local_work[0]  = 16;
328     local_work[1]  = 16;
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,
331                                                 1, 16);
332     if (err < 0)
333         return err;
334
335     cle = clEnqueueNDRangeKernel(ctx->command_queue, kernel, 2, NULL,
336                                  global_work, local_work,
337                                  0, NULL, NULL);
338     CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to enqueue kernel: %d.\n", cle);
339     return 0;
340 fail:
341     return err;
342 }
343
344 static int tonemap_opencl_filter_frame(AVFilterLink *inlink, AVFrame *input)
345 {
346     AVFilterContext    *avctx = inlink->dst;
347     AVFilterLink     *outlink = avctx->outputs[0];
348     TonemapOpenCLContext *ctx = avctx->priv;
349     AVFrame *output = NULL;
350     cl_int cle;
351     int err;
352     double peak = ctx->peak;
353
354     AVHWFramesContext *input_frames_ctx =
355         (AVHWFramesContext*)input->hw_frames_ctx->data;
356
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);
360
361     if (!input->hw_frames_ctx)
362         return AVERROR(EINVAL);
363
364     output = ff_get_video_buffer(outlink, outlink->w, outlink->h);
365     if (!output) {
366         err = AVERROR(ENOMEM);
367         goto fail;
368     }
369
370     err = av_frame_copy_props(output, input);
371     if (err < 0)
372         goto fail;
373
374     if (!peak)
375         peak = ff_determine_signal_peak(input);
376
377     if (ctx->trc != -1)
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;
385
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;
395
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);
401             goto fail;
402         }
403
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);
407             goto fail;
408         }
409
410         err = tonemap_opencl_init(avctx);
411         if (err < 0)
412             goto fail;
413     }
414
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;
419         break;
420     default:
421         err = AVERROR(ENOSYS);
422         goto fail;
423     }
424
425     cle = clFinish(ctx->command_queue);
426     CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to finish command queue: %d.\n", cle);
427
428     av_frame_free(&input);
429
430     ff_update_hdr_metadata(output, ctx->target_peak);
431
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);
435 #ifndef NDEBUG
436     {
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
444         if (ptr) {
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,
453                                     NULL, NULL);
454         }
455     }
456 #endif
457
458     return ff_filter_frame(outlink, output);
459
460 fail:
461     clFinish(ctx->command_queue);
462     av_frame_free(&input);
463     av_frame_free(&output);
464     return err;
465 }
466
467 static av_cold void tonemap_opencl_uninit(AVFilterContext *avctx)
468 {
469     TonemapOpenCLContext *ctx = avctx->priv;
470     cl_int cle;
471
472     if (ctx->util_mem)
473         clReleaseMemObject(ctx->util_mem);
474     if (ctx->kernel) {
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);
479     }
480
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);
486     }
487
488     ff_opencl_filter_uninit(avctx);
489 }
490
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 },
525     { NULL }
526 };
527
528 AVFILTER_DEFINE_CLASS(tonemap_opencl);
529
530 static const AVFilterPad tonemap_opencl_inputs[] = {
531     {
532         .name         = "default",
533         .type         = AVMEDIA_TYPE_VIDEO,
534         .filter_frame = &tonemap_opencl_filter_frame,
535         .config_props = &ff_opencl_filter_config_input,
536     },
537     { NULL }
538 };
539
540 static const AVFilterPad tonemap_opencl_outputs[] = {
541     {
542         .name         = "default",
543         .type         = AVMEDIA_TYPE_VIDEO,
544         .config_props = &tonemap_opencl_config_output,
545     },
546     { NULL }
547 };
548
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,
560 };