• Home
  • Line#
  • Scopes#
  • Navigate#
  • Raw
  • Download
1 /*
2  * This file is part of FFmpeg.
3  *
4  * FFmpeg is free software; you can redistribute it and/or
5  * modify it under the terms of the GNU Lesser General Public
6  * License as published by the Free Software Foundation; either
7  * version 2.1 of the License, or (at your option) any later version.
8  *
9  * FFmpeg is distributed in the hope that it will be useful,
10  * but WITHOUT ANY WARRANTY; without even the implied warranty of
11  * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU
12  * Lesser General Public License for more details.
13  *
14  * You should have received a copy of the GNU Lesser General Public
15  * License along with FFmpeg; if not, write to the Free Software
16  * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA
17  */
18 #include <float.h>
19 
20 #include "libavutil/avassert.h"
21 #include "libavutil/common.h"
22 #include "libavutil/imgutils.h"
23 #include "libavutil/opt.h"
24 #include "libavutil/pixdesc.h"
25 
26 #include "avfilter.h"
27 #include "internal.h"
28 #include "opencl.h"
29 #include "opencl_source.h"
30 #include "video.h"
31 #include "colorspace.h"
32 
33 // TODO:
34 // - separate peak-detection from tone-mapping kernel to solve
35 //    one-frame-delay issue.
36 // - more format support
37 
38 #define DETECTION_FRAMES 63
39 
40 enum TonemapAlgorithm {
41     TONEMAP_NONE,
42     TONEMAP_LINEAR,
43     TONEMAP_GAMMA,
44     TONEMAP_CLIP,
45     TONEMAP_REINHARD,
46     TONEMAP_HABLE,
47     TONEMAP_MOBIUS,
48     TONEMAP_MAX,
49 };
50 
51 typedef struct TonemapOpenCLContext {
52     OpenCLFilterContext ocf;
53 
54     enum AVColorSpace colorspace, colorspace_in, colorspace_out;
55     enum AVColorTransferCharacteristic trc, trc_in, trc_out;
56     enum AVColorPrimaries primaries, primaries_in, primaries_out;
57     enum AVColorRange range, range_in, range_out;
58     enum AVChromaLocation chroma_loc;
59 
60     enum TonemapAlgorithm tonemap;
61     enum AVPixelFormat    format;
62     double                peak;
63     double                param;
64     double                desat_param;
65     double                target_peak;
66     double                scene_threshold;
67     int                   initialised;
68     cl_kernel             kernel;
69     cl_command_queue      command_queue;
70     cl_mem                util_mem;
71 } TonemapOpenCLContext;
72 
73 static const char *const linearize_funcs[AVCOL_TRC_NB] = {
74     [AVCOL_TRC_SMPTE2084] = "eotf_st2084",
75     [AVCOL_TRC_ARIB_STD_B67] = "inverse_oetf_hlg",
76 };
77 
78 static const char *const delinearize_funcs[AVCOL_TRC_NB] = {
79     [AVCOL_TRC_BT709]     = "inverse_eotf_bt1886",
80     [AVCOL_TRC_BT2020_10] = "inverse_eotf_bt1886",
81 };
82 
83 static const char *const tonemap_func[TONEMAP_MAX] = {
84     [TONEMAP_NONE]     = "direct",
85     [TONEMAP_LINEAR]   = "linear",
86     [TONEMAP_GAMMA]    = "gamma",
87     [TONEMAP_CLIP]     = "clip",
88     [TONEMAP_REINHARD] = "reinhard",
89     [TONEMAP_HABLE]    = "hable",
90     [TONEMAP_MOBIUS]   = "mobius",
91 };
92 
get_rgb2rgb_matrix(enum AVColorPrimaries in,enum AVColorPrimaries out,double rgb2rgb[3][3])93 static int get_rgb2rgb_matrix(enum AVColorPrimaries in, enum AVColorPrimaries out,
94                               double rgb2rgb[3][3]) {
95     double rgb2xyz[3][3], xyz2rgb[3][3];
96 
97     const AVColorPrimariesDesc *in_primaries = av_csp_primaries_desc_from_id(in);
98     const AVColorPrimariesDesc *out_primaries = av_csp_primaries_desc_from_id(out);
99 
100     if (!in_primaries || !out_primaries)
101         return AVERROR(EINVAL);
102 
103     ff_fill_rgb2xyz_table(&out_primaries->prim, &out_primaries->wp, rgb2xyz);
104     ff_matrix_invert_3x3(rgb2xyz, xyz2rgb);
105     ff_fill_rgb2xyz_table(&in_primaries->prim, &in_primaries->wp, rgb2xyz);
106     ff_matrix_mul_3x3(rgb2rgb, rgb2xyz, xyz2rgb);
107 
108     return 0;
109 }
110 
111 #define OPENCL_SOURCE_NB 3
112 // Average light level for SDR signals. This is equal to a signal level of 0.5
113 // under a typical presentation gamma of about 2.0.
114 static const float sdr_avg = 0.25f;
115 
tonemap_opencl_init(AVFilterContext * avctx)116 static int tonemap_opencl_init(AVFilterContext *avctx)
117 {
118     TonemapOpenCLContext *ctx = avctx->priv;
119     int rgb2rgb_passthrough = 1;
120     double rgb2rgb[3][3], rgb2yuv[3][3], yuv2rgb[3][3];
121     const AVLumaCoefficients *luma_src, *luma_dst;
122     cl_int cle;
123     int err;
124     AVBPrint header;
125     const char *opencl_sources[OPENCL_SOURCE_NB];
126 
127     av_bprint_init(&header, 1024, AV_BPRINT_SIZE_AUTOMATIC);
128 
129     switch(ctx->tonemap) {
130     case TONEMAP_GAMMA:
131         if (isnan(ctx->param))
132             ctx->param = 1.8f;
133         break;
134     case TONEMAP_REINHARD:
135         if (!isnan(ctx->param))
136             ctx->param = (1.0f - ctx->param) / ctx->param;
137         break;
138     case TONEMAP_MOBIUS:
139         if (isnan(ctx->param))
140             ctx->param = 0.3f;
141         break;
142     }
143 
144     if (isnan(ctx->param))
145         ctx->param = 1.0f;
146 
147     // SDR peak is 1.0f
148     ctx->target_peak = 1.0f;
149     av_log(ctx, AV_LOG_DEBUG, "tone mapping transfer from %s to %s\n",
150            av_color_transfer_name(ctx->trc_in),
151            av_color_transfer_name(ctx->trc_out));
152     av_log(ctx, AV_LOG_DEBUG, "mapping colorspace from %s to %s\n",
153            av_color_space_name(ctx->colorspace_in),
154            av_color_space_name(ctx->colorspace_out));
155     av_log(ctx, AV_LOG_DEBUG, "mapping primaries from %s to %s\n",
156            av_color_primaries_name(ctx->primaries_in),
157            av_color_primaries_name(ctx->primaries_out));
158     av_log(ctx, AV_LOG_DEBUG, "mapping range from %s to %s\n",
159            av_color_range_name(ctx->range_in),
160            av_color_range_name(ctx->range_out));
161     // checking valid value just because of limited implementaion
162     // please remove when more functionalities are implemented
163     av_assert0(ctx->trc_out == AVCOL_TRC_BT709 ||
164                ctx->trc_out == AVCOL_TRC_BT2020_10);
165     av_assert0(ctx->trc_in == AVCOL_TRC_SMPTE2084||
166                ctx->trc_in == AVCOL_TRC_ARIB_STD_B67);
167     av_assert0(ctx->colorspace_in == AVCOL_SPC_BT2020_NCL ||
168                ctx->colorspace_in == AVCOL_SPC_BT709);
169     av_assert0(ctx->primaries_in == AVCOL_PRI_BT2020 ||
170                ctx->primaries_in == AVCOL_PRI_BT709);
171 
172     av_bprintf(&header, "__constant const float tone_param = %.4ff;\n",
173                ctx->param);
174     av_bprintf(&header, "__constant const float desat_param = %.4ff;\n",
175                ctx->desat_param);
176     av_bprintf(&header, "__constant const float target_peak = %.4ff;\n",
177                ctx->target_peak);
178     av_bprintf(&header, "__constant const float sdr_avg = %.4ff;\n", sdr_avg);
179     av_bprintf(&header, "__constant const float scene_threshold = %.4ff;\n",
180                ctx->scene_threshold);
181     av_bprintf(&header, "#define TONE_FUNC %s\n", tonemap_func[ctx->tonemap]);
182     av_bprintf(&header, "#define DETECTION_FRAMES %d\n", DETECTION_FRAMES);
183 
184     if (ctx->primaries_out != ctx->primaries_in) {
185         if ((err = get_rgb2rgb_matrix(ctx->primaries_in, ctx->primaries_out, rgb2rgb)) < 0)
186             goto fail;
187         rgb2rgb_passthrough = 0;
188     }
189     if (ctx->range_in == AVCOL_RANGE_JPEG)
190         av_bprintf(&header, "#define FULL_RANGE_IN\n");
191 
192     if (ctx->range_out == AVCOL_RANGE_JPEG)
193         av_bprintf(&header, "#define FULL_RANGE_OUT\n");
194 
195     av_bprintf(&header, "#define chroma_loc %d\n", (int)ctx->chroma_loc);
196 
197     if (rgb2rgb_passthrough)
198         av_bprintf(&header, "#define RGB2RGB_PASSTHROUGH\n");
199     else
200         ff_opencl_print_const_matrix_3x3(&header, "rgb2rgb", rgb2rgb);
201 
202 
203     luma_src = av_csp_luma_coeffs_from_avcsp(ctx->colorspace_in);
204     if (!luma_src) {
205         err = AVERROR(EINVAL);
206         av_log(avctx, AV_LOG_ERROR, "unsupported input colorspace %d (%s)\n",
207                ctx->colorspace_in, av_color_space_name(ctx->colorspace_in));
208         goto fail;
209     }
210 
211     luma_dst = av_csp_luma_coeffs_from_avcsp(ctx->colorspace_out);
212     if (!luma_dst) {
213         err = AVERROR(EINVAL);
214         av_log(avctx, AV_LOG_ERROR, "unsupported output colorspace %d (%s)\n",
215                ctx->colorspace_out, av_color_space_name(ctx->colorspace_out));
216         goto fail;
217     }
218 
219     ff_fill_rgb2yuv_table(luma_dst, rgb2yuv);
220     ff_opencl_print_const_matrix_3x3(&header, "yuv_matrix", rgb2yuv);
221 
222     ff_fill_rgb2yuv_table(luma_src, rgb2yuv);
223     ff_matrix_invert_3x3(rgb2yuv, yuv2rgb);
224     ff_opencl_print_const_matrix_3x3(&header, "rgb_matrix", yuv2rgb);
225 
226     av_bprintf(&header, "constant float3 luma_src = {%.4ff, %.4ff, %.4ff};\n",
227                av_q2d(luma_src->cr), av_q2d(luma_src->cg), av_q2d(luma_src->cb));
228     av_bprintf(&header, "constant float3 luma_dst = {%.4ff, %.4ff, %.4ff};\n",
229                av_q2d(luma_dst->cr), av_q2d(luma_dst->cg), av_q2d(luma_dst->cb));
230 
231     av_bprintf(&header, "#define linearize %s\n", linearize_funcs[ctx->trc_in]);
232     av_bprintf(&header, "#define delinearize %s\n",
233                delinearize_funcs[ctx->trc_out]);
234 
235     if (ctx->trc_in == AVCOL_TRC_ARIB_STD_B67)
236         av_bprintf(&header, "#define ootf_impl ootf_hlg\n");
237 
238     if (ctx->trc_out == AVCOL_TRC_ARIB_STD_B67)
239         av_bprintf(&header, "#define inverse_ootf_impl inverse_ootf_hlg\n");
240 
241     av_log(avctx, AV_LOG_DEBUG, "Generated OpenCL header:\n%s\n", header.str);
242     opencl_sources[0] = header.str;
243     opencl_sources[1] = ff_opencl_source_tonemap;
244     opencl_sources[2] = ff_opencl_source_colorspace_common;
245     err = ff_opencl_filter_load_program(avctx, opencl_sources, OPENCL_SOURCE_NB);
246 
247     av_bprint_finalize(&header, NULL);
248     if (err < 0)
249         goto fail;
250 
251     ctx->command_queue = clCreateCommandQueue(ctx->ocf.hwctx->context,
252                                               ctx->ocf.hwctx->device_id,
253                                               0, &cle);
254     CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create OpenCL "
255                      "command queue %d.\n", cle);
256 
257     ctx->kernel = clCreateKernel(ctx->ocf.program, "tonemap", &cle);
258     CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create kernel %d.\n", cle);
259 
260     ctx->util_mem =
261         clCreateBuffer(ctx->ocf.hwctx->context, 0,
262                        (2 * DETECTION_FRAMES + 7) * sizeof(unsigned),
263                        NULL, &cle);
264     CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create util buffer: %d.\n", cle);
265 
266     ctx->initialised = 1;
267     return 0;
268 
269 fail:
270     av_bprint_finalize(&header, NULL);
271     if (ctx->util_mem)
272         clReleaseMemObject(ctx->util_mem);
273     if (ctx->command_queue)
274         clReleaseCommandQueue(ctx->command_queue);
275     if (ctx->kernel)
276         clReleaseKernel(ctx->kernel);
277     return err;
278 }
279 
tonemap_opencl_config_output(AVFilterLink * outlink)280 static int tonemap_opencl_config_output(AVFilterLink *outlink)
281 {
282     AVFilterContext *avctx = outlink->src;
283     TonemapOpenCLContext *s = avctx->priv;
284     int ret;
285     if (s->format == AV_PIX_FMT_NONE)
286         av_log(avctx, AV_LOG_WARNING, "format not set, use default format NV12\n");
287     else {
288       if (s->format != AV_PIX_FMT_P010 &&
289           s->format != AV_PIX_FMT_NV12) {
290         av_log(avctx, AV_LOG_ERROR, "unsupported output format,"
291                "only p010/nv12 supported now\n");
292         return AVERROR(EINVAL);
293       }
294     }
295 
296     s->ocf.output_format = s->format == AV_PIX_FMT_NONE ? AV_PIX_FMT_NV12 : s->format;
297     ret = ff_opencl_filter_config_output(outlink);
298     if (ret < 0)
299         return ret;
300 
301     return 0;
302 }
303 
launch_kernel(AVFilterContext * avctx,cl_kernel kernel,AVFrame * output,AVFrame * input,float peak)304 static int launch_kernel(AVFilterContext *avctx, cl_kernel kernel,
305                          AVFrame *output, AVFrame *input, float peak) {
306     TonemapOpenCLContext *ctx = avctx->priv;
307     int err = AVERROR(ENOSYS);
308     size_t global_work[2];
309     size_t local_work[2];
310     cl_int cle;
311 
312     CL_SET_KERNEL_ARG(kernel, 0, cl_mem, &output->data[0]);
313     CL_SET_KERNEL_ARG(kernel, 1, cl_mem, &input->data[0]);
314     CL_SET_KERNEL_ARG(kernel, 2, cl_mem, &output->data[1]);
315     CL_SET_KERNEL_ARG(kernel, 3, cl_mem, &input->data[1]);
316     CL_SET_KERNEL_ARG(kernel, 4, cl_mem, &ctx->util_mem);
317     CL_SET_KERNEL_ARG(kernel, 5, cl_float, &peak);
318 
319     local_work[0]  = 16;
320     local_work[1]  = 16;
321     // Note the work size based on uv plane, as we process a 2x2 quad in one workitem
322     err = ff_opencl_filter_work_size_from_image(avctx, global_work, output,
323                                                 1, 16);
324     if (err < 0)
325         return err;
326 
327     cle = clEnqueueNDRangeKernel(ctx->command_queue, kernel, 2, NULL,
328                                  global_work, local_work,
329                                  0, NULL, NULL);
330     CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to enqueue kernel: %d.\n", cle);
331     return 0;
332 fail:
333     return err;
334 }
335 
tonemap_opencl_filter_frame(AVFilterLink * inlink,AVFrame * input)336 static int tonemap_opencl_filter_frame(AVFilterLink *inlink, AVFrame *input)
337 {
338     AVFilterContext    *avctx = inlink->dst;
339     AVFilterLink     *outlink = avctx->outputs[0];
340     TonemapOpenCLContext *ctx = avctx->priv;
341     AVFrame *output = NULL;
342     cl_int cle;
343     int err;
344     double peak = ctx->peak;
345 
346     AVHWFramesContext *input_frames_ctx =
347         (AVHWFramesContext*)input->hw_frames_ctx->data;
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 
356     output = ff_get_video_buffer(outlink, outlink->w, outlink->h);
357     if (!output) {
358         err = AVERROR(ENOMEM);
359         goto fail;
360     }
361 
362     err = av_frame_copy_props(output, input);
363     if (err < 0)
364         goto fail;
365 
366     if (!peak)
367         peak = ff_determine_signal_peak(input);
368 
369     if (ctx->trc != -1)
370         output->color_trc = ctx->trc;
371     if (ctx->primaries != -1)
372         output->color_primaries = ctx->primaries;
373     if (ctx->colorspace != -1)
374         output->colorspace = ctx->colorspace;
375     if (ctx->range != -1)
376         output->color_range = ctx->range;
377 
378     ctx->trc_in = input->color_trc;
379     ctx->trc_out = output->color_trc;
380     ctx->colorspace_in = input->colorspace;
381     ctx->colorspace_out = output->colorspace;
382     ctx->primaries_in = input->color_primaries;
383     ctx->primaries_out = output->color_primaries;
384     ctx->range_in = input->color_range;
385     ctx->range_out = output->color_range;
386     ctx->chroma_loc = output->chroma_location;
387 
388     if (!ctx->initialised) {
389         if (!(input->color_trc == AVCOL_TRC_SMPTE2084 ||
390             input->color_trc == AVCOL_TRC_ARIB_STD_B67)) {
391             av_log(ctx, AV_LOG_ERROR, "unsupported transfer function characteristic.\n");
392             err = AVERROR(ENOSYS);
393             goto fail;
394         }
395 
396         if (input_frames_ctx->sw_format != AV_PIX_FMT_P010) {
397             av_log(ctx, AV_LOG_ERROR, "unsupported format in tonemap_opencl.\n");
398             err = AVERROR(ENOSYS);
399             goto fail;
400         }
401 
402         err = tonemap_opencl_init(avctx);
403         if (err < 0)
404             goto fail;
405     }
406 
407     switch(input_frames_ctx->sw_format) {
408     case AV_PIX_FMT_P010:
409         err = launch_kernel(avctx, ctx->kernel, output, input, peak);
410         if (err < 0) goto fail;
411         break;
412     default:
413         err = AVERROR(ENOSYS);
414         goto fail;
415     }
416 
417     cle = clFinish(ctx->command_queue);
418     CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to finish command queue: %d.\n", cle);
419 
420     av_frame_free(&input);
421 
422     ff_update_hdr_metadata(output, ctx->target_peak);
423 
424     av_log(ctx, AV_LOG_DEBUG, "Tone-mapping output: %s, %ux%u (%"PRId64").\n",
425            av_get_pix_fmt_name(output->format),
426            output->width, output->height, output->pts);
427 #ifndef NDEBUG
428     {
429         uint32_t *ptr, *max_total_p, *avg_total_p, *frame_number_p;
430         float peak_detected, avg_detected;
431         unsigned map_size = (2 * DETECTION_FRAMES  + 7) * sizeof(unsigned);
432         ptr = (void *)clEnqueueMapBuffer(ctx->command_queue, ctx->util_mem,
433                                          CL_TRUE, CL_MAP_READ, 0, map_size,
434                                          0, NULL, NULL, &cle);
435         // For the layout of the util buffer, refer tonemap.cl
436         if (ptr) {
437             max_total_p = ptr + 2 * (DETECTION_FRAMES + 1) + 1;
438             avg_total_p = max_total_p + 1;
439             frame_number_p = avg_total_p + 2;
440             peak_detected = (float)*max_total_p / (REFERENCE_WHITE * (*frame_number_p));
441             avg_detected = (float)*avg_total_p / (REFERENCE_WHITE * (*frame_number_p));
442             av_log(ctx, AV_LOG_DEBUG, "peak %f, avg %f will be used for next frame\n",
443                    peak_detected, avg_detected);
444             clEnqueueUnmapMemObject(ctx->command_queue, ctx->util_mem, ptr, 0,
445                                     NULL, NULL);
446         }
447     }
448 #endif
449 
450     return ff_filter_frame(outlink, output);
451 
452 fail:
453     clFinish(ctx->command_queue);
454     av_frame_free(&input);
455     av_frame_free(&output);
456     return err;
457 }
458 
tonemap_opencl_uninit(AVFilterContext * avctx)459 static av_cold void tonemap_opencl_uninit(AVFilterContext *avctx)
460 {
461     TonemapOpenCLContext *ctx = avctx->priv;
462     cl_int cle;
463 
464     if (ctx->util_mem)
465         clReleaseMemObject(ctx->util_mem);
466     if (ctx->kernel) {
467         cle = clReleaseKernel(ctx->kernel);
468         if (cle != CL_SUCCESS)
469             av_log(avctx, AV_LOG_ERROR, "Failed to release "
470                    "kernel: %d.\n", cle);
471     }
472 
473     if (ctx->command_queue) {
474         cle = clReleaseCommandQueue(ctx->command_queue);
475         if (cle != CL_SUCCESS)
476             av_log(avctx, AV_LOG_ERROR, "Failed to release "
477                    "command queue: %d.\n", cle);
478     }
479 
480     ff_opencl_filter_uninit(avctx);
481 }
482 
483 #define OFFSET(x) offsetof(TonemapOpenCLContext, x)
484 #define FLAGS (AV_OPT_FLAG_FILTERING_PARAM | AV_OPT_FLAG_VIDEO_PARAM)
485 static const AVOption tonemap_opencl_options[] = {
486     { "tonemap",      "tonemap algorithm selection", OFFSET(tonemap), AV_OPT_TYPE_INT, {.i64 = TONEMAP_NONE}, TONEMAP_NONE, TONEMAP_MAX - 1, FLAGS, "tonemap" },
487     {     "none",     0, 0, AV_OPT_TYPE_CONST, {.i64 = TONEMAP_NONE},              0, 0, FLAGS, "tonemap" },
488     {     "linear",   0, 0, AV_OPT_TYPE_CONST, {.i64 = TONEMAP_LINEAR},            0, 0, FLAGS, "tonemap" },
489     {     "gamma",    0, 0, AV_OPT_TYPE_CONST, {.i64 = TONEMAP_GAMMA},             0, 0, FLAGS, "tonemap" },
490     {     "clip",     0, 0, AV_OPT_TYPE_CONST, {.i64 = TONEMAP_CLIP},              0, 0, FLAGS, "tonemap" },
491     {     "reinhard", 0, 0, AV_OPT_TYPE_CONST, {.i64 = TONEMAP_REINHARD},          0, 0, FLAGS, "tonemap" },
492     {     "hable",    0, 0, AV_OPT_TYPE_CONST, {.i64 = TONEMAP_HABLE},             0, 0, FLAGS, "tonemap" },
493     {     "mobius",   0, 0, AV_OPT_TYPE_CONST, {.i64 = TONEMAP_MOBIUS},            0, 0, FLAGS, "tonemap" },
494     { "transfer", "set transfer characteristic", OFFSET(trc), AV_OPT_TYPE_INT, {.i64 = AVCOL_TRC_BT709}, -1, INT_MAX, FLAGS, "transfer" },
495     { "t",        "set transfer characteristic", OFFSET(trc), AV_OPT_TYPE_INT, {.i64 = AVCOL_TRC_BT709}, -1, INT_MAX, FLAGS, "transfer" },
496     {     "bt709",            0,       0,                 AV_OPT_TYPE_CONST, {.i64 = AVCOL_TRC_BT709},         0, 0, FLAGS, "transfer" },
497     {     "bt2020",           0,       0,                 AV_OPT_TYPE_CONST, {.i64 = AVCOL_TRC_BT2020_10},     0, 0, FLAGS, "transfer" },
498     { "matrix", "set colorspace matrix", OFFSET(colorspace), AV_OPT_TYPE_INT, {.i64 = -1}, -1, INT_MAX, FLAGS, "matrix" },
499     { "m",      "set colorspace matrix", OFFSET(colorspace), AV_OPT_TYPE_INT, {.i64 = -1}, -1, INT_MAX, FLAGS, "matrix" },
500     {     "bt709",            0,       0,                 AV_OPT_TYPE_CONST, {.i64 = AVCOL_SPC_BT709},         0, 0, FLAGS, "matrix" },
501     {     "bt2020",           0,       0,                 AV_OPT_TYPE_CONST, {.i64 = AVCOL_SPC_BT2020_NCL},    0, 0, FLAGS, "matrix" },
502     { "primaries", "set color primaries", OFFSET(primaries), AV_OPT_TYPE_INT, {.i64 = -1}, -1, INT_MAX, FLAGS, "primaries" },
503     { "p",         "set color primaries", OFFSET(primaries), AV_OPT_TYPE_INT, {.i64 = -1}, -1, INT_MAX, FLAGS, "primaries" },
504     {     "bt709",            0,       0,                 AV_OPT_TYPE_CONST, {.i64 = AVCOL_PRI_BT709},         0, 0, FLAGS, "primaries" },
505     {     "bt2020",           0,       0,                 AV_OPT_TYPE_CONST, {.i64 = AVCOL_PRI_BT2020},        0, 0, FLAGS, "primaries" },
506     { "range",         "set color range", OFFSET(range), AV_OPT_TYPE_INT, {.i64 = -1}, -1, INT_MAX, FLAGS, "range" },
507     { "r",             "set color range", OFFSET(range), AV_OPT_TYPE_INT, {.i64 = -1}, -1, INT_MAX, FLAGS, "range" },
508     {     "tv",            0,       0,                 AV_OPT_TYPE_CONST, {.i64 = AVCOL_RANGE_MPEG},         0, 0, FLAGS, "range" },
509     {     "pc",            0,       0,                 AV_OPT_TYPE_CONST, {.i64 = AVCOL_RANGE_JPEG},         0, 0, FLAGS, "range" },
510     {     "limited",       0,       0,                 AV_OPT_TYPE_CONST, {.i64 = AVCOL_RANGE_MPEG},         0, 0, FLAGS, "range" },
511     {     "full",          0,       0,                 AV_OPT_TYPE_CONST, {.i64 = AVCOL_RANGE_JPEG},         0, 0, FLAGS, "range" },
512     { "format",    "output pixel format", OFFSET(format), AV_OPT_TYPE_PIXEL_FMT, {.i64 = AV_PIX_FMT_NONE}, AV_PIX_FMT_NONE, INT_MAX, FLAGS, "fmt" },
513     { "peak",      "signal peak override", OFFSET(peak), AV_OPT_TYPE_DOUBLE, {.dbl = 0}, 0, DBL_MAX, FLAGS },
514     { "param",     "tonemap parameter",   OFFSET(param), AV_OPT_TYPE_DOUBLE, {.dbl = NAN}, DBL_MIN, DBL_MAX, FLAGS },
515     { "desat",     "desaturation parameter",   OFFSET(desat_param), AV_OPT_TYPE_DOUBLE, {.dbl = 0.5}, 0, DBL_MAX, FLAGS },
516     { "threshold", "scene detection threshold",   OFFSET(scene_threshold), AV_OPT_TYPE_DOUBLE, {.dbl = 0.2}, 0, DBL_MAX, FLAGS },
517     { NULL }
518 };
519 
520 AVFILTER_DEFINE_CLASS(tonemap_opencl);
521 
522 static const AVFilterPad tonemap_opencl_inputs[] = {
523     {
524         .name         = "default",
525         .type         = AVMEDIA_TYPE_VIDEO,
526         .filter_frame = &tonemap_opencl_filter_frame,
527         .config_props = &ff_opencl_filter_config_input,
528     },
529 };
530 
531 static const AVFilterPad tonemap_opencl_outputs[] = {
532     {
533         .name         = "default",
534         .type         = AVMEDIA_TYPE_VIDEO,
535         .config_props = &tonemap_opencl_config_output,
536     },
537 };
538 
539 const AVFilter ff_vf_tonemap_opencl = {
540     .name           = "tonemap_opencl",
541     .description    = NULL_IF_CONFIG_SMALL("Perform HDR to SDR conversion with tonemapping."),
542     .priv_size      = sizeof(TonemapOpenCLContext),
543     .priv_class     = &tonemap_opencl_class,
544     .init           = &ff_opencl_filter_init,
545     .uninit         = &tonemap_opencl_uninit,
546     FILTER_INPUTS(tonemap_opencl_inputs),
547     FILTER_OUTPUTS(tonemap_opencl_outputs),
548     FILTER_SINGLE_PIXFMT(AV_PIX_FMT_OPENCL),
549     .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE,
550 };
551