From patchwork Wed Jun 29 19:36:00 2022 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Paul B Mahol X-Patchwork-Id: 36531 Delivered-To: ffmpegpatchwork2@gmail.com Received: by 2002:a05:6a20:8b27:b0:88:1bbf:7fd2 with SMTP id l39csp513721pzh; Wed, 29 Jun 2022 12:33:28 -0700 (PDT) X-Google-Smtp-Source: AGRyM1u3C+fR6TZ2UTHKFHD1s9NG5e+njcinEvv5zQ6NFqG5p6vKQDF5kWR1FdRyVQKEIHRbI2qU X-Received: by 2002:a05:6402:540c:b0:434:d965:f8a with SMTP id ev12-20020a056402540c00b00434d9650f8amr6503942edb.30.1656531208249; Wed, 29 Jun 2022 12:33:28 -0700 (PDT) ARC-Seal: i=1; a=rsa-sha256; t=1656531208; cv=none; d=google.com; s=arc-20160816; b=ynsHoFWbGxj6rN+75y4HVnUhUyQ8/lYubxeDw0+xnRLMWDdkuUU8MzsUu3sJYVHIEw ykWKRJkB+giPBczpnDv549wQouv/L929HvnanWIWLP9KYMji9JH42pESmeahxu2wbYRg VpNNRbziIwHLvO6GuFL9dTfSo1Dvt3X6IBnTA6E7aiiceQonPA9ny/OhhMldVwwnXWfW j102YZSagFcU8gFcJr3rZbSrG4eZR9wDrTHjzhj4Q4SVEgurZ5yRgszQsBDjta3tFiBs mwuez7xx56iY3/ttbrOT0bCMFeJUu71ckccuD5ROPjhFktBTlCQgC/si+4A+KK5IlktU k3eg== ARC-Message-Signature: i=1; a=rsa-sha256; c=relaxed/relaxed; d=google.com; s=arc-20160816; h=sender:errors-to:reply-to:list-subscribe:list-help:list-post :list-archive:list-unsubscribe:list-id:precedence:subject:to :message-id:date:from:mime-version:dkim-signature:delivered-to; bh=KbOyg3nz410ttzkLfSa/A2QmGY3kMxpIXBTtQTFQEzQ=; b=GLxbelyD4PnZudxeqoVYOgqqmbpSx9k2/dEEJL1efE1byHg/3OugA/tEWACikdfNGG X1VPH9SUtiNRW4FNcInVxjQ3zsYD8WTI+6Qa+NmeDJrqaHQ2bWlagHsF0zsiHHJX/BTM +EVW9qBS7T4gcWQtsuN/8R2yBbE1Gswr1VL93tO4AKB70G7U1VUPRlRzkefLt4DRM+Tx byT1GqHhS32kC2vu2SrGyB1fcStB53aHkAbcKx4UL96a0MGQI6W3sj2AaIKGbHADlESZ xHxsoDFbL3005hoDftEnQZurcotG5q73JgrT/j3kQBr+XUILqE1AGG27xCskanUFI117 YPvg== ARC-Authentication-Results: i=1; mx.google.com; dkim=neutral (body hash did not verify) header.i=@gmail.com header.s=20210112 header.b=FNVaFjoC; spf=pass (google.com: domain of ffmpeg-devel-bounces@ffmpeg.org designates 79.124.17.100 as permitted sender) smtp.mailfrom=ffmpeg-devel-bounces@ffmpeg.org; dmarc=fail (p=NONE sp=QUARANTINE dis=NONE) header.from=gmail.com Return-Path: Received: from ffbox0-bg.mplayerhq.hu (ffbox0-bg.ffmpeg.org. [79.124.17.100]) by mx.google.com with ESMTP id s24-20020a50ab18000000b0043574fb5ca5si22173572edc.266.2022.06.29.12.33.26; Wed, 29 Jun 2022 12:33:28 -0700 (PDT) Received-SPF: pass (google.com: domain of ffmpeg-devel-bounces@ffmpeg.org designates 79.124.17.100 as permitted sender) client-ip=79.124.17.100; Authentication-Results: mx.google.com; dkim=neutral (body hash did not verify) header.i=@gmail.com header.s=20210112 header.b=FNVaFjoC; spf=pass (google.com: domain of ffmpeg-devel-bounces@ffmpeg.org designates 79.124.17.100 as permitted sender) smtp.mailfrom=ffmpeg-devel-bounces@ffmpeg.org; dmarc=fail (p=NONE sp=QUARANTINE dis=NONE) header.from=gmail.com Received: from [127.0.1.1] (localhost [127.0.0.1]) by ffbox0-bg.mplayerhq.hu (Postfix) with ESMTP id B32C268B6A0; Wed, 29 Jun 2022 22:33:22 +0300 (EEST) X-Original-To: ffmpeg-devel@ffmpeg.org Delivered-To: ffmpeg-devel@ffmpeg.org Received: from mail-yw1-f170.google.com (mail-yw1-f170.google.com [209.85.128.170]) by ffbox0-bg.mplayerhq.hu (Postfix) with ESMTPS id 73CDC68B094 for ; Wed, 29 Jun 2022 22:33:15 +0300 (EEST) Received: by mail-yw1-f170.google.com with SMTP id 00721157ae682-31772f8495fso158346807b3.4 for ; Wed, 29 Jun 2022 12:33:15 -0700 (PDT) DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=gmail.com; s=20210112; h=mime-version:from:date:message-id:subject:to; bh=VM+yRVeaNT2TaODusAJ2ar43Dj6wll4FtG6bv6r9BrM=; b=FNVaFjoCZ4iEz7tLT1Kc+U6pBZY1b4u8spfFaTFKlKkyhS9mUfPrT+2y/jLxLk0E2S 4STTWTkJG9uYghLVDvXhmq/rrny9uVpxnTdbgkhpNefsaEk+b6H82dlO7ZccOpA7zBcZ nIveEmwNZqilMWW8Q2jkTboxxQoENpP9+l4RV5uLw7zQ7ntOe/IiVenthHuoZgbVWSO4 5ZdX/8PmzUYcPP0YqutniwEei4T21qP3tcLKepPnFhgBfeJXQh3JHV2j5S/IpFOss+Ko wWcTUsAD5czwQtbcCWa/p2HChcIXuM0A8+b7WSd4MbUrU1sHyTHUFKY2IaF0QDMQpr0E cwhw== X-Google-DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=1e100.net; s=20210112; h=x-gm-message-state:mime-version:from:date:message-id:subject:to; bh=VM+yRVeaNT2TaODusAJ2ar43Dj6wll4FtG6bv6r9BrM=; b=yY0ri91Td/FjlXXMQGVfz2QJtOHOhcjcAbCDftZHYhHJXXmmw7248RvXoNtHPlNgQO Sp/m+aZp57x6UQrFA3WgX0TEFO8iQzxXifp3p/dDexvhJKGjzugJjVKAh3AgQ6zOmWuO DFCYKyjN3jmjLl+53sxxxieO4mmcUZ5DjIB8xQ2wm17Uu1XeRgkXXEPSaF+8B5sd3xcT D5QiYQQzYA23ul2h3whQbIZmkn/uayVuBjBkWW1Wqb9I1VTxG6hAdgaRce6VIDa67Yxv n8/6nlI04gW4hcVYIzhOaXWYUS28cEKjrDgLgPO7Vs4EYXQvsYyiBkbt9wp3rdNM+B+s snug== X-Gm-Message-State: AJIora920j5BKI8gKXYIiSM8Rg9UWpLWNauk2Jh3jKCIORH7Ps/iMie2 hbovRQNc3Gm5eE/Vp+0Z5AdCZ6hwVwHaHQCt3NMX1YaOfks= X-Received: by 2002:a81:1c06:0:b0:318:27ed:8d41 with SMTP id c6-20020a811c06000000b0031827ed8d41mr5623106ywc.221.1656531194130; Wed, 29 Jun 2022 12:33:14 -0700 (PDT) MIME-Version: 1.0 From: Paul B Mahol Date: Wed, 29 Jun 2022 21:36:00 +0200 Message-ID: To: FFmpeg development discussions and patches X-Content-Filtered-By: Mailman/MimeDel 2.1.29 Subject: [FFmpeg-devel] [PATCH] avfilter: add remap_opencl filter X-BeenThere: ffmpeg-devel@ffmpeg.org X-Mailman-Version: 2.1.29 Precedence: list List-Id: FFmpeg development discussions and patches List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , Reply-To: FFmpeg development discussions and patches Errors-To: ffmpeg-devel-bounces@ffmpeg.org Sender: "ffmpeg-devel" X-TUID: 8lHg5wOIXJiA Hello, patches attached. From 011ec1b924adad0a46ff036ebed13d24bca034d9 Mon Sep 17 00:00:00 2001 From: Paul B Mahol Date: Wed, 29 Jun 2022 19:12:24 +0200 Subject: [PATCH 1/2] avfilter: add remap opencl filter Signed-off-by: Paul B Mahol --- libavfilter/Makefile | 2 + libavfilter/allfilters.c | 1 + libavfilter/opencl/remap.cl | 39 ++++ libavfilter/opencl_source.h | 1 + libavfilter/vf_remap_opencl.c | 329 ++++++++++++++++++++++++++++++++++ 5 files changed, 372 insertions(+) create mode 100644 libavfilter/opencl/remap.cl create mode 100644 libavfilter/vf_remap_opencl.c diff --git a/libavfilter/Makefile b/libavfilter/Makefile index b9ce1a715b..367eb92063 100644 --- a/libavfilter/Makefile +++ b/libavfilter/Makefile @@ -421,6 +421,8 @@ OBJS-$(CONFIG_READEIA608_FILTER) += vf_readeia608.o OBJS-$(CONFIG_READVITC_FILTER) += vf_readvitc.o OBJS-$(CONFIG_REALTIME_FILTER) += f_realtime.o OBJS-$(CONFIG_REMAP_FILTER) += vf_remap.o framesync.o +OBJS-$(CONFIG_REMAP_OPENCL_FILTER) += vf_remap_opencl.o framesync.o opencl.o \ + opencl/remap.o OBJS-$(CONFIG_REMOVEGRAIN_FILTER) += vf_removegrain.o OBJS-$(CONFIG_REMOVELOGO_FILTER) += bbox.o lswsutils.o lavfutils.o vf_removelogo.o OBJS-$(CONFIG_REPEATFIELDS_FILTER) += vf_repeatfields.o diff --git a/libavfilter/allfilters.c b/libavfilter/allfilters.c index 0152acbb81..05f0fa85db 100644 --- a/libavfilter/allfilters.c +++ b/libavfilter/allfilters.c @@ -400,6 +400,7 @@ extern const AVFilter ff_vf_readeia608; extern const AVFilter ff_vf_readvitc; extern const AVFilter ff_vf_realtime; extern const AVFilter ff_vf_remap; +extern const AVFilter ff_vf_remap_opencl; extern const AVFilter ff_vf_removegrain; extern const AVFilter ff_vf_removelogo; extern const AVFilter ff_vf_repeatfields; diff --git a/libavfilter/opencl/remap.cl b/libavfilter/opencl/remap.cl new file mode 100644 index 0000000000..8851cdc429 --- /dev/null +++ b/libavfilter/opencl/remap.cl @@ -0,0 +1,39 @@ +/* + * This file is part of FFmpeg. + * + * FFmpeg is free software; you can redistribute it and/or + * modify it under the terms of the GNU Lesser General Public + * License as published by the Free Software Foundation; either + * version 2.1 of the License, or (at your option) any later version. + * + * FFmpeg is distributed in the hope that it will be useful, + * but WITHOUT ANY WARRANTY; without even the implied warranty of + * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU + * Lesser General Public License for more details. + * + * You should have received a copy of the GNU Lesser General Public + * License along with FFmpeg; if not, write to the Free Software + * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA + */ + +const sampler_t linear_sampler = (CLK_NORMALIZED_COORDS_FALSE | + CLK_FILTER_LINEAR); + +const sampler_t nearest_sampler = (CLK_NORMALIZED_COORDS_FALSE | + CLK_FILTER_NEAREST); + +__kernel void remap(__write_only image2d_t dst, + __read_only image2d_t src, + __read_only image2d_t xmapi, + __read_only image2d_t ymapi) +{ + int2 p = (int2)(get_global_id(0), get_global_id(1)); + + float4 xmap = read_imagef(xmapi, nearest_sampler, p); + float4 ymap = read_imagef(ymapi, nearest_sampler, p); + float2 pos = (float2)(xmap.x, ymap.x); + pos.xy = pos.xy * 65535.f; + float4 val = read_imagef(src, linear_sampler, pos); + + write_imagef(dst, p, val); +} diff --git a/libavfilter/opencl_source.h b/libavfilter/opencl_source.h index 7e8133090e..9eac2dc516 100644 --- a/libavfilter/opencl_source.h +++ b/libavfilter/opencl_source.h @@ -28,6 +28,7 @@ extern const char *ff_opencl_source_neighbor; extern const char *ff_opencl_source_nlmeans; extern const char *ff_opencl_source_overlay; extern const char *ff_opencl_source_pad; +extern const char *ff_opencl_source_remap; extern const char *ff_opencl_source_tonemap; extern const char *ff_opencl_source_transpose; extern const char *ff_opencl_source_unsharp; diff --git a/libavfilter/vf_remap_opencl.c b/libavfilter/vf_remap_opencl.c new file mode 100644 index 0000000000..0282b6b4d0 --- /dev/null +++ b/libavfilter/vf_remap_opencl.c @@ -0,0 +1,329 @@ +/* + * Copyright (c) 2022 Paul B Mahol + * + * This file is part of FFmpeg. + * + * FFmpeg is free software; you can redistribute it and/or + * modify it under the terms of the GNU Lesser General Public + * License as published by the Free Software Foundation; either + * version 2.1 of the License, or (at your option) any later version. + * + * FFmpeg is distributed in the hope that it will be useful, + * but WITHOUT ANY WARRANTY; without even the implied warranty of + * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU + * Lesser General Public License for more details. + * + * You should have received a copy of the GNU Lesser General Public + * License along with FFmpeg; if not, write to the Free Software + * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301 USA + */ + +#include "libavutil/colorspace.h" +#include "libavutil/imgutils.h" +#include "libavutil/pixdesc.h" +#include "libavutil/opt.h" +#include "avfilter.h" +#include "drawutils.h" +#include "formats.h" +#include "framesync.h" +#include "internal.h" +#include "opencl.h" +#include "opencl_source.h" +#include "video.h" + +typedef struct RemapOpenCLContext { + OpenCLFilterContext ocf; + + int nb_planes; + int nb_components; + uint8_t fill_rgba[4]; + int fill_color[4]; + + int initialised; + cl_kernel kernel; + cl_command_queue command_queue; + + FFFrameSync fs; +} RemapOpenCLContext; + +#define OFFSET(x) offsetof(RemapOpenCLContext, x) +#define FLAGS AV_OPT_FLAG_FILTERING_PARAM|AV_OPT_FLAG_VIDEO_PARAM + +static const AVOption remap_opencl_options[] = { + { "fill", "set the color of the unmapped pixels", OFFSET(fill_rgba), AV_OPT_TYPE_COLOR, {.str="black"}, .flags = FLAGS }, + { NULL } +}; + +AVFILTER_DEFINE_CLASS(remap_opencl); + +static av_cold int remap_opencl_init(AVFilterContext *avctx) +{ + return ff_opencl_filter_init(avctx); +} + +static int remap_opencl_load(AVFilterContext *avctx, + enum AVPixelFormat main_format, + enum AVPixelFormat xmap_format, + enum AVPixelFormat ymap_format) +{ + RemapOpenCLContext *ctx = avctx->priv; + cl_int cle; + const char *source = ff_opencl_source_remap; + const char *kernel = "remap"; + const AVPixFmtDescriptor *main_desc, *xmap_desc, *ymap_desc; + int err, main_planes, xmap_planes, ymap_planes; + + main_desc = av_pix_fmt_desc_get(main_format); + xmap_desc = av_pix_fmt_desc_get(xmap_format); + ymap_desc = av_pix_fmt_desc_get(ymap_format); + + main_planes = xmap_planes = ymap_planes = 0; + for (int i = 0; i < main_desc->nb_components; i++) + main_planes = FFMAX(main_planes, + main_desc->comp[i].plane + 1); + for (int i = 0; i < xmap_desc->nb_components; i++) + xmap_planes = FFMAX(xmap_planes, + xmap_desc->comp[i].plane + 1); + for (int i = 0; i < ymap_desc->nb_components; i++) + ymap_planes = FFMAX(ymap_planes, + ymap_desc->comp[i].plane + 1); + + ctx->nb_planes = main_planes; + + err = ff_opencl_filter_load_program(avctx, &source, 1); + if (err < 0) + goto fail; + + ctx->command_queue = clCreateCommandQueue(ctx->ocf.hwctx->context, + ctx->ocf.hwctx->device_id, + 0, &cle); + CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create OpenCL " + "command queue %d.\n", cle); + + ctx->kernel = clCreateKernel(ctx->ocf.program, kernel, &cle); + CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to create kernel %d.\n", cle); + + ctx->initialised = 1; + return 0; + +fail: + if (ctx->command_queue) + clReleaseCommandQueue(ctx->command_queue); + if (ctx->kernel) + clReleaseKernel(ctx->kernel); + return err; +} + +static int remap_opencl_process_frame(FFFrameSync *fs) +{ + AVFilterContext *avctx = fs->parent; + AVFilterLink *outlink = avctx->outputs[0]; + RemapOpenCLContext *ctx = avctx->priv; + AVFrame *input_main, *input_xmap, *input_ymap; + AVFrame *output; + cl_mem mem; + cl_int cle; + size_t global_work[2]; + int kernel_arg = 0; + int err, plane; + + err = ff_framesync_get_frame(fs, 0, &input_main, 0); + if (err < 0) + return err; + err = ff_framesync_get_frame(fs, 1, &input_xmap, 0); + if (err < 0) + return err; + err = ff_framesync_get_frame(fs, 2, &input_ymap, 0); + if (err < 0) + return err; + + if (!ctx->initialised) { + AVHWFramesContext *main_fc = + (AVHWFramesContext*)input_main->hw_frames_ctx->data; + AVHWFramesContext *xmap_fc = + (AVHWFramesContext*)input_xmap->hw_frames_ctx->data; + AVHWFramesContext *ymap_fc = + (AVHWFramesContext*)input_ymap->hw_frames_ctx->data; + + err = remap_opencl_load(avctx, main_fc->sw_format, + xmap_fc->sw_format, + ymap_fc->sw_format); + if (err < 0) + return err; + } + + output = ff_get_video_buffer(outlink, outlink->w, outlink->h); + if (!output) { + err = AVERROR(ENOMEM); + goto fail; + } + + for (plane = 0; plane < ctx->nb_planes; plane++) { + kernel_arg = 0; + + mem = (cl_mem)output->data[plane]; + CL_SET_KERNEL_ARG(ctx->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); + kernel_arg++; + + mem = (cl_mem)input_xmap->data[0]; + CL_SET_KERNEL_ARG(ctx->kernel, kernel_arg, cl_mem, &mem); + kernel_arg++; + + mem = (cl_mem)input_ymap->data[0]; + CL_SET_KERNEL_ARG(ctx->kernel, kernel_arg, cl_mem, &mem); + kernel_arg++; + + err = ff_opencl_filter_work_size_from_image(avctx, global_work, + output, plane, 0); + if (err < 0) + goto fail; + + cle = clEnqueueNDRangeKernel(ctx->command_queue, ctx->kernel, 2, NULL, + global_work, NULL, 0, NULL, NULL); + CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to enqueue remap kernel " + "for plane %d: %d.\n", plane, cle); + } + + cle = clFinish(ctx->command_queue); + CL_FAIL_ON_ERROR(AVERROR(EIO), "Failed to finish command queue: %d.\n", cle); + + err = av_frame_copy_props(output, input_main); + + av_log(avctx, AV_LOG_DEBUG, "Filter output: %s, %ux%u (%"PRId64").\n", + av_get_pix_fmt_name(output->format), + output->width, output->height, output->pts); + + return ff_filter_frame(outlink, output); + +fail: + av_frame_free(&output); + return err; +} + +static int config_output(AVFilterLink *outlink) +{ + AVFilterContext *ctx = outlink->src; + RemapOpenCLContext *s = ctx->priv; + AVFilterLink *srclink = ctx->inputs[0]; + AVFilterLink *xlink = ctx->inputs[1]; + AVFilterLink *ylink = ctx->inputs[2]; + FFFrameSyncIn *in; + int ret; + + ret = ff_opencl_filter_config_output(outlink); + if (ret < 0) + return ret; + + if (xlink->w != ylink->w || xlink->h != ylink->h) { + av_log(ctx, AV_LOG_ERROR, "Second input link %s parameters " + "(size %dx%d) do not match the corresponding " + "third input link %s parameters (%dx%d)\n", + ctx->input_pads[1].name, xlink->w, xlink->h, + ctx->input_pads[2].name, ylink->w, ylink->h); + return AVERROR(EINVAL); + } + + outlink->w = xlink->w; + outlink->h = xlink->h; + outlink->sample_aspect_ratio = srclink->sample_aspect_ratio; + outlink->frame_rate = srclink->frame_rate; + + ret = ff_framesync_init(&s->fs, ctx, 3); + if (ret < 0) + return ret; + + in = s->fs.in; + in[0].time_base = srclink->time_base; + in[1].time_base = xlink->time_base; + in[2].time_base = ylink->time_base; + in[0].sync = 2; + in[0].before = EXT_STOP; + in[0].after = EXT_STOP; + in[1].sync = 1; + in[1].before = EXT_NULL; + in[1].after = EXT_INFINITY; + in[2].sync = 1; + in[2].before = EXT_NULL; + in[2].after = EXT_INFINITY; + s->fs.opaque = s; + s->fs.on_event = remap_opencl_process_frame; + + ret = ff_framesync_configure(&s->fs); + outlink->time_base = s->fs.time_base; + + return ret; +} + +static int activate(AVFilterContext *ctx) +{ + RemapOpenCLContext *s = ctx->priv; + return ff_framesync_activate(&s->fs); +} + +static av_cold void remap_opencl_uninit(AVFilterContext *avctx) +{ + RemapOpenCLContext *ctx = avctx->priv; + cl_int cle; + + if (ctx->kernel) { + cle = clReleaseKernel(ctx->kernel); + 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) + av_log(avctx, AV_LOG_ERROR, "Failed to release " + "command queue: %d.\n", cle); + } + + ff_opencl_filter_uninit(avctx); + + ff_framesync_uninit(&ctx->fs); +} + +static const AVFilterPad remap_opencl_inputs[] = { + { + .name = "source", + .type = AVMEDIA_TYPE_VIDEO, + .config_props = &ff_opencl_filter_config_input, + }, + { + .name = "xmap", + .type = AVMEDIA_TYPE_VIDEO, + .config_props = &ff_opencl_filter_config_input, + }, + { + .name = "ymap", + .type = AVMEDIA_TYPE_VIDEO, + .config_props = &ff_opencl_filter_config_input, + }, +}; + +static const AVFilterPad remap_opencl_outputs[] = { + { + .name = "default", + .type = AVMEDIA_TYPE_VIDEO, + .config_props = config_output, + }, +}; + +const AVFilter ff_vf_remap_opencl = { + .name = "remap_opencl", + .description = NULL_IF_CONFIG_SMALL("Remap pixels using OpenCL."), + .priv_size = sizeof(RemapOpenCLContext), + .init = remap_opencl_init, + .uninit = remap_opencl_uninit, + .activate = activate, + FILTER_INPUTS(remap_opencl_inputs), + FILTER_OUTPUTS(remap_opencl_outputs), + FILTER_SINGLE_PIXFMT(AV_PIX_FMT_OPENCL), + .priv_class = &remap_opencl_class, + .flags_internal = FF_FILTER_FLAG_HWFRAME_AWARE, +}; -- 2.36.1