> -----Original Message----- > From: ffmpeg-devel [mailto:ffmpeg-devel-boun...@ffmpeg.org] On Behalf Of > Mark Thompson > Sent: Monday, June 4, 2018 7:20 AM > To: ffmpeg-devel@ffmpeg.org > Subject: Re: [FFmpeg-devel] [PATCH] lavfi: add opencl tonemap filter. > > On 29/05/18 06:54, Ruiling Song wrote: > > This filter does HDR(HDR10/HLG) to SDR conversion with tone-mapping. > > > > An example command to use this filter with vaapi codecs: > > FFMPEG -init_hw_device vaapi=va:/dev/dri/renderD128 -init_hw_device \ > > opencl=ocl@va -hwaccel vaapi -hwaccel_device va -hwaccel_output_format \ > > vaapi -i INPUT -filter_hw_device ocl -filter_complex \ > > '[0:v]hwmap,tonemap_opencl=t=bt2020:tonemap=linear:format=p010[x1]; \ > > [x1]hwmap=derive_device=vaapi:reverse=1' -c:v hevc_vaapi -profile 2 > OUTPUT > > > > v2: > > add peak detection. > > > > Signed-off-by: Ruiling Song <ruiling.s...@intel.com> > > --- > > configure | 1 + > > libavfilter/Makefile | 2 + > > libavfilter/allfilters.c | 1 + > > libavfilter/colorspace_basic.c | 89 +++++ > > libavfilter/colorspace_basic.h | 40 ++ > > libavfilter/opencl/colorspace_basic.cl | 187 ++++++++++ > > libavfilter/opencl/tonemap.cl | 278 ++++++++++++++ > > libavfilter/opencl_source.h | 2 + > > libavfilter/vf_tonemap_opencl.c | 655 > +++++++++++++++++++++++++++++++++ > > 9 files changed, 1255 insertions(+) > > create mode 100644 libavfilter/colorspace_basic.c > > create mode 100644 libavfilter/colorspace_basic.h > > create mode 100644 libavfilter/opencl/colorspace_basic.cl > > create mode 100644 libavfilter/opencl/tonemap.cl > > create mode 100644 libavfilter/vf_tonemap_opencl.c > > This segfaults when run on CPU implementations (both AMD and Intel on > Windows) - can you check that? Maybe an out-of-bounds memory reference I am setting up the windows environment to check it, but could you show me how do you test it? So I can reproduce it.
> which doesn't get noticed on a GPU. (Apologies for the terrible report - I > can > only see it on opaque proprietary drivers, where it dies on some internal > thread > with no information at all. The filter unfortunately can't run on pocl > because of > lack of R/RG support there.) Already documented in TODO list. > > Still not sure why it fails on Mali (it doesn't feel like it uses a lot of > memory so > I'm not sure what's going wrong), but it does work well on AMD on Windows. > > What set of implementations have you tested on? I have only tested Beignet + vaapi use-case. > > > diff --git a/configure b/configure > > index e52f8f8..ee3586b 100755 > > --- a/configure > > +++ b/configure > > @@ -3401,6 +3401,7 @@ tinterlace_filter_deps="gpl" > > tinterlace_merge_test_deps="tinterlace_filter" > > tinterlace_pad_test_deps="tinterlace_filter" > > tonemap_filter_deps="const_nan" > > +tonemap_opencl_filter_deps="opencl" > > unsharp_opencl_filter_deps="opencl" > > uspp_filter_deps="gpl avcodec" > > vaguedenoiser_filter_deps="gpl" > > diff --git a/libavfilter/Makefile b/libavfilter/Makefile > > index c68ef05..0915656 100644 > > --- a/libavfilter/Makefile > > +++ b/libavfilter/Makefile > > @@ -352,6 +352,8 @@ OBJS-$(CONFIG_TINTERLACE_FILTER) += > vf_tinterlace.o > > OBJS-$(CONFIG_TLUT2_FILTER) += vf_lut2.o framesync.o > > OBJS-$(CONFIG_TMIX_FILTER) += vf_mix.o framesync.o > > OBJS-$(CONFIG_TONEMAP_FILTER) += vf_tonemap.o > > +OBJS-$(CONFIG_TONEMAP_OPENCL_FILTER) += vf_tonemap_opencl.o > colorspace_basic.o opencl.o \ > > + opencl/tonemap.o > > opencl/colorspace_basic.o > > OBJS-$(CONFIG_TRANSPOSE_FILTER) += vf_transpose.o > > OBJS-$(CONFIG_TRIM_FILTER) += trim.o > > OBJS-$(CONFIG_UNPREMULTIPLY_FILTER) += vf_premultiply.o > framesync.o > > diff --git a/libavfilter/allfilters.c b/libavfilter/allfilters.c > > index b44093d..6873bab 100644 > > --- a/libavfilter/allfilters.c > > +++ b/libavfilter/allfilters.c > > @@ -343,6 +343,7 @@ extern AVFilter ff_vf_tinterlace; > > extern AVFilter ff_vf_tlut2; > > extern AVFilter ff_vf_tmix; > > extern AVFilter ff_vf_tonemap; > > +extern AVFilter ff_vf_tonemap_opencl; > > extern AVFilter ff_vf_transpose; > > extern AVFilter ff_vf_trim; > > extern AVFilter ff_vf_unpremultiply; > > diff --git a/libavfilter/colorspace_basic.c b/libavfilter/colorspace_basic.c > > new file mode 100644 > > index 0000000..93f9f08 > > --- /dev/null > > +++ b/libavfilter/colorspace_basic.c > > The name of this file feels strange to me. It's common parts used by > colorspace- > related filters, so maybe just colorspace.c? Will fix it. > > > @@ -0,0 +1,89 @@ > > +/* > > + * This file is part of FFmpeg. > > + * > > + * FFmpeg is free software; you can redistribute it and/or > > + * modify it under the terms of the GNU Lesser General Public > > + * License as published by the Free Software Foundation; either > > + * version 2.1 of the License, or (at your option) any later version. > > + * > > + * FFmpeg is distributed in the hope that it will be useful, > > + * but WITHOUT ANY WARRANTY; without even the implied warranty of > > + * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU > > + * Lesser General Public License for more details. > > + * > > + * You should have received a copy of the GNU Lesser General Public > > + * License along with FFmpeg; if not, write to the Free Software > > + * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 > USA > > + */ > > + > > +#include "colorspace_basic.h" > > + > > + > > +void invert_matrix3x3(const double in[3][3], double out[3][3]) > > +{ > > + double m00 = in[0][0], m01 = in[0][1], m02 = in[0][2], > > + m10 = in[1][0], m11 = in[1][1], m12 = in[1][2], > > + m20 = in[2][0], m21 = in[2][1], m22 = in[2][2]; > > + int i, j; > > + double det; > > + > > + out[0][0] = (m11 * m22 - m21 * m12); > > + out[0][1] = -(m01 * m22 - m21 * m02); > > + out[0][2] = (m01 * m12 - m11 * m02); > > + out[1][0] = -(m10 * m22 - m20 * m12); > > + out[1][1] = (m00 * m22 - m20 * m02); > > + out[1][2] = -(m00 * m12 - m10 * m02); > > + out[2][0] = (m10 * m21 - m20 * m11); > > + out[2][1] = -(m00 * m21 - m20 * m01); > > + out[2][2] = (m00 * m11 - m10 * m01); > > + > > + det = m00 * out[0][0] + m10 * out[0][1] + m20 * out[0][2]; > > + det = 1.0 / det; > > + > > + for (i = 0; i < 3; i++) { > > + for (j = 0; j < 3; j++) > > + out[i][j] *= det; > > + } > > +} > > + > > +void mul3x3(double dst[3][3], const double src1[3][3], const double > src2[3][3]) > > +{ > > + int m, n; > > + > > + for (m = 0; m < 3; m++) > > + for (n = 0; n < 3; n++) > > + dst[m][n] = src2[m][0] * src1[0][n] + > > + src2[m][1] * src1[1][n] + > > + src2[m][2] * src1[2][n]; > > +} > > +/* > > + * see e.g. > http://www.brucelindbloom.com/index.html?Eqn_RGB_XYZ_Matrix.html > > + */ > > +void fill_rgb2xyz_table(const struct ColorPrimaries *coeffs, > > + const struct WhitePoint *wp, > > + double rgb2xyz[3][3]) > > +{ > > + double i[3][3], sr, sg, sb, zw; > > + > > + rgb2xyz[0][0] = coeffs->xr / coeffs->yr; > > + rgb2xyz[0][1] = coeffs->xg / coeffs->yg; > > + rgb2xyz[0][2] = coeffs->xb / coeffs->yb; > > + rgb2xyz[1][0] = rgb2xyz[1][1] = rgb2xyz[1][2] = 1.0; > > + rgb2xyz[2][0] = (1.0 - coeffs->xr - coeffs->yr) / coeffs->yr; > > + rgb2xyz[2][1] = (1.0 - coeffs->xg - coeffs->yg) / coeffs->yg; > > + rgb2xyz[2][2] = (1.0 - coeffs->xb - coeffs->yb) / coeffs->yb; > > + invert_matrix3x3(rgb2xyz, i); > > + zw = 1.0 - wp->xw - wp->yw; > > + sr = i[0][0] * wp->xw + i[0][1] * wp->yw + i[0][2] * zw; > > + sg = i[1][0] * wp->xw + i[1][1] * wp->yw + i[1][2] * zw; > > + sb = i[2][0] * wp->xw + i[2][1] * wp->yw + i[2][2] * zw; > > + rgb2xyz[0][0] *= sr; > > + rgb2xyz[0][1] *= sg; > > + rgb2xyz[0][2] *= sb; > > + rgb2xyz[1][0] *= sr; > > + rgb2xyz[1][1] *= sg; > > + rgb2xyz[1][2] *= sb; > > + rgb2xyz[2][0] *= sr; > > + rgb2xyz[2][1] *= sg; > > + rgb2xyz[2][2] *= sb; > > +} > > Since this is copied from vf_colorspace.c, please remove the static versions > from > there at the same time. Will remove it in a separate patch. I will send it out together with next version for review. > > Also, you should check whether any copyright statement needs to be propgated > into this file. I will use copyright from that file, sorry missed it. > > > diff --git a/libavfilter/colorspace_basic.h b/libavfilter/colorspace_basic.h > > new file mode 100644 > > index 0000000..5647ca6 > > --- /dev/null > > +++ b/libavfilter/colorspace_basic.h > > @@ -0,0 +1,40 @@ > > +/* > > + * This file is part of FFmpeg. > > + * > > + * FFmpeg is free software; you can redistribute it and/or > > + * modify it under the terms of the GNU Lesser General Public > > + * License as published by the Free Software Foundation; either > > + * version 2.1 of the License, or (at your option) any later version. > > + * > > + * FFmpeg is distributed in the hope that it will be useful, > > + * but WITHOUT ANY WARRANTY; without even the implied warranty of > > + * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU > > + * Lesser General Public License for more details. > > + * > > + * You should have received a copy of the GNU Lesser General Public > > + * License along with FFmpeg; if not, write to the Free Software > > + * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 > USA > > + */ > > + > > +#ifndef AVFILTER_COLORSPACE_BASIC_H > > +#define AVFILTER_COLORSPACE_BASIC_H > > + > > +#include "libavutil/common.h" > > + > > +struct LumaCoefficients { > > + double cr, cg, cb; > > +}; > > + > > +struct ColorPrimaries { > > + double xr, yr, xg, yg, xb, yb; > > +}; > > + > > +struct WhitePoint { > > + double xw, yw; > > +}; > > + > > +void invert_matrix3x3(const double in[3][3], double out[3][3]); > > +void mul3x3(double dst[3][3], const double src1[3][3], const double > src2[3][3]); > > +void fill_rgb2xyz_table(const struct ColorPrimaries *coeffs, > > + const struct WhitePoint *wp, double rgb2xyz[3][3]); > > +#endif > > diff --git a/libavfilter/opencl/colorspace_basic.cl > b/libavfilter/opencl/colorspace_basic.cl > > new file mode 100644 > > index 0000000..eaea253 > > --- /dev/null > > +++ b/libavfilter/opencl/colorspace_basic.cl > > @@ -0,0 +1,187 @@ > > +/* > > + * This file is part of FFmpeg. > > + * > > + * FFmpeg is free software; you can redistribute it and/or > > + * modify it under the terms of the GNU Lesser General Public > > + * License as published by the Free Software Foundation; either > > + * version 2.1 of the License, or (at your option) any later version. > > + * > > + * FFmpeg is distributed in the hope that it will be useful, > > + * but WITHOUT ANY WARRANTY; without even the implied warranty of > > + * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU > > + * Lesser General Public License for more details. > > + * > > + * You should have received a copy of the GNU Lesser General Public > > + * License along with FFmpeg; if not, write to the Free Software > > + * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 > USA > > + */ > > + > > +#define ST2084_MAX_LUMINANCE 10000.0f > > +#define REFERENCE_WHITE 100.0f > > +constant const float ST2084_M1 = 0.1593017578125f; > > +constant const float ST2084_M2 = 78.84375f; > > +constant const float ST2084_C1 = 0.8359375f; > > +constant const float ST2084_C2 = 18.8515625f; > > +constant const float ST2084_C3 = 18.6875f; > > + > > +__constant float yuv2rgb_bt2020[] = { > > + 1.0f, 0.0f, 1.4746f, > > + 1.0f, -0.16455f, -0.57135f, > > + 1.0f, 1.8814f, 0.0f > > +}; > > + > > +__constant float yuv2rgb_bt709[] = { > > + 1.0f, 0.0f, 1.5748f, > > + 1.0f, -0.18732f, -0.46812f, > > + 1.0f, 1.8556f, 0.0f > > +}; > > + > > +__constant float rgb2yuv_bt709[] = { > > + 0.2126f, 0.7152f, 0.0722f, > > + -0.11457f, -0.38543f, 0.5f, > > + 0.5f, -0.45415f, -0.04585f > > +}; > > These don't look like the matrices I would expect for BT.709. Can you explain > where they come from? (I think I must be missing some subtlety here.) we can dump these matrices (rgb2yuv and yuv2rgb) in vf_colorspace.c > > > + > > +__constant float rgb2yuv_bt2020[] ={ > > + 0.2627f, 0.678f, 0.0593f, > > + -0.1396f, -0.36037f, 0.5f, > > + 0.5f, -0.4598f, -0.0402f, > > +}; > > + > > + > > +float get_luma_dst(float3 c) { > > + return luma_dst.x * c.x + luma_dst.y * c.y + luma_dst.z * c.z; > > +} > > + > > +float get_luma_src(float3 c) { > > + return luma_src.x * c.x + luma_src.y * c.y + luma_src.z * c.z; > > +} > > + > > +float eotf_st2084(float x) { > > + float p = powr(x, 1.0f / ST2084_M2); > > + float a = max(p -ST2084_C1, 0.0f); > > + float b = max(ST2084_C2 - ST2084_C3 * p, 1e-6f); > > + float c = powr(a / b, 1.0f / ST2084_M1); > > + return x > 0.0f ? c * ST2084_MAX_LUMINANCE / REFERENCE_WHITE : 0.0f; > > +} > > + > > +__constant const float HLG_A = 0.17883277f; > > +__constant const float HLG_B = 0.28466892f; > > +__constant const float HLG_C = 0.55991073f; > > + > > +// linearizer for HLG > > +float inverse_oetf_hlg(float x) { > > + float a = 4.0f * x * x; > > + float b = exp((x - HLG_C) / HLG_A) + HLG_B; > > + return x < 0.5f ? a : b; > > +} > > + > > +// delinearizer for HLG > > +float oetf_hlg(float x) { > > + float a = 0.5f * sqrt(x); > > + float b = HLG_A * log(x - HLG_B) + HLG_C; > > + return x <= 1.0f ? a : b; > > +} > > + > > +float3 ootf_hlg(float3 c, float peak) { > > + float luma = get_luma_src(c); > > + float gamma = 1.2f + 0.42f * log10(peak * REFERENCE_WHITE / 1000.0f); > > + gamma = max(1.0f, gamma); > > + float factor = peak * powr(luma, gamma - 1.0f) / powr(12.0f, gamma); > > + return c * factor; > > +} > > + > > +float3 inverse_ootf_hlg(float3 c, float peak) { > > + float gamma = 1.2f + 0.42f * log10(peak * REFERENCE_WHITE / 1000.0f); > > + c *= powr(12.0f, gamma) / peak; > > + c /= powr(get_luma_dst(c), (gamma - 1.0f) / gamma); > > + return c; > > +} > > + > > +float inverse_eotf_bt1886(float c) { > > + return c < 0.0f ? 0.0f : powr(c, 1.0f / 2.4f); > > +} > > + > > +float oetf_bt709(float c) { > > + c = c < 0.0f ? 0.0f : c; > > + float r1 = 4.5f * c; > > + float r2 = 1.099f * powr(c, 0.45f) - 0.099f; > > + return c < 0.018f ? r1 : r2; > > +} > > +float inverse_oetf_bt709(float c) { > > + float r1 = c / 4.5f; > > + float r2 = powr((c + 0.099f) / 1.099f, 1.0f / 0.45f); > > + return c < 0.081f ? r1 : r2; > > +} > > + > > +float3 yuv2rgb(float y, float u, float v) { > > +#ifdef FULL_RANGE_IN > > + u -= 0.5f; v -= 0.5f; > > +#else > > + y = (y * 255.0f - 16.0f) / 219.0f; > > + u = (u * 255.0f - 128.0f) / 224.0f; > > + v = (v * 255.0f - 128.0f) / 224.0f; > > +#endif > > + float r = y * rgb_matrix[0] + u * rgb_matrix[1] + v * rgb_matrix[2]; > > + float g = y * rgb_matrix[3] + u * rgb_matrix[4] + v * rgb_matrix[5]; > > + float b = y * rgb_matrix[6] + u * rgb_matrix[7] + v * rgb_matrix[8]; > > + return (float3)(r, g, b); > > +} > > + > > +float3 yuv2lrgb(float3 yuv) { > > + float3 rgb = yuv2rgb(yuv.x, yuv.y, yuv.z); > > + float r = linearize(rgb.x); > > + float g = linearize(rgb.y); > > + float b = linearize(rgb.z); > > + return (float3)(r, g, b); > > +} > > + > > +float3 rgb2yuv(float r, float g, float b) { > > + float y = r*yuv_matrix[0] + g*yuv_matrix[1] + b*yuv_matrix[2]; > > + float u = r*yuv_matrix[3] + g*yuv_matrix[4] + b*yuv_matrix[5]; > > + float v = r*yuv_matrix[6] + g*yuv_matrix[7] + b*yuv_matrix[8]; > > +#ifdef FULL_RANGE_OUT > > + u += 0.5f; v += 0.5f; > > +#else > > + y = (219.0f * y + 16.0f) / 255.0f; > > + u = (224.0f * u + 128.0f) / 255.0f; > > + v = (224.0f * v + 128.0f) / 255.0f; > > +#endif > > + return (float3)(y, u, v); > > +} > > + > > +float3 lrgb2yuv(float3 c) { > > + float r = delinearize(c.x); > > + float g = delinearize(c.y); > > + float b = delinearize(c.z); > > + > > + return rgb2yuv(r, g, b); > > +} > > + > > +float3 lrgb2lrgb(float3 c) { > > +#ifdef RGB2RGB_PASSTHROUGH > > + return c; > > +#else > > + float r = c.x, g = c.y, b = c.z; > > + float rr = rgb2rgb[0] * r + rgb2rgb[1] * g + rgb2rgb[2] * b; > > + float gg = rgb2rgb[3] * r + rgb2rgb[4] * g + rgb2rgb[5] * b; > > + float bb = rgb2rgb[6] * r + rgb2rgb[7] * g + rgb2rgb[8] * b; > > + return (float3)(rr, gg, bb); > > +#endif > > +} > > + > > +float3 ootf(float3 c, float peak) { > > +#ifdef ootf_impl > > + return ootf_impl(c, peak); > > +#else > > + return c; > > +#endif > > +} > > + > > +float3 inverse_ootf(float3 c, float peak) { > > +#ifdef inverse_ootf_impl > > + return inverse_ootf_impl(c, peak); > > +#else > > + return c; > > +#endif > > +} > > diff --git a/libavfilter/opencl/tonemap.cl b/libavfilter/opencl/tonemap.cl > > new file mode 100644 > > index 0000000..f88bc40 > > --- /dev/null > > +++ b/libavfilter/opencl/tonemap.cl > > @@ -0,0 +1,278 @@ > > +/* > > + * This file is part of FFmpeg. > > + * > > + * FFmpeg is free software; you can redistribute it and/or > > + * modify it under the terms of the GNU Lesser General Public > > + * License as published by the Free Software Foundation; either > > + * version 2.1 of the License, or (at your option) any later version. > > + * > > + * FFmpeg is distributed in the hope that it will be useful, > > + * but WITHOUT ANY WARRANTY; without even the implied warranty of > > + * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU > > + * Lesser General Public License for more details. > > + * > > + * You should have received a copy of the GNU Lesser General Public > > + * License along with FFmpeg; if not, write to the Free Software > > + * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 > USA > > + */ > > + > > +#define REFERENCE_WHITE 100.0f > > +extern float3 lrgb2yuv(float3); > > +extern float3 yuv2lrgb(float3); > > +extern float3 lrgb2lrgb(float3); > > +extern float get_luma_src(float3); > > +extern float get_luma_dst(float3); > > +extern float3 ootf(float3 c, float peak); > > +extern float3 inverse_ootf(float3 c, float peak); > > +struct detection_result { > > + float peak; > > + float average; > > +}; > > + > > +float hable_f(float in) { > > + float a = 0.15f, b = 0.50f, c = 0.10f, d = 0.20f, e = 0.02f, f = 0.30f; > > + return (in * (in * a + b * c) + d * e) / (in * (in * a + b) + d * f) - > > e / f; > > +} > > + > > +float direct(float s, float peak) { > > + return s; > > +} > > + > > +float linear(float s, float peak) { > > + return s * tone_param / peak; > > +} > > + > > +float gamma(float s, float peak) { > > + float p = s > 0.05f ? s /peak : 0.05f / peak; > > + float v = pow(p, 1.0f / tone_param); > > + return s > 0.05f ? v : (s * v /0.05f); > > +} > > + > > +float clip(float s, float peak) { > > + return clamp(s * tone_param, 0.0f, 1.0f); > > +} > > + > > +float reinhard(float s, float peak) { > > + return s / (s + tone_param) * (peak + tone_param) / peak; > > +} > > + > > +float hable(float s, float peak) { > > + return hable_f(s)/hable_f(peak); > > +} > > + > > +float mobius(float s, float peak) { > > + float j = tone_param; > > + float a, b; > > + > > + if (s <= j) > > + return s; > > + > > + a = -j * j * (peak - 1.0f) / (j * j - 2.0f * j + peak); > > + b = (j * j - 2.0f * j * peak + peak) / max(peak - 1.0f, 1e-6f); > > + > > + return (b * b + 2.0f * b * j + j * j) / (b - a) * (s + a) / (s + b); > > +} > > + > > +// detect peak/average signal of a frame, the algorithm was ported from: > > +// libplacebo (https://github.com/haasn/libplacebo) > > +struct detection_result > > +detect_peak_avg(global uint *util_buf, __local uint *sum_wg, > > + float signal, float peak) { > > +// layout of the util buffer > > +// > > +// Name: : Size (units of 4-bytes) > > +// average buffer : detection_frames + 1 > > +// peak buffer : detection_frames + 1 > > +// workgroup counter : 1 > > +// total of peak : 1 > > +// total of average : 1 > > +// frame index : 1 > > +// frame number : 1 > > + global uint *avg_buf = util_buf; > > + global uint *peak_buf = avg_buf + DETECTION_FRAMES + 1; > > + global uint *counter_wg_p = peak_buf + DETECTION_FRAMES + 1; > > + global uint *max_total_p = counter_wg_p + 1; > > + global uint *avg_total_p = max_total_p + 1; > > + global uint *frame_idx_p = avg_total_p + 1; > > + global uint *scene_frame_num_p = frame_idx_p + 1; > > + > > + uint frame_idx = *frame_idx_p; > > + uint scene_frame_num = *scene_frame_num_p; > > + > > + size_t lidx = get_local_id(0); > > + size_t lidy = get_local_id(1); > > + size_t lsizex = get_local_size(0); > > + size_t lsizey = get_local_size(1); > > + uint num_wg = get_num_groups(0) * get_num_groups(1); > > + size_t group_idx = get_group_id(0); > > + size_t group_idy = get_group_id(1); > > + struct detection_result r = {peak, sdr_avg}; > > + if (lidx == 0 && lidy == 0) > > + *sum_wg = 0; > > + barrier(CLK_LOCAL_MEM_FENCE); > > + > > + // update workgroup sum > > + atomic_add(sum_wg, (uint)(signal * REFERENCE_WHITE)); > > + barrier(CLK_LOCAL_MEM_FENCE); > > + > > + // update frame peak/avg using work-group-average. > > + if (lidx == 0 && lidy == 0) { > > + uint avg_wg = *sum_wg / (lsizex * lsizey); > > + atomic_max(&peak_buf[frame_idx], avg_wg); > > + atomic_add(&avg_buf[frame_idx], avg_wg); > > + } > > + > > + if (scene_frame_num > 0) { > > + float peak = (float)*max_total_p / (REFERENCE_WHITE * > scene_frame_num); > > + float avg = (float)*avg_total_p / (REFERENCE_WHITE * > scene_frame_num); > > + r.peak = max(1.0f, peak); > > + r.average = max(0.25f, avg); > > + } > > + > > + if (lidx == 0 && lidy == 0 && atomic_add(counter_wg_p, 1) == num_wg - > > 1) > { > > + *counter_wg_p = 0; > > + avg_buf[frame_idx] /= num_wg; > > + > > + if (scene_threshold > 0.0f) { > > + uint cur_max = peak_buf[frame_idx]; > > + uint cur_avg = avg_buf[frame_idx]; > > + int diff = (int)(scene_frame_num * cur_avg) - > > (int)*avg_total_p; > > + > > + if (abs(diff) > scene_frame_num * scene_threshold * > REFERENCE_WHITE) { > > + for (uint i = 0; i < DETECTION_FRAMES + 1; i++) > > + avg_buf[i] = 0; > > + for (uint i = 0; i < DETECTION_FRAMES + 1; i++) > > + peak_buf[i] = 0; > > + *avg_total_p = *max_total_p = 0; > > + *scene_frame_num_p = 0; > > + avg_buf[frame_idx] = cur_avg; > > + peak_buf[frame_idx] = cur_max; > > + } > > + } > > + uint next = (frame_idx + 1) % (DETECTION_FRAMES + 1); > > + // add current frame, subtract next frame > > + *max_total_p += peak_buf[frame_idx] - peak_buf[next]; > > + *avg_total_p += avg_buf[frame_idx] - avg_buf[next]; > > + // reset next frame > > + peak_buf[next] = avg_buf[next] = 0; > > + *frame_idx_p = next; > > + *scene_frame_num_p = min(*scene_frame_num_p + 1, > (uint)DETECTION_FRAMES); > > + } > > + return r; > > +} > > + > > +float3 map_one_pixel_rgb(float3 rgb, float peak, float average) { > > + float sig = max(max(rgb.x, max(rgb.y, rgb.z)), 1e-6f); > > + > > + // Rescale the variables in order to bring it into a representation > > where > > + // 1.0 represents the dst_peak. This is because all of the tone mapping > > + // algorithms are defined in such a way that they map to the range > > [0.0, > 1.0]. > > + if (target_peak > 1.0f) { > > + sig *= 1.0f / target_peak; > > + peak *= 1.0f / target_peak; > > + } > > + > > + float sig_old = sig; > > + > > + // Scale the signal to compensate for differences in the average > > brightness > > + float slope = min(1.0f, sdr_avg / average); > > + sig *= slope; > > + peak *= slope; > > + > > + // Desaturate the color using a coefficient dependent on the signal > > level > > + if (desat_param > 0.0f) { > > + float luma = get_luma_dst(rgb); > > + float coeff = max(sig - 0.18f, 1e-6f) / max(sig, 1e-6f); > > + coeff = native_powr(coeff, 10.0f / desat_param); > > + rgb = mix(rgb, (float3)luma, (float3)coeff); > > + sig = mix(sig, luma * slope, coeff); > > + } > > + > > + sig = TONE_FUNC(sig, peak); > > + > > + sig = min(sig, 1.0f); > > + rgb *= (sig/sig_old); > > + return rgb; > > +} > > +// map from source space YUV to destination space RGB > > +float3 map_to_dst_space_from_yuv(float3 yuv, float peak) { > > + float3 c = yuv2lrgb(yuv); > > + c = ootf(c, peak); > > + c = lrgb2lrgb(c); > > + return c; > > +} > > + > > +// convert from rgb to yuv, with possible inverse-ootf > > +float3 convert_to_yuv(float3 c, float peak) { > > + c = inverse_ootf(c, peak); > > + return lrgb2yuv(c); > > +} > > + > > +__kernel void tonemap(__write_only image2d_t dst1, > > + __write_only image2d_t dst2, > > + __read_only image2d_t src1, > > + __read_only image2d_t src2, > > +#ifdef THIRD_PLANE > > + __write_only image2d_t dst3, > > + __read_only image2d_t src3, > > +#endif > > THIRD_PLANE isn't currently set anywhere. I think either make it work (add > the > pixel formats) or remove the unused code. Will fix it. > > > + global uint *util_buf, > > + float peak > > + ) > > +{ > > + __local uint sum_wg; > > + const sampler_t sampler = (CLK_NORMALIZED_COORDS_FALSE | > > + CLK_FILTER_NEAREST); > > + int xi = get_global_id(0); > > + int yi = get_global_id(1); > > + // each work item process four pixels > > + int x = 2 * xi; > > + int y = 2 * yi; > > + > > + float y0 = read_imagef(src1, sampler, (int2)(x, y)).x; > > + float y1 = read_imagef(src1, sampler, (int2)(x + 1, y)).x; > > + float y2 = read_imagef(src1, sampler, (int2)(x, y + 1)).x; > > + float y3 = read_imagef(src1, sampler, (int2)(x + 1, y + 1)).x; > > +#ifdef THIRD_PLANE > > + float u = read_imagef(src2, sampler, (int2)(xi, yi)).x; > > + float v = read_imagef(src3, sampler, (int2)(xi, yi)).x; > > + float2 uv = (float2)(u, v); > > +#else > > + float2 uv = read_imagef(src2, sampler, (int2)(xi, yi)).xy; > > +#endif > > + > > + float3 c0 = map_to_dst_space_from_yuv((float3)(y0, uv.x, uv.y), peak); > > + float3 c1 = map_to_dst_space_from_yuv((float3)(y1, uv.x, uv.y), peak); > > + float3 c2 = map_to_dst_space_from_yuv((float3)(y2, uv.x, uv.y), peak); > > + float3 c3 = map_to_dst_space_from_yuv((float3)(y3, uv.x, uv.y), peak); > > + > > + float sig0 = max(c0.x, max(c0.y, c0.z)); > > + float sig1 = max(c1.x, max(c1.y, c1.z)); > > + float sig2 = max(c2.x, max(c2.y, c2.z)); > > + float sig3 = max(c3.x, max(c3.y, c3.z)); > > + float sig = max(sig0, max(sig1, max(sig2, sig3))); > > + > > + struct detection_result r = detect_peak_avg(util_buf, &sum_wg, sig, > > peak); > > + > > + float3 c0_old = c0, c1_old = c1, c2_old = c2; > > + c0 = map_one_pixel_rgb(c0, r.peak, r.average); > > + c1 = map_one_pixel_rgb(c1, r.peak, r.average); > > + c2 = map_one_pixel_rgb(c2, r.peak, r.average); > > + c3 = map_one_pixel_rgb(c3, r.peak, r.average); > > + > > + float3 yuv0 = convert_to_yuv(c0, target_peak); > > + float3 yuv1 = convert_to_yuv(c1, target_peak); > > + float3 yuv2 = convert_to_yuv(c2, target_peak); > > + float3 yuv3 = convert_to_yuv(c3, target_peak); > > + > > + write_imagef(dst1, (int2)(x, y), (float4)(yuv0.x, 0.0f, 0.0f, 1.0f)); > > + write_imagef(dst1, (int2)(x+1, y), (float4)(yuv1.x, 0.0f, 0.0f, 1.0f)); > > + write_imagef(dst1, (int2)(x, y+1), (float4)(yuv2.x, 0.0f, 0.0f, 1.0f)); > > + write_imagef(dst1, (int2)(x+1, y+1), (float4)(yuv3.x, 0.0f, 0.0f, > > 1.0f)); > > +#ifdef THIRD_PLANE > > + write_imagef(dst2, (int2)(xi, yi), (float4)(yuv0.y, 0.0f, 0.0f, 1.0f)); > > + write_imagef(dst3, (int2)(xi, yi), (float4)(yuv0.z, 0.0f, 0.0f, 1.0f)); > > +#else > > + write_imagef(dst2, (int2)(xi, yi), (float4)(yuv0.y, yuv0.z, 0.0f, > > 1.0f)); > > +#endif > > +} > > I'm not sure if it makes any significant difference, but should the chroma > sampling location be taken into account here? This is reading as if a centre > value and writing as if top-left, but most things will probably be > centre-left. Sorry I didn't noticed this. will try to fix it. Seems it only applies to writing YUV? > > > diff --git a/libavfilter/opencl_source.h b/libavfilter/opencl_source.h > > index 4bb9969..c5b3f37 100644 > > --- a/libavfilter/opencl_source.h > > +++ b/libavfilter/opencl_source.h > > @@ -21,7 +21,9 @@ > > > > extern const char *ff_opencl_source_avgblur; > > extern const char *ff_opencl_source_convolution; > > +extern const char *ff_opencl_source_colorspace_basic; > > extern const char *ff_opencl_source_overlay; > > +extern const char *ff_opencl_source_tonemap; > > extern const char *ff_opencl_source_unsharp; > > > > #endif /* AVFILTER_OPENCL_SOURCE_H */ > > diff --git a/libavfilter/vf_tonemap_opencl.c > > b/libavfilter/vf_tonemap_opencl.c > > new file mode 100644 > > index 0000000..1a5bb20 > > --- /dev/null > > +++ b/libavfilter/vf_tonemap_opencl.c > > @@ -0,0 +1,655 @@ > > +/* > > + * This file is part of FFmpeg. > > + * > > + * FFmpeg is free software; you can redistribute it and/or > > + * modify it under the terms of the GNU Lesser General Public > > + * License as published by the Free Software Foundation; either > > + * version 2.1 of the License, or (at your option) any later version. > > + * > > + * FFmpeg is distributed in the hope that it will be useful, > > + * but WITHOUT ANY WARRANTY; without even the implied warranty of > > + * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU > > + * Lesser General Public License for more details. > > + * > > + * You should have received a copy of the GNU Lesser General Public > > + * License along with FFmpeg; if not, write to the Free Software > > + * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 > USA > > + */ > > +#include <float.h> > > + > > +#include "libavutil/avassert.h" > > +#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" > > + > > +#include "avfilter.h" > > +#include "internal.h" > > +#include "opencl.h" > > +#include "opencl_source.h" > > +#include "video.h" > > +#include "colorspace_basic.h" > > + > > +//#define DEBUG > > +// TODO: > > +// - seperate peak-detection from tone-mapping kernel to solve > > +// one-frame-delay issue. > > +// - import colorspace matrix generation from vf_colorspace.c > > +// - more format support > > + > > +#define DETECTION_FRAMES 63 > > +#define REFERENCE_WHITE 100.0f > > + > > +enum TonemapAlgorithm { > > + TONEMAP_NONE, > > + TONEMAP_LINEAR, > > + TONEMAP_GAMMA, > > + TONEMAP_CLIP, > > + TONEMAP_REINHARD, > > + TONEMAP_HABLE, > > + TONEMAP_MOBIUS, > > + TONEMAP_MAX, > > +}; > > + > > +typedef struct TonemapOpenCLContext { > > + OpenCLFilterContext ocf; > > + > > + enum AVColorSpace colorspace, colorspace_in, colorspace_out; > > + enum AVColorTransferCharacteristic trc, trc_in, trc_out; > > + enum AVColorPrimaries primaries, primaries_in, primaries_out; > > + enum AVColorRange range, range_in, range_out; > > + > > + enum TonemapAlgorithm tonemap; > > + enum AVPixelFormat format; > > + double peak; > > + double param; > > + double desat_param; > > + double target_peak; > > + double scene_threshold; > > + int initialised; > > + cl_kernel kernel; > > + cl_command_queue command_queue; > > + cl_mem util_mem; > > +} TonemapOpenCLContext; > > + > > +const char *yuv_coff[AVCOL_SPC_NB] = { > > + [AVCOL_SPC_BT709] = "rgb2yuv_bt709", > > + [AVCOL_SPC_BT2020_NCL] = "rgb2yuv_bt2020", > > +}; > > + > > +const char *rgb_coff[AVCOL_SPC_NB] = { > > + [AVCOL_SPC_BT709] = "yuv2rgb_bt709", > > + [AVCOL_SPC_BT2020_NCL] = "yuv2rgb_bt2020", > > +}; > > + > > +const char *linearize_funcs[AVCOL_TRC_NB] = { > > + [AVCOL_TRC_SMPTE2084] = "eotf_st2084", > > + [AVCOL_TRC_ARIB_STD_B67] = "inverse_oetf_hlg", > > +}; > > + > > +const char *delinearize_funcs[AVCOL_TRC_NB] = { > > + [AVCOL_TRC_BT709] = "inverse_eotf_bt1886", > > + [AVCOL_TRC_BT2020_10] = "inverse_eotf_bt1886", > > +}; > > + > > +static const struct LumaCoefficients luma_coefficients[AVCOL_SPC_NB] = { > > + [AVCOL_SPC_BT709] = { 0.2126, 0.7152, 0.0722 }, > > + [AVCOL_SPC_BT2020_NCL] = { 0.2627, 0.6780, 0.0593 }, > > +}; > > + > > +struct ColorPrimaries primaries_table[AVCOL_PRI_NB] = { > > + [AVCOL_PRI_BT709] = { 0.640, 0.330, 0.300, 0.600, 0.150, 0.060 }, > > + [AVCOL_PRI_BT2020] = { 0.708, 0.292, 0.170, 0.797, 0.131, 0.046 }, > > +}; > > + > > +struct WhitePoint whitepoint_table[AVCOL_PRI_NB] = { > > + [AVCOL_PRI_BT709] = { 0.3127, 0.3290 }, > > + [AVCOL_PRI_BT2020] = { 0.3127, 0.3290 }, > > +}; > > + > > +const char *tonemap_func[TONEMAP_MAX] = { > > + [TONEMAP_NONE] = "direct", > > + [TONEMAP_LINEAR] = "linear", > > + [TONEMAP_GAMMA] = "gamma", > > + [TONEMAP_CLIP] = "clip", > > + [TONEMAP_REINHARD] = "reinhard", > > + [TONEMAP_HABLE] = "hable", > > + [TONEMAP_MOBIUS] = "mobius", > > +}; > > + > > +static void get_rgb2rgb_matrix(enum AVColorPrimaries in, enum > AVColorPrimaries out, > > + 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); > > +} > > + > > +#define OPENCL_SOURCE_NB 3 > > +// Average light level for SDR signals. This is equal to a signal level of > > 0.5 > > +// under a typical presentation gamma of about 2.0. > > +static const float sdr_avg = 0.25f; > > + > > +static int tonemap_opencl_init(AVFilterContext *avctx) > > +{ > > + TonemapOpenCLContext *ctx = avctx->priv; > > + int rgb2rgb_passthrough = 1; > > + double rgb2rgb[3][3]; > > + struct LumaCoefficients luma_src, luma_dst; > > + cl_int cle; > > + int err; > > + AVBPrint header; > > + const char *opencl_sources[OPENCL_SOURCE_NB]; > > + > > + av_bprint_init(&header, 1024, AV_BPRINT_SIZE_AUTOMATIC); > > + > > + switch(ctx->tonemap) { > > + case TONEMAP_GAMMA: > > + if (isnan(ctx->param)) > > + ctx->param = 1.8f; > > + break; > > + case TONEMAP_REINHARD: > > + if (!isnan(ctx->param)) > > + ctx->param = (1.0f - ctx->param) / ctx->param; > > + break; > > + case TONEMAP_MOBIUS: > > + if (isnan(ctx->param)) > > + ctx->param = 0.3f; > > + break; > > + } > > + > > + if (isnan(ctx->param)) > > + ctx->param = 1.0f; > > + > > + // SDR peak is 1.0f > > + ctx->target_peak = 1.0f; > > + av_log(ctx, AV_LOG_DEBUG, "tone mapping transfer from %s to %s\n", > > + av_color_transfer_name(ctx->trc_in), > > + av_color_transfer_name(ctx->trc_out)); > > + av_log(ctx, AV_LOG_DEBUG, "mapping colorspace from %s to %s\n", > > + av_color_space_name(ctx->colorspace_in), > > + av_color_space_name(ctx->colorspace_out)); > > + av_log(ctx, AV_LOG_DEBUG, "mapping primaries from %s to %s\n", > > + av_color_primaries_name(ctx->primaries_in), > > + av_color_primaries_name(ctx->primaries_out)); > > + av_log(ctx, AV_LOG_DEBUG, "mapping range from %s to %s\n", > > + av_color_range_name(ctx->range_in), > > + av_color_range_name(ctx->range_out)); > > + // checking valid value just because of limited implementaion > > + // please remove when more functionalities are implemented > > + av_assert0(ctx->trc_out == AVCOL_TRC_BT709 || > > + ctx->trc_out == AVCOL_TRC_BT2020_10); > > + av_assert0(ctx->trc_in == AVCOL_TRC_SMPTE2084|| > > + ctx->trc_in == AVCOL_TRC_ARIB_STD_B67); > > + av_assert0(ctx->colorspace_in == AVCOL_SPC_BT2020_NCL || > > + ctx->colorspace_in == AVCOL_SPC_BT709); > > + av_assert0(ctx->primaries_in == AVCOL_PRI_BT2020 || > > + ctx->primaries_in == AVCOL_PRI_BT709); > > + > > + av_bprintf(&header, "__constant const float tone_param = %.4ff;\n", > > + ctx->param); > > + av_bprintf(&header, "__constant const float desat_param = %.4ff;\n", > > + ctx->desat_param); > > + av_bprintf(&header, "__constant const float target_peak = %.4ff;\n", > > + ctx->target_peak); > > + av_bprintf(&header, "__constant const float sdr_avg = %.4ff;\n", > > sdr_avg); > > + av_bprintf(&header, "__constant const float scene_threshold = > > %.4ff;\n", > > + ctx->scene_threshold); > > + av_bprintf(&header, "#define TONE_FUNC %s\n", tonemap_func[ctx- > >tonemap]); > > + av_bprintf(&header, "#define DETECTION_FRAMES %d\n", > DETECTION_FRAMES); > > + > > + if (ctx->primaries_out != ctx->primaries_in) { > > + get_rgb2rgb_matrix(ctx->primaries_in, ctx->primaries_out, rgb2rgb); > > + rgb2rgb_passthrough = 0; > > + } > > + if (ctx->range_in == AVCOL_RANGE_JPEG) > > + av_bprintf(&header, "#define FULL_RANGE_IN\n"); > > + > > + if (ctx->range_out == AVCOL_RANGE_JPEG) > > + av_bprintf(&header, "#define FULL_RANGE_OUT\n"); > > + > > + if (rgb2rgb_passthrough) > > + av_bprintf(&header, "#define RGB2RGB_PASSTHROUGH\n"); > > + else { > > + av_bprintf(&header, "__constant float rgb2rgb[9] = {\n"); > > + av_bprintf(&header, " %.4ff, %.4ff, %.4ff,\n", > > + rgb2rgb[0][0], rgb2rgb[0][1], rgb2rgb[0][2]); > > + av_bprintf(&header, " %.4ff, %.4ff, %.4ff,\n", > > + rgb2rgb[1][0], rgb2rgb[1][1], rgb2rgb[1][2]); > > + av_bprintf(&header, " %.4ff, %.4ff, %.4ff};\n", > > + rgb2rgb[2][0], rgb2rgb[2][1], rgb2rgb[2][2]); > > + } > > + > > + av_bprintf(&header, "#define rgb_matrix %s\n", > > + rgb_coff[ctx->colorspace_in]); > > + av_bprintf(&header, "#define yuv_matrix %s\n", > > + yuv_coff[ctx->colorspace_out]); > > + > > + luma_src = luma_coefficients[ctx->colorspace_in]; > > + luma_dst = luma_coefficients[ctx->colorspace_out]; > > + av_bprintf(&header, "constant float3 luma_src = {%.4ff, %.4ff, > > %.4ff};\n", > > + luma_src.cr, luma_src.cg, luma_src.cb); > > + av_bprintf(&header, "constant float3 luma_dst = {%.4ff, %.4ff, > > %.4ff};\n", > > + luma_dst.cr, luma_dst.cg, luma_dst.cb); > > + > > + av_bprintf(&header, "#define linearize %s\n", > > linearize_funcs[ctx->trc_in]); > > + av_bprintf(&header, "#define delinearize %s\n", > > + delinearize_funcs[ctx->trc_out]); > > + > > + if (ctx->trc_in == AVCOL_TRC_ARIB_STD_B67) > > + av_bprintf(&header, "#define ootf_impl ootf_hlg\n"); > > + > > + if (ctx->trc_out == AVCOL_TRC_ARIB_STD_B67) > > + av_bprintf(&header, "#define inverse_ootf_impl > > inverse_ootf_hlg\n"); > > + > > + av_log(avctx, AV_LOG_DEBUG, "Generated OpenCL header:\n%s\n", > header.str); > > + opencl_sources[0] = header.str; > > + opencl_sources[1] = ff_opencl_source_tonemap; > > + opencl_sources[2] = ff_opencl_source_colorspace_basic; > > + err = ff_opencl_filter_load_program(avctx, opencl_sources, > OPENCL_SOURCE_NB); > > + > > + av_bprint_finalize(&header, NULL); > > + if (err < 0) > > + goto fail; > > + > > + 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; > > + } > > + > > + 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; > > + } > > + > > + 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; > > + } > > + > > + ctx->initialised = 1; > > + return 0; > > + > > +fail: > > + if (ctx->util_mem) > > + clReleaseMemObject(ctx->util_mem); > > + if (ctx->command_queue) > > + clReleaseCommandQueue(ctx->command_queue); > > + if (ctx->kernel) > > + clReleaseKernel(ctx->kernel); > > + return err; > > +} > > + > > +static int tonemap_opencl_config_output(AVFilterLink *outlink) > > +{ > > + AVFilterContext *avctx = outlink->src; > > + TonemapOpenCLContext *s = avctx->priv; > > + int ret; > > + if (s->format == AV_PIX_FMT_NONE) > > + av_log(avctx, AV_LOG_WARNING, "format not set, use default format > NV12\n"); > > + else { > > + if (s->format != AV_PIX_FMT_P010 && > > + s->format != AV_PIX_FMT_NV12) { > > + av_log(avctx, AV_LOG_ERROR, "unsupported output format," > > + "only p010/nv12 supported now\n"); > > + return AVERROR(EINVAL); > > + } > > + } > > + > > + s->ocf.output_format = s->format == AV_PIX_FMT_NONE ? > AV_PIX_FMT_NV12 : s->format; > > + ret = ff_opencl_filter_config_output(outlink); > > + if (ret < 0) > > + return ret; > > + > > + return 0; > > +} > > + > > +static int launch_kernel(AVFilterContext *avctx, cl_kernel kernel, > > + AVFrame *output, AVFrame *input, float peak) { > > + TonemapOpenCLContext *ctx = avctx->priv; > > + int err = AVERROR(ENOSYS); > > + size_t global_work[2]; > > + size_t local_work[2]; > > + cl_int cle; > > + > > + cle = clSetKernelArg(kernel, 0, sizeof(cl_mem), &output->data[0]); > > + if (cle != CL_SUCCESS) { > > + av_log(avctx, AV_LOG_ERROR, "Failed to set kernel " > > + "destination image 1st plane: %d.\n", cle); > > + return AVERROR(EINVAL); > > + } > > + > > + cle = clSetKernelArg(kernel, 1, sizeof(cl_mem), &output->data[1]); > > + if (cle != CL_SUCCESS) { > > + av_log(avctx, AV_LOG_ERROR, "Failed to set kernel " > > + "destination image 2nd plane: %d.\n", cle); > > + return AVERROR(EINVAL); > > + } > > + > > + cle = clSetKernelArg(kernel, 2, sizeof(cl_mem), &input->data[0]); > > + if (cle != CL_SUCCESS) { > > + av_log(avctx, AV_LOG_ERROR, "Failed to set kernel " > > + "source image 1st plane: %d.\n", cle); > > + return AVERROR(EINVAL); > > + } > > + > > + cle = clSetKernelArg(kernel, 3, sizeof(cl_mem), &input->data[1]); > > + if (cle != CL_SUCCESS) { > > + av_log(avctx, AV_LOG_ERROR, "Failed to set kernel " > > + "source image 2nd plane: %d.\n", cle); > > + return AVERROR(EINVAL); > > + } > > + > > + cle = clSetKernelArg(kernel, 4, sizeof(cl_mem), &ctx->util_mem); > > + if (cle != CL_SUCCESS) { > > + av_log(avctx, AV_LOG_ERROR, "Failed to set kernel " > > + "source image 2nd plane: %d.\n", cle); > > + return AVERROR(EINVAL); > > + } > > + > > + cle = clSetKernelArg(kernel, 5, sizeof(cl_float), &peak); > > + if (cle != CL_SUCCESS) { > > + av_log(avctx, AV_LOG_ERROR, "Failed to set kernel " > > + "peak luma: %d.\n", cle); > > + return AVERROR(EINVAL); > > + } > > + > > + local_work[0] = 16; > > + local_work[1] = 16; > > + // Note the work size based on uv plane, as we process a 2x2 quad in > > one > workitem > > + err = ff_opencl_filter_work_size_from_image(avctx, global_work, output, > > + 1, 16); > > + if (err < 0) > > + return err; > > + > > + 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); > > + } > > + return 0; > > +} > > + > > +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; > > + AVFilterLink *outlink = avctx->outputs[0]; > > + TonemapOpenCLContext *ctx = avctx->priv; > > + AVFrame *output = NULL; > > + cl_int cle; > > + int err; > > + double peak = ctx->peak; > > + > > + AVHWFramesContext *input_frames_ctx = > > + (AVHWFramesContext*)input->hw_frames_ctx->data; > > + > > + av_log(ctx, AV_LOG_DEBUG, "Filter input: %s, %ux%u (%"PRId64").\n", > > + av_get_pix_fmt_name(input->format), > > + input->width, input->height, input->pts); > > + > > + if (!input->hw_frames_ctx) > > + return AVERROR(EINVAL); > > + > > + output = ff_get_video_buffer(outlink, outlink->w, outlink->h); > > + if (!output) { > > + err = AVERROR(ENOMEM); > > + goto fail; > > + } > > + > > + err = av_frame_copy_props(output, input); > > + if (err < 0) > > + goto fail; > > + > > + if (!peak) > > + peak = determine_signal_peak(input); > > + > > + if (ctx->trc != -1) > > + output->color_trc = ctx->trc; > > + if (ctx->primaries != -1) > > + output->color_primaries = ctx->primaries; > > + if (ctx->colorspace != -1) > > + output->colorspace = ctx->colorspace; > > + if (ctx->range != -1) > > + output->color_range = ctx->range; > > + > > + ctx->trc_in = input->color_trc; > > + ctx->trc_out = output->color_trc; > > + ctx->colorspace_in = input->colorspace; > > + ctx->colorspace_out = output->colorspace; > > + ctx->primaries_in = input->color_primaries; > > + ctx->primaries_out = output->color_primaries; > > + ctx->range_in = input->color_range; > > + ctx->range_out = output->color_range; > > + > > + if (!ctx->initialised) { > > + if (!(input->color_trc == AVCOL_TRC_SMPTE2084 || > > + input->color_trc == AVCOL_TRC_ARIB_STD_B67)) { > > + av_log(ctx, AV_LOG_ERROR, "unsupported transfer function > characteristic.\n"); > > + err = AVERROR(ENOSYS); > > + goto fail; > > + } > > + > > + if (input_frames_ctx->sw_format != AV_PIX_FMT_P010) { > > + av_log(ctx, AV_LOG_ERROR, "unsupported format in > tonemap_opencl.\n"); > > + err = AVERROR(ENOSYS); > > + goto fail; > > + } > > + > > + err = tonemap_opencl_init(avctx); > > + if (err < 0) > > + goto fail; > > + } > > + > > + switch(input_frames_ctx->sw_format) { > > + case AV_PIX_FMT_P010: > > + err = launch_kernel(avctx, ctx->kernel, output, input, peak); > > + if (err < 0) goto fail; > > + break; > > + default: > > + err = AVERROR(ENOSYS); > > + goto fail; > > + } > > + > > + 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; > > + } > > + > > + av_frame_free(&input); > > + > > + update_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), > > + output->width, output->height, output->pts); > > +#ifdef DEBUG > > I like the inclusion of this code, but I don't think it's a good idea to have > blocks > which are never built without source editing. Put it behind an option > instead? I noticed that I can use "#ifndef NDEBUG". It think it is more proper. > > > + { > > + uint32_t *ptr, *max_total_p, *avg_total_p, *frame_number_p; > > + float peak_detected, avg_detected; > > + > > + unsigned map_size = (2 * DETECTION_FRAMES + 7) * sizeof(unsigned); > > + ptr = (void *)clEnqueueMapBuffer(ctx->command_queue, ctx->util_mem, > > + CL_TRUE, CL_MAP_READ, 0, map_size, > > + 0, NULL, NULL, &cle); > > + // For the layout of the util buffer, refer tonemap.cl > > + if (ptr) { > > + max_total_p = ptr + 2 * (DETECTION_FRAMES + 1) + 1; > > + avg_total_p = max_total_p + 1; > > + frame_number_p = avg_total_p + 2; > > + peak_detected = (float)*max_total_p / (REFERENCE_WHITE * > (*frame_number_p)); > > + avg_detected = (float)*avg_total_p / (REFERENCE_WHITE * > (*frame_number_p)); > > + av_log(ctx, AV_LOG_DEBUG, "peak %f, avg %f will be used for > > next > frame\n", > > + peak_detected, avg_detected); > > + clEnqueueUnmapMemObject(ctx->command_queue, ctx->util_mem, > ptr, 0, > > + NULL, NULL); > > + } > > + } > > +#endif > > + > > + return ff_filter_frame(outlink, output); > > + > > +fail: > > + clFinish(ctx->command_queue); > > + av_frame_free(&input); > > + av_frame_free(&output); > > + return err; > > +} > > + > > +static av_cold void tonemap_opencl_uninit(AVFilterContext *avctx) > > +{ > > + TonemapOpenCLContext *ctx = avctx->priv; > > + cl_int cle; > > + > > + if (ctx->util_mem) > > + clReleaseMemObject(ctx->util_mem); > > + if (ctx->kernel) { > > + cle = clReleaseKernel(ctx->kernel); > > + if (cle != CL_SUCCESS) > > + av_log(avctx, AV_LOG_ERROR, "Failed to release " > > + "kernel: %d.\n", cle); > > + } > > + > > + if (ctx->command_queue) { > > + cle = clReleaseCommandQueue(ctx->command_queue); > > + if (cle != CL_SUCCESS) > > + av_log(avctx, AV_LOG_ERROR, "Failed to release " > > + "command queue: %d.\n", cle); > > + } > > + > > + ff_opencl_filter_uninit(avctx); > > +}9 > > + > > +#define OFFSET(x) offsetof(TonemapOpenCLContext, x) > > +#define FLAGS (AV_OPT_FLAG_FILTERING_PARAM | > AV_OPT_FLAG_VIDEO_PARAM) > > +static const AVOption tonemap_opencl_options[] = { > > + { "tonemap", "tonemap algorithm selection", OFFSET(tonemap), > AV_OPT_TYPE_INT, {.i64 = TONEMAP_NONE}, TONEMAP_NONE, > TONEMAP_MAX - 1, FLAGS, "tonemap" }, > > + { "none", 0, 0, AV_OPT_TYPE_CONST, {.i64 = TONEMAP_NONE}, > 0, 0, FLAGS, "tonemap" }, > > + { "linear", 0, 0, AV_OPT_TYPE_CONST, {.i64 = TONEMAP_LINEAR}, > 0, 0, FLAGS, "tonemap" }, > > + { "gamma", 0, 0, AV_OPT_TYPE_CONST, {.i64 = TONEMAP_GAMMA}, > 0, 0, FLAGS, "tonemap" }, > > + { "clip", 0, 0, AV_OPT_TYPE_CONST, {.i64 = TONEMAP_CLIP}, > > 0, > 0, FLAGS, "tonemap" }, > > + { "reinhard", 0, 0, AV_OPT_TYPE_CONST, {.i64 = TONEMAP_REINHARD}, > 0, 0, FLAGS, "tonemap" }, > > + { "hable", 0, 0, AV_OPT_TYPE_CONST, {.i64 = TONEMAP_HABLE}, > 0, 0, FLAGS, "tonemap" }, > > + { "mobius", 0, 0, AV_OPT_TYPE_CONST, {.i64 = TONEMAP_MOBIUS}, > 0, 0, FLAGS, "tonemap" }, > > + { "transfer", "set transfer characteristic", OFFSET(trc), > > AV_OPT_TYPE_INT, > {.i64 = AVCOL_TRC_BT709}, -1, INT_MAX, FLAGS, "transfer" }, > > + { "t", "set transfer characteristic", OFFSET(trc), > > AV_OPT_TYPE_INT, > {.i64 = AVCOL_TRC_BT709}, -1, INT_MAX, FLAGS, "transfer" }, > > + { "bt709", 0, 0, > > AV_OPT_TYPE_CONST, {.i64 = > AVCOL_TRC_BT709}, 0, 0, FLAGS, "transfer" }, > > + { "bt2020", 0, 0, > > AV_OPT_TYPE_CONST, {.i64 = > AVCOL_TRC_BT2020_10}, 0, 0, FLAGS, "transfer" }, > > + { "matrix", "set colorspace matrix", OFFSET(colorspace), > AV_OPT_TYPE_INT, {.i64 = -1}, -1, INT_MAX, FLAGS, "matrix" }, > > + { "m", "set colorspace matrix", OFFSET(colorspace), > > AV_OPT_TYPE_INT, > {.i64 = -1}, -1, INT_MAX, FLAGS, "matrix" }, > > + { "bt709", 0, 0, > > AV_OPT_TYPE_CONST, {.i64 = > AVCOL_SPC_BT709}, 0, 0, FLAGS, "matrix" }, > > + { "bt2020", 0, 0, > > AV_OPT_TYPE_CONST, {.i64 = > AVCOL_SPC_BT2020_NCL}, 0, 0, FLAGS, "matrix" }, > > + { "primaries", "set color primaries", OFFSET(primaries), > > AV_OPT_TYPE_INT, > {.i64 = -1}, -1, INT_MAX, FLAGS, "primaries" }, > > + { "p", "set color primaries", OFFSET(primaries), > > AV_OPT_TYPE_INT, > {.i64 = -1}, -1, INT_MAX, FLAGS, "primaries" }, > > + { "bt709", 0, 0, > > AV_OPT_TYPE_CONST, {.i64 = > AVCOL_PRI_BT709}, 0, 0, FLAGS, "primaries" }, > > + { "bt2020", 0, 0, > > AV_OPT_TYPE_CONST, {.i64 = > AVCOL_PRI_BT2020}, 0, 0, FLAGS, "primaries" }, > > + { "range", "set color range", OFFSET(range), AV_OPT_TYPE_INT, > > {.i64 = > -1}, -1, INT_MAX, FLAGS, "range" }, > > + { "r", "set color range", OFFSET(range), AV_OPT_TYPE_INT, > > {.i64 = - > 1}, -1, INT_MAX, FLAGS, "range" }, > > + { "tv", 0, 0, AV_OPT_TYPE_CONST, > > {.i64 = > AVCOL_RANGE_MPEG}, 0, 0, FLAGS, "range" }, > > + { "pc", 0, 0, AV_OPT_TYPE_CONST, > > {.i64 = > AVCOL_RANGE_JPEG}, 0, 0, FLAGS, "range" }, > > + { "limited", 0, 0, AV_OPT_TYPE_CONST, > > {.i64 = > AVCOL_RANGE_MPEG}, 0, 0, FLAGS, "range" }, > > + { "full", 0, 0, AV_OPT_TYPE_CONST, > > {.i64 = > AVCOL_RANGE_JPEG}, 0, 0, FLAGS, "range" }, > > + { "format", "output pixel format", OFFSET(format), > AV_OPT_TYPE_PIXEL_FMT, {.i64 = AV_PIX_FMT_NONE}, AV_PIX_FMT_NONE, > AV_PIX_FMT_GBRAP12LE, FLAGS, "fmt" }, > > Just make the top limit a large number to avoid putting something weird here > (e.g. libavcodec/options_table.h uses INT_MAX for this purpose). Will fix it. > > > + { "peak", "signal peak override", OFFSET(peak), > > AV_OPT_TYPE_DOUBLE, > {.dbl = 0}, 0, DBL_MAX, FLAGS }, > > + { "param", "tonemap parameter", OFFSET(param), > AV_OPT_TYPE_DOUBLE, {.dbl = NAN}, DBL_MIN, DBL_MAX, FLAGS }, > > You need a configure dependency on const_nan to use NAN. Will fix it, thanks! > > > + { "desat", "desaturation parameter", OFFSET(desat_param), > AV_OPT_TYPE_DOUBLE, {.dbl = 0.5}, 0, DBL_MAX, FLAGS }, > > + { "threshold", "scene detection threshold", OFFSET(scene_threshold), > AV_OPT_TYPE_DOUBLE, {.dbl = 0.2}, 0, DBL_MAX, FLAGS }, > > + { NULL } > > +}; > > + > > +AVFILTER_DEFINE_CLASS(tonemap_opencl); > > + > > +static const AVFilterPad tonemap_opencl_inputs[] = { > > + { > > + .name = "default", > > + .type = AVMEDIA_TYPE_VIDEO, > > + .filter_frame = &tonemap_opencl_filter_frame, > > + .config_props = &ff_opencl_filter_config_input, > > + }, > > + { NULL } > > +}; > > + > > +static const AVFilterPad tonemap_opencl_outputs[] = { > > + { > > + .name = "default", > > + .type = AVMEDIA_TYPE_VIDEO, > > + .config_props = &tonemap_opencl_config_output, > > + }, > > + { NULL } > > +}; > > + > > +AVFilter ff_vf_tonemap_opencl = { > > + .name = "tonemap_opencl", > > + .description = NULL_IF_CONFIG_SMALL("perform HDR to SDR conversion > with tonemapping"), > > + .priv_size = sizeof(TonemapOpenCLContext), > > + .priv_class = &tonemap_opencl_class, > > + .init = &ff_opencl_filter_init, > > + .uninit = &tonemap_opencl_uninit, > > + .query_formats = &ff_opencl_filter_query_formats, > > + .inputs = tonemap_opencl_inputs, > > + .outputs = tonemap_opencl_outputs, > > + .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE, > > +}; > > > > Thanks, > > - Mark Thanks for your careful review work! Ruiling > _______________________________________________ > ffmpeg-devel mailing list > ffmpeg-devel@ffmpeg.org > http://ffmpeg.org/mailman/listinfo/ffmpeg-devel _______________________________________________ ffmpeg-devel mailing list ffmpeg-devel@ffmpeg.org http://ffmpeg.org/mailman/listinfo/ffmpeg-devel