From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: from ffbox0-bg.ffmpeg.org (ffbox0-bg.ffmpeg.org [79.124.17.100]) by master.gitmailbox.com (Postfix) with ESMTPS id 75B254D237 for ; Mon, 30 Jun 2025 14:43:10 +0000 (UTC) Received: from [127.0.1.1] (localhost [127.0.0.1]) by ffbox0-bg.ffmpeg.org (Postfix) with ESMTP id 4E3F368E289; Mon, 30 Jun 2025 17:43:06 +0300 (EEST) Date: Mon, 30 Jun 2025 17:42:20 +0300 To: ffmpeg-devel@ffmpeg.org In-Reply-To: <20250501170711.60035-1-jendas1@yahoo.com> References: <20250501170711.60035-1-jendas1.ref@yahoo.com> <20250501170711.60035-1-jendas1@yahoo.com> MIME-Version: 1.0 Message-ID: List-Id: FFmpeg development discussions and patches List-Post: From: =?utf-8?q?Jan_Studen=C3=BD_via_ffmpeg-devel?= Precedence: list Cc: =?utf-8?Q?Jan_Studen=C3=BD?= X-Mailman-Version: 2.1.29 X-BeenThere: ffmpeg-devel@ffmpeg.org List-Subscribe: , List-Unsubscribe: , List-Archive: Reply-To: FFmpeg development discussions and patches List-Help: Subject: Re: [FFmpeg-devel] [PATCH v2] avfilter/vf_lut3d_opencl Initial support for OpenCL implementation of vf_lut3d. Content-Type: multipart/mixed; boundary="===============4016265456810750075==" Errors-To: ffmpeg-devel-bounces@ffmpeg.org Sender: "ffmpeg-devel" Archived-At: List-Archive: List-Post: --===============4016265456810750075== Content-Type: message/rfc822 Content-Disposition: inline Return-Path: X-Original-To: ffmpeg-devel@ffmpeg.org Delivered-To: ffmpeg-devel@ffmpeg.org Received: from sonic316-11.consmr.mail.bf2.yahoo.com (sonic316-11.consmr.mail.bf2.yahoo.com [74.6.130.121]) by ffbox0-bg.ffmpeg.org (Postfix) with ESMTPS id 3DFA768E1A8 for ; Mon, 30 Jun 2025 17:42:59 +0300 (EEST) DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=yahoo.com; s=s2048; t=1751294577; bh=OugzYcg06oc9HkpCH/RX3qCo63FdQD56TZcV0M7PVXw=; h=Date:From:To:In-Reply-To:References:Subject:From:Subject:Reply-To; b=UKlliqpgFvDm0OdsO9A0q8WubaHqPPhSc8kmtXtXwOF6gLr//kh/gYUtKVroJUqO779mS3a41Q4/v1Q3DplYjSYBpio8yBWSXFvbGo9ytMG/xNvOCWgFJes63khCL+FvHeGNH3Ln51n7A33Nw75Yoa7fTuOr16oUltiy9iMFAUwUBJU9YKgqFQ1FpYRrLs5Yu9css4wGbme6plpwCTylUmy+xNKoqunSmEA3Y09d84TGtmJDniZI/HQF3uuNChI7ftePNmkzKr1qkMp4fkNLJsyguu8G7+vg11NBd07gXFKW0Mj4IdszQX20DzAR1DZfFE87uEplUIEFfINmR1964Q== X-SONIC-DKIM-SIGN: v=1; a=rsa-sha256; c=relaxed/relaxed; d=yahoo.com; s=s2048; t=1751294577; bh=hWinvp2xmTCz6mFbmKcI4/1w6YsQ4kLcnKp+jMQFuJE=; h=X-Sonic-MF:Date:From:To:Subject:From:Subject; b=awLGqMBiie3k5flm8gzE8D5lE7ghJfH6srCrIyues4gqotd29jId9/K1u4kGOcvMc5BpCaxj3XEHwSIl0ca+gP7RfV7iTUQptPSQOD8oFXNLv0QflbuEFPCxaZW/ccp3xEQu6lNk9rsNvUpPSe3UA9Scl2kkcG7OuRdXQSjohokHIkMsxYq9mPWJeukmL7jAibShry/kKuIq84EgudMCebqvFPm+vWCbR0dMXVmd4og+wyNfG52YTJe6UF0ezSy3iYY+wmU/C181Ln31jP7VLtCguCQS26HAHDKdkH75s5L9b9JUsCzphIjIrrkE1mEPYffz+V0OxtTLtDydbkwXog== X-YMail-OSG: VTbAEaYVM1k0oEiZtEm6eg_415jKkfyjuCg4slaVzVIEbl3_EMBIJvVXVuc2Q6f Do9FCHCtvHSg8x5xfsgwlJf9Nx7BmVQq824D8LrbJ280yyXNtL_QDcJWld3dGebqsUPLA7J8Jnns g0pJuKkq1g_RwC5y76BPo6xM.ienrJljwHn5Sg.saaMD5ozJn4wxBVAM9t3TzjJizCEVnt4F8N2k PXQGaIEUTsZN.ywTbWSu11sWGp53_i9.bbWACWY8Go40KmYB8SHy.T12QRP6USKXG4BJpNfHxAgn NQHQ.y5DHWrx.zbyoPwpBQ2DRy8lcVa8jk62C.arnPZz6miMDc2P4dAtD0th_MIqB5V7e_zRjnvd dQ8r97xBya99yFq7hM9EdTvB9ZyQ0PVEws04SPtQBiyJAtub8XpJfO4kg.WmJfl8gH4De7sh4SZG HniCPtaU7eRkJF34D6c4WWEoiEAIZ.yu2KaaxGsRyn6L9D3Bj8XzXgBd4KDjdkEIqifqVLdLz7Xb uURctxL9BDcczAoTGxYLWhvWBNpo7kPLcSISZuo71pPlJqSaWbUzUEolqfME.lxA3g4R.BEGO9z5 KrD0Wndjju1IAtWX1IZI5dlcWCk7AKvVPzPD4D5Eiky49DSkFPaafoZTssuOQ2uKrG3ZUzhnI7i5 hGg5o6cuALLCKUACUZK2bI_126P1ZRespRV3y6MZxZFXMPNY6QeE23AbyfIBkwCVk5N9.Qzxi25a PiON9UmAutYzX.G4QXLV05yZTOSsvkMELtag6TtoxJvi2rukHDeqxYUY33GRAMEkt3XajaBWXuk9 K4_FwMvrTWDz8ymKKIBNIv5ZYmNYFv3Xj9OYvDVvdSR9fEwcL.r8Gx30G6rTesrwe27wMjRDxFKA KgJHycpkuIH0hNwnL9YCKqSFtDX2pHmrY_WSwwViB7JZj2cTJvJdnab2R1GCI02GPBX6wKySnt2P fLfR7nP1DsRnKgeiLIhI6..32nyosBQMySUntBJNtGNgWpiyUOdZ5uJX8ayR6f_rLTcjdw2Unn4N YC6pzX5RQ4hF6mSUPAClWO4Tx9od9CECVjys4YXnwfx1uMllaWpe88V8.wh9IsCRk7NRubhGsOFg qmJuiiMLEwnuepmjXhBN4hVb5cAaQ60P9aBzGK59Zx20eQrqv4HACHwKNGMhLaQruSTZgDF54hyq QM6kZLWzERmIjmiSPi2qkMEeZfLZWiP2qbkX1Na4tq2GKSqrDL6OUKOT4WpLgdG.kL1dgCSHTChW 2x65d5XYIDcTXzcOr0f8zUSw9i6tbdThgooIzPjLQz.ITqVkVicuWt5B3it4gIxavuGlYxxwnX3j wArGsIX8fWd6pGzEDO.aqjyDX3AIXWea.F1sh.Dul2p.pyBA9OSMKbFRKHURXi65bVrgS3mBo3k5 KV8_2J8In10zys0mbp0jTmU6usUKebYT7f5ZcHKimggyfO1kJOxLLKvAsiLvw8cJ7UuOt5zN9yJM .0XcD.GZ_bzkXSN_snnV9MKegMsu_bwZy5rg3vJ9eSpkZQ2OdrKr5KYGpCUSoJ9JbblFYh3uxJ6j ChWwO7WxxlOA0iQQ1bw.VpmU8Q_Tyr8NGFT1I5vGUT8jPMptmyCCgeALyQtyhseBs_wuwR_o2iMG 8jqx8Ua13zNEez4tYUATUdo4MmE5TeFPAP7oGJkUYThrCb8X3aszXQP9jhrZrsHr9Q3GUVa3895i 0GCls3MW_k1EoLJd4N2JP.EYAu6aqhTVJEGo2.eQlRrW0jmMwl8e4Q0kigrma8vka3zyXm9gc16W EmA9Zjzgy1MXJtjKYE.r24IfDY9NzsPnMvloEElTlE3jETn5T5F13rqxjkVx3.KptVdKHjIs6Olu ebPvPNTws8_uEsWAW6wgA196.6YXNd9OYzfqlggmTKMwTuSSUjrcvly9ERDJsCd2Ly477tAAEAVH Ptv1aib32wwkIQmu_CKZEk2TmKNYTCWtCJP1ORxkLVzZrs.BTA1gKnRhWOWF9ZSG8tqkd6h5xFfJ 63mF4Hq.I9vTDxNn2iUj_g.eGVImY.pGs5IhYhRm1sE0ikJNBcDawe..xQJy2eQiyXDdl4wz_mQ7 IKq8kPyWSKodDHuXAmDm4DBefh0WEAUH6i2HXqKEQxG.GIUU0siaYIpWij69z0xNIDV9ATIDsp3k cpgiyEWzfGpabLjcgBOFq10Ws99qzUDmP..Q63n9bNUyE_5GgQSrrPn.D.dU7uWYqidKJuqpxXkj 8KHx8ouoHapwpmQ-- X-Sonic-MF: X-Sonic-ID: cd334d46-79a8-441d-af1e-2fb2a0cb26fe Received: from sonic.gate.mail.ne1.yahoo.com by sonic316.consmr.mail.bf2.yahoo.com with HTTP; Mon, 30 Jun 2025 14:42:57 +0000 Received: by hermes--production-ir2-858bd4ff7b-nknbf (Yahoo Inc. Hermes SMTP Server) with ESMTPA ID 76348c873d8ba739513b1f2e74737fbf; Mon, 30 Jun 2025 14:42:55 +0000 (UTC) Date: Mon, 30 Jun 2025 17:42:20 +0300 From: =?utf-8?Q?Jan_Studen=C3=BD?= To: ffmpeg-devel@ffmpeg.org Message-ID: In-Reply-To: <20250501170711.60035-1-jendas1@yahoo.com> References: <20250501170711.60035-1-jendas1.ref@yahoo.com> <20250501170711.60035-1-jendas1@yahoo.com> Subject: Re: [PATCH v2] avfilter/vf_lut3d_opencl Initial support for OpenCL implementation of vf_lut3d. X-Readdle-Message-ID: d3d6a0f9-ae88-4b26-8f1f-f41096c8fe0d@Spark MIME-Version: 1.0 X-Mailer: WebService/1.1.24076 mail.backend.jedi.jws.acl:role.jedi.acl.token.atz.jws.hermes.yahoo Content-Type: text/plain; charset="utf-8" Content-Transfer-Encoding: quoted-printable Content-Disposition: inline X-Content-Filtered-By: Mailman/MimeDel 2.1.29 Hello =46=46mpeg maintainers, I submitted a patch nearly two months ago that adds OpenCL=E2=80=90based = acceleration to the lut3d filter (see below for details). Version 2 haven= =E2=80=99t received any feedback, so I=E2=80=99m sending this reminder as= I=E2=80=99d really appreciate someone taking a look at it. Regards, Jan Studen=C3=BD On May 1, 2025 at 20:11 +0300, Jan Studen=C3=BD , wr= ote: > The comile error is fixed by adding opencl dependency to configure. > > --- > configure =7C 1 + > libavfilter/Makefile =7C 1 + > libavfilter/allfilters.c =7C 1 + > libavfilter/opencl/lut3d.cl =7C 177 +++++++++++++ > libavfilter/opencl=5Fsource.h =7C 2 + > libavfilter/vf=5Flut3d=5Fopencl.c =7C 460 +++++++++++++++++++++++++++++= +++++ > 6 files changed, 642 insertions(+) > create mode 100644 libavfilter/opencl/lut3d.cl > create mode 100644 libavfilter/vf=5Flut3d=5Fopencl.c > > diff --git a/configure b/configure > index ee270b770c..9b9ea3b39b 100755 > --- a/configure > +++ b/configure > =40=40 -3934,6 +3934,7 =40=40 ladspa=5Ffilter=5Fdeps=3D=22ladspa libdl=22= > lcevc=5Ffilter=5Fdeps=3D=22liblcevc=5Fdec=22 > lensfun=5Ffilter=5Fdeps=3D=22liblensfun version3=22 > libplacebo=5Ffilter=5Fdeps=3D=22libplacebo vulkan=22 > +lut3d=5Fopencl=5Ffilter=5Fdeps=3D=22opencl=22 > lv2=5Ffilter=5Fdeps=3D=22lv2=22 > mcdeint=5Ffilter=5Fdeps=3D=22avcodec gpl=22 > metadata=5Ffilter=5Fdeps=3D=22avformat=22 > diff --git a/libavfilter/Makefile b/libavfilter/Makefile > index 7c0d879ec9..6524d0f91a 100644 > --- a/libavfilter/Makefile > +++ b/libavfilter/Makefile > =40=40 -378,6 +378,7 =40=40 OBJS-=24(CON=46IG=5FLUT1D=5F=46ILTER) +=3D = vf=5Flut3d.o > OBJS-=24(CON=46IG=5FLUT=5F=46ILTER) +=3D vf=5Flut.o > OBJS-=24(CON=46IG=5FLUT2=5F=46ILTER) +=3D vf=5Flut2.o framesync.o > OBJS-=24(CON=46IG=5FLUT3D=5F=46ILTER) +=3D vf=5Flut3d.o framesync.o > +OBJS-=24(CON=46IG=5FLUT3D=5FOPENCL=5F=46ILTER) +=3D vf=5Flut3d=5Fopenc= l.o opencl.o opencl/lut3d.o > OBJS-=24(CON=46IG=5FLUTRGB=5F=46ILTER) +=3D vf=5Flut.o > OBJS-=24(CON=46IG=5FLUTYUV=5F=46ILTER) +=3D vf=5Flut.o > OBJS-=24(CON=46IG=5FMASKEDCLAMP=5F=46ILTER) +=3D vf=5Fmaskedclamp.o fra= mesync.o > diff --git a/libavfilter/allfilters.c b/libavfilter/allfilters.c > index 740d9ab265..72c2f48ac4 100644 > --- a/libavfilter/allfilters.c > +++ b/libavfilter/allfilters.c > =40=40 -353,6 +353,7 =40=40 extern const =46=46=46ilter ff=5Fvf=5Flut; > extern const =46=46=46ilter ff=5Fvf=5Flut1d; > extern const =46=46=46ilter ff=5Fvf=5Flut2; > extern const =46=46=46ilter ff=5Fvf=5Flut3d; > +extern const =46=46=46ilter ff=5Fvf=5Flut3d=5Fopencl; > extern const =46=46=46ilter ff=5Fvf=5Flutrgb; > extern const =46=46=46ilter ff=5Fvf=5Flutyuv; > extern const =46=46=46ilter ff=5Fvf=5Fmaskedclamp; > diff --git a/libavfilter/opencl/lut3d.cl b/libavfilter/opencl/lut3d.cl > new file mode 100644 > index 0000000000..16dfecdc4e > --- /dev/null > +++ b/libavfilter/opencl/lut3d.cl > =40=40 -0,0 +1,177 =40=40 > +/* > + * Copyright (c) 2025 Jan Studeny > + * > + * This file is part of =46=46mpeg. > + * > + * =46=46mpeg 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 =46ree Software =46oundation; either > + * version 2.1 of the License, or (at your option) any later version. > + * > + * =46=46mpeg is distributed in the hope that it will be useful, > + * but WITHOUT ANY WARRANTY; without even the implied warranty of > + * MERCHANTABILITY or =46ITNESS =46OR A PARTICULAR PURPOSE. See the GN= U > + * Lesser General Public License for more details. > + * > + * You should have received a copy of the GNU Lesser General Public > + * License along with =46=46mpeg; if not, write to the =46ree Software= > + * =46oundation, Inc., 51 =46ranklin Street, =46ifth =46loor, Boston, = MA 02110-1301 USA > + */ > + > +typedef struct rgbvec =7B > + float r, g, b, a; > +=7D rgbvec; > + > +=23define MIN(X, Y) (((X) < (Y)) =3F (X) : (Y)) > + > +=23define NEAR(x) ((int)((x) + .5)) > +=23define PREV(x) ((int)(x)) > +=23define NEXT(x) (MIN((int)(x) + 1, lut=5Fedge=5Fsize - 1)) > + > +/** > + * Get the nearest defined point > + */ > +static rgbvec interp=5Fnearest(float4 px, =5F=5Fglobal const rgbvec *l= ut, int lut=5Fedge=5Fsize) > +=7B > + int r =3D NEAR(px=5B0=5D); > + int g =3D NEAR(px=5B1=5D); > + int b =3D NEAR(px=5B2=5D); > + int index =3D r * lut=5Fedge=5Fsize * lut=5Fedge=5Fsize + g * lut=5Fe= dge=5Fsize + b; > + return lut=5Bindex=5D; > +=7D > + > +static float lerpf(float v0, float v1, float f) > +=7B > + return v0 + (v1 - v0) * f; > +=7D > + > +static rgbvec lerp(const rgbvec *v0, const rgbvec *v1, float f) > +=7B > + rgbvec v =3D =7B > + lerpf(v0->r, v1->r, f), lerpf(v0->g, v1->g, f), lerpf(v0->b, v1->b, f= ) > + =7D; > + return v; > +=7D > +/** > + * Interpolate using the 8 vertices of a cube > + * =40see https://en.wikipedia.org/wiki/Trilinear=5Finterpolation > + */ > +static rgbvec interp=5Ftrilinear(float4 px, =5F=5Fglobal const rgbvec = *lut, int lut=5Fedge=5Fsize) > +=7B > + const int lutsize2 =3D lut=5Fedge=5Fsize * lut=5Fedge=5Fsize; > + const int lutsize =3D lut=5Fedge=5Fsize; > + > + const int prev=5B=5D =3D =7B PREV(px=5B0=5D), PREV(px=5B1=5D), PREV(p= x=5B2=5D) =7D; > + const int next=5B=5D =3D =7B NEXT(px=5B0=5D), NEXT(px=5B1=5D), NEXT(p= x=5B2=5D) =7D; > + > + const rgbvec d =3D =7B > + px=5B0=5D - prev=5B0=5D, > + px=5B1=5D - prev=5B1=5D, > + px=5B2=5D - prev=5B2=5D > + =7D; > + > + const rgbvec c000 =3D lut=5Bprev=5B0=5D * lutsize2 + prev=5B1=5D * lu= tsize + prev=5B2=5D=5D; > + const rgbvec c001 =3D lut=5Bprev=5B0=5D * lutsize2 + prev=5B1=5D * lu= tsize + next=5B2=5D=5D; > + const rgbvec c010 =3D lut=5Bprev=5B0=5D * lutsize2 + next=5B1=5D * lu= tsize + prev=5B2=5D=5D; > + const rgbvec c011 =3D lut=5Bprev=5B0=5D * lutsize2 + next=5B1=5D * lu= tsize + next=5B2=5D=5D; > + const rgbvec c100 =3D lut=5Bnext=5B0=5D * lutsize2 + prev=5B1=5D * lu= tsize + prev=5B2=5D=5D; > + const rgbvec c101 =3D lut=5Bnext=5B0=5D * lutsize2 + prev=5B1=5D * lu= tsize + next=5B2=5D=5D; > + const rgbvec c110 =3D lut=5Bnext=5B0=5D * lutsize2 + next=5B1=5D * lu= tsize + prev=5B2=5D=5D; > + const rgbvec c111 =3D lut=5Bnext=5B0=5D * lutsize2 + next=5B1=5D * lu= tsize + next=5B2=5D=5D; > + > + const rgbvec c00 =3D lerp(&c000, &c100, d.r); > + const rgbvec c10 =3D lerp(&c010, &c110, d.r); > + const rgbvec c01 =3D lerp(&c001, &c101, d.r); > + const rgbvec c11 =3D lerp(&c011, &c111, d.r); > + > + const rgbvec c0 =3D lerp(&c00, &c10, d.g); > + const rgbvec c1 =3D lerp(&c01, &c11, d.g); > + > + return lerp(&c0, &c1, d.b); > +=7D > + > +/** > + * Tetrahedral interpolation. Based on code found in Truelight Softwar= e Library paper. > + * =40see http://www.filmlight.ltd.uk/pdf/whitepapers/=46L-TL-TN-0057-= SoftwareLib.pdf > + */ > + > +static rgbvec interp=5Ftetrahedral(float4 px, =5F=5Fglobal const rgbve= c *lut, int lut=5Fedge=5Fsize) > +=7B > + const int lutsize2 =3D lut=5Fedge=5Fsize*lut=5Fedge=5Fsize; > + const int lutsize =3D lut=5Fedge=5Fsize; > + const int prev=5B=5D =3D =7BPREV(px=5B0=5D), PREV(px=5B1=5D), PREV(px= =5B2=5D)=7D; > + const int next=5B=5D =3D =7BNEXT(px=5B0=5D), NEXT(px=5B1=5D), NEXT(px= =5B2=5D)=7D; > + const rgbvec d =3D =7Bpx=5B0=5D - prev=5B0=5D, px=5B1=5D - prev=5B1=5D= , px=5B2=5D - prev=5B2=5D=7D; > + const rgbvec c000 =3D lut=5Bprev=5B0=5D * lutsize2 + prev=5B1=5D * lu= tsize + prev=5B2=5D=5D; > + const rgbvec c111 =3D lut=5Bnext=5B0=5D * lutsize2 + next=5B1=5D * lu= tsize + next=5B2=5D=5D; > + rgbvec c; > + if (d.r > d.g) =7B > + if (d.g > d.b) =7B > + const rgbvec c100 =3D lut=5Bnext=5B0=5D * lutsize2 + prev=5B1=5D * lu= tsize + prev=5B2=5D=5D; > + const rgbvec c110 =3D lut=5Bnext=5B0=5D * lutsize2 + next=5B1=5D * lu= tsize + prev=5B2=5D=5D; > + c.r =3D (1-d.r) * c000.r + (d.r-d.g) * c100.r + (d.g-d.b) * c110.r + = (d.b) * c111.r; > + c.g =3D (1-d.r) * c000.g + (d.r-d.g) * c100.g + (d.g-d.b) * c110.g + = (d.b) * c111.g; > + c.b =3D (1-d.r) * c000.b + (d.r-d.g) * c100.b + (d.g-d.b) * c110.b + = (d.b) * c111.b; > + =7D else if (d.r > d.b) =7B > + const rgbvec c100 =3D lut=5Bnext=5B0=5D * lutsize2 + prev=5B1=5D * lu= tsize + prev=5B2=5D=5D; > + const rgbvec c101 =3D lut=5Bnext=5B0=5D * lutsize2 + prev=5B1=5D * lu= tsize + next=5B2=5D=5D; > + c.r =3D (1-d.r) * c000.r + (d.r-d.b) * c100.r + (d.b-d.g) * c101.r + = (d.g) * c111.r; > + c.g =3D (1-d.r) * c000.g + (d.r-d.b) * c100.g + (d.b-d.g) * c101.g + = (d.g) * c111.g; > + c.b =3D (1-d.r) * c000.b + (d.r-d.b) * c100.b + (d.b-d.g) * c101.b + = (d.g) * c111.b; > + =7D else =7B > + const rgbvec c001 =3D lut=5Bprev=5B0=5D * lutsize2 + prev=5B1=5D * lu= tsize + next=5B2=5D=5D; > + const rgbvec c101 =3D lut=5Bnext=5B0=5D * lutsize2 + prev=5B1=5D * lu= tsize + next=5B2=5D=5D; > + c.r =3D (1-d.b) * c000.r + (d.b-d.r) * c001.r + (d.r-d.g) * c101.r + = (d.g) * c111.r; > + c.g =3D (1-d.b) * c000.g + (d.b-d.r) * c001.g + (d.r-d.g) * c101.g + = (d.g) * c111.g; > + c.b =3D (1-d.b) * c000.b + (d.b-d.r) * c001.b + (d.r-d.g) * c101.b + = (d.g) * c111.b; > + =7D > + =7D else =7B > + if (d.b > d.g) =7B > + const rgbvec c001 =3D lut=5Bprev=5B0=5D * lutsize2 + prev=5B1=5D * lu= tsize + next=5B2=5D=5D; > + const rgbvec c011 =3D lut=5Bprev=5B0=5D * lutsize2 + next=5B1=5D * lu= tsize + next=5B2=5D=5D; > + c.r =3D (1-d.b) * c000.r + (d.b-d.g) * c001.r + (d.g-d.r) * c011.r + = (d.r) * c111.r; > + c.g =3D (1-d.b) * c000.g + (d.b-d.g) * c001.g + (d.g-d.r) * c011.g + = (d.r) * c111.g; > + c.b =3D (1-d.b) * c000.b + (d.b-d.g) * c001.b + (d.g-d.r) * c011.b + = (d.r) * c111.b; > + =7D else if (d.b > d.r) =7B > + const rgbvec c010 =3D lut=5Bprev=5B0=5D * lutsize2 + next=5B1=5D * lu= tsize + prev=5B2=5D=5D; > + const rgbvec c011 =3D lut=5Bprev=5B0=5D * lutsize2 + next=5B1=5D * lu= tsize + next=5B2=5D=5D; > + c.r =3D (1-d.g) * c000.r + (d.g-d.b) * c010.r + (d.b-d.r) * c011.r + = (d.r) * c111.r; > + c.g =3D (1-d.g) * c000.g + (d.g-d.b) * c010.g + (d.b-d.r) * c011.g + = (d.r) * c111.g; > + c.b =3D (1-d.g) * c000.b + (d.g-d.b) * c010.b + (d.b-d.r) * c011.b + = (d.r) * c111.b; > + =7D else =7B > + const rgbvec c010 =3D lut=5Bprev=5B0=5D * lutsize2 + next=5B1=5D * lu= tsize + prev=5B2=5D=5D; > + const rgbvec c110 =3D lut=5Bnext=5B0=5D * lutsize2 + next=5B1=5D * lu= tsize + prev=5B2=5D=5D; > + c.r =3D (1-d.g) * c000.r + (d.g-d.r) * c010.r + (d.r-d.b) * c110.r + = (d.b) * c111.r; > + c.g =3D (1-d.g) * c000.g + (d.g-d.r) * c010.g + (d.r-d.b) * c110.g + = (d.b) * c111.g; > + c.b =3D (1-d.g) * c000.b + (d.g-d.r) * c010.b + (d.r-d.b) * c110.b + = (d.b) * c111.b; > + =7D > + =7D > + return c; > +=7D > + > +=23define LUT3D=5FKERNEL(INTERP=5F=46UNC) =5C > +=5F=5Fkernel void lut3d=5F=23=23INTERP=5F=46UNC( =5C > + =5F=5Fread=5Fonly image2d=5Ft src, =5C > + =5F=5Fwrite=5Fonly image2d=5Ft dst, =5C > + =5F=5Fglobal const rgbvec* lut, =5C > + int lut=5Fedge=5Fsize) =5C > +=7B =5C > + const sampler=5Ft sampler =3D (CLK=5FNORMALIZED=5FCOORDS=5F=46ALSE =7C= =5C > + CLK=5FADDRESS=5FCLAMP=5FTO=5FEDGE =7C =5C > + CLK=5F=46ILTER=5FNEAREST); =5C > + =5C > + int2 loc =3D (int2)(get=5Fglobal=5Fid(0), get=5Fglobal=5Fid(1)); =5C > + float4 px =3D read=5Fimagef(src, sampler, loc); =5C > + =5C > + for (int i =3D 0; i < 3; i++) =7B =5C > + px=5Bi=5D *=3D (lut=5Fedge=5Fsize - 1); =5C > + =7D =5C > + =5C > + rgbvec lutpx =3D INTERP=5F=46UNC(px, lut, lut=5Fedge=5Fsize); =5C > + =5C > + write=5Fimagef(dst, loc, (float4)(lutpx.r, lutpx.g, lutpx.b, 0.0f)); = =5C > +=7D > + > +LUT3D=5FKERNEL(interp=5Fnearest) > +LUT3D=5FKERNEL(interp=5Ftrilinear) > +LUT3D=5FKERNEL(interp=5Ftetrahedral) > diff --git a/libavfilter/opencl=5Fsource.h b/libavfilter/opencl=5Fsourc= e.h > index b6930fb686..d143286d21 100644 > --- a/libavfilter/opencl=5Fsource.h > +++ b/libavfilter/opencl=5Fsource.h > =40=40 -26,6 +26,7 =40=40 extern const char *ff=5Fsource=5Fconvolution=5F= cl; > extern const char *ff=5Fsource=5Fdeshake=5Fcl; > extern const char *ff=5Fsource=5Fneighbor=5Fcl; > extern const char *ff=5Fsource=5Fnlmeans=5Fcl; > +extern const char *ff=5Fsource=5Flut3d=5Fcl; > extern const char *ff=5Fsource=5Foverlay=5Fcl; > extern const char *ff=5Fsource=5Fpad=5Fcl; > extern const char *ff=5Fsource=5Fremap=5Fcl; > =40=40 -34,4 +35,5 =40=40 extern const char *ff=5Fsource=5Ftranspose=5F= cl; > extern const char *ff=5Fsource=5Funsharp=5Fcl; > extern const char *ff=5Fsource=5Fxfade=5Fcl; > > + > =23endif /* AV=46ILTER=5FOPENCL=5FSOURCE=5FH */ > diff --git a/libavfilter/vf=5Flut3d=5Fopencl.c b/libavfilter/vf=5Flut3d= =5Fopencl.c > new file mode 100644 > index 0000000000..bb7d10ed37 > --- /dev/null > +++ b/libavfilter/vf=5Flut3d=5Fopencl.c > =40=40 -0,0 +1,460 =40=40 > +/* > + * Copyright (c) 2025 Jan Studeny > + * > + * This file is part of =46=46mpeg. > + * > + * =46=46mpeg 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 =46ree Software =46oundation; either > + * version 2.1 of the License, or (at your option) any later version. > + * > + * =46=46mpeg is distributed in the hope that it will be useful, > + * but WITHOUT ANY WARRANTY; without even the implied warranty of > + * MERCHANTABILITY or =46ITNESS =46OR A PARTICULAR PURPOSE. See the GN= U > + * Lesser General Public License for more details. > + * > + * You should have received a copy of the GNU Lesser General Public > + * License along with =46=46mpeg; if not, write to the =46ree Software= > + * =46oundation, Inc., 51 =46ranklin Street, =46ifth =46loor, Boston, = MA 02110-1301 USA > + */ > + > +=23include =22config=5Fcomponents.h=22 > + > +=23include =22libavutil/avassert.h=22 > +=23include =22libavutil/common.h=22 > +=23include =22libavutil/imgutils.h=22 > +=23include =22libavutil/mem.h=22 > +=23include =22libavutil/opt.h=22 > +=23include =22libavutil/pixdesc.h=22 > +=23include =22libavutil/avstring.h=22 > + > +=23include =22libavutil/file=5Fopen.h=22 > + > +=23include =22avfilter.h=22 > +=23include =22filters.h=22 > +=23include =22opencl.h=22 > +=23include =22drawutils.h=22 > +=23include =22opencl=5Fsource.h=22 > +=23include =22video.h=22 > + > + > +=23define MAX=5FLINE=5FSIZE 512 > + > +enum interp=5Fmode =7B > + INTERPOLATE=5FNEAREST, > + INTERPOLATE=5FTRILINEAR, > + INTERPOLATE=5FTETRAHEDRAL, > + INTERPOLATE=5FPYRAMID, > + INTERPOLATE=5FPRISM, > + NB=5FINTERP=5FMODE > +=7D; > + > +typedef struct rgbvec =7B > + cl=5Ffloat r, g, b, a; > +=7D rgbvec; > + > +=23define MAX=5FLEVEL 256 > + > + > +typedef struct LUT3DOpenCLContext =7B > + OpenCL=46ilterContext ocf; > + > + int initialised; > + cl=5Fkernel kernel; > + cl=5Fcommand=5Fqueue command=5Fqueue; > + cl=5Fmem lut3d=5Fbuf; > + > + struct rgbvec *lut; > + int lutsize; > + int lutsize2; > + struct rgbvec scale; > + int interpolation; /// + char *file; > +=7D LUT3DOpenCLContext; > + > +static int allocate=5F3dlut(AV=46ilterContext *ctx, int lutsize) > +=7B > + LUT3DOpenCLContext *lut3d =3D ctx->priv; > + if (lutsize < 2 =7C=7C lutsize > MAX=5FLEVEL) =7B > + av=5Flog(ctx, AV=5FLOG=5FERROR, =22Too large or invalid 3D LUT size=5C= n=22); > + return AVERROR(EINVAL); > + =7D > + > + av=5Ffreep(&lut3d->lut); > + lut3d->lut =3D av=5Fmalloc=5Farray(lutsize * lutsize * lutsize, sizeo= f(*lut3d->lut)); > + if (=21lut3d->lut) > + return AVERROR(ENOMEM); > + > + lut3d->lutsize =3D lutsize; > + lut3d->lutsize2 =3D lutsize * lutsize; > + return 0; > +=7D > + > +static int set=5Fidentity=5Fmatrix(AV=46ilterContext *ctx, int size) > +=7B > + LUT3DOpenCLContext *lut3d =3D ctx->priv; > + int ret, i, j, k; > + const int size2 =3D size * size; > + const float c =3D 1. / (size - 1); > + > + ret =3D allocate=5F3dlut(ctx, size); > + if (ret < 0) > + return ret; > + > + for (k =3D 0; k < size; k++) =7B > + for (j =3D 0; j < size; j++) =7B > + for (i =3D 0; i < size; i++) =7B > + struct rgbvec *vec =3D &lut3d->lut=5Bk * size2 + j * size + i=5D; > + vec->r =3D k * c; > + vec->g =3D j * c; > + vec->b =3D i * c; > + =7D > + =7D > + =7D > + > + return 0; > +=7D > + > +static int skip=5Fline(const char *p) > +=7B > + while (*p && av=5Fisspace(*p)) > + p++; > + return =21*p =7C=7C *p =3D=3D '=23'; > +=7D > + > +=23define NEXT=5FLINE(loop=5Fcond) do =7B =5C > + if (=21fgets(line, sizeof(line), f)) =7B =5C > + av=5Flog(ctx, AV=5FLOG=5FERROR, =22Unexpected EO=46=5Cn=22); =5C > + return AVERROR=5FINVALIDDATA; =5C > + =7D =5C > +=7D while (loop=5Fcond) > + > +static int parse=5Fcube(AV=46ilterContext *ctx, =46ILE *f) > +=7B > + LUT3DOpenCLContext *lut3d =3D ctx->priv; > + char line=5BMAX=5FLINE=5FSIZE=5D; > + > + while (fgets(line, sizeof(line), f)) =7B > + if (=21strncmp(line, =22LUT=5F3D=5FSIZE=22, 11)) =7B > + int ret, i, j, k; > + const int size =3D strtol(line + 12, NULL, 0); > + const int size2 =3D size * size; > + > + ret =3D allocate=5F3dlut(ctx, size); > + if (ret < 0) > + return ret; > + > + for (k =3D 0; k < size; k++) =7B > + for (j =3D 0; j < size; j++) =7B > + for (i =3D 0; i < size; i++) =7B > + struct rgbvec *vec =3D &lut3d->lut=5Bi * size2 + j * size + k=5D; > + > + do =7B > +try=5Fagain: > + NEXT=5FLINE(0); > + if (=21strncmp(line, =22DOMAIN=5F=22, 7)) =7B > + av=5Flog(ctx, AV=5FLOG=5FERROR, =22Min/max not supported in this form= at=5Cn=22); > + return AVERROR=5FINVALIDDATA; > + =7D else if (=21strncmp(line, =22TITLE=22, 5)) =7B > + goto try=5Fagain; > + =7D > + =7D while (skip=5Fline(line)); > + if (av=5Fsscanf(line, =22%f %f %f=22, &vec->r, &vec->g, &vec->b) =21=3D= 3) > + return AVERROR=5FINVALIDDATA; > + =7D > + =7D > + =7D > + break; > + =7D > + =7D > + > + return 0; > +=7D > + > +static int lut3d=5Fopencl=5Finit=5Fdevice(AV=46ilterContext *avctx) > +=7B > + int err; > + LUT3DOpenCLContext *ctx =3D avctx->priv; > + cl=5Fint cle; > + > + > + size=5Ft n =3D ctx->lutsize; > + size=5Ft total =3D n * n * n; > + > + > + cl=5Fmem lut3d=5Fbuf =3D clCreateBuffer(ctx->ocf.hwctx->context, > + CL=5FMEM=5FREAD=5FONLY =7C > + CL=5FMEM=5FCOPY=5FHOST=5FPTR =7C > + CL=5FMEM=5FHOST=5FNO=5FACCESS, > + sizeof(rgbvec) * total, > + ctx->lut, &cle); > + > + if (=21lut3d=5Fbuf) =7B > + av=5Flog(avctx, AV=5FLOG=5FERROR, =22=46ailed to create buffer: =22 > + =22%d.=5Cn=22, cle); > + return AVERROR(EIO); > + =7D > + ctx->lut3d=5Fbuf =3D lut3d=5Fbuf; > + > + av=5Flog(avctx, AV=5FLOG=5FDEBUG, =22LUT3D data loaded onto host=5Cn=22= ); > + > + > + > + > + err =3D ff=5Fopencl=5Ffilter=5Fload=5Fprogram(avctx, &ff=5Fsource=5Fl= ut3d=5Fcl, 1); > + if (err < 0) > + return err; > + > + ctx->command=5Fqueue =3D clCreateCommandQueue(ctx->ocf.hwctx->context= , > + ctx->ocf.hwctx->device=5Fid, > + 0, &cle); > + CL=5F=46AIL=5FON=5FERROR(AVERROR(EIO), =22=46ailed to create OpenCL =22= > + =22command queue %d.=5Cn=22, cle); > + > + const char *kernel=5Fname; > + switch (ctx->interpolation) =7B > + case INTERPOLATE=5FNEAREST: kernel=5Fname =3D =22lut3d=5Finterp=5Fnea= rest=22; break; > + case INTERPOLATE=5FTRILINEAR: kernel=5Fname =3D =22lut3d=5Finterp=5Ft= rilinear=22; break; > + case INTERPOLATE=5FTETRAHEDRAL: kernel=5Fname =3D =22lut3d=5Finterp=5F= tetrahedral=22; break; > + default: > + av=5Fassert0(0); > + =7D > + ctx->kernel =3D clCreateKernel(ctx->ocf.program, kernel=5Fname, &cle)= ; > + CL=5F=46AIL=5FON=5FERROR(AVERROR(EIO), =22=46ailed to create =22 > + =22kernel %d.=5Cn=22, cle); > + > + ctx->initialised =3D 1; > + return 0; > + fail: > + if (ctx->command=5Fqueue) > + clReleaseCommandQueue(ctx->command=5Fqueue); > + if (ctx->kernel) > + clReleaseKernel(ctx->kernel); > + return err; > +=7D > + > +static int lut3d=5Fopencl=5Finit(AV=46ilterContext *avctx) > +=7B > + > + av=5Flog(avctx, AV=5FLOG=5FDEBUG, =22Starting intialization of LUT3D = OpenCL=5Cn=22); > + LUT3DOpenCLContext *ctx =3D avctx->priv; > + int err =3D 0; > + > + ff=5Fopencl=5Ffilter=5Finit(avctx); > + > + av=5Flog(avctx, AV=5FLOG=5FDEBUG, =22LUT3D OpenCL filter initialized=5C= n=22); > + > + > + =46ILE *f; > + const char *ext; > + > + if (=21ctx->file) =7B > + return set=5Fidentity=5Fmatrix(avctx, 32); > + =7D > + else =7B > + ext =3D strrchr(ctx->file, '.'); > + if (=21ext) =7B > + av=5Flog(avctx, AV=5FLOG=5FERROR, =22Unable to guess the format from = the extension=5Cn=22); > + err =3D AVERROR=5FINVALIDDATA; > + return err; > + =7D > + ext++; > + if (=21av=5Fstrcasecmp(ext, =22cube=22)) =7B > + f =3D avpriv=5Ffopen=5Futf8(ctx->file, =22r=22); > + if (=21f) =7B > + err =3D AVERROR(errno); > + av=5Flog(avctx, AV=5FLOG=5FERROR, =22%s: %s=5Cn=22, ctx->file, av=5Fe= rr2str(err)); > + return err; > + =7D > + err =3D parse=5Fcube(avctx, f); > + fclose(f); > + =7D else =7B > + av=5Flog(avctx, AV=5FLOG=5FERROR, =22Unrecognized '.%s' file type=5Cn= =22, ext); > + err =3D AVERROR(EINVAL); > + return err; > + =7D > + if (=21err && =21ctx->lutsize) =7B > + av=5Flog(avctx, AV=5FLOG=5FERROR, =223D LUT is empty=5Cn=22); > + err =3D AVERROR=5FINVALIDDATA; > + return err; > + =7D > + > + =7D > + av=5Flog(avctx, AV=5FLOG=5FDEBUG, =22LUT3D OpenCL data loaded=5Cn=22)= ; > + return err; > +=7D > + > +static int lut3d=5Fopencl=5Ffilter=5Fframe(AV=46ilterLink *inlink, AV=46= rame *input) > +=7B > + AV=46ilterContext *avctx =3D inlink->dst; > + AV=46ilterLink *outlink =3D avctx->outputs=5B0=5D; > + LUT3DOpenCLContext *ctx =3D avctx->priv; > + AV=46rame *output =3D NULL; > + cl=5Fint cle; > + size=5Ft global=5Fwork=5B2=5D; > + cl=5Fmem src, dst; > + int err, p; > + > + av=5Flog(ctx, AV=5FLOG=5FDEBUG, =22=46ilter input: %s, %ux%u (%=22PRI= d64=22).=5Cn=22, > + av=5Fget=5Fpix=5Ffmt=5Fname(input->format), > + input->width, input->height, input->pts); > + > + if (=21input->hw=5Fframes=5Fctx) > + return AVERROR(EINVAL); > + > + if (=21ctx->initialised) =7B > + AVHW=46ramesContext *input=5Fctx =3D > + (AVHW=46ramesContext*)input->hw=5Fframes=5Fctx->data; > + int fmt =3D input=5Fctx->sw=5Fformat; > + > + // Make sure the input is a format we support > + if (fmt =21=3D AV=5FPIX=5F=46MT=5FARGB && > + fmt =21=3D AV=5FPIX=5F=46MT=5FRGBA && > + fmt =21=3D AV=5FPIX=5F=46MT=5FABGR && > + fmt =21=3D AV=5FPIX=5F=46MT=5FBGRA > + ) =7B > + av=5Flog(avctx, AV=5FLOG=5FERROR, =22unsupported (non-RGB) format in = lut3d=5Fopencl.=5Cn=22); > + err =3D AVERROR(ENOSYS); > + goto fail; > + =7D > + > + > + err =3D lut3d=5Fopencl=5Finit=5Fdevice(avctx); > + if (err < 0) > + goto fail; > + =7D > + > + output =3D ff=5Fget=5Fvideo=5Fbuffer(outlink, outlink->w, outlink->h)= ; > + if (=21output) =7B > + err =3D AVERROR(ENOMEM); > + goto fail; > + =7D > + > + > + for (p =3D 0; p < =46=46=5FARRAY=5FELEMS(output->data); p++) =7B > + src =3D (cl=5Fmem) input->data=5Bp=5D; > + dst =3D (cl=5Fmem)output->data=5Bp=5D; > + > + if (=21dst) > + break; > + > + CL=5FSET=5FKERNEL=5FARG(ctx->kernel, 0, cl=5Fmem, &src); > + CL=5FSET=5FKERNEL=5FARG(ctx->kernel, 1, cl=5Fmem, &dst); > + CL=5FSET=5FKERNEL=5FARG(ctx->kernel, 2, cl=5Fmem, &ctx->lut3d=5Fbuf);= > + CL=5FSET=5FKERNEL=5FARG(ctx->kernel, 3, cl=5Fint, &ctx->lutsize); > + > + err =3D ff=5Fopencl=5Ffilter=5Fwork=5Fsize=5Ffrom=5Fimage(avctx, glob= al=5Fwork, output, p, 0); > + if (err < 0) > + goto fail; > + > + av=5Flog(avctx, AV=5FLOG=5FDEBUG, =22Run kernel on plane %d =22 > + =22(%=22SIZE=5FSPECI=46IER=22x%=22SIZE=5FSPECI=46IER=22).=5Cn=22, > + p, global=5Fwork=5B0=5D, global=5Fwork=5B1=5D); > + > + cle =3D clEnqueueNDRangeKernel(ctx->command=5Fqueue, ctx->kernel, 2, = NULL, > + global=5Fwork, NULL, > + 0, NULL, NULL); > + CL=5F=46AIL=5FON=5FERROR(AVERROR(EIO), =22=46ailed to enqueue =22 > + =22kernel: %d.=5Cn=22, cle); > + =7D > + > + cle =3D cl=46inish(ctx->command=5Fqueue); > + CL=5F=46AIL=5FON=5FERROR(AVERROR(EIO), =22=46ailed to finish command = queue: %d.=5Cn=22, cle); > + > + err =3D av=5Fframe=5Fcopy=5Fprops(output, input); > + if (err < 0) > + goto fail; > + > + av=5Fframe=5Ffree(&input); > + > + av=5Flog(ctx, AV=5FLOG=5FDEBUG, =22=46ilter output: %s, %ux%u (%=22PR= Id64=22).=5Cn=22, > + av=5Fget=5Fpix=5Ffmt=5Fname(output->format), > + output->width, output->height, output->pts); > + > + return ff=5Ffilter=5Fframe(outlink, output); > + > +fail: > + cl=46inish(ctx->command=5Fqueue); > + av=5Fframe=5Ffree(&input); > + av=5Fframe=5Ffree(&output); > + return err; > +=7D > + > +static av=5Fcold void lut3d=5Fopencl=5Funinit(AV=46ilterContext *avctx= ) > +=7B > + LUT3DOpenCLContext *ctx =3D avctx->priv; > + cl=5Fint cle; > + > + clReleaseMemObject(ctx->lut3d=5Fbuf); > + > + if (ctx->kernel) =7B > + cle =3D clReleaseKernel(ctx->kernel); > + if (cle =21=3D CL=5FSUCCESS) > + av=5Flog(avctx, AV=5FLOG=5FERROR, =22=46ailed to release =22 > + =22kernel: %d.=5Cn=22, cle); > + =7D > + > + if (ctx->command=5Fqueue) =7B > + cle =3D clReleaseCommandQueue(ctx->command=5Fqueue); > + if (cle =21=3D CL=5FSUCCESS) > + av=5Flog(avctx, AV=5FLOG=5FERROR, =22=46ailed to release =22 > + =22command queue: %d.=5Cn=22, cle); > + =7D > + > + av=5Ffreep(&ctx->lut); > + > + ff=5Fopencl=5Ffilter=5Funinit(avctx); > +=7D > + > +static const AV=46ilterPad lut3d=5Fopencl=5Finputs=5B=5D =3D =7B > + =7B > + .name =3D =22default=22, > + .type =3D AVMEDIA=5FTYPE=5FVIDEO, > + .filter=5Fframe =3D &lut3d=5Fopencl=5Ffilter=5Fframe, > + .config=5Fprops =3D &ff=5Fopencl=5Ffilter=5Fconfig=5Finput, > + =7D, > +=7D; > + > +static const AV=46ilterPad lut3d=5Fopencl=5Foutputs=5B=5D =3D =7B > + =7B > + .name =3D =22default=22, > + .type =3D AVMEDIA=5FTYPE=5FVIDEO, > + .config=5Fprops =3D &ff=5Fopencl=5Ffilter=5Fconfig=5Foutput, > + =7D, > +=7D; > + > +=23define O=46=46SET(x) offsetof(LUT3DOpenCLContext, x) > +=23define =46LAGS (AV=5FOPT=5F=46LAG=5F=46ILTERING=5FPARAM =7C AV=5FOP= T=5F=46LAG=5FVIDEO=5FPARAM) > +=23define T=46LAGS AV=5FOPT=5F=46LAG=5F=46ILTERING=5FPARAM=7CAV=5FOPT=5F= =46LAG=5FVIDEO=5FPARAM=7CAV=5FOPT=5F=46LAG=5FRUNTIME=5FPARAM > + > + > + > +=23if CON=46IG=5FLUT3D=5FOPENCL=5F=46ILTER > + > + > +static const AVOption lut3d=5Fopencl=5Foptions=5B=5D =3D =7B > + =7B =22file=22, =22set 3D LUT file name=22, O=46=46SET(file), AV=5FOP= T=5FTYPE=5FSTRING, =7B.str=3DNULL=7D, .flags =3D =46LAGS =7D, > + =7B =22interp=22, =22select interpolation mode=22, O=46=46SET(interpo= lation), AV=5FOPT=5FTYPE=5FINT, =7B.i64=3DINTERPOLATE=5FTETRAHEDRAL=7D, 0= , NB=5FINTERP=5FMODE-1, T=46LAGS, .unit =3D =22interp=5Fmode=22 =7D, > + =7B =22nearest=22, =22use values from the nearest defined points=22, = 0, AV=5FOPT=5FTYPE=5FCONST, =7B.i64=3DINTERPOLATE=5FNEAREST=7D, 0, 0, T=46= LAGS, .unit =3D =22interp=5Fmode=22 =7D, > + =7B =22trilinear=22, =22interpolate values using the 8 points definin= g a cube=22, 0, AV=5FOPT=5FTYPE=5FCONST, =7B.i64=3DINTERPOLATE=5FTRILINEA= R=7D, 0, 0, T=46LAGS, .unit =3D =22interp=5Fmode=22 =7D, > + =7B =22tetrahedral=22, =22interpolate values using a tetrahedron=22, = 0, AV=5FOPT=5FTYPE=5FCONST, =7B.i64=3DINTERPOLATE=5FTETRAHEDRAL=7D, 0, 0,= T=46LAGS, .unit =3D =22interp=5Fmode=22 =7D, =5C > + =7B NULL =7D > +=7D; > + > +AV=46ILTER=5FDE=46INE=5FCLASS(lut3d=5Fopencl); > + > +const =46=46=46ilter ff=5Fvf=5Flut3d=5Fopencl =3D =7B > + .p.name =3D =22lut3d=5Fopencl=22, > + .p.description =3D NULL=5FI=46=5FCON=46IG=5FSMALL(=22Adjust colors us= ing a 3D LUT.=22), > + .p.priv=5Fclass =3D &lut3d=5Fopencl=5Fclass, > + .p.flags =3D AV=46ILTER=5F=46LAG=5FHWDEVICE, > + .priv=5Fsize =3D sizeof(LUT3DOpenCLContext), > + .init =3D &lut3d=5Fopencl=5Finit, > + .uninit =3D &lut3d=5Fopencl=5Funinit, > + =46ILTER=5FINPUTS(lut3d=5Fopencl=5Finputs), > + =46ILTER=5FOUTPUTS(lut3d=5Fopencl=5Foutputs), > + =46ILTER=5FSINGLE=5FPIX=46MT(AV=5FPIX=5F=46MT=5FOPENCL), > + .flags=5Finternal =3D =46=46=5F=46ILTER=5F=46LAG=5FHW=46RAME=5FAWARE,= > +=7D; > + > +=23endif /* CON=46IG=5FLUT3D=5FOPENCL=5F=46ILTER */ > -- > 2.39.5 (Apple Git-154) > > --===============4016265456810750075== Content-Type: text/plain; charset="us-ascii" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit Content-Disposition: inline _______________________________________________ ffmpeg-devel mailing list ffmpeg-devel@ffmpeg.org https://ffmpeg.org/mailman/listinfo/ffmpeg-devel To unsubscribe, visit link above, or email ffmpeg-devel-request@ffmpeg.org with subject "unsubscribe". --===============4016265456810750075==--