From b1e95436b8e7cabb9636408a32d5e9224916f4b3 Mon Sep 17 00:00:00 2001 From: nyanmisaka Date: Mon, 27 Mar 2023 02:21:55 +0800 Subject: [PATCH 1/7] Add a tone-mapping mode RGB in CUDA & OpenCL Signed-off-by: nyanmisaka --- .../patches/0005-add-cuda-tonemap-impl.patch | 199 ++++++------- ...-and-code-refactor-to-opencl-tonemap.patch | 261 +++++++++++------- 2 files changed, 249 insertions(+), 211 deletions(-) diff --git a/debian/patches/0005-add-cuda-tonemap-impl.patch b/debian/patches/0005-add-cuda-tonemap-impl.patch index f920e314a2..4fa08718c1 100644 --- a/debian/patches/0005-add-cuda-tonemap-impl.patch +++ b/debian/patches/0005-add-cuda-tonemap-impl.patch @@ -985,7 +985,7 @@ Index: jellyfin-ffmpeg/libavfilter/cuda/tonemap.cu =================================================================== --- /dev/null +++ jellyfin-ffmpeg/libavfilter/cuda/tonemap.cu -@@ -0,0 +1,449 @@ +@@ -0,0 +1,418 @@ +/* + * This file is part of FFmpeg. + * @@ -1123,24 +1123,12 @@ Index: jellyfin-ffmpeg/libavfilter/cuda/tonemap.cu +} + +static __inline__ __device__ -+float3 map_one_pixel_rgb(float3 rgb, const FFCUDAFrame& src, const FFCUDAFrame& dst) { ++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; + float peak = src.peak; + float dst_peak = 1.0f; + -+/* -+ // 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 (dst.peak > 1.0f) { -+ float scale = 1.0f / dst.peak; -+ sig *= scale; -+ peak *= scale; -+ } -+*/ -+ -+ float sig_old = sig; -+ + // Desaturate the color using a coefficient dependent on the signal level + if (desat_param > 0.0f) { + float luma = get_luma_dst(rgb, luma_dst); @@ -1156,6 +1144,36 @@ Index: jellyfin-ffmpeg/libavfilter/cuda/tonemap.cu + return rgb; +} + ++static __inline__ __device__ ++float3 map_one_pixel_rgb_mode_rgb(float3 rgb, const FFCUDAFrame& src, const FFCUDAFrame& dst) { ++ float3 sig; ++ sig.x = max(rgb.x, FLOAT_EPS); ++ sig.y = max(rgb.y, FLOAT_EPS); ++ sig.z = max(rgb.z, FLOAT_EPS); ++ float3 sig_old = sig; ++ float peak = src.peak; ++ float dst_peak = 1.0f; ++ ++ // Desaturate the color using a coefficient dependent on the signal level ++ if (desat_param > 0.0f) { ++ float sig_max = max(max(rgb.x, max(rgb.y, rgb.z)), FLOAT_EPS); ++ float luma = get_luma_dst(rgb, luma_dst); ++ float coeff = max(sig_max - 0.18f, FLOAT_EPS) / max(sig_max, FLOAT_EPS); ++ coeff = __powf(coeff, 10.0f / desat_param); ++ rgb = mix(rgb, make_float3(luma, luma, luma), make_float3(coeff, coeff, coeff)); ++ } ++ ++ sig.x = map(sig.x, peak, dst_peak); ++ sig.y = map(sig.y, peak, dst_peak); ++ sig.z = map(sig.z, peak, dst_peak); ++ sig.x = min(sig.x, 1.0f); ++ sig.y = min(sig.y, 1.0f); ++ sig.z = min(sig.z, 1.0f); ++ rgb = rgb * (sig / sig_old); ++ ++ return rgb; ++} ++ +// Map from source space YUV to destination space RGB +static __inline__ __device__ +float3 map_to_dst_space_from_yuv(float3 yuv) { @@ -1325,11 +1343,17 @@ Index: jellyfin-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 _TONEMAP \ -+ c0 = map_one_pixel_rgb(c0, src, dst); \ -+ c1 = map_one_pixel_rgb(c1, src, dst); \ -+ c2 = map_one_pixel_rgb(c2, src, dst); \ -+ c3 = map_one_pixel_rgb(c3, src, dst); ++#define _TONEMAP_MAX \ ++ c0 = map_one_pixel_rgb_mode_max(c0, src, dst); \ ++ c1 = map_one_pixel_rgb_mode_max(c1, src, dst); \ ++ c2 = map_one_pixel_rgb_mode_max(c2, src, dst); \ ++ c3 = map_one_pixel_rgb_mode_max(c3, src, dst); ++ ++#define _TONEMAP_RGB \ ++ c0 = map_one_pixel_rgb_mode_rgb(c0, src, dst); \ ++ c1 = map_one_pixel_rgb_mode_rgb(c1, src, dst); \ ++ c2 = map_one_pixel_rgb_mode_rgb(c2, src, dst); \ ++ c3 = map_one_pixel_rgb_mode_rgb(c3, src, dst); + +#define _RGB2YUV \ + yuv0 = lrgb2yuv(c0); \ @@ -1347,99 +1371,44 @@ Index: jellyfin-ffmpeg/libavfilter/cuda/tonemap.cu +#define _WRITER \ + write_2x2_flt(dst, x, y, yuv0, yuv1, yuv2, yuv3); + -+__global__ void tonemap(FFCUDAFrame src, FFCUDAFrame dst, -+ cudaTextureObject_t ditherTex, float *unused) -+{ -+ _READER -+ _YUV2RGB -+ _TONEMAP -+ _RGB2YUV -+ _WRITER -+} -+ -+__global__ void tonemap_d(FFCUDAFrame src, FFCUDAFrame dst, -+ cudaTextureObject_t ditherTex, float *unused) -+{ -+ _READER -+ _YUV2RGB -+ _TONEMAP -+ _RGB2YUV -+ _DITHER -+ _WRITER -+} -+ -+__global__ void tonemap_dovi(FFCUDAFrame src, FFCUDAFrame dst, -+ cudaTextureObject_t ditherTex, float *doviBuf) -+{ -+ _READER -+ _RESHAPE -+ _YCC2RGB -+ _TONEMAP -+ _RGB2YUV -+ _WRITER -+} -+ -+__global__ void tonemap_dovi_d(FFCUDAFrame src, FFCUDAFrame dst, -+ cudaTextureObject_t ditherTex, float *doviBuf) -+{ -+ _READER -+ _RESHAPE -+ _YCC2RGB -+ _TONEMAP -+ _RGB2YUV -+ _DITHER -+ _WRITER ++#define TONEMAP_VARIANT(NAME, READER, RESHAPE, YUV2RGB, TONEMAP, RGB2YUV, DITHER, WRITER) \ ++__global__ void tonemap ## NAME( \ ++ FFCUDAFrame src, FFCUDAFrame dst, \ ++ cudaTextureObject_t ditherTex, float *doviBuf) \ ++{ \ ++ READER \ ++ RESHAPE \ ++ YUV2RGB \ ++ TONEMAP \ ++ RGB2YUV \ ++ DITHER \ ++ WRITER \ +} + -+__global__ void tonemap_dovi_f(FFCUDAFrame src, FFCUDAFrame dst, -+ cudaTextureObject_t ditherTex, float *doviBuf) -+{ -+ _READER -+ _RESHAPE -+ _YCC2RGB_F -+ _TONEMAP -+ _RGB2YUV -+ _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) + -+__global__ void tonemap_dovi_d_f(FFCUDAFrame src, FFCUDAFrame dst, -+ cudaTextureObject_t ditherTex, float *doviBuf) -+{ -+ _READER -+ _RESHAPE -+ _YCC2RGB_F -+ _TONEMAP -+ _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) + -+__global__ void tonemap_dovi_pq(FFCUDAFrame src, FFCUDAFrame dst, -+ cudaTextureObject_t ditherTex, float *doviBuf) -+{ -+ _READER -+ _RESHAPE -+ _YCC2RGB -+ _RGB2YUV -+ _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) + -+__global__ void tonemap_dovi_pq_f(FFCUDAFrame src, FFCUDAFrame dst, -+ cudaTextureObject_t ditherTex, float *doviBuf) -+{ -+ _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: jellyfin-ffmpeg/libavfilter/cuda/tonemap.h =================================================================== --- /dev/null +++ jellyfin-ffmpeg/libavfilter/cuda/tonemap.h -@@ -0,0 +1,34 @@ +@@ -0,0 +1,40 @@ +/* + * This file is part of FFmpeg. + * @@ -1470,7 +1439,13 @@ Index: jellyfin-ffmpeg/libavfilter/cuda/tonemap.h + TONEMAP_HABLE, + TONEMAP_MOBIUS, + TONEMAP_BT2390, -+ TONEMAP_MAX, ++ TONEMAP_COUNT, ++}; ++ ++enum TonemapMode { ++ TONEMAP_MODE_MAX, ++ TONEMAP_MODE_RGB, ++ TONEMAP_MODE_COUNT, +}; + +#endif /* AVFILTER_CUDA_TONEMAP_H */ @@ -1569,7 +1544,7 @@ Index: jellyfin-ffmpeg/libavfilter/vf_tonemap_cuda.c =================================================================== --- /dev/null +++ jellyfin-ffmpeg/libavfilter/vf_tonemap_cuda.c -@@ -0,0 +1,1088 @@ +@@ -0,0 +1,1096 @@ +/* + * This file is part of FFmpeg. + * @@ -1676,6 +1651,7 @@ Index: jellyfin-ffmpeg/libavfilter/vf_tonemap_cuda.c + float *dovi_pbuf; + + enum TonemapAlgorithm tonemap; ++ enum TonemapMode tonemap_mode; + int apply_dovi; + int tradeoff; + int init_with_dovi; @@ -2081,6 +2057,7 @@ Index: jellyfin-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 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, @@ -2312,13 +2289,16 @@ Index: jellyfin-ffmpeg/libavfilter/vf_tonemap_cuda.c + goto fail2; + + ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_tm, s->cu_module, -+ d ? "tonemap_d" : "tonemap")); ++ rgb ? (d ? "tonemap_rgb_d" : "tonemap_rgb") ++ : (d ? "tonemap_d" : "tonemap"))); + 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"))); ++ s->tradeoff == 1 ? (rgb ? (d ? "tonemap_dovi_rgb_d_f" : "tonemap_dovi_rgb_f") ++ : (d ? "tonemap_dovi_d_f" : "tonemap_dovi_f")) ++ : (rgb ? (d ? "tonemap_dovi_rgb_d" : "tonemap_dovi_rgb") ++ : (d ? "tonemap_dovi_d" : "tonemap_dovi")))); + if (ret < 0) + goto fail2; + @@ -2577,7 +2557,7 @@ Index: jellyfin-ffmpeg/libavfilter/vf_tonemap_cuda.c +#define OFFSET(x) offsetof(TonemapCUDAContext, x) +#define FLAGS (AV_OPT_FLAG_FILTERING_PARAM|AV_OPT_FLAG_VIDEO_PARAM) +static const AVOption options[] = { -+ { "tonemap", "Tonemap algorithm selection", OFFSET(tonemap), AV_OPT_TYPE_INT, {.i64 = TONEMAP_NONE}, TONEMAP_NONE, TONEMAP_MAX - 1, FLAGS, "tonemap" }, ++ { "tonemap", "Tonemap algorithm selection", OFFSET(tonemap), AV_OPT_TYPE_INT, {.i64 = TONEMAP_NONE}, TONEMAP_NONE, TONEMAP_COUNT - 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" }, @@ -2586,6 +2566,9 @@ Index: jellyfin-ffmpeg/libavfilter/vf_tonemap_cuda.c + { "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" }, + { "bt2390", 0, 0, AV_OPT_TYPE_CONST, {.i64 = TONEMAP_BT2390}, 0, 0, FLAGS, "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, "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" }, + { "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 53b373d39a..658fa7407b 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 @@ -2,7 +2,7 @@ Index: jellyfin-ffmpeg/libavfilter/opencl.c =================================================================== --- jellyfin-ffmpeg.orig/libavfilter/opencl.c +++ jellyfin-ffmpeg/libavfilter/opencl.c -@@ -169,7 +169,7 @@ int ff_opencl_filter_load_program(AVFilt +@@ -170,7 +170,7 @@ int ff_opencl_filter_load_program(AVFilt } cle = clBuildProgram(ctx->program, 1, &ctx->hwctx->device_id, @@ -11,7 +11,7 @@ Index: jellyfin-ffmpeg/libavfilter/opencl.c if (cle != CL_SUCCESS) { av_log(avctx, AV_LOG_ERROR, "Failed to build program: %d.\n", cle); -@@ -330,7 +330,7 @@ void ff_opencl_print_const_matrix_3x3(AV +@@ -331,7 +331,7 @@ void ff_opencl_print_const_matrix_3x3(AV av_bprintf(buf, "__constant float %s[9] = {\n", name_str); for (i = 0; i < 3; i++) { for (j = 0; j < 3; j++) @@ -384,7 +384,7 @@ Index: jellyfin-ffmpeg/libavfilter/opencl/tonemap.cl float j = tone_param; float a, b; -@@ -71,202 +77,312 @@ float mobius(float s, float peak) { +@@ -71,202 +77,335 @@ float mobius(float s, float peak) { return s; a = -j * j * (peak - 1.0f) / (j * j - 2.0f * j + peak); @@ -442,13 +442,6 @@ Index: jellyfin-ffmpeg/libavfilter/opencl/tonemap.cl - 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); -- } +float bt2390(float s, float peak_inv_pq, float target_peak_inv_pq) { + float peak_pq = peak_inv_pq; + float scale = peak_pq > 0.0f ? (1.0f / peak_pq) : 1.0f; @@ -464,6 +457,37 @@ Index: jellyfin-ffmpeg/libavfilter/opencl/tonemap.cl + (tb3 - 2.0f * tb2 + tb) * (1.0f - ks) + + (-2.0f * tb3 + 3.0f * tb2) * max_lum; + float sig = mix(pb, s_pq, s_pq < ks); ++ ++ return eotf_st2084(sig * peak_pq); ++} ++ ++void map_four_pixels_rgb(float4 *r4, float4 *g4, float4 *b4, float peak) { ++#ifdef TONE_MODE_RGB ++ float4 sig_r = fmax(*r4, FLOAT_EPS), sig_ro = sig_r; ++ float4 sig_g = fmax(*g4, FLOAT_EPS), sig_go = sig_g; ++ float4 sig_b = fmax(*b4, FLOAT_EPS), sig_bo = sig_b; ++#else ++ float4 sig = fmax(fmax(*r4, fmax(*g4, *b4)), FLOAT_EPS); ++ float4 sig_o = sig; ++#endif + +- 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); ++ // Desaturate the color using a coefficient dependent on the signal level ++ if (desat_param > 0.0f) { ++#ifdef TONE_MODE_RGB ++ float4 sig = fmax(fmax(*r4, fmax(*g4, *b4)), FLOAT_EPS); ++#endif ++ float4 luma = get_luma_dst4(*r4, *g4, *b4); ++ float4 coeff = fmax(sig - 0.18f, FLOAT_EPS) / fmax(sig, FLOAT_EPS); ++ coeff = native_powr(coeff, 10.0f / desat_param); ++ *r4 = mix(*r4, luma, coeff); ++ *g4 = mix(*g4, luma, coeff); ++ *b4 = mix(*b4, luma, coeff); + } - if (lidx == 0 && lidy == 0 && atomic_add(counter_wg_p, 1) == num_wg - 1) { - *counter_wg_p = 0; @@ -496,77 +520,51 @@ Index: jellyfin-ffmpeg/libavfilter/opencl/tonemap.cl - (uint)DETECTION_FRAMES); - } - return r; -+ return eotf_st2084(sig * peak_pq); - } - --float3 map_one_pixel_rgb(float3 rgb, float peak, float average) { -- float sig = max(max(rgb.x, max(rgb.y, rgb.z)), 1e-6f); -- -+/* -+float rescale_peak(float sig) { - // 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) { -+ if (target_peak > 1.0f) - sig *= 1.0f / target_peak; -- peak *= 1.0f / target_peak; -- } - -- float sig_old = sig; -+ return sig; ++#define MAP_FOUR_PIXELS(sig, peak, target_peak) \ ++{ \ ++ sig.x = TONE_FUNC(sig.x, peak, target_peak); \ ++ sig.y = TONE_FUNC(sig.y, peak, target_peak); \ ++ sig.z = TONE_FUNC(sig.z, peak, target_peak); \ ++ sig.w = TONE_FUNC(sig.w, peak, target_peak); \ +} -+*/ - -- // Scale the signal to compensate for differences in the average brightness -- float slope = min(1.0f, sdr_avg / average); -- sig *= slope; -- peak *= slope; -+void map_four_pixels_rgb(float4 *r4, float4 *g4, float4 *b4, float peak) { -+ float4 sig = fmax(fmax(*r4, fmax(*g4, *b4)), FLOAT_EPS); -+ float4 sig_old = sig; - - // 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); -+ float4 luma = get_luma_dst4(*r4, *g4, *b4); -+ float4 coeff = fmax(sig - 0.18f, FLOAT_EPS) / fmax(sig, FLOAT_EPS); - coeff = native_powr(coeff, 10.0f / desat_param); -- rgb = mix(rgb, (float3)luma, (float3)coeff); -- sig = mix(sig, luma * slope, coeff); -+ *r4 = mix(*r4, luma, coeff); -+ *g4 = mix(*g4, luma, coeff); -+ *b4 = mix(*b4, luma, 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) { ++ +#ifdef TONE_FUNC_BT2390 + float src_peak_delin_pq = inverse_eotf_st2084(peak); + float dst_peak_delin_pq = inverse_eotf_st2084(1.0f); -+ sig.x = TONE_FUNC(sig.x, src_peak_delin_pq, dst_peak_delin_pq); -+ sig.y = TONE_FUNC(sig.y, src_peak_delin_pq, dst_peak_delin_pq); -+ sig.z = TONE_FUNC(sig.z, src_peak_delin_pq, dst_peak_delin_pq); -+ sig.w = TONE_FUNC(sig.w, src_peak_delin_pq, dst_peak_delin_pq); ++ #ifdef TONE_MODE_RGB ++ 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) ++ #else ++ MAP_FOUR_PIXELS(sig, src_peak_delin_pq, dst_peak_delin_pq) ++ #endif +#else -+ sig.x = TONE_FUNC(sig.x, peak, 1.0f); -+ sig.y = TONE_FUNC(sig.y, peak, 1.0f); -+ sig.z = TONE_FUNC(sig.z, peak, 1.0f); -+ sig.w = TONE_FUNC(sig.w, peak, 1.0f); ++ #ifdef TONE_MODE_RGB ++ MAP_FOUR_PIXELS(sig_r, peak, 1.0f) ++ MAP_FOUR_PIXELS(sig_g, peak, 1.0f) ++ MAP_FOUR_PIXELS(sig_b, peak, 1.0f) ++ #else ++ MAP_FOUR_PIXELS(sig, peak, 1.0f) ++ #endif +#endif -+ sig = fmin(sig, 1.0f); + -+ float4 factor = sig / sig_old; ++#ifdef TONE_MODE_RGB ++ sig_r = fmin(sig_r, 1.0f); ++ sig_g = fmin(sig_g, 1.0f); ++ sig_b = fmin(sig_b, 1.0f); ++ float4 factor_r = sig_r / sig_ro; ++ float4 factor_g = sig_g / sig_go; ++ float4 factor_b = sig_b / sig_bo; ++ *r4 *= factor_r; ++ *g4 *= factor_g; ++ *b4 *= factor_b; ++#else ++ sig = fmin(sig, 1.0f); ++ float4 factor = sig / sig_o; + *r4 *= factor; + *g4 *= factor; + *b4 *= factor; ++#endif +} + +// Map from source space YUV to destination space RGB @@ -576,18 +574,25 @@ Index: jellyfin-ffmpeg/libavfilter/opencl/tonemap.cl + c = lms2rgb(c.x, c.y, c.z); + c = rgb2lrgb(c); +#else - float3 c = yuv2lrgb(yuv); -- c = ootf(c, peak); - c = lrgb2lrgb(c); ++ float3 c = yuv2lrgb(yuv); ++ c = lrgb2lrgb(c); +#endif - return c; ++ return c; } +-float3 map_one_pixel_rgb(float3 rgb, float peak, float average) { +- float sig = max(max(rgb.x, max(rgb.y, rgb.z)), 1e-6f); +#ifdef DOVI_RESHAPE +float reshape_poly(float s, float4 coeffs) { + return (coeffs.z * s + coeffs.y) * s + coeffs.x; +} -+ + +- // 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 reshape_mmr(float3 sig, + float4 coeffs, + __global float4 *dovi_mmr, @@ -616,8 +621,9 @@ Index: jellyfin-ffmpeg/libavfilter/opencl/tonemap.cl + s += dot(dovi_mmr[mmr_idx + 4].xyz, sig2 * sig); + s += dot(dovi_mmr[mmr_idx + 5], sigX2 * sigX); + } -+ } -+ + } + +- float sig_old = sig; + return s; +} + @@ -666,20 +672,43 @@ Index: jellyfin-ffmpeg/libavfilter/opencl/tonemap.cl + (float4)(s >= dovi_pivots[5])), + (float4)(s >= dovi_pivots[3])); + } -+ + +- // Scale the signal to compensate for differences in the average brightness +- float slope = min(1.0f, sdr_avg / average); +- sig *= slope; +- peak *= slope; + int has_mmr_poly = dovi_has_mmr && dovi_has_poly; -+ + +- // 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); +- } + if ((has_mmr_poly && coeffs.w == 0.0f) || (!has_mmr_poly && dovi_has_poly)) + s = reshape_poly(s, coeffs); + else + s = reshape_mmr(sig, coeffs, dovi_mmr, + dovi_mmr_single, dovi_min_order, dovi_max_order); -+ + +- sig = TONE_FUNC(sig, peak); + sig_arr[i] = clamp(s, dovi_lo, dovi_hi); + } -+ + +- 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; + return (float3)(sig_arr[0], sig_arr[1], sig_arr[2]); -+} + } +#endif + +__constant sampler_t sampler = (CLK_NORMALIZED_COORDS_FALSE | @@ -693,7 +722,7 @@ Index: jellyfin-ffmpeg/libavfilter/opencl/tonemap.cl +__constant sampler_t d_sampler = (CLK_NORMALIZED_COORDS_TRUE | + CLK_ADDRESS_REPEAT | + CLK_FILTER_NEAREST); -+ + __kernel void tonemap(__write_only image2d_t dst1, __read_only image2d_t src1, __write_only image2d_t dst2, @@ -906,15 +935,23 @@ Index: jellyfin-ffmpeg/libavfilter/vf_tonemap_opencl.c enum TonemapAlgorithm { TONEMAP_NONE, -@@ -45,6 +56,7 @@ enum TonemapAlgorithm { +@@ -45,7 +56,14 @@ enum TonemapAlgorithm { TONEMAP_REINHARD, TONEMAP_HABLE, TONEMAP_MOBIUS, +- TONEMAP_MAX, + TONEMAP_BT2390, - TONEMAP_MAX, ++ TONEMAP_COUNT, ++}; ++ ++enum TonemapMode { ++ TONEMAP_MODE_MAX, ++ TONEMAP_MODE_RGB, ++ TONEMAP_MODE_COUNT, }; -@@ -56,23 +68,43 @@ typedef struct TonemapOpenCLContext { + typedef struct TonemapOpenCLContext { +@@ -56,23 +74,44 @@ typedef struct TonemapOpenCLContext { enum AVColorPrimaries primaries, primaries_in, primaries_out; enum AVColorRange range, range_in, range_out; enum AVChromaLocation chroma_loc; @@ -936,6 +973,7 @@ Index: jellyfin-ffmpeg/libavfilter/vf_tonemap_opencl.c + cl_mem dovi_buf; enum TonemapAlgorithm tonemap; ++ enum TonemapMode tonemap_mode; enum AVPixelFormat format; + int apply_dovi; + double ref_white; @@ -962,7 +1000,16 @@ Index: jellyfin-ffmpeg/libavfilter/vf_tonemap_opencl.c }; static const char *const delinearize_funcs[AVCOL_TRC_NB] = { -@@ -88,8 +120,54 @@ static const char *const tonemap_func[TO +@@ -80,7 +119,7 @@ static const char *const delinearize_fun + [AVCOL_TRC_BT2020_10] = "inverse_eotf_bt1886", + }; + +-static const char *const tonemap_func[TONEMAP_MAX] = { ++static const char *const tonemap_func[TONEMAP_COUNT] = { + [TONEMAP_NONE] = "direct", + [TONEMAP_LINEAR] = "linear", + [TONEMAP_GAMMA] = "gamma", +@@ -88,8 +127,54 @@ static const char *const tonemap_func[TO [TONEMAP_REINHARD] = "reinhard", [TONEMAP_HABLE] = "hable", [TONEMAP_MOBIUS] = "mobius", @@ -1017,7 +1064,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 +186,149 @@ static int get_rgb2rgb_matrix(enum AVCol +@@ -108,23 +193,149 @@ static int get_rgb2rgb_matrix(enum AVCol return 0; } @@ -1176,7 +1223,7 @@ Index: jellyfin-ffmpeg/libavfilter/vf_tonemap_opencl.c switch(ctx->tonemap) { case TONEMAP_GAMMA: -@@ -144,48 +348,143 @@ static int tonemap_opencl_init(AVFilterC +@@ -144,48 +355,148 @@ static int tonemap_opencl_init(AVFilterC if (isnan(ctx->param)) ctx->param = 1.0f; @@ -1314,6 +1361,11 @@ Index: jellyfin-ffmpeg/libavfilter/vf_tonemap_opencl.c + if (ctx->tonemap == TONEMAP_BT2390) + av_bprintf(&header, "#define TONE_FUNC_BT2390\n"); + ++ if (ctx->tonemap_mode == TONEMAP_MODE_RGB) ++ av_bprintf(&header, "#define TONE_MODE_RGB\n"); ++ else ++ av_bprintf(&header, "#define TONE_MODE_MAX\n"); ++ + if (ctx->in_planes > 2) + av_bprintf(&header, "#define NON_SEMI_PLANAR_IN\n"); + @@ -1335,7 +1387,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 +498,41 @@ static int tonemap_opencl_init(AVFilterC +@@ -199,19 +510,41 @@ static int tonemap_opencl_init(AVFilterC else ff_opencl_print_const_matrix_3x3(&header, "rgb2rgb", rgb2rgb); @@ -1384,7 +1436,7 @@ Index: jellyfin-ffmpeg/libavfilter/vf_tonemap_opencl.c ctx->colorspace_out, av_color_space_name(ctx->colorspace_out)); goto fail; } -@@ -219,24 +540,23 @@ static int tonemap_opencl_init(AVFilterC +@@ -219,24 +552,23 @@ static int tonemap_opencl_init(AVFilterC ff_fill_rgb2yuv_table(luma_dst, rgb2yuv); ff_opencl_print_const_matrix_3x3(&header, "yuv_matrix", rgb2yuv); @@ -1424,7 +1476,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 +574,170 @@ static int tonemap_opencl_init(AVFilterC +@@ -254,46 +586,170 @@ static int tonemap_opencl_init(AVFilterC CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create OpenCL " "command queue %d.\n", cle); @@ -1614,7 +1666,7 @@ Index: jellyfin-ffmpeg/libavfilter/vf_tonemap_opencl.c ret = ff_opencl_filter_config_output(outlink); if (ret < 0) return ret; -@@ -308,13 +752,42 @@ static int launch_kernel(AVFilterContext +@@ -308,13 +764,42 @@ static int launch_kernel(AVFilterContext size_t global_work[2]; size_t local_work[2]; cl_int cle; @@ -1659,7 +1711,7 @@ Index: jellyfin-ffmpeg/libavfilter/vf_tonemap_opencl.c local_work[0] = 16; local_work[1] = 16; -@@ -338,13 +811,10 @@ static int tonemap_opencl_filter_frame(A +@@ -338,13 +823,10 @@ static int tonemap_opencl_filter_frame(A AVFilterContext *avctx = inlink->dst; AVFilterLink *outlink = avctx->outputs[0]; TonemapOpenCLContext *ctx = avctx->priv; @@ -1674,7 +1726,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 +833,6 @@ static int tonemap_opencl_filter_frame(A +@@ -363,9 +845,6 @@ static int tonemap_opencl_filter_frame(A if (err < 0) goto fail; @@ -1684,7 +1736,7 @@ Index: jellyfin-ffmpeg/libavfilter/vf_tonemap_opencl.c if (ctx->trc != -1) output->color_trc = ctx->trc; if (ctx->primaries != -1) -@@ -385,72 +852,92 @@ static int tonemap_opencl_filter_frame(A +@@ -385,72 +864,92 @@ static int tonemap_opencl_filter_frame(A ctx->range_out = output->color_range; ctx->chroma_loc = output->chroma_location; @@ -1817,14 +1869,13 @@ Index: jellyfin-ffmpeg/libavfilter/vf_tonemap_opencl.c av_frame_free(&input); av_frame_free(&output); return err; -@@ -458,24 +945,9 @@ fail: +@@ -458,24 +957,9 @@ fail: static av_cold void tonemap_opencl_uninit(AVFilterContext *avctx) { - TonemapOpenCLContext *ctx = avctx->priv; - cl_int cle; -+ tonemap_opencl_uninit_common(avctx); - +- - if (ctx->util_mem) - clReleaseMemObject(ctx->util_mem); - if (ctx->kernel) { @@ -1833,7 +1884,8 @@ Index: jellyfin-ffmpeg/libavfilter/vf_tonemap_opencl.c - av_log(avctx, AV_LOG_ERROR, "Failed to release " - "kernel: %d.\n", cle); - } -- ++ tonemap_opencl_uninit_common(avctx); + - if (ctx->command_queue) { - cle = clReleaseCommandQueue(ctx->command_queue); - if (cle != CL_SUCCESS) @@ -1844,7 +1896,7 @@ Index: jellyfin-ffmpeg/libavfilter/vf_tonemap_opencl.c ff_opencl_filter_uninit(avctx); } -@@ -483,37 +955,44 @@ static av_cold void tonemap_opencl_unini +@@ -483,37 +967,47 @@ 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[] = { @@ -1879,7 +1931,7 @@ Index: jellyfin-ffmpeg/libavfilter/vf_tonemap_opencl.c - { "param", "tonemap parameter", OFFSET(param), AV_OPT_TYPE_DOUBLE, {.dbl = NAN}, DBL_MIN, DBL_MAX, FLAGS }, - { "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 }, -+ { "tonemap", "Tonemap algorithm selection", OFFSET(tonemap), AV_OPT_TYPE_INT, { .i64 = TONEMAP_NONE }, TONEMAP_NONE, TONEMAP_MAX - 1, FLAGS, "tonemap" }, ++ { "tonemap", "Tonemap algorithm selection", OFFSET(tonemap), AV_OPT_TYPE_INT, { .i64 = TONEMAP_NONE }, TONEMAP_NONE, TONEMAP_COUNT - 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" }, @@ -1888,6 +1940,9 @@ Index: jellyfin-ffmpeg/libavfilter/vf_tonemap_opencl.c + { "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" }, + { "bt2390", 0, 0, AV_OPT_TYPE_CONST, { .i64 = TONEMAP_BT2390 }, 0, 0, FLAGS, "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, "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" }, + { "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" }, From 1fdfde52d399fe613008cd33c190f6b21dab19af Mon Sep 17 00:00:00 2001 From: nyanmisaka Date: Tue, 4 Apr 2023 15:19:42 +0800 Subject: [PATCH 2/7] Update dependencies Signed-off-by: nyanmisaka --- builder/scripts.d/10-mingw.sh | 2 +- builder/scripts.d/20-libiconv.sh | 3 ++- builder/scripts.d/20-libxml2.sh | 2 +- builder/scripts.d/25-freetype.sh | 2 +- builder/scripts.d/35-fontconfig.sh | 2 +- builder/scripts.d/45-harfbuzz.sh | 2 +- builder/scripts.d/45-opencl.sh | 2 +- builder/scripts.d/45-x11/10-xcbproto.sh | 2 +- builder/scripts.d/45-x11/10-xproto.sh | 2 +- builder/scripts.d/45-x11/20-libxau.sh | 2 +- builder/scripts.d/45-x11/30-libxcb.sh | 2 +- builder/scripts.d/45-x11/40-libx11.sh | 2 +- builder/scripts.d/45-x11/60-libxcursor.sh | 2 +- builder/scripts.d/50-dav1d.sh | 2 +- builder/scripts.d/50-libass.sh | 2 +- builder/scripts.d/50-libvpx.sh | 2 +- builder/scripts.d/50-libwebp.sh | 2 +- builder/scripts.d/50-svtav1.sh | 2 +- builder/scripts.d/50-vaapi/30-libpciaccess.sh | 2 +- builder/scripts.d/50-vaapi/40-libdrm.sh | 2 +- builder/scripts.d/50-vaapi/50-libva.sh | 2 +- builder/scripts.d/50-vulkan/45-vulkan.sh | 2 +- builder/scripts.d/50-vulkan/50-shaderc.sh | 2 +- builder/scripts.d/50-vulkan/55-spirv-cross.sh | 2 +- builder/scripts.d/50-vulkan/60-libplacebo.sh | 2 +- builder/scripts.d/50-x265.sh | 2 +- docker-build-win64.sh | 8 ++++++-- docker-build.sh | 14 +++++++------- 28 files changed, 40 insertions(+), 35 deletions(-) diff --git a/builder/scripts.d/10-mingw.sh b/builder/scripts.d/10-mingw.sh index 083ed7ba38..71293ef3de 100755 --- a/builder/scripts.d/10-mingw.sh +++ b/builder/scripts.d/10-mingw.sh @@ -1,7 +1,7 @@ #!/bin/bash SCRIPT_REPO="https://github.com/mirror/mingw-w64.git" -SCRIPT_COMMIT="0f2264e7b8fedbe225921367e82aeb97ddfed46b" +SCRIPT_COMMIT="eff726c461e09f35eeaed125a3570fa5f807f02b" ffbuild_enabled() { [[ $TARGET == win* ]] || return -1 diff --git a/builder/scripts.d/20-libiconv.sh b/builder/scripts.d/20-libiconv.sh index 220634f472..6b933c1d7d 100755 --- a/builder/scripts.d/20-libiconv.sh +++ b/builder/scripts.d/20-libiconv.sh @@ -1,7 +1,8 @@ #!/bin/bash SCRIPT_REPO="https://git.savannah.gnu.org/git/libiconv.git" -SCRIPT_COMMIT="c593e206b2d4bc689950c742a0fb00b8013756a0" +SCRIPT_COMMIT="v1.17" +SCRIPT_TAGFILTER="v?.*" ffbuild_enabled() { return 0 diff --git a/builder/scripts.d/20-libxml2.sh b/builder/scripts.d/20-libxml2.sh index a3dcf50542..2503a54c60 100755 --- a/builder/scripts.d/20-libxml2.sh +++ b/builder/scripts.d/20-libxml2.sh @@ -1,7 +1,7 @@ #!/bin/bash SCRIPT_REPO="https://github.com/GNOME/libxml2.git" -SCRIPT_COMMIT="b1319c902f6e44d08f8cb33f1fc28847f2bc8aeb" +SCRIPT_COMMIT="d7d0bc6581e332f49c9ff628f548eced03c65189" ffbuild_enabled() { return 0 diff --git a/builder/scripts.d/25-freetype.sh b/builder/scripts.d/25-freetype.sh index 36ce2cd7e2..7041d6eee1 100755 --- a/builder/scripts.d/25-freetype.sh +++ b/builder/scripts.d/25-freetype.sh @@ -1,7 +1,7 @@ #!/bin/bash SCRIPT_REPO="https://gitlab.freedesktop.org/freetype/freetype.git" -SCRIPT_COMMIT="4f0a55d15ed1c9af267ed8223cc2f04307a3d656" +SCRIPT_COMMIT="6d7b8b22c4f4ca6cde1998997ae4ea2aa6e57d26" ffbuild_enabled() { return 0 diff --git a/builder/scripts.d/35-fontconfig.sh b/builder/scripts.d/35-fontconfig.sh index 108f97a17c..a1e8be2e56 100755 --- a/builder/scripts.d/35-fontconfig.sh +++ b/builder/scripts.d/35-fontconfig.sh @@ -1,7 +1,7 @@ #!/bin/bash SCRIPT_REPO="https://gitlab.freedesktop.org/fontconfig/fontconfig.git" -SCRIPT_COMMIT="04546f18768e1ce57c24743035118b3eabbd4181" +SCRIPT_COMMIT="c2666a6d9a6ed18b1bfcef8176e25f62993e24db" ffbuild_enabled() { return 0 diff --git a/builder/scripts.d/45-harfbuzz.sh b/builder/scripts.d/45-harfbuzz.sh index 4f322db00d..0d23c85105 100755 --- a/builder/scripts.d/45-harfbuzz.sh +++ b/builder/scripts.d/45-harfbuzz.sh @@ -1,7 +1,7 @@ #!/bin/bash SCRIPT_REPO="https://github.com/harfbuzz/harfbuzz.git" -SCRIPT_COMMIT="79233a149209e3da199bb4e2f74271668502c574" +SCRIPT_COMMIT="04a47932a3844f7e73e3af8b05fb98c8b54fb779" ffbuild_enabled() { return 0 diff --git a/builder/scripts.d/45-opencl.sh b/builder/scripts.d/45-opencl.sh index 511cf9e620..35582ad0df 100755 --- a/builder/scripts.d/45-opencl.sh +++ b/builder/scripts.d/45-opencl.sh @@ -1,7 +1,7 @@ #!/bin/bash SCRIPT_REPO="https://github.com/KhronosGroup/OpenCL-Headers.git" -SCRIPT_COMMIT="e3e85862d6905eba9f4b5ed02c52effe673721d3" +SCRIPT_COMMIT="8c4f0111ccf2350a51655d8bbe862c7405f2c0db" SCRIPT_REPO2="https://github.com/KhronosGroup/OpenCL-ICD-Loader.git" SCRIPT_COMMIT2="ece91448a958099b9c277f050fca9df96a2ea718" diff --git a/builder/scripts.d/45-x11/10-xcbproto.sh b/builder/scripts.d/45-x11/10-xcbproto.sh index 4fa10ec6c8..632f64268d 100755 --- a/builder/scripts.d/45-x11/10-xcbproto.sh +++ b/builder/scripts.d/45-x11/10-xcbproto.sh @@ -1,7 +1,7 @@ #!/bin/bash SCRIPT_REPO="https://gitlab.freedesktop.org/xorg/proto/xcbproto.git" -SCRIPT_COMMIT="15d140d7867e8e654ce917b8d6d1dbd45b3de3b8" +SCRIPT_COMMIT="74c03b4edfe61ff5f30cc0b61f7d2deff025c6a3" ffbuild_enabled() { [[ $TARGET != linux* ]] && return -1 diff --git a/builder/scripts.d/45-x11/10-xproto.sh b/builder/scripts.d/45-x11/10-xproto.sh index 7b2f0fb6e7..b7360dd167 100755 --- a/builder/scripts.d/45-x11/10-xproto.sh +++ b/builder/scripts.d/45-x11/10-xproto.sh @@ -1,7 +1,7 @@ #!/bin/bash SCRIPT_REPO="https://gitlab.freedesktop.org/xorg/proto/xorgproto.git" -SCRIPT_COMMIT="cf35a91fe57d4721a173d2bc428dd07dcc4674f9" +SCRIPT_COMMIT="fca42f2e5a5da961b231dbbc10f87bb95588d5b1" ffbuild_enabled() { [[ $TARGET != linux* ]] && return -1 diff --git a/builder/scripts.d/45-x11/20-libxau.sh b/builder/scripts.d/45-x11/20-libxau.sh index 4ab0d1e850..4c2e8beebb 100755 --- a/builder/scripts.d/45-x11/20-libxau.sh +++ b/builder/scripts.d/45-x11/20-libxau.sh @@ -1,7 +1,7 @@ #!/bin/bash SCRIPT_REPO="https://gitlab.freedesktop.org/xorg/lib/libxau.git" -SCRIPT_COMMIT="c52f54e9533046a52edf84bcc02abedc2dbcb1a7" +SCRIPT_COMMIT="df1bf4fe528a5a9eef420f78efb225e4696ac467" ffbuild_enabled() { [[ $TARGET != linux* ]] && return -1 diff --git a/builder/scripts.d/45-x11/30-libxcb.sh b/builder/scripts.d/45-x11/30-libxcb.sh index fd2f43dd33..726f98f469 100755 --- a/builder/scripts.d/45-x11/30-libxcb.sh +++ b/builder/scripts.d/45-x11/30-libxcb.sh @@ -1,7 +1,7 @@ #!/bin/bash SCRIPT_REPO="https://gitlab.freedesktop.org/xorg/lib/libxcb.git" -SCRIPT_COMMIT="18e109d755c5ce18157fdabb6de8ee6845b348ff" +SCRIPT_COMMIT="8935793f1f3751a6aa9d78955c7d6236177986de" ffbuild_enabled() { [[ $TARGET != linux* ]] && return -1 diff --git a/builder/scripts.d/45-x11/40-libx11.sh b/builder/scripts.d/45-x11/40-libx11.sh index dba70794c1..8152f4f81a 100755 --- a/builder/scripts.d/45-x11/40-libx11.sh +++ b/builder/scripts.d/45-x11/40-libx11.sh @@ -1,7 +1,7 @@ #!/bin/bash SCRIPT_REPO="https://gitlab.freedesktop.org/xorg/lib/libx11.git" -SCRIPT_COMMIT="ca99e338a9b8aad300933b1336f9e3c091392213" +SCRIPT_COMMIT="eb166af8f9f4184108fe8e0611a09af03d12059e" ffbuild_enabled() { [[ $TARGET != linux* ]] && return -1 diff --git a/builder/scripts.d/45-x11/60-libxcursor.sh b/builder/scripts.d/45-x11/60-libxcursor.sh index 8e9c9673b5..18e4057008 100755 --- a/builder/scripts.d/45-x11/60-libxcursor.sh +++ b/builder/scripts.d/45-x11/60-libxcursor.sh @@ -1,7 +1,7 @@ #!/bin/bash SCRIPT_REPO="https://gitlab.freedesktop.org/xorg/lib/libxcursor.git" -SCRIPT_COMMIT="81dc4a481b64499ab7c355ee43c91e4fe0767545" +SCRIPT_COMMIT="a353f02a7ac4504ad5edb5407278c7a03c507aef" ffbuild_enabled() { [[ $TARGET != linux* ]] && return -1 diff --git a/builder/scripts.d/50-dav1d.sh b/builder/scripts.d/50-dav1d.sh index cef3bcf84b..3a615d5eb0 100755 --- a/builder/scripts.d/50-dav1d.sh +++ b/builder/scripts.d/50-dav1d.sh @@ -1,7 +1,7 @@ #!/bin/bash SCRIPT_REPO="https://code.videolan.org/videolan/dav1d.git" -SCRIPT_COMMIT="16c943484e63da7ed8a8a8d85af88995369f23cd" +SCRIPT_COMMIT="922bd82b4e0b4df1e23d8d4325417be812dfa23a" ffbuild_enabled() { return 0 diff --git a/builder/scripts.d/50-libass.sh b/builder/scripts.d/50-libass.sh index ce74a4d53f..798bb98cfe 100755 --- a/builder/scripts.d/50-libass.sh +++ b/builder/scripts.d/50-libass.sh @@ -1,7 +1,7 @@ #!/bin/bash SCRIPT_REPO="https://github.com/libass/libass.git" -SCRIPT_COMMIT="218dacece7d24b45e4637ced4dc56564de29919d" +SCRIPT_COMMIT="9b3c0d5b350e13eb69dc87b52751ec0ad44280af" ffbuild_enabled() { return 0 diff --git a/builder/scripts.d/50-libvpx.sh b/builder/scripts.d/50-libvpx.sh index 009d427643..dbeda7f042 100755 --- a/builder/scripts.d/50-libvpx.sh +++ b/builder/scripts.d/50-libvpx.sh @@ -1,7 +1,7 @@ #!/bin/bash SCRIPT_REPO="https://chromium.googlesource.com/webm/libvpx" -SCRIPT_COMMIT="6788c75055899796c13787ed44cc6f1cc45e09d7" +SCRIPT_COMMIT="0f42bd3fb81cd79dfaef5c6aef26c643c48be909" ffbuild_enabled() { return 0 diff --git a/builder/scripts.d/50-libwebp.sh b/builder/scripts.d/50-libwebp.sh index 8097ead140..aba5ffadc2 100755 --- a/builder/scripts.d/50-libwebp.sh +++ b/builder/scripts.d/50-libwebp.sh @@ -1,7 +1,7 @@ #!/bin/bash SCRIPT_REPO="https://chromium.googlesource.com/webm/libwebp" -SCRIPT_COMMIT="0825faa4c10d622f2747c5b2752e7b6a6848bf3f" +SCRIPT_COMMIT="349f4353dd638f2f85e7dd2d9df9a68f54d9919c" ffbuild_enabled() { return 0 diff --git a/builder/scripts.d/50-svtav1.sh b/builder/scripts.d/50-svtav1.sh index 8d7ec41860..3feb503c2d 100755 --- a/builder/scripts.d/50-svtav1.sh +++ b/builder/scripts.d/50-svtav1.sh @@ -1,7 +1,7 @@ #!/bin/bash SCRIPT_REPO="https://gitlab.com/AOMediaCodec/SVT-AV1.git" -SCRIPT_COMMIT="72fd3d479b315fbf6bd1d7e80dec7afe42037819" +SCRIPT_COMMIT="f338431ff69b258686b575ef9a0520fa857399fe" ffbuild_enabled() { [[ $TARGET == win32 ]] && return -1 diff --git a/builder/scripts.d/50-vaapi/30-libpciaccess.sh b/builder/scripts.d/50-vaapi/30-libpciaccess.sh index d284296b2c..ff55e48f67 100755 --- a/builder/scripts.d/50-vaapi/30-libpciaccess.sh +++ b/builder/scripts.d/50-vaapi/30-libpciaccess.sh @@ -1,7 +1,7 @@ #!/bin/bash SCRIPT_REPO="https://gitlab.freedesktop.org/xorg/lib/libpciaccess.git" -SCRIPT_COMMIT="833c86ce15cee2a84a37ae71015f236fd32615d9" +SCRIPT_COMMIT="8980a39004b10298a7db1f1b0b711a19e8b73aee" ffbuild_enabled() { [[ $TARGET != linux* ]] && return -1 diff --git a/builder/scripts.d/50-vaapi/40-libdrm.sh b/builder/scripts.d/50-vaapi/40-libdrm.sh index 524466a71c..49bea6ef31 100755 --- a/builder/scripts.d/50-vaapi/40-libdrm.sh +++ b/builder/scripts.d/50-vaapi/40-libdrm.sh @@ -1,7 +1,7 @@ #!/bin/bash SCRIPT_REPO="https://gitlab.freedesktop.org/mesa/drm.git" -SCRIPT_COMMIT="d1681af05471176215ad3d437249c38768dc959f" +SCRIPT_COMMIT="28d9a3c4fb4c99aafc31b288b3f735e19e728d64" ffbuild_enabled() { [[ $TARGET != linux* ]] && return -1 diff --git a/builder/scripts.d/50-vaapi/50-libva.sh b/builder/scripts.d/50-vaapi/50-libva.sh index c38106e8db..87ac41e9e1 100755 --- a/builder/scripts.d/50-vaapi/50-libva.sh +++ b/builder/scripts.d/50-vaapi/50-libva.sh @@ -1,7 +1,7 @@ #!/bin/bash SCRIPT_REPO="https://github.com/intel/libva.git" -SCRIPT_COMMIT="0fa448dd00c509f08cf53b9a95a55922ff09d5cf" +SCRIPT_COMMIT="3aadcf0e8035e320b6be1fedb4d5fb5182b63377" ffbuild_enabled() { [[ $TARGET != linux* ]] && return -1 diff --git a/builder/scripts.d/50-vulkan/45-vulkan.sh b/builder/scripts.d/50-vulkan/45-vulkan.sh index 1ef7d963fb..a69706aea2 100755 --- a/builder/scripts.d/50-vulkan/45-vulkan.sh +++ b/builder/scripts.d/50-vulkan/45-vulkan.sh @@ -1,7 +1,7 @@ #!/bin/bash SCRIPT_REPO="https://github.com/KhronosGroup/Vulkan-Headers.git" -SCRIPT_COMMIT="v1.3.240" +SCRIPT_COMMIT="v1.3.246" SCRIPT_TAGFILTER="v?.*.*" ffbuild_enabled() { diff --git a/builder/scripts.d/50-vulkan/50-shaderc.sh b/builder/scripts.d/50-vulkan/50-shaderc.sh index c9720278f2..2e1a27eaed 100755 --- a/builder/scripts.d/50-vulkan/50-shaderc.sh +++ b/builder/scripts.d/50-vulkan/50-shaderc.sh @@ -1,7 +1,7 @@ #!/bin/bash SCRIPT_REPO="https://github.com/google/shaderc.git" -SCRIPT_COMMIT="f1268e6d36cfdcd647d5c7032a6c61a0aad8487b" +SCRIPT_COMMIT="4a8f5e537f20bbcfe4b11f1fe45314f1dcbfddf6" ffbuild_enabled() { return 0 diff --git a/builder/scripts.d/50-vulkan/55-spirv-cross.sh b/builder/scripts.d/50-vulkan/55-spirv-cross.sh index 6efd2ca6a0..405ca5c6b0 100755 --- a/builder/scripts.d/50-vulkan/55-spirv-cross.sh +++ b/builder/scripts.d/50-vulkan/55-spirv-cross.sh @@ -1,7 +1,7 @@ #!/bin/bash SCRIPT_REPO="https://github.com/KhronosGroup/SPIRV-Cross.git" -SCRIPT_COMMIT="d26c233e1c2629fec1ae1b6fdf538958e5d52bff" +SCRIPT_COMMIT="3327924addfcffe8f1a4119cb09c1ad353fe1190" ffbuild_enabled() { return 0 diff --git a/builder/scripts.d/50-vulkan/60-libplacebo.sh b/builder/scripts.d/50-vulkan/60-libplacebo.sh index 808e740f3a..a7c85d17ea 100755 --- a/builder/scripts.d/50-vulkan/60-libplacebo.sh +++ b/builder/scripts.d/50-vulkan/60-libplacebo.sh @@ -1,7 +1,7 @@ #!/bin/bash SCRIPT_REPO="https://code.videolan.org/videolan/libplacebo.git" -SCRIPT_COMMIT="2394aea167a5995860387a2511f8972143cd5b82" +SCRIPT_COMMIT="ed29e541a55acf28022738440b2a925386292551" ffbuild_enabled() { return 0 diff --git a/builder/scripts.d/50-x265.sh b/builder/scripts.d/50-x265.sh index 15ace41630..36ab24b1ff 100755 --- a/builder/scripts.d/50-x265.sh +++ b/builder/scripts.d/50-x265.sh @@ -1,7 +1,7 @@ #!/bin/bash SCRIPT_REPO="https://bitbucket.org/multicoreware/x265_git.git" -SCRIPT_COMMIT="38cf1c379b5af08856bb2fdd65f65a1f99384886" +SCRIPT_COMMIT="c07d076cf20ab41d8077a804dda70fdbc3b97386" ffbuild_enabled() { [[ $VARIANT == lgpl* ]] && return -1 diff --git a/docker-build-win64.sh b/docker-build-win64.sh index 6168c46f11..b14c8e99c1 100755 --- a/docker-build-win64.sh +++ b/docker-build-win64.sh @@ -320,7 +320,7 @@ make install popd # LIBWEBP -git clone -b v1.2.3 --depth=1 https://chromium.googlesource.com/webm/libwebp +git clone -b v1.3.0 --depth=1 https://chromium.googlesource.com/webm/libwebp pushd libwebp ./autogen.sh ./configure \ @@ -373,11 +373,14 @@ make install popd # X265 -git clone -b 3.5 --depth=1 https://bitbucket.org/multicoreware/x265_git.git +git clone --depth=1 https://bitbucket.org/multicoreware/x265_git.git pushd x265_git +# Wa for https://bitbucket.org/multicoreware/x265_git/issues/624 +rm -rf .git x265_conf=" -DCMAKE_TOOLCHAIN_FILE=${FF_CMAKE_TOOLCHAIN} -DCMAKE_INSTALL_PREFIX=${FF_DEPS_PREFIX} + -DCMAKE_BUILD_TYPE=Release -DCMAKE_ASM_NASM_FLAGS=-w-macro-params-legacy -DENABLE_ASSEMBLY=ON -DENABLE_SHARED=OFF @@ -429,6 +432,7 @@ SAVE END EOF make install +echo "Libs.private: -lstdc++" >> ${FF_DEPS_PREFIX}/lib/pkgconfig/x265.pc popd popd diff --git a/docker-build.sh b/docker-build.sh index 84c7d89e50..d92254a726 100755 --- a/docker-build.sh +++ b/docker-build.sh @@ -251,7 +251,7 @@ prepare_extra_amd64() { # Provides MSDK runtime (libmfxhw64.so.1) for 11th Gen Rocket Lake and older # Provides MFX dispatcher (libmfx.so.1) for FFmpeg pushd ${SOURCE_DIR} - git clone -b intel-mediasdk-23.1.4 --depth=1 https://github.com/Intel-Media-SDK/MediaSDK.git + git clone -b intel-mediasdk-23.1.5 --depth=1 https://github.com/Intel-Media-SDK/MediaSDK.git pushd MediaSDK sed -i 's|MFX_PLUGINS_CONF_DIR "/plugins.cfg"|"/usr/lib/jellyfin-ffmpeg/lib/mfx/plugins.cfg"|g' api/mfx_dispatch/linux/mfxloader.cpp mkdir build && pushd build @@ -292,7 +292,7 @@ prepare_extra_amd64() { # Provides VPL runtime (libmfx-gen.so.1.2) for 11th Gen Tiger Lake and newer # Both MSDK and VPL runtime can be loaded by MFX dispatcher (libmfx.so.1) pushd ${SOURCE_DIR} - git clone -b intel-onevpl-23.1.4 --depth=1 https://github.com/oneapi-src/oneVPL-intel-gpu.git + git clone -b intel-onevpl-23.1.5 --depth=1 https://github.com/oneapi-src/oneVPL-intel-gpu.git pushd oneVPL-intel-gpu mkdir build && pushd build cmake -DCMAKE_INSTALL_PREFIX=${TARGET_DIR} .. @@ -306,7 +306,7 @@ prepare_extra_amd64() { # Full Feature Build: ENABLE_KERNELS=ON(Default) ENABLE_NONFREE_KERNELS=ON(Default) # Free Kernel Build: ENABLE_KERNELS=ON ENABLE_NONFREE_KERNELS=OFF pushd ${SOURCE_DIR} - git clone -b intel-media-23.1.4 --depth=1 https://github.com/intel/media-driver.git + git clone -b intel-media-23.1.5 --depth=1 https://github.com/intel/media-driver.git pushd media-driver # Possible fix for TGLx timeout caused by 'HCP Scalability Decode' under heavy load wget -q -O - https://github.com/intel/media-driver/commit/284750bf.patch | git apply @@ -329,7 +329,7 @@ prepare_extra_amd64() { # Vulkan Headers pushd ${SOURCE_DIR} - git clone -b v1.3.240 --depth=1 https://github.com/KhronosGroup/Vulkan-Headers.git + git clone -b v1.3.246 --depth=1 https://github.com/KhronosGroup/Vulkan-Headers.git pushd Vulkan-Headers mkdir build && pushd build cmake \ @@ -342,7 +342,7 @@ prepare_extra_amd64() { # Vulkan ICD Loader pushd ${SOURCE_DIR} - git clone -b v1.3.240 --depth=1 https://github.com/KhronosGroup/Vulkan-Loader.git + git clone -b v1.3.246 --depth=1 https://github.com/KhronosGroup/Vulkan-Loader.git pushd Vulkan-Loader mkdir build && pushd build cmake \ @@ -392,7 +392,7 @@ prepare_extra_amd64() { pushd ${SOURCE_DIR} git clone -b main https://gitlab.freedesktop.org/mesa/mesa.git pushd mesa - git reset --hard "f39ffc69" + git reset --hard "1995762d" popd # disable the broken hevc packed header MESA_VA_PIC="mesa/src/gallium/frontends/va/picture.c" @@ -453,7 +453,7 @@ prepare_extra_amd64() { # LIBPLACEBO pushd ${SOURCE_DIR} - git clone -b v5.229.2 --recursive --depth=1 https://github.com/haasn/libplacebo.git + git clone -b v5.264.1 --recursive --depth=1 https://github.com/haasn/libplacebo.git sed -i 's|env: python_env,||g' libplacebo/src/vulkan/meson.build meson setup libplacebo placebo_build \ --prefix=${TARGET_DIR} \ From 80b91ed0608eac099ad902ea50fda4ce3bc6522a Mon Sep 17 00:00:00 2001 From: nyanmisaka Date: Wed, 29 Mar 2023 21:56:48 +0800 Subject: [PATCH 3/7] Bump version to 6.0-2 Signed-off-by: nyanmisaka --- build.yaml | 2 +- debian/changelog | 7 +++++++ 2 files changed, 8 insertions(+), 1 deletion(-) diff --git a/build.yaml b/build.yaml index 43219f1feb..cee62a3583 100644 --- a/build.yaml +++ b/build.yaml @@ -1,7 +1,7 @@ --- # We just wrap `build` so this is really it name: "jellyfin-ffmpeg" -version: "6.0-1" +version: "6.0-2" packages: - buster-amd64 - buster-armhf diff --git a/debian/changelog b/debian/changelog index e46c6f8f4b..7d929e92fb 100644 --- a/debian/changelog +++ b/debian/changelog @@ -1,3 +1,10 @@ +jellyfin-ffmpeg (6.0-2) unstable; urgency=medium + + * Add a tone-mapping mode RGB in CUDA & OpenCL + * Update dependencies + + -- nyanmisaka Tue, 29 Mar 2023 21:55:14 +0800 + jellyfin-ffmpeg (6.0-1) unstable; urgency=medium * New upstream version 6.0 From efbd38922324ec9492996b60269bbf0cc2860a53 Mon Sep 17 00:00:00 2001 From: nyanmisaka Date: Thu, 30 Mar 2023 12:43:39 +0800 Subject: [PATCH 4/7] Enable libopenmpt Signed-off-by: nyanmisaka --- builder/scripts.d/50-openmpt.sh | 69 +++++++++++++++++++++++++++++++++ debian/changelog | 1 + debian/control | 2 + debian/rules | 1 + docker-build-win64.sh | 20 ++++++++++ 5 files changed, 93 insertions(+) create mode 100755 builder/scripts.d/50-openmpt.sh diff --git a/builder/scripts.d/50-openmpt.sh b/builder/scripts.d/50-openmpt.sh new file mode 100755 index 0000000000..df2a27ed0b --- /dev/null +++ b/builder/scripts.d/50-openmpt.sh @@ -0,0 +1,69 @@ +#!/bin/bash + +SCRIPT_REPO="https://source.openmpt.org/svn/openmpt/trunk/OpenMPT" +SCRIPT_REV="18952" + +ffbuild_enabled() { + return 0 +} + +ffbuild_dockerbuild() { + retry-tool sh -c "rm -rf openmpt && svn checkout '${SCRIPT_REPO}@${SCRIPT}' openmpt" + cd openmpt + + local myconf=( + PREFIX="$FFBUILD_PREFIX" + CXXSTDLIB_PCLIBSPRIVATE="-lstdc++" + VERBOSE=2 + STATIC_LIB=1 + SHARED_LIB=0 + DYNLINK=0 + EXAMPLES=0 + OPENMPT123=0 + IN_OPENMPT=0 + XMP_OPENMPT=0 + DEBUG=0 + OPTIMIZE=1 + TEST=0 + MODERN=1 + FORCE_DEPS=1 + NO_MINIMP3=0 + NO_ZLIB=0 + NO_OGG=0 + NO_VORBIS=0 + NO_VORBISFILE=0 + NO_MPG123=1 + NO_SDL2=1 + NO_PULSEAUDIO=1 + NO_SNDFILE=1 + NO_PORTAUDIO=1 + NO_PORTAUDIOCPP=1 + NO_FLAC=1 + ) + + if [[ $TARGET == win* ]]; then + myconf+=( + CONFIG=mingw64-"$TARGET" + ) + export CPPFLAGS="$CPPFLAGS -DMPT_WITH_MINGWSTDTHREADS" + elif [[ $TARGET == linux* ]]; then + myconf+=( + CONFIG=gcc + TOOLCHAIN_PREFIX="$FFBUILD_CROSS_PREFIX" + ) + else + echo "Unknown target" + return -1 + fi + + make -j$(nproc) "${myconf[@]}" all install + rm -r "$FFBUILD_PREFIX"/share/doc/libopenmpt +} + +ffbuild_configure() { + echo --enable-libopenmpt +} + +ffbuild_unconfigure() { + echo --disable-libopenmpt +} diff --git a/debian/changelog b/debian/changelog index 7d929e92fb..4c3cdaf252 100644 --- a/debian/changelog +++ b/debian/changelog @@ -2,6 +2,7 @@ jellyfin-ffmpeg (6.0-2) unstable; urgency=medium * Add a tone-mapping mode RGB in CUDA & OpenCL * Update dependencies + * Enable libopenmpt -- nyanmisaka Tue, 29 Mar 2023 21:55:14 +0800 diff --git a/debian/control b/debian/control index 7f78c16515..26a3e5cbfa 100644 --- a/debian/control +++ b/debian/control @@ -35,6 +35,8 @@ Build-Depends: libtheora-dev, # --enable-libvorbis libvorbis-dev, +# --enable-libopenmpt + libopenmpt-dev, # --enable-libwebp libwebp-dev, # --enable-libvpx diff --git a/debian/rules b/debian/rules index daef3b0c00..078737195b 100755 --- a/debian/rules +++ b/debian/rules @@ -33,6 +33,7 @@ CONFIG := --prefix=${TARGET_DIR} \ --enable-libopus \ --enable-libtheora \ --enable-libvorbis \ + --enable-libopenmpt \ --enable-libdav1d \ --enable-libwebp \ --enable-libvpx \ diff --git a/docker-build-win64.sh b/docker-build-win64.sh index b14c8e99c1..f161a7aa66 100755 --- a/docker-build-win64.sh +++ b/docker-build-win64.sh @@ -319,6 +319,25 @@ make -j$(nproc) make install popd +# OPENMPT +mkdir mpt +pushd mpt +mpt_ver="0.6.9" +mpt_link="https://lib.openmpt.org/files/libopenmpt/src/libopenmpt-${mpt_ver}+release.autotools.tar.gz" +wget ${mpt_link} -O mpt.tar.gz +tar xaf mpt.tar.gz +pushd libopenmpt-${mpt_ver}+release.autotools +./configure \ + --prefix=${FF_DEPS_PREFIX} \ + --host=${FF_TOOLCHAIN} \ + --enable-static \ + --disable-{shared,examples,tests,openmpt123} \ + --without-{mpg123,portaudio,portaudiocpp,sndfile,flac} +make -j$(nproc) +make install +popd +popd + # LIBWEBP git clone -b v1.3.0 --depth=1 https://chromium.googlesource.com/webm/libwebp pushd libwebp @@ -595,6 +614,7 @@ fi --enable-libopus \ --enable-libtheora \ --enable-libvorbis \ + --enable-libopenmpt \ --enable-libwebp \ --enable-libvpx \ --enable-libzimg \ From 9bdac4e1f553dad62428e6c494d95b5116952158 Mon Sep 17 00:00:00 2001 From: nyanmisaka Date: Fri, 31 Mar 2023 00:25:25 +0800 Subject: [PATCH 5/7] Fix vaSyncSurface for AMD in some cases Signed-off-by: nyanmisaka --- .../0044-add-fixes-for-vaapi-drm-vulkan-tearing.patch | 5 ++--- 1 file changed, 2 insertions(+), 3 deletions(-) diff --git a/debian/patches/0044-add-fixes-for-vaapi-drm-vulkan-tearing.patch b/debian/patches/0044-add-fixes-for-vaapi-drm-vulkan-tearing.patch index 3427abd9a2..01c40d5fc7 100644 --- a/debian/patches/0044-add-fixes-for-vaapi-drm-vulkan-tearing.patch +++ b/debian/patches/0044-add-fixes-for-vaapi-drm-vulkan-tearing.patch @@ -2,7 +2,7 @@ Index: jellyfin-ffmpeg/libavutil/hwcontext_vaapi.c =================================================================== --- jellyfin-ffmpeg.orig/libavutil/hwcontext_vaapi.c +++ jellyfin-ffmpeg/libavutil/hwcontext_vaapi.c -@@ -1319,8 +1319,17 @@ static int vaapi_map_to_drm_esh(AVHWFram +@@ -1319,8 +1319,16 @@ static int vaapi_map_to_drm_esh(AVHWFram surface_id = (VASurfaceID)(uintptr_t)src->data[3]; export_flags = VA_EXPORT_SURFACE_SEPARATE_LAYERS; @@ -12,9 +12,8 @@ Index: jellyfin-ffmpeg/libavutil/hwcontext_vaapi.c + + vas = vaSyncSurface(hwctx->display, surface_id); + if (vas != VA_STATUS_SUCCESS) { -+ av_log(hwfc, AV_LOG_ERROR, "Failed to sync surface " ++ av_log(hwfc, AV_LOG_WARNING, "Failed to sync surface " + "%#x: %d (%s).\n", surface_id, vas, vaErrorStr(vas)); -+ return AVERROR(EIO); + } + } + From 2256bab44fc809e2a1f3a44e6d012d81925ef6d2 Mon Sep 17 00:00:00 2001 From: nyanmisaka Date: Sat, 1 Apr 2023 18:55:31 +0800 Subject: [PATCH 6/7] Enable AENC in VPL runtime Signed-off-by: nyanmisaka --- docker-build.sh | 8 +++++++- 1 file changed, 7 insertions(+), 1 deletion(-) diff --git a/docker-build.sh b/docker-build.sh index d92254a726..a2aa290521 100755 --- a/docker-build.sh +++ b/docker-build.sh @@ -295,7 +295,13 @@ prepare_extra_amd64() { git clone -b intel-onevpl-23.1.5 --depth=1 https://github.com/oneapi-src/oneVPL-intel-gpu.git pushd oneVPL-intel-gpu mkdir build && pushd build - cmake -DCMAKE_INSTALL_PREFIX=${TARGET_DIR} .. + cmake -DCMAKE_INSTALL_PREFIX=${TARGET_DIR} \ + -DCMAKE_INSTALL_LIBDIR=${TARGET_DIR}/lib \ + -DCMAKE_BUILD_TYPE=Release \ + -DBUILD_RUNTIME=ON \ + -DBUILD_{TESTS,TOOLS}=OFF \ + -DMFX_ENABLE_{KERNELS,ENCTOOLS,AENC}=ON \ + .. make -j$(nproc) && make install && make install DESTDIR=${SOURCE_DIR}/intel echo "intel${TARGET_DIR}/lib/libmfx-gen* usr/lib/jellyfin-ffmpeg/lib" >> ${DPKG_INSTALL_LIST} popd From 9747d76d672462f93467128106cf4f3af36dd7bd Mon Sep 17 00:00:00 2001 From: nyanmisaka Date: Mon, 3 Apr 2023 00:14:36 +0800 Subject: [PATCH 7/7] Revert libplacebo to ad-hoc mode Signed-off-by: nyanmisaka --- ...050-revert-libplacebo-to-ad-hoc-mode.patch | 143 ++++++++++++++++++ debian/patches/series | 1 + 2 files changed, 144 insertions(+) create mode 100644 debian/patches/0050-revert-libplacebo-to-ad-hoc-mode.patch diff --git a/debian/patches/0050-revert-libplacebo-to-ad-hoc-mode.patch b/debian/patches/0050-revert-libplacebo-to-ad-hoc-mode.patch new file mode 100644 index 0000000000..4c89d4b356 --- /dev/null +++ b/debian/patches/0050-revert-libplacebo-to-ad-hoc-mode.patch @@ -0,0 +1,143 @@ +Index: jellyfin-ffmpeg/libavfilter/vf_libplacebo.c +=================================================================== +--- jellyfin-ffmpeg.orig/libavfilter/vf_libplacebo.c ++++ jellyfin-ffmpeg/libavfilter/vf_libplacebo.c +@@ -62,6 +62,7 @@ static const struct pl_tone_map_function + typedef struct LibplaceboContext { + /* lavfi vulkan*/ + FFVulkanContext vkctx; ++ int initialized; + + /* libplacebo */ + pl_log log; +@@ -258,25 +259,10 @@ static int init_vulkan(AVFilterContext * + { + int err = 0; + LibplaceboContext *s = avctx->priv; +- const AVHWDeviceContext *avhwctx; +- const AVVulkanDeviceContext *hwctx; ++ const AVVulkanDeviceContext *hwctx = s->vkctx.hwctx; + uint8_t *buf = NULL; + size_t buf_len; + +- if (!avctx->hw_device_ctx) { +- av_log(s, AV_LOG_ERROR, "Missing vulkan hwdevice for vf_libplacebo.\n"); +- return AVERROR(EINVAL); +- } +- +- avhwctx = (AVHWDeviceContext *) avctx->hw_device_ctx->data; +- if (avhwctx->type != AV_HWDEVICE_TYPE_VULKAN) { +- av_log(s, AV_LOG_ERROR, "Expected vulkan hwdevice for vf_libplacebo, got %s.\n", +- av_hwdevice_get_type_name(avhwctx->type)); +- return AVERROR(EINVAL); +- } +- +- hwctx = avhwctx->hwctx; +- + /* Import libavfilter vulkan context into libplacebo */ + s->vulkan = pl_vulkan_import(s->log, pl_vulkan_import_params( + .instance = hwctx->inst, +@@ -325,6 +311,7 @@ static int init_vulkan(AVFilterContext * + fail: + if (buf) + av_file_unmap(buf, buf_len); ++ s->initialized = 1; + return err; + } + +@@ -340,6 +327,7 @@ static void libplacebo_uninit(AVFilterCo + pl_vulkan_destroy(&s->vulkan); + pl_log_destroy(&s->log); + ff_vk_uninit(&s->vkctx); ++ s->initialized = 0; + s->gpu = NULL; + } + +@@ -500,6 +488,8 @@ static int filter_frame(AVFilterLink *li + } + + pl_log_level_update(s->log, get_log_level()); ++ if (!s->initialized) ++ RET(init_vulkan(ctx)); + + RET(av_frame_copy_props(out, in)); + out->width = outlink->w; +@@ -551,69 +541,6 @@ fail: + return err; + } + +-static int libplacebo_query_format(AVFilterContext *ctx) +-{ +- int err; +- LibplaceboContext *s = ctx->priv; +- const AVPixFmtDescriptor *desc = NULL; +- AVFilterFormats *infmts = NULL, *outfmts = NULL; +- +- RET(init_vulkan(ctx)); +- +- while ((desc = av_pix_fmt_desc_next(desc))) { +- enum AVPixelFormat pixfmt = av_pix_fmt_desc_get_id(desc); +- +-#if PL_API_VER < 232 +- // Older libplacebo can't handle >64-bit pixel formats, so safe-guard +- // this to prevent triggering an assertion +- if (av_get_bits_per_pixel(desc) > 64) +- continue; +-#endif +- +- if (!pl_test_pixfmt(s->gpu, pixfmt)) +- continue; +- +- RET(ff_add_format(&infmts, pixfmt)); +- +- /* Filter for supported output pixel formats */ +- if (desc->flags & AV_PIX_FMT_FLAG_BE) +- continue; /* BE formats are not supported by pl_download_avframe */ +- +- /* Mask based on user specified format */ +- if (s->out_format != AV_PIX_FMT_NONE) { +- if (pixfmt == AV_PIX_FMT_VULKAN && av_vkfmt_from_pixfmt(s->out_format)) { +- /* OK */ +- } else if (pixfmt == s->out_format) { +- /* OK */ +- } else { +- continue; /* Not OK */ +- } +- } +- +- RET(ff_add_format(&outfmts, pixfmt)); +- } +- +- if (!infmts || !outfmts) { +- if (s->out_format) { +- av_log(s, AV_LOG_ERROR, "Invalid output format '%s'!\n", +- av_get_pix_fmt_name(s->out_format)); +- } +- err = AVERROR(EINVAL); +- goto fail; +- } +- +- RET(ff_formats_ref(infmts, &ctx->inputs[0]->outcfg.formats)); +- RET(ff_formats_ref(outfmts, &ctx->outputs[0]->incfg.formats)); +- return 0; +- +-fail: +- if (infmts && !infmts->refcount) +- ff_formats_unref(&infmts); +- if (outfmts && !outfmts->refcount) +- ff_formats_unref(&outfmts); +- return err; +-} +- + static int libplacebo_config_input(AVFilterLink *inlink) + { + AVFilterContext *avctx = inlink->dst; +@@ -881,7 +808,7 @@ const AVFilter ff_vf_libplacebo = { + .process_command = &ff_filter_process_command, + FILTER_INPUTS(libplacebo_inputs), + FILTER_OUTPUTS(libplacebo_outputs), +- FILTER_QUERY_FUNC(libplacebo_query_format), ++ FILTER_SINGLE_PIXFMT(AV_PIX_FMT_VULKAN), + .priv_class = &libplacebo_class, + .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE, + }; diff --git a/debian/patches/series b/debian/patches/series index 00b6cdc5fd..b7ed5eac55 100644 --- a/debian/patches/series +++ b/debian/patches/series @@ -47,3 +47,4 @@ 0047-update-the-default-params-of-qsv-encoders.patch 0048-fix-the-empty-output-in-webvtt-transcoding.patch 0049-add-queued-key-fixes-from-upstream.patch +0050-revert-libplacebo-to-ad-hoc-mode.patch