From d859373cd77f70f1538033adbaa6698f230dee43 Mon Sep 17 00:00:00 2001 From: nyanmisaka Date: Sat, 25 Apr 2026 21:28:19 +0800 Subject: [PATCH 1/4] Backport some fixes from new branch Signed-off-by: nyanmisaka --- .../0002-update-cuda-func-header.patch | 3 +- ...-enhanced-cuda-pixfmt-converter-impl.patch | 69 +++++++++++++++++-- .../patches/0004-add-cuda-tonemap-impl.patch | 24 +++---- ...-and-code-refactor-to-opencl-tonemap.patch | 66 +++++++++--------- .../0054-add-cuda-transpose-filter-impl.patch | 41 +++++------ ...flip-feat-to-opencl-transpose-filter.patch | 36 +++++++++- 6 files changed, 164 insertions(+), 75 deletions(-) diff --git a/debian/patches/0002-update-cuda-func-header.patch b/debian/patches/0002-update-cuda-func-header.patch index 3d58a7b17..007e8a5e7 100644 --- a/debian/patches/0002-update-cuda-func-header.patch +++ b/debian/patches/0002-update-cuda-func-header.patch @@ -165,7 +165,7 @@ Index: FFmpeg/compat/cuda/cuda_runtime.h { float4 ret; asm("tex.2d.v4.f32.f32 {%0, %1, %2, %3}, [%4, {%5, %6}];" : -@@ -156,37 +147,41 @@ inline __device__ float4 tex2D(c +@@ -156,37 +147,42 @@ inline __device__ float4 tex2D(c } template<> @@ -216,6 +216,7 @@ Index: FFmpeg/compat/cuda/cuda_runtime.h +static __inline__ __device__ float fabs(float a) { return __builtin_fabsf(a); } +static __inline__ __device__ double fabs(double a) { return __builtin_fabs(a); } +static __inline__ __device__ float sqrtf(float a) { return __builtin_sqrtf(a); } ++static __inline__ __device__ float rintf(float a) { return __builtin_rintf(a); } + +static __inline__ __device__ float __saturatef(float a) { return __nvvm_saturate_f(a); } +static __inline__ __device__ float __sinf(float a) { return __nvvm_sin_approx_f(a); } diff --git a/debian/patches/0003-add-enhanced-cuda-pixfmt-converter-impl.patch b/debian/patches/0003-add-enhanced-cuda-pixfmt-converter-impl.patch index c8ea2c543..92b1394e0 100644 --- a/debian/patches/0003-add-enhanced-cuda-pixfmt-converter-impl.patch +++ b/debian/patches/0003-add-enhanced-cuda-pixfmt-converter-impl.patch @@ -1,3 +1,35 @@ +Index: FFmpeg/libavfilter/cuda/vector_helpers.cuh +=================================================================== +--- FFmpeg.orig/libavfilter/cuda/vector_helpers.cuh ++++ FFmpeg/libavfilter/cuda/vector_helpers.cuh +@@ -109,4 +109,27 @@ inline __device__ float4 lerp_scalar ++inline __device__ T saturate_rintf(T a, float factor) { ++ return rintf(__saturatef(a) * factor); ++} ++ ++template<> ++inline __device__ float2 saturate_rintf(float2 a, float factor) { ++ return make_float2( ++ saturate_rintf(a.x, factor), ++ saturate_rintf(a.y, factor) ++ ); ++} ++ ++template<> ++inline __device__ float4 saturate_rintf(float4 a, float factor) { ++ return make_float4( ++ saturate_rintf(a.x, factor), ++ saturate_rintf(a.y, factor), ++ saturate_rintf(a.z, factor), ++ saturate_rintf(a.w, factor) ++ ); ++} ++ + #endif Index: FFmpeg/libavfilter/dither_matrix.h =================================================================== --- /dev/null @@ -332,7 +364,7 @@ Index: FFmpeg/libavfilter/vf_scale_cuda.c CHECK_CU(cu->cuCtxPopCurrent(&dummy)); } -@@ -275,6 +290,68 @@ static av_cold int init_processing_chain +@@ -275,6 +290,69 @@ static av_cold int init_processing_chain return 0; } @@ -362,7 +394,8 @@ Index: FFmpeg/libavfilter/vf_scale_cuda.c + #define CU_TRSF_NORMALIZED_COORDINATES 2 +#endif + CUDA_TEXTURE_DESC tex_desc = { -+ .addressMode = { CU_TR_ADDRESS_MODE_WRAP }, ++ .addressMode = { CU_TR_ADDRESS_MODE_WRAP, ++ CU_TR_ADDRESS_MODE_WRAP }, + .filterMode = CU_TR_FILTER_MODE_POINT, + .flags = CU_TRSF_NORMALIZED_COORDINATES, + }; @@ -401,7 +434,7 @@ Index: FFmpeg/libavfilter/vf_scale_cuda.c static av_cold int cudascale_load_functions(AVFilterContext *ctx) { CUDAScaleContext *s = ctx->priv; -@@ -383,6 +460,11 @@ static av_cold int cudascale_config_prop +@@ -383,6 +461,11 @@ static av_cold int cudascale_config_prop s->hwctx = device_hwctx; s->cu_stream = s->hwctx->stream; @@ -413,7 +446,7 @@ Index: FFmpeg/libavfilter/vf_scale_cuda.c if (inlink->sample_aspect_ratio.num) { outlink->sample_aspect_ratio = av_mul_q((AVRational){outlink->h*inlink->w, outlink->w*inlink->h}, -@@ -418,11 +500,15 @@ static int call_resize_kernel(AVFilterCo +@@ -418,11 +501,15 @@ static int call_resize_kernel(AVFilterCo (CUdeviceptr)out_frame->data[2], (CUdeviceptr)out_frame->data[3] }; @@ -430,7 +463,7 @@ Index: FFmpeg/libavfilter/vf_scale_cuda.c }; return CHECK_CU(cu->cuLaunchKernel(func, -@@ -446,6 +532,7 @@ static int scalecuda_resize(AVFilterCont +@@ -446,6 +533,7 @@ static int scalecuda_resize(AVFilterCont for (i = 0; i < s->in_planes; i++) { CUDA_TEXTURE_DESC tex_desc = { @@ -615,7 +648,29 @@ Index: FFmpeg/libavfilter/vf_scale_cuda.cu const float pi = 3.141592654f; float4 res = make_float4( -@@ -1147,7 +1198,9 @@ __device__ static inline T Subsample_Bic +@@ -1129,12 +1180,15 @@ __device__ static inline T Subsample_Bic + #define PIX(x, y) tex2D(tex, (x), (y)) + + return from_floatN( +- apply_coeffs(coeffsY, +- apply_coeffs(coeffsX, PIX(px - 1, py - 1), PIX(px, py - 1), PIX(px + 1, py - 1), PIX(px + 2, py - 1)), +- apply_coeffs(coeffsX, PIX(px - 1, py ), PIX(px, py ), PIX(px + 1, py ), PIX(px + 2, py )), +- apply_coeffs(coeffsX, PIX(px - 1, py + 1), PIX(px, py + 1), PIX(px + 1, py + 1), PIX(px + 2, py + 1)), +- apply_coeffs(coeffsX, PIX(px - 1, py + 2), PIX(px, py + 2), PIX(px + 1, py + 2), PIX(px + 2, py + 2)) +- ) * factor ++ saturate_rintf( ++ apply_coeffs(coeffsY, ++ apply_coeffs(coeffsX, PIX(px - 1, py - 1), PIX(px, py - 1), PIX(px + 1, py - 1), PIX(px + 2, py - 1)), ++ apply_coeffs(coeffsX, PIX(px - 1, py ), PIX(px, py ), PIX(px + 1, py ), PIX(px + 2, py )), ++ apply_coeffs(coeffsX, PIX(px - 1, py + 1), PIX(px, py + 1), PIX(px + 1, py + 1), PIX(px + 2, py + 1)), ++ apply_coeffs(coeffsX, PIX(px - 1, py + 2), PIX(px, py + 2), PIX(px + 1, py + 2), PIX(px + 2, py + 2)) ++ ), ++ factor ++ ) + ); + + #undef PIX +@@ -1147,7 +1201,9 @@ __device__ static inline T Subsample_Bic cudaTextureObject_t src_tex_2, cudaTextureObject_t src_tex_3, \ T *dst_0, T *dst_1, T *dst_2, T *dst_3, \ int dst_width, int dst_height, int dst_pitch, \ @@ -626,7 +681,7 @@ Index: FFmpeg/libavfilter/vf_scale_cuda.cu #define SUBSAMPLE(Convert, T) \ cudaTextureObject_t src_tex[4] = \ -@@ -1159,7 +1212,9 @@ __device__ static inline T Subsample_Bic +@@ -1159,7 +1215,9 @@ __device__ static inline T Subsample_Bic Convert( \ src_tex, dst, xo, yo, \ dst_width, dst_height, dst_pitch, \ diff --git a/debian/patches/0004-add-cuda-tonemap-impl.patch b/debian/patches/0004-add-cuda-tonemap-impl.patch index 5866b8a35..997be64d2 100644 --- a/debian/patches/0004-add-cuda-tonemap-impl.patch +++ b/debian/patches/0004-add-cuda-tonemap-impl.patch @@ -743,7 +743,7 @@ Index: FFmpeg/libavfilter/cuda/host_util.c + #define CU_TRSF_NORMALIZED_COORDINATES 2 +#endif + CUDA_TEXTURE_DESC tex_desc = { -+ .addressMode = { CU_TR_ADDRESS_MODE_CLAMP }, ++ .addressMode = { CU_TR_ADDRESS_MODE_CLAMP, CU_TR_ADDRESS_MODE_CLAMP }, + .filterMode = i == 0 ? CU_TR_FILTER_MODE_POINT : CU_TR_FILTER_MODE_LINEAR, + .flags = i == 0 ? 0 : CU_TRSF_NORMALIZED_COORDINATES, + }; @@ -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 + 1) / (float)(frame.width + 1); -+ float ncoord_y = (float)(y + 1) / (float)(frame.height + 1); ++ float ncoord_x = (float)x / (float)frame.width; ++ float ncoord_y = (float)y / (float)frame.height; + + float px_y = tex2D(frame.tex[0], x, y); + float2 px_uv = tex2D(frame.tex[1], ncoord_x, ncoord_y); @@ -2264,7 +2264,8 @@ Index: FFmpeg/libavfilter/vf_tonemap_cuda.c + #define CU_TRSF_NORMALIZED_COORDINATES 2 +#endif + CUDA_TEXTURE_DESC tex_desc = { -+ .addressMode = { CU_TR_ADDRESS_MODE_WRAP }, ++ .addressMode = { CU_TR_ADDRESS_MODE_WRAP, ++ CU_TR_ADDRESS_MODE_WRAP }, + .filterMode = CU_TR_FILTER_MODE_POINT, + .flags = CU_TRSF_NORMALIZED_COORDINATES, + }; @@ -2559,6 +2560,9 @@ Index: FFmpeg/libavfilter/vf_tonemap_cuda.c + extern const unsigned char ff_tonemap_ptx_data[]; + extern const unsigned int ff_tonemap_ptx_len; + ++ if (s->tonemap_mode == TONEMAP_MODE_AUTO) ++ s->tonemap_mode = TONEMAP_MODE_ITP; ++ + switch(s->tonemap) { + case TONEMAP_GAMMA: + if (isnan(s->param)) @@ -2774,14 +2778,6 @@ Index: FFmpeg/libavfilter/vf_tonemap_cuda.c + av_log(ctx, AV_LOG_DEBUG, "Disabled dovi tradeoff on high perf GPU.\n"); + } + -+ if (s->tonemap_mode == TONEMAP_MODE_AUTO) { -+ if (s->tradeoff) { -+ s->tonemap_mode = TONEMAP_MODE_LUM; -+ } else { -+ s->tonemap_mode = TONEMAP_MODE_ITP; -+ } -+ } -+ + if (s->cu_module) { + ret = CHECK_CU(cu->cuModuleUnload(s->cu_module)); + if (ret < 0) @@ -2952,6 +2948,10 @@ Index: FFmpeg/libavfilter/vf_tonemap_cuda.c + BLOCKX, BLOCKY, 1, 0, s->hwctx->stream, (s->lut_buffer ? args2 : args), NULL)); + +fail: ++ for (int i = 0; i < src.planes; i++) ++ if (src.tex[i]) ++ CHECK_CU(cu->cuTexObjectDestroy(src.tex[i])); ++ + return ret; +} + diff --git a/debian/patches/0007-add-bt2390-eetf-and-code-refactor-to-opencl-tonemap.patch b/debian/patches/0007-add-bt2390-eetf-and-code-refactor-to-opencl-tonemap.patch index a7e337f97..020909f3f 100644 --- a/debian/patches/0007-add-bt2390-eetf-and-code-refactor-to-opencl-tonemap.patch +++ b/debian/patches/0007-add-bt2390-eetf-and-code-refactor-to-opencl-tonemap.patch @@ -40,7 +40,7 @@ Index: FFmpeg/libavfilter/opencl.c int ff_opencl_filter_load_program(AVFilterContext *avctx, const char **program_source_array, int nb_strings) -@@ -171,8 +195,42 @@ int ff_opencl_filter_load_program(AVFilt +@@ -171,8 +195,40 @@ int ff_opencl_filter_load_program(AVFilt return AVERROR(EIO); } @@ -50,8 +50,8 @@ Index: FFmpeg/libavfilter/opencl.c + char *device_vendor = check_opencl_device_str(ctx->hwctx->device_id, CL_DEVICE_VENDOR); + char *device_name = check_opencl_device_str(ctx->hwctx->device_id, CL_DEVICE_NAME); + -+ if (strstr(device_vendor, "ARM") && -+ (strstr(device_name, "Mali") || strstr(device_name, "Immortalis"))) { ++ if (device_vendor && strstr(device_vendor, "ARM") && ++ device_name && strstr(device_name, "Mali")) { + av_freep(&device_vendor); + av_freep(&device_name); + @@ -71,10 +71,8 @@ Index: FFmpeg/libavfilter/opencl.c + return AVERROR(EIO); + } + } -+ if (device_vendor) -+ av_freep(&device_vendor); -+ if (device_name) -+ av_freep(&device_name); ++ av_freep(&device_vendor); ++ av_freep(&device_name); + } +#endif + @@ -84,7 +82,7 @@ Index: FFmpeg/libavfilter/opencl.c if (cle != CL_SUCCESS) { av_log(avctx, AV_LOG_ERROR, "Failed to build program: %d.\n", cle); -@@ -333,7 +391,7 @@ void ff_opencl_print_const_matrix_3x3(AV +@@ -333,7 +389,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++) @@ -1956,7 +1954,7 @@ Index: FFmpeg/libavfilter/vf_tonemap_opencl.c }; static int get_rgb2rgb_matrix(enum AVColorPrimaries in, enum AVColorPrimaries out, -@@ -108,90 +173,458 @@ static int get_rgb2rgb_matrix(enum AVCol +@@ -108,90 +173,453 @@ static int get_rgb2rgb_matrix(enum AVCol return 0; } @@ -2173,9 +2171,8 @@ Index: FFmpeg/libavfilter/vf_tonemap_opencl.c + int i, j, err; - av_bprint_init(&header, 1024, AV_BPRINT_SIZE_AUTOMATIC); -+ if (ctx->tonemap_mode == TONEMAP_MODE_AUTO) { ++ if (ctx->tonemap_mode == TONEMAP_MODE_AUTO) + ctx->tonemap_mode = TONEMAP_MODE_ITP; -+ } switch(ctx->tonemap) { case TONEMAP_GAMMA: @@ -2257,10 +2254,8 @@ Index: FFmpeg/libavfilter/vf_tonemap_opencl.c + if (!strstr(device_vendor, "ARM") && + !strstr(device_name, "Mali")) + ctx->tradeoff = 0; -+ if (device_vendor) -+ av_freep(&device_vendor); -+ if (device_name) -+ av_freep(&device_name); ++ av_freep(&device_vendor); ++ av_freep(&device_name); + } else { + ctx->tradeoff = 0; + } @@ -2272,7 +2267,6 @@ Index: FFmpeg/libavfilter/vf_tonemap_opencl.c + // for low perf device, only do reshaping for pure dovi + if (ctx->tradeoff && ctx->dovi && !ctx->is_pure_dovi) { + av_freep(&ctx->dovi); -+ ctx->dovi = NULL; + ctx->apply_dovi = 0; + } + @@ -2307,8 +2301,7 @@ Index: FFmpeg/libavfilter/vf_tonemap_opencl.c + dovi_buf_flags |= CL_MEM_FORCE_HOST_MEMORY_INTEL; + } + -+ if (device_exts) -+ av_freep(&device_exts); ++ av_freep(&device_exts); + + av_log(ctx, AV_LOG_DEBUG, "Tonemapping transfer from %s to %s\n", av_color_transfer_name(ctx->trc_in), @@ -2444,7 +2437,7 @@ Index: FFmpeg/libavfilter/vf_tonemap_opencl.c av_bprintf(&header, "#define chroma_loc %d\n", (int)ctx->chroma_loc); if (rgb2rgb_passthrough) -@@ -199,19 +632,41 @@ static int tonemap_opencl_init(AVFilterC +@@ -199,19 +627,41 @@ static int tonemap_opencl_init(AVFilterC else ff_opencl_print_const_matrix_3x3(&header, "rgb2rgb", rgb2rgb); @@ -2493,7 +2486,7 @@ Index: FFmpeg/libavfilter/vf_tonemap_opencl.c ctx->colorspace_out, av_color_space_name(ctx->colorspace_out)); goto fail; } -@@ -219,24 +674,13 @@ static int tonemap_opencl_init(AVFilterC +@@ -219,24 +669,13 @@ static int tonemap_opencl_init(AVFilterC ff_fill_rgb2yuv_table(luma_dst, rgb2yuv); ff_opencl_print_const_matrix_3x3(&header, "yuv_matrix", rgb2yuv); @@ -2523,7 +2516,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,224 @@ static int tonemap_opencl_init(AVFilterC +@@ -254,46 +693,231 @@ static int tonemap_opencl_init(AVFilterC CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create OpenCL " "command queue %d.\n", cle); @@ -2562,6 +2555,10 @@ Index: FFmpeg/libavfilter/vf_tonemap_opencl.c + + cle = clWaitForEvents(1, &event); + CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to wait for event completion: %d.\n", cle); ++ if (event) { ++ clReleaseEvent(event); ++ event = NULL; ++ } + } + + if (ctx->tradeoff) { @@ -2610,6 +2607,9 @@ Index: FFmpeg/libavfilter/vf_tonemap_opencl.c - clReleaseMemObject(ctx->util_mem); - if (ctx->command_queue) - clReleaseCommandQueue(ctx->command_queue); ++ av_freep(&device_vendor); ++ av_freep(&device_name); ++ av_freep(&device_exts); + if (event) + clReleaseEvent(event); if (ctx->kernel) @@ -2770,7 +2770,7 @@ Index: FFmpeg/libavfilter/vf_tonemap_opencl.c ret = ff_opencl_filter_config_output(outlink); if (ret < 0) return ret; -@@ -308,13 +930,49 @@ static int launch_kernel(AVFilterContext +@@ -308,13 +932,49 @@ static int launch_kernel(AVFilterContext size_t global_work[2]; size_t local_work[2]; cl_int cle; @@ -2780,6 +2780,11 @@ Index: FFmpeg/libavfilter/vf_tonemap_opencl.c + err = AVERROR(EIO); + goto fail; + } ++ ++ if (ctx->out_planes > 2 && !output->data[2]) { ++ 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]); @@ -2787,11 +2792,6 @@ Index: FFmpeg/libavfilter/vf_tonemap_opencl.c - 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; -+ } -+ + if (ctx->in_planes > 2 && !input->data[2]) { + err = AVERROR(EIO); + goto fail; @@ -2826,7 +2826,7 @@ Index: FFmpeg/libavfilter/vf_tonemap_opencl.c local_work[0] = 16; local_work[1] = 16; -@@ -338,12 +996,10 @@ static int tonemap_opencl_filter_frame(A +@@ -338,12 +998,10 @@ static int tonemap_opencl_filter_frame(A AVFilterContext *avctx = inlink->dst; AVFilterLink *outlink = avctx->outputs[0]; TonemapOpenCLContext *ctx = avctx->priv; @@ -2840,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 +1007,6 @@ static int tonemap_opencl_filter_frame(A +@@ -351,7 +1009,6 @@ static int tonemap_opencl_filter_frame(A if (!input->hw_frames_ctx) return AVERROR(EINVAL); @@ -2848,7 +2848,7 @@ Index: FFmpeg/libavfilter/vf_tonemap_opencl.c output = ff_get_video_buffer(outlink, outlink->w, outlink->h); if (!output) { -@@ -363,17 +1018,65 @@ static int tonemap_opencl_filter_frame(A +@@ -363,17 +1020,65 @@ static int tonemap_opencl_filter_frame(A if (err < 0) goto fail; @@ -2920,7 +2920,7 @@ Index: FFmpeg/libavfilter/vf_tonemap_opencl.c ctx->trc_in = input->color_trc; ctx->trc_out = output->color_trc; -@@ -385,72 +1088,50 @@ static int tonemap_opencl_filter_frame(A +@@ -385,72 +1090,50 @@ static int tonemap_opencl_filter_frame(A ctx->range_out = output->color_range; ctx->chroma_loc = output->chroma_location; @@ -3016,7 +3016,7 @@ Index: FFmpeg/libavfilter/vf_tonemap_opencl.c av_frame_free(&input); av_frame_free(&output); return err; -@@ -458,62 +1139,101 @@ fail: +@@ -458,62 +1141,101 @@ fail: static av_cold void tonemap_opencl_uninit(AVFilterContext *avctx) { @@ -3165,7 +3165,7 @@ Index: FFmpeg/libavfilter/vf_tonemap_opencl.c { NULL } }; -@@ -541,11 +1261,12 @@ const AVFilter ff_vf_tonemap_opencl = { +@@ -541,11 +1263,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, diff --git a/debian/patches/0054-add-cuda-transpose-filter-impl.patch b/debian/patches/0054-add-cuda-transpose-filter-impl.patch index ff9157684..4f4e7e94c 100644 --- a/debian/patches/0054-add-cuda-transpose-filter-impl.patch +++ b/debian/patches/0054-add-cuda-transpose-filter-impl.patch @@ -40,7 +40,7 @@ Index: FFmpeg/libavfilter/vf_transpose_cuda.c =================================================================== --- /dev/null +++ FFmpeg/libavfilter/vf_transpose_cuda.c -@@ -0,0 +1,482 @@ +@@ -0,0 +1,483 @@ +/* + * Copyright (C) 2024 NyanMisaka + * @@ -111,6 +111,7 @@ Index: FFmpeg/libavfilter/vf_transpose_cuda.c + CUfunction cu_func_ushort; + CUstream cu_stream; + ++ int flip_wh; + int passthrough; ///< PassthroughType, landscape passthrough mode enabled + int dir; ///< TransposeDir +} TransposeCUDAContext; @@ -250,7 +251,7 @@ Index: FFmpeg/libavfilter/vf_transpose_cuda.c + TransposeCUDAContext *s = ctx->priv; + CUcontext dummy, cuda_ctx; + CudaFunctions *cu; -+ int flip_wh, ret; ++ int ret = 0; + + if ((inlink->w >= inlink->h && s->passthrough == TRANSPOSE_PT_TYPE_LANDSCAPE) || + (inlink->w <= inlink->h && s->passthrough == TRANSPOSE_PT_TYPE_PORTRAIT)) { @@ -275,16 +276,16 @@ Index: FFmpeg/libavfilter/vf_transpose_cuda.c + case TRANSPOSE_CLOCK_FLIP: + outlink->w = inlink->h; + outlink->h = inlink->w; -+ flip_wh = 1; ++ s->flip_wh = 1; + break; + default: + outlink->w = inlink->w; + outlink->h = inlink->h; -+ flip_wh = 0; ++ s->flip_wh = 0; + break; + } + -+ if (flip_wh && inlink->sample_aspect_ratio.num) ++ if (s->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; @@ -302,31 +303,26 @@ Index: FFmpeg/libavfilter/vf_transpose_cuda.c + + ret = ff_cuda_load_module(ctx, s->hwctx, &s->cu_module, + ff_vf_transpose_cuda_ptx_data, ff_vf_transpose_cuda_ptx_len); -+ if (ret < 0) { -+ CHECK_CU(cu->cuCtxPopCurrent(&dummy)); -+ return ret; -+ } ++ if (ret < 0) ++ goto exit; + + ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_uchar, + s->cu_module, "Transpose_Cuda_uchar")); -+ if (ret < 0) { -+ CHECK_CU(cu->cuCtxPopCurrent(&dummy)); -+ return ret; -+ } ++ if (ret < 0) ++ goto exit; + + ret = CHECK_CU(cu->cuModuleGetFunction(&s->cu_func_ushort, + s->cu_module, "Transpose_Cuda_ushort")); -+ if (ret < 0) { -+ CHECK_CU(cu->cuCtxPopCurrent(&dummy)); -+ return ret; -+ } -+ -+ CHECK_CU(cu->cuCtxPopCurrent(&dummy)); ++ if (ret < 0) ++ goto exit; + + av_log(ctx, AV_LOG_VERBOSE, + "w:%d h:%d dir:%d -> w:%d h:%d\n", + inlink->w, inlink->h, s->dir, outlink->w, outlink->h); -+ return 0; ++exit: ++ CHECK_CU(cu->cuCtxPopCurrent(&dummy)); ++ ++ return ret; +} + +static int call_transpose_kernel(AVFilterContext *ctx, @@ -418,6 +414,11 @@ Index: FFmpeg/libavfilter/vf_transpose_cuda.c + if (ret < 0) + return ret; + ++ if (s->flip_wh && in->sample_aspect_ratio.num) ++ out->sample_aspect_ratio = av_inv_q(in->sample_aspect_ratio); ++ else ++ out->sample_aspect_ratio = in->sample_aspect_ratio; ++ + return 0; +} + diff --git a/debian/patches/0055-add-flip-feat-to-opencl-transpose-filter.patch b/debian/patches/0055-add-flip-feat-to-opencl-transpose-filter.patch index 96c4ab733..22af650c2 100644 --- a/debian/patches/0055-add-flip-feat-to-opencl-transpose-filter.patch +++ b/debian/patches/0055-add-flip-feat-to-opencl-transpose-filter.patch @@ -19,7 +19,15 @@ Index: FFmpeg/libavfilter/vf_transpose_opencl.c =================================================================== --- FFmpeg.orig/libavfilter/vf_transpose_opencl.c +++ FFmpeg/libavfilter/vf_transpose_opencl.c -@@ -103,8 +103,20 @@ static int transpose_opencl_config_outpu +@@ -32,6 +32,7 @@ + typedef struct TransposeOpenCLContext { + OpenCLFilterContext ocf; + int initialised; ++ int flip_wh; + int passthrough; ///< PassthroughType, landscape passthrough mode enabled + int dir; ///< TransposeDir + cl_kernel kernel; +@@ -103,23 +104,34 @@ static int transpose_opencl_config_outpu return AVERROR(EINVAL); } @@ -32,17 +40,25 @@ Index: FFmpeg/libavfilter/vf_transpose_opencl.c + case TRANSPOSE_CLOCK_FLIP: + s->ocf.output_width = inlink->h; + s->ocf.output_height = inlink->w; ++ s->flip_wh = 1; + break; + default: + s->ocf.output_width = inlink->w; + s->ocf.output_height = inlink->h; ++ s->flip_wh = 0; + break; + } + ret = ff_opencl_filter_config_output(outlink); if (ret < 0) return ret; -@@ -116,10 +128,8 @@ static int transpose_opencl_config_outpu + +- if (inlink->sample_aspect_ratio.num) +- outlink->sample_aspect_ratio = av_div_q((AVRational) { 1, 1 }, +- inlink->sample_aspect_ratio); ++ if (s->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; av_log(avctx, AV_LOG_VERBOSE, @@ -55,6 +71,22 @@ Index: FFmpeg/libavfilter/vf_transpose_opencl.c return 0; } +@@ -163,12 +175,10 @@ static int transpose_opencl_filter_frame + if (err < 0) + goto fail; + +- if (input->sample_aspect_ratio.num == 0) { ++ if (ctx->flip_wh && input->sample_aspect_ratio.num) ++ output->sample_aspect_ratio = av_inv_q(input->sample_aspect_ratio); ++ else + output->sample_aspect_ratio = input->sample_aspect_ratio; +- } else { +- output->sample_aspect_ratio.num = input->sample_aspect_ratio.den; +- output->sample_aspect_ratio.den = input->sample_aspect_ratio.num; +- } + + if (!ctx->initialised) { + err = transpose_opencl_init(avctx); @@ -237,11 +247,14 @@ static av_cold void transpose_opencl_uni #define OFFSET(x) offsetof(TransposeOpenCLContext, x) #define FLAGS (AV_OPT_FLAG_FILTERING_PARAM | AV_OPT_FLAG_VIDEO_PARAM) From d58d7cf30af954336439450a9b59f98f6795416f Mon Sep 17 00:00:00 2001 From: nyanmisaka Date: Sat, 25 Apr 2026 21:48:08 +0800 Subject: [PATCH 2/4] Update build scripts Signed-off-by: nyanmisaka --- builder/scripts.d/20-libiconv.sh | 14 +++++++------- docker-build.sh | 8 ++++---- 2 files changed, 11 insertions(+), 11 deletions(-) diff --git a/builder/scripts.d/20-libiconv.sh b/builder/scripts.d/20-libiconv.sh index b886fd36d..d3f512eca 100755 --- a/builder/scripts.d/20-libiconv.sh +++ b/builder/scripts.d/20-libiconv.sh @@ -1,9 +1,14 @@ #!/bin/bash +#SCRIPT_ORIG="https://git.savannah.gnu.org/git/libiconv.git" SCRIPT_REPO="https://skia.googlesource.com/third_party/libiconv" SCRIPT_COMMIT="v1.18" SCRIPT_TAGFILTER="v?.*" +SCRIPT_ORIG2="git://git.savannah.gnu.org/gnulib.git" +SCRIPT_REPO2="https://github.com/coreutils/gnulib.git" +SCRIPT_COMMIT2="e9c1d94f58eaacee919bb2015da490b980a5eedf" + ffbuild_enabled() { return 0 } @@ -15,14 +20,9 @@ ffbuild_dockerbuild() { git-mini-clone "$SCRIPT_REPO" "$SCRIPT_COMMIT" iconv cd iconv - cat < ./.gitmodules -[subcheckout "gnulib"] - url = https://github.com/coreutils/gnulib.git - path = gnulib -EOF - + sed -i "s|${SCRIPT_ORIG2}|${SCRIPT_REPO2}|g" ./.gitmodules ./gitsub.sh pull - ./gitsub.sh checkout gnulib e9c1d94f58eaacee919bb2015da490b980a5eedf + ./gitsub.sh checkout gnulib "$SCRIPT_COMMIT2" # No automake 1.17 packaged anywhere yet. sed -i 's/-1.17/-1.16/' Makefile.devel diff --git a/docker-build.sh b/docker-build.sh index 30d7253fc..278759e5d 100755 --- a/docker-build.sh +++ b/docker-build.sh @@ -325,10 +325,10 @@ prepare_extra_amd64() { pushd ${SOURCE_DIR} mkdir libdrm pushd libdrm - libdrm_ver="2.4.131" - libdrm_link="https://dri.freedesktop.org/libdrm/libdrm-${libdrm_ver}.tar.xz" - wget ${libdrm_link} -O libdrm.tar.xz - tar xaf libdrm.tar.xz + libdrm_ver="libdrm-2.4.131" + libdrm_link="https://gitlab.freedesktop.org/mesa/libdrm/-/archive/${libdrm_ver}/libdrm-${libdrm_ver}.tar.gz" + wget ${libdrm_link} -O libdrm.tar.gz + tar xaf libdrm.tar.gz meson setup libdrm-${libdrm_ver} drm_build \ --prefix=${TARGET_DIR} \ --libdir=lib \ From 09157f8c5c0ebd12c504afdd9273e44454f67778 Mon Sep 17 00:00:00 2001 From: nyanmisaka Date: Sat, 25 Apr 2026 21:33:50 +0800 Subject: [PATCH 3/4] Enable Ubuntu 26.04 LTS (Resolute) in release builds Signed-off-by: nyanmisaka --- .github/workflows/publish.yaml | 5 +++-- 1 file changed, 3 insertions(+), 2 deletions(-) diff --git a/.github/workflows/publish.yaml b/.github/workflows/publish.yaml index 26a8cfb20..9b2d47f80 100644 --- a/.github/workflows/publish.yaml +++ b/.github/workflows/publish.yaml @@ -22,7 +22,7 @@ jobs: uses: ./.github/workflows/_meta.yaml with: distro: 'ubuntu' - codenames: '["jammy", "noble"]' + codenames: '["jammy", "noble", "resolute"]' architectures: '["amd64", "arm64"]' release: true secrets: @@ -86,7 +86,8 @@ jobs: {distro: 'debian', codename: 'bookworm'}, {distro: 'debian', codename: 'trixie'}, {distro: 'ubuntu', codename: 'jammy'}, - {distro: 'ubuntu', codename: 'noble'} + {distro: 'ubuntu', codename: 'noble'}, + {distro: 'ubuntu', codename: 'resolute'} ] steps: - name: Import packages into reprepro From 0eda89a0afc44dcc06d8b7ab1b6a871f3d978f38 Mon Sep 17 00:00:00 2001 From: nyanmisaka Date: Sat, 25 Apr 2026 21:43:11 +0800 Subject: [PATCH 4/4] Bump version to 7.1.3-6 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 0d186f9fb..7d6b20a35 100644 --- a/build.yaml +++ b/build.yaml @@ -1,7 +1,7 @@ --- # We just wrap `build` so this is really it name: "jellyfin-ffmpeg" -version: "7.1.3-5" +version: "7.1.3-6" packages: - bullseye-amd64 - bullseye-arm64 diff --git a/debian/changelog b/debian/changelog index 9eb9c8bdc..061432e49 100644 --- a/debian/changelog +++ b/debian/changelog @@ -1,3 +1,10 @@ +jellyfin-ffmpeg (7.1.3-6) unstable; urgency=medium + + * Enable Ubuntu 26.04 LTS (Resolute) in release builds + * Backport some fixes from new branch + + -- nyanmisaka Sat, 25 Apr 2026 21:42:18 +0800 + jellyfin-ffmpeg (7.1.3-5) unstable; urgency=medium * ffprobe: avoid runaway only_first_vframe probing