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