jellyfin/jellyfin-ffmpeg
1
0
Fork 0
mirror of https://github.com/jellyfin/jellyfin-ffmpeg.git synced 2026-09-03 05:09:58 +03:00
Code Issues Packages Projects Releases Wiki Activity

Merge pull request #430 from jellyfin/tonemap-7.0

Browse Source 
avfilter/tonemap*: add ICtCp based tonemap method
This commit is contained in:
gnattu
2024-08-10 16:05:07 +08:00
committed by GitHub
parent 587d4e0db1 16c55918a6
commit 24345c4812
3 changed files with 722 additions and 278 deletions
Download Patch File Download Diff File Expand all files Collapse all files
debian/patches/0004-add-cuda-tonemap-impl.patch
Vendored
+234 -71
View File
@@ -28,7 +28,7 @@ Index: FFmpeg/configure
+ nvccflags_default="--cuda-gpu-arch=sm_30 -O2 -ffast-math"
NVCC_C=""
fi
@@ -6711,7 +6713,7 @@ fi
if enabled cuda_nvcc; then
nvccflags="$nvccflags -ptx"
@@ -37,7 +37,7 @@ Index: FFmpeg/configure
+ nvccflags="$nvccflags -S -nocudalib -nocudainc --cuda-device-only -Wno-c++11-narrowing -std=c++14 -include ${source_link}/compat/cuda/cuda_runtime.h"
check_nvcc cuda_llvm
fi
Index: FFmpeg/ffbuild/common.mak
===================================================================
--- FFmpeg.orig/ffbuild/common.mak
@@ -47,7 +47,7 @@ Index: FFmpeg/ffbuild/common.mak
$(call PREPEND,CXXFLAGS, CPPFLAGS CFLAGS)
X86ASMFLAGS += $(IFLAGS:%=%/) -I$(<D)/ -Pconfig.asm
+NVCCFLAGS += $(IFLAGS)
HOSTCCFLAGS = $(IFLAGS) $(HOSTCPPFLAGS) $(HOSTCFLAGS)
LDFLAGS := $(ALLFFLIBS:%=$(LD_PATH)lib%) $(LDFLAGS)
Index: FFmpeg/libavfilter/Makefile
@@ -82,7 +82,7 @@ Index: FFmpeg/libavfilter/colorspace.c
@@ -51,6 +51,18 @@ void ff_matrix_invert_3x3(const double i
}
}
+void ff_matrix_transpose_3x3(const double in[3][3], double out[3][3])
+{
+ int i, j;
@@ -262,7 +262,7 @@ Index: FFmpeg/libavfilter/colorspace.h
#include "libavutil/frame.h"
#include "libavutil/pixfmt.h"
+#include "libavutil/dovi_meta.h"
#define REFERENCE_WHITE 100.0f
+#define REFERENCE_WHITE_ALT 203.0f
+#define ST2084_MAX_LUMINANCE 10000.0f
@@ -294,7 +294,7 @@ Index: FFmpeg/libavfilter/colorspace.h
+ float mmr_coeffs[8][3 /* order */][7];
+ } comp[3];
+};
void ff_matrix_invert_3x3(const double in[3][3], double out[3][3]);
+void ff_matrix_transpose_3x3(const double in[3][3], double out[3][3]);
void ff_matrix_mul_3x3(double dst[3][3],
@@ -303,7 +303,7 @@ Index: FFmpeg/libavfilter/colorspace.h
@@ -38,4 +70,19 @@ void ff_fill_rgb2yuv_table(const AVLumaC
double ff_determine_signal_peak(AVFrame *in);
void ff_update_hdr_metadata(AVFrame *in, double peak);
+double ff_determine_dovi_signal_peak(const AVDOVIMetadata *data);
+void ff_map_dovi_metadata(struct DoviMetadata *out, const AVDOVIMetadata *data);
+
@@ -324,7 +324,7 @@ Index: FFmpeg/libavfilter/cuda/colorspace_common.h
===================================================================
--- /dev/null
+++ FFmpeg/libavfilter/cuda/colorspace_common.h
@@ -0,0 +1,267 @@
@@ -0,0 +1,293 @@
+/*
+ * This file is part of FFmpeg.
+ *
@@ -561,7 +561,7 @@ Index: FFmpeg/libavfilter/cuda/colorspace_common.h
+ float r = linearize(c.x);
+ float g = linearize(c.y);
+ float b = linearize(c.z);
+ return lrgb2lrgb(make_float3(r, g, b));
+ return make_float3(r, g, b);
+}
+
+static __inline__ __device__ float3 ycc2rgb(float y, float cb, float cr) {
@@ -581,7 +581,7 @@ Index: FFmpeg/libavfilter/cuda/colorspace_common.h
+ rr = inverse_eotf_st2084_common(rr);
+ gg = inverse_eotf_st2084_common(gg);
+ bb = inverse_eotf_st2084_common(bb);
+ return make_float3(rr, gg, bb);
+ return rgb2lrgb(make_float3(rr, gg, bb));
+}
+
+static __inline__ __device__ float3 lms2rgb_fast(float r, float g, float b) {
@@ -591,6 +591,32 @@ Index: FFmpeg/libavfilter/cuda/colorspace_common.h
+ return make_float3(rr, gg, bb);
+}
+
+static __inline__ __device__ float3 lrgb2ictcp(float r, float g, float b) {
+ float l = 0.412109375000000f * r + 0.523925781250000f * g + 0.063964843750000f * b;
+ float m = 0.166748046875000f * r + 0.720458984375000f * g + 0.112792968750000f * b;
+ float s = 0.024169921875000f * r + 0.075439453125000f * g + 0.900390625000000f * b;
+ l = inverse_eotf_st2084(l);
+ m = inverse_eotf_st2084(m);
+ s = inverse_eotf_st2084(s);
+ float i = 0.5f * l + 0.5f * m;
+ float ct = 1.613769531250000f * l - 3.323486328125000f * m + 1.709716796875000f * s;
+ float cp = 4.378173828125000f * l - 4.245605468750000f * m - 0.132568359375000f * s;
+ return make_float3(i, ct, cp);
+}
+
+static __inline__ __device__ float3 ictcp2lrgb(float i, float ct, float cp) {
+ float ll = i + 0.008609037037933f * ct + 0.111029625003026f * cp;
+ float mm = i - 0.008609037037933f * ct - 0.111029625003026f * cp;
+ float ss = i + 0.560031335710679f * ct - 0.320627174987319f * cp;
+ ll = eotf_st2084(ll);
+ mm = eotf_st2084(mm);
+ ss = eotf_st2084(ss);
+ float r = 3.436606694333079f * ll - 2.506452118656270f * mm + 0.069845424323191f * ss;
+ float g = -0.791329555598929f * ll + 1.983600451792291f * mm - 0.192270896193362f * ss;
+ float b = -0.025949899690593f * ll - 0.098913714711726f * mm + 1.124863614402319f * ss;
+ return make_float3(r, g, b);
+}
+
+#endif /* AVFILTER_CUDA_COLORSPACE_COMMON_H */
Index: FFmpeg/libavfilter/cuda/host_util.c
===================================================================
@@ -981,7 +1007,7 @@ Index: FFmpeg/libavfilter/cuda/tonemap.cu
===================================================================
--- /dev/null
+++ FFmpeg/libavfilter/cuda/tonemap.cu
@@ -0,0 +1,455 @@
@@ -0,0 +1,563 @@
+/*
+ * This file is part of FFmpeg.
+ *
@@ -1018,6 +1044,15 @@ Index: FFmpeg/libavfilter/cuda/tonemap.cu
+#define dot4(a, b) ((a).w * (b).w + ((a).z * (b).z + ((a).y * (b).y + (a).x * (b).x)))
+
+static __inline__ __device__
+float3 clamp3(const float3 a, const float min_val, const float max_val) {
+ float3 result;
+ result.x = clamp(a.x, min_val, max_val);
+ result.y = clamp(a.y, min_val, max_val);
+ result.z = clamp(a.z, min_val, max_val);
+ return result;
+}
+
+static __inline__ __device__
+float get_dithered_y(float y, float d) {
+ return floor(y * dither_quantization + d + 0.5f / (dither_size * dither_size)) * 1.0f / dither_quantization;
+}
@@ -1075,11 +1110,11 @@ Index: FFmpeg/libavfilter/cuda/tonemap.cu
+}
+
+static __inline__ __device__
+float bt2390(float s, float peak, float dst_peak) {
+float bt2390_common(float s, float peak, float dst_peak) {
+ float peak_pq = inverse_eotf_st2084(peak);
+ float scale = peak_pq > 0.0f ? (1.0f / peak_pq) : 1.0f;
+
+ float s_pq = inverse_eotf_st2084(s) * scale;
+ float s_pq = s * scale;
+ float max_lum = inverse_eotf_st2084(dst_peak) * scale;
+
+ float ks = 1.5f * max_lum - 0.5f;
@@ -1091,7 +1126,13 @@ Index: FFmpeg/libavfilter/cuda/tonemap.cu
+ (-2.0f * tb3 + 3.0f * tb2) * max_lum;
+ float sig = mix(pb, s_pq, s_pq < ks);
+
+ return eotf_st2084(sig * peak_pq);
+ return sig * peak_pq;
+}
+
+static __inline__ __device__
+float bt2390(float s, float peak, float dst_peak) {
+ float s_pq = inverse_eotf_st2084(s);
+ return eotf_st2084(bt2390_common(s_pq, peak, dst_peak));
+}
+
+static __inline__ __device__
@@ -1119,6 +1160,17 @@ Index: FFmpeg/libavfilter/cuda/tonemap.cu
+}
+
+static __inline__ __device__
+float map_itp(float s, float peak, float dst_peak)
+{
+ switch (tonemap_func) {
+ default:
+ return inverse_eotf_st2084(map(eotf_st2084(s), peak, dst_peak));
+ case TONEMAP_BT2390:
+ return bt2390_common(s, peak, dst_peak);
+ }
+}
+
+static __inline__ __device__
+float3 map_one_pixel_rgb_mode_max(float3 rgb, const FFCUDAFrame& src, const FFCUDAFrame& dst) {
+ float sig = max(max(rgb.x, max(rgb.y, rgb.z)), FLOAT_EPS);
+ float sig_old = sig;
@@ -1171,7 +1223,7 @@ Index: FFmpeg/libavfilter/cuda/tonemap.cu
+}
+
+static __inline__ __device__
+float3 map_one_pixel_rgb_mode_rl(float3 rgb, const FFCUDAFrame& src, const FFCUDAFrame& dst) {
+float3 map_one_pixel_rgb_mode_lum(float3 rgb, const FFCUDAFrame& src, const FFCUDAFrame& dst) {
+ float sig = max((rgb.x * 0.2627f + rgb.y * 0.678f + rgb.z * 0.0593f), FLOAT_EPS);
+ float peak = src.peak;
+ sig = min(sig, peak);
@@ -1180,21 +1232,41 @@ Index: FFmpeg/libavfilter/cuda/tonemap.cu
+
+ // Desaturate the color using a coefficient dependent on the signal level
+ if (desat_param > 0.0f) {
+ float luma = get_luma_dst(rgb, luma_dst);
+ float coeff = max(sig - 0.18f, FLOAT_EPS) / max(sig, FLOAT_EPS);
+ coeff = __powf(coeff, 10.0f / desat_param);
+ rgb = mix(rgb, make_float3(luma, luma, luma), make_float3(coeff, coeff, coeff));
+ rgb = mix(rgb, make_float3(sig, sig, sig), make_float3(coeff, coeff, coeff));
+ }
+
+ sig = map(sig, peak, dst_peak);
+ rgb = rgb * (sig / sig_old);
+ rgb.x = clamp(rgb.x, 0.0f, 1.0f);
+ rgb.y = clamp(rgb.y, 0.0f, 1.0f);
+ rgb.z = clamp(rgb.z, 0.0f, 1.0f);
+
+ return rgb;
+}
+
+static __inline__ __device__
+float3 map_one_pixel_itp_mode(float3 rgb, const FFCUDAFrame& src, const FFCUDAFrame& dst) {
+ float3 ictcp = lrgb2ictcp(rgb.x, rgb.y, rgb.z);
+ float peak = src.peak;
+ float dst_peak = 1.0f;
+ ictcp.x = max(ictcp.x, FLOAT_EPS);
+ float i_o = ictcp.x;
+
+ if (desat_param > 0.0f) {
+ float p = eotf_st2084(ictcp.x) - (dst_peak - desat_param) * 0.5f;
+ float coeff = __expf(-(p * p) / (2.0f * peak));
+ ictcp.y *= coeff;
+ ictcp.z *= coeff;
+ }
+
+ ictcp.x = map_itp(ictcp.x, peak, dst_peak);
+ ictcp.x = min(ictcp.x, 1.0f);
+ float factor = min(ictcp.x / i_o, i_o / ictcp.x);
+ ictcp.y *= factor;
+ ictcp.z *= factor;
+
+ return ictcp2lrgb(ictcp.x, ictcp.y, ictcp.z);
+}
+
+// Map from source space YUV to destination space RGB
+static __inline__ __device__
+float3 map_to_dst_space_from_yuv(float3 yuv) {
@@ -1203,17 +1275,39 @@ Index: FFmpeg/libavfilter/cuda/tonemap.cu
+}
+
+static __inline__ __device__
+float3 map_to_src_space_from_yuv(float3 yuv) {
+ float3 c = yuv2lrgb(yuv);
+ return c;
+}
+
+static __inline__ __device__
+float3 map_to_dst_space_from_yuv_dovi(float3 yuv) {
+ float3 c = ycc2rgb(yuv.x, yuv.y, yuv.z);
+ c = lms2rgb(c.x, c.y, c.z);
+ return rgb2lrgb(c);
+ c = lrgb2lrgb(c);
+ return c;
+}
+
+static __inline__ __device__
+float3 map_to_dst_space_from_yuv_dovi_fast(float3 yuv) {
+ float3 c = ycc2rgb(yuv.x, yuv.y, yuv.z);
+ c = lms2rgb_fast(c.x, c.y, c.z);
+ return rgb2lrgb(c);
+ c = lrgb2lrgb(c);
+ return c;
+}
+
+static __inline__ __device__
+float3 map_to_src_space_from_yuv_dovi(float3 yuv) {
+ float3 c = ycc2rgb(yuv.x, yuv.y, yuv.z);
+ c = lms2rgb(c.x, c.y, c.z);
+ return c;
+}
+
+static __inline__ __device__
+float3 map_to_src_space_from_yuv_dovi_fast(float3 yuv) {
+ float3 c = ycc2rgb(yuv.x, yuv.y, yuv.z);
+ c = lms2rgb_fast(c.x, c.y, c.z);
+ return c;
+}
+
+static __inline__ __device__
@@ -1352,6 +1446,12 @@ Index: FFmpeg/libavfilter/cuda/tonemap.cu
+ float3 c2 = map_to_dst_space_from_yuv(yuv2); \
+ float3 c3 = map_to_dst_space_from_yuv(yuv3);
+
+#define _YUV2RGB_S \
+ float3 c0 = map_to_src_space_from_yuv(yuv0); \
+ float3 c1 = map_to_src_space_from_yuv(yuv1); \
+ float3 c2 = map_to_src_space_from_yuv(yuv2); \
+ float3 c3 = map_to_src_space_from_yuv(yuv3);
+
+#define _YCC2RGB \
+ float3 c0 = map_to_dst_space_from_yuv_dovi(yuv0); \
+ float3 c1 = map_to_dst_space_from_yuv_dovi(yuv1); \
@@ -1364,6 +1464,18 @@ Index: FFmpeg/libavfilter/cuda/tonemap.cu
+ float3 c2 = map_to_dst_space_from_yuv_dovi_fast(yuv2); \
+ float3 c3 = map_to_dst_space_from_yuv_dovi_fast(yuv3);
+
+#define _YCC2RGB_S \
+ float3 c0 = map_to_src_space_from_yuv_dovi(yuv0); \
+ float3 c1 = map_to_src_space_from_yuv_dovi(yuv1); \
+ float3 c2 = map_to_src_space_from_yuv_dovi(yuv2); \
+ float3 c3 = map_to_src_space_from_yuv_dovi(yuv3);
+
+#define _YCC2RGB_FS \
+ float3 c0 = map_to_src_space_from_yuv_dovi_fast(yuv0); \
+ float3 c1 = map_to_src_space_from_yuv_dovi_fast(yuv1); \
+ float3 c2 = map_to_src_space_from_yuv_dovi_fast(yuv2); \
+ float3 c3 = map_to_src_space_from_yuv_dovi_fast(yuv3);
+
+#define _TONEMAP_MAX \
+ c0 = map_one_pixel_rgb_mode_max(c0, src, dst); \
+ c1 = map_one_pixel_rgb_mode_max(c1, src, dst); \
@@ -1376,11 +1488,17 @@ Index: FFmpeg/libavfilter/cuda/tonemap.cu
+ c2 = map_one_pixel_rgb_mode_rgb(c2, src, dst); \
+ c3 = map_one_pixel_rgb_mode_rgb(c3, src, dst);
+
+#define _TONEMAP_RL \
+ c0 = map_one_pixel_rgb_mode_rl(c0, src, dst); \
+ c1 = map_one_pixel_rgb_mode_rl(c1, src, dst); \
+ c2 = map_one_pixel_rgb_mode_rl(c2, src, dst); \
+ c3 = map_one_pixel_rgb_mode_rl(c3, src, dst);
+#define _TONEMAP_LUM \
+ c0 = map_one_pixel_rgb_mode_lum(c0, src, dst); \
+ c1 = map_one_pixel_rgb_mode_lum(c1, src, dst); \
+ c2 = map_one_pixel_rgb_mode_lum(c2, src, dst); \
+ c3 = map_one_pixel_rgb_mode_lum(c3, src, dst);
+
+#define _TONEMAP_ITP \
+ c0 = map_one_pixel_itp_mode(c0, src, dst); \
+ c1 = map_one_pixel_itp_mode(c1, src, dst); \
+ c2 = map_one_pixel_itp_mode(c2, src, dst); \
+ c3 = map_one_pixel_itp_mode(c3, src, dst);
+
+#define _RGB2YUV \
+ yuv0 = lrgb2yuv(c0); \
@@ -1388,6 +1506,16 @@ Index: FFmpeg/libavfilter/cuda/tonemap.cu
+ yuv2 = lrgb2yuv(c2); \
+ yuv3 = lrgb2yuv(c3);
+
+#define _RGB2YUV_S \
+ c0 = clamp3(lrgb2lrgb(c0), 0.0f, 1.0f); \
+ c1 = clamp3(lrgb2lrgb(c1), 0.0f, 1.0f); \
+ c2 = clamp3(lrgb2lrgb(c2), 0.0f, 1.0f); \
+ c3 = clamp3(lrgb2lrgb(c3), 0.0f, 1.0f); \
+ yuv0 = lrgb2yuv(c0); \
+ yuv1 = lrgb2yuv(c1); \
+ yuv2 = lrgb2yuv(c2); \
+ yuv3 = lrgb2yuv(c3);
+
+#define _DITHER \
+ float d = read_dither(ditherTex, dither_size, xi, yi); \
+ yuv0.x = get_dithered_y(yuv0.x, d); \
@@ -1412,36 +1540,42 @@ Index: FFmpeg/libavfilter/cuda/tonemap.cu
+ WRITER \
+}
+
+TONEMAP_VARIANT(, _READER, , _YUV2RGB, _TONEMAP_MAX, _RGB2YUV, , _WRITER)
+TONEMAP_VARIANT(_d, _READER, , _YUV2RGB, _TONEMAP_MAX, _RGB2YUV, _DITHER, _WRITER)
+TONEMAP_VARIANT(_rgb, _READER, , _YUV2RGB, _TONEMAP_RGB, _RGB2YUV, , _WRITER)
+TONEMAP_VARIANT(_rgb_d, _READER, , _YUV2RGB, _TONEMAP_RGB, _RGB2YUV, _DITHER, _WRITER)
+TONEMAP_VARIANT(_rl, _READER, , _YUV2RGB, _TONEMAP_RL, _RGB2YUV, , _WRITER)
+TONEMAP_VARIANT(_rl_d, _READER, , _YUV2RGB, _TONEMAP_RL, _RGB2YUV, _DITHER, _WRITER)
+TONEMAP_VARIANT(, _READER, , _YUV2RGB, _TONEMAP_MAX, _RGB2YUV, , _WRITER)
+TONEMAP_VARIANT(_d, _READER, , _YUV2RGB, _TONEMAP_MAX, _RGB2YUV, _DITHER, _WRITER)
+TONEMAP_VARIANT(_rgb, _READER, , _YUV2RGB, _TONEMAP_RGB, _RGB2YUV, , _WRITER)
+TONEMAP_VARIANT(_rgb_d, _READER, , _YUV2RGB, _TONEMAP_RGB, _RGB2YUV, _DITHER, _WRITER)
+TONEMAP_VARIANT(_lum, _READER, , _YUV2RGB_S, _TONEMAP_LUM, _RGB2YUV_S, , _WRITER)
+TONEMAP_VARIANT(_lum_d, _READER, , _YUV2RGB_S, _TONEMAP_LUM, _RGB2YUV_S, _DITHER, _WRITER)
+TONEMAP_VARIANT(_itp, _READER, , _YUV2RGB_S, _TONEMAP_ITP, _RGB2YUV_S, , _WRITER)
+TONEMAP_VARIANT(_itp_d, _READER, , _YUV2RGB_S, _TONEMAP_ITP, _RGB2YUV_S, _DITHER, _WRITER)
+
+TONEMAP_VARIANT(_dovi, _READER, _RESHAPE, _YCC2RGB, _TONEMAP_MAX, _RGB2YUV, , _WRITER)
+TONEMAP_VARIANT(_dovi_d, _READER, _RESHAPE, _YCC2RGB, _TONEMAP_MAX, _RGB2YUV, _DITHER, _WRITER)
+TONEMAP_VARIANT(_dovi_rgb, _READER, _RESHAPE, _YCC2RGB, _TONEMAP_RGB, _RGB2YUV, , _WRITER)
+TONEMAP_VARIANT(_dovi_rgb_d, _READER, _RESHAPE, _YCC2RGB, _TONEMAP_RGB, _RGB2YUV, _DITHER, _WRITER)
+TONEMAP_VARIANT(_dovi_rl, _READER, _RESHAPE, _YCC2RGB, _TONEMAP_RL, _RGB2YUV, , _WRITER)
+TONEMAP_VARIANT(_dovi_rl_d, _READER, _RESHAPE, _YCC2RGB, _TONEMAP_RL, _RGB2YUV, _DITHER, _WRITER)
+TONEMAP_VARIANT(_dovi, _READER, _RESHAPE, _YCC2RGB, _TONEMAP_MAX, _RGB2YUV, , _WRITER)
+TONEMAP_VARIANT(_dovi_d, _READER, _RESHAPE, _YCC2RGB, _TONEMAP_MAX, _RGB2YUV, _DITHER, _WRITER)
+TONEMAP_VARIANT(_dovi_rgb, _READER, _RESHAPE, _YCC2RGB, _TONEMAP_RGB, _RGB2YUV, , _WRITER)
+TONEMAP_VARIANT(_dovi_rgb_d, _READER, _RESHAPE, _YCC2RGB, _TONEMAP_RGB, _RGB2YUV, _DITHER, _WRITER)
+TONEMAP_VARIANT(_dovi_lum, _READER, _RESHAPE, _YCC2RGB_S, _TONEMAP_LUM, _RGB2YUV_S, , _WRITER)
+TONEMAP_VARIANT(_dovi_lum_d, _READER, _RESHAPE, _YCC2RGB_S, _TONEMAP_LUM, _RGB2YUV_S, _DITHER, _WRITER)
+TONEMAP_VARIANT(_dovi_itp, _READER, _RESHAPE, _YCC2RGB_S, _TONEMAP_ITP, _RGB2YUV_S, , _WRITER)
+TONEMAP_VARIANT(_dovi_itp_d, _READER, _RESHAPE, _YCC2RGB_S, _TONEMAP_ITP, _RGB2YUV_S, _DITHER, _WRITER)
+
+TONEMAP_VARIANT(_dovi_f, _READER, _RESHAPE, _YCC2RGB_F, _TONEMAP_MAX, _RGB2YUV, , _WRITER)
+TONEMAP_VARIANT(_dovi_d_f, _READER, _RESHAPE, _YCC2RGB_F, _TONEMAP_MAX, _RGB2YUV, _DITHER, _WRITER)
+TONEMAP_VARIANT(_dovi_rgb_f, _READER, _RESHAPE, _YCC2RGB_F, _TONEMAP_RGB, _RGB2YUV, , _WRITER)
+TONEMAP_VARIANT(_dovi_rgb_d_f, _READER, _RESHAPE, _YCC2RGB_F, _TONEMAP_RGB, _RGB2YUV, _DITHER, _WRITER)
+TONEMAP_VARIANT(_dovi_rl_f, _READER, _RESHAPE, _YCC2RGB_F, _TONEMAP_RL, _RGB2YUV, , _WRITER)
+TONEMAP_VARIANT(_dovi_rl_d_f, _READER, _RESHAPE, _YCC2RGB_F, _TONEMAP_RL, _RGB2YUV, _DITHER, _WRITER)
+TONEMAP_VARIANT(_dovi_f, _READER, _RESHAPE, _YCC2RGB_F, _TONEMAP_MAX, _RGB2YUV, , _WRITER)
+TONEMAP_VARIANT(_dovi_d_f, _READER, _RESHAPE, _YCC2RGB_F, _TONEMAP_MAX, _RGB2YUV, _DITHER, _WRITER)
+TONEMAP_VARIANT(_dovi_rgb_f, _READER, _RESHAPE, _YCC2RGB_F, _TONEMAP_RGB, _RGB2YUV, , _WRITER)
+TONEMAP_VARIANT(_dovi_rgb_d_f, _READER, _RESHAPE, _YCC2RGB_F, _TONEMAP_RGB, _RGB2YUV, _DITHER, _WRITER)
+TONEMAP_VARIANT(_dovi_lum_f, _READER, _RESHAPE, _YCC2RGB_FS, _TONEMAP_LUM, _RGB2YUV_S, , _WRITER)
+TONEMAP_VARIANT(_dovi_lum_d_f, _READER, _RESHAPE, _YCC2RGB_FS, _TONEMAP_LUM, _RGB2YUV_S, _DITHER, _WRITER)
+TONEMAP_VARIANT(_dovi_itp_f, _READER, _RESHAPE, _YCC2RGB_FS, _TONEMAP_ITP, _RGB2YUV_S, , _WRITER)
+TONEMAP_VARIANT(_dovi_itp_d_f, _READER, _RESHAPE, _YCC2RGB_FS, _TONEMAP_ITP, _RGB2YUV_S, _DITHER, _WRITER)
+
+TONEMAP_VARIANT(_dovi_pq, _READER, _RESHAPE, _YCC2RGB, , _RGB2YUV, , _WRITER)
+TONEMAP_VARIANT(_dovi_pq_f, _READER, _RESHAPE, _YCC2RGB_F, , _RGB2YUV, , _WRITER)
+TONEMAP_VARIANT(_dovi_pq, _READER, _RESHAPE, _YCC2RGB, , _RGB2YUV, , _WRITER)
+TONEMAP_VARIANT(_dovi_pq_f, _READER, _RESHAPE, _YCC2RGB_F, , _RGB2YUV, , _WRITER)
+
+}
Index: FFmpeg/libavfilter/cuda/tonemap.h
===================================================================
--- /dev/null
+++ FFmpeg/libavfilter/cuda/tonemap.h
@@ -0,0 +1,41 @@
@@ -0,0 +1,43 @@
+/*
+ * This file is part of FFmpeg.
+ *
@@ -1479,6 +1613,8 @@ Index: FFmpeg/libavfilter/cuda/tonemap.h
+ TONEMAP_MODE_MAX,
+ TONEMAP_MODE_RGB,
+ TONEMAP_MODE_LUM,
+ TONEMAP_MODE_ITP,
+ TONEMAP_MODE_AUTO,
+ TONEMAP_MODE_COUNT,
+};
+
@@ -1578,7 +1714,7 @@ Index: FFmpeg/libavfilter/vf_tonemap_cuda.c
===================================================================
--- /dev/null
+++ FFmpeg/libavfilter/vf_tonemap_cuda.c
@@ -0,0 +1,1100 @@
@@ -0,0 +1,1127 @@
+/*
+ * This file is part of FFmpeg.
+ *
@@ -2090,8 +2226,6 @@ Index: FFmpeg/libavfilter/vf_tonemap_cuda.c
+ enum AVColorSpace in_spc = s->in_spc, out_spc = s->out_spc;
+ enum AVColorPrimaries in_pri = s->in_pri, out_pri = s->out_pri;
+ enum AVColorRange in_range = s->in_range, out_range = s->out_range;
+ int rgb = s->tonemap_mode == TONEMAP_MODE_RGB;
+ int max = s->tonemap_mode == TONEMAP_MODE_MAX;
+ int d = s->in_desc->comp[0].depth > s->out_desc->comp[0].depth && s->ditherTex;
+ char info_log[4096], error_log[4096];
+ CUjit_option options[] = { CU_JIT_INFO_LOG_BUFFER,
@@ -2289,6 +2423,14 @@ Index: FFmpeg/libavfilter/vf_tonemap_cuda.c
+ av_log(ctx, AV_LOG_DEBUG, "Disabled dovi tradeoff on high perf GPU.\n");
+ }
+
+ if (s->tonemap_mode == TONEMAP_MODE_AUTO) {
+ if (s->tradeoff) {
+ s->tonemap_mode = TONEMAP_MODE_LUM;
+ } else {
+ s->tonemap_mode = TONEMAP_MODE_ITP;
+ }
+ }
+
+ if (s->cu_module) {
+ ret = CHECK_CU(cu->cuModuleUnload(s->cu_module));
+ if (ret < 0)
@@ -2322,22 +2464,41 @@ Index: FFmpeg/libavfilter/vf_tonemap_cuda.c
+ if (ret < 0)
+ goto fail2;
+
+ ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_tm, s->cu_module,
+ rgb ? (d ? "tonemap_rgb_d" : "tonemap_rgb")
+ : (max ? (d ? "tonemap_d" : "tonemap")
+ : (d ? "tonemap_rl_d" : "tonemap_rl"))));
+ if (ret < 0)
+ goto fail2;
+
+ ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_dovi, s->cu_module,
+ s->tradeoff == 1 ? (rgb ? (d ? "tonemap_dovi_rgb_d_f" : "tonemap_dovi_rgb_f")
+ : (max ? (d ? "tonemap_dovi_d_f" : "tonemap_dovi_f")
+ : (d ? "tonemap_dovi_rl_d_f" : "tonemap_dovi_rl_f")))
+ : (rgb ? (d ? "tonemap_dovi_rgb_d" : "tonemap_dovi_rgb")
+ : (max ? (d ? "tonemap_dovi_d" : "tonemap_dovi")
+ : (d ? "tonemap_dovi_rl_d" : "tonemap_dovi_rl")))));
+ if (ret < 0)
+ goto fail2;
+ switch (s->tonemap_mode) {
+ default:
+ case TONEMAP_MODE_MAX:
+ ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_tm, s->cu_module, d ? "tonemap_max_d" : "tonemap_max"));
+ if (ret < 0) goto fail2;
+ ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_dovi, s->cu_module,
+ s->tradeoff == 1 ? (d ? "tonemap_dovi_d_f" : "tonemap_dovi_f")
+ : (d ? "tonemap_dovi_d" : "tonemap_dovi")));
+ if (ret < 0) goto fail2;
+ break;
+ case TONEMAP_MODE_RGB:
+ ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_tm, s->cu_module, d ? "tonemap_rgb_d" : "tonemap_rgb"));
+ if (ret < 0) goto fail2;
+ ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_dovi, s->cu_module,
+ s->tradeoff == 1 ? (d ? "tonemap_dovi_rgb_d_f" : "tonemap_dovi_rgb_f")
+ : (d ? "tonemap_dovi_rgb_d" : "tonemap_dovi_rgb")));
+ if (ret < 0) goto fail2;
+ break;
+ case TONEMAP_MODE_LUM:
+ ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_tm, s->cu_module, d ? "tonemap_lum_d" : "tonemap_lum"));
+ if (ret < 0) goto fail2;
+ ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_dovi, s->cu_module,
+ s->tradeoff == 1 ? (d ? "tonemap_dovi_lum_d_f" : "tonemap_dovi_lum_f")
+ : (d ? "tonemap_dovi_lum_d" : "tonemap_dovi_lum")));
+ if (ret < 0) goto fail2;
+ break;
+ case TONEMAP_MODE_ITP:
+ ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_tm, s->cu_module, d ? "tonemap_itp_d" : "tonemap_itp"));
+ if (ret < 0) goto fail2;
+ ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_dovi, s->cu_module,
+ s->tradeoff == 1 ? (d ? "tonemap_dovi_itp_d_f" : "tonemap_dovi_itp_f")
+ : (d ? "tonemap_dovi_itp_d" : "tonemap_dovi_itp")));
+ if (ret < 0) goto fail2;
+ break;
+ }
+
+ ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_dovi_pq, s->cu_module,
+ s->tradeoff == 1 ? "tonemap_dovi_pq_f"
@@ -2603,10 +2764,12 @@ Index: FFmpeg/libavfilter/vf_tonemap_cuda.c
+ { "hable", 0, 0, AV_OPT_TYPE_CONST, {.i64 = TONEMAP_HABLE}, 0, 0, FLAGS, .unit = "tonemap" },
+ { "mobius", 0, 0, AV_OPT_TYPE_CONST, {.i64 = TONEMAP_MOBIUS}, 0, 0, FLAGS, .unit = "tonemap" },
+ { "bt2390", 0, 0, AV_OPT_TYPE_CONST, {.i64 = TONEMAP_BT2390}, 0, 0, FLAGS, .unit = "tonemap" },
+ { "tonemap_mode", "Tonemap mode selection", OFFSET(tonemap_mode), AV_OPT_TYPE_INT, {.i64 = TONEMAP_MODE_MAX}, TONEMAP_MODE_MAX, TONEMAP_MODE_COUNT - 1, FLAGS, .unit = "tonemap_mode" },
+ { "max", "Brightest channel based tonemap", 0, AV_OPT_TYPE_CONST, {.i64 = TONEMAP_MODE_MAX}, 0, 0, FLAGS, .unit = "tonemap_mode" },
+ { "rgb", "Per-channel based tonemap", 0, AV_OPT_TYPE_CONST, {.i64 = TONEMAP_MODE_RGB}, 0, 0, FLAGS, .unit = "tonemap_mode" },
+ { "lum", "Relative luminance based tonemap", 0, AV_OPT_TYPE_CONST, {.i64 = TONEMAP_MODE_LUM}, 0, 0, FLAGS, .unit = "tonemap_mode" },
+ { "tonemap_mode", "Tonemap mode selection", OFFSET(tonemap_mode), AV_OPT_TYPE_INT, {.i64 = TONEMAP_MODE_AUTO}, TONEMAP_MODE_MAX, TONEMAP_MODE_COUNT - 1, FLAGS, .unit = "tonemap_mode" },
+ { "max", "Brightest channel based tonemap", 0, AV_OPT_TYPE_CONST, {.i64 = TONEMAP_MODE_MAX}, 0, 0, FLAGS, .unit = "tonemap_mode" },
+ { "rgb", "Per-channel based tonemap", 0, AV_OPT_TYPE_CONST, {.i64 = TONEMAP_MODE_RGB}, 0, 0, FLAGS, .unit = "tonemap_mode" },
+ { "lum", "Relative luminance based tonemap", 0, AV_OPT_TYPE_CONST, {.i64 = TONEMAP_MODE_LUM}, 0, 0, FLAGS, .unit = "tonemap_mode" },
+ { "itp", "ICtCp intensity based tonemap", 0, AV_OPT_TYPE_CONST, {.i64 = TONEMAP_MODE_ITP}, 0, 0, FLAGS, .unit = "tonemap_mode" },
+ { "auto", "Select based on GPU spec", 0, AV_OPT_TYPE_CONST, {.i64 = TONEMAP_MODE_AUTO}, 0, 0, FLAGS, .unit = "tonemap_mode" },
+ { "transfer", "Set transfer characteristic", OFFSET(trc), AV_OPT_TYPE_INT, {.i64 = AVCOL_TRC_BT709}, -1, INT_MAX, FLAGS, .unit = "transfer" },
+ { "t", "Set transfer characteristic", OFFSET(trc), AV_OPT_TYPE_INT, {.i64 = AVCOL_TRC_BT709}, -1, INT_MAX, FLAGS, .unit = "transfer" },
+ { "bt709", 0, 0, AV_OPT_TYPE_CONST, {.i64 = AVCOL_TRC_BT709}, 0, 0, FLAGS, .unit = "transfer" },
debian/patches/0007-add-bt2390-eetf-and-code-refactor-to-opencl-tonemap.patch
Vendored
+323 -179
View File
File diff suppressed because it is too large Load Diff
debian/patches/0052-add-vf-tonemap-videotoolbox-filter.patch
Vendored
+165 -28
View File
@@ -40,7 +40,7 @@ Index: FFmpeg/libavfilter/metal/vf_tonemap_videotoolbox.metal
===================================================================
--- /dev/null
+++ FFmpeg/libavfilter/metal/vf_tonemap_videotoolbox.metal
@@ -0,0 +1,725 @@
@@ -0,0 +1,850 @@
+/*
+ * Copyright (c) 2024 Gnattu OC <gnattuoc@me.com>
+ *
@@ -121,6 +121,8 @@ Index: FFmpeg/libavfilter/metal/vf_tonemap_videotoolbox.metal
+constant float3 luma_dst [[function_constant(35)]];
+constant short linearize_type [[function_constant(36)]];
+constant short delinearize_type [[function_constant(37)]];
+constant bool map_in_src_space [[function_constant(38)]];
+constant bool is_tone_mode_itp [[function_constant(39)]];
+
+enum AVChromaLocation {
+ AVCHROMA_LOC_UNSPECIFIED,
@@ -167,6 +169,14 @@ Index: FFmpeg/libavfilter/metal/vf_tonemap_videotoolbox.metal
+ return eotf_st2084_common(x) * pq_max_lum_div_ref_white;
+}
+
+float4 eotf_st2084x4(float4 x) {
+ x.x = eotf_st2084_common(x.x);
+ x.y = eotf_st2084_common(x.y);
+ x.z = eotf_st2084_common(x.z);
+ x.w = eotf_st2084_common(x.w);
+ return x * pq_max_lum_div_ref_white;
+}
+
+// delinearizer for PQ/ST2084
+float inverse_eotf_st2084_common(float x) {
+ x = fmax(x, 0.0f);
@@ -181,6 +191,15 @@ Index: FFmpeg/libavfilter/metal/vf_tonemap_videotoolbox.metal
+ return inverse_eotf_st2084_common(x);
+}
+
+float4 inverse_eotf_st2084x4(float4 x) {
+ x *= ref_white_div_pq_max_lum;
+ x.x = inverse_eotf_st2084_common(x.x);
+ x.y = inverse_eotf_st2084_common(x.y);
+ x.z = inverse_eotf_st2084_common(x.z);
+ x.w = inverse_eotf_st2084_common(x.w);
+ return x;
+}
+
+float ootf_1_2(float x) {
+ return x > 0.0f ? powr(x, 1.2f) : x;
+}
@@ -213,12 +232,27 @@ Index: FFmpeg/libavfilter/metal/vf_tonemap_videotoolbox.metal
+ return oetf_arib_b67(inverse_ootf_1_2(x / 5.0f));
+}
+
+float4 oetf_arib_b67x4(float4 x) {
+ x.x = oetf_arib_b67(x.x);
+ x.y = oetf_arib_b67(x.y);
+ x.z = oetf_arib_b67(x.z);
+ x.w = oetf_arib_b67(x.w);
+ return x;
+}
+
+float4 inverse_oetf_arib_b67x4(float4 x) {
+ x.x = inverse_oetf_arib_b67(x.x);
+ x.y = inverse_oetf_arib_b67(x.y);
+ x.z = inverse_oetf_arib_b67(x.z);
+ x.w = inverse_oetf_arib_b67(x.w);
+ return x;
+}
+
+// delinearizer for BT709, BT2020-10
+float inverse_eotf_bt1886(float x) {
+ return x > 0.0f ? powr(x, 1.0f / 2.4f) : 0.0f;
+}
+
+
+float linearize(float x) {
+ if (linearize_type == 1) {
+ return eotf_st2084(x);
@@ -322,7 +356,7 @@ Index: FFmpeg/libavfilter/metal/vf_tonemap_videotoolbox.metal
+ float r = linearize(c.x);
+ float g = linearize(c.y);
+ float b = linearize(c.z);
+ return lrgb2lrgb(float3(r, g, b));
+ return float3(r, g, b);
+}
+
+float3 ycc2rgb(float y, float cb, float cr) {
@@ -345,6 +379,30 @@ Index: FFmpeg/libavfilter/metal/vf_tonemap_videotoolbox.metal
+ return float3(rr, gg, bb);
+}
+
+// The following assumes bt2020
+void lrgb2ictcp(float4 r4, float4 g4, float4 b4, thread float4* i4, thread float4* ct4, thread float4* cp4) {
+ float4 l4 = 0.412109375000000f * r4 + 0.523925781250000f * g4 + 0.063964843750000f * b4;
+ float4 m4 = 0.166748046875000f * r4 + 0.720458984375000f * g4 + 0.112792968750000f * b4;
+ float4 s4 = 0.024169921875000f * r4 + 0.075439453125000f * g4 + 0.900390625000000f * b4;
+ l4 = inverse_eotf_st2084x4(l4);
+ m4 = inverse_eotf_st2084x4(m4);
+ s4 = inverse_eotf_st2084x4(s4);
+ *i4 = 0.5f * l4 + 0.5f * m4;
+ *ct4 = 1.613769531250000f * l4 - 3.323486328125000f * m4 + 1.709716796875000f * s4;
+ *cp4 = 4.378173828125000f * l4 - 4.245605468750000f * m4 - 0.132568359375000f * s4;
+}
+
+void ictcp2lrgb(float4 i4, float4 ct4, float4 cp4, thread float4* r4, thread float4* g4, thread float4* b4) {
+ float4 ll4 = i4 + 0.008609037037933f * ct4 + 0.111029625003026f * cp4;
+ float4 mm4 = i4 - 0.008609037037933f * ct4 - 0.111029625003026f * cp4;
+ float4 ss4 = i4 + 0.560031335710679f * ct4 - 0.320627174987319f * cp4;
+ ll4 = eotf_st2084x4(ll4);
+ mm4 = eotf_st2084x4(mm4);
+ ss4 = eotf_st2084x4(ss4);
+ *r4 = 3.436606694333079f * ll4 - 2.506452118656270f * mm4 + 0.069845424323191f * ss4;
+ *g4 = -0.791329555598929f * ll4 + 1.983600451792291f * mm4 - 0.192270896193362f * ss4;
+ *b4 = -0.025949899690593f * ll4 - 0.098913714711726f * mm4 + 1.124863614402319f * ss4;
+}
+
+//------------
+// Tonemapping methods
@@ -408,7 +466,7 @@ Index: FFmpeg/libavfilter/metal/vf_tonemap_videotoolbox.metal
+ float peak_pq = peak_inv_pq;
+ float scale = peak_pq > 0.0f ? (1.0f / peak_pq) : 1.0f;
+
+ float s_pq = inverse_eotf_st2084(s) * scale;
+ float s_pq = s * scale;
+ float max_lum = target_peak_inv_pq * scale;
+
+ float ks = 1.5f * max_lum - 0.5f;
@@ -420,7 +478,7 @@ Index: FFmpeg/libavfilter/metal/vf_tonemap_videotoolbox.metal
+ (-2.0f * tb3 + 3.0f * tb2) * max_lum;
+ float sig = mix(pb, s_pq, s_pq < ks);
+
+ return eotf_st2084(sig * peak_pq);
+ return sig * peak_pq;
+}
+
+float tonemap(float s, float peak, float target_peak) {
@@ -455,7 +513,7 @@ Index: FFmpeg/libavfilter/metal/vf_tonemap_videotoolbox.metal
+ return floor(y * dither_quantization + d + 0.5f / dither_size2) * 1.0f / dither_quantization;
+}
+
+void map_four_pixels_rgb(thread float4 *r4, thread float4 *g4, thread float4 *b4, float peak) {
+void map_four_pixels(thread float4 *r4, thread float4 *g4, thread float4 *b4, float peak) {
+#define MAP_FOUR_PIXELS(sig, peak, target_peak) \
+{ \
+ sig.x = tonemap(sig.x, peak, target_peak); \
@@ -467,14 +525,14 @@ Index: FFmpeg/libavfilter/metal/vf_tonemap_videotoolbox.metal
+ float4 sig_r = fmax(*r4, FLOAT_EPS);
+ float4 sig_g = fmax(*g4, FLOAT_EPS);
+ float4 sig_b = fmax(*b4, FLOAT_EPS);
+ if (is_tone_func_bt2390) {
+ sig_r = fmin(sig_r, peak);
+ sig_g = fmin(sig_g, peak);
+ sig_b = fmin(sig_b, peak);
+ }
+ float4 sig_ro = sig_r;
+ float4 sig_go = sig_g;
+ float4 sig_bo = sig_b;
+ if (is_tone_func_bt2390) {
+ sig_r = inverse_eotf_st2084x4(fmin(sig_r, peak));
+ sig_g = inverse_eotf_st2084x4(fmin(sig_g, peak));
+ sig_b = inverse_eotf_st2084x4(fmin(sig_b, peak));
+ }
+ // Desaturate the color using a coefficient dependent on the signal level
+ if (desat_param > 0.0f) {
+ float4 sig = fmax(fmax(*r4, fmax(*g4, *b4)), FLOAT_EPS);
@@ -491,6 +549,9 @@ Index: FFmpeg/libavfilter/metal/vf_tonemap_videotoolbox.metal
+ MAP_FOUR_PIXELS(sig_r, src_peak_delin_pq, dst_peak_delin_pq)
+ MAP_FOUR_PIXELS(sig_g, src_peak_delin_pq, dst_peak_delin_pq)
+ MAP_FOUR_PIXELS(sig_b, src_peak_delin_pq, dst_peak_delin_pq)
+ sig_r = fmin(eotf_st2084x4(sig_r), peak);
+ sig_g = fmin(eotf_st2084x4(sig_g), peak);
+ sig_b = fmin(eotf_st2084x4(sig_b), peak);
+ } else {
+ MAP_FOUR_PIXELS(sig_r, peak, 1.0f)
+ MAP_FOUR_PIXELS(sig_g, peak, 1.0f)
@@ -505,6 +566,30 @@ Index: FFmpeg/libavfilter/metal/vf_tonemap_videotoolbox.metal
+ *r4 *= factor_r;
+ *g4 *= factor_g;
+ *b4 *= factor_b;
+ } else if (is_tone_mode_itp) {
+ float4 i4_o, i4, ct4 , cp4;
+ lrgb2ictcp(*r4, *g4, *b4, &i4, &ct4, &cp4);
+ i4 = fmax(i4, FLOAT_EPS);
+ i4_o = i4;
+ if (desat_param > 0.0f) {
+ float4 coeff = exp(-pow(eotf_st2084x4(i4) - (target_peak - desat_param) * 0.5f, 2) / (2.0f * peak));
+ ct4 *= coeff;
+ cp4 *= coeff;
+ }
+ if (is_tone_func_bt2390) {
+ float src_peak_delin_pq = inverse_eotf_st2084(peak);
+ float dst_peak_delin_pq = inverse_eotf_st2084(1.0f);
+ MAP_FOUR_PIXELS(i4, src_peak_delin_pq, dst_peak_delin_pq)
+ } else {
+ i4 = eotf_st2084x4(i4);
+ MAP_FOUR_PIXELS(i4, peak, 1.0f)
+ i4 = inverse_eotf_st2084x4(i4);
+ }
+ i4 = fmin(i4, 1.0f);
+ float4 factor = min(i4/i4_o, i4_o/i4);
+ ct4 *= factor;
+ cp4 *= factor;
+ ictcp2lrgb(i4, ct4, cp4, r4, g4, b4);
+ } else {
+ float4 sig;
+ if (is_tone_mode_max) {
@@ -517,7 +602,12 @@ Index: FFmpeg/libavfilter/metal/vf_tonemap_videotoolbox.metal
+ }
+ float4 sig_o = sig;
+ if (desat_param > 0.0f) {
+ float4 luma = get_luma_dst4(*r4, *g4, *b4);
+ float4 luma;
+ if (is_tone_mode_max) {
+ luma = get_luma_dst4(*r4, *g4, *b4);
+ } else {
+ luma = sig;
+ }
+ float4 coeff = fmax(sig - 0.18f, FLOAT_EPS) / fmax(sig, FLOAT_EPS);
+ coeff = powr(coeff, 10.0f / desat_param);
+ *r4 = mix(*r4, luma, coeff);
@@ -527,7 +617,9 @@ Index: FFmpeg/libavfilter/metal/vf_tonemap_videotoolbox.metal
+ if (is_tone_func_bt2390) {
+ float src_peak_delin_pq = inverse_eotf_st2084(peak);
+ float dst_peak_delin_pq = inverse_eotf_st2084(1.0f);
+ sig = inverse_eotf_st2084x4(sig);
+ MAP_FOUR_PIXELS(sig, src_peak_delin_pq, dst_peak_delin_pq)
+ sig = fmin(eotf_st2084x4(sig), peak);
+ } else {
+ MAP_FOUR_PIXELS(sig, peak, 1.0f)
+ sig = fmin(sig, 1.0f);
@@ -539,13 +631,26 @@ Index: FFmpeg/libavfilter/metal/vf_tonemap_videotoolbox.metal
+ }
+}
+
+// Map from source space YUV to source space RGB
+float3 map_to_src_space_from_yuv(float3 yuv) {
+ if (dovi_reshape) {
+ float3 c = ycc2rgb(yuv.x, yuv.y, yuv.z);
+ c = lms2rgb(c.x, c.y, c.z);
+ c = rgb2lrgb(c);
+ return c;
+ } else {
+ float3 c = yuv2lrgb(yuv);
+ return c;
+ }
+}
+
+// Map from source space YUV to destination space RGB
+float3 map_to_dst_space_from_yuv(float3 yuv) {
+ if (dovi_reshape) {
+ float3 c = ycc2rgb(yuv.x, yuv.y, yuv.z);
+ c = lms2rgb(c.x, c.y, c.z);
+ c = rgb2lrgb(c);
+ return c;
+ return lrgb2lrgb(c);
+ } else {
+ float3 c = yuv2lrgb(yuv);
+ c = lrgb2lrgb(c);
@@ -722,22 +827,42 @@ Index: FFmpeg/libavfilter/metal/vf_tonemap_videotoolbox.metal
+ yuv3 = reshape_dovi_yuv(yuv3, dovi_params, dovi_pivots, dovi_coeffs, dovi_mmr);
+ }
+
+ float3 c0 = map_to_dst_space_from_yuv(yuv0);
+ float3 c1 = map_to_dst_space_from_yuv(yuv1);
+ float3 c2 = map_to_dst_space_from_yuv(yuv2);
+ float3 c3 = map_to_dst_space_from_yuv(yuv3);
+ float3 c0, c1, c2, c3;
+
+ if (map_in_src_space) {
+ c0 = map_to_src_space_from_yuv(yuv0);
+ c1 = map_to_src_space_from_yuv(yuv1);
+ c2 = map_to_src_space_from_yuv(yuv2);
+ c3 = map_to_src_space_from_yuv(yuv3);
+ } else {
+ c0 = map_to_dst_space_from_yuv(yuv0);
+ c1 = map_to_dst_space_from_yuv(yuv1);
+ c2 = map_to_dst_space_from_yuv(yuv2);
+ c3 = map_to_dst_space_from_yuv(yuv3);
+ }
+
+ if(!skip_tonemap) {
+ float4 r4 = float4(c0.x, c1.x, c2.x, c3.x);
+ float4 g4 = float4(c0.y, c1.y, c2.y, c3.y);
+ float4 b4 = float4(c0.z, c1.z, c2.z, c3.z);
+ map_four_pixels_rgb(&r4, &g4, &b4, *peak);
+ map_four_pixels(&r4, &g4, &b4, *peak);
+ c0 = float3(r4.x, g4.x, b4.x);
+ c1 = float3(r4.y, g4.y, b4.y);
+ c2 = float3(r4.z, g4.z, b4.z);
+ c3 = float3(r4.w, g4.w, b4.w);
+ }
+
+ if (map_in_src_space) {
+ c0 = lrgb2lrgb(c0);
+ c1 = lrgb2lrgb(c1);
+ c2 = lrgb2lrgb(c2);
+ c3 = lrgb2lrgb(c3);
+ c0 = clamp(c0, 0.0f, 1.0f);
+ c1 = clamp(c1, 0.0f, 1.0f);
+ c2 = clamp(c2, 0.0f, 1.0f);
+ c3 = clamp(c3, 0.0f, 1.0f);
+ }
+
+ float y0 = lrgb2y(c0);
+ float y1 = lrgb2y(c1);
+ float y2 = lrgb2y(c2);
@@ -770,7 +895,7 @@ Index: FFmpeg/libavfilter/vf_tonemap_videotoolbox.m
===================================================================
--- /dev/null
+++ FFmpeg/libavfilter/vf_tonemap_videotoolbox.m
@@ -0,0 +1,1136 @@
@@ -0,0 +1,1148 @@
+/*
+ * Copyright (c) 2024 Gnattu OC <gnattuoc@me.com>
+ *
@@ -806,6 +931,7 @@ Index: FFmpeg/libavfilter/vf_tonemap_videotoolbox.m
+#include "colorspace.h"
+#include "dither_matrix.h"
+#include "metal/utils.h"
+#include "libavutil/hwcontext_videotoolbox.h"
+
+#define params_cnt 8
+#define pivots_cnt (7+1)
@@ -840,6 +966,7 @@ Index: FFmpeg/libavfilter/vf_tonemap_videotoolbox.m
+ TONEMAP_MODE_MAX,
+ TONEMAP_MODE_RGB,
+ TONEMAP_MODE_LUM,
+ TONEMAP_MODE_ITP,
+ TONEMAP_MODE_COUNT,
+};
+
@@ -1106,6 +1233,7 @@ Index: FFmpeg/libavfilter/vf_tonemap_videotoolbox.m
+ bool is_tone_func_bt2390;
+ bool is_tone_mode_rgb;
+ bool is_tone_mode_max;
+ bool is_tone_mode_itp;
+ bool is_non_semi_planar_in;
+ bool is_non_semi_planar_out;
+ bool enable_dither;
@@ -1116,6 +1244,7 @@ Index: FFmpeg/libavfilter/vf_tonemap_videotoolbox.m
+ int chroma_loc;
+ bool skip_tonemap;
+ bool dovi_reshape;
+ bool map_in_src_space;
+
+ int i, j, err;
+ NSError* ns_error = nil;
@@ -1243,6 +1372,7 @@ Index: FFmpeg/libavfilter/vf_tonemap_videotoolbox.m
+ is_tone_func_bt2390 = ctx->tonemap == TONEMAP_BT2390;
+ is_tone_mode_rgb = ctx->tonemap_mode == TONEMAP_MODE_RGB;
+ is_tone_mode_max = ctx->tonemap_mode == TONEMAP_MODE_MAX;
+ is_tone_mode_itp = ctx->tonemap_mode == TONEMAP_MODE_ITP;
+ is_non_semi_planar_in = ctx->in_planes > 2;
+ is_non_semi_planar_out = ctx->out_planes > 2;
+ enable_dither = ctx->in_desc->comp[0].depth > ctx->out_desc->comp[0].depth;
@@ -1253,6 +1383,7 @@ Index: FFmpeg/libavfilter/vf_tonemap_videotoolbox.m
+ chroma_loc = (int)ctx->chroma_loc;
+ skip_tonemap = ctx->trc_out == AVCOL_TRC_SMPTE2084;
+ dovi_reshape = !!ctx->dovi;
+ map_in_src_space = !is_tone_mode_rgb && !is_tone_mode_max;
+
+ [constant_values setConstantValue:&ref_white type:MTLDataTypeFloat withName:@"ref_white"];
+ [constant_values setConstantValue:&tone_param type:MTLDataTypeFloat withName:@"tone_param"];
@@ -1266,6 +1397,7 @@ Index: FFmpeg/libavfilter/vf_tonemap_videotoolbox.m
+ [constant_values setConstantValue:&is_tone_func_bt2390 type:MTLDataTypeBool withName:@"is_tone_func_bt2390"];
+ [constant_values setConstantValue:&is_tone_mode_rgb type:MTLDataTypeBool withName:@"is_tone_mode_rgb"];
+ [constant_values setConstantValue:&is_tone_mode_max type:MTLDataTypeBool withName:@"is_tone_mode_max"];
+ [constant_values setConstantValue:&is_tone_mode_itp type:MTLDataTypeBool withName:@"is_tone_mode_itp"];
+
+ [constant_values setConstantValue:&is_non_semi_planar_in type:MTLDataTypeBool withName:@"is_non_semi_planar_in"];
+ [constant_values setConstantValue:&is_non_semi_planar_out type:MTLDataTypeBool withName:@"is_non_semi_planar_out"];
@@ -1443,6 +1575,8 @@ Index: FFmpeg/libavfilter/vf_tonemap_videotoolbox.m
+ [command_buffer waitUntilCompleted];
+ }
+
+ [constant_values setConstantValue:&map_in_src_space type:MTLDataTypeBool withName:@"map_in_src_space"];
+
+ ctx->mtl_function = [ctx->mtl_library newFunctionWithName:@"tonemap" constantValues:constant_values error:&ns_error];
+ if (ns_error) {
+ av_log(ctx, AV_LOG_ERROR, "Failed to create Metal function: %s\n", ns_error.description.UTF8String);
@@ -1787,14 +1921,11 @@ Index: FFmpeg/libavfilter/vf_tonemap_videotoolbox.m
+ CVBufferPropagateAttachments((CVPixelBufferRef)input->data[3], (CVPixelBufferRef)output->data[3]);
+ av_frame_free(&input);
+
+ if (ctx->trc_out != AVCOL_TRC_SMPTE2084) {
+ av_frame_remove_side_data(output, AV_FRAME_DATA_MASTERING_DISPLAY_METADATA);
+ av_frame_remove_side_data(output, AV_FRAME_DATA_CONTENT_LIGHT_LEVEL);
+ } else {
+ {
+ CGColorSpaceRef colorspace = NULL;
+ CFStringRef colormatrix = kCVImageBufferYCbCrMatrix_ITU_R_2020;
+ CFStringRef colorpri = kCVImageBufferColorPrimaries_ITU_R_2020;
+ CFStringRef colortrc = kCVImageBufferTransferFunction_SMPTE_ST_2084_PQ;
+ CFStringRef colormatrix = av_map_videotoolbox_color_matrix_from_av(ctx->colorspace_out);
+ CFStringRef colorpri = av_map_videotoolbox_color_primaries_from_av(ctx->primaries_out);
+ CFStringRef colortrc = av_map_videotoolbox_color_trc_from_av(ctx->trc_out);
+ CFMutableDictionaryRef attachments = CFDictionaryCreateMutable(NULL, 4,
+ &kCFTypeDictionaryKeyCallBacks,
+ &kCFTypeDictionaryValueCallBacks);
@@ -1817,7 +1948,12 @@ Index: FFmpeg/libavfilter/vf_tonemap_videotoolbox.m
+ attachments,
+ kCVAttachmentMode_ShouldPropagate);
+ CFRelease(attachments);
+ ff_update_hdr_metadata(output, 100.0f);
+ if (ctx->trc_out != AVCOL_TRC_SMPTE2084) {
+ av_frame_remove_side_data(output, AV_FRAME_DATA_MASTERING_DISPLAY_METADATA);
+ av_frame_remove_side_data(output, AV_FRAME_DATA_CONTENT_LIGHT_LEVEL);
+ } else {
+ ff_update_hdr_metadata(output, 100.0f);
+ }
+ }
+
+ av_frame_remove_side_data(output, AV_FRAME_DATA_DOVI_RPU_BUFFER);
@@ -1845,10 +1981,11 @@ Index: FFmpeg/libavfilter/vf_tonemap_videotoolbox.m
+ { "hable", 0, 0, AV_OPT_TYPE_CONST, { .i64 = TONEMAP_HABLE }, 0, 0, FLAGS, .unit = "tonemap" },
+ { "mobius", 0, 0, AV_OPT_TYPE_CONST, { .i64 = TONEMAP_MOBIUS }, 0, 0, FLAGS, .unit = "tonemap" },
+ { "bt2390", 0, 0, AV_OPT_TYPE_CONST, { .i64 = TONEMAP_BT2390 }, 0, 0, FLAGS, .unit = "tonemap" },
+ { "tonemap_mode", "Tonemap mode selection", OFFSET(tonemap_mode), AV_OPT_TYPE_INT, { .i64 = TONEMAP_MODE_MAX }, TONEMAP_MODE_MAX, TONEMAP_MODE_COUNT - 1, FLAGS, .unit = "tonemap_mode" },
+ { "tonemap_mode", "Tonemap mode selection", OFFSET(tonemap_mode), AV_OPT_TYPE_INT, { .i64 = TONEMAP_MODE_ITP }, TONEMAP_MODE_MAX, TONEMAP_MODE_COUNT - 1, FLAGS, .unit = "tonemap_mode" },
+ { "max", "Brightest channel based tonemap", 0, AV_OPT_TYPE_CONST, { .i64 = TONEMAP_MODE_MAX }, 0, 0, FLAGS, .unit = "tonemap_mode" },
+ { "rgb", "Per-channel based tonemap", 0, AV_OPT_TYPE_CONST, { .i64 = TONEMAP_MODE_RGB }, 0, 0, FLAGS, .unit = "tonemap_mode" },
+ { "lum", "Relative luminance based tonemap", 0, AV_OPT_TYPE_CONST, { .i64 = TONEMAP_MODE_LUM }, 0, 0, FLAGS, .unit = "tonemap_mode" },
+ { "itp", "ICtCp intensity based tonemap", 0, AV_OPT_TYPE_CONST, { .i64 = TONEMAP_MODE_ITP }, 0, 0, FLAGS, .unit = "tonemap_mode" },
+ { "transfer", "Set transfer characteristic", OFFSET(trc), AV_OPT_TYPE_INT, { .i64 = AVCOL_TRC_BT709 }, -1, INT_MAX, FLAGS, .unit = "transfer" },
+ { "t", "Set transfer characteristic", OFFSET(trc), AV_OPT_TYPE_INT, { .i64 = AVCOL_TRC_BT709 }, -1, INT_MAX, FLAGS, .unit = "transfer" },
+ { "bt709", 0, 0, AV_OPT_TYPE_CONST, { .i64 = AVCOL_TRC_BT709 }, 0, 0, FLAGS, .unit = "transfer" },