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

Backport more fixes from new branch

Browse Source 
Signed-off-by: nyanmisaka <nst799610810@gmail.com>
This commit is contained in:
nyanmisaka
2026-03-22 19:46:27 +08:00
parent 283c706780
commit d2ffb0d9a7
4 changed files with 59 additions and 48 deletions
Download Patch File Download Diff File Expand all files Collapse all files
debian/patches/0004-add-cuda-tonemap-impl.patch
Vendored
+2 -2
View File
@@ -1121,8 +1121,8 @@ Index: FFmpeg/libavfilter/cuda/pixfmt.h
+
+static __inline__ __device__ float3 read_tex_px_flt(const FFCUDAFrame& frame, int x, int y)
+{
+ float ncoord_x = (float)(x + 0.5f) / (float)frame.width;
+ float ncoord_y = (float)(y + 0.5f) / (float)frame.height;
+ float ncoord_x = (float)(x + 1) / (float)(frame.width + 1);
+ float ncoord_y = (float)(y + 1) / (float)(frame.height + 1);
+
+ float px_y = tex2D<float>(frame.tex[0], x, y);
+ float2 px_uv = tex2D<float2>(frame.tex[1], ncoord_x, ncoord_y);
debian/patches/0007-add-bt2390-eetf-and-code-refactor-to-opencl-tonemap.patch
Vendored
+46 -36
View File
@@ -1368,10 +1368,10 @@ Index: FFmpeg/libavfilter/opencl/tonemap.cl
+ return;
+
+ float2 src1_sz_recip = native_recip(convert_float2(src1_sz));
+ float2 ncoords_yuv0 = (convert_float2((int2)(x, y)) + 0.5f) * src1_sz_recip;
+ float2 ncoords_yuv1 = (convert_float2((int2)(x + 1, y)) + 0.5f) * src1_sz_recip;
+ float2 ncoords_yuv2 = (convert_float2((int2)(x, y + 1)) + 0.5f) * src1_sz_recip;
+ float2 ncoords_yuv3 = (convert_float2((int2)(x + 1, y + 1)) + 0.5f) * src1_sz_recip;
+ float2 ncoords_yuv0 = convert_float2((int2)(x, y)) * src1_sz_recip;
+ float2 ncoords_yuv1 = convert_float2((int2)(x + 1, y)) * src1_sz_recip;
+ float2 ncoords_yuv2 = convert_float2((int2)(x, y + 1)) * src1_sz_recip;
+ float2 ncoords_yuv3 = convert_float2((int2)(x + 1, y + 1)) * src1_sz_recip;
+
+ float3 yuv0, yuv1, yuv2, yuv3;
+
@@ -1651,10 +1651,10 @@ Index: FFmpeg/libavfilter/opencl/tonemap.cl
+ return;
+
+ float2 src1_sz_recip = native_recip(convert_float2(src1_sz));
+ float2 ncoords_yuv0 = (convert_float2((int2)(x, y)) + 0.5f) * src1_sz_recip;
+ float2 ncoords_yuv1 = (convert_float2((int2)(x + 1, y)) + 0.5f) * src1_sz_recip;
+ float2 ncoords_yuv2 = (convert_float2((int2)(x, y + 1)) + 0.5f) * src1_sz_recip;
+ float2 ncoords_yuv3 = (convert_float2((int2)(x + 1, y + 1)) + 0.5f) * src1_sz_recip;
+ float2 ncoords_yuv0 = convert_float2((int2)(x, y)) * src1_sz_recip;
+ float2 ncoords_yuv1 = convert_float2((int2)(x + 1, y)) * src1_sz_recip;
+ float2 ncoords_yuv2 = convert_float2((int2)(x, y + 1)) * src1_sz_recip;
+ float2 ncoords_yuv3 = convert_float2((int2)(x + 1, y + 1)) * src1_sz_recip;
+
+ float3 yuv0, yuv1, yuv2, yuv3;
+
@@ -1832,15 +1832,16 @@ Index: FFmpeg/libavfilter/vf_tonemap_opencl.c
#include "video.h"
#include "colorspace.h"
+#include "dither_matrix.h"
+
+#define OPENCL_SOURCE_NB 3
+
+#define REF_WHITE_SCALE (REFERENCE_WHITE / REFERENCE_WHITE_ALT)
-// TODO:
-// - separate peak-detection from tone-mapping kernel to solve
-// one-frame-delay issue.
-// - more format support
+#define OPENCL_SOURCE_NB 3
-#define DETECTION_FRAMES 63
+#define REF_WHITE_SCALE (REFERENCE_WHITE / REFERENCE_WHITE_ALT)
+
+static const enum AVPixelFormat supported_formats[] = {
+ AV_PIX_FMT_YUV420P,
+ AV_PIX_FMT_YUV420P16,
@@ -1849,8 +1850,7 @@ Index: FFmpeg/libavfilter/vf_tonemap_opencl.c
+ AV_PIX_FMT_P010,
+ AV_PIX_FMT_P016,
+};
-#define DETECTION_FRAMES 63
+
+static const int colorspaces_out[] = {
+ AVCOL_SPC_UNSPECIFIED,
+ AVCOL_SPC_BT709,
@@ -2523,7 +2523,7 @@ Index: 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 +698,216 @@ static int tonemap_opencl_init(AVFilterC
@@ -254,46 +698,224 @@ static int tonemap_opencl_init(AVFilterC
CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create OpenCL "
"command queue %d.\n", cle);
@@ -2608,20 +2608,22 @@ Index: FFmpeg/libavfilter/vf_tonemap_opencl.c
av_bprint_finalize(&header, NULL);
- if (ctx->util_mem)
- clReleaseMemObject(ctx->util_mem);
+ if (ctx->dovi_buf)
+ clReleaseMemObject(ctx->dovi_buf);
if (ctx->command_queue)
clReleaseCommandQueue(ctx->command_queue);
- if (ctx->command_queue)
- clReleaseCommandQueue(ctx->command_queue);
+ if (event)
+ clReleaseEvent(event);
if (ctx->kernel)
clReleaseKernel(ctx->kernel);
+ if (ctx->lut_generation_kernel)
+ clReleaseKernel(ctx->lut_generation_kernel);
+ if (event)
+ clReleaseEvent(event);
+ if (ctx->dither_image)
+ clReleaseMemObject(ctx->dither_image);
+ if (ctx->lut_buffer)
+ clReleaseMemObject(ctx->lut_buffer);
+ if (ctx->dovi_buf)
+ clReleaseMemObject(ctx->dovi_buf);
+ if (ctx->command_queue)
+ clReleaseCommandQueue(ctx->command_queue);
return err;
}
@@ -2662,6 +2664,14 @@ Index: FFmpeg/libavfilter/vf_tonemap_opencl.c
+ "lut_generation_kernel: %d.\n", cle);
+ }
+
+ if (ctx->ocf.program) {
+ cle = clReleaseProgram(ctx->ocf.program);
+ if (cle != CL_SUCCESS)
+ av_log(avctx, AV_LOG_ERROR, "Failed to release "
+ "program: %d.\n", cle);
+ ctx->ocf.program = NULL;
+ }
+
+ if (ctx->dither_image) {
+ cle = clReleaseMemObject(ctx->dither_image);
+ if (cle != CL_SUCCESS)
@@ -2760,7 +2770,7 @@ Index: FFmpeg/libavfilter/vf_tonemap_opencl.c
ret = ff_opencl_filter_config_output(outlink);
if (ret < 0)
return ret;
@@ -308,13 +922,49 @@ static int launch_kernel(AVFilterContext
@@ -308,13 +930,49 @@ static int launch_kernel(AVFilterContext
size_t global_work[2];
size_t local_work[2];
cl_int cle;
@@ -2770,7 +2780,13 @@ Index: FFmpeg/libavfilter/vf_tonemap_opencl.c
+ err = AVERROR(EIO);
+ goto fail;
+ }
+
- CL_SET_KERNEL_ARG(kernel, 0, cl_mem, &output->data[0]);
- 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);
+ if (ctx->out_planes > 2 && !output->data[2]) {
+ err = AVERROR(EIO);
+ goto fail;
@@ -2805,18 +2821,12 @@ Index: FFmpeg/libavfilter/vf_tonemap_opencl.c
+ if (ctx->dovi_buf) {
+ CL_SET_KERNEL_ARG(kernel, idx_arg++, cl_mem, &ctx->dovi_buf);
+ }
- CL_SET_KERNEL_ARG(kernel, 0, cl_mem, &output->data[0]);
- 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, idx_arg++, cl_float, &peak);
local_work[0] = 16;
local_work[1] = 16;
@@ -338,12 +988,10 @@ static int tonemap_opencl_filter_frame(A
@@ -338,12 +996,10 @@ static int tonemap_opencl_filter_frame(A
AVFilterContext *avctx = inlink->dst;
AVFilterLink *outlink = avctx->outputs[0];
TonemapOpenCLContext *ctx = avctx->priv;
@@ -2830,7 +2840,7 @@ Index: 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),
@@ -351,7 +999,6 @@ static int tonemap_opencl_filter_frame(A
@@ -351,7 +1007,6 @@ static int tonemap_opencl_filter_frame(A
if (!input->hw_frames_ctx)
return AVERROR(EINVAL);
@@ -2838,7 +2848,7 @@ Index: FFmpeg/libavfilter/vf_tonemap_opencl.c
output = ff_get_video_buffer(outlink, outlink->w, outlink->h);
if (!output) {
@@ -363,17 +1010,65 @@ static int tonemap_opencl_filter_frame(A
@@ -363,17 +1018,65 @@ static int tonemap_opencl_filter_frame(A
if (err < 0)
goto fail;
@@ -2910,7 +2920,7 @@ Index: FFmpeg/libavfilter/vf_tonemap_opencl.c
ctx->trc_in = input->color_trc;
ctx->trc_out = output->color_trc;
@@ -385,72 +1080,50 @@ static int tonemap_opencl_filter_frame(A
@@ -385,72 +1088,50 @@ static int tonemap_opencl_filter_frame(A
ctx->range_out = output->color_range;
ctx->chroma_loc = output->chroma_location;
@@ -3006,7 +3016,7 @@ Index: FFmpeg/libavfilter/vf_tonemap_opencl.c
av_frame_free(&input);
av_frame_free(&output);
return err;
@@ -458,62 +1131,101 @@ fail:
@@ -458,62 +1139,101 @@ fail:
static av_cold void tonemap_opencl_uninit(AVFilterContext *avctx)
{
@@ -3155,7 +3165,7 @@ Index: FFmpeg/libavfilter/vf_tonemap_opencl.c
{ NULL }
};
@@ -541,11 +1253,12 @@ const AVFilter ff_vf_tonemap_opencl = {
@@ -541,11 +1261,12 @@ const AVFilter ff_vf_tonemap_opencl = {
.description = NULL_IF_CONFIG_SMALL("Perform HDR to SDR conversion with tonemapping."),
.priv_size = sizeof(TonemapOpenCLContext),
.priv_class = &tonemap_opencl_class,
debian/patches/0050-add-vf-tonemap-videotoolbox-filter.patch
Vendored
+4 -4
View File
@@ -857,10 +857,10 @@ Index: FFmpeg/libavfilter/metal/vf_tonemap_videotoolbox.metal
+ if (xi >= dst2_sz.x || yi >= dst2_sz.y)
+ return;
+
+ float2 ncoords_yuv0 = (float2(int2(x, y)) + 0.5f) / float2(src1_sz);
+ float2 ncoords_yuv1 = (float2(int2(x + 1, y)) + 0.5f) / float2(src1_sz);
+ float2 ncoords_yuv2 = (float2(int2(x, y + 1)) + 0.5f) / float2(src1_sz);
+ float2 ncoords_yuv3 = (float2(int2(x + 1, y + 1)) + 0.5f) / float2(src1_sz);
+ float2 ncoords_yuv0 = float2(int2(x, y)) / float2(src1_sz);
+ float2 ncoords_yuv1 = float2(int2(x + 1, y)) / float2(src1_sz);
+ float2 ncoords_yuv2 = float2(int2(x, y + 1)) / float2(src1_sz);
+ float2 ncoords_yuv3 = float2(int2(x + 1, y + 1)) / float2(src1_sz);
+
+ float3 yuv0, yuv1, yuv2, yuv3;
+
debian/patches/0054-add-cuda-transpose-filter-impl.patch
Vendored
+7 -6
View File
@@ -2,7 +2,7 @@ Index: FFmpeg/configure
===================================================================
--- FFmpeg.orig/configure
+++ FFmpeg/configure
@@ -3317,6 +3317,8 @@ thumbnail_cuda_filter_deps="ffnvcodec"
@@ -3321,6 +3321,8 @@ thumbnail_cuda_filter_deps="ffnvcodec"
thumbnail_cuda_filter_deps_any="cuda_nvcc cuda_llvm"
tonemap_cuda_filter_deps="ffnvcodec const_nan"
tonemap_cuda_filter_deps_any="cuda_nvcc cuda_llvm"
@@ -40,7 +40,7 @@ Index: FFmpeg/libavfilter/vf_transpose_cuda.c
===================================================================
--- /dev/null
+++ FFmpeg/libavfilter/vf_transpose_cuda.c
@@ -0,0 +1,481 @@
@@ -0,0 +1,482 @@
+/*
+ * Copyright (C) 2024 NyanMisaka
+ *
@@ -250,7 +250,7 @@ Index: FFmpeg/libavfilter/vf_transpose_cuda.c
+ TransposeCUDAContext *s = ctx->priv;
+ CUcontext dummy, cuda_ctx;
+ CudaFunctions *cu;
+ int ret;
+ int flip_wh, ret;
+
+ if ((inlink->w >= inlink->h && s->passthrough == TRANSPOSE_PT_TYPE_LANDSCAPE) ||
+ (inlink->w <= inlink->h && s->passthrough == TRANSPOSE_PT_TYPE_PORTRAIT)) {
@@ -275,16 +275,17 @@ Index: FFmpeg/libavfilter/vf_transpose_cuda.c
+ case TRANSPOSE_CLOCK_FLIP:
+ outlink->w = inlink->h;
+ outlink->h = inlink->w;
+ flip_wh = 1;
+ break;
+ default:
+ outlink->w = inlink->w;
+ outlink->h = inlink->h;
+ flip_wh = 0;
+ break;
+ }
+
+ if (inlink->sample_aspect_ratio.num)
+ outlink->sample_aspect_ratio = av_div_q((AVRational) { 1, 1 },
+ inlink->sample_aspect_ratio);
+ if (flip_wh && inlink->sample_aspect_ratio.num)
+ outlink->sample_aspect_ratio = av_inv_q(inlink->sample_aspect_ratio);
+ else
+ outlink->sample_aspect_ratio = inlink->sample_aspect_ratio;
+