PR #24745 opened by mantonx
URL: https://code.ffmpeg.org/FFmpeg/FFmpeg/pulls/24745
Patch URL: https://code.ffmpeg.org/FFmpeg/FFmpeg/pulls/24745.patch
overlay_opencl gives wrong output in two cases:
1. With a subsampled main and a separate-alpha overlay (yuv420p main,
yuva420p overlay, the pair the documentation shows), the chroma
planes read alpha from the wrong position whenever the overlay is
not at (0,0). Patch 1 fixes the coordinate.
2. With an nv12 or p010 main and a yuva420p or yuv420p overlay, the
filter takes alpha from the overlay's V plane and blends the
interleaved UV plane against the overlay's U plane only. The result
is a dim, green-tinted overlay, blended over only a quarter of its
area. NV12 is what VAAPI and QSV frames map to in OpenCL, so this is
the natural way to put a logo over mapped hardware frames. Patch 2
adds kernels for this case and rejects layouts that none of the
kernels handle, which were previously blended incorrectly without
an error.
Test: lavfi sources only, OpenCL result against the CPU overlay filter,
reporting the maximum per-plane difference over the whole frame. The
filtergraph is:
color=c=0x203040:s=320x240,format=nv12,split[mg][mc];
[email protected]:s=128x64,format=yuva420p,split[og][oc];
[mg]hwupload[mgh]; [og]hwupload[ogh];
[mgh][ogh]overlay_opencl=x=64:y=32,hwdownload,format=nv12,format=yuv420p[g];
[mc][oc]overlay=x=64:y=32:format=yuv420,format=yuv420p[c];
[g][c]blend=all_mode=difference,signalstats
(run with -init_hw_device opencl=ocl -filter_hw_device ocl, then read
YMAX/UMAX/VMAX)
main overlay pos before (Y/U/V) after
nv12 yuva420p 64,32 52 / 39 / 117 0 / 0 / 0
nv12 yuv420p 64,32 86 / 65 / 140 0 / 0 / 0
nv12 yuva420p 0,0 52 / 39 / 117 0 / 0 / 0
yuv420p yuva420p 64,32 0 / 39 / 34 0 / 0 / 0
yuv420p yuv420p 64,32 0 / 0 / 0 0 / 0 / 0
p010 yuva420p 64,32 - 0 / 1 / 1 (vs 8-bit CPU)
The same results hold with testsrc2 for both inputs (overlay alpha a
uniform 60%), so the reads are aligned inside the overlay as well as at
its edges. Tested on NVIDIA OpenCL, and on Intel Arc (A380) with Intel
compute-runtime OpenCL; results below.
There is no FATE test: FATE has no OpenCL (or other hwcontext filter)
tests, and a test would need an OpenCL device at run time, which FATE
clients are not required to have and which the build system cannot
detect at configure time.
Not changed: with a subsampled main, the chroma planes use the alpha
of the top-left luma sample of each block, where the CPU filter
averages the block. With a uniform alpha the results are identical;
along alpha edges they can differ.
Intel Arc A380 (compute-runtime OpenCL), the same test run against n9.0.1 and
against this branch:
| main | overlay | case | n9.0.1 (Y/U/V) | this branch |
|---|---|---|---|---|
| nv12 | yuva420p | solid, 64,32 | 52 / 39 / 117 | 0 / 0 / 0 |
| nv12 | yuva420p | testsrc2, 0,0 | 94 / 130 / 226 | 0 / 0 / 0 |
| nv12 | yuv420p | solid, 64,32 | 86 / 65 / 140 | 0 / 0 / 0 |
| nv12 | yuv420p | testsrc2, 0,0 | 156 / 216 / 226 | 0 / 0 / 0 |
| yuv420p | yuva420p | solid, 64,32 | 0 / 39 / 34 | 0 / 0 / 0 |
| yuv420p | yuva420p | testsrc2, 0,0 | 0 / 0 / 0 | 0 / 0 / 0 |
This was found while putting a channel logo over VAAPI-decoded frames mapped to
OpenCL, where stock `overlay_opencl` produced a dim, green-tinted logo.
The fix was developed with AI assistance (Claude) and verified on the hardware
above.
>From 399743b674d103353b72661a6074701a0ee5c55a Mon Sep 17 00:00:00 2001
From: Matthew Panton <[email protected]>
Date: Sun, 27 Sep 2026 16:45:56 -0400
Subject: [PATCH 1/2] avfilter/vf_overlay_opencl: fix alpha coordinates on
subsampled planes
overlay_external_alpha computed the alpha plane coordinate as
loc * alpha_adj - overlay_position. Both loc and overlay_position are in
the coordinates of the plane being blended, so for a subsampled plane
the scale was applied to the position as well as the offset. This is
only correct with the overlay at (0,0). Anywhere else the chroma planes
take alpha from the wrong pixels, partly outside the alpha image.
Scale the offset into the overlay instead.
With a yuv420p main and a 128x64 yuva420p overlay at (64,32), the
maximum difference from the CPU overlay filter goes from 39 (U) and
34 (V) to 0.
---
libavfilter/opencl/overlay.cl | 4 ++--
1 file changed, 2 insertions(+), 2 deletions(-)
diff --git a/libavfilter/opencl/overlay.cl b/libavfilter/opencl/overlay.cl
index 8c783d0edc..6764c40686 100644
--- a/libavfilter/opencl/overlay.cl
+++ b/libavfilter/opencl/overlay.cl
@@ -94,8 +94,8 @@ __kernel void overlay_external_alpha(__write_only image2d_t
dst,
float4 in_main = read_imagef(main, sampler, loc);
float4 in_overlay = read_imagef(overlay, sampler, loc - loc_overlay);
- int2 loc_alpha = (int2)(loc.x * alpha_adj_x,
- loc.y * alpha_adj_y) - loc_overlay;
+ int2 loc_alpha = (int2)((loc.x - x_position) * alpha_adj_x,
+ (loc.y - y_position) * alpha_adj_y);
float4 in_alpha = read_imagef(alpha, sampler, loc_alpha);
float4 val = in_overlay * in_alpha.x + in_main * (1.0f - in_alpha.x);
--
2.52.0
>From 7f340cc12ec7e909631b5b99c98ea0a595e61118 Mon Sep 17 00:00:00 2001
From: Matthew Panton <[email protected]>
Date: Sun, 27 Sep 2026 16:48:24 -0400
Subject: [PATCH 2/2] avfilter/vf_overlay_opencl: support a planar overlay on a
semi-planar main
With an nv12 (or p010) main and a yuva420p or yuv420p overlay, the plane
counts differ, so the filter picked overlay_external_alpha and took the
alpha from overlay->data[nb_planes]. With a two-plane main that is the
overlay's V plane, not its alpha, so a yuv420p overlay was also blended
as if it had alpha. The main's interleaved UV plane was then blended
against the overlay's U plane only, pulling V towards zero.
NV12 is the layout a VAAPI or QSV surface maps to in OpenCL, so a
yuva420p logo over mapped hardware frames came out dim, green-tinted
and blended over only a quarter of its area.
Add kernels which read U and V from the overlay's two chroma planes and
write them to the main's interleaved plane, and take the alpha from the
plane the overlay's pixel format places it in. Layouts that none of the
kernels handle are now rejected instead of being blended incorrectly.
With a 320x240 nv12 main and a 128x64 yuva420p overlay at (64,32),
the maximum difference from the CPU overlay filter goes from 52 (Y),
39 (U) and 117 (V) to 0 in every plane, and likewise for a yuv420p
overlay (from 86, 65 and 140).
---
doc/filters.texi | 10 ++++-
libavfilter/opencl/overlay.cl | 68 ++++++++++++++++++++++++++++
libavfilter/vf_overlay_opencl.c | 80 +++++++++++++++++++++++++++------
3 files changed, 144 insertions(+), 14 deletions(-)
diff --git a/doc/filters.texi b/doc/filters.texi
index 1a9bbeff7e..70d654c667 100644
--- a/doc/filters.texi
+++ b/doc/filters.texi
@@ -28400,7 +28400,10 @@ Non-local Means denoise filter through OpenCL, this
filter accepts same options
Overlay one video on top of another.
It takes two inputs and has one output. The first input is the "main" video on
which the second input is overlaid.
-This filter requires same memory layout for all the inputs. So, format
conversion may be needed.
+The overlay must have the same memory layout as the main input, optionally
with an additional alpha plane.
+As an exception, a semi-planar YUV main input (such as nv12 or p010) also
accepts a planar YUV overlay
+with the same chroma subsampling, with or without an alpha plane (such as
yuv420p or yuva420p).
+Other combinations need a format conversion.
The filter accepts the following options:
@@ -28429,6 +28432,11 @@ The inputs have same memory layout for color channels
, the overlay has addition
@example
-i INPUT -i LOGO -filter_complex "[0:v]hwupload[a], [1:v]format=yuva420p,
hwupload[b], [a][b]overlay_opencl, hwdownload" OUTPUT
@end example
+@item
+Overlay a yuva420p LOGO on an nv12 INPUT.
+@example
+-i INPUT -i LOGO -filter_complex "[0:v]format=nv12, hwupload[a],
[1:v]format=yuva420p, hwupload[b], [a][b]overlay_opencl, hwdownload" OUTPUT
+@end example
@end itemize
diff --git a/libavfilter/opencl/overlay.cl b/libavfilter/opencl/overlay.cl
index 6764c40686..72168dca4a 100644
--- a/libavfilter/opencl/overlay.cl
+++ b/libavfilter/opencl/overlay.cl
@@ -102,3 +102,71 @@ __kernel void overlay_external_alpha(__write_only
image2d_t dst,
write_imagef(dst, loc, val);
}
}
+
+// The _uv variants blend into an interleaved chroma plane (as in NV12)
+// from an overlay which carries U and V in two separate planes.
+
+__kernel void overlay_no_alpha_uv(__write_only image2d_t dst,
+ __read_only image2d_t main,
+ __read_only image2d_t overlay_u,
+ __read_only image2d_t overlay_v,
+ int x_position,
+ int y_position)
+{
+ const sampler_t sampler = (CLK_NORMALIZED_COORDS_FALSE |
+ CLK_FILTER_NEAREST);
+
+ int2 overlay_size = get_image_dim(overlay_u);
+ int2 loc = (int2)(get_global_id(0), get_global_id(1));
+
+ if (loc.x < x_position ||
+ loc.y < y_position ||
+ loc.x >= overlay_size.x + x_position ||
+ loc.y >= overlay_size.y + y_position) {
+ float4 val = read_imagef(main, sampler, loc);
+ write_imagef(dst, loc, val);
+ } else {
+ int2 loc_overlay = loc - (int2)(x_position, y_position);
+ float4 in_u = read_imagef(overlay_u, sampler, loc_overlay);
+ float4 in_v = read_imagef(overlay_v, sampler, loc_overlay);
+ float4 val = (float4)(in_u.x, in_v.x, 0.0f, 1.0f);
+ write_imagef(dst, loc, val);
+ }
+}
+
+__kernel void overlay_external_alpha_uv(__write_only image2d_t dst,
+ __read_only image2d_t main,
+ __read_only image2d_t overlay_u,
+ __read_only image2d_t overlay_v,
+ __read_only image2d_t alpha,
+ int x_position,
+ int y_position,
+ int alpha_adj_x,
+ int alpha_adj_y)
+{
+ const sampler_t sampler = (CLK_NORMALIZED_COORDS_FALSE |
+ CLK_FILTER_NEAREST);
+
+ int2 overlay_size = get_image_dim(overlay_u);
+ int2 loc = (int2)(get_global_id(0), get_global_id(1));
+
+ if (loc.x < x_position ||
+ loc.y < y_position ||
+ loc.x >= overlay_size.x + x_position ||
+ loc.y >= overlay_size.y + y_position) {
+ float4 val = read_imagef(main, sampler, loc);
+ write_imagef(dst, loc, val);
+ } else {
+ int2 loc_overlay = loc - (int2)(x_position, y_position);
+ float4 in_main = read_imagef(main, sampler, loc);
+ float4 in_u = read_imagef(overlay_u, sampler, loc_overlay);
+ float4 in_v = read_imagef(overlay_v, sampler, loc_overlay);
+ float4 in_overlay = (float4)(in_u.x, in_v.x, 0.0f, 1.0f);
+
+ int2 loc_alpha = loc_overlay * (int2)(alpha_adj_x, alpha_adj_y);
+ float4 in_alpha = read_imagef(alpha, sampler, loc_alpha);
+
+ float4 val = in_overlay * in_alpha.x + in_main * (1.0f - in_alpha.x);
+ write_imagef(dst, loc, val);
+ }
+}
diff --git a/libavfilter/vf_overlay_opencl.c b/libavfilter/vf_overlay_opencl.c
index e930aced04..1bbb03e5d1 100644
--- a/libavfilter/vf_overlay_opencl.c
+++ b/libavfilter/vf_overlay_opencl.c
@@ -32,6 +32,7 @@ typedef struct OverlayOpenCLContext {
int initialised;
cl_kernel kernel;
+ cl_kernel kernel_uv;
cl_command_queue command_queue;
FFFrameSync fs;
@@ -40,6 +41,7 @@ typedef struct OverlayOpenCLContext {
int x_subsample;
int y_subsample;
int alpha_separate;
+ int alpha_plane;
int x_position;
int y_position;
@@ -52,9 +54,10 @@ static int overlay_opencl_load(AVFilterContext *avctx,
OverlayOpenCLContext *ctx = avctx->priv;
cl_int cle;
const char *source = ff_source_overlay_cl;
- const char *kernel;
+ const char *kernel, *kernel_uv = NULL;
const AVPixFmtDescriptor *main_desc, *overlay_desc;
int err, i, main_planes, overlay_planes;
+ int main_semiplanar, overlay_planar;
main_desc = av_pix_fmt_desc_get(main_format);
overlay_desc = av_pix_fmt_desc_get(overlay_format);
@@ -79,18 +82,47 @@ static int overlay_opencl_load(AVFilterContext *avctx,
ctx->x_subsample, ctx->y_subsample);
}
+ main_semiplanar = !(main_desc->flags & AV_PIX_FMT_FLAG_RGB) &&
+ main_planes == 2 &&
+ main_desc->comp[1].plane == 1 &&
+ main_desc->comp[2].plane == 1;
+ overlay_planar = !(overlay_desc->flags & AV_PIX_FMT_FLAG_RGB) &&
+ overlay_desc->comp[1].plane == 1 &&
+ overlay_desc->comp[2].plane == 2 &&
+ overlay_desc->log2_chroma_w == main_desc->log2_chroma_w
&&
+ overlay_desc->log2_chroma_h == main_desc->log2_chroma_h;
+
if (main_planes == overlay_planes) {
if (main_desc->nb_components == overlay_desc->nb_components)
kernel = "overlay_no_alpha";
else
kernel = "overlay_internal_alpha";
ctx->alpha_separate = 0;
- } else {
+ } else if (main_semiplanar && overlay_planar) {
+ // The main chroma plane interleaves U and V, which the overlay
+ // carries in two separate planes.
+ ctx->alpha_separate = !!(overlay_desc->flags & AV_PIX_FMT_FLAG_ALPHA);
+ ctx->alpha_plane = overlay_desc->comp[3].plane;
+ if (ctx->alpha_separate) {
+ kernel = "overlay_external_alpha";
+ kernel_uv = "overlay_external_alpha_uv";
+ } else {
+ kernel = "overlay_no_alpha";
+ kernel_uv = "overlay_no_alpha_uv";
+ }
+ } else if (overlay_planes == main_planes + 1 &&
+ overlay_desc->flags & AV_PIX_FMT_FLAG_ALPHA) {
kernel = "overlay_external_alpha";
ctx->alpha_separate = 1;
+ ctx->alpha_plane = overlay_desc->comp[3].plane;
+ } else {
+ av_log(avctx, AV_LOG_ERROR, "Unsupported overlay format %s "
+ "on main format %s.\n", overlay_desc->name, main_desc->name);
+ return AVERROR(ENOSYS);
}
- av_log(avctx, AV_LOG_DEBUG, "Using kernel %s.\n", kernel);
+ av_log(avctx, AV_LOG_DEBUG, "Using kernel %s%s%s.\n", kernel,
+ kernel_uv ? " and " : "", kernel_uv ? kernel_uv : "");
err = ff_opencl_filter_load_program(avctx, &source, 1);
if (err < 0)
@@ -105,6 +137,11 @@ static int overlay_opencl_load(AVFilterContext *avctx,
ctx->kernel = clCreateKernel(ctx->ocf.program, kernel, &cle);
CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create kernel %d.\n", cle);
+ if (kernel_uv) {
+ ctx->kernel_uv = clCreateKernel(ctx->ocf.program, kernel_uv, &cle);
+ CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create kernel %d.\n", cle);
+ }
+
ctx->initialised = 1;
return 0;
@@ -113,6 +150,8 @@ fail:
clReleaseCommandQueue(ctx->command_queue);
if (ctx->kernel)
clReleaseKernel(ctx->kernel);
+ if (ctx->kernel_uv)
+ clReleaseKernel(ctx->kernel_uv);
return err;
}
@@ -155,41 +194,49 @@ static int overlay_opencl_blend(FFFrameSync *fs)
}
for (plane = 0; plane < ctx->nb_planes; plane++) {
+ cl_kernel kernel = plane > 0 && ctx->kernel_uv ? ctx->kernel_uv
+ : ctx->kernel;
kernel_arg = 0;
mem = (cl_mem)output->data[plane];
- CL_SET_KERNEL_ARG(ctx->kernel, kernel_arg, cl_mem, &mem);
+ CL_SET_KERNEL_ARG(kernel, kernel_arg, cl_mem, &mem);
kernel_arg++;
mem = (cl_mem)input_main->data[plane];
- CL_SET_KERNEL_ARG(ctx->kernel, kernel_arg, cl_mem, &mem);
+ CL_SET_KERNEL_ARG(kernel, kernel_arg, cl_mem, &mem);
kernel_arg++;
mem = (cl_mem)input_overlay->data[plane];
- CL_SET_KERNEL_ARG(ctx->kernel, kernel_arg, cl_mem, &mem);
+ CL_SET_KERNEL_ARG(kernel, kernel_arg, cl_mem, &mem);
kernel_arg++;
+ if (kernel == ctx->kernel_uv) {
+ mem = (cl_mem)input_overlay->data[plane + 1];
+ CL_SET_KERNEL_ARG(kernel, kernel_arg, cl_mem, &mem);
+ kernel_arg++;
+ }
+
if (ctx->alpha_separate) {
- mem = (cl_mem)input_overlay->data[ctx->nb_planes];
- CL_SET_KERNEL_ARG(ctx->kernel, kernel_arg, cl_mem, &mem);
+ mem = (cl_mem)input_overlay->data[ctx->alpha_plane];
+ CL_SET_KERNEL_ARG(kernel, kernel_arg, cl_mem, &mem);
kernel_arg++;
}
x = ctx->x_position / (plane == 0 ? 1 : ctx->x_subsample);
y = ctx->y_position / (plane == 0 ? 1 : ctx->y_subsample);
- CL_SET_KERNEL_ARG(ctx->kernel, kernel_arg, cl_int, &x);
+ CL_SET_KERNEL_ARG(kernel, kernel_arg, cl_int, &x);
kernel_arg++;
- CL_SET_KERNEL_ARG(ctx->kernel, kernel_arg, cl_int, &y);
+ CL_SET_KERNEL_ARG(kernel, kernel_arg, cl_int, &y);
kernel_arg++;
if (ctx->alpha_separate) {
cl_int alpha_adj_x = plane == 0 ? 1 : ctx->x_subsample;
cl_int alpha_adj_y = plane == 0 ? 1 : ctx->y_subsample;
- CL_SET_KERNEL_ARG(ctx->kernel, kernel_arg, cl_int, &alpha_adj_x);
+ CL_SET_KERNEL_ARG(kernel, kernel_arg, cl_int, &alpha_adj_x);
kernel_arg++;
- CL_SET_KERNEL_ARG(ctx->kernel, kernel_arg, cl_int, &alpha_adj_y);
+ CL_SET_KERNEL_ARG(kernel, kernel_arg, cl_int, &alpha_adj_y);
kernel_arg++;
}
@@ -198,7 +245,7 @@ static int overlay_opencl_blend(FFFrameSync *fs)
if (err < 0)
goto fail;
- cle = clEnqueueNDRangeKernel(ctx->command_queue, ctx->kernel, 2, NULL,
+ cle = clEnqueueNDRangeKernel(ctx->command_queue, kernel, 2, NULL,
global_work, NULL, 0, NULL, NULL);
CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to enqueue overlay kernel "
"for plane %d: %d.\n", plane, cle);
@@ -265,6 +312,13 @@ static av_cold void overlay_opencl_uninit(AVFilterContext
*avctx)
"kernel: %d.\n", cle);
}
+ if (ctx->kernel_uv) {
+ cle = clReleaseKernel(ctx->kernel_uv);
+ if (cle != CL_SUCCESS)
+ av_log(avctx, AV_LOG_ERROR, "Failed to release "
+ "kernel: %d.\n", cle);
+ }
+
if (ctx->command_queue) {
cle = clReleaseCommandQueue(ctx->command_queue);
if (cle != CL_SUCCESS)
--
2.52.0
_______________________________________________
ffmpeg-devel mailing list -- [email protected]
To unsubscribe send an email to [email protected]
