tonemap.c

Go to the documentation of this file.

// Generated from libavfilter/opencl/tonemap.cl
const char *ff_source_tonemap_cl =
"/*\n"
" * This file is part of FFmpeg.\n"
" *\n"
" * FFmpeg is free software; you can redistribute it and/or\n"
" * modify it under the terms of the GNU Lesser General Public\n"
" * License as published by the Free Software Foundation; either\n"
" * version 2.1 of the License, or (at your option) any later version.\n"
" *\n"
" * FFmpeg is distributed in the hope that it will be useful,\n"
" * but WITHOUT ANY WARRANTY; without even the implied warranty of\n"
" * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the GNU\n"
" * Lesser General Public License for more details.\n"
" *\n"
" * You should have received a copy of the GNU Lesser General Public\n"
" * License along with FFmpeg; if not, write to the Free Software\n"
" * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA\n"
" */\n"
"\n"
"#define REFERENCE_WHITE 100.0f\n"
"extern float3 lrgb2yuv(float3);\n"
"extern float  lrgb2y(float3);\n"
"extern float3 yuv2lrgb(float3);\n"
"extern float3 lrgb2lrgb(float3);\n"
"extern float  get_luma_src(float3);\n"
"extern float  get_luma_dst(float3);\n"
"extern float3 ootf(float3 c, float peak);\n"
"extern float3 inverse_ootf(float3 c, float peak);\n"
"extern float3 get_chroma_sample(float3, float3, float3, float3);\n"
"\n"
"struct detection_result {\n"
"    float peak;\n"
"    float average;\n"
"};\n"
"\n"
"float hable_f(float in) {\n"
"    float a = 0.15f, b = 0.50f, c = 0.10f, d = 0.20f, e = 0.02f, f = 0.30f;\n"
"    return (in * (in * a + b * c) + d * e) / (in * (in * a + b) + d * f) - e / f;\n"
"}\n"
"\n"
"float direct(float s, float peak) {\n"
"    return s;\n"
"}\n"
"\n"
"float linear(float s, float peak) {\n"
"    return s * tone_param / peak;\n"
"}\n"
"\n"
"float gamma(float s, float peak) {\n"
"    float p = s > 0.05f ? s /peak : 0.05f / peak;\n"
"    float v = powr(p, 1.0f / tone_param);\n"
"    return s > 0.05f ? v : (s * v /0.05f);\n"
"}\n"
"\n"
"float clip(float s, float peak) {\n"
"    return clamp(s * tone_param, 0.0f, 1.0f);\n"
"}\n"
"\n"
"float reinhard(float s, float peak) {\n"
"    return s / (s + tone_param) * (peak + tone_param) / peak;\n"
"}\n"
"\n"
"float hable(float s, float peak) {\n"
"    return hable_f(s)/hable_f(peak);\n"
"}\n"
"\n"
"float mobius(float s, float peak) {\n"
"    float j = tone_param;\n"
"    float a, b;\n"
"\n"
"    if (s <= j)\n"
"        return s;\n"
"\n"
"    a = -j * j * (peak - 1.0f) / (j * j - 2.0f * j + peak);\n"
"    b = (j * j - 2.0f * j * peak + peak) / max(peak - 1.0f, 1e-6f);\n"
"\n"
"    return (b * b + 2.0f * b * j + j * j) / (b - a) * (s + a) / (s + b);\n"
"}\n"
"\n"
"// detect peak/average signal of a frame, the algorithm was ported from:\n"
"// libplacebo (https://github.com/haasn/libplacebo)\n"
"struct detection_result\n"
"detect_peak_avg(global uint *util_buf, __local uint *sum_wg,\n"
"            float signal, float peak) {\n"
"// layout of the util buffer\n"
"//\n"
"// Name:             : Size (units of 4-bytes)\n"
"// average buffer    : detection_frames + 1\n"
"// peak buffer       : detection_frames + 1\n"
"// workgroup counter : 1\n"
"// total of peak     : 1\n"
"// total of average  : 1\n"
"// frame index       : 1\n"
"// frame number      : 1\n"
"    global uint *avg_buf = util_buf;\n"
"    global uint *peak_buf = avg_buf + DETECTION_FRAMES + 1;\n"
"    global uint *counter_wg_p = peak_buf + DETECTION_FRAMES + 1;\n"
"    global uint *max_total_p = counter_wg_p + 1;\n"
"    global uint *avg_total_p = max_total_p + 1;\n"
"    global uint *frame_idx_p = avg_total_p + 1;\n"
"    global uint *scene_frame_num_p = frame_idx_p + 1;\n"
"\n"
"    uint frame_idx = *frame_idx_p;\n"
"    uint scene_frame_num = *scene_frame_num_p;\n"
"\n"
"    size_t lidx = get_local_id(0);\n"
"    size_t lidy = get_local_id(1);\n"
"    size_t lsizex = get_local_size(0);\n"
"    size_t lsizey = get_local_size(1);\n"
"    uint num_wg = get_num_groups(0) * get_num_groups(1);\n"
"    size_t group_idx = get_group_id(0);\n"
"    size_t group_idy = get_group_id(1);\n"
"    struct detection_result r = {peak, sdr_avg};\n"
"    if (lidx == 0 && lidy == 0)\n"
"        *sum_wg = 0;\n"
"    barrier(CLK_LOCAL_MEM_FENCE);\n"
"\n"
"    // update workgroup sum\n"
"    atomic_add(sum_wg, (uint)(signal * REFERENCE_WHITE));\n"
"    barrier(CLK_LOCAL_MEM_FENCE);\n"
"\n"
"    // update frame peak/avg using work-group-average.\n"
"    if (lidx == 0 && lidy == 0) {\n"
"        uint avg_wg = *sum_wg / (lsizex * lsizey);\n"
"        atomic_max(&peak_buf[frame_idx], avg_wg);\n"
"        atomic_add(&avg_buf[frame_idx], avg_wg);\n"
"    }\n"
"\n"
"    if (scene_frame_num > 0) {\n"
"        float peak = (float)*max_total_p / (REFERENCE_WHITE * scene_frame_num);\n"
"        float avg = (float)*avg_total_p / (REFERENCE_WHITE * scene_frame_num);\n"
"        r.peak = max(1.0f, peak);\n"
"        r.average = max(0.25f, avg);\n"
"    }\n"
"\n"
"    if (lidx == 0 && lidy == 0 && atomic_add(counter_wg_p, 1) == num_wg - 1) {\n"
"        *counter_wg_p = 0;\n"
"        avg_buf[frame_idx] /= num_wg;\n"
"\n"
"        if (scene_threshold > 0.0f) {\n"
"            uint cur_max = peak_buf[frame_idx];\n"
"            uint cur_avg = avg_buf[frame_idx];\n"
"            int diff = (int)(scene_frame_num * cur_avg) - (int)*avg_total_p;\n"
"\n"
"            if (abs(diff) > scene_frame_num * scene_threshold * REFERENCE_WHITE) {\n"
"                for (uint i = 0; i < DETECTION_FRAMES + 1; i++)\n"
"                  avg_buf[i] = 0;\n"
"                for (uint i = 0; i < DETECTION_FRAMES + 1; i++)\n"
"                  peak_buf[i] = 0;\n"
"                *avg_total_p = *max_total_p = 0;\n"
"                *scene_frame_num_p = 0;\n"
"                avg_buf[frame_idx] = cur_avg;\n"
"                peak_buf[frame_idx] = cur_max;\n"
"            }\n"
"        }\n"
"        uint next = (frame_idx + 1) % (DETECTION_FRAMES + 1);\n"
"        // add current frame, subtract next frame\n"
"        *max_total_p += peak_buf[frame_idx] - peak_buf[next];\n"
"        *avg_total_p += avg_buf[frame_idx] - avg_buf[next];\n"
"        // reset next frame\n"
"        peak_buf[next] = avg_buf[next] = 0;\n"
"        *frame_idx_p = next;\n"
"        *scene_frame_num_p = min(*scene_frame_num_p + 1,\n"
"                                 (uint)DETECTION_FRAMES);\n"
"    }\n"
"    return r;\n"
"}\n"
"\n"
"float3 map_one_pixel_rgb(float3 rgb, float peak, float average) {\n"
"    float sig = max(max(rgb.x, max(rgb.y, rgb.z)), 1e-6f);\n"
"\n"
"    // Rescale the variables in order to bring it into a representation where\n"
"    // 1.0 represents the dst_peak. This is because all of the tone mapping\n"
"    // algorithms are defined in such a way that they map to the range [0.0, 1.0].\n"
"    if (target_peak > 1.0f) {\n"
"        sig *= 1.0f / target_peak;\n"
"        peak *= 1.0f / target_peak;\n"
"    }\n"
"\n"
"    float sig_old = sig;\n"
"\n"
"    // Scale the signal to compensate for differences in the average brightness\n"
"    float slope = min(1.0f, sdr_avg / average);\n"
"    sig *= slope;\n"
"    peak *= slope;\n"
"\n"
"    // Desaturate the color using a coefficient dependent on the signal level\n"
"    if (desat_param > 0.0f) {\n"
"        float luma = get_luma_dst(rgb);\n"
"        float coeff = max(sig - 0.18f, 1e-6f) / max(sig, 1e-6f);\n"
"        coeff = native_powr(coeff, 10.0f / desat_param);\n"
"        rgb = mix(rgb, (float3)luma, (float3)coeff);\n"
"        sig = mix(sig, luma * slope, coeff);\n"
"    }\n"
"\n"
"    sig = TONE_FUNC(sig, peak);\n"
"\n"
"    sig = min(sig, 1.0f);\n"
"    rgb *= (sig/sig_old);\n"
"    return rgb;\n"
"}\n"
"// map from source space YUV to destination space RGB\n"
"float3 map_to_dst_space_from_yuv(float3 yuv, float peak) {\n"
"    float3 c = yuv2lrgb(yuv);\n"
"    c = ootf(c, peak);\n"
"    c = lrgb2lrgb(c);\n"
"    return c;\n"
"}\n"
"\n"
"__kernel void tonemap(__write_only image2d_t dst1,\n"
"                      __read_only  image2d_t src1,\n"
"                      __write_only image2d_t dst2,\n"
"                      __read_only  image2d_t src2,\n"
"                      global uint *util_buf,\n"
"                      float peak\n"
"                      )\n"
"{\n"
"    __local uint sum_wg;\n"
"    const sampler_t sampler = (CLK_NORMALIZED_COORDS_FALSE |\n"
"                               CLK_ADDRESS_CLAMP_TO_EDGE   |\n"
"                               CLK_FILTER_NEAREST);\n"
"    int xi = get_global_id(0);\n"
"    int yi = get_global_id(1);\n"
"    // each work item process four pixels\n"
"    int x = 2 * xi;\n"
"    int y = 2 * yi;\n"
"\n"
"    float y0 = read_imagef(src1, sampler, (int2)(x,     y)).x;\n"
"    float y1 = read_imagef(src1, sampler, (int2)(x + 1, y)).x;\n"
"    float y2 = read_imagef(src1, sampler, (int2)(x,     y + 1)).x;\n"
"    float y3 = read_imagef(src1, sampler, (int2)(x + 1, y + 1)).x;\n"
"    float2 uv = read_imagef(src2, sampler, (int2)(xi,     yi)).xy;\n"
"\n"
"    float3 c0 = map_to_dst_space_from_yuv((float3)(y0, uv.x, uv.y), peak);\n"
"    float3 c1 = map_to_dst_space_from_yuv((float3)(y1, uv.x, uv.y), peak);\n"
"    float3 c2 = map_to_dst_space_from_yuv((float3)(y2, uv.x, uv.y), peak);\n"
"    float3 c3 = map_to_dst_space_from_yuv((float3)(y3, uv.x, uv.y), peak);\n"
"\n"
"    float sig0 = max(c0.x, max(c0.y, c0.z));\n"
"    float sig1 = max(c1.x, max(c1.y, c1.z));\n"
"    float sig2 = max(c2.x, max(c2.y, c2.z));\n"
"    float sig3 = max(c3.x, max(c3.y, c3.z));\n"
"    float sig = max(sig0, max(sig1, max(sig2, sig3)));\n"
"\n"
"    struct detection_result r = detect_peak_avg(util_buf, &sum_wg, sig, peak);\n"
"\n"
"    float3 c0_old = c0, c1_old = c1, c2_old = c2;\n"
"    c0 = map_one_pixel_rgb(c0, r.peak, r.average);\n"
"    c1 = map_one_pixel_rgb(c1, r.peak, r.average);\n"
"    c2 = map_one_pixel_rgb(c2, r.peak, r.average);\n"
"    c3 = map_one_pixel_rgb(c3, r.peak, r.average);\n"
"\n"
"    c0 = inverse_ootf(c0, target_peak);\n"
"    c1 = inverse_ootf(c1, target_peak);\n"
"    c2 = inverse_ootf(c2, target_peak);\n"
"    c3 = inverse_ootf(c3, target_peak);\n"
"\n"
"    y0 = lrgb2y(c0);\n"
"    y1 = lrgb2y(c1);\n"
"    y2 = lrgb2y(c2);\n"
"    y3 = lrgb2y(c3);\n"
"    float3 chroma_c = get_chroma_sample(c0, c1, c2, c3);\n"
"    float3 chroma = lrgb2yuv(chroma_c);\n"
"\n"
"    if (xi < get_image_width(dst2) && yi < get_image_height(dst2)) {\n"
"        write_imagef(dst1, (int2)(x, y), (float4)(y0, 0.0f, 0.0f, 1.0f));\n"
"        write_imagef(dst1, (int2)(x+1, y), (float4)(y1, 0.0f, 0.0f, 1.0f));\n"
"        write_imagef(dst1, (int2)(x, y+1), (float4)(y2, 0.0f, 0.0f, 1.0f));\n"
"        write_imagef(dst1, (int2)(x+1, y+1), (float4)(y3, 0.0f, 0.0f, 1.0f));\n"
"        write_imagef(dst2, (int2)(xi, yi),\n"
"                     (float4)(chroma.y, chroma.z, 0.0f, 1.0f));\n"
"    }\n"
"}\n"
;