On 26/01/2020 18:28, Paul B Mahol wrote:
> Signed-off-by: Paul B Mahol <one...@gmail.com>
> ---
> configure | 1 +
> doc/filters.texi | 97 ++++++++
> libavfilter/Makefile | 1 +
> libavfilter/allfilters.c | 1 +
> libavfilter/opencl/xfade.cl | 136 +++++++++++
> libavfilter/opencl_source.h | 1 +
> libavfilter/vf_xfade_opencl.c | 420 ++++++++++++++++++++++++++++++++++
> 7 files changed, 657 insertions(+)
> create mode 100644 libavfilter/opencl/xfade.cl
> create mode 100644 libavfilter/vf_xfade_opencl.c
>
> ...
> +
> +void slide(__write_only image2d_t dst,
> + __read_only image2d_t src1,
> + __read_only image2d_t src2,
> + float progress,
> + int2 direction)
> +{
> + const sampler_t sampler = (CLK_NORMALIZED_COORDS_FALSE | \
> + CLK_FILTER_NEAREST);
From the Mali driver:
<source>:87:21: error: non-kernel function variable cannot be declared in
constant address space
const sampler_t sampler = (CLK_NORMALIZED_COORDS_FALSE | \
^
error: Compiler frontend failed (error code 60)
OpenCL 1.2 ยง6.9:
"The sampler type (sampler_t) can only be used as the type of a function
argument or a
variable declared in the program scope or the outermost scope of a kernel
function."
I think just make it global, since you're now using the same sampler in every
kernel?
> + int w = get_image_dim(src1).x;
> + int h = get_image_dim(src1).y;
> + int2 wh = (int2)(w, h);
> + int2 uv = (int2)(get_global_id(0), get_global_id(1));
> + int2 pi = (int2)(progress * w, progress * h);
> + int2 p = uv + pi * direction;
> + int2 f = p % wh;
> +
> + f = f + (int2)(w, h) * (int2)(f.x < 0, f.y < 0);
> + float4 val1 = read_imagef(src1, sampler, f);
> + float4 val2 = read_imagef(src2, sampler, f);
> + write_imagef(dst, uv, mix(val1, val2, (p.y >= 0) * (h > p.y) * (p.x >=
> 0) * (w > p.x)));
> +}
> +
> ...
> +
> +AVFilter ff_vf_xfade_opencl = {
> + .name = "xfade_opencl",
> + .description = NULL_IF_CONFIG_SMALL("Cross fade one video with
> another video."),
> + .priv_size = sizeof(XFadeOpenCLContext),
> + .priv_class = &xfade_opencl_class,
> + .init = &ff_opencl_filter_init,
> + .uninit = &xfade_opencl_uninit,
> + .query_formats = &ff_opencl_filter_query_formats,
> + .activate = &xfade_opencl_activate,
> + .inputs = xfade_opencl_inputs,
> + .outputs = xfade_opencl_outputs,
> + .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE,
> +};
>
No other comments from me, so LGTM with that fixed.
Thanks,
- Mark
_______________________________________________
ffmpeg-devel mailing list
ffmpeg-devel@ffmpeg.org
https://ffmpeg.org/mailman/listinfo/ffmpeg-devel
To unsubscribe, visit link above, or email
ffmpeg-devel-requ...@ffmpeg.org with subject "unsubscribe".
