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

Merge pull request #697 from nyanmisaka/enable-ubu2604-in-release

Browse Source
This commit is contained in:
Joshua M. Boniface
2026-04-29 23:41:47 -04:00
committed by GitHub
parent f4360d7dee 0eda89a0af
commit f6f3061ba5
11 changed files with 186 additions and 89 deletions
Download Patch File Download Diff File Expand all files Collapse all files
.github/workflows/publish.yaml
+3 -2
View File
@@ -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
build.yaml
+1 -1
View File
@@ -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
builder/scripts.d/20-libiconv.sh
+7 -7
View File
@@ -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 <<EOF > ./.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
debian/changelog
Vendored
+7
View File
@@ -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 <nst799610810@gmail.com> Sat, 25 Apr 2026 21:42:18 +0800
jellyfin-ffmpeg (7.1.3-5) unstable; urgency=medium
* ffprobe: avoid runaway only_first_vframe probing
debian/patches/0002-update-cuda-func-header.patch
Vendored
+2 -1
View File
@@ -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<float4>(c
@@ -156,37 +147,42 @@ inline __device__ float4 tex2D<float4>(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); }
debian/patches/0003-add-enhanced-cuda-pixfmt-converter-impl.patch
Vendored
+62 -7
View File
@@ -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<flo
);
}
+template<typename T>
+inline __device__ T saturate_rintf(T a, float factor) {
+ return rintf(__saturatef(a) * factor);
+}
+
+template<>
+inline __device__ float2 saturate_rintf<float2>(float2 a, float factor) {
+ return make_float2(
+ saturate_rintf(a.x, factor),
+ saturate_rintf(a.y, factor)
+ );
+}
+
+template<>
+inline __device__ float4 saturate_rintf<float4>(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<floatT>(tex, (x), (y))
return from_floatN<T, floatT>(
- apply_coeffs<floatT>(coeffsY,
- apply_coeffs<floatT>(coeffsX, PIX(px - 1, py - 1), PIX(px, py - 1), PIX(px + 1, py - 1), PIX(px + 2, py - 1)),
- apply_coeffs<floatT>(coeffsX, PIX(px - 1, py ), PIX(px, py ), PIX(px + 1, py ), PIX(px + 2, py )),
- apply_coeffs<floatT>(coeffsX, PIX(px - 1, py + 1), PIX(px, py + 1), PIX(px + 1, py + 1), PIX(px + 2, py + 1)),
- apply_coeffs<floatT>(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<floatT>(coeffsY,
+ apply_coeffs<floatT>(coeffsX, PIX(px - 1, py - 1), PIX(px, py - 1), PIX(px + 1, py - 1), PIX(px + 2, py - 1)),
+ apply_coeffs<floatT>(coeffsX, PIX(px - 1, py ), PIX(px, py ), PIX(px + 1, py ), PIX(px + 2, py )),
+ apply_coeffs<floatT>(coeffsX, PIX(px - 1, py + 1), PIX(px, py + 1), PIX(px + 1, py + 1), PIX(px + 2, py + 1)),
+ apply_coeffs<floatT>(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, \
debian/patches/0004-add-cuda-tonemap-impl.patch
Vendored
+12 -12
View File
@@ -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<float>(frame.tex[0], x, y);
+ float2 px_uv = tex2D<float2>(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;
+}
+
debian/patches/0007-add-bt2390-eetf-and-code-refactor-to-opencl-tonemap.patch
Vendored
+33 -33
View File
@@ -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,
debian/patches/0054-add-cuda-transpose-filter-impl.patch
Vendored
+21 -20
View File
@@ -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;
+}
+
debian/patches/0055-add-flip-feat-to-opencl-transpose-filter.patch
Vendored
+34 -2
View File
@@ -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)
docker-build.sh
+4 -4
View File
@@ -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 \