3 "#line 1 \"libavfilter/opencl/tonemap.cl\"\n" 5 " * This file is part of FFmpeg.\n" 7 " * FFmpeg is free software; you can redistribute it and/or\n" 8 " * modify it under the terms of the GNU Lesser General Public\n" 9 " * License as published by the Free Software Foundation; either\n" 10 " * version 2.1 of the License, or (at your option) any later version.\n" 12 " * FFmpeg is distributed in the hope that it will be useful,\n" 13 " * but WITHOUT ANY WARRANTY; without even the implied warranty of\n" 14 " * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU\n" 15 " * Lesser General Public License for more details.\n" 17 " * You should have received a copy of the GNU Lesser General Public\n" 18 " * License along with FFmpeg; if not, write to the Free Software\n" 19 " * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA\n" 22 "#define REFERENCE_WHITE 100.0f\n" 23 "extern float3 lrgb2yuv(float3);\n" 24 "extern float lrgb2y(float3);\n" 25 "extern float3 yuv2lrgb(float3);\n" 26 "extern float3 lrgb2lrgb(float3);\n" 27 "extern float get_luma_src(float3);\n" 28 "extern float get_luma_dst(float3);\n" 29 "extern float3 ootf(float3 c, float peak);\n" 30 "extern float3 inverse_ootf(float3 c, float peak);\n" 31 "extern float3 get_chroma_sample(float3, float3, float3, float3);\n" 33 "struct detection_result {\n" 38 "float hable_f(float in) {\n" 39 " float a = 0.15f, b = 0.50f, c = 0.10f, d = 0.20f, e = 0.02f, f = 0.30f;\n" 40 " return (in * (in * a + b * c) + d * e) / (in * (in * a + b) + d * f) - e / f;\n" 43 "float direct(float s, float peak) {\n" 47 "float linear(float s, float peak) {\n" 48 " return s * tone_param / peak;\n" 51 "float gamma(float s, float peak) {\n" 52 " float p = s > 0.05f ? s /peak : 0.05f / peak;\n" 53 " float v = powr(p, 1.0f / tone_param);\n" 54 " return s > 0.05f ? v : (s * v /0.05f);\n" 57 "float clip(float s, float peak) {\n" 58 " return clamp(s * tone_param, 0.0f, 1.0f);\n" 61 "float reinhard(float s, float peak) {\n" 62 " return s / (s + tone_param) * (peak + tone_param) / peak;\n" 65 "float hable(float s, float peak) {\n" 66 " return hable_f(s)/hable_f(peak);\n" 69 "float mobius(float s, float peak) {\n" 70 " float j = tone_param;\n" 76 " a = -j * j * (peak - 1.0f) / (j * j - 2.0f * j + peak);\n" 77 " b = (j * j - 2.0f * j * peak + peak) / max(peak - 1.0f, 1e-6f);\n" 79 " return (b * b + 2.0f * b * j + j * j) / (b - a) * (s + a) / (s + b);\n" 82 "// detect peak/average signal of a frame, the algorithm was ported from:\n" 83 "// libplacebo (https://github.com/haasn/libplacebo)\n" 84 "struct detection_result\n" 85 "detect_peak_avg(global uint *util_buf, __local uint *sum_wg,\n" 86 " float signal, float peak) {\n" 87 "// layout of the util buffer\n" 89 "// Name: : Size (units of 4-bytes)\n" 90 "// average buffer : detection_frames + 1\n" 91 "// peak buffer : detection_frames + 1\n" 92 "// workgroup counter : 1\n" 93 "// total of peak : 1\n" 94 "// total of average : 1\n" 95 "// frame index : 1\n" 96 "// frame number : 1\n" 97 " global uint *avg_buf = util_buf;\n" 98 " global uint *peak_buf = avg_buf + DETECTION_FRAMES + 1;\n" 99 " global uint *counter_wg_p = peak_buf + DETECTION_FRAMES + 1;\n" 100 " global uint *max_total_p = counter_wg_p + 1;\n" 101 " global uint *avg_total_p = max_total_p + 1;\n" 102 " global uint *frame_idx_p = avg_total_p + 1;\n" 103 " global uint *scene_frame_num_p = frame_idx_p + 1;\n" 105 " uint frame_idx = *frame_idx_p;\n" 106 " uint scene_frame_num = *scene_frame_num_p;\n" 108 " size_t lidx = get_local_id(0);\n" 109 " size_t lidy = get_local_id(1);\n" 110 " size_t lsizex = get_local_size(0);\n" 111 " size_t lsizey = get_local_size(1);\n" 112 " uint num_wg = get_num_groups(0) * get_num_groups(1);\n" 113 " size_t group_idx = get_group_id(0);\n" 114 " size_t group_idy = get_group_id(1);\n" 115 " struct detection_result r = {peak, sdr_avg};\n" 116 " if (lidx == 0 && lidy == 0)\n" 118 " barrier(CLK_LOCAL_MEM_FENCE);\n" 120 " // update workgroup sum\n" 121 " atomic_add(sum_wg, (uint)(signal * REFERENCE_WHITE));\n" 122 " barrier(CLK_LOCAL_MEM_FENCE);\n" 124 " // update frame peak/avg using work-group-average.\n" 125 " if (lidx == 0 && lidy == 0) {\n" 126 " uint avg_wg = *sum_wg / (lsizex * lsizey);\n" 127 " atomic_max(&peak_buf[frame_idx], avg_wg);\n" 128 " atomic_add(&avg_buf[frame_idx], avg_wg);\n" 131 " if (scene_frame_num > 0) {\n" 132 " float peak = (float)*max_total_p / (REFERENCE_WHITE * scene_frame_num);\n" 133 " float avg = (float)*avg_total_p / (REFERENCE_WHITE * scene_frame_num);\n" 134 " r.peak = max(1.0f, peak);\n" 135 " r.average = max(0.25f, avg);\n" 138 " if (lidx == 0 && lidy == 0 && atomic_add(counter_wg_p, 1) == num_wg - 1) {\n" 139 " *counter_wg_p = 0;\n" 140 " avg_buf[frame_idx] /= num_wg;\n" 142 " if (scene_threshold > 0.0f) {\n" 143 " uint cur_max = peak_buf[frame_idx];\n" 144 " uint cur_avg = avg_buf[frame_idx];\n" 145 " int diff = (int)(scene_frame_num * cur_avg) - (int)*avg_total_p;\n" 147 " if (abs(diff) > scene_frame_num * scene_threshold * REFERENCE_WHITE) {\n" 148 " for (uint i = 0; i < DETECTION_FRAMES + 1; i++)\n" 150 " for (uint i = 0; i < DETECTION_FRAMES + 1; i++)\n" 151 " peak_buf[i] = 0;\n" 152 " *avg_total_p = *max_total_p = 0;\n" 153 " *scene_frame_num_p = 0;\n" 154 " avg_buf[frame_idx] = cur_avg;\n" 155 " peak_buf[frame_idx] = cur_max;\n" 158 " uint next = (frame_idx + 1) % (DETECTION_FRAMES + 1);\n" 159 " // add current frame, subtract next frame\n" 160 " *max_total_p += peak_buf[frame_idx] - peak_buf[next];\n" 161 " *avg_total_p += avg_buf[frame_idx] - avg_buf[next];\n" 162 " // reset next frame\n" 163 " peak_buf[next] = avg_buf[next] = 0;\n" 164 " *frame_idx_p = next;\n" 165 " *scene_frame_num_p = min(*scene_frame_num_p + 1,\n" 166 " (uint)DETECTION_FRAMES);\n" 171 "float3 map_one_pixel_rgb(float3 rgb, float peak, float average) {\n" 172 " float sig = max(max(rgb.x, max(rgb.y, rgb.z)), 1e-6f);\n" 174 " // Rescale the variables in order to bring it into a representation where\n" 175 " // 1.0 represents the dst_peak. This is because all of the tone mapping\n" 176 " // algorithms are defined in such a way that they map to the range [0.0, 1.0].\n" 177 " if (target_peak > 1.0f) {\n" 178 " sig *= 1.0f / target_peak;\n" 179 " peak *= 1.0f / target_peak;\n" 182 " float sig_old = sig;\n" 184 " // Scale the signal to compensate for differences in the average brightness\n" 185 " float slope = min(1.0f, sdr_avg / average);\n" 189 " // Desaturate the color using a coefficient dependent on the signal level\n" 190 " if (desat_param > 0.0f) {\n" 191 " float luma = get_luma_dst(rgb);\n" 192 " float coeff = max(sig - 0.18f, 1e-6f) / max(sig, 1e-6f);\n" 193 " coeff = native_powr(coeff, 10.0f / desat_param);\n" 194 " rgb = mix(rgb, (float3)luma, (float3)coeff);\n" 195 " sig = mix(sig, luma * slope, coeff);\n" 198 " sig = TONE_FUNC(sig, peak);\n" 200 " sig = min(sig, 1.0f);\n" 201 " rgb *= (sig/sig_old);\n" 204 "// map from source space YUV to destination space RGB\n" 205 "float3 map_to_dst_space_from_yuv(float3 yuv, float peak) {\n" 206 " float3 c = yuv2lrgb(yuv);\n" 207 " c = ootf(c, peak);\n" 208 " c = lrgb2lrgb(c);\n" 212 "__kernel void tonemap(__write_only image2d_t dst1,\n" 213 " __read_only image2d_t src1,\n" 214 " __write_only image2d_t dst2,\n" 215 " __read_only image2d_t src2,\n" 216 " global uint *util_buf,\n" 220 " __local uint sum_wg;\n" 221 " const sampler_t sampler = (CLK_NORMALIZED_COORDS_FALSE |\n" 222 " CLK_ADDRESS_CLAMP_TO_EDGE |\n" 223 " CLK_FILTER_NEAREST);\n" 224 " int xi = get_global_id(0);\n" 225 " int yi = get_global_id(1);\n" 226 " // each work item process four pixels\n" 230 " float y0 = read_imagef(src1, sampler, (int2)(x, y)).x;\n" 231 " float y1 = read_imagef(src1, sampler, (int2)(x + 1, y)).x;\n" 232 " float y2 = read_imagef(src1, sampler, (int2)(x, y + 1)).x;\n" 233 " float y3 = read_imagef(src1, sampler, (int2)(x + 1, y + 1)).x;\n" 234 " float2 uv = read_imagef(src2, sampler, (int2)(xi, yi)).xy;\n" 236 " float3 c0 = map_to_dst_space_from_yuv((float3)(y0, uv.x, uv.y), peak);\n" 237 " float3 c1 = map_to_dst_space_from_yuv((float3)(y1, uv.x, uv.y), peak);\n" 238 " float3 c2 = map_to_dst_space_from_yuv((float3)(y2, uv.x, uv.y), peak);\n" 239 " float3 c3 = map_to_dst_space_from_yuv((float3)(y3, uv.x, uv.y), peak);\n" 241 " float sig0 = max(c0.x, max(c0.y, c0.z));\n" 242 " float sig1 = max(c1.x, max(c1.y, c1.z));\n" 243 " float sig2 = max(c2.x, max(c2.y, c2.z));\n" 244 " float sig3 = max(c3.x, max(c3.y, c3.z));\n" 245 " float sig = max(sig0, max(sig1, max(sig2, sig3)));\n" 247 " struct detection_result r = detect_peak_avg(util_buf, &sum_wg, sig, peak);\n" 249 " float3 c0_old = c0, c1_old = c1, c2_old = c2;\n" 250 " c0 = map_one_pixel_rgb(c0, r.peak, r.average);\n" 251 " c1 = map_one_pixel_rgb(c1, r.peak, r.average);\n" 252 " c2 = map_one_pixel_rgb(c2, r.peak, r.average);\n" 253 " c3 = map_one_pixel_rgb(c3, r.peak, r.average);\n" 255 " c0 = inverse_ootf(c0, target_peak);\n" 256 " c1 = inverse_ootf(c1, target_peak);\n" 257 " c2 = inverse_ootf(c2, target_peak);\n" 258 " c3 = inverse_ootf(c3, target_peak);\n" 260 " y0 = lrgb2y(c0);\n" 261 " y1 = lrgb2y(c1);\n" 262 " y2 = lrgb2y(c2);\n" 263 " y3 = lrgb2y(c3);\n" 264 " float3 chroma_c = get_chroma_sample(c0, c1, c2, c3);\n" 265 " float3 chroma = lrgb2yuv(chroma_c);\n" 267 " if (xi < get_image_width(dst2) && yi < get_image_height(dst2)) {\n" 268 " write_imagef(dst1, (int2)(x, y), (float4)(y0, 0.0f, 0.0f, 1.0f));\n" 269 " write_imagef(dst1, (int2)(x+1, y), (float4)(y1, 0.0f, 0.0f, 1.0f));\n" 270 " write_imagef(dst1, (int2)(x, y+1), (float4)(y2, 0.0f, 0.0f, 1.0f));\n" 271 " write_imagef(dst1, (int2)(x+1, y+1), (float4)(y3, 0.0f, 0.0f, 1.0f));\n" 272 " write_imagef(dst2, (int2)(xi, yi),\n" 273 " (float4)(chroma.y, chroma.z, 0.0f, 1.0f));\n" const char * ff_opencl_source_tonemap