]> git.sesse.net Git - ffmpeg/blobdiff - libavfilter/vf_tonemap_opencl.c
libavfilter/vf_sr.c: Removes uint8 -> float and float -> uint8 conversions.
[ffmpeg] / libavfilter / vf_tonemap_opencl.c
index 6063456d2864bafd4d985ec8c50091505bf75638..5da2333169b888fffad39cab5c53b5bfa39c2124 100644 (file)
@@ -21,7 +21,6 @@
 #include "libavutil/bprint.h"
 #include "libavutil/common.h"
 #include "libavutil/imgutils.h"
-#include "libavutil/mastering_display_metadata.h"
 #include "libavutil/mem.h"
 #include "libavutil/opt.h"
 #include "libavutil/pixdesc.h"
@@ -40,7 +39,6 @@
 // - more format support
 
 #define DETECTION_FRAMES 63
-#define REFERENCE_WHITE 100.0f
 
 enum TonemapAlgorithm {
     TONEMAP_NONE,
@@ -124,10 +122,10 @@ static void get_rgb2rgb_matrix(enum AVColorPrimaries in, enum AVColorPrimaries o
                                double rgb2rgb[3][3]) {
     double rgb2xyz[3][3], xyz2rgb[3][3];
 
-    fill_rgb2xyz_table(&primaries_table[out], &whitepoint_table[out], rgb2xyz);
-    invert_matrix3x3(rgb2xyz, xyz2rgb);
-    fill_rgb2xyz_table(&primaries_table[in], &whitepoint_table[in], rgb2xyz);
-    mul3x3(rgb2rgb, rgb2xyz, xyz2rgb);
+    ff_fill_rgb2xyz_table(&primaries_table[out], &whitepoint_table[out], rgb2xyz);
+    ff_matrix_invert_3x3(rgb2xyz, xyz2rgb);
+    ff_fill_rgb2xyz_table(&primaries_table[in], &whitepoint_table[in], rgb2xyz);
+    ff_matrix_mul_3x3(rgb2rgb, rgb2xyz, xyz2rgb);
 }
 
 #define OPENCL_SOURCE_NB 3
@@ -262,29 +260,17 @@ static int tonemap_opencl_init(AVFilterContext *avctx)
     ctx->command_queue = clCreateCommandQueue(ctx->ocf.hwctx->context,
                                               ctx->ocf.hwctx->device_id,
                                               0, &cle);
-    if (!ctx->command_queue) {
-        av_log(avctx, AV_LOG_ERROR, "Failed to create OpenCL "
-               "command queue: %d.\n", cle);
-        err = AVERROR(EIO);
-        goto fail;
-    }
+    CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create OpenCL "
+                     "command queue %d.\n", cle);
 
     ctx->kernel = clCreateKernel(ctx->ocf.program, "tonemap", &cle);
-    if (!ctx->kernel) {
-        av_log(avctx, AV_LOG_ERROR, "Failed to create kernel: %d.\n", cle);
-        err = AVERROR(EIO);
-        goto fail;
-    }
+    CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create kernel %d.\n", cle);
 
     ctx->util_mem =
         clCreateBuffer(ctx->ocf.hwctx->context, 0,
                        (2 * DETECTION_FRAMES + 7) * sizeof(unsigned),
                        NULL, &cle);
-    if (cle != CL_SUCCESS) {
-        av_log(avctx, AV_LOG_ERROR, "Failed to create util buffer: %d.\n", cle);
-        err = AVERROR(EIO);
-        goto fail;
-    }
+    CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create util buffer: %d.\n", cle);
 
     ctx->initialised = 1;
     return 0;
@@ -349,57 +335,12 @@ static int launch_kernel(AVFilterContext *avctx, cl_kernel kernel,
     cle = clEnqueueNDRangeKernel(ctx->command_queue, kernel, 2, NULL,
                                  global_work, local_work,
                                  0, NULL, NULL);
-    if (cle != CL_SUCCESS) {
-        av_log(avctx, AV_LOG_ERROR, "Failed to enqueue kernel: %d.\n",
-               cle);
-        return AVERROR(EIO);
-    }
+    CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to enqueue kernel: %d.\n", cle);
     return 0;
 fail:
     return err;
 }
 
-static double determine_signal_peak(AVFrame *in)
-{
-    AVFrameSideData *sd = av_frame_get_side_data(in, AV_FRAME_DATA_CONTENT_LIGHT_LEVEL);
-    double peak = 0;
-
-    if (sd) {
-        AVContentLightMetadata *clm = (AVContentLightMetadata *)sd->data;
-        peak = clm->MaxCLL / REFERENCE_WHITE;
-    }
-
-    sd = av_frame_get_side_data(in, AV_FRAME_DATA_MASTERING_DISPLAY_METADATA);
-    if (!peak && sd) {
-        AVMasteringDisplayMetadata *metadata = (AVMasteringDisplayMetadata *)sd->data;
-        if (metadata->has_luminance)
-            peak = av_q2d(metadata->max_luminance) / REFERENCE_WHITE;
-    }
-
-    // For untagged source, use peak of 10000 if SMPTE ST.2084
-    // otherwise assume HLG with reference display peak 1000.
-    if (!peak)
-        peak = in->color_trc == AVCOL_TRC_SMPTE2084 ? 100.0f : 10.0f;
-
-    return peak;
-}
-
-static void update_metadata(AVFrame *in, double peak) {
-    AVFrameSideData *sd = av_frame_get_side_data(in, AV_FRAME_DATA_CONTENT_LIGHT_LEVEL);
-
-    if (sd) {
-        AVContentLightMetadata *clm = (AVContentLightMetadata *)sd->data;
-        clm->MaxCLL = (unsigned)(peak * REFERENCE_WHITE);
-    }
-
-    sd = av_frame_get_side_data(in, AV_FRAME_DATA_MASTERING_DISPLAY_METADATA);
-    if (sd) {
-        AVMasteringDisplayMetadata *metadata = (AVMasteringDisplayMetadata *)sd->data;
-        if (metadata->has_luminance)
-            metadata->max_luminance =av_d2q(peak * REFERENCE_WHITE, 10000);
-    }
-}
-
 static int tonemap_opencl_filter_frame(AVFilterLink *inlink, AVFrame *input)
 {
     AVFilterContext    *avctx = inlink->dst;
@@ -431,7 +372,7 @@ static int tonemap_opencl_filter_frame(AVFilterLink *inlink, AVFrame *input)
         goto fail;
 
     if (!peak)
-        peak = determine_signal_peak(input);
+        peak = ff_determine_signal_peak(input);
 
     if (ctx->trc != -1)
         output->color_trc = ctx->trc;
@@ -482,16 +423,11 @@ static int tonemap_opencl_filter_frame(AVFilterLink *inlink, AVFrame *input)
     }
 
     cle = clFinish(ctx->command_queue);
-    if (cle != CL_SUCCESS) {
-        av_log(avctx, AV_LOG_ERROR, "Failed to finish command queue: %d.\n",
-               cle);
-        err = AVERROR(EIO);
-        goto fail;
-    }
+    CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to finish command queue: %d.\n", cle);
 
     av_frame_free(&input);
 
-    update_metadata(output, ctx->target_peak);
+    ff_update_hdr_metadata(output, ctx->target_peak);
 
     av_log(ctx, AV_LOG_DEBUG, "Tone-mapping output: %s, %ux%u (%"PRId64").\n",
            av_get_pix_fmt_name(output->format),