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

cleanup for adding windows patches

Browse Source
This commit is contained in:
nyanmisaka
2021-12-18 00:34:27 +08:00
parent d1819de544
commit 86b92d1804
10 changed files with 0 additions and 6320 deletions
Download Patch File Download Diff File Expand all files Collapse all files
debian/patches/0001_fix-segment-muxer.patch
Vendored
-37
View File
@@ -1,37 +0,0 @@
Index: jellyfin-ffmpeg/libavformat/segment.c
===================================================================
--- jellyfin-ffmpeg.orig/libavformat/segment.c
+++ jellyfin-ffmpeg/libavformat/segment.c
@@ -87,6 +87,7 @@ typedef struct SegmentContext {
int64_t last_val; ///< remember last time for wrap around detection
int cut_pending;
int header_written; ///< whether we've already called avformat_write_header
+ int64_t start_pts; ///< pts of the very first packet processed, used to compute correct segment length
char *entry_prefix; ///< prefix to add to list entry filenames
int list_type; ///< set the list type
@@ -712,6 +713,7 @@ static int seg_init(AVFormatContext *s)
if ((ret = parse_frames(s, &seg->frames, &seg->nb_frames, seg->frames_str)) < 0)
return ret;
} else {
+ seg->start_pts = -1;
if (seg->use_clocktime) {
if (seg->time <= 0) {
av_log(s, AV_LOG_ERROR, "Invalid negative segment_time with segment_atclocktime option set\n");
@@ -889,7 +891,15 @@ calc_times:
seg->cut_pending = 1;
seg->last_val = wrapped_val;
} else {
- end_pts = seg->time * (seg->segment_count + 1);
+ if (seg->start_pts != -1) {
+ end_pts = seg->start_pts + seg->time * (seg->segment_count + 1);
+ } else if (pkt->stream_index == seg->reference_stream_index && pkt->pts != AV_NOPTS_VALUE) {
+ // this is the first packet of the reference stream we see, initialize start point
+ seg->start_pts = av_rescale_q(pkt->pts, st->time_base, AV_TIME_BASE_Q);
+ seg->cur_entry.start_time = (double)pkt->pts * av_q2d(st->time_base);
+ seg->cur_entry.start_pts = seg->start_pts;
+ end_pts = seg->start_pts + seg->time * (seg->segment_count + 1);
+ }
}
}
debian/patches/0002-lavfi-add-a-filter-for-uploading-normal-frames-to-VAAPI.patch
Vendored
-241
View File
@@ -1,241 +0,0 @@
#From c1fb9225a1b8e26875cb9b4e2b3ae2f4d68c5630 Mon Sep 17 00:00:00 2001
#From: nyanmisaka <nst799610810@gmail.com>
#Date: Sun, 24 Jan 2021 19:58:04 +0800
#Subject: [PATCH] lavfi: add a filter for uploading normal frames to VAAPI
Index: jellyfin-ffmpeg/configure
===================================================================
--- jellyfin-ffmpeg.orig/configure
+++ jellyfin-ffmpeg/configure
@@ -3577,6 +3577,7 @@ fspp_filter_deps="gpl"
headphone_filter_select="fft"
histeq_filter_deps="gpl"
hqdn3d_filter_deps="gpl"
+hwupload_vaapi_filter_deps="vaapi"
interlace_filter_deps="gpl"
kerndeint_filter_deps="gpl"
ladspa_filter_deps="ladspa libdl"
Index: jellyfin-ffmpeg/libavfilter/Makefile
===================================================================
--- jellyfin-ffmpeg.orig/libavfilter/Makefile
+++ jellyfin-ffmpeg/libavfilter/Makefile
@@ -297,6 +297,7 @@ OBJS-$(CONFIG_HUE_FILTER)
OBJS-$(CONFIG_HWDOWNLOAD_FILTER) += vf_hwdownload.o
OBJS-$(CONFIG_HWMAP_FILTER) += vf_hwmap.o
OBJS-$(CONFIG_HWUPLOAD_CUDA_FILTER) += vf_hwupload_cuda.o
+OBJS-$(CONFIG_HWUPLOAD_VAAPI_FILTER) += vf_hwupload_vaapi.o
OBJS-$(CONFIG_HWUPLOAD_FILTER) += vf_hwupload.o
OBJS-$(CONFIG_HYSTERESIS_FILTER) += vf_hysteresis.o framesync.o
OBJS-$(CONFIG_IDENTITY_FILTER) += vf_identity.o
Index: jellyfin-ffmpeg/libavfilter/allfilters.c
===================================================================
--- jellyfin-ffmpeg.orig/libavfilter/allfilters.c
+++ jellyfin-ffmpeg/libavfilter/allfilters.c
@@ -282,6 +282,7 @@ extern AVFilter ff_vf_hwdownload;
extern AVFilter ff_vf_hwmap;
extern AVFilter ff_vf_hwupload;
extern AVFilter ff_vf_hwupload_cuda;
+extern AVFilter ff_vf_hwupload_vaapi;
extern AVFilter ff_vf_hysteresis;
extern AVFilter ff_vf_identity;
extern AVFilter ff_vf_idet;
Index: jellyfin-ffmpeg/libavfilter/vf_hwupload_vaapi.c
===================================================================
--- /dev/null
+++ jellyfin-ffmpeg/libavfilter/vf_hwupload_vaapi.c
@@ -0,0 +1,196 @@
+/*
+ * This file is part of FFmpeg.
+ *
+ * FFmpeg is free software; you can redistribute it and/or
+ * modify it under the terms of the GNU Lesser General Public
+ * License as published by the Free Software Foundation; either
+ * version 2.1 of the License, or (at your option) any later version.
+ *
+ * FFmpeg is distributed in the hope that it will be useful,
+ * but WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU
+ * Lesser General Public License for more details.
+ *
+ * You should have received a copy of the GNU Lesser General Public
+ * License along with FFmpeg; if not, write to the Free Software
+ * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA
+ */
+
+#include "libavutil/buffer.h"
+#include "libavutil/hwcontext.h"
+#include "libavutil/log.h"
+#include "libavutil/opt.h"
+
+#include "avfilter.h"
+#include "formats.h"
+#include "internal.h"
+#include "video.h"
+
+typedef struct VaapiUploadContext {
+ const AVClass *class;
+ int device_idx;
+
+ AVBufferRef *hwdevice;
+ AVBufferRef *hwframe;
+} VaapiUploadContext;
+
+static av_cold int vaapiupload_init(AVFilterContext *ctx)
+{
+ VaapiUploadContext *s = ctx->priv;
+ return av_hwdevice_ctx_create(&s->hwdevice, AV_HWDEVICE_TYPE_VAAPI, NULL, NULL, 0);
+}
+
+static av_cold void vaapiupload_uninit(AVFilterContext *ctx)
+{
+ VaapiUploadContext *s = ctx->priv;
+
+ av_buffer_unref(&s->hwframe);
+ av_buffer_unref(&s->hwdevice);
+}
+
+static int vaapiupload_query_formats(AVFilterContext *ctx)
+{
+ int ret;
+
+ static const enum AVPixelFormat input_pix_fmts[] = {
+ AV_PIX_FMT_NV12, AV_PIX_FMT_YUV420P, AV_PIX_FMT_YUV422P,
+ AV_PIX_FMT_UYVY422, AV_PIX_FMT_YUYV422, AV_PIX_FMT_Y210,
+ AV_PIX_FMT_YUV411P, AV_PIX_FMT_YUV440P, AV_PIX_FMT_YUV444P,
+ AV_PIX_FMT_GRAY8, AV_PIX_FMT_P010, AV_PIX_FMT_BGRA,
+ AV_PIX_FMT_BGR0, AV_PIX_FMT_RGBA, AV_PIX_FMT_RGB0,
+ AV_PIX_FMT_ABGR, AV_PIX_FMT_0BGR, AV_PIX_FMT_ARGB,
+ AV_PIX_FMT_0RGB, AV_PIX_FMT_NONE,
+ };
+ static const enum AVPixelFormat output_pix_fmts[] = {
+ AV_PIX_FMT_VAAPI, AV_PIX_FMT_NONE,
+ };
+ AVFilterFormats *in_fmts = ff_make_format_list(input_pix_fmts);
+ AVFilterFormats *out_fmts;
+
+ ret = ff_formats_ref(in_fmts, &ctx->inputs[0]->outcfg.formats);
+ if (ret < 0)
+ return ret;
+
+ out_fmts = ff_make_format_list(output_pix_fmts);
+
+ ret = ff_formats_ref(out_fmts, &ctx->outputs[0]->incfg.formats);
+ if (ret < 0)
+ return ret;
+
+ return 0;
+}
+
+static int vaapiupload_config_output(AVFilterLink *outlink)
+{
+ AVFilterContext *ctx = outlink->src;
+ AVFilterLink *inlink = ctx->inputs[0];
+ VaapiUploadContext *s = ctx->priv;
+
+ AVHWFramesContext *hwframe_ctx;
+ int ret;
+
+ av_buffer_unref(&s->hwframe);
+ s->hwframe = av_hwframe_ctx_alloc(s->hwdevice);
+ if (!s->hwframe)
+ return AVERROR(ENOMEM);
+
+ hwframe_ctx = (AVHWFramesContext*)s->hwframe->data;
+ hwframe_ctx->format = AV_PIX_FMT_VAAPI;
+ if (inlink->hw_frames_ctx) {
+ AVHWFramesContext *in_hwframe_ctx = (AVHWFramesContext*)inlink->hw_frames_ctx->data;
+ hwframe_ctx->sw_format = in_hwframe_ctx->sw_format;
+ } else {
+ hwframe_ctx->sw_format = inlink->format;
+ }
+ hwframe_ctx->width = inlink->w;
+ hwframe_ctx->height = inlink->h;
+
+ ret = av_hwframe_ctx_init(s->hwframe);
+ if (ret < 0)
+ return ret;
+
+ outlink->hw_frames_ctx = av_buffer_ref(s->hwframe);
+ if (!outlink->hw_frames_ctx)
+ return AVERROR(ENOMEM);
+
+ return 0;
+}
+
+static int vaapiupload_filter_frame(AVFilterLink *link, AVFrame *in)
+{
+ AVFilterContext *ctx = link->dst;
+ AVFilterLink *outlink = ctx->outputs[0];
+
+ AVFrame *out = NULL;
+ int ret;
+
+ out = ff_get_video_buffer(outlink, outlink->w, outlink->h);
+ if (!out) {
+ ret = AVERROR(ENOMEM);
+ goto fail;
+ }
+
+ out->width = in->width;
+ out->height = in->height;
+
+ ret = av_hwframe_transfer_data(out, in, 0);
+ if (ret < 0) {
+ av_log(ctx, AV_LOG_ERROR, "Error transferring data to the GPU\n");
+ goto fail;
+ }
+
+ ret = av_frame_copy_props(out, in);
+ if (ret < 0)
+ goto fail;
+
+ av_frame_free(&in);
+
+ return ff_filter_frame(ctx->outputs[0], out);
+fail:
+ av_frame_free(&in);
+ av_frame_free(&out);
+ return ret;
+}
+
+static const AVClass vaapiupload_class = {
+ .class_name = "vaapiupload",
+ .item_name = av_default_item_name,
+ .option = NULL,
+ .version = LIBAVUTIL_VERSION_INT,
+};
+
+static const AVFilterPad vaapiupload_inputs[] = {
+ {
+ .name = "default",
+ .type = AVMEDIA_TYPE_VIDEO,
+ .filter_frame = vaapiupload_filter_frame,
+ },
+ { NULL }
+};
+
+static const AVFilterPad vaapiupload_outputs[] = {
+ {
+ .name = "default",
+ .type = AVMEDIA_TYPE_VIDEO,
+ .config_props = vaapiupload_config_output,
+ },
+ { NULL }
+};
+
+AVFilter ff_vf_hwupload_vaapi = {
+ .name = "hwupload_vaapi",
+ .description = NULL_IF_CONFIG_SMALL("Upload a system memory frame to a VAAPI device."),
+
+ .init = vaapiupload_init,
+ .uninit = vaapiupload_uninit,
+
+ .query_formats = vaapiupload_query_formats,
+
+ .priv_size = sizeof(VaapiUploadContext),
+ .priv_class = &vaapiupload_class,
+
+ .inputs = vaapiupload_inputs,
+ .outputs = vaapiupload_outputs,
+
+ .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE,
+};
debian/patches/0003-fix-for-the-broken-tonemap_vaapi-filter.patch
Vendored
-259
View File
@@ -1,259 +0,0 @@
# Fix for the broken tonemap_vaapi filter
# avfilter/tonemap_vaapi: pass filter parameters to VA parameter buffer
# avfilter: Add H2H support in tonemap_vaapi
Index: jellyfin-ffmpeg/libavfilter/vf_tonemap_vaapi.c
===================================================================
--- jellyfin-ffmpeg.orig/libavfilter/vf_tonemap_vaapi.c
+++ jellyfin-ffmpeg/libavfilter/vf_tonemap_vaapi.c
@@ -41,7 +41,11 @@ typedef struct HDRVAAPIContext {
enum AVColorTransferCharacteristic color_transfer;
enum AVColorSpace color_matrix;
+ char *master_display;
+ char *content_light;
+
VAHdrMetaDataHDR10 in_metadata;
+ VAHdrMetaDataHDR10 out_metadata;
AVFrameSideData *src_display;
AVFrameSideData *src_light;
@@ -148,6 +152,107 @@ static int tonemap_vaapi_save_metadata(A
return 0;
}
+static int tonemap_vaapi_update_sidedata(AVFilterContext *avctx, AVFrame *output_frame)
+{
+ HDRVAAPIContext *ctx = avctx->priv;
+ AVFrameSideData *metadata;
+ AVMasteringDisplayMetadata *hdr_meta;
+ AVFrameSideData *metadata_lt;
+ AVContentLightMetadata *hdr_meta_lt;
+
+ int i;
+ const int mapping[3] = {1, 2, 0}; //green, blue, red
+ const int chroma_den = 50000;
+ const int luma_den = 10000;
+
+ metadata = av_frame_get_side_data(output_frame,
+ AV_FRAME_DATA_MASTERING_DISPLAY_METADATA);
+ if (metadata) {
+ av_frame_remove_side_data(output_frame,
+ AV_FRAME_DATA_MASTERING_DISPLAY_METADATA);
+ metadata = av_frame_new_side_data(output_frame,
+ AV_FRAME_DATA_MASTERING_DISPLAY_METADATA,
+ sizeof(AVMasteringDisplayMetadata));
+ } else {
+ metadata = av_frame_new_side_data(output_frame,
+ AV_FRAME_DATA_MASTERING_DISPLAY_METADATA,
+ sizeof(AVMasteringDisplayMetadata));
+ }
+
+ hdr_meta = (AVMasteringDisplayMetadata *)metadata->data;
+
+ for (i = 0; i < 3; i++) {
+ const int j = mapping[i];
+ hdr_meta->display_primaries[j][0].num = ctx->out_metadata.display_primaries_x[i];
+ hdr_meta->display_primaries[j][0].den = chroma_den;
+
+ hdr_meta->display_primaries[j][1].num = ctx->out_metadata.display_primaries_y[i];
+ hdr_meta->display_primaries[j][1].den = chroma_den;
+ }
+
+ hdr_meta->white_point[0].num = ctx->out_metadata.white_point_x;
+ hdr_meta->white_point[0].den = chroma_den;
+
+ hdr_meta->white_point[1].num = ctx->out_metadata.white_point_y;
+ hdr_meta->white_point[1].den = chroma_den;
+ hdr_meta->has_primaries = 1;
+
+ hdr_meta->max_luminance.num = ctx->out_metadata.max_display_mastering_luminance;
+ hdr_meta->max_luminance.den = luma_den;
+
+ hdr_meta->min_luminance.num = ctx->out_metadata.min_display_mastering_luminance;
+ hdr_meta->min_luminance.den = luma_den;
+ hdr_meta->has_luminance = 1;
+
+ av_log(avctx, AV_LOG_DEBUG,
+ "Mastering Display Metadata(out luminance):\n");
+ av_log(avctx, AV_LOG_DEBUG,
+ "min_luminance=%u, max_luminance=%u\n",
+ ctx->out_metadata.min_display_mastering_luminance,
+ ctx->out_metadata.max_display_mastering_luminance);
+
+ av_log(avctx, AV_LOG_DEBUG,
+ "Mastering Display Metadata(out primaries):\n");
+ av_log(avctx, AV_LOG_DEBUG,
+ "G(%u,%u) B(%u,%u) R(%u,%u) WP(%u,%u)\n",
+ ctx->out_metadata.display_primaries_x[0],
+ ctx->out_metadata.display_primaries_y[0],
+ ctx->out_metadata.display_primaries_x[1],
+ ctx->out_metadata.display_primaries_y[1],
+ ctx->out_metadata.display_primaries_x[2],
+ ctx->out_metadata.display_primaries_y[2],
+ ctx->out_metadata.white_point_x,
+ ctx->out_metadata.white_point_y);
+
+ metadata_lt = av_frame_get_side_data(output_frame,
+ AV_FRAME_DATA_CONTENT_LIGHT_LEVEL);
+ if (metadata_lt) {
+ av_frame_remove_side_data(output_frame,
+ AV_FRAME_DATA_CONTENT_LIGHT_LEVEL);
+ metadata_lt = av_frame_new_side_data(output_frame,
+ AV_FRAME_DATA_CONTENT_LIGHT_LEVEL,
+ sizeof(AVContentLightMetadata));
+ } else {
+ metadata_lt = av_frame_new_side_data(output_frame,
+ AV_FRAME_DATA_CONTENT_LIGHT_LEVEL,
+ sizeof(AVContentLightMetadata));
+ }
+
+ hdr_meta_lt = (AVContentLightMetadata *)metadata_lt->data;
+
+ hdr_meta_lt->MaxCLL = FFMIN(ctx->out_metadata.max_content_light_level, 65535);
+ hdr_meta_lt->MaxFALL = FFMIN(ctx->out_metadata.max_pic_average_light_level, 65535);
+
+ av_log(avctx, AV_LOG_DEBUG,
+ "Mastering Content Light Level (out):\n");
+ av_log(avctx, AV_LOG_DEBUG,
+ "MaxCLL(%u) MaxFALL(%u)\n",
+ ctx->out_metadata.max_content_light_level,
+ ctx->out_metadata.max_pic_average_light_level);
+
+ return 0;
+}
+
static int tonemap_vaapi_set_filter_params(AVFilterContext *avctx, AVFrame *input_frame)
{
VAAPIVPPContext *vpp_ctx = avctx->priv;
@@ -210,15 +315,26 @@ static int tonemap_vaapi_build_filter_pa
return AVERROR(EINVAL);
}
- for (i = 0; i < num_query_caps; i++) {
- if (VA_TONE_MAPPING_HDR_TO_SDR & hdr_cap[i].caps_flag)
- break;
- }
-
- if (i >= num_query_caps) {
- av_log(avctx, AV_LOG_ERROR,
- "VAAPI driver doesn't support HDR to SDR\n");
- return AVERROR(EINVAL);
+ if (ctx->color_transfer == AVCOL_TRC_SMPTE2084) {
+ for (i = 0; i < num_query_caps; i++) {
+ if (VA_TONE_MAPPING_HDR_TO_HDR & hdr_cap[i].caps_flag)
+ break;
+ }
+ if (i >= num_query_caps) {
+ av_log(avctx, AV_LOG_ERROR,
+ "VAAPI driver doesn't support HDR to HDR\n");
+ return AVERROR(EINVAL);
+ }
+ } else {
+ for (i = 0; i < num_query_caps; i++) {
+ if (VA_TONE_MAPPING_HDR_TO_SDR & hdr_cap[i].caps_flag)
+ break;
+ }
+ if (i >= num_query_caps) {
+ av_log(avctx, AV_LOG_ERROR,
+ "VAAPI driver doesn't support HDR to SDR\n");
+ return AVERROR(EINVAL);
+ }
}
hdrtm_param.type = VAProcFilterHighDynamicRangeToneMapping;
@@ -243,6 +359,8 @@ static int tonemap_vaapi_filter_frame(AV
VAProcPipelineParameterBuffer params;
int err;
+ VAHdrMetaData out_hdr_metadata;
+
av_log(avctx, AV_LOG_DEBUG, "Filter input: %s, %ux%u (%"PRId64").\n",
av_get_pix_fmt_name(input_frame->format),
input_frame->width, input_frame->height, input_frame->pts);
@@ -291,11 +409,26 @@ static int tonemap_vaapi_filter_frame(AV
if (ctx->color_matrix != AVCOL_SPC_UNSPECIFIED)
output_frame->colorspace = ctx->color_matrix;
+ if (output_frame->color_trc == AVCOL_TRC_SMPTE2084) {
+ err = tonemap_vaapi_update_sidedata(avctx, output_frame);
+ if (err < 0)
+ goto fail;
+
+ out_hdr_metadata.metadata_type = VAProcHighDynamicRangeMetadataHDR10;
+ out_hdr_metadata.metadata = &ctx->out_metadata;
+ out_hdr_metadata.metadata_size = sizeof(VAHdrMetaDataHDR10);
+
+ params.output_hdr_metadata = &out_hdr_metadata;
+ }
+
err = ff_vaapi_vpp_init_params(avctx, &params,
input_frame, output_frame);
if (err < 0)
goto fail;
+ params.filters = &vpp_ctx->filter_buffers[0];
+ params.num_filters = vpp_ctx->nb_filter_buffers;
+
err = ff_vaapi_vpp_render_picture(avctx, &params, output_frame);
if (err < 0)
goto fail;
@@ -355,6 +488,46 @@ static av_cold int tonemap_vaapi_init(AV
STRING_OPTION(color_transfer, color_transfer, AVCOL_TRC_UNSPECIFIED);
STRING_OPTION(color_matrix, color_space, AVCOL_SPC_UNSPECIFIED);
+ if (ctx->color_transfer == AVCOL_TRC_SMPTE2084) {
+ if (!ctx->master_display) {
+ av_log(avctx, AV_LOG_ERROR,
+ "Option mastering-display input invalid\n");
+ return AVERROR(EINVAL);
+ }
+
+ if (10 != sscanf(ctx->master_display,
+ "G(%hu|%hu)B(%hu|%hu)R(%hu|%hu)WP(%hu|%hu)L(%u|%u)",
+ &ctx->out_metadata.display_primaries_x[0],
+ &ctx->out_metadata.display_primaries_y[0],
+ &ctx->out_metadata.display_primaries_x[1],
+ &ctx->out_metadata.display_primaries_y[1],
+ &ctx->out_metadata.display_primaries_x[2],
+ &ctx->out_metadata.display_primaries_y[2],
+ &ctx->out_metadata.white_point_x,
+ &ctx->out_metadata.white_point_y,
+ &ctx->out_metadata.min_display_mastering_luminance,
+ &ctx->out_metadata.max_display_mastering_luminance)) {
+ av_log(avctx, AV_LOG_ERROR,
+ "Option mastering-display input invalid\n");
+ return AVERROR(EINVAL);
+ }
+
+ if (!ctx->content_light) {
+ av_log(avctx, AV_LOG_ERROR,
+ "Option content-light input invalid\n");
+ return AVERROR(EINVAL);
+ }
+
+ if (2 != sscanf(ctx->content_light,
+ "CLL(%hu)FALL(%hu)",
+ &ctx->out_metadata.max_content_light_level,
+ &ctx->out_metadata.max_pic_average_light_level)) {
+ av_log(avctx, AV_LOG_ERROR,
+ "Option content-light input invalid\n");
+ return AVERROR(EINVAL);
+ }
+ }
+
return 0;
}
@@ -380,10 +553,11 @@ static const AVOption tonemap_vaapi_opti
{ "t", "Output color transfer characteristics set",
OFFSET(color_transfer_string), AV_OPT_TYPE_STRING,
{ .str = NULL }, .flags = FLAGS, "transfer" },
+ { "display", "set master display", OFFSET(master_display), AV_OPT_TYPE_STRING, {.str=NULL}, CHAR_MIN, CHAR_MAX, FLAGS },
+ { "light", "set content light", OFFSET(content_light), AV_OPT_TYPE_STRING, {.str=NULL}, CHAR_MIN, CHAR_MAX, FLAGS },
{ NULL }
};
-
AVFILTER_DEFINE_CLASS(tonemap_vaapi);
static const AVFilterPad tonemap_vaapi_inputs[] = {
debian/patches/0004-cuda-format-converter-impl.patch
Vendored
-1438
View File
File diff suppressed because it is too large Load Diff
debian/patches/0005-cuda-tonemap-impl.patch
Vendored
-1824
View File
File diff suppressed because it is too large Load Diff
debian/patches/0006-bt2390-and-fix-for-peak-detection-in-opencl-tonemap.patch
Vendored
-755
View File
@@ -1,755 +0,0 @@
Index: jellyfin-ffmpeg/libavfilter/opencl/colorspace_common.cl
===================================================================
--- jellyfin-ffmpeg.orig/libavfilter/opencl/colorspace_common.cl
+++ jellyfin-ffmpeg/libavfilter/opencl/colorspace_common.cl
@@ -16,8 +16,23 @@
* Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA
*/
+#define BT709_ALPHA 1.09929682680944f
+#define BT709_BETA 0.018053968510807f
+
#define ST2084_MAX_LUMINANCE 10000.0f
-#define REFERENCE_WHITE 100.0f
+#define REFERENCE_WHITE 203.0f
+
+#define ST2084_M1 0.1593017578125f
+#define ST2084_M2 78.84375f
+#define ST2084_C1 0.8359375f
+#define ST2084_C2 18.8515625f
+#define ST2084_C3 18.6875f
+
+#define ARIB_B67_A 0.17883277f
+#define ARIB_B67_B 0.28466892f
+#define ARIB_B67_C 0.55991073f
+
+#define FLOAT_EPS 1.175494351e-38f
#if chroma_loc == 1
#define chroma_sample(a,b,c,d) (((a) + (c)) * 0.5f)
@@ -33,12 +48,6 @@
#define chroma_sample(a,b,c,d) (((a) + (b) + (c) + (d)) * 0.25f)
#endif
-constant const float ST2084_M1 = 0.1593017578125f;
-constant const float ST2084_M2 = 78.84375f;
-constant const float ST2084_C1 = 0.8359375f;
-constant const float ST2084_C2 = 18.8515625f;
-constant const float ST2084_C3 = 18.6875f;
-
float get_luma_dst(float3 c) {
return luma_dst.x * c.x + luma_dst.y * c.y + luma_dst.z * c.z;
}
@@ -51,61 +60,87 @@ float3 get_chroma_sample(float3 a, float
return chroma_sample(a, b, c, d);
}
+// linearizer for PQ/ST2084
float eotf_st2084(float x) {
- float p = powr(x, 1.0f / ST2084_M2);
- float a = max(p -ST2084_C1, 0.0f);
- float b = max(ST2084_C2 - ST2084_C3 * p, 1e-6f);
- float c = powr(a / b, 1.0f / ST2084_M1);
- return x > 0.0f ? c * ST2084_MAX_LUMINANCE / REFERENCE_WHITE : 0.0f;
-}
-
-__constant const float HLG_A = 0.17883277f;
-__constant const float HLG_B = 0.28466892f;
-__constant const float HLG_C = 0.55991073f;
-
-// linearizer for HLG
-float inverse_oetf_hlg(float x) {
- float a = 4.0f * x * x;
- float b = exp((x - HLG_C) / HLG_A) + HLG_B;
- return x < 0.5f ? a : b;
-}
-
-// delinearizer for HLG
-float oetf_hlg(float x) {
- float a = 0.5f * sqrt(x);
- float b = HLG_A * log(x - HLG_B) + HLG_C;
- return x <= 1.0f ? a : b;
-}
-
-float3 ootf_hlg(float3 c, float peak) {
- float luma = get_luma_src(c);
- float gamma = 1.2f + 0.42f * log10(peak * REFERENCE_WHITE / 1000.0f);
- gamma = max(1.0f, gamma);
- float factor = peak * powr(luma, gamma - 1.0f) / powr(12.0f, gamma);
- return c * factor;
-}
-
-float3 inverse_ootf_hlg(float3 c, float peak) {
- float gamma = 1.2f + 0.42f * log10(peak * REFERENCE_WHITE / 1000.0f);
- c *= powr(12.0f, gamma) / peak;
- c /= powr(get_luma_dst(c), (gamma - 1.0f) / gamma);
- return c;
+ if (x > 0.0f) {
+ float xpow = powr(x, 1.0f / ST2084_M2);
+ float num = max(xpow - ST2084_C1, 0.0f);
+ float den = max(ST2084_C2 - ST2084_C3 * xpow, FLOAT_EPS);
+ x = powr(num / den, 1.0f / ST2084_M1);
+ return x * ST2084_MAX_LUMINANCE / REFERENCE_WHITE;
+ } else {
+ return 0.0f;
+ }
+}
+
+// delinearizer for PQ/ST2084
+float inverse_eotf_st2084(float x) {
+ if (x > 0.0f) {
+ x *= REFERENCE_WHITE / ST2084_MAX_LUMINANCE;
+ float xpow = powr(x, ST2084_M1);
+#if 0
+ // Original formulation from SMPTE ST 2084:2014 publication.
+ float num = ST2084_C1 + ST2084_C2 * xpow;
+ float den = 1.0f + ST2084_C3 * xpow;
+ return powr(num / den, ST2084_M2);
+#else
+ // More stable arrangement that avoids some cancellation error.
+ float num = (ST2084_C1 - 1.0f) + (ST2084_C2 - ST2084_C3) * xpow;
+ float den = 1.0f + ST2084_C3 * xpow;
+ return powr(1.0f + num / den, ST2084_M2);
+#endif
+ } else {
+ return 0.0f;
+ }
+}
+
+float ootf_1_2(float x) {
+ return x < 0.0f ? x : powr(x, 1.2f);
+}
+
+float inverse_ootf_1_2(float x) {
+ return x < 0.0f ? x : powr(x, 1.0f / 1.2f);
+}
+
+float oetf_arib_b67(float x) {
+ x = max(x, 0.0f);
+ return x <= (1.0f / 12.0f)
+ ? sqrt(3.0f * x)
+ : (ARIB_B67_A * log(12.0f * x - ARIB_B67_B) + ARIB_B67_C);
+}
+
+float inverse_oetf_arib_b67(float x) {
+ x = max(x, 0.0f);
+ return x <= 0.5f
+ ? (x * x) * (1.0f / 3.0f)
+ : (exp((x - ARIB_B67_C) / ARIB_B67_A) + ARIB_B67_B) * (1.0f / 12.0f);
}
-float inverse_eotf_bt1886(float c) {
- return c < 0.0f ? 0.0f : powr(c, 1.0f / 2.4f);
+// linearizer for HLG/ARIB-B67
+float eotf_arib_b67(float x) {
+ return ootf_1_2(inverse_oetf_arib_b67(x));
}
-float oetf_bt709(float c) {
- c = c < 0.0f ? 0.0f : c;
- float r1 = 4.5f * c;
- float r2 = 1.099f * powr(c, 0.45f) - 0.099f;
- return c < 0.018f ? r1 : r2;
-}
-float inverse_oetf_bt709(float c) {
- float r1 = c / 4.5f;
- float r2 = powr((c + 0.099f) / 1.099f, 1.0f / 0.45f);
- return c < 0.081f ? r1 : r2;
+// delinearizer for HLG/ARIB-B67
+float inverse_eotf_arib_b67(float x) {
+ return oetf_arib_b67(inverse_ootf_1_2(x));
+}
+
+float inverse_eotf_bt1886(float x) {
+ return x < 0.0f ? 0.0f : powr(x, 1.0f / 2.4f);
+}
+
+float oetf_bt709(float x) {
+ x = max(0.0f, x);
+ return x < BT709_BETA
+ ? (x * 4.5f)
+ : (BT709_ALPHA * powr(x, 0.45f) - (BT709_ALPHA - 1.0f));
+}
+
+float inverse_oetf_bt709(float x) {
+ return x < (4.5f * BT709_BETA)
+ ? (x / 4.5f)
+ : (powr((x + (BT709_ALPHA - 1.0f)) / BT709_ALPHA, 1.0f / 0.45f));
}
float3 yuv2rgb(float y, float u, float v) {
@@ -187,19 +222,3 @@ float3 lrgb2lrgb(float3 c) {
return (float3)(rr, gg, bb);
#endif
}
-
-float3 ootf(float3 c, float peak) {
-#ifdef ootf_impl
- return ootf_impl(c, peak);
-#else
- return c;
-#endif
-}
-
-float3 inverse_ootf(float3 c, float peak) {
-#ifdef inverse_ootf_impl
- return inverse_ootf_impl(c, peak);
-#else
- return c;
-#endif
-}
Index: jellyfin-ffmpeg/libavfilter/opencl/tonemap.cl
===================================================================
--- jellyfin-ffmpeg.orig/libavfilter/opencl/tonemap.cl
+++ jellyfin-ffmpeg/libavfilter/opencl/tonemap.cl
@@ -16,54 +16,50 @@
* Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA
*/
-#define REFERENCE_WHITE 100.0f
+#define FLOAT_EPS 1.175494351e-38f
+
extern float3 lrgb2yuv(float3);
extern float lrgb2y(float3);
extern float3 yuv2lrgb(float3);
extern float3 lrgb2lrgb(float3);
extern float get_luma_src(float3);
extern float get_luma_dst(float3);
-extern float3 ootf(float3 c, float peak);
-extern float3 inverse_ootf(float3 c, float peak);
+extern float eotf_st2084(float);
+extern float inverse_eotf_st2084(float);
extern float3 get_chroma_sample(float3, float3, float3, float3);
-struct detection_result {
- float peak;
- float average;
-};
-
float hable_f(float in) {
float a = 0.15f, b = 0.50f, c = 0.10f, d = 0.20f, e = 0.02f, f = 0.30f;
return (in * (in * a + b * c) + d * e) / (in * (in * a + b) + d * f) - e / f;
}
-float direct(float s, float peak) {
+float direct(float s, float peak, float target_peak) {
return s;
}
-float linear(float s, float peak) {
+float linear(float s, float peak, float target_peak) {
return s * tone_param / peak;
}
-float gamma(float s, float peak) {
- float p = s > 0.05f ? s /peak : 0.05f / peak;
+float gamma(float s, float peak, float target_peak) {
+ float p = s > 0.05f ? s / peak : 0.05f / peak;
float v = powr(p, 1.0f / tone_param);
- return s > 0.05f ? v : (s * v /0.05f);
+ return s > 0.05f ? v : (s * v / 0.05f);
}
-float clip(float s, float peak) {
+float clip(float s, float peak, float target_peak) {
return clamp(s * tone_param, 0.0f, 1.0f);
}
-float reinhard(float s, float peak) {
+float reinhard(float s, float peak, float target_peak) {
return s / (s + tone_param) * (peak + tone_param) / peak;
}
-float hable(float s, float peak) {
- return hable_f(s)/hable_f(peak);
+float hable(float s, float peak, float target_peak) {
+ return hable_f(s) / hable_f(peak);
}
-float mobius(float s, float peak) {
+float mobius(float s, float peak, float target_peak) {
float j = tone_param;
float a, b;
@@ -71,102 +67,32 @@ float mobius(float s, float peak) {
return s;
a = -j * j * (peak - 1.0f) / (j * j - 2.0f * j + peak);
- b = (j * j - 2.0f * j * peak + peak) / max(peak - 1.0f, 1e-6f);
+ b = (j * j - 2.0f * j * peak + peak) / max(peak - 1.0f, FLOAT_EPS);
return (b * b + 2.0f * b * j + j * j) / (b - a) * (s + a) / (s + b);
}
-// detect peak/average signal of a frame, the algorithm was ported from:
-// libplacebo (https://github.com/haasn/libplacebo)
-struct detection_result
-detect_peak_avg(global uint *util_buf, __local uint *sum_wg,
- float signal, float peak) {
-// layout of the util buffer
-//
-// Name: : Size (units of 4-bytes)
-// average buffer : detection_frames + 1
-// peak buffer : detection_frames + 1
-// workgroup counter : 1
-// total of peak : 1
-// total of average : 1
-// frame index : 1
-// frame number : 1
- global uint *avg_buf = util_buf;
- global uint *peak_buf = avg_buf + DETECTION_FRAMES + 1;
- global uint *counter_wg_p = peak_buf + DETECTION_FRAMES + 1;
- global uint *max_total_p = counter_wg_p + 1;
- global uint *avg_total_p = max_total_p + 1;
- global uint *frame_idx_p = avg_total_p + 1;
- global uint *scene_frame_num_p = frame_idx_p + 1;
-
- uint frame_idx = *frame_idx_p;
- uint scene_frame_num = *scene_frame_num_p;
-
- size_t lidx = get_local_id(0);
- size_t lidy = get_local_id(1);
- size_t lsizex = get_local_size(0);
- size_t lsizey = get_local_size(1);
- uint num_wg = get_num_groups(0) * get_num_groups(1);
- size_t group_idx = get_group_id(0);
- size_t group_idy = get_group_id(1);
- struct detection_result r = {peak, sdr_avg};
- if (lidx == 0 && lidy == 0)
- *sum_wg = 0;
- barrier(CLK_LOCAL_MEM_FENCE);
-
- // update workgroup sum
- atomic_add(sum_wg, (uint)(signal * REFERENCE_WHITE));
- barrier(CLK_LOCAL_MEM_FENCE);
-
- // update frame peak/avg using work-group-average.
- if (lidx == 0 && lidy == 0) {
- uint avg_wg = *sum_wg / (lsizex * lsizey);
- atomic_max(&peak_buf[frame_idx], avg_wg);
- atomic_add(&avg_buf[frame_idx], avg_wg);
- }
-
- if (scene_frame_num > 0) {
- float peak = (float)*max_total_p / (REFERENCE_WHITE * scene_frame_num);
- float avg = (float)*avg_total_p / (REFERENCE_WHITE * scene_frame_num);
- r.peak = max(1.0f, peak);
- r.average = max(0.25f, avg);
- }
+float bt2390(float s, float peak, float target_peak) {
+ float peak_pq = inverse_eotf_st2084(peak);
+ float scale = 1.0f / peak_pq;
+
+ float s_pq = inverse_eotf_st2084(s) * scale;
+ float maxLum = inverse_eotf_st2084(target_peak) * scale;
+
+ float ks = 1.5f * maxLum - 0.5f;
+ float tb = (s_pq - ks) / (1.0f - ks);
+ float tb2 = tb * tb;
+ float tb3 = tb2 * tb;
+ float pb = (2.0f * tb3 - 3.0f * tb2 + 1.0f) * ks +
+ (tb3 - 2.0f * tb2 + tb) * (1.0f - ks) +
+ (-2.0f * tb3 + 3.0f * tb2) * maxLum;
+ float sig = (s_pq < ks) ? s_pq : pb;
- if (lidx == 0 && lidy == 0 && atomic_add(counter_wg_p, 1) == num_wg - 1) {
- *counter_wg_p = 0;
- avg_buf[frame_idx] /= num_wg;
-
- if (scene_threshold > 0.0f) {
- uint cur_max = peak_buf[frame_idx];
- uint cur_avg = avg_buf[frame_idx];
- int diff = (int)(scene_frame_num * cur_avg) - (int)*avg_total_p;
-
- if (abs(diff) > scene_frame_num * scene_threshold * REFERENCE_WHITE) {
- for (uint i = 0; i < DETECTION_FRAMES + 1; i++)
- avg_buf[i] = 0;
- for (uint i = 0; i < DETECTION_FRAMES + 1; i++)
- peak_buf[i] = 0;
- *avg_total_p = *max_total_p = 0;
- *scene_frame_num_p = 0;
- avg_buf[frame_idx] = cur_avg;
- peak_buf[frame_idx] = cur_max;
- }
- }
- uint next = (frame_idx + 1) % (DETECTION_FRAMES + 1);
- // add current frame, subtract next frame
- *max_total_p += peak_buf[frame_idx] - peak_buf[next];
- *avg_total_p += avg_buf[frame_idx] - avg_buf[next];
- // reset next frame
- peak_buf[next] = avg_buf[next] = 0;
- *frame_idx_p = next;
- *scene_frame_num_p = min(*scene_frame_num_p + 1,
- (uint)DETECTION_FRAMES);
- }
- return r;
+ 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);
+float3 map_one_pixel_rgb(float3 rgb, float peak) {
+ float sig = max(max(rgb.x, max(rgb.y, rgb.z)), FLOAT_EPS);
// 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
@@ -178,30 +104,24 @@ float3 map_one_pixel_rgb(float3 rgb, flo
float sig_old = sig;
- // Scale the signal to compensate for differences in the average brightness
- float slope = min(1.0f, sdr_avg / average);
- sig *= slope;
- peak *= slope;
-
// Desaturate the color using a coefficient dependent on the signal level
if (desat_param > 0.0f) {
float luma = get_luma_dst(rgb);
- float coeff = max(sig - 0.18f, 1e-6f) / max(sig, 1e-6f);
- coeff = native_powr(coeff, 10.0f / desat_param);
+ float coeff = max(sig - 0.18f, FLOAT_EPS) / max(sig, FLOAT_EPS);
+ coeff = powr(coeff, 10.0f / desat_param);
rgb = mix(rgb, (float3)luma, (float3)coeff);
- sig = mix(sig, luma * slope, coeff);
}
- sig = TONE_FUNC(sig, peak);
+ sig = TONE_FUNC(sig, peak, target_peak);
sig = min(sig, 1.0f);
- rgb *= (sig/sig_old);
+ rgb *= (sig / sig_old);
return rgb;
}
-// map from source space YUV to destination space 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;
}
@@ -210,7 +130,6 @@ __kernel void tonemap(__write_only image
__read_only image2d_t src1,
__write_only image2d_t dst2,
__read_only image2d_t src2,
- global uint *util_buf,
float peak
)
{
@@ -241,23 +160,17 @@ __kernel void tonemap(__write_only image
float sig3 = max(c3.x, max(c3.y, c3.z));
float sig = max(sig0, max(sig1, max(sig2, sig3)));
- struct detection_result r = detect_peak_avg(util_buf, &sum_wg, sig, peak);
-
float3 c0_old = c0, c1_old = c1, c2_old = c2;
- c0 = map_one_pixel_rgb(c0, r.peak, r.average);
- c1 = map_one_pixel_rgb(c1, r.peak, r.average);
- c2 = map_one_pixel_rgb(c2, r.peak, r.average);
- c3 = map_one_pixel_rgb(c3, r.peak, r.average);
-
- c0 = inverse_ootf(c0, target_peak);
- c1 = inverse_ootf(c1, target_peak);
- c2 = inverse_ootf(c2, target_peak);
- c3 = inverse_ootf(c3, target_peak);
+ c0 = map_one_pixel_rgb(c0, peak);
+ c1 = map_one_pixel_rgb(c1, peak);
+ c2 = map_one_pixel_rgb(c2, peak);
+ c3 = map_one_pixel_rgb(c3, peak);
y0 = lrgb2y(c0);
y1 = lrgb2y(c1);
y2 = lrgb2y(c2);
y3 = lrgb2y(c3);
+
float3 chroma_c = get_chroma_sample(c0, c1, c2, c3);
float3 chroma = lrgb2yuv(chroma_c);
Index: jellyfin-ffmpeg/libavfilter/vf_tonemap_opencl.c
===================================================================
--- jellyfin-ffmpeg.orig/libavfilter/vf_tonemap_opencl.c
+++ jellyfin-ffmpeg/libavfilter/vf_tonemap_opencl.c
@@ -15,6 +15,7 @@
* License along with FFmpeg; if not, write to the Free Software
* Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA
*/
+
#include <float.h>
#include "libavutil/avassert.h"
@@ -31,13 +32,6 @@
#include "video.h"
#include "colorspace.h"
-// TODO:
-// - separate peak-detection from tone-mapping kernel to solve
-// one-frame-delay issue.
-// - more format support
-
-#define DETECTION_FRAMES 63
-
enum TonemapAlgorithm {
TONEMAP_NONE,
TONEMAP_LINEAR,
@@ -46,6 +40,7 @@ enum TonemapAlgorithm {
TONEMAP_REINHARD,
TONEMAP_HABLE,
TONEMAP_MOBIUS,
+ TONEMAP_BT2390,
TONEMAP_MAX,
};
@@ -68,12 +63,11 @@ typedef struct TonemapOpenCLContext {
int initialised;
cl_kernel kernel;
cl_command_queue command_queue;
- cl_mem util_mem;
} TonemapOpenCLContext;
static const char *const linearize_funcs[AVCOL_TRC_NB] = {
- [AVCOL_TRC_SMPTE2084] = "eotf_st2084",
- [AVCOL_TRC_ARIB_STD_B67] = "inverse_oetf_hlg",
+ [AVCOL_TRC_SMPTE2084] = "eotf_st2084",
+ [AVCOL_TRC_ARIB_STD_B67] = "eotf_arib_b67",
};
static const char *const delinearize_funcs[AVCOL_TRC_NB] = {
@@ -99,6 +93,7 @@ static const char *const tonemap_func[TO
[TONEMAP_REINHARD] = "reinhard",
[TONEMAP_HABLE] = "hable",
[TONEMAP_MOBIUS] = "mobius",
+ [TONEMAP_BT2390] = "bt2390",
};
static void get_rgb2rgb_matrix(enum AVColorPrimaries in, enum AVColorPrimaries out,
@@ -112,9 +107,6 @@ static void get_rgb2rgb_matrix(enum AVCo
}
#define OPENCL_SOURCE_NB 3
-// Average light level for SDR signals. This is equal to a signal level of 0.5
-// under a typical presentation gamma of about 2.0.
-static const float sdr_avg = 0.25f;
static int tonemap_opencl_init(AVFilterContext *avctx)
{
@@ -127,7 +119,7 @@ static int tonemap_opencl_init(AVFilterC
AVBPrint header;
const char *opencl_sources[OPENCL_SOURCE_NB];
- av_bprint_init(&header, 1024, AV_BPRINT_SIZE_AUTOMATIC);
+ av_bprint_init(&header, 2048, AV_BPRINT_SIZE_UNLIMITED);
switch(ctx->tonemap) {
case TONEMAP_GAMMA:
@@ -149,18 +141,20 @@ static int tonemap_opencl_init(AVFilterC
// SDR peak is 1.0f
ctx->target_peak = 1.0f;
- av_log(ctx, AV_LOG_DEBUG, "tone mapping transfer from %s to %s\n",
+
+ av_log(ctx, AV_LOG_DEBUG, "Tonemapping transfer from %s to %s\n",
av_color_transfer_name(ctx->trc_in),
av_color_transfer_name(ctx->trc_out));
- av_log(ctx, AV_LOG_DEBUG, "mapping colorspace from %s to %s\n",
+ av_log(ctx, AV_LOG_DEBUG, "Mapping colorspace from %s to %s\n",
av_color_space_name(ctx->colorspace_in),
av_color_space_name(ctx->colorspace_out));
- av_log(ctx, AV_LOG_DEBUG, "mapping primaries from %s to %s\n",
+ av_log(ctx, AV_LOG_DEBUG, "Mapping primaries from %s to %s\n",
av_color_primaries_name(ctx->primaries_in),
av_color_primaries_name(ctx->primaries_out));
- av_log(ctx, AV_LOG_DEBUG, "mapping range from %s to %s\n",
+ av_log(ctx, AV_LOG_DEBUG, "Mapping range from %s to %s\n",
av_color_range_name(ctx->range_in),
av_color_range_name(ctx->range_out));
+
// checking valid value just because of limited implementaion
// please remove when more functionalities are implemented
av_assert0(ctx->trc_out == AVCOL_TRC_BT709 ||
@@ -178,11 +172,9 @@ static int tonemap_opencl_init(AVFilterC
ctx->desat_param);
av_bprintf(&header, "__constant const float target_peak = %.4ff;\n",
ctx->target_peak);
- av_bprintf(&header, "__constant const float sdr_avg = %.4ff;\n", sdr_avg);
av_bprintf(&header, "__constant const float scene_threshold = %.4ff;\n",
ctx->scene_threshold);
av_bprintf(&header, "#define TONE_FUNC %s\n", tonemap_func[ctx->tonemap]);
- av_bprintf(&header, "#define DETECTION_FRAMES %d\n", DETECTION_FRAMES);
if (ctx->primaries_out != ctx->primaries_in) {
get_rgb2rgb_matrix(ctx->primaries_in, ctx->primaries_out, rgb2rgb);
@@ -196,6 +188,16 @@ static int tonemap_opencl_init(AVFilterC
av_bprintf(&header, "#define chroma_loc %d\n", (int)ctx->chroma_loc);
+ av_bprintf(&header, "#define powr native_powr\n");
+
+ av_bprintf(&header, "#define exp native_exp\n");
+
+ av_bprintf(&header, "#define log native_log\n");
+
+ av_bprintf(&header, "#define log10 native_log10\n");
+
+ av_bprintf(&header, "#define sqrt native_sqrt\n");
+
if (rgb2rgb_passthrough)
av_bprintf(&header, "#define RGB2RGB_PASSTHROUGH\n");
else
@@ -205,7 +207,7 @@ static int tonemap_opencl_init(AVFilterC
luma_src = ff_get_luma_coefficients(ctx->colorspace_in);
if (!luma_src) {
err = AVERROR(EINVAL);
- av_log(avctx, AV_LOG_ERROR, "unsupported input colorspace %d (%s)\n",
+ av_log(avctx, AV_LOG_ERROR, "Unsupported input colorspace %d (%s)\n",
ctx->colorspace_in, av_color_space_name(ctx->colorspace_in));
goto fail;
}
@@ -213,7 +215,7 @@ static int tonemap_opencl_init(AVFilterC
luma_dst = ff_get_luma_coefficients(ctx->colorspace_out);
if (!luma_dst) {
err = AVERROR(EINVAL);
- av_log(avctx, AV_LOG_ERROR, "unsupported output colorspace %d (%s)\n",
+ av_log(avctx, AV_LOG_ERROR, "Unsupported output colorspace %d (%s)\n",
ctx->colorspace_out, av_color_space_name(ctx->colorspace_out));
goto fail;
}
@@ -225,21 +227,16 @@ static int tonemap_opencl_init(AVFilterC
ff_matrix_invert_3x3(rgb2yuv, yuv2rgb);
ff_opencl_print_const_matrix_3x3(&header, "rgb_matrix", yuv2rgb);
- av_bprintf(&header, "constant float3 luma_src = {%.4ff, %.4ff, %.4ff};\n",
+ av_bprintf(&header, "__constant float3 luma_src = {%.4ff, %.4ff, %.4ff};\n",
luma_src->cr, luma_src->cg, luma_src->cb);
- av_bprintf(&header, "constant float3 luma_dst = {%.4ff, %.4ff, %.4ff};\n",
+ av_bprintf(&header, "__constant float3 luma_dst = {%.4ff, %.4ff, %.4ff};\n",
luma_dst->cr, luma_dst->cg, luma_dst->cb);
- av_bprintf(&header, "#define linearize %s\n", linearize_funcs[ctx->trc_in]);
+ av_bprintf(&header, "#define linearize %s\n",
+ linearize_funcs[ctx->trc_in]);
av_bprintf(&header, "#define delinearize %s\n",
delinearize_funcs[ctx->trc_out]);
- if (ctx->trc_in == AVCOL_TRC_ARIB_STD_B67)
- av_bprintf(&header, "#define ootf_impl ootf_hlg\n");
-
- if (ctx->trc_out == AVCOL_TRC_ARIB_STD_B67)
- av_bprintf(&header, "#define inverse_ootf_impl inverse_ootf_hlg\n");
-
av_log(avctx, AV_LOG_DEBUG, "Generated OpenCL header:\n%s\n", header.str);
opencl_sources[0] = header.str;
opencl_sources[1] = ff_opencl_source_tonemap;
@@ -259,19 +256,11 @@ static int tonemap_opencl_init(AVFilterC
ctx->kernel = clCreateKernel(ctx->ocf.program, "tonemap", &cle);
CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create kernel %d.\n", cle);
- ctx->util_mem =
- clCreateBuffer(ctx->ocf.hwctx->context, 0,
- (2 * DETECTION_FRAMES + 7) * sizeof(unsigned),
- NULL, &cle);
- CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create util buffer: %d.\n", cle);
-
ctx->initialised = 1;
return 0;
fail:
av_bprint_finalize(&header, NULL);
- if (ctx->util_mem)
- clReleaseMemObject(ctx->util_mem);
if (ctx->command_queue)
clReleaseCommandQueue(ctx->command_queue);
if (ctx->kernel)
@@ -285,11 +274,11 @@ static int tonemap_opencl_config_output(
TonemapOpenCLContext *s = avctx->priv;
int ret;
if (s->format == AV_PIX_FMT_NONE)
- av_log(avctx, AV_LOG_WARNING, "format not set, use default format NV12\n");
+ av_log(avctx, AV_LOG_WARNING, "Format not set, use default format NV12\n");
else {
if (s->format != AV_PIX_FMT_P010 &&
s->format != AV_PIX_FMT_NV12) {
- av_log(avctx, AV_LOG_ERROR, "unsupported output format,"
+ av_log(avctx, AV_LOG_ERROR, "Unsupported output format,"
"only p010/nv12 supported now\n");
return AVERROR(EINVAL);
}
@@ -315,8 +304,7 @@ static int launch_kernel(AVFilterContext
CL_SET_KERNEL_ARG(kernel, 1, cl_mem, &input->data[0]);
CL_SET_KERNEL_ARG(kernel, 2, cl_mem, &output->data[1]);
CL_SET_KERNEL_ARG(kernel, 3, cl_mem, &input->data[1]);
- CL_SET_KERNEL_ARG(kernel, 4, cl_mem, &ctx->util_mem);
- CL_SET_KERNEL_ARG(kernel, 5, cl_float, &peak);
+ CL_SET_KERNEL_ARG(kernel, 4, cl_float, &peak);
local_work[0] = 16;
local_work[1] = 16;
@@ -390,13 +378,15 @@ static int tonemap_opencl_filter_frame(A
if (!ctx->initialised) {
if (!(input->color_trc == AVCOL_TRC_SMPTE2084 ||
input->color_trc == AVCOL_TRC_ARIB_STD_B67)) {
- av_log(ctx, AV_LOG_ERROR, "unsupported transfer function characteristic.\n");
+ av_log(ctx, AV_LOG_ERROR, "Unsupported transfer function characteristic: %s\n",
+ av_color_transfer_name(input->color_trc));
err = AVERROR(ENOSYS);
goto fail;
}
if (input_frames_ctx->sw_format != AV_PIX_FMT_P010) {
- av_log(ctx, AV_LOG_ERROR, "unsupported format in tonemap_opencl.\n");
+ av_log(ctx, AV_LOG_ERROR, "Unsupported input format: %s\n",
+ av_get_pix_fmt_name(input_frames_ctx->sw_format));
err = AVERROR(ENOSYS);
goto fail;
}
@@ -423,31 +413,9 @@ static int tonemap_opencl_filter_frame(A
ff_update_hdr_metadata(output, ctx->target_peak);
- av_log(ctx, AV_LOG_DEBUG, "Tone-mapping output: %s, %ux%u (%"PRId64").\n",
+ av_log(ctx, AV_LOG_DEBUG, "Tonemapping output: %s, %ux%u (%"PRId64").\n",
av_get_pix_fmt_name(output->format),
output->width, output->height, output->pts);
-#ifndef NDEBUG
- {
- uint32_t *ptr, *max_total_p, *avg_total_p, *frame_number_p;
- float peak_detected, avg_detected;
- unsigned map_size = (2 * DETECTION_FRAMES + 7) * sizeof(unsigned);
- ptr = (void *)clEnqueueMapBuffer(ctx->command_queue, ctx->util_mem,
- CL_TRUE, CL_MAP_READ, 0, map_size,
- 0, NULL, NULL, &cle);
- // For the layout of the util buffer, refer tonemap.cl
- if (ptr) {
- max_total_p = ptr + 2 * (DETECTION_FRAMES + 1) + 1;
- avg_total_p = max_total_p + 1;
- frame_number_p = avg_total_p + 2;
- peak_detected = (float)*max_total_p / (REFERENCE_WHITE * (*frame_number_p));
- avg_detected = (float)*avg_total_p / (REFERENCE_WHITE * (*frame_number_p));
- av_log(ctx, AV_LOG_DEBUG, "peak %f, avg %f will be used for next frame\n",
- peak_detected, avg_detected);
- clEnqueueUnmapMemObject(ctx->command_queue, ctx->util_mem, ptr, 0,
- NULL, NULL);
- }
- }
-#endif
return ff_filter_frame(outlink, output);
@@ -463,8 +431,6 @@ static av_cold void tonemap_opencl_unini
TonemapOpenCLContext *ctx = avctx->priv;
cl_int cle;
- if (ctx->util_mem)
- clReleaseMemObject(ctx->util_mem);
if (ctx->kernel) {
cle = clReleaseKernel(ctx->kernel);
if (cle != CL_SUCCESS)
@@ -493,6 +459,7 @@ static const AVOption tonemap_opencl_opt
{ "reinhard", 0, 0, AV_OPT_TYPE_CONST, {.i64 = TONEMAP_REINHARD}, 0, 0, FLAGS, "tonemap" },
{ "hable", 0, 0, AV_OPT_TYPE_CONST, {.i64 = TONEMAP_HABLE}, 0, 0, FLAGS, "tonemap" },
{ "mobius", 0, 0, AV_OPT_TYPE_CONST, {.i64 = TONEMAP_MOBIUS}, 0, 0, FLAGS, "tonemap" },
+ { "bt2390", 0, 0, AV_OPT_TYPE_CONST, {.i64 = TONEMAP_BT2390}, 0, 0, FLAGS, "tonemap" },
{ "transfer", "set transfer characteristic", OFFSET(trc), AV_OPT_TYPE_INT, {.i64 = AVCOL_TRC_BT709}, -1, INT_MAX, FLAGS, "transfer" },
{ "t", "set transfer characteristic", OFFSET(trc), AV_OPT_TYPE_INT, {.i64 = AVCOL_TRC_BT709}, -1, INT_MAX, FLAGS, "transfer" },
{ "bt709", 0, 0, AV_OPT_TYPE_CONST, {.i64 = AVCOL_TRC_BT709}, 0, 0, FLAGS, "transfer" },
debian/patches/0007-fix-for-fmp4-in-hlsenc.patch
Vendored
-24
View File
@@ -1,24 +0,0 @@
Index: jellyfin-ffmpeg/libavformat/hlsenc.c
===================================================================
--- jellyfin-ffmpeg.orig/libavformat/hlsenc.c
+++ jellyfin-ffmpeg/libavformat/hlsenc.c
@@ -2672,14 +2672,13 @@ static int hls_write_packet(AVFormatCont
vs->packets_written++;
if (oc->pb) {
- int64_t keyframe_pre_pos = avio_tell(oc->pb);
ret = ff_write_chained(oc, stream_index, pkt, s, 0);
- if ((st->codecpar->codec_type == AVMEDIA_TYPE_VIDEO) &&
- (pkt->flags & AV_PKT_FLAG_KEY) && !keyframe_pre_pos) {
- av_write_frame(oc, NULL); /* Flush any buffered data */
- vs->video_keyframe_size = avio_tell(oc->pb) - keyframe_pre_pos;
+ vs->video_keyframe_size += pkt->size;
+ if ((st->codecpar->codec_type == AVMEDIA_TYPE_VIDEO) && (pkt->flags & AV_PKT_FLAG_KEY)) {
+ vs->video_keyframe_size = avio_tell(oc->pb);
+ } else {
+ vs->video_keyframe_pos = avio_tell(vs->out);
}
- vs->video_keyframe_pos = vs->start_pos;
if (hls->ignore_io_errors)
ret = 0;
}
debian/patches/0008-fix-nvdec-exceeded-32-surfaces-error.patch
Vendored
-17
View File
@@ -1,17 +0,0 @@
Index: jellyfin-ffmpeg/libavcodec/nvdec.c
===================================================================
--- jellyfin-ffmpeg.orig/libavcodec/nvdec.c
+++ jellyfin-ffmpeg/libavcodec/nvdec.c
@@ -303,8 +303,10 @@ static int nvdec_init_hwframes(AVCodecCo
frames_ctx = (AVHWFramesContext*)(*out_frames_ref)->data;
if (dummy) {
- // Copied from ff_decode_get_hw_frames_ctx for compatibility
- frames_ctx->initial_pool_size += 3;
+ // The function above guarantees 1 work surface, We must guarantee 4 work surfaces.
+ // (the absolute minimum), so add the missing count without exceeding the maximum
+ // recommended for nvdec.
+ frames_ctx->initial_pool_size = FFMIN(frames_ctx->initial_pool_size + 3, 32);
frames_ctx->free = nvdec_free_dummy;
frames_ctx->pool = av_buffer_pool_init(0, nvdec_alloc_dummy);
debian/patches/0009-fix-for-nvenc-from-upstream.patch
Vendored
-1716
View File
File diff suppressed because it is too large Load Diff
debian/patches/series
Vendored
-9
View File
@@ -1,9 +0,0 @@
0001_fix-segment-muxer.patch
0002-lavfi-add-a-filter-for-uploading-normal-frames-to-VAAPI.patch
0003-fix-for-the-broken-tonemap_vaapi-filter.patch
0004-cuda-format-converter-impl.patch
0005-cuda-tonemap-impl.patch
0006-bt2390-and-fix-for-peak-detection-in-opencl-tonemap.patch
0007-fix-for-fmp4-in-hlsenc.patch
0008-fix-nvdec-exceeded-32-surfaces-error.patch
0009-fix-for-nvenc-from-upstream.patch