From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: from ffbox0-bg.mplayerhq.hu (ffbox0-bg.ffmpeg.org [79.124.17.100]) by master.gitmailbox.com (Postfix) with ESMTP id A5A2241252 for ; Sun, 11 Sep 2022 07:29:08 +0000 (UTC) Received: from [127.0.1.1] (localhost [127.0.0.1]) by ffbox0-bg.mplayerhq.hu (Postfix) with ESMTP id 3BC3C68BAEB; Sun, 11 Sep 2022 10:29:05 +0300 (EEST) Received: from NAM10-BN7-obe.outbound.protection.outlook.com (mail-bn7nam10on2062.outbound.protection.outlook.com [40.107.92.62]) by ffbox0-bg.mplayerhq.hu (Postfix) with ESMTPS id 50B3768B889 for ; Sun, 11 Sep 2022 10:28:58 +0300 (EEST) ARC-Seal: i=1; a=rsa-sha256; s=arcselector9901; d=microsoft.com; cv=none; b=nb8Ohu1PcEyVp+DOoRnta+6sWVxamhzAbKZqUSryeboE3ddWKOnL+dXF57ejGtoq6kQgUJuqeng9w3tXVSunfmjno4ebGD8/02KCqUxT01JhVWWk7M5SKdjdvII4mcRQGxFE8z5dAk7htUMbuJ+grysBaC+KghJRU5l1rRWdDwk7DEAIZMa87AxekD6iShHzEarCJA6uJNIxZs7OfZ3QHHimySLtlFDOt4mLAksGvsFyqXKg3bMaYZtxcKqPLhBStH9socMi7i9DaSsHM7g05UqOP6HaGHkUePWD1gFwqpQjyZ9vsgatSH5lFFteOla2n5WofwuRUprcyZdqa/eJxg== ARC-Message-Signature: i=1; a=rsa-sha256; c=relaxed/relaxed; d=microsoft.com; s=arcselector9901; h=From:Date:Subject:Message-ID:Content-Type:MIME-Version:X-MS-Exchange-AntiSpam-MessageData-ChunkCount:X-MS-Exchange-AntiSpam-MessageData-0:X-MS-Exchange-AntiSpam-MessageData-1; bh=xQbC6M1UKLryRKhk5qNM7SDAE134+K5b94okEGoMCCk=; b=RmlFPve06LusCqIODoNDkgw5zlfrMMDQqtnY5fXfEEbrWz9py1zmcGjhGz9oz0de2OBsOdCxKBXliJcEg8GBhYQu6cXPrMrvW3f3PQ6a9S25dFbNFV97LGq705Ilp2K4fkjTYf2N9UkwN8EGSuhuT0lSZPW/u+5XbENXB1J3bSki613o+TaQtfaBbgccrMSwtnoz/M2fpCUfnfV4y+4EIriDX2HBnHwfoy4iJEQGHdsDEp6vhSvW+pXFB2sQkvnOASoXJ+cfssDhBH3S3hjfVmVCr+TLZVe0Ng6qIe0claq+C/BtBrZ17xESDfal+sc3GN7IVS/jX1aqU9GTuYaLGw== ARC-Authentication-Results: i=1; mx.microsoft.com 1; spf=pass smtp.mailfrom=nvidia.com; dmarc=pass action=none header.from=nvidia.com; dkim=pass header.d=nvidia.com; arc=none DKIM-Signature: v=1; a=rsa-sha256; c=relaxed/relaxed; d=Nvidia.com; s=selector2; h=From:Date:Subject:Message-ID:Content-Type:MIME-Version:X-MS-Exchange-SenderADCheck; bh=xQbC6M1UKLryRKhk5qNM7SDAE134+K5b94okEGoMCCk=; b=lpb/6nKjakUvNXW/OPlmWSUlPeGfXkKRLvgZc9q9m90jbhDLblgs/ZDbf0/tG93fsrPy9Lb3H6uHc6zzpjG5iKf98ST6Gui6hiFsp6RQWx/e8vl9ZMud857lRWOfO55A3+jxSNgx77DOhpMSyrUyjaXPsowsDnlfob8TzbKGC3QR9y+Af53w7YzMlTdSOyfkh7LHsQ2V5TYcaToe1dNkNzJlWlBHBjZ0j2wIOcANOYxfpKJcGCJJD5/bQpuNbjWHnuRWKMWv3XLwduVVcfne6vmUZKxkX0yHUXLmB/6hTrqPH/6YtQY9LflEaND+pKoSJ5ersTFqH+AnBOxGrGhg3A== Received: from PH7PR12MB5831.namprd12.prod.outlook.com (2603:10b6:510:1d6::13) by PH7PR12MB7308.namprd12.prod.outlook.com (2603:10b6:510:20c::9) with Microsoft SMTP Server (version=TLS1_2, cipher=TLS_ECDHE_RSA_WITH_AES_256_GCM_SHA384) id 15.20.5612.20; Sun, 11 Sep 2022 07:28:52 +0000 Received: from PH7PR12MB5831.namprd12.prod.outlook.com ([fe80::dc19:2d25:8ca7:60ca]) by PH7PR12MB5831.namprd12.prod.outlook.com ([fe80::dc19:2d25:8ca7:60ca%4]) with mapi id 15.20.5612.020; Sun, 11 Sep 2022 07:28:52 +0000 From: Roman Arzumanyan To: FFmpeg development discussions and patches , timo Thread-Topic: [FFmpeg-devel] [PATCH] libavfilter/vf_colorrange_cuda: CUDA-accelerated video filter for MPEG and JPEG color range conversions Thread-Index: AQHYxOyhXyavxE/bjU2gpL6pjmNvca3YpVcAgAEwUWg= Date: Sun, 11 Sep 2022 07:28:52 +0000 Message-ID: References: <75e2d875-56ac-c55e-ea6d-632a83dd87ae@rothenpieler.org> In-Reply-To: <75e2d875-56ac-c55e-ea6d-632a83dd87ae@rothenpieler.org> Accept-Language: ru-RU, en-US Content-Language: ru-RU X-MS-Has-Attach: yes X-MS-TNEF-Correlator: msip_labels: authentication-results: dkim=none (message not signed) header.d=none;dmarc=none action=none header.from=nvidia.com; x-ms-publictraffictype: Email x-ms-traffictypediagnostic: PH7PR12MB5831:EE_|PH7PR12MB7308:EE_ x-ms-office365-filtering-correlation-id: 37a2461f-c49a-4e95-d10f-08da93c746c5 x-ms-exchange-senderadcheck: 1 x-ms-exchange-antispam-relay: 0 x-microsoft-antispam: BCL:0; x-microsoft-antispam-message-info: VKMfsw9IbnF5JfY+9nnjhBEJkaZDhr4hEyeP7AVWtpNKOrJWzHO25pgf3K5XGXkgBrRlmr9WXRoCOVRzjWOC42VjulNuAwlAQjZyCOFuqpEaila+cvSmh3Rs7UzOsMGgd+kE/Gl+pEVts36YzT2X4+amRqsGGyGm9Vhh3oS3pUEOeZ/KRxaPye7WO7hvlyaFMPu/Xqv2WKr7lYMIKUsN+wYFDCjTPhhMaaBhytObJknKq6WewZTR/U70qRd1VCfFrk1Lc8csrZpQCxi65ouFH3wbuxQEQlKfaABJjeX/6AJ0ATARACVv9oZ9PWl1tl+z1sqDdvs+zO3f60Gk1j3d2jiF+jsY4a+SG2qwr6QR/RtdfWHrjQIAMGLiE/9NapLenGQtjV/fyUhPtx8DzknSlrdGCCtPanE2gpGSkchvJ9Bcyquuhu3+TaiX816MhqsLRLM6u/M7cdHtCdsX7Y6hBLKytElTl8Wv4YVrcKx2+o3SEfLglygCsncCXmFwnm1TrXwOVjWLoUz0tWdLMT/BMq99RsTzaxrAp2KLhHnHF98IxXA42v4GbbUeOEpUaKvygEm1j2YlKWzg/DQAY6L8WKW16lHWI67NmyPbUv/9PzDkX2nQjaB2vUpJNXxO3y7ZVw/1dElfS/nTpd3E88PWoa0UPiTX1iMEMp3MxNw4U77CaWdc/cLJFMVnXMsD3T/EfyDSoKKfby2hLpR37StiPHecFUSrL884b/mm7xZ3WkGkYx9/pMPgYud07B0llOayt7rL+EjzGTE4Kx7gX/uSY8BVgLgVhPoreL6xh6aFQlw= x-forefront-antispam-report: CIP:255.255.255.255; CTRY:; LANG:en; SCL:1; SRV:; IPV:NLI; SFV:NSPM; H:PH7PR12MB5831.namprd12.prod.outlook.com; PTR:; CAT:NONE; SFS:(13230016)(4636009)(346002)(39860400002)(366004)(396003)(136003)(376002)(6506007)(26005)(7696005)(19627405001)(71200400001)(45080400002)(53546011)(99936003)(83380400001)(9686003)(186003)(966005)(107886003)(41300700001)(478600001)(166002)(110136005)(55016003)(316002)(8676002)(76116006)(91956017)(8936002)(4326008)(52536014)(2906002)(66556008)(30864003)(66446008)(5660300002)(38100700002)(122000001)(54906003)(86362001)(38070700005)(33656002)(64756008)(66946007)(66476007); DIR:OUT; SFP:1101; x-ms-exchange-antispam-messagedata-chunkcount: 1 x-ms-exchange-antispam-messagedata-0: =?koi8-r?Q?JmXHiMlwElL4id4KHuiCLegOIpFU+hCxMRWm69UGeAXMz/C4g+X88vdV1JcR0c?= =?koi8-r?Q?v3ga4qv72Ffjcta5Xg0u4eNnGBW16j4eOVll5mRPZt5HqHQ29N0x9N/V9bflJA?= =?koi8-r?Q?nFFXjfjvx1JZRNohVEFAtQPY0ciiGkM+9438ybErTLishTi1JT0ANuWn9kvQCh?= =?koi8-r?Q?wJNFZ9CTTKCTvdJkfqcSbgDuJs6Q9bHkHgXnMMf9v3l3GhDr+yfgMZS2iDCYIe?= =?koi8-r?Q?eKi9Iu+JJ5WEEi7rFoYuc+ov+Wg00m+/0OhO8TZshgOxHeSm53YSA0Obg7aG7V?= =?koi8-r?Q?G3HnWNFPp+wVs4CefeEEXdkWiSR1WAqlyB7oz8EhHY3669SZx8tmZq/A+gAiYX?= =?koi8-r?Q?/M6JmVm0LNaBjPUCghnIn7q6ToqBO4oaNnX+rJGBQmwBAMebVWmYxjvLe9JSiZ?= =?koi8-r?Q?R+nBS/9hHKwmRg1uyY8zCq6giTEMhZQfMXLELGmEcjFLyd9lU6Y1oqY+UbjsQC?= =?koi8-r?Q?TBB2jzhou3egAyyqwKsUS7THzEaunJm4wwkJgLokHEXUgsJXv+7K1Qaa/5yOxc?= =?koi8-r?Q?17RDLtN34BEwiAuwlcbgUBTUnCz3EkrWWjGFI501sB35dYjce3mTI4ssyx51GH?= =?koi8-r?Q?oPn0WRe+ua2Oi944ItdwVx9VdY4un3k6Xo8IMWPW40UheesP2/7np7L3gdMB3x?= =?koi8-r?Q?3lqeYkDSAfHNVe1PNLBLeHTqCqU2Y7Z0LBEUe1C6tRKn0MPfG6SqNzXYdWHimR?= =?koi8-r?Q?O5sXFiFq26dgS6lTe04rKINfWMQj7a9YheNkmDlZ2BoCYrlf6hGswWvmNA6F6H?= =?koi8-r?Q?JouhqyDjF5eRc1687poI6xsdzWqHb+Yp6AfWzrJDxZ2TjLcASLixLwxm1iUNxu?= =?koi8-r?Q?SiXvSok74S9m/jJMPz+4/9rjVepmBWoMFdoxjCGI2+wfptvSQ4DJWTNlWCGulB?= =?koi8-r?Q?Z/iKI9cxY1WFB3247hwMVRsPnbxapjLZ54czYsN5zcRTD4MFewgGJDbkv3IoTn?= =?koi8-r?Q?iSYYOS4FUrZv3VTEuNznYKdXUcZ1YsnvPemqfufee4em+F/u2fBWlFc0l8aDKw?= =?koi8-r?Q?X5fhjmtKgYpSjr/qzRSjX4zlMwul4pwXdpMFe++GeYLjnzauu//+KebU/LX68S?= =?koi8-r?Q?6jEdmaFhzzYHE3/iRjgeOLrAWFcpggaEtPJOig3/uhb6tG5I7Bxv/YIUYpGZOR?= =?koi8-r?Q?qXhPN/kyzXRsTrdarESojp8Jnmnl+0NgcB9tJpW/+Rl8fT05C7lFbEW7iN56Kk?= =?koi8-r?Q?1a9TdOetq8kiDOGanexFW6ujqMqrn0EeqEkWeRBTJ0R0FIxfZl4m8QZwRYGOGp?= =?koi8-r?Q?TbmowibRSUuCIM54VDxnuffFTzFu7ciJ/5ba1Kg4N9xahVpkh0e0KksGSI0XN5?= =?koi8-r?Q?iQpEhpaSIfGx1B/xLNvuGL/p53Io5zFoDV0j+BKiYqy7TouOa36gKid2/8FRMK?= =?koi8-r?Q?MDtbbyxEkemFTDZXuzudW2DFBTPOK05X5v1aAJ7I1P649KPybDK6APfPmO1sC1?= =?koi8-r?Q?PPfPFlbf9d3F4l/wyCgbWPaS/jvfAxrnT/cJQQ0wppwtM/gZ8jS3if58pGt3AK?= =?koi8-r?Q?vLdqxp20/NOzZnjLGdKzXCkj0tgF7AYdnyrxDKxtWOg0Svd7Dn?= Content-Type: multipart/mixed; boundary="_004_PH7PR12MB58312279BA34EDEDC705A0CFD2459PH7PR12MB5831namp_" MIME-Version: 1.0 X-OriginatorOrg: Nvidia.com X-MS-Exchange-CrossTenant-AuthAs: Internal X-MS-Exchange-CrossTenant-AuthSource: PH7PR12MB5831.namprd12.prod.outlook.com X-MS-Exchange-CrossTenant-Network-Message-Id: 37a2461f-c49a-4e95-d10f-08da93c746c5 X-MS-Exchange-CrossTenant-originalarrivaltime: 11 Sep 2022 07:28:52.1333 (UTC) X-MS-Exchange-CrossTenant-fromentityheader: Hosted X-MS-Exchange-CrossTenant-id: 43083d15-7273-40c1-b7db-39efd9ccc17a X-MS-Exchange-CrossTenant-mailboxtype: HOSTED X-MS-Exchange-CrossTenant-userprincipalname: rSA2uY0wfaptuBUQOwAc6CH+4aYWXJey7syZsLYphz1CZvY66/80yRznf6krcZPt8C/6OCC80EEq4js51USDvw== X-MS-Exchange-Transport-CrossTenantHeadersStamped: PH7PR12MB7308 X-Content-Filtered-By: Mailman/MimeDel 2.1.29 Subject: Re: [FFmpeg-devel] [PATCH] libavfilter/vf_colorrange_cuda: CUDA-accelerated video filter for MPEG and JPEG color range conversions 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 Cc: Yogender Gupta , Sven Middelberg , Hermann Held Errors-To: ffmpeg-devel-bounces@ffmpeg.org Sender: "ffmpeg-devel" Archived-At: List-Archive: List-Post: --_004_PH7PR12MB58312279BA34EDEDC705A0CFD2459PH7PR12MB5831namp_ Content-Type: text/plain; charset="koi8-r" Content-Transfer-Encoding: quoted-printable Thanks for the detailed review, Timo. Please find fixed patch in attachement. ________________________________ =EF=D4: ffmpeg-devel =CF=D4 =C9=CD=C5=CE= =C9 Timo Rothenpieler =EF=D4=D0=D2=C1=D7=CC=C5=CE=CF: 10 =D3=C5=CE=D4=D1=C2=D2=D1 2022 =C7. 16:16 =EB=CF=CD=D5: FFmpeg development discussions and patches ; Roman Arzumanyan =EB=CF=D0=C9=D1: Yogender Gupta ; Sven Middelberg ; Hermann Held =F4=C5=CD=C1: Re: [FFmpeg-devel] [PATCH] libavfilter/vf_colorrange_cuda: CU= DA-accelerated video filter for MPEG and JPEG color range conversions External email: Use caution opening links or attachments On 10.09.2022 10:16, Roman Arzumanyan wrote: > From 2b15d8a609a12d97b1ba7500c7f8771b336e2fdf Mon Sep 17 00:00:00 2001 > From: Roman Arzumanyan > Date: Sat, 10 Sep 2022 11:05:56 +0300 > Subject: [PATCH] libavfilter/vf_colorrange_cuda CUDA-accelerated color ra= nge > conversion filter We could also call this colorspace_cuda, since it does overlap with what the colorspace software filter does, just not nearly to the same degree of feature-completeness. That's fine in my book though, and if someone cares enough, the other features of the colorspace filter can be added over time. > --- > configure | 2 + > libavfilter/Makefile | 3 + > libavfilter/allfilters.c | 1 + > libavfilter/vf_colorrange_cuda.c | 432 ++++++++++++++++++++++++++++++ > libavfilter/vf_colorrange_cuda.cu | 93 +++++++ > 5 files changed, 531 insertions(+) > create mode 100644 libavfilter/vf_colorrange_cuda.c > create mode 100644 libavfilter/vf_colorrange_cuda.cu > > diff --git a/configure b/configure > index 9d6457d81b..e5f9738ad1 100755 > --- a/configure > +++ b/configure > @@ -3155,6 +3155,8 @@ transpose_npp_filter_deps=3D"ffnvcodec libnpp" > overlay_cuda_filter_deps=3D"ffnvcodec" > overlay_cuda_filter_deps_any=3D"cuda_nvcc cuda_llvm" > sharpen_npp_filter_deps=3D"ffnvcodec libnpp" > +colorrange_cuda_filter_deps=3D"ffnvcodec" > +colorrange_cuda_filter_deps_any=3D"cuda_nvcc cuda_llvm" Typically should be sorted in by alphapetical ordering. > amf_deps_any=3D"libdl LoadLibrary" > nvenc_deps=3D"ffnvcodec" > diff --git a/libavfilter/Makefile b/libavfilter/Makefile > index 30cc329fb6..784e154d81 100644 > --- a/libavfilter/Makefile > +++ b/libavfilter/Makefile > @@ -230,6 +230,9 @@ OBJS-$(CONFIG_COLORMAP_FILTER) +=3D vf_= colormap.o > OBJS-$(CONFIG_COLORMATRIX_FILTER) +=3D vf_colormatrix.o > OBJS-$(CONFIG_COLORSPACE_FILTER) +=3D vf_colorspace.o colors= pacedsp.o > OBJS-$(CONFIG_COLORTEMPERATURE_FILTER) +=3D vf_colortemperature.o > +OBJS-$(CONFIG_COLORRANGE_CUDA_FILTER) +=3D vf_colorrange_cuda.o \ > + vf_colorrange_cuda.ptx.o= \ > + cuda/load_helper.o Same here on alphabetical ordering, should be between colormatrix and colorspace. > OBJS-$(CONFIG_CONVOLUTION_FILTER) +=3D vf_convolution.o > OBJS-$(CONFIG_CONVOLUTION_OPENCL_FILTER) +=3D vf_convolution_opencl.= o opencl.o \ > opencl/convolution.o > diff --git a/libavfilter/allfilters.c b/libavfilter/allfilters.c > index 5ebacfde27..5e9cbe57ec 100644 > --- a/libavfilter/allfilters.c > +++ b/libavfilter/allfilters.c > @@ -213,6 +213,7 @@ extern const AVFilter ff_vf_colormap; > extern const AVFilter ff_vf_colormatrix; > extern const AVFilter ff_vf_colorspace; > extern const AVFilter ff_vf_colortemperature; > +extern const AVFilter ff_vf_colorrange_cuda; > extern const AVFilter ff_vf_convolution; > extern const AVFilter ff_vf_convolution_opencl; > extern const AVFilter ff_vf_convolve; > diff --git a/libavfilter/vf_colorrange_cuda.c b/libavfilter/vf_colorrange= _cuda.c > new file mode 100644 > index 0000000000..949e7d3bbf > --- /dev/null > +++ b/libavfilter/vf_colorrange_cuda.c > @@ -0,0 +1,432 @@ > +/* > + * Copyright (c) 2022, NVIDIA CORPORATION. All rights reserved. > + * > + * Permission is hereby granted, free of charge, to any person obtaining= a > + * copy of this software and associated documentation files (the "Softwa= re"), > + * to deal in the Software without restriction, including without limita= tion > + * the rights to use, copy, modify, merge, publish, distribute, sublicen= se, > + * and/or sell copies of the Software, and to permit persons to whom the > + * Software is furnished to do so, subject to the following conditions: > + * > + * The above copyright notice and this permission notice shall be includ= ed in > + * all copies or substantial portions of the Software. > + * > + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRE= SS OR > + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILI= TY, > + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SH= ALL > + * THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR = OTHER > + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISI= NG > + * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER > + * DEALINGS IN THE SOFTWARE. > + */ > + > +#include > + > +#include "libavutil/avstring.h" > +#include "libavutil/common.h" > +#include "libavutil/cuda_check.h" > +#include "libavutil/hwcontext.h" > +#include "libavutil/hwcontext_cuda_internal.h" > +#include "libavutil/internal.h" > +#include "libavutil/opt.h" > +#include "libavutil/pixdesc.h" > + > +#include "avfilter.h" > +#include "formats.h" > +#include "internal.h" > +#include "scale_eval.h" > +#include "video.h" > + > +#include "cuda/load_helper.h" > + > +static const enum AVPixelFormat supported_formats[] =3D { > + AV_PIX_FMT_NV12, > + AV_PIX_FMT_YUV420P, > + AV_PIX_FMT_YUV444P, > +}; > + > +#define DIV_UP(a, b) (((a) + (b)-1) / (b)) > +#define BLOCKX 32 > +#define BLOCKY 16 > + > +#define CHECK_CU(x) FF_CUDA_CHECK_DL(ctx, s->hwctx->internal->cuda_dl, x= ) > + > +typedef struct CUDAConvContext { > + const AVClass* class; > + > + AVCUDADeviceContext* hwctx; > + AVBufferRef* frames_ctx; > + AVFrame* own_frame; > + AVFrame* tmp_frame; > + > + CUcontext cu_ctx; > + CUstream cu_stream; > + CUmodule cu_module; > + CUfunction cu_convert[AVCOL_RANGE_NB]; > + > + enum AVPixelFormat pix_fmt; > + enum AVColorRange range; > + > + int num_planes; > +} CUDAConvContext; > + > +static av_cold int cudaconv_init(AVFilterContext* ctx) > +{ > + CUDAConvContext* s =3D ctx->priv; > + > + s->own_frame =3D av_frame_alloc(); > + if (!s->own_frame) > + return AVERROR(ENOMEM); > + > + s->tmp_frame =3D av_frame_alloc(); > + if (!s->tmp_frame) > + return AVERROR(ENOMEM); > + > + return 0; > +} > + > +static av_cold void cudaconv_uninit(AVFilterContext* ctx) > +{ > + CUDAConvContext* s =3D ctx->priv; > + > + if (s->hwctx && s->cu_module) { > + CudaFunctions* cu =3D s->hwctx->internal->cuda_dl; > + CUcontext dummy; > + > + CHECK_CU(cu->cuCtxPushCurrent(s->hwctx->cuda_ctx)); > + CHECK_CU(cu->cuModuleUnload(s->cu_module)); > + s->cu_module =3D NULL; > + CHECK_CU(cu->cuCtxPopCurrent(&dummy)); > + } > + > + av_frame_free(&s->own_frame); > + av_buffer_unref(&s->frames_ctx); > + av_frame_free(&s->tmp_frame); > +} > + > +static av_cold int init_hwframe_ctx(CUDAConvContext* s, AVBufferRef* dev= ice_ctx, > + int width, int height) > +{ > + AVBufferRef* out_ref =3D NULL; > + AVHWFramesContext* out_ctx; > + int ret; > + > + out_ref =3D av_hwframe_ctx_alloc(device_ctx); > + if (!out_ref) > + return AVERROR(ENOMEM); > + > + out_ctx =3D (AVHWFramesContext*)out_ref->data; > + > + out_ctx->format =3D AV_PIX_FMT_CUDA; > + out_ctx->sw_format =3D s->pix_fmt; > + out_ctx->width =3D FFALIGN(width, 32); > + out_ctx->height =3D FFALIGN(height, 32); > + > + ret =3D av_hwframe_ctx_init(out_ref); > + if (ret < 0) > + goto fail; > + > + av_frame_unref(s->own_frame); > + ret =3D av_hwframe_get_buffer(out_ref, s->own_frame, 0); > + if (ret < 0) > + goto fail; > + > + s->own_frame->width =3D width; > + s->own_frame->height =3D height; > + > + av_buffer_unref(&s->frames_ctx); > + s->frames_ctx =3D out_ref; > + > + return 0; > +fail: > + av_buffer_unref(&out_ref); > + return ret; > +} > + > +static int format_is_supported(enum AVPixelFormat fmt) > +{ > + for (int i =3D 0; i < FF_ARRAY_ELEMS(supported_formats); i++) > + if (fmt =3D=3D supported_formats[i]) > + return 1; > + > + return 0; > +} > + > +static av_cold int init_processing_chain(AVFilterContext* ctx, int width= , > + int height) > +{ > + CUDAConvContext* s =3D ctx->priv; > + AVHWFramesContext* in_frames_ctx; > + > + int ret; > + > + if (!ctx->inputs[0]->hw_frames_ctx) { > + av_log(ctx, AV_LOG_ERROR, "No hw context provided on input\n"); > + return AVERROR(EINVAL); > + } > + > + in_frames_ctx =3D (AVHWFramesContext*)ctx->inputs[0]->hw_frames_ctx-= >data; > + s->pix_fmt =3D in_frames_ctx->sw_format; > + > + if (!format_is_supported(s->pix_fmt)) { > + av_log(ctx, AV_LOG_ERROR, "Unsupported pixel format: %s\n", > + av_get_pix_fmt_name(s->pix_fmt)); > + return AVERROR(ENOSYS); > + } > + > + s->num_planes =3D av_pix_fmt_count_planes(s->pix_fmt); > + > + ret =3D init_hwframe_ctx(s, in_frames_ctx->device_ref, width, height= ); > + if (ret < 0) > + return ret; > + > + ctx->outputs[0]->hw_frames_ctx =3D av_buffer_ref(s->frames_ctx); > + if (!ctx->outputs[0]->hw_frames_ctx) > + return AVERROR(ENOMEM); > + > + return 0; > +} > + > +static av_cold int cudaconv_load_functions(AVFilterContext* ctx) > +{ > + CUDAConvContext* s =3D ctx->priv; > + CUcontext dummy, cuda_ctx =3D s->hwctx->cuda_ctx; > + CudaFunctions* cu =3D s->hwctx->internal->cuda_dl; > + int ret; > + > + extern const unsigned char ff_vf_colorrange_cuda_ptx_data[]; > + extern const unsigned int ff_vf_colorrange_cuda_ptx_len; > + > + ret =3D CHECK_CU(cu->cuCtxPushCurrent(cuda_ctx)); > + if (ret < 0) > + return ret; > + > + ret =3D ff_cuda_load_module(ctx, s->hwctx, &s->cu_module, > + ff_vf_colorrange_cuda_ptx_data, > + ff_vf_colorrange_cuda_ptx_len); > + if (ret < 0) > + goto fail; > + > + ret =3D CHECK_CU(cu->cuModuleGetFunction( > + &s->cu_convert[AVCOL_RANGE_MPEG], s->cu_module, > + "to_mpeg_cuda")); > + > + if (ret < 0) > + goto fail; > + > + ret =3D CHECK_CU(cu->cuModuleGetFunction( > + &s->cu_convert[AVCOL_RANGE_JPEG], s->cu_module, > + "to_jpeg_cuda")); > + > + if (ret < 0) > + goto fail; > + > +fail: > + CHECK_CU(cu->cuCtxPopCurrent(&dummy)); > + return ret; > +} > + > +static av_cold int cudaconv_config_props(AVFilterLink* outlink) > +{ > + AVFilterContext* ctx =3D outlink->src; > + AVFilterLink* inlink =3D outlink->src->inputs[0]; > + CUDAConvContext* s =3D ctx->priv; > + AVHWFramesContext* frames_ctx =3D > + (AVHWFramesContext*)inlink->hw_frames_ctx->data; > + AVCUDADeviceContext* device_hwctx =3D frames_ctx->device_ctx->hwctx; > + int ret; > + > + s->hwctx =3D device_hwctx; > + s->cu_stream =3D s->hwctx->stream; > + > + outlink->w =3D inlink->w; > + outlink->h =3D inlink->h; > + > + ret =3D init_processing_chain(ctx, inlink->w, inlink->h); > + if (ret < 0) > + return ret; > + > + if (inlink->sample_aspect_ratio.num) { > + outlink->sample_aspect_ratio =3D av_mul_q( > + (AVRational){outlink->h * inlink->w, outlink->w * inlink->h}= , > + inlink->sample_aspect_ratio); > + } else { > + outlink->sample_aspect_ratio =3D inlink->sample_aspect_ratio; > + } > + > + ret =3D cudaconv_load_functions(ctx); > + if (ret < 0) > + return ret; > + > + return ret; > +} > + > +static int conv_cuda_convert(AVFilterContext* ctx, AVFrame* out, AVFrame= * in) > +{ > + CUDAConvContext* s =3D ctx->priv; > + CudaFunctions* cu =3D s->hwctx->internal->cuda_dl; > + CUcontext dummy, cuda_ctx =3D s->hwctx->cuda_ctx; > + int ret; > + > + ret =3D CHECK_CU(cu->cuCtxPushCurrent(cuda_ctx)); > + if (ret < 0) > + return ret; > + > + out->color_range =3D s->range; > + > + for (int i =3D 0; i < s->num_planes; i++) { > + int width =3D in->width, height =3D in->height, comp_id =3D (i >= 0); > + > + switch (s->pix_fmt) { > + case AV_PIX_FMT_YUV444P: > + break; > + case AV_PIX_FMT_YUV420P: > + width =3D comp_id ? in->width / 2 : in->width; > + case AV_PIX_FMT_NV12: > + height =3D comp_id ? in->height / 2 : in->height; > + break; > + default: > + return AVERROR(ENOSYS); > + } > + > + if (in->color_range !=3D out->color_range) { > + void* args[] =3D {&in->data[i], &out->data[i], &in->linesize= [i], > + &comp_id}; > + ret =3D CHECK_CU(cu->cuLaunchKernel( > + s->cu_convert[out->color_range], DIV_UP(width, BLOCKX), What happens if the user specifies a color range that's not mpeg or jpeg? Like, UNSPECIFIED, which is even the default. The AVOption absolutely allows that, and I see no check that verifies a kernel for that conversion exists, so this would end up passing a NULL Kernel to cuLaunchKernel. Should be an easy enough check at init time, after loading the kernels. No Kernel for the given color range? Error. > + DIV_UP(height, BLOCKY), 1, BLOCKX, BLOCKY, 1, 0, s->cu_s= tream, > + args, NULL)); > + } else { > + av_hwframe_transfer_data(out, in, 0); > + } > + } > + > + CHECK_CU(cu->cuCtxPopCurrent(&dummy)); > + return ret; > +} > + > +static int cudaconv_conv(AVFilterContext* ctx, AVFrame* out, AVFrame* in= ) > +{ > + CUDAConvContext* s =3D ctx->priv; > + AVFilterLink* outlink =3D ctx->outputs[0]; > + AVFrame* src =3D in; > + int ret; > + > + ret =3D conv_cuda_convert(ctx, s->own_frame, src); > + if (ret < 0) > + return ret; > + > + src =3D s->own_frame; > + ret =3D av_hwframe_get_buffer(src->hw_frames_ctx, s->tmp_frame, 0); > + if (ret < 0) > + return ret; > + > + av_frame_move_ref(out, s->own_frame); > + av_frame_move_ref(s->own_frame, s->tmp_frame); > + > + s->own_frame->width =3D outlink->w; > + s->own_frame->height =3D outlink->h; > + > + ret =3D av_frame_copy_props(out, in); > + if (ret < 0) > + return ret; > + > + return 0; > +} > + > +static int cudaconv_filter_frame(AVFilterLink* link, AVFrame* in) > +{ > + AVFilterContext* ctx =3D link->dst; > + CUDAConvContext* s =3D ctx->priv; > + AVFilterLink* outlink =3D ctx->outputs[0]; > + CudaFunctions* cu =3D s->hwctx->internal->cuda_dl; > + > + AVFrame* out =3D NULL; > + CUcontext dummy; > + int ret =3D 0; > + > + out =3D av_frame_alloc(); > + if (!out) { > + ret =3D AVERROR(ENOMEM); > + goto fail; > + } > + > + ret =3D CHECK_CU(cu->cuCtxPushCurrent(s->hwctx->cuda_ctx)); > + if (ret < 0) > + goto fail; > + > + ret =3D cudaconv_conv(ctx, out, in); > + > + CHECK_CU(cu->cuCtxPopCurrent(&dummy)); > + if (ret < 0) > + goto fail; > + > + av_reduce(&out->sample_aspect_ratio.num, &out->sample_aspect_ratio.d= en, > + (int64_t)in->sample_aspect_ratio.num * outlink->h * link->= w, > + (int64_t)in->sample_aspect_ratio.den * outlink->w * link->= h, > + INT_MAX); > + > + av_frame_free(&in); > + return ff_filter_frame(outlink, out); > +fail: > + av_frame_free(&in); > + av_frame_free(&out); > + return ret; > +} > + > +static AVFrame* cudaconv_get_video_buffer(AVFilterLink* inlink, int w, i= nt h) > +{ > + return ff_default_get_video_buffer(inlink, w, h); > +} This function can be removed entirely, since ff_default_get_video_buffer is what's called by default anyway. > +#define OFFSET(x) offsetof(CUDAConvContext, x) > +#define FLAGS (AV_OPT_FLAG_FILTERING_PARAM | AV_OPT_FLAG_VIDEO_PARAM) > +static const AVOption options[] =3D { > + {"range", "Output video range", OFFSET(range), AV_OPT_TYPE_INT, {.i6= 4 =3D AVCOL_RANGE_UNSPECIFIED}, AVCOL_RANGE_UNSPECIFIED, AVCOL_RANGE_NB - 1= , FLAGS, "range"}, > + {"mpeg", "limited range", 0, AV_OPT_TYPE_CONST, {.i64 =3D AVCOL_= RANGE_MPEG}, 0, 0, FLAGS, "range"}, > + {"jpeg", "full range", 0, AV_OPT_TYPE_CONST, {.i64 =3D AVCOL_= RANGE_JPEG}, 0, 0, FLAGS, "range"}, > + {NULL}, > +}; > + > +static const AVClass cudaconv_class =3D { > + .class_name =3D "cudaconv", All the mentions of cudaconv in this file should be renamed to match the filter name. It doesn't overly matter for functionality, but the class name does end up in logs, and the function names are purely for neatness. > + .item_name =3D av_default_item_name, > + .option =3D options, > + .version =3D LIBAVUTIL_VERSION_INT, > +}; > + > +static const AVFilterPad cudaconv_inputs[] =3D { > + { > + .name =3D "default", > + .type =3D AVMEDIA_TYPE_VIDEO, > + .filter_frame =3D cudaconv_filter_frame, > + .get_buffer.video =3D cudaconv_get_video_buffer, > + }, > +}; > + > +static const AVFilterPad cudaconv_outputs[] =3D { > + { > + .name =3D "default", > + .type =3D AVMEDIA_TYPE_VIDEO, > + .config_props =3D cudaconv_config_props, > + }, > +}; > + > +const AVFilter ff_vf_colorrange_cuda =3D { > + .name =3D "colorrange_cuda", > + .description =3D > + NULL_IF_CONFIG_SMALL("CUDA accelerated video color range convert= er"), > + > + .init =3D cudaconv_init, > + .uninit =3D cudaconv_uninit, > + > + .priv_size =3D sizeof(CUDAConvContext), > + .priv_class =3D &cudaconv_class, > + > + FILTER_INPUTS(cudaconv_inputs), > + FILTER_OUTPUTS(cudaconv_outputs), > + > + FILTER_SINGLE_PIXFMT(AV_PIX_FMT_CUDA), > + > + .flags_internal =3D FF_FILTER_FLAG_HWFRAME_AWARE, > +}; > diff --git a/libavfilter/vf_colorrange_cuda.cu b/libavfilter/vf_colorrang= e_cuda.cu > new file mode 100644 > index 0000000000..6f617493f8 > --- /dev/null > +++ b/libavfilter/vf_colorrange_cuda.cu > @@ -0,0 +1,93 @@ > +/* > + * Copyright (c) 2022, NVIDIA CORPORATION. All rights reserved. > + * > + * Permission is hereby granted, free of charge, to any person obtaining= a > + * copy of this software and associated documentation files (the "Softwa= re"), > + * to deal in the Software without restriction, including without limita= tion > + * the rights to use, copy, modify, merge, publish, distribute, sublicen= se, > + * and/or sell copies of the Software, and to permit persons to whom the > + * Software is furnished to do so, subject to the following conditions: > + * > + * The above copyright notice and this permission notice shall be includ= ed in > + * all copies or substantial portions of the Software. > + * > + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRE= SS OR > + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILI= TY, > + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SH= ALL > + * THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR = OTHER > + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISI= NG > + * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER > + * DEALINGS IN THE SOFTWARE. > + */ > + > +extern "C" { > +#define MPEG_LUMA_MIN (16) > +#define MPEG_CHROMA_MIN (16) > +#define MPEG_LUMA_MAX (235) > +#define MPEG_CHROMA_MAX (240) > + > +#define JPEG_LUMA_MIN (0) > +#define JPEG_CHROMA_MIN (1) > +#define JPEG_LUMA_MAX (255) > +#define JPEG_CHROMA_MAX (255) > + > +__device__ int mpeg_min[] =3D {MPEG_LUMA_MIN, MPEG_CHROMA_MIN}; > +__device__ int mpeg_max[] =3D {MPEG_LUMA_MAX, MPEG_CHROMA_MAX}; > + > +__device__ int jpeg_min[] =3D {JPEG_LUMA_MIN, JPEG_CHROMA_MIN}; > +__device__ int jpeg_max[] =3D {JPEG_LUMA_MAX, JPEG_CHROMA_MAX}; > + > +__device__ int clamp(int val, int min, int max) > +{ > + if (val < min) > + return min; > + else if (val > max) > + return max; > + else > + return val; > +} > + > +__global__ void to_jpeg_cuda(const unsigned char* src, unsigned char* ds= t, > + int pitch, int comp_id) > +{ > + int x =3D blockIdx.x * blockDim.x + threadIdx.x; > + int y =3D blockIdx.y * blockDim.y + threadIdx.y; > + int src_, dst_; > + > + // 8 bit -> 15 bit for better precision; > + src_ =3D static_cast(src[x + y * pitch]) << 7; > + > + // Conversion; > + dst_ =3D comp_id ? (min(src_, 30775) * 4663 - 9289992) >> 12 // c= hroma > + : (min(src_, 30189) * 19077 - 39057361) >> 14; // lum= a > + > + // Dither replacement; > + dst_ =3D dst_ + 64; > + > + // Back to 8 bit; > + dst_ =3D clamp(dst_ >> 7, jpeg_min[comp_id], jpeg_max[comp_id]); > + dst[x + y * pitch] =3D static_cast(dst_); > +} > + > +__global__ void to_mpeg_cuda(const unsigned char* src, unsigned char* ds= t, > + int pitch, int comp_id) > +{ > + int x =3D blockIdx.x * blockDim.x + threadIdx.x; > + int y =3D blockIdx.y * blockDim.y + threadIdx.y; > + int src_, dst_; > + > + // 8 bit -> 15 bit for better precision; > + src_ =3D static_cast(src[x + y * pitch]) << 7; > + > + // Conversion; > + dst_ =3D comp_id ? (src_ * 1799 + 4081085) >> 11 // chroma > + : (src_ * 14071 + 33561947) >> 14; // luma > + > + // Dither replacement; > + dst_ =3D dst_ + 64; > + > + // Back to 8 bit; > + dst_ =3D clamp(dst_ >> 7, mpeg_min[comp_id], mpeg_max[comp_id]); > + dst[x + y * pitch] =3D static_cast(dst_); > +} > +} > \ No newline at end of file > -- > 2.25.1 > Looks good to me on first glance otherwise, will give it a test soon _______________________________________________ ffmpeg-devel mailing list ffmpeg-devel@ffmpeg.org https://nam11.safelinks.protection.outlook.com/?url=3Dhttps%3A%2F%2Fffmpeg.= org%2Fmailman%2Flistinfo%2Fffmpeg-devel&data=3D05%7C01%7Crarzumanyan%40= nvidia.com%7Cd8c621ebee5e4c822cb708da932ec42e%7C43083d15727340c1b7db39efd9c= cc17a%7C0%7C0%7C637984126353127576%7CUnknown%7CTWFpbGZsb3d8eyJWIjoiMC4wLjAw= MDAiLCJQIjoiV2luMzIiLCJBTiI6Ik1haWwiLCJXVCI6Mn0%3D%7C3000%7C%7C%7C&sdat= a=3Dq%2FsiAKkOCBYVKn%2FVVsB2%2Fohu1%2FRw0YwHyuExmrcAlwY%3D&reserved=3D0 To unsubscribe, visit link above, or email ffmpeg-devel-request@ffmpeg.org with subject "unsubscribe". --_004_PH7PR12MB58312279BA34EDEDC705A0CFD2459PH7PR12MB5831namp_ Content-Type: text/x-patch; name="0001-libavfilter-vf_colorrange_cuda-CUDA-accelerated-colo.patch" Content-Description: 0001-libavfilter-vf_colorrange_cuda-CUDA-accelerated-colo.patch Content-Disposition: attachment; filename="0001-libavfilter-vf_colorrange_cuda-CUDA-accelerated-colo.patch"; size=19217; creation-date="Sun, 11 Sep 2022 07:27:28 GMT"; modification-date="Sun, 11 Sep 2022 07:27:34 GMT" Content-Transfer-Encoding: base64 RnJvbSA5YTI0OTBjOWVmMjgzOTkwMTc4OTJjMjAxNzhmMjMyZDE4NTZlNDE3IE1vbiBTZXAgMTcg MDA6MDA6MDAgMjAwMQpGcm9tOiBSb21hbiBBcnp1bWFueWFuIDxyYXJ6dW1hbnlhbkBudmlkaWEu Y29tPgpEYXRlOiBTYXQsIDEwIFNlcCAyMDIyIDExOjA1OjU2ICswMzAwClN1YmplY3Q6IFtQQVRD SF0gbGliYXZmaWx0ZXIvdmZfY29sb3JyYW5nZV9jdWRhIENVREEtYWNjZWxlcmF0ZWQgY29sb3Ig cmFuZ2UKIGNvbnZlcnNpb24gZmlsdGVyCgotLS0KIGNvbmZpZ3VyZSAgICAgICAgICAgICAgICAg ICAgICAgICB8ICAgMiArCiBsaWJhdmZpbHRlci9NYWtlZmlsZSAgICAgICAgICAgICAgfCAgIDMg KwogbGliYXZmaWx0ZXIvYWxsZmlsdGVycy5jICAgICAgICAgIHwgICAxICsKIGxpYmF2ZmlsdGVy L3ZmX2NvbG9ycmFuZ2VfY3VkYS5jICB8IDQ0MSArKysrKysrKysrKysrKysrKysrKysrKysrKysr KysKIGxpYmF2ZmlsdGVyL3ZmX2NvbG9ycmFuZ2VfY3VkYS5jdSB8ICA5MyArKysrKysrCiA1IGZp bGVzIGNoYW5nZWQsIDU0MCBpbnNlcnRpb25zKCspCiBjcmVhdGUgbW9kZSAxMDA2NDQgbGliYXZm aWx0ZXIvdmZfY29sb3JyYW5nZV9jdWRhLmMKIGNyZWF0ZSBtb2RlIDEwMDY0NCBsaWJhdmZpbHRl ci92Zl9jb2xvcnJhbmdlX2N1ZGEuY3UKCmRpZmYgLS1naXQgYS9jb25maWd1cmUgYi9jb25maWd1 cmUKaW5kZXggOWQ2NDU3ZDgxYi4uY2ExYjU0YWUyMyAxMDA3NTUKLS0tIGEvY29uZmlndXJlCisr KyBiL2NvbmZpZ3VyZQpAQCAtMzE0NCw2ICszMTQ0LDggQEAgdjRsMl9tMm1fZGVwcz0ibGludXhf dmlkZW9kZXYyX2ggc2VtX3RpbWVkd2FpdCIKIAogY2hyb21ha2V5X2N1ZGFfZmlsdGVyX2RlcHM9 ImZmbnZjb2RlYyIKIGNocm9tYWtleV9jdWRhX2ZpbHRlcl9kZXBzX2FueT0iY3VkYV9udmNjIGN1 ZGFfbGx2bSIKK2NvbG9ycmFuZ2VfY3VkYV9maWx0ZXJfZGVwcz0iZmZudmNvZGVjIgorY29sb3Jy YW5nZV9jdWRhX2ZpbHRlcl9kZXBzX2FueT0iY3VkYV9udmNjIGN1ZGFfbGx2bSIKIGh3dXBsb2Fk X2N1ZGFfZmlsdGVyX2RlcHM9ImZmbnZjb2RlYyIKIHNjYWxlX25wcF9maWx0ZXJfZGVwcz0iZmZu dmNvZGVjIGxpYm5wcCIKIHNjYWxlMnJlZl9ucHBfZmlsdGVyX2RlcHM9ImZmbnZjb2RlYyBsaWJu cHAiCmRpZmYgLS1naXQgYS9saWJhdmZpbHRlci9NYWtlZmlsZSBiL2xpYmF2ZmlsdGVyL01ha2Vm aWxlCmluZGV4IDMwY2MzMjlmYjYuLmQ5NWM2MDRkZWEgMTAwNjQ0Ci0tLSBhL2xpYmF2ZmlsdGVy L01ha2VmaWxlCisrKyBiL2xpYmF2ZmlsdGVyL01ha2VmaWxlCkBAIC0yMjgsNiArMjI4LDkgQEAg T0JKUy0kKENPTkZJR19DT0xPUkhPTERfRklMVEVSKSAgICAgICAgICAgICAgKz0gdmZfY29sb3Jr ZXkubwogT0JKUy0kKENPTkZJR19DT0xPUkxFVkVMU19GSUxURVIpICAgICAgICAgICAgKz0gdmZf Y29sb3JsZXZlbHMubwogT0JKUy0kKENPTkZJR19DT0xPUk1BUF9GSUxURVIpICAgICAgICAgICAg ICAgKz0gdmZfY29sb3JtYXAubwogT0JKUy0kKENPTkZJR19DT0xPUk1BVFJJWF9GSUxURVIpICAg ICAgICAgICAgKz0gdmZfY29sb3JtYXRyaXgubworT0JKUy0kKENPTkZJR19DT0xPUlJBTkdFX0NV REFfRklMVEVSKSAgICAgICAgKz0gdmZfY29sb3JyYW5nZV9jdWRhLm8gXAorICAgICAgICAgICAg ICAgICAgICAgICAgICAgICAgICAgICAgICAgICAgICAgICAgdmZfY29sb3JyYW5nZV9jdWRhLnB0 eC5vIFwKKyAgICAgICAgICAgICAgICAgICAgICAgICAgICAgICAgICAgICAgICAgICAgICAgIGN1 ZGEvbG9hZF9oZWxwZXIubwogT0JKUy0kKENPTkZJR19DT0xPUlNQQUNFX0ZJTFRFUikgICAgICAg ICAgICAgKz0gdmZfY29sb3JzcGFjZS5vIGNvbG9yc3BhY2Vkc3AubwogT0JKUy0kKENPTkZJR19D T0xPUlRFTVBFUkFUVVJFX0ZJTFRFUikgICAgICAgKz0gdmZfY29sb3J0ZW1wZXJhdHVyZS5vCiBP QkpTLSQoQ09ORklHX0NPTlZPTFVUSU9OX0ZJTFRFUikgICAgICAgICAgICArPSB2Zl9jb252b2x1 dGlvbi5vCmRpZmYgLS1naXQgYS9saWJhdmZpbHRlci9hbGxmaWx0ZXJzLmMgYi9saWJhdmZpbHRl ci9hbGxmaWx0ZXJzLmMKaW5kZXggNWViYWNmZGUyNy4uNWU5Y2JlNTdlYyAxMDA2NDQKLS0tIGEv bGliYXZmaWx0ZXIvYWxsZmlsdGVycy5jCisrKyBiL2xpYmF2ZmlsdGVyL2FsbGZpbHRlcnMuYwpA QCAtMjEzLDYgKzIxMyw3IEBAIGV4dGVybiBjb25zdCBBVkZpbHRlciBmZl92Zl9jb2xvcm1hcDsK IGV4dGVybiBjb25zdCBBVkZpbHRlciBmZl92Zl9jb2xvcm1hdHJpeDsKIGV4dGVybiBjb25zdCBB VkZpbHRlciBmZl92Zl9jb2xvcnNwYWNlOwogZXh0ZXJuIGNvbnN0IEFWRmlsdGVyIGZmX3ZmX2Nv bG9ydGVtcGVyYXR1cmU7CitleHRlcm4gY29uc3QgQVZGaWx0ZXIgZmZfdmZfY29sb3JyYW5nZV9j dWRhOwogZXh0ZXJuIGNvbnN0IEFWRmlsdGVyIGZmX3ZmX2NvbnZvbHV0aW9uOwogZXh0ZXJuIGNv bnN0IEFWRmlsdGVyIGZmX3ZmX2NvbnZvbHV0aW9uX29wZW5jbDsKIGV4dGVybiBjb25zdCBBVkZp bHRlciBmZl92Zl9jb252b2x2ZTsKZGlmZiAtLWdpdCBhL2xpYmF2ZmlsdGVyL3ZmX2NvbG9ycmFu Z2VfY3VkYS5jIGIvbGliYXZmaWx0ZXIvdmZfY29sb3JyYW5nZV9jdWRhLmMKbmV3IGZpbGUgbW9k ZSAxMDA2NDQKaW5kZXggMDAwMDAwMDAwMC4uYzAwZmZmZWFlMwotLS0gL2Rldi9udWxsCisrKyBi L2xpYmF2ZmlsdGVyL3ZmX2NvbG9ycmFuZ2VfY3VkYS5jCkBAIC0wLDAgKzEsNDQxIEBACisvKgor ICogQ29weXJpZ2h0IChjKSAyMDIyLCBOVklESUEgQ09SUE9SQVRJT04uIEFsbCByaWdodHMgcmVz ZXJ2ZWQuCisgKgorICogUGVybWlzc2lvbiBpcyBoZXJlYnkgZ3JhbnRlZCwgZnJlZSBvZiBjaGFy Z2UsIHRvIGFueSBwZXJzb24gb2J0YWluaW5nIGEKKyAqIGNvcHkgb2YgdGhpcyBzb2Z0d2FyZSBh bmQgYXNzb2NpYXRlZCBkb2N1bWVudGF0aW9uIGZpbGVzICh0aGUgIlNvZnR3YXJlIiksCisgKiB0 byBkZWFsIGluIHRoZSBTb2Z0d2FyZSB3aXRob3V0IHJlc3RyaWN0aW9uLCBpbmNsdWRpbmcgd2l0 aG91dCBsaW1pdGF0aW9uCisgKiB0aGUgcmlnaHRzIHRvIHVzZSwgY29weSwgbW9kaWZ5LCBtZXJn ZSwgcHVibGlzaCwgZGlzdHJpYnV0ZSwgc3VibGljZW5zZSwKKyAqIGFuZC9vciBzZWxsIGNvcGll cyBvZiB0aGUgU29mdHdhcmUsIGFuZCB0byBwZXJtaXQgcGVyc29ucyB0byB3aG9tIHRoZQorICog U29mdHdhcmUgaXMgZnVybmlzaGVkIHRvIGRvIHNvLCBzdWJqZWN0IHRvIHRoZSBmb2xsb3dpbmcg Y29uZGl0aW9uczoKKyAqCisgKiBUaGUgYWJvdmUgY29weXJpZ2h0IG5vdGljZSBhbmQgdGhpcyBw ZXJtaXNzaW9uIG5vdGljZSBzaGFsbCBiZSBpbmNsdWRlZCBpbgorICogYWxsIGNvcGllcyBvciBz dWJzdGFudGlhbCBwb3J0aW9ucyBvZiB0aGUgU29mdHdhcmUuCisgKgorICogVEhFIFNPRlRXQVJF IElTIFBST1ZJREVEICJBUyBJUyIsIFdJVEhPVVQgV0FSUkFOVFkgT0YgQU5ZIEtJTkQsIEVYUFJF U1MgT1IKKyAqIElNUExJRUQsIElOQ0xVRElORyBCVVQgTk9UIExJTUlURUQgVE8gVEhFIFdBUlJB TlRJRVMgT0YgTUVSQ0hBTlRBQklMSVRZLAorICogRklUTkVTUyBGT1IgQSBQQVJUSUNVTEFSIFBV UlBPU0UgQU5EIE5PTklORlJJTkdFTUVOVC4gIElOIE5PIEVWRU5UIFNIQUxMCisgKiBUSEUgQVVU SE9SUyBPUiBDT1BZUklHSFQgSE9MREVSUyBCRSBMSUFCTEUgRk9SIEFOWSBDTEFJTSwgREFNQUdF UyBPUiBPVEhFUgorICogTElBQklMSVRZLCBXSEVUSEVSIElOIEFOIEFDVElPTiBPRiBDT05UUkFD VCwgVE9SVCBPUiBPVEhFUldJU0UsIEFSSVNJTkcKKyAqIEZST00sIE9VVCBPRiBPUiBJTiBDT05O RUNUSU9OIFdJVEggVEhFIFNPRlRXQVJFIE9SIFRIRSBVU0UgT1IgT1RIRVIKKyAqIERFQUxJTkdT IElOIFRIRSBTT0ZUV0FSRS4KKyAqLworCisjaW5jbHVkZSA8c3RyaW5nLmg+CisKKyNpbmNsdWRl ICJsaWJhdnV0aWwvYXZzdHJpbmcuaCIKKyNpbmNsdWRlICJsaWJhdnV0aWwvY29tbW9uLmgiCisj aW5jbHVkZSAibGliYXZ1dGlsL2N1ZGFfY2hlY2suaCIKKyNpbmNsdWRlICJsaWJhdnV0aWwvaHdj b250ZXh0LmgiCisjaW5jbHVkZSAibGliYXZ1dGlsL2h3Y29udGV4dF9jdWRhX2ludGVybmFsLmgi CisjaW5jbHVkZSAibGliYXZ1dGlsL2ludGVybmFsLmgiCisjaW5jbHVkZSAibGliYXZ1dGlsL29w dC5oIgorI2luY2x1ZGUgImxpYmF2dXRpbC9waXhkZXNjLmgiCisKKyNpbmNsdWRlICJhdmZpbHRl ci5oIgorI2luY2x1ZGUgImZvcm1hdHMuaCIKKyNpbmNsdWRlICJpbnRlcm5hbC5oIgorI2luY2x1 ZGUgInNjYWxlX2V2YWwuaCIKKyNpbmNsdWRlICJ2aWRlby5oIgorCisjaW5jbHVkZSAiY3VkYS9s b2FkX2hlbHBlci5oIgorCitzdGF0aWMgY29uc3QgZW51bSBBVlBpeGVsRm9ybWF0IHN1cHBvcnRl ZF9mb3JtYXRzW10gPSB7CisgICAgQVZfUElYX0ZNVF9OVjEyLAorICAgIEFWX1BJWF9GTVRfWVVW NDIwUCwKKyAgICBBVl9QSVhfRk1UX1lVVjQ0NFAsCit9OworCisjZGVmaW5lIERJVl9VUChhLCBi KSAoKChhKSArIChiKS0xKSAvIChiKSkKKyNkZWZpbmUgQkxPQ0tYIDMyCisjZGVmaW5lIEJMT0NL WSAxNgorCisjZGVmaW5lIENIRUNLX0NVKHgpIEZGX0NVREFfQ0hFQ0tfREwoY3R4LCBzLT5od2N0 eC0+aW50ZXJuYWwtPmN1ZGFfZGwsIHgpCisKK3R5cGVkZWYgc3RydWN0IENVREFDb252Q29udGV4 dCB7CisgICAgY29uc3QgQVZDbGFzcyogY2xhc3M7CisKKyAgICBBVkNVREFEZXZpY2VDb250ZXh0 KiBod2N0eDsKKyAgICBBVkJ1ZmZlclJlZiogZnJhbWVzX2N0eDsKKyAgICBBVkZyYW1lKiBvd25f ZnJhbWU7CisgICAgQVZGcmFtZSogdG1wX2ZyYW1lOworCisgICAgQ1Vjb250ZXh0IGN1X2N0eDsK KyAgICBDVXN0cmVhbSBjdV9zdHJlYW07CisgICAgQ1Vtb2R1bGUgY3VfbW9kdWxlOworICAgIENV ZnVuY3Rpb24gY3VfY29udmVydFtBVkNPTF9SQU5HRV9OQl07CisKKyAgICBlbnVtIEFWUGl4ZWxG b3JtYXQgcGl4X2ZtdDsKKyAgICBlbnVtIEFWQ29sb3JSYW5nZSByYW5nZTsKKworICAgIGludCBu dW1fcGxhbmVzOworfSBDVURBQ29udkNvbnRleHQ7CisKK3N0YXRpYyBhdl9jb2xkIGludCBjdWRh Y29udl9pbml0KEFWRmlsdGVyQ29udGV4dCogY3R4KQoreworICAgIENVREFDb252Q29udGV4dCog cyA9IGN0eC0+cHJpdjsKKworICAgIHMtPm93bl9mcmFtZSA9IGF2X2ZyYW1lX2FsbG9jKCk7Cisg ICAgaWYgKCFzLT5vd25fZnJhbWUpCisgICAgICAgIHJldHVybiBBVkVSUk9SKEVOT01FTSk7CisK KyAgICBzLT50bXBfZnJhbWUgPSBhdl9mcmFtZV9hbGxvYygpOworICAgIGlmICghcy0+dG1wX2Zy YW1lKQorICAgICAgICByZXR1cm4gQVZFUlJPUihFTk9NRU0pOworCisgICAgcmV0dXJuIDA7Cit9 CisKK3N0YXRpYyBhdl9jb2xkIHZvaWQgY3VkYWNvbnZfdW5pbml0KEFWRmlsdGVyQ29udGV4dCog Y3R4KQoreworICAgIENVREFDb252Q29udGV4dCogcyA9IGN0eC0+cHJpdjsKKworICAgIGlmIChz LT5od2N0eCAmJiBzLT5jdV9tb2R1bGUpIHsKKyAgICAgICAgQ3VkYUZ1bmN0aW9ucyogY3UgPSBz LT5od2N0eC0+aW50ZXJuYWwtPmN1ZGFfZGw7CisgICAgICAgIENVY29udGV4dCBkdW1teTsKKwor ICAgICAgICBDSEVDS19DVShjdS0+Y3VDdHhQdXNoQ3VycmVudChzLT5od2N0eC0+Y3VkYV9jdHgp KTsKKyAgICAgICAgQ0hFQ0tfQ1UoY3UtPmN1TW9kdWxlVW5sb2FkKHMtPmN1X21vZHVsZSkpOwor ICAgICAgICBzLT5jdV9tb2R1bGUgPSBOVUxMOworICAgICAgICBDSEVDS19DVShjdS0+Y3VDdHhQ b3BDdXJyZW50KCZkdW1teSkpOworICAgIH0KKworICAgIGF2X2ZyYW1lX2ZyZWUoJnMtPm93bl9m cmFtZSk7CisgICAgYXZfYnVmZmVyX3VucmVmKCZzLT5mcmFtZXNfY3R4KTsKKyAgICBhdl9mcmFt ZV9mcmVlKCZzLT50bXBfZnJhbWUpOworfQorCitzdGF0aWMgYXZfY29sZCBpbnQgaW5pdF9od2Zy YW1lX2N0eChDVURBQ29udkNvbnRleHQqIHMsIEFWQnVmZmVyUmVmKiBkZXZpY2VfY3R4LAorICAg ICAgICAgICAgICAgICAgICAgICAgICAgICAgICAgICAgaW50IHdpZHRoLCBpbnQgaGVpZ2h0KQor eworICAgIEFWQnVmZmVyUmVmKiBvdXRfcmVmID0gTlVMTDsKKyAgICBBVkhXRnJhbWVzQ29udGV4 dCogb3V0X2N0eDsKKyAgICBpbnQgcmV0OworCisgICAgb3V0X3JlZiA9IGF2X2h3ZnJhbWVfY3R4 X2FsbG9jKGRldmljZV9jdHgpOworICAgIGlmICghb3V0X3JlZikKKyAgICAgICAgcmV0dXJuIEFW RVJST1IoRU5PTUVNKTsKKworICAgIG91dF9jdHggPSAoQVZIV0ZyYW1lc0NvbnRleHQqKW91dF9y ZWYtPmRhdGE7CisKKyAgICBvdXRfY3R4LT5mb3JtYXQgPSBBVl9QSVhfRk1UX0NVREE7CisgICAg b3V0X2N0eC0+c3dfZm9ybWF0ID0gcy0+cGl4X2ZtdDsKKyAgICBvdXRfY3R4LT53aWR0aCA9IEZG QUxJR04od2lkdGgsIDMyKTsKKyAgICBvdXRfY3R4LT5oZWlnaHQgPSBGRkFMSUdOKGhlaWdodCwg MzIpOworCisgICAgcmV0ID0gYXZfaHdmcmFtZV9jdHhfaW5pdChvdXRfcmVmKTsKKyAgICBpZiAo cmV0IDwgMCkKKyAgICAgICAgZ290byBmYWlsOworCisgICAgYXZfZnJhbWVfdW5yZWYocy0+b3du X2ZyYW1lKTsKKyAgICByZXQgPSBhdl9od2ZyYW1lX2dldF9idWZmZXIob3V0X3JlZiwgcy0+b3du X2ZyYW1lLCAwKTsKKyAgICBpZiAocmV0IDwgMCkKKyAgICAgICAgZ290byBmYWlsOworCisgICAg cy0+b3duX2ZyYW1lLT53aWR0aCA9IHdpZHRoOworICAgIHMtPm93bl9mcmFtZS0+aGVpZ2h0ID0g aGVpZ2h0OworCisgICAgYXZfYnVmZmVyX3VucmVmKCZzLT5mcmFtZXNfY3R4KTsKKyAgICBzLT5m cmFtZXNfY3R4ID0gb3V0X3JlZjsKKworICAgIHJldHVybiAwOworZmFpbDoKKyAgICBhdl9idWZm ZXJfdW5yZWYoJm91dF9yZWYpOworICAgIHJldHVybiByZXQ7Cit9CisKK3N0YXRpYyBpbnQgZm9y bWF0X2lzX3N1cHBvcnRlZChlbnVtIEFWUGl4ZWxGb3JtYXQgZm10KQoreworICAgIGZvciAoaW50 IGkgPSAwOyBpIDwgRkZfQVJSQVlfRUxFTVMoc3VwcG9ydGVkX2Zvcm1hdHMpOyBpKyspCisgICAg ICAgIGlmIChmbXQgPT0gc3VwcG9ydGVkX2Zvcm1hdHNbaV0pCisgICAgICAgICAgICByZXR1cm4g MTsKKworICAgIHJldHVybiAwOworfQorCitzdGF0aWMgYXZfY29sZCBpbnQgaW5pdF9wcm9jZXNz aW5nX2NoYWluKEFWRmlsdGVyQ29udGV4dCogY3R4LCBpbnQgd2lkdGgsCisgICAgICAgICAgICAg ICAgICAgICAgICAgICAgICAgICAgICAgICAgIGludCBoZWlnaHQpCit7CisgICAgQ1VEQUNvbnZD b250ZXh0KiBzID0gY3R4LT5wcml2OworICAgIEFWSFdGcmFtZXNDb250ZXh0KiBpbl9mcmFtZXNf Y3R4OworCisgICAgaW50IHJldDsKKworICAgIGlmICghY3R4LT5pbnB1dHNbMF0tPmh3X2ZyYW1l c19jdHgpIHsKKyAgICAgICAgYXZfbG9nKGN0eCwgQVZfTE9HX0VSUk9SLCAiTm8gaHcgY29udGV4 dCBwcm92aWRlZCBvbiBpbnB1dFxuIik7CisgICAgICAgIHJldHVybiBBVkVSUk9SKEVJTlZBTCk7 CisgICAgfQorCisgICAgaW5fZnJhbWVzX2N0eCA9IChBVkhXRnJhbWVzQ29udGV4dCopY3R4LT5p bnB1dHNbMF0tPmh3X2ZyYW1lc19jdHgtPmRhdGE7CisgICAgcy0+cGl4X2ZtdCA9IGluX2ZyYW1l c19jdHgtPnN3X2Zvcm1hdDsKKworICAgIGlmICghZm9ybWF0X2lzX3N1cHBvcnRlZChzLT5waXhf Zm10KSkgeworICAgICAgICBhdl9sb2coY3R4LCBBVl9MT0dfRVJST1IsICJVbnN1cHBvcnRlZCBw aXhlbCBmb3JtYXQ6ICVzXG4iLAorICAgICAgICAgICAgICAgYXZfZ2V0X3BpeF9mbXRfbmFtZShz LT5waXhfZm10KSk7CisgICAgICAgIHJldHVybiBBVkVSUk9SKEVJTlZBTCk7CisgICAgfQorCisg ICAgaWYgKChBVkNPTF9SQU5HRV9NUEVHICE9IHMtPnJhbmdlKSAmJiAoQVZDT0xfUkFOR0VfSlBF RyAhPSBzLT5yYW5nZSkpIHsKKyAgICAgICAgYXZfbG9nKGN0eCwgQVZfTE9HX0VSUk9SLCAiVW5z dXBwb3J0ZWQgY29sb3IgcmFuZ2VcbiIpOworICAgICAgICByZXR1cm4gQVZFUlJPUihFSU5WQUwp OworICAgIH0KKworICAgIHMtPm51bV9wbGFuZXMgPSBhdl9waXhfZm10X2NvdW50X3BsYW5lcyhz LT5waXhfZm10KTsKKworICAgIHJldCA9IGluaXRfaHdmcmFtZV9jdHgocywgaW5fZnJhbWVzX2N0 eC0+ZGV2aWNlX3JlZiwgd2lkdGgsIGhlaWdodCk7CisgICAgaWYgKHJldCA8IDApCisgICAgICAg IHJldHVybiByZXQ7CisKKyAgICBjdHgtPm91dHB1dHNbMF0tPmh3X2ZyYW1lc19jdHggPSBhdl9i dWZmZXJfcmVmKHMtPmZyYW1lc19jdHgpOworICAgIGlmICghY3R4LT5vdXRwdXRzWzBdLT5od19m cmFtZXNfY3R4KQorICAgICAgICByZXR1cm4gQVZFUlJPUihFTk9NRU0pOworCisgICAgcmV0dXJu IDA7Cit9CisKK3N0YXRpYyBhdl9jb2xkIGludCBjdWRhY29udl9sb2FkX2Z1bmN0aW9ucyhBVkZp bHRlckNvbnRleHQqIGN0eCkKK3sKKyAgICBDVURBQ29udkNvbnRleHQqIHMgPSBjdHgtPnByaXY7 CisgICAgQ1Vjb250ZXh0IGR1bW15LCBjdWRhX2N0eCA9IHMtPmh3Y3R4LT5jdWRhX2N0eDsKKyAg ICBDdWRhRnVuY3Rpb25zKiBjdSA9IHMtPmh3Y3R4LT5pbnRlcm5hbC0+Y3VkYV9kbDsKKyAgICBp bnQgcmV0OworCisgICAgZXh0ZXJuIGNvbnN0IHVuc2lnbmVkIGNoYXIgZmZfdmZfY29sb3JyYW5n ZV9jdWRhX3B0eF9kYXRhW107CisgICAgZXh0ZXJuIGNvbnN0IHVuc2lnbmVkIGludCBmZl92Zl9j b2xvcnJhbmdlX2N1ZGFfcHR4X2xlbjsKKworICAgIHJldCA9IENIRUNLX0NVKGN1LT5jdUN0eFB1 c2hDdXJyZW50KGN1ZGFfY3R4KSk7CisgICAgaWYgKHJldCA8IDApCisgICAgICAgIHJldHVybiBy ZXQ7CisKKyAgICByZXQgPSBmZl9jdWRhX2xvYWRfbW9kdWxlKGN0eCwgcy0+aHdjdHgsICZzLT5j dV9tb2R1bGUsCisgICAgICAgICAgICAgICAgICAgICAgICAgICAgICBmZl92Zl9jb2xvcnJhbmdl X2N1ZGFfcHR4X2RhdGEsCisgICAgICAgICAgICAgICAgICAgICAgICAgICAgICBmZl92Zl9jb2xv cnJhbmdlX2N1ZGFfcHR4X2xlbik7CisgICAgaWYgKHJldCA8IDApCisgICAgICAgIGdvdG8gZmFp bDsKKworICAgIHJldCA9IENIRUNLX0NVKGN1LT5jdU1vZHVsZUdldEZ1bmN0aW9uKAorICAgICAg ICAmcy0+Y3VfY29udmVydFtBVkNPTF9SQU5HRV9NUEVHXSwgcy0+Y3VfbW9kdWxlLAorICAgICAg ICAidG9fbXBlZ19jdWRhIikpOworCisgICAgaWYgKHJldCA8IDApCisgICAgICAgIGdvdG8gZmFp bDsKKworICAgIHJldCA9IENIRUNLX0NVKGN1LT5jdU1vZHVsZUdldEZ1bmN0aW9uKAorICAgICAg ICAmcy0+Y3VfY29udmVydFtBVkNPTF9SQU5HRV9KUEVHXSwgcy0+Y3VfbW9kdWxlLAorICAgICAg ICAidG9fanBlZ19jdWRhIikpOworCisgICAgaWYgKHJldCA8IDApCisgICAgICAgIGdvdG8gZmFp bDsKKworZmFpbDoKKyAgICBDSEVDS19DVShjdS0+Y3VDdHhQb3BDdXJyZW50KCZkdW1teSkpOwor ICAgIHJldHVybiByZXQ7Cit9CisKK3N0YXRpYyBhdl9jb2xkIGludCBjdWRhY29udl9jb25maWdf cHJvcHMoQVZGaWx0ZXJMaW5rKiBvdXRsaW5rKQoreworICAgIEFWRmlsdGVyQ29udGV4dCogY3R4 ID0gb3V0bGluay0+c3JjOworICAgIEFWRmlsdGVyTGluayogaW5saW5rID0gb3V0bGluay0+c3Jj LT5pbnB1dHNbMF07CisgICAgQ1VEQUNvbnZDb250ZXh0KiBzID0gY3R4LT5wcml2OworICAgIEFW SFdGcmFtZXNDb250ZXh0KiBmcmFtZXNfY3R4ID0KKyAgICAgICAgKEFWSFdGcmFtZXNDb250ZXh0 KilpbmxpbmstPmh3X2ZyYW1lc19jdHgtPmRhdGE7CisgICAgQVZDVURBRGV2aWNlQ29udGV4dCog ZGV2aWNlX2h3Y3R4ID0gZnJhbWVzX2N0eC0+ZGV2aWNlX2N0eC0+aHdjdHg7CisgICAgaW50IHJl dDsKKworICAgIHMtPmh3Y3R4ID0gZGV2aWNlX2h3Y3R4OworICAgIHMtPmN1X3N0cmVhbSA9IHMt Pmh3Y3R4LT5zdHJlYW07CisKKyAgICBvdXRsaW5rLT53ID0gaW5saW5rLT53OworICAgIG91dGxp bmstPmggPSBpbmxpbmstPmg7CisKKyAgICByZXQgPSBpbml0X3Byb2Nlc3NpbmdfY2hhaW4oY3R4 LCBpbmxpbmstPncsIGlubGluay0+aCk7CisgICAgaWYgKHJldCA8IDApCisgICAgICAgIHJldHVy biByZXQ7CisKKyAgICBpZiAoaW5saW5rLT5zYW1wbGVfYXNwZWN0X3JhdGlvLm51bSkgeworICAg ICAgICBvdXRsaW5rLT5zYW1wbGVfYXNwZWN0X3JhdGlvID0gYXZfbXVsX3EoCisgICAgICAgICAg ICAoQVZSYXRpb25hbCl7b3V0bGluay0+aCAqIGlubGluay0+dywgb3V0bGluay0+dyAqIGlubGlu ay0+aH0sCisgICAgICAgICAgICBpbmxpbmstPnNhbXBsZV9hc3BlY3RfcmF0aW8pOworICAgIH0g ZWxzZSB7CisgICAgICAgIG91dGxpbmstPnNhbXBsZV9hc3BlY3RfcmF0aW8gPSBpbmxpbmstPnNh bXBsZV9hc3BlY3RfcmF0aW87CisgICAgfQorCisgICAgcmV0ID0gY3VkYWNvbnZfbG9hZF9mdW5j dGlvbnMoY3R4KTsKKyAgICBpZiAocmV0IDwgMCkKKyAgICAgICAgcmV0dXJuIHJldDsKKworICAg IHJldHVybiByZXQ7Cit9CisKK3N0YXRpYyBpbnQgY29udl9jdWRhX2NvbnZlcnQoQVZGaWx0ZXJD b250ZXh0KiBjdHgsIEFWRnJhbWUqIG91dCwgQVZGcmFtZSogaW4pCit7CisgICAgQ1VEQUNvbnZD b250ZXh0KiBzID0gY3R4LT5wcml2OworICAgIEN1ZGFGdW5jdGlvbnMqIGN1ID0gcy0+aHdjdHgt PmludGVybmFsLT5jdWRhX2RsOworICAgIENVY29udGV4dCBkdW1teSwgY3VkYV9jdHggPSBzLT5o d2N0eC0+Y3VkYV9jdHg7CisgICAgaW50IHJldDsKKworICAgIHJldCA9IENIRUNLX0NVKGN1LT5j dUN0eFB1c2hDdXJyZW50KGN1ZGFfY3R4KSk7CisgICAgaWYgKHJldCA8IDApCisgICAgICAgIHJl dHVybiByZXQ7CisKKyAgICBvdXQtPmNvbG9yX3JhbmdlID0gcy0+cmFuZ2U7CisKKyAgICBmb3Ig KGludCBpID0gMDsgaSA8IHMtPm51bV9wbGFuZXM7IGkrKykgeworICAgICAgICBpbnQgd2lkdGgg PSBpbi0+d2lkdGgsIGhlaWdodCA9IGluLT5oZWlnaHQsIGNvbXBfaWQgPSAoaSA+IDApOworCisg ICAgICAgIHN3aXRjaCAocy0+cGl4X2ZtdCkgeworICAgICAgICBjYXNlIEFWX1BJWF9GTVRfWVVW NDQ0UDoKKyAgICAgICAgICAgIGJyZWFrOworICAgICAgICBjYXNlIEFWX1BJWF9GTVRfWVVWNDIw UDoKKyAgICAgICAgICAgIHdpZHRoID0gY29tcF9pZCA/IGluLT53aWR0aCAvIDIgOiBpbi0+d2lk dGg7CisgICAgICAgIGNhc2UgQVZfUElYX0ZNVF9OVjEyOgorICAgICAgICAgICAgaGVpZ2h0ID0g Y29tcF9pZCA/IGluLT5oZWlnaHQgLyAyIDogaW4tPmhlaWdodDsKKyAgICAgICAgICAgIGJyZWFr OworICAgICAgICBkZWZhdWx0OgorICAgICAgICAgICAgYXZfbG9nKGN0eCwgQVZfTE9HX0VSUk9S LCAiVW5zdXBwb3J0ZWQgcGl4ZWwgZm9ybWF0OiAlc1xuIiwKKyAgICAgICAgICAgICAgICAgICBh dl9nZXRfcGl4X2ZtdF9uYW1lKHMtPnBpeF9mbXQpKTsKKyAgICAgICAgICAgIHJldHVybiBBVkVS Uk9SKEVJTlZBTCk7CisgICAgICAgIH0KKworICAgICAgICBpZiAoIXMtPmN1X2NvbnZlcnRbb3V0 LT5jb2xvcl9yYW5nZV0pIHsKKyAgICAgICAgICAgIGF2X2xvZyhjdHgsIEFWX0xPR19FUlJPUiwg IlVuc3VwcG9ydGVkIGNvbG9yIHJhbmdlXG4iKTsKKyAgICAgICAgICAgIHJldHVybiBBVkVSUk9S KEVJTlZBTCk7CisgICAgICAgIH0KKworICAgICAgICBpZiAoaW4tPmNvbG9yX3JhbmdlICE9IG91 dC0+Y29sb3JfcmFuZ2UpIHsKKyAgICAgICAgICAgIHZvaWQqIGFyZ3NbXSA9IHsmaW4tPmRhdGFb aV0sICZvdXQtPmRhdGFbaV0sICZpbi0+bGluZXNpemVbaV0sCisgICAgICAgICAgICAgICAgICAg ICAgICAgICAgJmNvbXBfaWR9OworICAgICAgICAgICAgcmV0ID0gQ0hFQ0tfQ1UoY3UtPmN1TGF1 bmNoS2VybmVsKAorICAgICAgICAgICAgICAgIHMtPmN1X2NvbnZlcnRbb3V0LT5jb2xvcl9yYW5n ZV0sIERJVl9VUCh3aWR0aCwgQkxPQ0tYKSwKKyAgICAgICAgICAgICAgICBESVZfVVAoaGVpZ2h0 LCBCTE9DS1kpLCAxLCBCTE9DS1gsIEJMT0NLWSwgMSwgMCwgcy0+Y3Vfc3RyZWFtLAorICAgICAg ICAgICAgICAgIGFyZ3MsIE5VTEwpKTsKKyAgICAgICAgfSBlbHNlIHsKKyAgICAgICAgICAgIHJl dCA9IGF2X2h3ZnJhbWVfdHJhbnNmZXJfZGF0YShvdXQsIGluLCAwKTsKKyAgICAgICAgICAgIGlm IChyZXQgPCAwKQorICAgICAgICAgICAgICAgIHJldHVybiByZXQ7CisgICAgICAgIH0KKyAgICB9 CisKKyAgICBDSEVDS19DVShjdS0+Y3VDdHhQb3BDdXJyZW50KCZkdW1teSkpOworICAgIHJldHVy biByZXQ7Cit9CisKK3N0YXRpYyBpbnQgY3VkYWNvbnZfY29udihBVkZpbHRlckNvbnRleHQqIGN0 eCwgQVZGcmFtZSogb3V0LCBBVkZyYW1lKiBpbikKK3sKKyAgICBDVURBQ29udkNvbnRleHQqIHMg PSBjdHgtPnByaXY7CisgICAgQVZGaWx0ZXJMaW5rKiBvdXRsaW5rID0gY3R4LT5vdXRwdXRzWzBd OworICAgIEFWRnJhbWUqIHNyYyA9IGluOworICAgIGludCByZXQ7CisKKyAgICByZXQgPSBjb252 X2N1ZGFfY29udmVydChjdHgsIHMtPm93bl9mcmFtZSwgc3JjKTsKKyAgICBpZiAocmV0IDwgMCkK KyAgICAgICAgcmV0dXJuIHJldDsKKworICAgIHNyYyA9IHMtPm93bl9mcmFtZTsKKyAgICByZXQg PSBhdl9od2ZyYW1lX2dldF9idWZmZXIoc3JjLT5od19mcmFtZXNfY3R4LCBzLT50bXBfZnJhbWUs IDApOworICAgIGlmIChyZXQgPCAwKQorICAgICAgICByZXR1cm4gcmV0OworCisgICAgYXZfZnJh bWVfbW92ZV9yZWYob3V0LCBzLT5vd25fZnJhbWUpOworICAgIGF2X2ZyYW1lX21vdmVfcmVmKHMt Pm93bl9mcmFtZSwgcy0+dG1wX2ZyYW1lKTsKKworICAgIHMtPm93bl9mcmFtZS0+d2lkdGggPSBv dXRsaW5rLT53OworICAgIHMtPm93bl9mcmFtZS0+aGVpZ2h0ID0gb3V0bGluay0+aDsKKworICAg IHJldCA9IGF2X2ZyYW1lX2NvcHlfcHJvcHMob3V0LCBpbik7CisgICAgaWYgKHJldCA8IDApCisg ICAgICAgIHJldHVybiByZXQ7CisKKyAgICByZXR1cm4gMDsKK30KKworc3RhdGljIGludCBjdWRh Y29udl9maWx0ZXJfZnJhbWUoQVZGaWx0ZXJMaW5rKiBsaW5rLCBBVkZyYW1lKiBpbikKK3sKKyAg ICBBVkZpbHRlckNvbnRleHQqIGN0eCA9IGxpbmstPmRzdDsKKyAgICBDVURBQ29udkNvbnRleHQq IHMgPSBjdHgtPnByaXY7CisgICAgQVZGaWx0ZXJMaW5rKiBvdXRsaW5rID0gY3R4LT5vdXRwdXRz WzBdOworICAgIEN1ZGFGdW5jdGlvbnMqIGN1ID0gcy0+aHdjdHgtPmludGVybmFsLT5jdWRhX2Rs OworCisgICAgQVZGcmFtZSogb3V0ID0gTlVMTDsKKyAgICBDVWNvbnRleHQgZHVtbXk7CisgICAg aW50IHJldCA9IDA7CisKKyAgICBvdXQgPSBhdl9mcmFtZV9hbGxvYygpOworICAgIGlmICghb3V0 KSB7CisgICAgICAgIHJldCA9IEFWRVJST1IoRU5PTUVNKTsKKyAgICAgICAgZ290byBmYWlsOwor ICAgIH0KKworICAgIHJldCA9IENIRUNLX0NVKGN1LT5jdUN0eFB1c2hDdXJyZW50KHMtPmh3Y3R4 LT5jdWRhX2N0eCkpOworICAgIGlmIChyZXQgPCAwKQorICAgICAgICBnb3RvIGZhaWw7CisKKyAg ICByZXQgPSBjdWRhY29udl9jb252KGN0eCwgb3V0LCBpbik7CisKKyAgICBDSEVDS19DVShjdS0+ Y3VDdHhQb3BDdXJyZW50KCZkdW1teSkpOworICAgIGlmIChyZXQgPCAwKQorICAgICAgICBnb3Rv IGZhaWw7CisKKyAgICBhdl9yZWR1Y2UoJm91dC0+c2FtcGxlX2FzcGVjdF9yYXRpby5udW0sICZv dXQtPnNhbXBsZV9hc3BlY3RfcmF0aW8uZGVuLAorICAgICAgICAgICAgICAoaW50NjRfdClpbi0+ c2FtcGxlX2FzcGVjdF9yYXRpby5udW0gKiBvdXRsaW5rLT5oICogbGluay0+dywKKyAgICAgICAg ICAgICAgKGludDY0X3QpaW4tPnNhbXBsZV9hc3BlY3RfcmF0aW8uZGVuICogb3V0bGluay0+dyAq IGxpbmstPmgsCisgICAgICAgICAgICAgIElOVF9NQVgpOworCisgICAgYXZfZnJhbWVfZnJlZSgm aW4pOworICAgIHJldHVybiBmZl9maWx0ZXJfZnJhbWUob3V0bGluaywgb3V0KTsKK2ZhaWw6Cisg ICAgYXZfZnJhbWVfZnJlZSgmaW4pOworICAgIGF2X2ZyYW1lX2ZyZWUoJm91dCk7CisgICAgcmV0 dXJuIHJldDsKK30KKworI2RlZmluZSBPRkZTRVQoeCkgb2Zmc2V0b2YoQ1VEQUNvbnZDb250ZXh0 LCB4KQorI2RlZmluZSBGTEFHUyAoQVZfT1BUX0ZMQUdfRklMVEVSSU5HX1BBUkFNIHwgQVZfT1BU X0ZMQUdfVklERU9fUEFSQU0pCitzdGF0aWMgY29uc3QgQVZPcHRpb24gb3B0aW9uc1tdID0gewor ICAgIHsicmFuZ2UiLCAiT3V0cHV0IHZpZGVvIHJhbmdlIiwgT0ZGU0VUKHJhbmdlKSwgQVZfT1BU X1RZUEVfSU5ULCB7Lmk2NCA9IEFWQ09MX1JBTkdFX1VOU1BFQ0lGSUVEfSwgQVZDT0xfUkFOR0Vf VU5TUEVDSUZJRUQsIEFWQ09MX1JBTkdFX05CIC0gMSwgRkxBR1MsICJyYW5nZSJ9LAorICAgICAg ICB7Im1wZWciLCAibGltaXRlZCByYW5nZSIsIDAsIEFWX09QVF9UWVBFX0NPTlNULCB7Lmk2NCA9 IEFWQ09MX1JBTkdFX01QRUd9LCAwLCAwLCBGTEFHUywgInJhbmdlIn0sCisgICAgICAgIHsianBl ZyIsICJmdWxsIHJhbmdlIiwgICAgMCwgQVZfT1BUX1RZUEVfQ09OU1QsIHsuaTY0ID0gQVZDT0xf UkFOR0VfSlBFR30sIDAsIDAsIEZMQUdTLCAicmFuZ2UifSwKKyAgICB7TlVMTH0sCit9OworCitz dGF0aWMgY29uc3QgQVZDbGFzcyBjdWRhY29udl9jbGFzcyA9IHsKKyAgICAuY2xhc3NfbmFtZSA9 ICJjb2xvcnJhbmdlX2N1ZGEiLAorICAgIC5pdGVtX25hbWUgPSBhdl9kZWZhdWx0X2l0ZW1fbmFt ZSwKKyAgICAub3B0aW9uID0gb3B0aW9ucywKKyAgICAudmVyc2lvbiA9IExJQkFWVVRJTF9WRVJT SU9OX0lOVCwKK307CisKK3N0YXRpYyBjb25zdCBBVkZpbHRlclBhZCBjdWRhY29udl9pbnB1dHNb XSA9IHsKKyAgICB7CisgICAgICAgIC5uYW1lID0gImRlZmF1bHQiLAorICAgICAgICAudHlwZSA9 IEFWTUVESUFfVFlQRV9WSURFTywKKyAgICAgICAgLmZpbHRlcl9mcmFtZSA9IGN1ZGFjb252X2Zp bHRlcl9mcmFtZSwKKyAgICAgICAgLmdldF9idWZmZXIudmlkZW8gPSBmZl9kZWZhdWx0X2dldF92 aWRlb19idWZmZXIsCisgICAgfSwKK307CisKK3N0YXRpYyBjb25zdCBBVkZpbHRlclBhZCBjdWRh Y29udl9vdXRwdXRzW10gPSB7CisgICAgeworICAgICAgICAubmFtZSA9ICJkZWZhdWx0IiwKKyAg ICAgICAgLnR5cGUgPSBBVk1FRElBX1RZUEVfVklERU8sCisgICAgICAgIC5jb25maWdfcHJvcHMg PSBjdWRhY29udl9jb25maWdfcHJvcHMsCisgICAgfSwKK307CisKK2NvbnN0IEFWRmlsdGVyIGZm X3ZmX2NvbG9ycmFuZ2VfY3VkYSA9IHsKKyAgICAubmFtZSA9ICJjb2xvcnJhbmdlX2N1ZGEiLAor ICAgIC5kZXNjcmlwdGlvbiA9CisgICAgICAgIE5VTExfSUZfQ09ORklHX1NNQUxMKCJDVURBIGFj Y2VsZXJhdGVkIHZpZGVvIGNvbG9yIHJhbmdlIGNvbnZlcnRlciIpLAorCisgICAgLmluaXQgPSBj dWRhY29udl9pbml0LAorICAgIC51bmluaXQgPSBjdWRhY29udl91bmluaXQsCisKKyAgICAucHJp dl9zaXplID0gc2l6ZW9mKENVREFDb252Q29udGV4dCksCisgICAgLnByaXZfY2xhc3MgPSAmY3Vk YWNvbnZfY2xhc3MsCisKKyAgICBGSUxURVJfSU5QVVRTKGN1ZGFjb252X2lucHV0cyksCisgICAg RklMVEVSX09VVFBVVFMoY3VkYWNvbnZfb3V0cHV0cyksCisKKyAgICBGSUxURVJfU0lOR0xFX1BJ WEZNVChBVl9QSVhfRk1UX0NVREEpLAorCisgICAgLmZsYWdzX2ludGVybmFsID0gRkZfRklMVEVS X0ZMQUdfSFdGUkFNRV9BV0FSRSwKK307CmRpZmYgLS1naXQgYS9saWJhdmZpbHRlci92Zl9jb2xv cnJhbmdlX2N1ZGEuY3UgYi9saWJhdmZpbHRlci92Zl9jb2xvcnJhbmdlX2N1ZGEuY3UKbmV3IGZp bGUgbW9kZSAxMDA2NDQKaW5kZXggMDAwMDAwMDAwMC4uNmY2MTc0OTNmOAotLS0gL2Rldi9udWxs CisrKyBiL2xpYmF2ZmlsdGVyL3ZmX2NvbG9ycmFuZ2VfY3VkYS5jdQpAQCAtMCwwICsxLDkzIEBA CisvKgorICogQ29weXJpZ2h0IChjKSAyMDIyLCBOVklESUEgQ09SUE9SQVRJT04uIEFsbCByaWdo dHMgcmVzZXJ2ZWQuCisgKgorICogUGVybWlzc2lvbiBpcyBoZXJlYnkgZ3JhbnRlZCwgZnJlZSBv ZiBjaGFyZ2UsIHRvIGFueSBwZXJzb24gb2J0YWluaW5nIGEKKyAqIGNvcHkgb2YgdGhpcyBzb2Z0 d2FyZSBhbmQgYXNzb2NpYXRlZCBkb2N1bWVudGF0aW9uIGZpbGVzICh0aGUgIlNvZnR3YXJlIiks CisgKiB0byBkZWFsIGluIHRoZSBTb2Z0d2FyZSB3aXRob3V0IHJlc3RyaWN0aW9uLCBpbmNsdWRp bmcgd2l0aG91dCBsaW1pdGF0aW9uCisgKiB0aGUgcmlnaHRzIHRvIHVzZSwgY29weSwgbW9kaWZ5 LCBtZXJnZSwgcHVibGlzaCwgZGlzdHJpYnV0ZSwgc3VibGljZW5zZSwKKyAqIGFuZC9vciBzZWxs IGNvcGllcyBvZiB0aGUgU29mdHdhcmUsIGFuZCB0byBwZXJtaXQgcGVyc29ucyB0byB3aG9tIHRo ZQorICogU29mdHdhcmUgaXMgZnVybmlzaGVkIHRvIGRvIHNvLCBzdWJqZWN0IHRvIHRoZSBmb2xs b3dpbmcgY29uZGl0aW9uczoKKyAqCisgKiBUaGUgYWJvdmUgY29weXJpZ2h0IG5vdGljZSBhbmQg dGhpcyBwZXJtaXNzaW9uIG5vdGljZSBzaGFsbCBiZSBpbmNsdWRlZCBpbgorICogYWxsIGNvcGll cyBvciBzdWJzdGFudGlhbCBwb3J0aW9ucyBvZiB0aGUgU29mdHdhcmUuCisgKgorICogVEhFIFNP RlRXQVJFIElTIFBST1ZJREVEICJBUyBJUyIsIFdJVEhPVVQgV0FSUkFOVFkgT0YgQU5ZIEtJTkQs IEVYUFJFU1MgT1IKKyAqIElNUExJRUQsIElOQ0xVRElORyBCVVQgTk9UIExJTUlURUQgVE8gVEhF IFdBUlJBTlRJRVMgT0YgTUVSQ0hBTlRBQklMSVRZLAorICogRklUTkVTUyBGT1IgQSBQQVJUSUNV TEFSIFBVUlBPU0UgQU5EIE5PTklORlJJTkdFTUVOVC4gIElOIE5PIEVWRU5UIFNIQUxMCisgKiBU SEUgQVVUSE9SUyBPUiBDT1BZUklHSFQgSE9MREVSUyBCRSBMSUFCTEUgRk9SIEFOWSBDTEFJTSwg REFNQUdFUyBPUiBPVEhFUgorICogTElBQklMSVRZLCBXSEVUSEVSIElOIEFOIEFDVElPTiBPRiBD T05UUkFDVCwgVE9SVCBPUiBPVEhFUldJU0UsIEFSSVNJTkcKKyAqIEZST00sIE9VVCBPRiBPUiBJ TiBDT05ORUNUSU9OIFdJVEggVEhFIFNPRlRXQVJFIE9SIFRIRSBVU0UgT1IgT1RIRVIKKyAqIERF QUxJTkdTIElOIFRIRSBTT0ZUV0FSRS4KKyAqLworCitleHRlcm4gIkMiIHsKKyNkZWZpbmUgTVBF R19MVU1BX01JTiAgICgxNikKKyNkZWZpbmUgTVBFR19DSFJPTUFfTUlOICgxNikKKyNkZWZpbmUg TVBFR19MVU1BX01BWCAgICgyMzUpCisjZGVmaW5lIE1QRUdfQ0hST01BX01BWCAoMjQwKQorCisj ZGVmaW5lIEpQRUdfTFVNQV9NSU4gICAoMCkKKyNkZWZpbmUgSlBFR19DSFJPTUFfTUlOICgxKQor I2RlZmluZSBKUEVHX0xVTUFfTUFYICAgKDI1NSkKKyNkZWZpbmUgSlBFR19DSFJPTUFfTUFYICgy NTUpCisKK19fZGV2aWNlX18gaW50IG1wZWdfbWluW10gPSB7TVBFR19MVU1BX01JTiwgTVBFR19D SFJPTUFfTUlOfTsKK19fZGV2aWNlX18gaW50IG1wZWdfbWF4W10gPSB7TVBFR19MVU1BX01BWCwg TVBFR19DSFJPTUFfTUFYfTsKKworX19kZXZpY2VfXyBpbnQganBlZ19taW5bXSA9IHtKUEVHX0xV TUFfTUlOLCBKUEVHX0NIUk9NQV9NSU59OworX19kZXZpY2VfXyBpbnQganBlZ19tYXhbXSA9IHtK UEVHX0xVTUFfTUFYLCBKUEVHX0NIUk9NQV9NQVh9OworCitfX2RldmljZV9fIGludCBjbGFtcChp bnQgdmFsLCBpbnQgbWluLCBpbnQgbWF4KQoreworICAgIGlmICh2YWwgPCBtaW4pCisgICAgICAg IHJldHVybiBtaW47CisgICAgZWxzZSBpZiAodmFsID4gbWF4KQorICAgICAgICByZXR1cm4gbWF4 OworICAgIGVsc2UKKyAgICAgICAgcmV0dXJuIHZhbDsKK30KKworX19nbG9iYWxfXyB2b2lkIHRv X2pwZWdfY3VkYShjb25zdCB1bnNpZ25lZCBjaGFyKiBzcmMsIHVuc2lnbmVkIGNoYXIqIGRzdCwK KyAgICAgICAgICAgICAgICAgICAgICAgICAgICAgaW50IHBpdGNoLCBpbnQgY29tcF9pZCkKK3sK KyAgICBpbnQgeCA9IGJsb2NrSWR4LnggKiBibG9ja0RpbS54ICsgdGhyZWFkSWR4Lng7CisgICAg aW50IHkgPSBibG9ja0lkeC55ICogYmxvY2tEaW0ueSArIHRocmVhZElkeC55OworICAgIGludCBz cmNfLCBkc3RfOworCisgICAgLy8gOCBiaXQgLT4gMTUgYml0IGZvciBiZXR0ZXIgcHJlY2lzaW9u OworICAgIHNyY18gPSBzdGF0aWNfY2FzdDxpbnQ+KHNyY1t4ICsgeSAqIHBpdGNoXSkgPDwgNzsK KworICAgIC8vIENvbnZlcnNpb247CisgICAgZHN0XyA9IGNvbXBfaWQgPyAobWluKHNyY18sIDMw Nzc1KSAqIDQ2NjMgLSA5Mjg5OTkyKSA+PiAxMiAgICAvLyBjaHJvbWEKKyAgICAgICAgICAgICAg ICAgICA6IChtaW4oc3JjXywgMzAxODkpICogMTkwNzcgLSAzOTA1NzM2MSkgPj4gMTQ7IC8vIGx1 bWEKKworICAgIC8vIERpdGhlciByZXBsYWNlbWVudDsKKyAgICBkc3RfID0gZHN0XyArIDY0Owor CisgICAgLy8gQmFjayB0byA4IGJpdDsKKyAgICBkc3RfID0gY2xhbXAoZHN0XyA+PiA3LCBqcGVn X21pbltjb21wX2lkXSwganBlZ19tYXhbY29tcF9pZF0pOworICAgIGRzdFt4ICsgeSAqIHBpdGNo XSA9IHN0YXRpY19jYXN0PHVuc2lnbmVkIGNoYXI+KGRzdF8pOworfQorCitfX2dsb2JhbF9fIHZv aWQgdG9fbXBlZ19jdWRhKGNvbnN0IHVuc2lnbmVkIGNoYXIqIHNyYywgdW5zaWduZWQgY2hhciog ZHN0LAorICAgICAgICAgICAgICAgICAgICAgICAgICAgICBpbnQgcGl0Y2gsIGludCBjb21wX2lk KQoreworICAgIGludCB4ID0gYmxvY2tJZHgueCAqIGJsb2NrRGltLnggKyB0aHJlYWRJZHgueDsK KyAgICBpbnQgeSA9IGJsb2NrSWR4LnkgKiBibG9ja0RpbS55ICsgdGhyZWFkSWR4Lnk7CisgICAg aW50IHNyY18sIGRzdF87CisKKyAgICAvLyA4IGJpdCAtPiAxNSBiaXQgZm9yIGJldHRlciBwcmVj aXNpb247CisgICAgc3JjXyA9IHN0YXRpY19jYXN0PGludD4oc3JjW3ggKyB5ICogcGl0Y2hdKSA8 PCA3OworCisgICAgLy8gQ29udmVyc2lvbjsKKyAgICBkc3RfID0gY29tcF9pZCA/IChzcmNfICog MTc5OSArIDQwODEwODUpID4+IDExICAgIC8vIGNocm9tYQorICAgICAgICAgICAgICAgICAgIDog KHNyY18gKiAxNDA3MSArIDMzNTYxOTQ3KSA+PiAxNDsgLy8gbHVtYQorCisgICAgLy8gRGl0aGVy IHJlcGxhY2VtZW50OworICAgIGRzdF8gPSBkc3RfICsgNjQ7CisKKyAgICAvLyBCYWNrIHRvIDgg Yml0OworICAgIGRzdF8gPSBjbGFtcChkc3RfID4+IDcsIG1wZWdfbWluW2NvbXBfaWRdLCBtcGVn X21heFtjb21wX2lkXSk7CisgICAgZHN0W3ggKyB5ICogcGl0Y2hdID0gc3RhdGljX2Nhc3Q8dW5z aWduZWQgY2hhcj4oZHN0Xyk7Cit9Cit9ClwgTm8gbmV3bGluZSBhdCBlbmQgb2YgZmlsZQotLSAK Mi4yNS4xCgo= --_004_PH7PR12MB58312279BA34EDEDC705A0CFD2459PH7PR12MB5831namp_ 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". --_004_PH7PR12MB58312279BA34EDEDC705A0CFD2459PH7PR12MB5831namp_--