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