diff --git a/debian/patches/0005-add-cuda-tonemap-impl.patch b/debian/patches/0005-add-cuda-tonemap-impl.patch index a0a56afb533..e784474dcf4 100644 --- a/debian/patches/0005-add-cuda-tonemap-impl.patch +++ b/debian/patches/0005-add-cuda-tonemap-impl.patch @@ -1,7 +1,7 @@ -Index: jellyfin-ffmpeg/configure +Index: FFmpeg/configure =================================================================== ---- jellyfin-ffmpeg.orig/configure -+++ jellyfin-ffmpeg/configure +--- FFmpeg.orig/configure ++++ FFmpeg/configure @@ -3143,6 +3143,8 @@ scale_cuda_filter_deps="ffnvcodec" scale_cuda_filter_deps_any="cuda_nvcc cuda_llvm" thumbnail_cuda_filter_deps="ffnvcodec" @@ -38,10 +38,10 @@ Index: jellyfin-ffmpeg/configure check_nvcc cuda_llvm fi -Index: jellyfin-ffmpeg/ffbuild/common.mak +Index: FFmpeg/ffbuild/common.mak =================================================================== ---- jellyfin-ffmpeg.orig/ffbuild/common.mak -+++ jellyfin-ffmpeg/ffbuild/common.mak +--- FFmpeg.orig/ffbuild/common.mak ++++ FFmpeg/ffbuild/common.mak @@ -44,6 +44,7 @@ ASFLAGS := $(CPPFLAGS) $(ASFLAGS) # end up in CXXFLAGS. $(call PREPEND,CXXFLAGS, CPPFLAGS CFLAGS) @@ -50,10 +50,10 @@ Index: jellyfin-ffmpeg/ffbuild/common.mak HOSTCCFLAGS = $(IFLAGS) $(HOSTCPPFLAGS) $(HOSTCFLAGS) LDFLAGS := $(ALLFFLIBS:%=$(LD_PATH)lib%) $(LDFLAGS) -Index: jellyfin-ffmpeg/libavfilter/Makefile +Index: FFmpeg/libavfilter/Makefile =================================================================== ---- jellyfin-ffmpeg.orig/libavfilter/Makefile -+++ jellyfin-ffmpeg/libavfilter/Makefile +--- FFmpeg.orig/libavfilter/Makefile ++++ FFmpeg/libavfilter/Makefile @@ -509,6 +509,8 @@ OBJS-$(CONFIG_TMIX_FILTER) OBJS-$(CONFIG_TONEMAP_FILTER) += vf_tonemap.o OBJS-$(CONFIG_TONEMAP_OPENCL_FILTER) += vf_tonemap_opencl.o opencl.o \ @@ -63,10 +63,10 @@ Index: jellyfin-ffmpeg/libavfilter/Makefile OBJS-$(CONFIG_TONEMAP_VAAPI_FILTER) += vf_tonemap_vaapi.o vaapi_vpp.o OBJS-$(CONFIG_TPAD_FILTER) += vf_tpad.o OBJS-$(CONFIG_TRANSPOSE_FILTER) += vf_transpose.o -Index: jellyfin-ffmpeg/libavfilter/allfilters.c +Index: FFmpeg/libavfilter/allfilters.c =================================================================== ---- jellyfin-ffmpeg.orig/libavfilter/allfilters.c -+++ jellyfin-ffmpeg/libavfilter/allfilters.c +--- FFmpeg.orig/libavfilter/allfilters.c ++++ FFmpeg/libavfilter/allfilters.c @@ -478,6 +478,7 @@ extern const AVFilter ff_vf_tmedian; extern const AVFilter ff_vf_tmidequalizer; extern const AVFilter ff_vf_tmix; @@ -75,10 +75,10 @@ Index: jellyfin-ffmpeg/libavfilter/allfilters.c extern const AVFilter ff_vf_tonemap_opencl; extern const AVFilter ff_vf_tonemap_vaapi; extern const AVFilter ff_vf_tpad; -Index: jellyfin-ffmpeg/libavfilter/colorspace.c +Index: FFmpeg/libavfilter/colorspace.c =================================================================== ---- jellyfin-ffmpeg.orig/libavfilter/colorspace.c -+++ jellyfin-ffmpeg/libavfilter/colorspace.c +--- FFmpeg.orig/libavfilter/colorspace.c ++++ FFmpeg/libavfilter/colorspace.c @@ -51,6 +51,18 @@ void ff_matrix_invert_3x3(const double i } } @@ -253,10 +253,10 @@ Index: jellyfin-ffmpeg/libavfilter/colorspace.c +float inverse_eotf_bt1886(float x) { + return x > 0.0f ? powf(x, 1.0f / 2.4f) : 0.0f; +} -Index: jellyfin-ffmpeg/libavfilter/colorspace.h +Index: FFmpeg/libavfilter/colorspace.h =================================================================== ---- jellyfin-ffmpeg.orig/libavfilter/colorspace.h -+++ jellyfin-ffmpeg/libavfilter/colorspace.h +--- FFmpeg.orig/libavfilter/colorspace.h ++++ FFmpeg/libavfilter/colorspace.h @@ -23,10 +23,42 @@ #include "libavutil/csp.h" #include "libavutil/frame.h" @@ -320,10 +320,10 @@ Index: jellyfin-ffmpeg/libavfilter/colorspace.h +float inverse_eotf_bt1886(float x); + #endif -Index: jellyfin-ffmpeg/libavfilter/cuda/colorspace_common.h +Index: FFmpeg/libavfilter/cuda/colorspace_common.h =================================================================== --- /dev/null -+++ jellyfin-ffmpeg/libavfilter/cuda/colorspace_common.h ++++ FFmpeg/libavfilter/cuda/colorspace_common.h @@ -0,0 +1,267 @@ +/* + * This file is part of FFmpeg. @@ -592,10 +592,10 @@ Index: jellyfin-ffmpeg/libavfilter/cuda/colorspace_common.h +} + +#endif /* AVFILTER_CUDA_COLORSPACE_COMMON_H */ -Index: jellyfin-ffmpeg/libavfilter/cuda/host_util.c +Index: FFmpeg/libavfilter/cuda/host_util.c =================================================================== --- /dev/null -+++ jellyfin-ffmpeg/libavfilter/cuda/host_util.c ++++ FFmpeg/libavfilter/cuda/host_util.c @@ -0,0 +1,77 @@ +/* + * This file is part of FFmpeg. @@ -674,10 +674,10 @@ Index: jellyfin-ffmpeg/libavfilter/cuda/host_util.c + + return ret; +} -Index: jellyfin-ffmpeg/libavfilter/cuda/host_util.h +Index: FFmpeg/libavfilter/cuda/host_util.h =================================================================== --- /dev/null -+++ jellyfin-ffmpeg/libavfilter/cuda/host_util.h ++++ FFmpeg/libavfilter/cuda/host_util.h @@ -0,0 +1,30 @@ +/* + * This file is part of FFmpeg. @@ -709,10 +709,10 @@ Index: jellyfin-ffmpeg/libavfilter/cuda/host_util.h + FFCUDAFrame *dst, const AVFrame *src, const AVPixFmtDescriptor *src_desc); + +#endif /* AVFILTER_CUDA_HOST_UTIL_H */ -Index: jellyfin-ffmpeg/libavfilter/cuda/pixfmt.h +Index: FFmpeg/libavfilter/cuda/pixfmt.h =================================================================== --- /dev/null -+++ jellyfin-ffmpeg/libavfilter/cuda/pixfmt.h ++++ FFmpeg/libavfilter/cuda/pixfmt.h @@ -0,0 +1,225 @@ +/* + * This file is part of FFmpeg. @@ -939,10 +939,10 @@ Index: jellyfin-ffmpeg/libavfilter/cuda/pixfmt.h +} + +#endif /* AVFILTER_CUDA_PIXFMT_H */ -Index: jellyfin-ffmpeg/libavfilter/cuda/shared.h +Index: FFmpeg/libavfilter/cuda/shared.h =================================================================== --- /dev/null -+++ jellyfin-ffmpeg/libavfilter/cuda/shared.h ++++ FFmpeg/libavfilter/cuda/shared.h @@ -0,0 +1,33 @@ +/* + * This file is part of FFmpeg. @@ -977,11 +977,11 @@ Index: jellyfin-ffmpeg/libavfilter/cuda/shared.h +} FFCUDAFrame; + +#endif /* AVFILTER_CUDA_SHARED_H */ -Index: jellyfin-ffmpeg/libavfilter/cuda/tonemap.cu +Index: FFmpeg/libavfilter/cuda/tonemap.cu =================================================================== --- /dev/null -+++ jellyfin-ffmpeg/libavfilter/cuda/tonemap.cu -@@ -0,0 +1,418 @@ ++++ FFmpeg/libavfilter/cuda/tonemap.cu +@@ -0,0 +1,455 @@ +/* + * This file is part of FFmpeg. + * @@ -1170,6 +1170,31 @@ Index: jellyfin-ffmpeg/libavfilter/cuda/tonemap.cu + return rgb; +} + ++static __inline__ __device__ ++float3 map_one_pixel_rgb_mode_rl(float3 rgb, const FFCUDAFrame& src, const FFCUDAFrame& dst) { ++ float sig = max((rgb.x * 0.2126f + rgb.y * 0.7152f + rgb.z * 0.0722f), FLOAT_EPS); ++ float peak = src.peak; ++ sig = min(sig, peak); ++ float sig_old = sig; ++ float dst_peak = 1.0f; ++ ++ // 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)); ++ } ++ ++ 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; ++} ++ +// Map from source space YUV to destination space RGB +static __inline__ __device__ +float3 map_to_dst_space_from_yuv(float3 yuv) { @@ -1351,6 +1376,12 @@ Index: jellyfin-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 _RGB2YUV \ + yuv0 = lrgb2yuv(c0); \ + yuv1 = lrgb2yuv(c1); \ @@ -1385,26 +1416,32 @@ Index: jellyfin-ffmpeg/libavfilter/cuda/tonemap.cu +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(_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_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_pq, _READER, _RESHAPE, _YCC2RGB, , _RGB2YUV, , _WRITER) +TONEMAP_VARIANT(_dovi_pq_f, _READER, _RESHAPE, _YCC2RGB_F, , _RGB2YUV, , _WRITER) + +} -Index: jellyfin-ffmpeg/libavfilter/cuda/tonemap.h +Index: FFmpeg/libavfilter/cuda/tonemap.h =================================================================== --- /dev/null -+++ jellyfin-ffmpeg/libavfilter/cuda/tonemap.h -@@ -0,0 +1,40 @@ ++++ FFmpeg/libavfilter/cuda/tonemap.h +@@ -0,0 +1,41 @@ +/* + * This file is part of FFmpeg. + * @@ -1441,14 +1478,15 @@ Index: jellyfin-ffmpeg/libavfilter/cuda/tonemap.h +enum TonemapMode { + TONEMAP_MODE_MAX, + TONEMAP_MODE_RGB, ++ TONEMAP_MODE_RELATIVE_LUMINANCE, + TONEMAP_MODE_COUNT, +}; + +#endif /* AVFILTER_CUDA_TONEMAP_H */ -Index: jellyfin-ffmpeg/libavfilter/cuda/util.h +Index: FFmpeg/libavfilter/cuda/util.h =================================================================== --- /dev/null -+++ jellyfin-ffmpeg/libavfilter/cuda/util.h ++++ FFmpeg/libavfilter/cuda/util.h @@ -0,0 +1,86 @@ +/* + * This file is part of FFmpeg. @@ -1536,11 +1574,11 @@ Index: jellyfin-ffmpeg/libavfilter/cuda/util.h +} + +#endif /* AVFILTER_CUDA_UTIL_H */ -Index: jellyfin-ffmpeg/libavfilter/vf_tonemap_cuda.c +Index: FFmpeg/libavfilter/vf_tonemap_cuda.c =================================================================== --- /dev/null -+++ jellyfin-ffmpeg/libavfilter/vf_tonemap_cuda.c -@@ -0,0 +1,1096 @@ ++++ FFmpeg/libavfilter/vf_tonemap_cuda.c +@@ -0,0 +1,1101 @@ +/* + * This file is part of FFmpeg. + * @@ -2054,6 +2092,7 @@ Index: jellyfin-ffmpeg/libavfilter/vf_tonemap_cuda.c + 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, @@ -2286,15 +2325,18 @@ Index: jellyfin-ffmpeg/libavfilter/vf_tonemap_cuda.c + + ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_tm, s->cu_module, + rgb ? (d ? "tonemap_rgb_d" : "tonemap_rgb") -+ : (d ? "tonemap_d" : "tonemap"))); ++ : (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") -+ : (d ? "tonemap_dovi_d_f" : "tonemap_dovi_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") -+ : (d ? "tonemap_dovi_d" : "tonemap_dovi")))); ++ : (max ? (d ? "tonemap_dovi_d" : "tonemap_dovi") ++ : (d ? "tonemap_dovi_rl_d" : "tonemap_dovi_rl"))))); + if (ret < 0) + goto fail2; + @@ -2565,6 +2607,7 @@ Index: jellyfin-ffmpeg/libavfilter/vf_tonemap_cuda.c + { "tonemap_mode", "Tonemap mode selection", OFFSET(tonemap_mode), AV_OPT_TYPE_INT, {.i64 = TONEMAP_MODE_MAX}, TONEMAP_MODE_MAX, TONEMAP_MODE_COUNT - 1, FLAGS, "tonemap_mode" }, + { "max", 0, 0, AV_OPT_TYPE_CONST, {.i64 = TONEMAP_MODE_MAX}, 0, 0, FLAGS, "tonemap_mode" }, + { "rgb", 0, 0, AV_OPT_TYPE_CONST, {.i64 = TONEMAP_MODE_RGB}, 0, 0, FLAGS, "tonemap_mode" }, ++ { "relative_luminance", 0, 0, AV_OPT_TYPE_CONST, {.i64 = TONEMAP_MODE_RELATIVE_LUMINANCE}, 0, 0, FLAGS, "tonemap_mode" }, + { "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" }, diff --git a/debian/patches/0008-add-bt2390-eetf-and-code-refactor-to-opencl-tonemap.patch b/debian/patches/0008-add-bt2390-eetf-and-code-refactor-to-opencl-tonemap.patch index 63f40955de6..84ce736fe62 100644 --- a/debian/patches/0008-add-bt2390-eetf-and-code-refactor-to-opencl-tonemap.patch +++ b/debian/patches/0008-add-bt2390-eetf-and-code-refactor-to-opencl-tonemap.patch @@ -1,7 +1,7 @@ -Index: jellyfin-ffmpeg/libavfilter/opencl.c +Index: FFmpeg/libavfilter/opencl.c =================================================================== ---- jellyfin-ffmpeg.orig/libavfilter/opencl.c -+++ jellyfin-ffmpeg/libavfilter/opencl.c +--- FFmpeg.orig/libavfilter/opencl.c ++++ FFmpeg/libavfilter/opencl.c @@ -170,7 +170,7 @@ int ff_opencl_filter_load_program(AVFilt } @@ -20,10 +20,10 @@ Index: jellyfin-ffmpeg/libavfilter/opencl.c av_bprintf(buf, "\n"); } av_bprintf(buf, "};\n"); -Index: jellyfin-ffmpeg/libavfilter/opencl.h +Index: FFmpeg/libavfilter/opencl.h =================================================================== ---- jellyfin-ffmpeg.orig/libavfilter/opencl.h -+++ jellyfin-ffmpeg/libavfilter/opencl.h +--- FFmpeg.orig/libavfilter/opencl.h ++++ FFmpeg/libavfilter/opencl.h @@ -206,17 +206,17 @@ do { } while(0) @@ -61,10 +61,10 @@ Index: jellyfin-ffmpeg/libavfilter/opencl.h * Create a buffer with the given information. * * The buffer variable in the context structure must be named . -Index: jellyfin-ffmpeg/libavfilter/opencl/colorspace_common.cl +Index: FFmpeg/libavfilter/opencl/colorspace_common.cl =================================================================== ---- jellyfin-ffmpeg.orig/libavfilter/opencl/colorspace_common.cl -+++ jellyfin-ffmpeg/libavfilter/opencl/colorspace_common.cl +--- FFmpeg.orig/libavfilter/opencl/colorspace_common.cl ++++ FFmpeg/libavfilter/opencl/colorspace_common.cl @@ -17,7 +17,17 @@ */ @@ -298,10 +298,10 @@ Index: jellyfin-ffmpeg/libavfilter/opencl/colorspace_common.cl + return (float3)(rr, gg, bb); } +#endif -Index: jellyfin-ffmpeg/libavfilter/opencl/tonemap.cl +Index: FFmpeg/libavfilter/opencl/tonemap.cl =================================================================== ---- jellyfin-ffmpeg.orig/libavfilter/opencl/tonemap.cl -+++ jellyfin-ffmpeg/libavfilter/opencl/tonemap.cl +--- FFmpeg.orig/libavfilter/opencl/tonemap.cl ++++ FFmpeg/libavfilter/opencl/tonemap.cl @@ -16,54 +16,60 @@ * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA */ @@ -384,7 +384,7 @@ Index: jellyfin-ffmpeg/libavfilter/opencl/tonemap.cl float j = tone_param; float a, b; -@@ -71,202 +77,335 @@ float mobius(float s, float peak) { +@@ -71,202 +77,339 @@ float mobius(float s, float peak) { return s; a = -j * j * (peak - 1.0f) / (j * j - 2.0f * j + peak); @@ -467,7 +467,11 @@ Index: jellyfin-ffmpeg/libavfilter/opencl/tonemap.cl + float4 sig_g = fmax(*g4, FLOAT_EPS), sig_go = sig_g; + float4 sig_b = fmax(*b4, FLOAT_EPS), sig_bo = sig_b; +#else ++ #ifdef TONE_MODE_MAX + float4 sig = fmax(fmax(*r4, fmax(*g4, *b4)), FLOAT_EPS); ++ #else ++ float4 sig = fmax((*r4 * 0.2126f + *g4 * 0.7152f + *b4 * 0.0722f), FLOAT_EPS); ++ #endif + float4 sig_o = sig; +#endif @@ -886,10 +890,10 @@ Index: jellyfin-ffmpeg/libavfilter/opencl/tonemap.cl + write_imagef(dst2, (int2)(xi, yi), (float4)(chroma.y, chroma.z, 0.0f, 1.0f)); +#endif } -Index: jellyfin-ffmpeg/libavfilter/vf_tonemap_opencl.c +Index: FFmpeg/libavfilter/vf_tonemap_opencl.c =================================================================== ---- jellyfin-ffmpeg.orig/libavfilter/vf_tonemap_opencl.c -+++ jellyfin-ffmpeg/libavfilter/vf_tonemap_opencl.c +--- FFmpeg.orig/libavfilter/vf_tonemap_opencl.c ++++ FFmpeg/libavfilter/vf_tonemap_opencl.c @@ -1,4 +1,4 @@ -/* + /* @@ -935,7 +939,7 @@ Index: jellyfin-ffmpeg/libavfilter/vf_tonemap_opencl.c enum TonemapAlgorithm { TONEMAP_NONE, -@@ -45,7 +56,14 @@ enum TonemapAlgorithm { +@@ -45,7 +56,15 @@ enum TonemapAlgorithm { TONEMAP_REINHARD, TONEMAP_HABLE, TONEMAP_MOBIUS, @@ -947,11 +951,12 @@ Index: jellyfin-ffmpeg/libavfilter/vf_tonemap_opencl.c +enum TonemapMode { + TONEMAP_MODE_MAX, + TONEMAP_MODE_RGB, ++ TONEMAP_MODE_RELATIVE_LUMINANCE, + TONEMAP_MODE_COUNT, }; typedef struct TonemapOpenCLContext { -@@ -56,23 +74,44 @@ typedef struct TonemapOpenCLContext { +@@ -56,23 +75,44 @@ typedef struct TonemapOpenCLContext { enum AVColorPrimaries primaries, primaries_in, primaries_out; enum AVColorRange range, range_in, range_out; enum AVChromaLocation chroma_loc; @@ -1000,7 +1005,7 @@ Index: jellyfin-ffmpeg/libavfilter/vf_tonemap_opencl.c }; static const char *const delinearize_funcs[AVCOL_TRC_NB] = { -@@ -80,7 +119,7 @@ static const char *const delinearize_fun +@@ -80,7 +120,7 @@ static const char *const delinearize_fun [AVCOL_TRC_BT2020_10] = "inverse_eotf_bt1886", }; @@ -1009,7 +1014,7 @@ Index: jellyfin-ffmpeg/libavfilter/vf_tonemap_opencl.c [TONEMAP_NONE] = "direct", [TONEMAP_LINEAR] = "linear", [TONEMAP_GAMMA] = "gamma", -@@ -88,8 +127,54 @@ static const char *const tonemap_func[TO +@@ -88,8 +128,54 @@ static const char *const tonemap_func[TO [TONEMAP_REINHARD] = "reinhard", [TONEMAP_HABLE] = "hable", [TONEMAP_MOBIUS] = "mobius", @@ -1064,7 +1069,7 @@ Index: jellyfin-ffmpeg/libavfilter/vf_tonemap_opencl.c static int get_rgb2rgb_matrix(enum AVColorPrimaries in, enum AVColorPrimaries out, double rgb2rgb[3][3]) { double rgb2xyz[3][3], xyz2rgb[3][3]; -@@ -108,23 +193,149 @@ static int get_rgb2rgb_matrix(enum AVCol +@@ -108,23 +194,149 @@ static int get_rgb2rgb_matrix(enum AVCol return 0; } @@ -1223,7 +1228,7 @@ Index: jellyfin-ffmpeg/libavfilter/vf_tonemap_opencl.c switch(ctx->tonemap) { case TONEMAP_GAMMA: -@@ -144,48 +355,148 @@ static int tonemap_opencl_init(AVFilterC +@@ -144,48 +356,148 @@ static int tonemap_opencl_init(AVFilterC if (isnan(ctx->param)) ctx->param = 1.0f; @@ -1363,7 +1368,7 @@ Index: jellyfin-ffmpeg/libavfilter/vf_tonemap_opencl.c + + if (ctx->tonemap_mode == TONEMAP_MODE_RGB) + av_bprintf(&header, "#define TONE_MODE_RGB\n"); -+ else ++ else if (ctx->tonemap_mode == TONEMAP_MODE_MAX) + av_bprintf(&header, "#define TONE_MODE_MAX\n"); + + if (ctx->in_planes > 2) @@ -1387,7 +1392,7 @@ Index: jellyfin-ffmpeg/libavfilter/vf_tonemap_opencl.c if (ctx->range_in == AVCOL_RANGE_JPEG) av_bprintf(&header, "#define FULL_RANGE_IN\n"); -@@ -199,19 +510,41 @@ static int tonemap_opencl_init(AVFilterC +@@ -199,19 +511,41 @@ static int tonemap_opencl_init(AVFilterC else ff_opencl_print_const_matrix_3x3(&header, "rgb2rgb", rgb2rgb); @@ -1436,7 +1441,7 @@ Index: jellyfin-ffmpeg/libavfilter/vf_tonemap_opencl.c ctx->colorspace_out, av_color_space_name(ctx->colorspace_out)); goto fail; } -@@ -219,24 +552,23 @@ static int tonemap_opencl_init(AVFilterC +@@ -219,24 +553,23 @@ static int tonemap_opencl_init(AVFilterC ff_fill_rgb2yuv_table(luma_dst, rgb2yuv); ff_opencl_print_const_matrix_3x3(&header, "yuv_matrix", rgb2yuv); @@ -1476,7 +1481,7 @@ Index: jellyfin-ffmpeg/libavfilter/vf_tonemap_opencl.c av_log(avctx, AV_LOG_DEBUG, "Generated OpenCL header:\n%s\n", header.str); opencl_sources[0] = header.str; -@@ -254,46 +586,171 @@ static int tonemap_opencl_init(AVFilterC +@@ -254,46 +587,171 @@ static int tonemap_opencl_init(AVFilterC CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create OpenCL " "command queue %d.\n", cle); @@ -1667,7 +1672,7 @@ Index: jellyfin-ffmpeg/libavfilter/vf_tonemap_opencl.c ret = ff_opencl_filter_config_output(outlink); if (ret < 0) return ret; -@@ -308,13 +765,46 @@ static int launch_kernel(AVFilterContext +@@ -308,13 +766,46 @@ static int launch_kernel(AVFilterContext size_t global_work[2]; size_t local_work[2]; cl_int cle; @@ -1716,7 +1721,7 @@ Index: jellyfin-ffmpeg/libavfilter/vf_tonemap_opencl.c local_work[0] = 16; local_work[1] = 16; -@@ -338,13 +828,10 @@ static int tonemap_opencl_filter_frame(A +@@ -338,13 +829,10 @@ static int tonemap_opencl_filter_frame(A AVFilterContext *avctx = inlink->dst; AVFilterLink *outlink = avctx->outputs[0]; TonemapOpenCLContext *ctx = avctx->priv; @@ -1731,7 +1736,7 @@ Index: jellyfin-ffmpeg/libavfilter/vf_tonemap_opencl.c av_log(ctx, AV_LOG_DEBUG, "Filter input: %s, %ux%u (%"PRId64").\n", av_get_pix_fmt_name(input->format), -@@ -363,9 +850,6 @@ static int tonemap_opencl_filter_frame(A +@@ -363,9 +851,6 @@ static int tonemap_opencl_filter_frame(A if (err < 0) goto fail; @@ -1741,7 +1746,7 @@ Index: jellyfin-ffmpeg/libavfilter/vf_tonemap_opencl.c if (ctx->trc != -1) output->color_trc = ctx->trc; if (ctx->primaries != -1) -@@ -385,72 +869,92 @@ static int tonemap_opencl_filter_frame(A +@@ -385,72 +870,92 @@ static int tonemap_opencl_filter_frame(A ctx->range_out = output->color_range; ctx->chroma_loc = output->chroma_location; @@ -1874,7 +1879,7 @@ Index: jellyfin-ffmpeg/libavfilter/vf_tonemap_opencl.c av_frame_free(&input); av_frame_free(&output); return err; -@@ -458,24 +962,9 @@ fail: +@@ -458,24 +963,9 @@ fail: static av_cold void tonemap_opencl_uninit(AVFilterContext *avctx) { @@ -1901,7 +1906,7 @@ Index: jellyfin-ffmpeg/libavfilter/vf_tonemap_opencl.c ff_opencl_filter_uninit(avctx); } -@@ -483,37 +972,47 @@ static av_cold void tonemap_opencl_unini +@@ -483,37 +973,48 @@ static av_cold void tonemap_opencl_unini #define OFFSET(x) offsetof(TonemapOpenCLContext, x) #define FLAGS (AV_OPT_FLAG_FILTERING_PARAM | AV_OPT_FLAG_VIDEO_PARAM) static const AVOption tonemap_opencl_options[] = { @@ -1948,6 +1953,7 @@ Index: jellyfin-ffmpeg/libavfilter/vf_tonemap_opencl.c + { "tonemap_mode", "Tonemap mode selection", OFFSET(tonemap_mode), AV_OPT_TYPE_INT, { .i64 = TONEMAP_MODE_MAX }, TONEMAP_MODE_MAX, TONEMAP_MODE_COUNT - 1, FLAGS, "tonemap_mode" }, + { "max", 0, 0, AV_OPT_TYPE_CONST, { .i64 = TONEMAP_MODE_MAX }, 0, 0, FLAGS, "tonemap_mode" }, + { "rgb", 0, 0, AV_OPT_TYPE_CONST, { .i64 = TONEMAP_MODE_RGB }, 0, 0, FLAGS, "tonemap_mode" }, ++ { "relative_luminance", 0, 0, AV_OPT_TYPE_CONST, { .i64 = TONEMAP_MODE_RELATIVE_LUMINANCE }, 0, 0, FLAGS, "tonemap_mode" }, + { "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" },