Coverage for cuda/core/_utils/cuda_utils.pyx: 95.20%
229 statements
« prev ^ index » next coverage.py v7.14.1, created at 2026-06-13 01:38 +0000
« prev ^ index » next coverage.py v7.14.1, created at 2026-06-13 01:38 +0000
1# SPDX-FileCopyrightText: Copyright (c) 2024-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
2#
3# SPDX-License-Identifier: Apache-2.0
5import functools
6from functools import partial
7import multiprocessing
8import platform
9import warnings
10from collections import namedtuple
11from collections.abc import Sequence
12from contextlib import ExitStack
13from typing import Any, Callable
15from cuda.bindings import driver as driver, nvrtc as nvrtc, runtime as runtime
17# Module-level annotations that reference `driver`, `nvrtc`, and `runtime` so
18# that stubgen-pyx keeps these imports in the generated `.pyi` (it would
19# otherwise trim them as unused). These names are not assigned, so they only
20# affect __annotations__ and have no runtime cost.
21_keep_driver_in_stub: 'driver.CUresult'
22_keep_nvrtc_in_stub: 'nvrtc.nvrtcResult'
23_keep_runtime_in_stub: 'runtime.cudaError_t'
25from cuda.bindings.nvvm import nvvmError
26from cuda.bindings.nvjitlink import nvJitLinkError
28from cpython.buffer cimport PyObject_GetBuffer, PyBuffer_Release, Py_buffer, PyBUF_SIMPLE
30from cuda.bindings cimport cynvrtc, cynvvm, cynvjitlink
32from cuda.core._utils.driver_cu_result_explanations import DRIVER_CU_RESULT_EXPLANATIONS
33from cuda.core._utils.runtime_cuda_error_explanations import RUNTIME_CUDA_ERROR_EXPLANATIONS
36class CUDAError(Exception):
37 pass
40class NVRTCError(CUDAError):
41 pass
45ComputeCapability = namedtuple("ComputeCapability", ("major", "minor"))
48def cast_to_3_tuple(label: str, cfg: int | tuple[int, ...]) -> tuple[int, int, int]:
49 cfg_orig = cfg 2# $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdbAbBbCbDbEbFbmfnfof`hhdpf0gidqfjdrfkdsfldtfmdXbufvfwfxfyfh i zfGbHbIbAfJbKbLbMbNbd e j f g 6 eb7 8 fbgbhbibjbObkblbmbnbobpb9 ! zcAcnend#c$cqbod#h^E_E`E{hxg=B|hBcpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdoepe}h{E~hCcBfCfaiDfEfbiFfGfqere=eseteHfIfJfKfLfMfNfOfPf1gQfRf2gSfTfueve?eUfVfWfwexeXfyezeYfZf0f1f2fRd3f4f3g%c4g5g@e[e6g7g8g9g!gSd'c#g]eTd$g%g'g5fUdAeBe6f7fCeDe^e8f9fci!f#fEeFe_eGeHe$f%f'f(f)f*f+f,f-f(g.f/f)g:f;fIeJe`e=f?f@fKeLe[fMeNe]f^f_f`f{fVd|f}f*g(c+g,g{e|e-g.g/g:g;gWd)c=g}eXd?g@g[g~fYdOePeagbgQeRe~ecgdgdiegfgeigghgfiigjggikglghimgngiiogpgjiqgrgki$hVEWE?B|E%h
50 if isinstance(cfg, int): 2# $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdbAbBbCbDbEbFbmfnfof`hhdpf0gidqfjdrfkdsfldtfmdXbufvfwfxfyfh i zfGbHbIbAfJbKbLbMbNbd e j f g 6 eb7 8 fbgbhbibjbObkblbmbnbobpb9 ! zcAcnend#c$cqbod#h^E_E`E{hxg=B|hBcpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdoepe}h{E~hCcBfCfaiDfEfbiFfGfqere=eseteHfIfJfKfLfMfNfOfPf1gQfRf2gSfTfueve?eUfVfWfwexeXfyezeYfZf0f1f2fRd3f4f3g%c4g5g@e[e6g7g8g9g!gSd'c#g]eTd$g%g'g5fUdAeBe6f7fCeDe^e8f9fci!f#fEeFe_eGeHe$f%f'f(f)f*f+f,f-f(g.f/f)g:f;fIeJe`e=f?f@fKeLe[fMeNe]f^f_f`f{fVd|f}f*g(c+g,g{e|e-g.g/g:g;gWd)c=g}eXd?g@g[g~fYdOePeagbgQeRe~ecgdgdiegfgeigghgfiigjggikglghimgngiiogpgjiqgrgki$hVEWE?B|E%h
51 cfg = (cfg,) 2# $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdbAbBbCbDbEbFbmfnfofpfqfrfsftfXbufvfwfxfyfh i zfGbHbIbAfJbKbLbMbNbd e j f g 6 eb7 8 fbgbhbibjbObkblbmbnbobpb9 ! zcAcnend#c$cqbod#h^E_E`Exg=BpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdoepe{ECcBfCfDfEfFfGfqereseteHfIfJfKfLfMfNfOfPfQfRfSfTfueveUfVfWfwexeXfyezeYfZf0f1f2fRd3f4f%c'c5fUdAeBe6f7fCeDe8f9f!f#fEeFeGeHe$f%f'f(f)f*f+f,f-f.f/f:f;fIeJe=f?f@fKeLe[fMeNe]f^f_f`f{fVd|f}f(c)c~fYdOePeagbgQeRecgdgegfggghgigjgkglgmgngogpgqgrg$hVE
52 else:
53 common = "must be an int, or a tuple with up to 3 ints" 2`hhd0gidjdkdldmd#h{hxg|hBc}h~haibi=e1g2g?e3g%c4g5g@e[e6g7g8g9g!gSd'c#g]eTd$g%g'g^eci_e(g)g`e*g(c+g,g{e|e-g.g/g:g;gWd)c=g}eXd?g@g[g~edieifigihiiijiki$hWE?B|E%h
54 if not isinstance(cfg, tuple): 2`hhd0gidjdkdldmd#h{hxg|hBc}h~haibi=e1g2g?e3g%c4g5g@e[e6g7g8g9g!gSd'c#g]eTd$g%g'g^eci_e(g)g`e*g(c+g,g{e|e-g.g/g:g;gWd)c=g}eXd?g@g[g~edieifigihiiijiki$hWE?B|E%h
55 raise ValueError(f"{label} {common} (got {type(cfg)})") 2|E
56 if len(cfg) > 3: 2`hhd0gidjdkdldmd#h{hxg|hBc}h~haibi=e1g2g?e3g%c4g5g@e[e6g7g8g9g!gSd'c#g]eTd$g%g'g^eci_e(g)g`e*g(c+g,g{e|e-g.g/g:g;gWd)c=g}eXd?g@g[g~edieifigihiiijiki$hWE?B%h
57 raise ValueError(f"{label} {common} (got tuple with length {len(cfg)})") 2WE
58 if any(not isinstance(val, int) for val in cfg): 2`hhd0gidjdkdldmd#h{hxg|hBc}h~haibi=e1g2g?e3g%c4g5g@e[e6g7g8g9g!gSd'c#g]eTd$g%g'g^eci_e(g)g`e*g(c+g,g{e|e-g.g/g:g;gWd)c=g}eXd?g@g[g~edieifigihiiijiki$h?B%h
59 raise ValueError(f"{label} {common} (got {cfg})") 2?B
60 if any(val < 1 for val in cfg): 2# $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdbAbBbCbDbEbFbmfnfof`hhdpf0gidqfjdrfkdsfldtfmdXbufvfwfxfyfh i zfGbHbIbAfJbKbLbMbNbd e j f g 6 eb7 8 fbgbhbibjbObkblbmbnbobpb9 ! zcAcnend#c$cqbod#h^E_E`E{hxg=B|hBcpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdoepe}h{E~hCcBfCfaiDfEfbiFfGfqere=eseteHfIfJfKfLfMfNfOfPf1gQfRf2gSfTfueve?eUfVfWfwexeXfyezeYfZf0f1f2fRd3f4f3g%c4g5g@e[e6g7g8g9g!gSd'c#g]eTd$g%g'g5fUdAeBe6f7fCeDe^e8f9fci!f#fEeFe_eGeHe$f%f'f(f)f*f+f,f-f(g.f/f)g:f;fIeJe`e=f?f@fKeLe[fMeNe]f^f_f`f{fVd|f}f*g(c+g,g{e|e-g.g/g:g;gWd)c=g}eXd?g@g[g~fYdOePeagbgQeRe~ecgdgdiegfgeigghgfiigjggikglghimgngiiogpgjiqgrgki$hVE%h
61 plural_s = "" if len(cfg) == 1 else "s" 2xgVE%h
62 raise ValueError(f"{label} value{plural_s} must be >= 1 (got {cfg_orig})") 2xgVE%h
63 return cfg + (1,) * (3 - len(cfg)) 2# $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdbAbBbCbDbEbFbmfnfof`hhdpf0gidqfjdrfkdsfldtfmdXbufvfwfxfyfh i zfGbHbIbAfJbKbLbMbNbd e j f g 6 eb7 8 fbgbhbibjbObkblbmbnbobpb9 ! zcAcnend#c$cqbod#h^E_E`E{hxg=B|hBcpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdoepe}h{E~hCcBfCfaiDfEfbiFfGfqere=eseteHfIfJfKfLfMfNfOfPf1gQfRf2gSfTfueve?eUfVfWfwexeXfyezeYfZf0f1f2fRd3f4f3g%c4g5g@e[e6g7g8g9g!gSd'c#g]eTd$g%g'g5fUdAeBe6f7fCeDe^e8f9fci!f#fEeFe_eGeHe$f%f'f(f)f*f+f,f-f(g.f/f)g:f;fIeJe`e=f?f@fKeLe[fMeNe]f^f_f`f{fVd|f}f*g(c+g,g{e|e-g.g/g:g;gWd)c=g}eXd?g@g[g~fYdOePeagbgQeRe~ecgdgdiegfgeigghgfiigjggikglghimgngiiogpgjiqgrgki$h
66cdef int HANDLE_RETURN(cydriver.CUresult err) except?-1 nogil:
67 if err != cydriver.CUresult.CUDA_SUCCESS: 2a 4LZd5L0d6L7L# 8L$ 9L!L% #L' $L( %L) 'L* (L+ )L, *L- +L. ,L/ -L: .L; /L= :L? ;L@ =L[ ?L] @L^ [L_ ]L` ^L{ _L| `L} {L~ |Lab}Lbb~LcbaMdbbMcMdMeMfMgMhMiMjMkM]glMmMnMoMpMqMrMsMtMuMvMwMxMyMzMAMBMCMDMEMFMGMHMIMJMKMLMSeMMNMOMPMTeQMUeRMVeSMWeTMAbUMXeVMYeWMZeXMYMZM0M1MBb2MCb3MDb4MEb5MFb6MDc7MEc8MFc9MGc!MHc#MIc$MJc%MKc'M(M)M*M+M,Myg-Mzg.MAg/M:M;M=M?M@M[Mmf]Mnf^M_M`M{Mof|M}E~E}M~MaNbNcNhddNaFbFlimieNfNgNhNiNjNkNpflNmNcFnNdFoNpNqNrNsNtNuNvNwN0gidxNyNzNeFANfFBNniCNoiDNENFNGNHNINJNKNLNMNNNONPNQNRNSNTNUNVNWNXNYNZN0N1N2Nqf3NgFhF4N5N6N7N8Njd9NiFjFpiqi!N#N$N%N'N(N)Nrf*N+N,NkF-NlF.N/N:N;N=N?N@N[N]N^N_Nkd`N{N|NmF}NnF~NriaOsibOcOdOeOfOgOhOiOjOsfkOlOmOoFnOpFoOpOqOrOsOtOuOvOwOxOyOldzOAOBOqFCOrFDOtiEOuiFOGOHOIOJOKOLOMONOtfOOPOQOsFROtFSOTOUOVOWOXOYOZO0O1O2Omd3O4O5OuF6OvF7Ovi8Owi9O!O#O$O%O'O(O)O*O+O,O-O.O/O:O;O=O?O@O[O]O^O_O`O{O|O}O~OaPbPcPdPePfPgPhPiPjPkPlPmPnP^goPpPqPrPsPtPuPvPwPxPyPzPAPBPCP_gDPEPFPGPHP@BIPXbJPKPLP`gMP0eNPOPPPQPRPSPTPUPVPWPXPYPZP0P1P2P3P4P5P6P7P8P9Puf!Pvf#Pwf$Pxf%P1e'P2e(Pyf)Paf*PwF+PxF,P*c-P+c.PyF/PzF:Ph ;Pi =P[B?Pzf@PGb[PHb]PIb^PAf_P3e`P{g{PJb|PKb}PLb~PMbaQNbbQcQdQeQfQgQhQiQjQd kQ|glQ4emQ}gnQ,coQ-cpQ.cqQ1drQ2dsQ3dtQe uQj vQf wQg xQ6 yQebzQ7 AQ8 BQfbCQgbDQhbEQ4dFQ5dGQ6dHQ7dIQ8dJQ9dKQLQMQNQibOQjbPQObQQkbRQlbSQmbTQnbUQobVQpbWQAFBFbfCFXQ~gahbhchdhehYQfhZQDFBg0QEFcfFFGF5e1QdfefffHFIFgfJFKFCg2Q3QghhhihhfLFifMFNFOFPFQFRFSFTFUFVFWF6e4QaLXFbLYF]BZF^B0FcL1FdL2F_B3F`B4F9 5F! 6Fjh5Qkh6Q'h7Q(h8Qlh9Qmh!Q)h#Q*h$Qxi%Q'Q(QeL7FfL8F9F!F#FgL$F)Q%F'F(FhL)F*Q*F+F,FiL-F+Q.F/F:FjL;F,Q=F?F@FkL[F-Q]F^F_FlL`F.Q{F|FmL}Fsg~FnLaGoLbGpLcGqLdG!deG/cfG/QYbgG:chG+b:QLc;Q;ciG=cjGtgkGrLlGsLmGtLnGQb=QRb?QSb@QTb[QUb]QVb^QDg_Q`Q{Q|Q}Q~QaRbREg7e8ecRdReRfRgRhRiRjRkRlRFgmRnRoRpRqRrRsRtRuRvRwRxRyRzRARBRCRDRERFRGRHRIRJRKRLRMRNRORPRQRRRSRTRURVRWRXRYRZR0R1R2R3R4R5R6R7R8R9R!R#R$R%R'R(R)R*R+R,R-R.R/R:R;R=R?R@R[R]R^R_R`R{R|R}R~RaSbScSdSeSfSgShSiSjSkSlSmSnSoSpSqSrSsStSuSvSwSxSySzSASBSCSDSESFSGSHSISJSKSLSMSNSOSPSQSRSSSTSUSVSWSXSYSZS0S1S2S3S4S5S6S7S8S9S!S#S$S%S'S(S)S*Snh+S,S-SohGgHgph{Bzc|B.SoG/SpG:SqG;SrG=SsG?StG@SuG[SAc]SvG^SwG_SxG`SyG{S}B|S~B}SaC~SneaTzGAGBGCGDGbTcTdTeTfTgTEGFGGGHGhTIGJGiTKGLGjTkTlTmTXEnToTYEqhMGZErhndyish#cbC$ccCpT0Eth1Euh2EvhqT3EwhrT4ExhNG5EyhsTtTuTvTwTxT6EyT7EzT8EAT9EBT!ECT#EDT$EETFTGTHTITJTKTLTMTqb,b-bodNTOT=BPTBcQTpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdoepe.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b~bacbcccdcecfcgchcicjckclcmcncocdCpcRTSTTTIgqcrcJgUTKgLgMgOGPGVTNgOgPgQgzhAhWTXTrbYTZTsb0T1T2Ttb3TZb4T5T6T7Tk 8T9Tl !Tm #Tn $To %Tp 'Tq (Tr )Ts *Tt +T,Tu -Tv .Tw /Tx :T;Ty =Tz ?T@T[T]T^Tub_T`Tvb{T|T}Twb~T0baUbUcUdUA eUfUB gUC hUD iUE jUF kUG lUH mUI nUJ oUpUK qUL rUM sUN tUuUO vUP wUxUyUzUAUxbBUCUybDUEUFUzbGU1bHUIUJUKUQ LUMUR NUS OUT PUU QUV RUW SUX TUY UUZ VUWU0 XU1 YU2 ZU3 0U1U4 2U5 3U4U5URgQGRGSG6U?c7UTG8UUG9U!U#U$U9eeCfC%UVG'UWG(UgC)UhC*UiC+UjC,UMc-U!e.U@c/U#e:U;UXG=UYG?UZG@U0G[U1G]U2G^U3G_U4G`U5G{U6G|U7G}U8G~U9GaV!GbV#GcV$GdVkCeV%GfVlCgV'GhVmCiV(GjVnCkV)GlVoCmV*GnVpCoV+GpVqCqV,GrVrCsV-GtV.GuV/GvV:GwV;GxV=GyV?GzV@GAV[GBV]GCV^GDV_GEV`GFV{GGV|GHV}GIV~GJV%EKVziLV'EMV(ENVAiOV)EPVaHQVRVSgTg$ebHcHdHSV[cTVeHfHgHUVhHVViHWVjHXVsCYVtCZVNcOcPbPc0VQc1V2Vsc3VuC4VvC#d5VwCxC6VyC7VzC8VACCc9VBCWbCC!VDC#VEC$VFC%VGC'VHC(VIC)VJC*VKC+VLC,VMCNCOCPCRcScTcQCRCSCUcVcWc-VTC.VUC/VVC:VWC;VXC=VYC?VZC@V0C[V1C]V2C^V3C_V4C`V5C{V6C|V7C}V8C~V9CaW!CbW#CcW$CdW%CeW'CfW(CgW)C*C+ChW,CiW-CjW.CkW/ClW:CmW;CnW=CoW?CpW@CqW[CrW]CsW^CtW_CuW`CvW{CwW|CxW}CyW~CzWaDAWbDBWcDdDCWeDDWfDEWgDFWhDGWiDHWjDIWkDJWlDKWmDLWnDMWoDNWpDOWqDPWrDQWsDRWtDuDugSW%eTW'eUW]cVW$dWWXWkHBfYWZW0W1WlHmHnH2W3W4W5W6WvDCf7WoHwD8WpH9W!WqHxDrH#WsHDf$W%W'W(WtHuHvH)W*W+W,W-WyDEf.WwHzD/WxH:W;WFfyHzH=W?W@W[WGf]WAH^W_WBHqeCHDH*E+EEHFHGHHHIHBire=eJHCiDiKHLHMHseteHfIfJfKfLf`WMf{WNHOH|W}W~WaXNfbXPHcXdXeXfXQHOfgXhXRHSHiXjXkXlXmXADPf1gnXTHBDoXpXqXrXUHQfsXtXuXVHWHXHvXwXxXyXzXCDRf2gAXYHDDBXZHCXDXEXSf0H1HFXGXHXIXTfJX2HKXLX3H,Eue4H5H6H7H8H9H!H#H$H%HEDve?e'HEiFD(H)H*H+H,H-H.H/H:H;H=H?H@H[HUf]H^HMXNXOXPX_HQXRXSXVfTXUX`H{HVXWXXXYXWfZX|H0X1X2X3X4X5X6X7X8X9X!X#XGDweHDIDJDKDLDMDNDODPDxeQDBhRDSDTDXfyezeYfZf0f1f2fRd3f4f3g%c4g5g@e[e6g7g8g9g!gSd'c#g]eTd$g%g'g5f}H~H$X%X'X(X)XUdFiGiaIbIcIdIeIfIUDAeVDWDXDYDZD0D1D2D3D4DBe5DCh6D7D8D*X6f+XgIhI,X-X.X/X7f:XiI;X=XjI-ECekIlImInIoIpIqIrIsItIuI9DDe^evIHi!DwIxIyI?X@XzI8f[X]X^X_XAIBICI`X{X|X}X~X#D9faYDI$DbYEIcYdY!fFIGIeYfYgYhY#fiYHIjYkYIIEeJIKI.E/ELIMINIOIPIIiFe_eQIJiKiRISITIGeHe$f%f'f(f)flY*fmYUIVInYoYpYqY+frYWIsYtYuYvYXI,fwYxYYIZIyYzYAYBYCY%D-f(gDY0I'DEYFYGYHY1I.fIYJYKY2I3I4ILYMYNYOYPY(D/f)gQY5I)DRY6ISYTYUY:f7I8IVYWYXYYY;fZY9I0Y1Y!I:EIe#I$I%I'I(I)I*I+I,I-I*DJe`e.ILi+D/I:I;I=I?I@I[I]I^I_I`I{I|I}I=f~IaJ2Y3Y4Y5YbJ6Y7Y8Y?f9Y!YcJdJ#Y$Y%Y'Y@f(YeJ)Y*Y+Y,Y-Y.Y/Y:Y;Y=Y?Y@Y,DKe-D.D/D:D;D=D?D@D[DLe]DDh^D_D`D[fMeNe]f^f_f`f{fVd|f}f*g(c+g,g{e|e-g.g/g:g;gWd)c=g}eXd?g@g[g~ffJgJ[Y]Y^Y_Y`YYdMiNihJiJjJkJlJmJ{DOe|D}D~DaEbEcEdEeEfEgEPehEEhiEjEkE{Yag|YnJoJ}Y~YaZbZbgcZpJdZeZqJ;EQerJsJtJuJvJwJxJyJzJAJBJlERe~eCJOimEDJEJFJfZgZhZGJ(ePiQinEoEiZHJcgjZkZlZmZIJJJKJnZoZpZqZrZpEdgsZLJqErEvgsEtEwgtZMJuZvZwZNJegxZyZzZAZOJPJQJBZCZDZEZFZuEfgGZRJvEHZSJIZJZKZTJggLZMZNZOZUJVJWJPZQZRZSZTZwEhgUZXJxEVZYJWZXZYZZJigZZ0Z1Z2Z0J1J2J3Z4Z5Z6Z7ZyEjg8Z3JzE9Z4J!Z#Z$Z5Jkg%Z'Z(Z)Z6J7J8J*Z+Z,Z-Z.ZAElg/Z9JBE:Z!J;Z=Z?Z#Jmg@Z[Z]Z^Z$J%J'J_Z`Z{Z|Z}ZCEng~Z(JDEa0)Jb0c0d0*Joge0f0g0h0+J,J-Ji0j0k0l0m0EEpgn0.JFEGEo0/Jp0q0r0:Jqgs0t0u0v0;J=J?Jw0x0y0z0A0HErgB0@JIEJEC0[JD0E0XcF0YcG0=EH0ZcI0J0K0L0M0N0O0P0Q0R0S0T0U0V0W0X0Y0Z000102030405060708090!0#0$0%0'0(0)0*0+0,0c -00c.0KE2bLE%d/0:0;0=0?0@0[0]0^0_0`0{0|0}0RiSi]JMENEOE~0PE^J_JQE?E1c+ha1b1c1d1e1f1g1h1i1j1k12cl1m1n1o1p1q1r1s1t1u1v1w1x1`Jy1z1A1B1C1D1E1F1G1H1REI1J1{JK1|JL1M1}JN1O1P1Q1R1S1T1U1V1W1X1Y1Z101112131415161718191!1#1$1%1'1(1)1*1+1,1-1.1/1:1;1=1?1@1[1]1^1_1`1{1|1}1~1a2b2c2d2e2f2g2h2i2j2k2l2m2n2o2p2q2r2s2t2u2v2w2x2~JaKy2z2bKcKdKeKfKA2gKhKiKB2jKC2kKD2lKmKE2nKF2oKG2pKH2qKI2J2rKsKtKuKvKK2wKL2xKM2N2O2yKP2Q2R2S2T2U2V2W2X2Y2Z202122232425262728292!2#2$2%2'2(2)2*2+2,2-2.2/2:2;2=2?2@2
68 return _check_driver_error(err) 2]g^g_g`g|g4e}gWb%d
69 return 0 2a 4LZd5L0d6L7L# 8L$ 9L!L% #L' $L( %L) 'L* (L+ )L, *L- +L. ,L/ -L: .L; /L= :L? ;L@ =L[ ?L] @L^ [L_ ]L` ^L{ _L| `L} {L~ |Lab}Lbb~LcbaMdbbMcMdMeMfMgMhMiMjMkM]glMmMnMoMpMqMrMsMtMuMvMwMxMyMzMAMBMCMDMEMFMGMHMIMJMKMLMSeMMNMOMPMTeQMUeRMVeSMWeTMAbUMXeVMYeWMZeXMYMZM0M1MBb2MCb3MDb4MEb5MFb6MDc7MEc8MFc9MGc!MHc#MIc$MJc%MKc'M(M)M*M+M,Myg-Mzg.MAg/M:M;M=M?M@M[Mmf]Mnf^M_M`M{Mof|M}E~E}M~MaNbNcNhddNaFbFlimieNfNgNhNiNjNkNpflNmNcFnNdFoNpNqNrNsNtNuNvNwN0gidxNyNzNeFANfFBNniCNoiDNENFNGNHNINJNKNLNMNNNONPNQNRNSNTNUNVNWNXNYNZN0N1N2Nqf3NgFhF4N5N6N7N8Njd9NiFjFpiqi!N#N$N%N'N(N)Nrf*N+N,NkF-NlF.N/N:N;N=N?N@N[N]N^N_Nkd`N{N|NmF}NnF~NriaOsibOcOdOeOfOgOhOiOjOsfkOlOmOoFnOpFoOpOqOrOsOtOuOvOwOxOyOldzOAOBOqFCOrFDOtiEOuiFOGOHOIOJOKOLOMONOtfOOPOQOsFROtFSOTOUOVOWOXOYOZO0O1O2Omd3O4O5OuF6OvF7Ovi8Owi9O!O#O$O%O'O(O)O*O+O,O-O.O/O:O;O=O?O@O[O]O^O_O`O{O|O}O~OaPbPcPdPePfPgPhPiPjPkPlPmPnP^goPpPqPrPsPtPuPvPwPxPyPzPAPBPCP_gDPEPFPGPHP@BIPXbJPKPLP`gMP0eNPOPPPQPRPSPTPUPVPWPXPYPZP0P1P2P3P4P5P6P7P8P9Puf!Pvf#Pwf$Pxf%P1e'P2e(Pyf)Paf*PwF+PxF,P*c-P+c.PyF/PzF:Ph ;Pi =P[B?Pzf@PGb[PHb]PIb^PAf_P3e`P{g{PJb|PKb}PLb~PMbaQNbbQcQdQeQfQgQhQiQjQd kQ|glQ4emQ}gnQ,coQ-cpQ.cqQ1drQ2dsQ3dtQe uQj vQf wQg xQ6 yQebzQ7 AQ8 BQfbCQgbDQhbEQ4dFQ5dGQ6dHQ7dIQ8dJQ9dKQLQMQNQibOQjbPQObQQkbRQlbSQmbTQnbUQobVQpbWQAFBFbfCFXQ~gahbhchdhehYQfhZQDFBg0QEFcfFFGF5e1QdfefffHFIFgfJFKFCg2Q3QghhhihhfLFifMFNFOFPFQFRFSFTFUFVFWF6e4QaLXFbLYF]BZF^B0FcL1FdL2F_B3F`B4F9 5F! 6Fjh5Qkh6Q'h7Q(h8Qlh9Qmh!Q)h#Q*h$Qxi%Q'Q(QeL7FfL8F9F!F#FgL$F)Q%F'F(FhL)F*Q*F+F,FiL-F+Q.F/F:FjL;F,Q=F?F@FkL[F-Q]F^F_FlL`F.Q{F|FmL}Fsg~FnLaGoLbGpLcGqLdG!deG/cfG/QYbgG:chG+b:QLc;Q;ciG=cjGtgkGrLlGsLmGtLnGQb=QRb?QSb@QTb[QUb]QVb^QDg_Q`Q{Q|Q}Q~QaRbREg7e8ecRdReRfRgRhRiRjRkRlRFgmRnRoRpRqRrRsRtRuRvRwRxRyRzRARBRCRDRERFRGRHRIRJRKRLRMRNRORPRQRRRSRTRURVRWRXRYRZR0R1R2R3R4R5R6R7R8R9R!R#R$R%R'R(R)R*R+R,R-R.R/R:R;R=R?R@R[R]R^R_R`R{R|R}R~RaSbScSdSeSfSgShSiSjSkSlSmSnSoSpSqSrSsStSuSvSwSxSySzSASBSCSDSESFSGSHSISJSKSLSMSNSOSPSQSRSSSTSUSVSWSXSYSZS0S1S2S3S4S5S6S7S8S9S!S#S$S%S'S(S)S*Snh+S,S-SohGgHgph{Bzc|B.SoG/SpG:SqG;SrG=SsG?StG@SuG[SAc]SvG^SwG_SxG`SyG{S}B|S~B}SaC~SneaTzGAGBGCGDGbTcTdTeTfTgTEGFGGGHGhTIGJGiTKGLGjTkTlTmTXEnToTYEqhMGZErhndyish#cbC$ccCpT0Eth1Euh2EvhqT3EwhrT4ExhNG5EyhsTtTuTvTwTxT6EyT7EzT8EAT9EBT!ECT#EDT$EETFTGTHTITJTKTLTMTqb,b-bodNTOT=BPTBcQTpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdoepe.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b~bacbcccdcecfcgchcicjckclcmcncocdCpcRTSTTTIgqcrcJgUTKgLgMgOGPGVTNgOgPgQgzhAhWTXTrbYTZTsb0T1T2Ttb3TZb4T5T6T7Tk 8T9Tl !Tm #Tn $To %Tp 'Tq (Tr )Ts *Tt +T,Tu -Tv .Tw /Tx :T;Ty =Tz ?T@T[T]T^Tub_T`Tvb{T|T}Twb~T0baUbUcUdUA eUfUB gUC hUD iUE jUF kUG lUH mUI nUJ oUpUK qUL rUM sUN tUuUO vUP wUxUyUzUAUxbBUCUybDUEUFUzbGU1bHUIUJUKUQ LUMUR NUS OUT PUU QUV RUW SUX TUY UUZ VUWU0 XU1 YU2 ZU3 0U1U4 2U5 3U4U5URgQGRGSG6U?c7UTG8UUG9U!U#U$U9eeCfC%UVG'UWG(UgC)UhC*UiC+UjC,UMc-U!e.U@c/U#e:U;UXG=UYG?UZG@U0G[U1G]U2G^U3G_U4G`U5G{U6G|U7G}U8G~U9GaV!GbV#GcV$GdVkCeV%GfVlCgV'GhVmCiV(GjVnCkV)GlVoCmV*GnVpCoV+GpVqCqV,GrVrCsV-GtV.GuV/GvV:GwV;GxV=GyV?GzV@GAV[GBV]GCV^GDV_GEV`GFV{GGV|GHV}GIV~GJV%EKVziLV'EMV(ENVAiOV)EPVaHQVRVSgTg$ebHcHdHSV[cTVeHfHgHUVhHVViHWVjHXVsCYVtCZVNcOcPbPc0VQc1V2Vsc3VuC4VvC#d5VwCxC6VyC7VzC8VACCc9VBCWbCC!VDC#VEC$VFC%VGC'VHC(VIC)VJC*VKC+VLC,VMCNCOCPCRcScTcQCRCSCUcVcWc-VTC.VUC/VVC:VWC;VXC=VYC?VZC@V0C[V1C]V2C^V3C_V4C`V5C{V6C|V7C}V8C~V9CaW!CbW#CcW$CdW%CeW'CfW(CgW)C*C+ChW,CiW-CjW.CkW/ClW:CmW;CnW=CoW?CpW@CqW[CrW]CsW^CtW_CuW`CvW{CwW|CxW}CyW~CzWaDAWbDBWcDdDCWeDDWfDEWgDFWhDGWiDHWjDIWkDJWlDKWmDLWnDMWoDNWpDOWqDPWrDQWsDRWtDuDugSW%eTW'eUW]cVW$dWWXWkHBfYWZW0W1WlHmHnH2W3W4W5W6WvDCf7WoHwD8WpH9W!WqHxDrH#WsHDf$W%W'W(WtHuHvH)W*W+W,W-WyDEf.WwHzD/WxH:W;WFfyHzH=W?W@W[WGf]WAH^W_WBHqeCHDH*E+EEHFHGHHHIHBire=eJHCiDiKHLHMHseteHfIfJfKfLf`WMf{WNHOH|W}W~WaXNfbXPHcXdXeXfXQHOfgXhXRHSHiXjXkXlXmXADPf1gnXTHBDoXpXqXrXUHQfsXtXuXVHWHXHvXwXxXyXzXCDRf2gAXYHDDBXZHCXDXEXSf0H1HFXGXHXIXTfJX2HKXLX3H,Eue4H5H6H7H8H9H!H#H$H%HEDve?e'HEiFD(H)H*H+H,H-H.H/H:H;H=H?H@H[HUf]H^HMXNXOXPX_HQXRXSXVfTXUX`H{HVXWXXXYXWfZX|H0X1X2X3X4X5X6X7X8X9X!X#XGDweHDIDJDKDLDMDNDODPDxeQDBhRDSDTDXfyezeYfZf0f1f2fRd3f4f3g%c4g5g@e[e6g7g8g9g!gSd'c#g]eTd$g%g'g5f}H~H$X%X'X(X)XUdFiGiaIbIcIdIeIfIUDAeVDWDXDYDZD0D1D2D3D4DBe5DCh6D7D8D*X6f+XgIhI,X-X.X/X7f:XiI;X=XjI-ECekIlImInIoIpIqIrIsItIuI9DDe^evIHi!DwIxIyI?X@XzI8f[X]X^X_XAIBICI`X{X|X}X~X#D9faYDI$DbYEIcYdY!fFIGIeYfYgYhY#fiYHIjYkYIIEeJIKI.E/ELIMINIOIPIIiFe_eQIJiKiRISITIGeHe$f%f'f(f)flY*fmYUIVInYoYpYqY+frYWIsYtYuYvYXI,fwYxYYIZIyYzYAYBYCY%D-f(gDY0I'DEYFYGYHY1I.fIYJYKY2I3I4ILYMYNYOYPY(D/f)gQY5I)DRY6ISYTYUY:f7I8IVYWYXYYY;fZY9I0Y1Y!I:EIe#I$I%I'I(I)I*I+I,I-I*DJe`e.ILi+D/I:I;I=I?I@I[I]I^I_I`I{I|I}I=f~IaJ2Y3Y4Y5YbJ6Y7Y8Y?f9Y!YcJdJ#Y$Y%Y'Y@f(YeJ)Y*Y+Y,Y-Y.Y/Y:Y;Y=Y?Y@Y,DKe-D.D/D:D;D=D?D@D[DLe]DDh^D_D`D[fMeNe]f^f_f`f{fVd|f}f*g(c+g,g{e|e-g.g/g:g;gWd)c=g}eXd?g@g[g~ffJgJ[Y]Y^Y_Y`YYdMiNihJiJjJkJlJmJ{DOe|D}D~DaEbEcEdEeEfEgEPehEEhiEjEkE{Yag|YnJoJ}Y~YaZbZbgcZpJdZeZqJ;EQerJsJtJuJvJwJxJyJzJAJBJlERe~eCJOimEDJEJFJfZgZhZGJ(ePiQinEoEiZHJcgjZkZlZmZIJJJKJnZoZpZqZrZpEdgsZLJqErEvgsEtEwgtZMJuZvZwZNJegxZyZzZAZOJPJQJBZCZDZEZFZuEfgGZRJvEHZSJIZJZKZTJggLZMZNZOZUJVJWJPZQZRZSZTZwEhgUZXJxEVZYJWZXZYZZJigZZ0Z1Z2Z0J1J2J3Z4Z5Z6Z7ZyEjg8Z3JzE9Z4J!Z#Z$Z5Jkg%Z'Z(Z)Z6J7J8J*Z+Z,Z-Z.ZAElg/Z9JBE:Z!J;Z=Z?Z#Jmg@Z[Z]Z^Z$J%J'J_Z`Z{Z|Z}ZCEng~Z(JDEa0)Jb0c0d0*Joge0f0g0h0+J,J-Ji0j0k0l0m0EEpgn0.JFEGEo0/Jp0q0r0:Jqgs0t0u0v0;J=J?Jw0x0y0z0A0HErgB0@JIEJEC0[JD0E0XcF0YcG0=EH0ZcI0J0K0L0M0N0O0P0Q0R0S0T0U0V0W0X0Y0Z000102030405060708090!0#0$0%0'0(0)0*0+0,0c -00c.0KE2bLE%d/0:0;0=0?0@0[0]0^0_0`0{0|0}0RiSi]JMENEOE~0PE^J_JQE?E1c+ha1b1c1d1e1f1g1h1i1j1k12cl1m1n1o1p1q1r1s1t1u1v1w1x1`Jy1z1A1B1C1D1E1F1G1H1REI1J1{JK1|JL1M1}JN1O1P1Q1R1S1T1U1V1W1X1Y1Z101112131415161718191!1#1$1%1'1(1)1*1+1,1-1.1/1:1;1=1?1@1[1]1^1_1`1{1|1}1~1a2b2c2d2e2f2g2h2i2j2k2l2m2n2o2p2q2r2s2t2u2v2w2x2~JaKy2z2bKcKdKeKfKA2gKhKiKB2jKC2kKD2lKmKE2nKF2oKG2pKH2qKI2J2rKsKtKuKvKK2wKL2xKM2N2O2yKP2Q2R2S2T2U2V2W2X2Y2Z202122232425262728292!2#2$2%2'2(2)2*2+2,2-2.2/2:2;2=2?2@2
72cdef int HANDLE_RETURN_NVRTC(cynvrtc.nvrtcProgram prog, cynvrtc.nvrtcResult err) except?-1 nogil:
73 """Handle NVRTC result codes, raising NVRTCError with program log on failure."""
74 if err == cynvrtc.nvrtcResult.NVRTC_SUCCESS: 2# $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdbAbBbCbDbEbFbmfnfofhdpfidqfjdrfkdsfldtfmdXbufvfwfxfyfh i [BzfGbHbIbAfJbKbLbMbNbd e j f g 6 eb7 8 fbgbhbibjbObkblbmbnbobpb9 ! zcAcneyibCcCqbodBcpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdoepeb .b/b:b;b=b?b@b[b]b^b_b`b{b|b}b~bacbcccdcecfcgchcicjckclcmcncocdCpcscuCvC#dwCxCyCzCACCcBCWbCCDCECFCGCHCICJCKCLCMCNCOCPCRcScTcQCRCSCUcVcWcTCUCVCWCXCYCZC0C1C2C3C4C5C6C7C8C9C!C#C$C%C'C(C)C*C+C,C-C.C/C:C;C=C?C@C[C]C^C_C`C{C|C}C~CaDbDcDdDeDfDgDhDiDjDkDlDmDnDoDpDqDrDsDtDuDBfvDCfwDxDDfyDEfzDFfGfqeBireDiseteHfIfJfKfLfMfNfOfADPfBDQfCDRfDDSfTfueEDveFDUfVfWfGDweHDIDJDKDLDMDNDODPDxeQDBhRDSDTDXfyezeYfZf0f1f2fRd3f4f%cSd'cTd5fUdUDAeVDWDXDYDZD0D1D2D3D4DBe5DCh6D7D8D6f7fCe9DDe!D8f#D9f$D!f#fEeIiFeKiGeHe$f%f'f(f)f*f+f,f%D-f'D.f(D/f)D:f;fIe*DJe+D=f?f@f,DKe-D.D/D:D;D=D?D@D[DLe]DDh^D_D`D[fMeNe]f^f_f`f{fVd|f}f(cWd)cXd~fYd{DOe|D}D~DaEbEcEdEeEfEgEPehEEhiEjEkEagbgQelERemEnEoEcgpEdgqErEsEeguEfgvEggwEhgxEigyEjgzEkgAElgBEmgCEngDEogEEpgFEGEqgHErgIEJEXcYc'dZc(d)d*d^c_c`c+d,d-d.d3c4c{c/d|c:d;d=d?d}c@d~cadbdcd[d]d^d_d`d{d|d}d~dddedfdc 0cMENEOEgdPEQE1c2caebecedeeefegezKAKBKCKheDKiejeEK
75 return 0 2# $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdbAbBbCbDbEbFbmfnfofhdpfidqfjdrfkdsfldtfmdXbufvfwfxfyfh i [BzfGbHbIbAfJbKbLbMbNbd e j f g 6 eb7 8 fbgbhbibjbObkblbmbnbobpb9 ! zcAcneyibCcCqbodBcpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdoepeb .b/b:b;b=b?b@b[b]b^b_b`b{b|b}b~bacbcccdcecfcgchcicjckclcmcncocdCpcscuCvC#dwCxCyCzCACCcBCWbCCDCECFCGCHCICJCKCLCMCNCOCPCRcScTcQCRCSCUcVcWcTCUCVCWCXCYCZC0C1C2C3C4C5C6C7C8C9C!C#C$C%C'C(C)C*C+C,C-C.C/C:C;C=C?C@C[C]C^C_C`C{C|C}C~CaDbDcDdDeDfDgDhDiDjDkDlDmDnDoDpDqDrDsDtDuDBfvDCfwDxDDfyDEfzDFfGfqeBireDiseteHfIfJfKfLfMfNfOfADPfBDQfCDRfDDSfTfueEDveFDUfVfWfGDweHDIDJDKDLDMDNDODPDxeQDBhRDSDTDXfyezeYfZf0f1f2fRd3f4f%cSd'cTd5fUdUDAeVDWDXDYDZD0D1D2D3D4DBe5DCh6D7D8D6f7fCe9DDe!D8f#D9f$D!f#fEeIiFeKiGeHe$f%f'f(f)f*f+f,f%D-f'D.f(D/f)D:f;fIe*DJe+D=f?f@f,DKe-D.D/D:D;D=D?D@D[DLe]DDh^D_D`D[fMeNe]f^f_f`f{fVd|f}f(cWd)cXd~fYd{DOe|D}D~DaEbEcEdEeEfEgEPehEEhiEjEkEagbgQelERemEnEoEcgpEdgqErEsEeguEfgvEggwEhgxEigyEjgzEkgAElgBEmgCEngDEogEEpgFEGEqgHErgIEJEXcYc'dZc(d)d*d^c_c`c+d,d-d.d3c4c{c/d|c:d;d=d?d}c@d~cadbdcd[d]d^d_d`d{d|d}d~dddedfdc 0cMENEOEgdPEQE1c2caebecedeeefegezKAKBKCKheDKiejeEK
76 with gil: 1c
77 _raise_nvrtc_error(prog, err) 1c
80cdef int _raise_nvrtc_error(cynvrtc.nvrtcProgram prog, cynvrtc.nvrtcResult err) except -1:
81 """Build error message with program log and raise NVRTCError."""
82 cdef const char* err_str = cynvrtc.nvrtcGetErrorString(err) 1c
83 cdef size_t logsize = 0 1c
84 if prog != NULL: 1c
85 cynvrtc.nvrtcGetProgramLogSize(prog, &logsize) 1c
86 cdef bytes log_bytes
87 cdef str log_str = "" 1c
88 if logsize > 1 and prog != NULL: 1c
89 log_bytes = b" " * logsize 1c
90 if cynvrtc.nvrtcGetProgramLog(prog, <char*>log_bytes) == cynvrtc.nvrtcResult.NVRTC_SUCCESS: 1c
91 log_str = log_bytes.decode("utf-8", errors="backslashreplace") 1c
92 err_msg = f"{err}: {err_str.decode()}" if err_str != NULL else f"NVRTC error {err}" 1c
93 if log_str: 1c
94 err_msg += f", compilation log:\n\n{log_str}" 1c
95 raise NVRTCError(err_msg) 1c
98cdef int HANDLE_RETURN_NVVM(cynvvm.nvvmProgram prog, cynvvm.nvvmResult err) except?-1 nogil:
99 """Handle NVVM result codes, raising nvvmError with program log on failure."""
100 if err == cynvvm.nvvmResult.NVVM_SUCCESS: 2tEFKGKKE2bLE%dHKIKJKKKLKMKNKOKPKQKRKSKTKUKRiSi
101 return 0 2tEFKGKKE2bLE%dHKIKJKKKLKMKNKOKPKQKRKSKTKUKRiSi
102 with gil: 22b
103 _raise_nvvm_error(prog, err) 22b
106cdef int _raise_nvvm_error(cynvvm.nvvmProgram prog, cynvvm.nvvmResult err) except -1:
107 """Raise nvvmError annotated with the program log."""
108 cdef size_t logsize = 0 22b
109 if prog != NULL: 22b
110 cynvvm.nvvmGetProgramLogSize(prog, &logsize) 22b
111 cdef bytes log_bytes
112 cdef str log_str = "" 22b
113 if logsize > 1 and prog != NULL: 22b
114 log_bytes = b" " * logsize 22b
115 if cynvvm.nvvmGetProgramLog(prog, <char*>log_bytes) == cynvvm.nvvmResult.NVVM_SUCCESS: 2a 2b
116 log_str = log_bytes.decode("utf-8", errors="backslashreplace") 22b
117 cdef object exc = nvvmError(err) 22b
118 if log_str: 22b
119 exc.args = (exc.args[0] + f"\nNVVM program log: {log_str}", *exc.args[1:]) 22b
120 raise exc 22b
123cdef int HANDLE_RETURN_NVJITLINK(
124 cynvjitlink.nvJitLinkHandle handle, cynvjitlink.nvJitLinkResult err) except?-1 nogil:
125 """Handle nvJitLink result codes, raising nvJitLinkError with error log on failure."""
126 if err == cynvjitlink.nvJitLinkResult.NVJITLINK_SUCCESS: 2b VKWKXKYKZK0K1K2K3K4K5K@E[E6K7K8K9K!K#K$K%K'K(K)K*K+Kjf,K-K.K/K:K+h;K=K?K@K[K]K^K_K`K{K
127 return 0 2b VKWKXKYKZK0K1K2K3K4K5K@E[E6K7K8K9K!K#K$K%K'K(K)K*K+K,K-K.K/K:K+h;K=K?K@K[K]K^K_K`K{K
128 with gil: 2b jf
129 _raise_nvjitlink_error(handle, err) 2b jf
132cdef int _raise_nvjitlink_error(
133 cynvjitlink.nvJitLinkHandle handle, cynvjitlink.nvJitLinkResult err) except -1:
134 """Raise nvJitLinkError annotated with the error log."""
135 cdef size_t logsize = 0 2b jf
136 if handle != NULL: 2b jf
137 cynvjitlink.nvJitLinkGetErrorLogSize(handle, &logsize) 1b
138 cdef bytes log_bytes
139 cdef str log_str = "" 2b jf
140 if logsize > 1 and handle != NULL: 2b jf
141 log_bytes = b" " * logsize 1b
142 if cynvjitlink.nvJitLinkGetErrorLog(handle, <char*>log_bytes) == \ 1b
143 cynvjitlink.nvJitLinkResult.NVJITLINK_SUCCESS: 1b
144 log_str = log_bytes.decode("utf-8", errors="backslashreplace") 1b
145 cdef object exc = nvJitLinkError(err) 2b jf
146 if log_str: 2b jf
147 exc.args = (exc.args[0] + f"\nnvJitLink error log: {log_str}", *exc.args[1:]) 1b
148 raise exc 2b jf
151cdef object _RUNTIME_SUCCESS = runtime.cudaError_t.cudaSuccess
152cdef object _NVRTC_SUCCESS = nvrtc.nvrtcResult.NVRTC_SUCCESS
155cpdef inline int _check_driver_error(cydriver.CUresult error) except?-1 nogil:
156 if error == cydriver.CUresult.CUDA_SUCCESS: 2a ZdTi0dUikfVi# Wi$ XilfYi% Zi' 0i( 1i) 2i* 3i+ 4i, 5i- 6i. 7i/ 8i: 9i; !i= #i? $i@ %i[ 'i] (i^ )i_ *i` +i{ ,i| -i} .i~ /iab:ibb;icb=idb?i@i[i]i^i]g_i`i{i|i}i~iajbjcjdjejfjgjhjijSejjkjljTemjUenjVeojWepjAbqjXerjYesjZetjujvjBbwjCbxjDbyjEbzjFbAjDcBjEcCjFcDjGcEjHcFjIcGjJcHjKcIjJj,hKj-hLj.hMjygNjzgOjAgPjQjRjSjTjUjVjWjXjYjZj0j1j2j3j4j5j6j7j8j9j!j#j$j%j'j(j)j*j+j,j-j.j/j:j;j=j?j@j[j]j^j_j`j{j|j}j~jakbkckdkekfkgkhkikjkkklkmknkokpkqkrksktkukvkwkxkykzkAkBkCkDkEkFkGkHkIkJkKkLkMkNkOkPkQkRkSkTkUkVkWkXkYkZk0k1k2k3k4k5k6k7k8k9k!k#k$k%k'k(k)k*k+k,k-k.k/k:k;k=k?k@k[k]k^k_k`k{k|k}k~kalblcldlelflglhliljlklllmlnlol^gplqlrlsltlulvlwl_gxlylzlAlXbBlCl`gDl0eElFlGlHlIlJlKlLlMlNlOlPlQlRlSlTl1eUl2eVlWlafXlYlZl*c0l+c1l2l3lh 4li 5l6l7lGb8lHb9lIb!l#l3e$l{g%lJb'lKb(lLb)lMb*lNb+l,l-l.l/ld :l|g;l4e=l}g?l,c-c.c1d2d3de j f g 6 eb7 8 fbgbhb4d5d6d7d8d9d@lib[ljb]lOb^lkb_llb`lmb{lnb|lob}lpb~lambfbmcm~gah/hbhch:hdheh;hfh=hdmBgemcf?hfm5egmdfef@hffhmgf[himCgjmgh]hhhih^hhfkmiflmmmnmompmqm6erm9 ! sg!d/cYb:c+bLc;c=ctgQbRbSbTbUbVbsmtmumvmwmxmymSE_hzmAmBmCmDmEmFmGmHmImJmKmLmMmnhNmOmPmohQmphzcRmSmTmUmVmWmXmYmAcZm0m1m2m3m4m5m6m7m8mqhrhndsh#c9m$c!mthuhvh#mwhxhyh$m%m'm(m)m*m+m,m-m.m/m:m;m=m?m@m[m]m^mqb,b-b_m`m{m|mpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdoepe}m~manbncndnenfngnhninjnknlnmnnnonpnqnrnsntnunvnwnxnynznAnBnCnDnEnFnGnHnIgqcrcJgKgLgMgNgOgPgQgzhAhke5crbtcsb6cuctbZb)e*e+ek 9bl m n o p q r s t !bu v w x 3by z 4b#ble7cubvcvb8cwcwb0b,e-e.eA $bB C D E F G H I J %bK L M N 5bO P 6b'bme9cxbxcyb!cyczb1b/e:e;eQ (bR S T U V W X Y Z )b0 1 2 3 7b4 5 8b*bRgInJnKn9eLnMnNnOnPnQnRnSnTnUnVnWnXnYnZn0n1n2n3n4n5n6n7n8n9n!n#n$n%n'n(n)n*n+n,n-n.n/n:n;n=n?n@n[n]n^n_n`n{n|n}n~naobocodoeofogohoiojokolomoSgTg$enooopoqorososCtotCuoNcOcPbPcvoQcwoxoyo#dzoAoBoCoDoEoCcFoGoWbHoIoJoKoLoMoNoOoPoQoRoSoToUoVoWoXoYoZo0o1o2o3o4o5o6o7o8o9o!o#o$o%o'o(o)o*o+o,o-o.o/o:o;o=o?o@o[o]o^o_o`o{o|o}o~oapbpcpdpepfpgphpipjpkplpmpnpopppqprpsptpupvpwpxpypzpApBpCpDpEpFpGpHpIpJpKpLpMpNpOpPpQpRpSpTpUpVpWpXpYpZp0p1p2p3p4p5p6p7p8p9p!p#p$p%p'p(p)p*p+p,p-p.p/p:p;p=p?p@p[p]p^p_p`p{p|p}p~paqbqcqdqeqfqgqhqiqjqkqlqmqnqoqpqqqrqsqtquqvqwqxqyqzqAqBqCqDqEqFqGqHqIqJqKqLqMqNqOqPqQqRqSqTqUqVqWqXqYqZq0q1q2q3q4q5q6q7q8q9q!q#q$q%q'q(q)q*q+q,q-q.q/q:q;q=q?q@q[q]q^q_q`q{q|q}q~qarbrcrdrerfrgrhrirjrkrlrmrnrorprqrrrsrtrurvrwrxryrzrArBrCrDrErFrGrHrIrJrKrLrMrNrOrPrQrRrSrTrUrVrWrXrYrZr0r1r2r3r4r5r6r7r8r9r!r#r$r%r'r(r)r*r+r,r-r.r/r:r;r=r?r@r[r]r^r_r`r{r|r}r~rasbscsdsesfsgshsisjskslsmsnsospsqsrssstsusvswsxsyszsAsBsCsDsEsFsGsHsIsJsKsLsMsNsOsPsQsRsSsTsUsVsWsXsYsZs0s1s2s3s4s5s6s7s8s9s!s#s$s%s's(s)s*s+s,s-s.s/s:s;s=s?s@s[s]s^s_s`s{s|s}s~satbtctdtetftgthtitjtktltmtntotptqtrtstttutvtwtxtytztAtBtCtDtEtFtGtHtItJtKtLtMtNtOtPtQtRtStTtUtVtWtXtYtZt0t1t2t3t4t5t6t7t8t9t!t#t$t%t't(t)t*t+t,t-t.t/t:t;t=t?t@t[t]t^t_t`t{t|t}t~taubucudueufuguhuiujukulumunuoupuqurusutuuuvuwuxuyuzuAuBuCuDuEuFuGuHuIuJuKuLuMuNuOuPuQuRuSuTuUuVuWuXuYuZu0u1u2u3u4u5u6u7u8u9u!u#u$u%u'u(u)u*u+u,u-u.u/u:u;u=u?u@u[u]u^u_u`u{u|u}u~uavbvcvdvevfvgvhvivjvkvlvmvnvovpvqvrvsvtvuvvvwvxvyvzvAvBvCvDvEvFvGvHvIvJvKvLvMvNvOvPvQvRvSvTvUvVvWvXvYvZv0v1v2v3v4v5v6v7v8v9v!v#v$v%v'v(v)v*v+v,v-v.v/v:v;v=v?v@v[v]v^v_v`v{v|v}v~vawbwcwdwewfwgwhwiwjwkwlwmwnwowpwqwrwswtwuwvwwwxwywzwAwBwCwDwEwFwGwHwIwJwKwLwMwNwOwPwQwRwSwTwUwVwWwXwYwZw0w1w2w3w4w5w6w7w8w9w!w#w$w%w'w(w)w*w+w,w-w.w/w:w;w=w?w@w[w]w^w_w`w{w(e|w}w~waxbxcxdxexfxgxhxixjxkxlxmxnxoxpxqxrxsxtxuxvxwxxxyxzxAxBxCxDxExFxGxHxIxJxKxLxMxNxOxPxQxRxSxTxUxVxWxXxYxZx0x1x2x3x4x5x6x7x8x9x!x#x$x%x'x(x)x*x+x,x-x.x/x:x;x=x?x@x[x]x^x_x`x{x|x}x~xaybycydyeyfygyhyiyjykylymynyoypyqyrysytyuyvywyxyyyzyAyByCyDyEyFyGyHyIyJyKyLyMyNyOyPyQyRySyTyUyVyWyXyYyZy0y1y2y3y4y5y6y7y8y9y!y#y$y%y'y(y)y*y+y,y-y.y/y:y;y=y?y@y[y]y^y_y`y{y|y}y~yazbzczdzezfzgzhzizjzkzlzmznzozpzqzrzsztzuzvzwzxzyzzzAzBzCzDzEzFzGzHzIzJzKzLzMzNzOzPzQzRzSzTzUzVzWzXzYzZz0z1z2z3z4z5z6z7z8z9z!z#z%d$z%z'z(z)z*z+z,z-z.z/z:z;z=z?z@z[z]z^z_z`z{z|z}z~zaAbAcAdAeAfAgAhAiAjAkAlAmAnAoApAqArAsAtAuAvAwAxAyAzAAABACADAEAFAGAHAIAJAKALAMANAOAPAQARASATAUAVAWAXAYAZA0A1A2A3A4A5A6A7A8A9A!A#A$A%A'A(A)A*A+A,A-A.A/A:A;A=A?A@A[A]A^A_A`A{A|A}A~AaBbBcBdBeBfBgBhBiBjBkBlBmBnBoBpBqBrBsBtBuBvBwBxByBzBABBBCBDBEBFBGBHBIBJBKBLBMBNBOBPBQBRBSBTBUBVBWBXBYBZB0B1B2B3B4B5B6B7B8B9B!B#B$B%B'B(B)B*B+B,B-B.B/B
157 return 0 2a ZdTi0dUikfVi# Wi$ XilfYi% Zi' 0i( 1i) 2i* 3i+ 4i, 5i- 6i. 7i/ 8i: 9i; !i= #i? $i@ %i[ 'i] (i^ )i_ *i` +i{ ,i| -i} .i~ /iab:ibb;icb=idb?i@i[i]i^i_i`i{i|i}i~iajbjcjdjejfjgjhjijSejjkjljTemjUenjVeojWepjAbqjXerjYesjZetjujvjBbwjCbxjDbyjEbzjFbAjDcBjEcCjFcDjGcEjHcFjIcGjJcHjKcIjJj,hKj-hLj.hMjygNjzgOjAgPjQjRjSjTjUjVjWjXjYjZj0j1j2j3j4j5j6j7j8j9j!j#j$j%j'j(j)j*j+j,j-j.j/j:j;j=j?j@j[j]j^j_j`j{j|j}j~jakbkckdkekfkgkhkikjkkklkmknkokpkqkrksktkukvkwkxkykzkAkBkCkDkEkFkGkHkIkJkKkLkMkNkOkPkQkRkSkTkUkVkWkXkYkZk0k1k2k3k4k5k6k7k8k9k!k#k$k%k'k(k)k*k+k,k-k.k/k:k;k=k?k@k[k]k^k_k`k{k|k}k~kalblcldlelflglhliljlklllmlnlolplqlrlsltlulvlwlxlylzlAlXbBlClDl0eElFlGlHlIlJlKlLlMlNlOlPlQlRlSlTl1eUl2eVlWlafXlYlZl*c0l+c1l2l3lh 4li 5l6l7lGb8lHb9lIb!l#l3e$l{g%lJb'lKb(lLb)lMb*lNb+l,l-l.l/ld :l;l4e=l?l,c-c.c1d2d3de j f g 6 eb7 8 fbgbhb4d5d6d7d8d9d@lib[ljb]lOb^lkb_llb`lmb{lnb|lob}lpb~lambfbmcm~gah/hbhch:hdheh;hfh=hdmBgemcf?hfm5egmdfef@hffhmgf[himCgjmgh]hhhih^hhfkmiflmmmnmompmqm6erm9 ! sg!d/cYb:c+bLc;c=ctgQbRbSbTbUbVbsmtmumvmwmxmym_hzmAmBmCmDmEmFmGmHmImJmKmLmMmnhNmOmPmohQmphzcRmSmTmUmVmWmXmYmAcZm0m1m2m3m4m5m6m7m8mqhrhndsh#c9m$c!mthuhvh#mwhxhyh$m%m'm(m)m*m+m,m-m.m/m:m;m=m?m@m[m]m^mqb,b-b_m`m{m|mpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdoepe}m~manbncndnenfngnhninjnknlnmnnnonpnqnrnsntnunvnwnxnynznAnBnCnDnEnFnGnHnIgqcrcJgKgLgMgNgOgPgQgzhAhke5crbtcsb6cuctbZb)e*e+ek 9bl m n o p q r s t !bu v w x 3by z 4b#ble7cubvcvb8cwcwb0b,e-e.eA $bB C D E F G H I J %bK L M N 5bO P 6b'bme9cxbxcyb!cyczb1b/e:e;eQ (bR S T U V W X Y Z )b0 1 2 3 7b4 5 8b*bRgInJnKn9eLnMnNnOnPnQnRnSnTnUnVnWnXnYnZn0n1n2n3n4n5n6n7n8n9n!n#n$n%n'n(n)n*n+n,n-n.n/n:n;n=n?n@n[n]n^n_n`n{n|n}n~naobocodoeofogohoiojokolomoSgTg$enooopoqorososCtotCuoNcOcPbPcvoQcwoxoyo#dzoAoBoCoDoEoCcFoGoWbHoIoJoKoLoMoNoOoPoQoRoSoToUoVoWoXoYoZo0o1o2o3o4o5o6o7o8o9o!o#o$o%o'o(o)o*o+o,o-o.o/o:o;o=o?o@o[o]o^o_o`o{o|o}o~oapbpcpdpepfpgphpipjpkplpmpnpopppqprpsptpupvpwpxpypzpApBpCpDpEpFpGpHpIpJpKpLpMpNpOpPpQpRpSpTpUpVpWpXpYpZp0p1p2p3p4p5p6p7p8p9p!p#p$p%p'p(p)p*p+p,p-p.p/p:p;p=p?p@p[p]p^p_p`p{p|p}p~paqbqcqdqeqfqgqhqiqjqkqlqmqnqoqpqqqrqsqtquqvqwqxqyqzqAqBqCqDqEqFqGqHqIqJqKqLqMqNqOqPqQqRqSqTqUqVqWqXqYqZq0q1q2q3q4q5q6q7q8q9q!q#q$q%q'q(q)q*q+q,q-q.q/q:q;q=q?q@q[q]q^q_q`q{q|q}q~qarbrcrdrerfrgrhrirjrkrlrmrnrorprqrrrsrtrurvrwrxryrzrArBrCrDrErFrGrHrIrJrKrLrMrNrOrPrQrRrSrTrUrVrWrXrYrZr0r1r2r3r4r5r6r7r8r9r!r#r$r%r'r(r)r*r+r,r-r.r/r:r;r=r?r@r[r]r^r_r`r{r|r}r~rasbscsdsesfsgshsisjskslsmsnsospsqsrssstsusvswsxsyszsAsBsCsDsEsFsGsHsIsJsKsLsMsNsOsPsQsRsSsTsUsVsWsXsYsZs0s1s2s3s4s5s6s7s8s9s!s#s$s%s's(s)s*s+s,s-s.s/s:s;s=s?s@s[s]s^s_s`s{s|s}s~satbtctdtetftgthtitjtktltmtntotptqtrtstttutvtwtxtytztAtBtCtDtEtFtGtHtItJtKtLtMtNtOtPtQtRtStTtUtVtWtXtYtZt0t1t2t3t4t5t6t7t8t9t!t#t$t%t't(t)t*t+t,t-t.t/t:t;t=t?t@t[t]t^t_t`t{t|t}t~taubucudueufuguhuiujukulumunuoupuqurusutuuuvuwuxuyuzuAuBuCuDuEuFuGuHuIuJuKuLuMuNuOuPuQuRuSuTuUuVuWuXuYuZu0u1u2u3u4u5u6u7u8u9u!u#u$u%u'u(u)u*u+u,u-u.u/u:u;u=u?u@u[u]u^u_u`u{u|u}u~uavbvcvdvevfvgvhvivjvkvlvmvnvovpvqvrvsvtvuvvvwvxvyvzvAvBvCvDvEvFvGvHvIvJvKvLvMvNvOvPvQvRvSvTvUvVvWvXvYvZv0v1v2v3v4v5v6v7v8v9v!v#v$v%v'v(v)v*v+v,v-v.v/v:v;v=v?v@v[v]v^v_v`v{v|v}v~vawbwcwdwewfwgwhwiwjwkwlwmwnwowpwqwrwswtwuwvwwwxwywzwAwBwCwDwEwFwGwHwIwJwKwLwMwNwOwPwQwRwSwTwUwVwWwXwYwZw0w1w2w3w4w5w6w7w8w9w!w#w$w%w'w(w)w*w+w,w-w.w/w:w;w=w?w@w[w]w^w_w`w{w(e|w}w~waxbxcxdxexfxgxhxixjxkxlxmxnxoxpxqxrxsxtxuxvxwxxxyxzxAxBxCxDxExFxGxHxIxJxKxLxMxNxOxPxQxRxSxTxUxVxWxXxYxZx0x1x2x3x4x5x6x7x8x9x!x#x$x%x'x(x)x*x+x,x-x.x/x:x;x=x?x@x[x]x^x_x`x{x|x}x~xaybycydyeyfygyhyiyjykylymynyoypyqyrysytyuyvywyxyyyzyAyByCyDyEyFyGyHyIyJyKyLyMyNyOyPyQyRySyTyUyVyWyXyYyZy0y1y2y3y4y5y6y7y8y9y!y#y$y%y'y(y)y*y+y,y-y.y/y:y;y=y?y@y[y]y^y_y`y{y|y}y~yazbzczdzezfzgzhzizjzkzlzmznzozpzqzrzsztzuzvzwzxzyzzzAzBzCzDzEzFzGzHzIzJzKzLzMzNzOzPzQzRzSzTzUzVzWzXzYzZz0z1z2z3z4z5z6z7z8z9z!z#z$z%z'z(z)z*z+z,z-z.z/z:z;z=z?z@z[z]z^z_z`z{z|z}z~zaAbAcAdAeAfAgAhAiAjAkAlAmAnAoApAqArAsAtAuAvAwAxAyAzAAABACADAEAFAGAHAIAJAKALAMANAOAPAQARASATAUAVAWAXAYAZA0A1A2A3A4A5A6A7A8A9A!A#A$A%A'A(A)A*A+A,A-A.A/A:A;A=A?A@A[A]A^A_A`A{A|A}A~AaBbBcBdBeBfBgBhBiBjBkBlBmBnBoBpBqBrBsBtBuBvBwBxByBzBABBBCBDBEBFBGBHBIBJBKBLBMBNBOBPBQBRBSBTBUBVBWBXBYBZB0B1B2B3B4B5B6B7B8B9B!B#B$B%B'B(B)B*B+B,B-B.B/B
158 cdef const char* name
159 name_err = cydriver.cuGetErrorName(error, &name) 2]g^g_g`g|g4e}gSE_hPbWb%d
160 if name_err != cydriver.CUresult.CUDA_SUCCESS: 2]g^g_g`g|g4e}gSE_hPbWb%d
161 raise CUDAError(f"UNEXPECTED ERROR CODE: {error}") 2_h
162 with gil: 2]g^g_g`g|g4e}gSE_hPbWb%d
163 # TODO: consider lower this to Cython
164 expl = DRIVER_CU_RESULT_EXPLANATIONS.get(int(error)) 2]g^g_g`g|g4e}gSE_hPbWb%d
165 if expl is not None: 2]g^g_g`g|g4e}gSE_hPbWb%d
166 raise CUDAError(f"{name.decode()}: {expl}") 2]g^g_g`g|g4e}gSE_hPbWb%d
167 cdef const char* desc
168 desc_err = cydriver.cuGetErrorString(error, &desc)
169 if desc_err != cydriver.CUresult.CUDA_SUCCESS:
170 raise CUDAError(f"{name.decode()}")
171 raise CUDAError(f"{name.decode()}: {desc.decode()}")
174cpdef inline int _check_runtime_error(error) except?-1:
175 if error == _RUNTIME_SUCCESS: 2QbRbSbTbUbVbDg|KFhEg7e8eFgGgHg,b-bqcrcrbsbtbZbk l m n o p q r s t u v w x y z ubvbwb0bA B C D E F G H I J K L M N O P xbybzb1bQ R S T U V W X Y Z 0 1 2 3 4 5 ?cMc!e@c#e[c
176 return 0 2QbRbSbTbUbVbDgFhEg7e8eFgGgHg,b-bqcrcrbsbtbZbk l m n o p q r s t u v w x y z ubvbwb0bA B C D E F G H I J K L M N O P xbybzb1bQ R S T U V W X Y Z 0 1 2 3 4 5 ?cMc!e@c#e[c
177 # `_check_error()` reaches this path only for `runtime.cudaError_t` values.
178 # Use the enum name directly because Windows hybrid cudart can lag that table.
179 name = error.name 2|KFh
180 expl = RUNTIME_CUDA_ERROR_EXPLANATIONS.get(int(error)) 2|KFh
181 if expl is not None: 2|KFh
182 raise CUDAError(f"{name}: {expl}") 2|KFh
183 desc_err, desc = runtime.cudaGetErrorString(error) 2Fh
184 if desc_err != _RUNTIME_SUCCESS: 2Fh
185 raise CUDAError(f"{name}")
186 desc = desc.decode() 2Fh
187 raise CUDAError(f"{name}: {desc}") 2Fh
190cpdef inline int _check_nvrtc_error(error, handle=None) except?-1:
191 if error == _NVRTC_SUCCESS: 2a :B}Kb .b/b:b;b=b?b@b[b]b^b_b`b{b|b}b~bacbcccdcecfcgchcicjckclcmcncocpcscRcScTcGhUcHhVcIhWcvgwgXcYc'dZc(d)d*d^c_c`c+d,d-d.d3c4c{c/d|c:d;d=d?d}c@d~cadbdcd[d]d^d_d`d{d|d}d~dddedfdc 0cgd1c2cUgVgWgXgYgZgJhKhLhMhNhOhPhQhRhShThUhVhWhXhYhZh0h1h2h3h4h5h6h7h8h9haebecedeeefegeheieje
192 return 0 2a }Kb .b/b:b;b=b?b@b[b]b^b_b`b{b|b}b~bacbcccdcecfcgchcicjckclcmcncocpcscRcScTcGhUcHhVcIhWcvgwgXcYc'dZc(d)d*d^c_c`c+d,d-d.d3c4c{c/d|c:d;d=d?d}c@d~cadbdcd[d]d^d_d`d{d|d}d~dddedfdc 0cgd1c2cUgVgWgXgYgZgJhKhLhMhNhOhPhQhRhShThUhVhWhXhYhZh0h1h2h3h4h5h6h7h8h9haebecedeeefegeheieje
193 err = f"{error}: {nvrtc.nvrtcGetErrorString(error)[1].decode()}" 2:B}K
194 if handle is not None: 2:B}K
195 _, logsize = nvrtc.nvrtcGetProgramLogSize(handle) 2:B
196 log = b" " * logsize 2:B
197 _ = nvrtc.nvrtcGetProgramLog(handle, log) 2:B
198 err += f", compilation log:\n\n{log.decode('utf-8', errors='backslashreplace')}" 2:B
199 raise NVRTCError(err) 2:B}K
202cdef inline int _check_error(error, handle=None) except?-1:
203 if isinstance(error, driver.CUresult): 2a ZdTi0dUikfVi# Wi$ XilfYi% Zi' 0i( 1i) 2i* 3i+ 4i, 5i- 6i. 7i/ 8i: 9i; !i= #i? $i@ %i[ 'i] (i^ )i_ *i` +i{ ,i| -i} .i~ /iab:ibb;icb=idb?i@i[i]i^i_i`i{i|i}i~iajbjcjdjejfjgjhjijSejjkjljTemjUenjVeojWepjAbqjXerjYesjZetjujvjBbwjCbxjDbyjEbzjFbAjDcBjEcCjFcDjGcEjHcFjIcGjJcHjKcIjJj,hKj-hLj.hMjygNjzgOjAgPjQjRjSjTjUjVjWjXjYjZj0j1j2j3j4j5j6j7j8j9j!j#j$j%j'j(j)j*j+j,j-j.j/j:j;j=j?j@j[j]j^j_j`j{j|j}j~jakbkckdkekfkgkhkikjkkklkmknkokpkqkrksktkukvkwkxkykzkAkBkCkDkEkFkGkHkIkJkKkLkMkNkOkPkQkRkSkTkUkVkWkXkYkZk0k1k2k3k4k5k6k7k8k9k!k#k$k%k'k(k)k*k+k,k-k.k/k:k;k=k?k@k[k]k^k_k`k{k|k}k~kalblcldlelflglhliljlklllmlnlolplqlrlsltlulvlwlxlylzlAlXbBlClDl0eElFlGlHlIlJlKlLlMlNlOlPlQlRlSlTl1eUl2eVlWlafXlYlZl*c0l+c1l2l3lh 4li 5l6l7lGb8lHb9lIb!l#l3e$l{g%lJb'lKb(lLb)lMb*lNb+l,l-l.l/ld :l;l=l?l,c-c.c1d2d3de j f g 6 eb7 8 fbgbhb4d5d6d7d8d9d@lib[ljb]lOb^lkb_llb`lmb{lnb|lob}lpb~lambfbmcm~gah/hbhch:hdheh;hfh=hdmBgemcf?hfm5egmdfef@hffhmgf[himCgjmgh]hhhih^hhfkmiflmmmnmompmqm6erm9 ! sg!d/cYb:c+bLc;c=ctgQbRbSbTbUbVbDgsmtmumvmwmxmymzmEg7eAm8eBmCmDmEmFmGmHmImFgJmKmLmMmnhNmOmPmohGgHgQmphRmSmTmUmVmWmXmYmZm0m1m2m3m4m5m6m7m8mqhrhsh9m!mthuhvh#mwhxhyh$m%m'm(m)m*m+m,m-m.m/m:m;m=m?m@m[m]m^mqb,b-b_m`m{m|mb .b}m/b~m:ban;bbn=bcn?bdn@ben[bfn]bgn^bhn_bin`bjn{bkn|bln}bmn~bnnaconbcpnccqndcrnecsnfctngcunhcvnicwnjcxnkcynlcznmcAnncBnocCnDnpcEnFnGnHnIgqcrcJgKgLgMgNgOgPgQgzhAhke5crbtcsb6cuctbZb)e*e+ek 9bl m n o p q r s t !bu v w x 3by z 4b#ble7cubvcvb8cwcwb0b,e-e.eA $bB C D E F G H I J %bK L M N 5bO P 6b'bme9cxbxcyb!cyczb1b/e:e;eQ (bR S T U V W X Y Z )b0 1 2 3 7b4 5 8b*bRg?cInJnKn9eLnMnNnOnPnQnRnMcSn!eTn@cUn#eVnWnXnYnZn0n1n2n3n4n5n6n7n8n9n!n#n$n%n'n(n)n*n+n,n-n.n/n:n;n=n?n@n[n]n^n_n`n{n|n}n~naobocodoeofogohoiojokolomoSgTg$eno[coopoqorosotouovoscwoxoyo#dzoAoBoCoDoEoFoGoWbHoIoJoKoLoMoNoOoPoQoRoSoToUoRcVoScWoTcXoYoZo0oGhUc1oHhVc2oIhWc3o4o5o6o7o8o9o!o#o$o%o'o(o)o*o+o,o-o.o/o:o;o=o?o@o[o]o^o_o`o{o|o}o~oapbpcpdpepfpgphpipjpkplpmpnpopppqprpsptpupvpwpxpypzpApBpCpDpEpFpGpHpIpJpKpLpMpNpOpPpQpRpSpTpUpVpWpXpYpZp0p1p2p3p4p5p6p7p8p9p!p#p$p%p'p(p)p*p+p,p-p.p/p:p;p=p?p@p[p]p^p_p`p{p|p}p~paqbqcqdqeqfqgqhqiqjqkqlqmqnqoqpqqqrqsqtquqvqwqxqyqzqAqBqCqDqEqFqGqHqIqJqKqLqMqNqOqPqQqRqSqTqUqVqWqXqYqZq0q1q2q3q4q5q6q7q8q9q!q#q$q%q'q(q)q*q+q,q-q.q/q:q;q=q?q@q[q]q^q_q`q{q|q}q~qarbrcrdrerfrgrhrirjrkrlrmrnrorprqrrrsrtrurvrwrxryrzrArBrCrDrErFrGrHrIrJrKrLrMrNrOrPrQrRrSrTrUrVrWrXrYrZr0r1r2r3r4r5r6r7r8r9r!r#r$r%r'r(r)r*r+r,r-r.r/r:r;r=r?r@r[r]r^r_r`r{r|r}r~rasbscsdsesfsgshsisjskslsmsnsospsqsrssstsusvswsxsyszsAsBsCsDsEsFsGsHsIsJsKsLsMsNsOsPsQsRsSsTsUsVsWsXsYsZs0s1s2s3s4s5s6s7s8s9s!s#s$s%s's(s)s*s+s,s-s.s/s:s;s=s?s@s[s]s^s_s`s{s|s}s~satbtctdtetftgthtitjtktltmtntotptqtrtstttutvtwtxtytztAtBtCtDtEtFtGtHtItJtKtLtMtNtOtPtQtRtStTtUtVtWtXtYtZt0t1t2t3t4t5t6t7t8t9t!t#t$t%t't(t)t*t+t,t-t.t/t:t;t=t?t@t[t]t^t_t`t{t|t}t~taubucudueufuguhuiujukulumunuoupuqurusutuuuvuwuxuyuzuAuBuCuDuEuFuGuHuIuJuKuLuMuNuOuPuQuRuSuTuUuVuWuXuYuZu0u1u2u3u4u5u6u7u8u9u!u#u$u%u'u(u)u*u+u,u-u.u/u:u;u=u?u@u[u]u^u_u`u{u|u}u~uavbvcvdvevfvgvhvivjvkvlvmvnvovpvqvrvsvtvuvvvwvxvyvzvAvBvCvDvEvFvGvHvIvJvKvLvMvNvOvPvQvRvSvTvUvVvWvXvYvZv0v1v2v3v4v5v6v7v8v9v!v#v$v%v'v(v)v*v+v,v-v.v/v:v;v=v?v@v[v]v^v_v`v{v|v}v~vawbwcwdwewfwgwhwiwjwkwlwmwnwowpwqwrwswtwuwvwwwxwywzwAwBwCwDwEwFwGwHwIwJwKwLwMwNwOwPwQwRwSwTwUwVwWwXwYwZw0w1w2w3w4w5w6w7w8w9w!w#w$w%w'w(w)w*w+w,w-w.w/w:w;w=w?w@w[w]w^w_w`w{w(e|w}w~waxbxcxdxexfxgxhxixjxkxlxmxnxoxpxqxrxsxtxvguxvxwxwgxxyxzxAxBxCxDxExFxGxHxIxJxKxLxMxNxOxPxQxRxSxTxUxVxWxXxYxZx0x1x2x3x4x5x6x7x8x9x!x#x$x%x'x(x)x*x+x,x-x.x/x:x;x=x?x@x[x]x^x_x`x{x|x}x~xaybycydyeyfygyhyiyjykylymynyoypyqyrysytyuyvywyxyyyzyAyByCyDyEyFyGyHyIyJyKyLyMyNyOyPyQyRySyTyUyVyWyXyYyZy0y1y2y3y4y5y6y7y8y9y!y#y$y%y'y(y)y*y+y,y-y.y/y:y;y=y?y@y[y]y^y_y`y{y|y}y~yazbzczdzezfzgzhzizjzkzlzmznzozpzqzrzsztzuzXcvzYcwz'dxzZcyz(dzz)dAz*dBz^cCz_cDz`cEz+dFz,dGz-dHz.dIz3cJz4cKz{cLz/dMz|cNz:dOz;dPz=dQz?dRz}cSz@dTz~cUzadVzbdWzcdXz[dYz]dZz^d0z_d1z`d2z{d3z|d4z}d5z~d6zdd7zed8zfd9zc !z0c#z$z%z'z(z)z*z+z,z-z.z/z:z;z=zgd?z1c@z[z]z^z_z`z{z|z}z~zaA2cbAUgVgWgXgYgZgJhKhLhMhNhOhPhQhRhShThUhVhWhcAXhYhZh0h1h2h3h4h5h6hdAeAfAgA7h8h9haebecedeeehAfegeheiejeiAjAkAlAmAnAoApAqArAsAtAuAvAwAxAyAzAAABACADAEAFAGAHAIAJAKALAMANAOAPAQARASATAUAVAWAXAYAZA0A1A2A3A4A5A6A7A8A9A!A#A$A%A'A(A)A*A+A,A-A.A/A:A;A=A?A@A[A]A^A_A`A{A|A}A~AaBbBcBdBeBfBgBhBiBjBkBlBmBnBoBpBqBrBsBtBuBvBwBxByBzBABBBCBDBEBFBGBHBIBJBKBLBMBNBOBPBQBRBSBTBUBVBWBXBYBZB0B1B2B3B4B5B6B7B8B9B!B#B$B%B'B(B)B*B+B,B-B.B/B
204 return _check_driver_error(error) 2a ZdTi0dUikfVi# Wi$ XilfYi% Zi' 0i( 1i) 2i* 3i+ 4i, 5i- 6i. 7i/ 8i: 9i; !i= #i? $i@ %i[ 'i] (i^ )i_ *i` +i{ ,i| -i} .i~ /iab:ibb;icb=idb?i@i[i]i^i_i`i{i|i}i~iajbjcjdjejfjgjhjijSejjkjljTemjUenjVeojWepjAbqjXerjYesjZetjujvjBbwjCbxjDbyjEbzjFbAjDcBjEcCjFcDjGcEjHcFjIcGjJcHjKcIjJj,hKj-hLj.hMjygNjzgOjAgPjQjRjSjTjUjVjWjXjYjZj0j1j2j3j4j5j6j7j8j9j!j#j$j%j'j(j)j*j+j,j-j.j/j:j;j=j?j@j[j]j^j_j`j{j|j}j~jakbkckdkekfkgkhkikjkkklkmknkokpkqkrksktkukvkwkxkykzkAkBkCkDkEkFkGkHkIkJkKkLkMkNkOkPkQkRkSkTkUkVkWkXkYkZk0k1k2k3k4k5k6k7k8k9k!k#k$k%k'k(k)k*k+k,k-k.k/k:k;k=k?k@k[k]k^k_k`k{k|k}k~kalblcldlelflglhliljlklllmlnlolplqlrlsltlulvlwlxlylzlAlXbBlClDl0eElFlGlHlIlJlKlLlMlNlOlPlQlRlSlTl1eUl2eVlWlafXlYlZl*c0l+c1l2l3lh 4li 5l6l7lGb8lHb9lIb!l#l3e$l{g%lJb'lKb(lLb)lMb*lNb+l,l-l.l/ld :l;l=l?l,c-c.c1d2d3de j f g 6 eb7 8 fbgbhb4d5d6d7d8d9d@lib[ljb]lOb^lkb_llb`lmb{lnb|lob}lpb~lambfbmcm~gah/hbhch:hdheh;hfh=hdmBgemcf?hfm5egmdfef@hffhmgf[himCgjmgh]hhhih^hhfkmiflmmmnmompmqm6erm9 ! sg!d/cYb:c+bLc;c=ctgQbRbSbTbUbVbsmtmumvmwmxmymzmAmBmCmDmEmFmGmHmImJmKmLmMmnhNmOmPmohQmphRmSmTmUmVmWmXmYmZm0m1m2m3m4m5m6m7m8mqhrhsh9m!mthuhvh#mwhxhyh$m%m'm(m)m*m+m,m-m.m/m:m;m=m?m@m[m]m^mqb,b-b_m`m{m|m}m~manbncndnenfngnhninjnknlnmnnnonpnqnrnsntnunvnwnxnynznAnBnCnDnEnFnGnHnIgqcrcJgKgLgMgNgOgPgQgzhAhke5crbtcsb6cuctbZb)e*e+ek 9bl m n o p q r s t !bu v w x 3by z 4b#ble7cubvcvb8cwcwb0b,e-e.eA $bB C D E F G H I J %bK L M N 5bO P 6b'bme9cxbxcyb!cyczb1b/e:e;eQ (bR S T U V W X Y Z )b0 1 2 3 7b4 5 8b*bRgInJnKn9eLnMnNnOnPnQnRnSnTnUnVnWnXnYnZn0n1n2n3n4n5n6n7n8n9n!n#n$n%n'n(n)n*n+n,n-n.n/n:n;n=n?n@n[n]n^n_n`n{n|n}n~naobocodoeofogohoiojokolomoSgTg$enooopoqorosotouovowoxoyo#dzoAoBoCoDoEoFoGoWbHoIoJoKoLoMoNoOoPoQoRoSoToUoVoWoXoYoZo0o1o2o3o4o5o6o7o8o9o!o#o$o%o'o(o)o*o+o,o-o.o/o:o;o=o?o@o[o]o^o_o`o{o|o}o~oapbpcpdpepfpgphpipjpkplpmpnpopppqprpsptpupvpwpxpypzpApBpCpDpEpFpGpHpIpJpKpLpMpNpOpPpQpRpSpTpUpVpWpXpYpZp0p1p2p3p4p5p6p7p8p9p!p#p$p%p'p(p)p*p+p,p-p.p/p:p;p=p?p@p[p]p^p_p`p{p|p}p~paqbqcqdqeqfqgqhqiqjqkqlqmqnqoqpqqqrqsqtquqvqwqxqyqzqAqBqCqDqEqFqGqHqIqJqKqLqMqNqOqPqQqRqSqTqUqVqWqXqYqZq0q1q2q3q4q5q6q7q8q9q!q#q$q%q'q(q)q*q+q,q-q.q/q:q;q=q?q@q[q]q^q_q`q{q|q}q~qarbrcrdrerfrgrhrirjrkrlrmrnrorprqrrrsrtrurvrwrxryrzrArBrCrDrErFrGrHrIrJrKrLrMrNrOrPrQrRrSrTrUrVrWrXrYrZr0r1r2r3r4r5r6r7r8r9r!r#r$r%r'r(r)r*r+r,r-r.r/r:r;r=r?r@r[r]r^r_r`r{r|r}r~rasbscsdsesfsgshsisjskslsmsnsospsqsrssstsusvswsxsyszsAsBsCsDsEsFsGsHsIsJsKsLsMsNsOsPsQsRsSsTsUsVsWsXsYsZs0s1s2s3s4s5s6s7s8s9s!s#s$s%s's(s)s*s+s,s-s.s/s:s;s=s?s@s[s]s^s_s`s{s|s}s~satbtctdtetftgthtitjtktltmtntotptqtrtstttutvtwtxtytztAtBtCtDtEtFtGtHtItJtKtLtMtNtOtPtQtRtStTtUtVtWtXtYtZt0t1t2t3t4t5t6t7t8t9t!t#t$t%t't(t)t*t+t,t-t.t/t:t;t=t?t@t[t]t^t_t`t{t|t}t~taubucudueufuguhuiujukulumunuoupuqurusutuuuvuwuxuyuzuAuBuCuDuEuFuGuHuIuJuKuLuMuNuOuPuQuRuSuTuUuVuWuXuYuZu0u1u2u3u4u5u6u7u8u9u!u#u$u%u'u(u)u*u+u,u-u.u/u:u;u=u?u@u[u]u^u_u`u{u|u}u~uavbvcvdvevfvgvhvivjvkvlvmvnvovpvqvrvsvtvuvvvwvxvyvzvAvBvCvDvEvFvGvHvIvJvKvLvMvNvOvPvQvRvSvTvUvVvWvXvYvZv0v1v2v3v4v5v6v7v8v9v!v#v$v%v'v(v)v*v+v,v-v.v/v:v;v=v?v@v[v]v^v_v`v{v|v}v~vawbwcwdwewfwgwhwiwjwkwlwmwnwowpwqwrwswtwuwvwwwxwywzwAwBwCwDwEwFwGwHwIwJwKwLwMwNwOwPwQwRwSwTwUwVwWwXwYwZw0w1w2w3w4w5w6w7w8w9w!w#w$w%w'w(w)w*w+w,w-w.w/w:w;w=w?w@w[w]w^w_w`w{w(e|w}w~waxbxcxdxexfxgxhxixjxkxlxmxnxoxpxqxrxsxtxuxvxwxxxyxzxAxBxCxDxExFxGxHxIxJxKxLxMxNxOxPxQxRxSxTxUxVxWxXxYxZx0x1x2x3x4x5x6x7x8x9x!x#x$x%x'x(x)x*x+x,x-x.x/x:x;x=x?x@x[x]x^x_x`x{x|x}x~xaybycydyeyfygyhyiyjykylymynyoypyqyrysytyuyvywyxyyyzyAyByCyDyEyFyGyHyIyJyKyLyMyNyOyPyQyRySyTyUyVyWyXyYyZy0y1y2y3y4y5y6y7y8y9y!y#y$y%y'y(y)y*y+y,y-y.y/y:y;y=y?y@y[y]y^y_y`y{y|y}y~yazbzczdzezfzgzhzizjzkzlzmznzozpzqzrzsztzuzvzwzxzyzzzAzBzCzDzEzFzGzHzIzJzKzLzMzNzOzPzQzRzSzTzUzVzWzXzYzZz0z1z2z3z4z5z6z7z8z9z!z#z$z%z'z(z)z*z+z,z-z.z/z:z;z=z?z@z[z]z^z_z`z{z|z}z~zaAbAcAdAeAfAgAhAiAjAkAlAmAnAoApAqArAsAtAuAvAwAxAyAzAAABACADAEAFAGAHAIAJAKALAMANAOAPAQARASATAUAVAWAXAYAZA0A1A2A3A4A5A6A7A8A9A!A#A$A%A'A(A)A*A+A,A-A.A/A:A;A=A?A@A[A]A^A_A`A{A|A}A~AaBbBcBdBeBfBgBhBiBjBkBlBmBnBoBpBqBrBsBtBuBvBwBxByBzBABBBCBDBEBFBGBHBIBJBKBLBMBNBOBPBQBRBSBTBUBVBWBXBYBZB0B1B2B3B4B5B6B7B8B9B!B#B$B%B'B(B)B*B+B,B-B.B/B
205 elif isinstance(error, runtime.cudaError_t): 2a QbRbSbTbUbVbDgEg7e8eFgGgHg,b-bb .b/b:b;b=b?b@b[b]b^b_b`b{b|b}b~bacbcccdcecfcgchcicjckclcmcncocpcqcrcrbsbtbZbk l m n o p q r s t u v w x y z ubvbwb0bA B C D E F G H I J K L M N O P xbybzb1bQ R S T U V W X Y Z 0 1 2 3 4 5 ?cMc!e@c#e[cscRcScTcGhUcHhVcIhWcvgwgXcYc'dZc(d)d*d^c_c`c+d,d-d.d3c4c{c/d|c:d;d=d?d}c@d~cadbdcd[d]d^d_d`d{d|d}d~dddedfdc 0cgd1c2cUgVgWgXgYgZgJhKhLhMhNhOhPhQhRhShThUhVhWhXhYhZh0h1h2h3h4h5h6h7h8h9haebecedeeefegeheieje
206 return _check_runtime_error(error) 2QbRbSbTbUbVbDgEg7e8eFgGgHg,b-bqcrcrbsbtbZbk l m n o p q r s t u v w x y z ubvbwb0bA B C D E F G H I J K L M N O P xbybzb1bQ R S T U V W X Y Z 0 1 2 3 4 5 ?cMc!e@c#e[c
207 elif isinstance(error, nvrtc.nvrtcResult): 2a b .b/b:b;b=b?b@b[b]b^b_b`b{b|b}b~bacbcccdcecfcgchcicjckclcmcncocpcscRcScTcGhUcHhVcIhWcvgwgXcYc'dZc(d)d*d^c_c`c+d,d-d.d3c4c{c/d|c:d;d=d?d}c@d~cadbdcd[d]d^d_d`d{d|d}d~dddedfdc 0cgd1c2cUgVgWgXgYgZgJhKhLhMhNhOhPhQhRhShThUhVhWhXhYhZh0h1h2h3h4h5h6h7h8h9haebecedeeefegeheieje
208 return _check_nvrtc_error(error, handle=handle) 2a b .b/b:b;b=b?b@b[b]b^b_b`b{b|b}b~bacbcccdcecfcgchcicjckclcmcncocpcscRcScTcGhUcHhVcIhWcvgwgXcYc'dZc(d)d*d^c_c`c+d,d-d.d3c4c{c/d|c:d;d=d?d}c@d~cadbdcd[d]d^d_d`d{d|d}d~dddedfdc 0cgd1c2cUgVgWgXgYgZgJhKhLhMhNhOhPhQhRhShThUhVhWhXhYhZh0h1h2h3h4h5h6h7h8h9haebecedeeefegeheieje
209 else:
210 raise RuntimeError(f"Unknown error type: {error}")
213def handle_return(result: tuple[Any, ...], handle: object = None) -> Any:
214 _check_error(result[0], handle=handle) 2a ZdTi0dUikfVi# Wi$ XilfYi% Zi' 0i( 1i) 2i* 3i+ 4i, 5i- 6i. 7i/ 8i: 9i; !i= #i? $i@ %i[ 'i] (i^ )i_ *i` +i{ ,i| -i} .i~ /iab:ibb;icb=idb?i@i[i]i^i_i`i{i|i}i~iajbjcjdjejfjgjhjijSejjkjljTemjUenjVeojWepjAbqjXerjYesjZetjujvjBbwjCbxjDbyjEbzjFbAjDcBjEcCjFcDjGcEjHcFjIcGjJcHjKcIjJj,hKj-hLj.hMjygNjzgOjAgPjQjRjSjTjUjVjWjXjYjZj0j1j2j3j4j5j6j7j8j9j!j#j$j%j'j(j)j*j+j,j-j.j/j:j;j=j?j@j[j]j^j_j`j{j|j}j~jakbkckdkekfkgkhkikjkkklkmknkokpkqkrksktkukvkwkxkykzkAkBkCkDkEkFkGkHkIkJkKkLkMkNkOkPkQkRkSkTkUkVkWkXkYkZk0k1k2k3k4k5k6k7k8k9k!k#k$k%k'k(k)k*k+k,k-k.k/k:k;k=k?k@k[k]k^k_k`k{k|k}k~kalblcldlelflglhliljlklllmlnlolplqlrlsltlulvlwlxlylzlAlXbBlClDl0eElFlGlHlIlJlKlLlMlNlOlPlQlRlSlTl1eUl2eVlWlafXlYlZl*c0l+c1l2l3lh 4li 5l6l7lGb8lHb9lIb!l#l3e$l{g%lJb'lKb(lLb)lMb*lNb+l,l-l.l/ld :l;l=l?l,c-c.c1d2d3de j f g 6 eb7 8 fbgbhb4d5d6d7d8d9d@lib[ljb]lOb^lkb_llb`lmb{lnb|lob}lpb~lambfbmcm~gah/hbhch:hdheh;hfh=hdmBgemcf?hfm5egmdfef@hffhmgf[himCgjmgh]hhhih^hhfkmiflmmmnmompmqm6erm9 ! sg!d/cYb:c+bLc;c=ctgQbRbSbTbUbVbDgsmtmumvmwmxmymzmEg7eAm8eBmCmDmEmFmGmHmImFgJmKmLmMmnhNmOmPmohGgHgQmphRmSmTmUmVmWmXmYmZm0m1m2m3m4m5m6m7m8mqhrhsh9m!mthuhvh#mwhxhyh$m%m'm(m)m*m+m,m-m.m/m:m;m=m?m@m[m]m^mqb,b-b_m`m{m|mb .b}m/b~m:ban;bbn=bcn?bdn@ben[bfn]bgn^bhn_bin`bjn{bkn|bln}bmn~bnnaconbcpnccqndcrnecsnfctngcunhcvnicwnjcxnkcynlcznmcAnncBnocCnDnpcEnFnGnHnIgqcrcJgKgLgMgNgOgPgQgzhAhke5crbtcsb6cuctbZb)e*e+ek 9bl m n o p q r s t !bu v w x 3by z 4b#ble7cubvcvb8cwcwb0b,e-e.eA $bB C D E F G H I J %bK L M N 5bO P 6b'bme9cxbxcyb!cyczb1b/e:e;eQ (bR S T U V W X Y Z )b0 1 2 3 7b4 5 8b*bRg?cInJnKn9eLnMnNnOnPnQnRnMcSn!eTn@cUn#eVnWnXnYnZn0n1n2n3n4n5n6n7n8n9n!n#n$n%n'n(n)n*n+n,n-n.n/n:n;n=n?n@n[n]n^n_n`n{n|n}n~naobocodoeofogohoiojokolomoSgTg$eno[coopoqorosotouovoscwoxoyo#dzoAoBoCoDoEoFoGoWbHoIoJoKoLoMoNoOoPoQoRoSoToUoRcVoScWoTcXoYoZo0oGhUc1oHhVc2oIhWc3o4o5o6o7o8o9o!o#o$o%o'o(o)o*o+o,o-o.o/o:o;o=o?o@o[o]o^o_o`o{o|o}o~oapbpcpdpepfpgphpipjpkplpmpnpopppqprpsptpupvpwpxpypzpApBpCpDpEpFpGpHpIpJpKpLpMpNpOpPpQpRpSpTpUpVpWpXpYpZp0p1p2p3p4p5p6p7p8p9p!p#p$p%p'p(p)p*p+p,p-p.p/p:p;p=p?p@p[p]p^p_p`p{p|p}p~paqbqcqdqeqfqgqhqiqjqkqlqmqnqoqpqqqrqsqtquqvqwqxqyqzqAqBqCqDqEqFqGqHqIqJqKqLqMqNqOqPqQqRqSqTqUqVqWqXqYqZq0q1q2q3q4q5q6q7q8q9q!q#q$q%q'q(q)q*q+q,q-q.q/q:q;q=q?q@q[q]q^q_q`q{q|q}q~qarbrcrdrerfrgrhrirjrkrlrmrnrorprqrrrsrtrurvrwrxryrzrArBrCrDrErFrGrHrIrJrKrLrMrNrOrPrQrRrSrTrUrVrWrXrYrZr0r1r2r3r4r5r6r7r8r9r!r#r$r%r'r(r)r*r+r,r-r.r/r:r;r=r?r@r[r]r^r_r`r{r|r}r~rasbscsdsesfsgshsisjskslsmsnsospsqsrssstsusvswsxsyszsAsBsCsDsEsFsGsHsIsJsKsLsMsNsOsPsQsRsSsTsUsVsWsXsYsZs0s1s2s3s4s5s6s7s8s9s!s#s$s%s's(s)s*s+s,s-s.s/s:s;s=s?s@s[s]s^s_s`s{s|s}s~satbtctdtetftgthtitjtktltmtntotptqtrtstttutvtwtxtytztAtBtCtDtEtFtGtHtItJtKtLtMtNtOtPtQtRtStTtUtVtWtXtYtZt0t1t2t3t4t5t6t7t8t9t!t#t$t%t't(t)t*t+t,t-t.t/t:t;t=t?t@t[t]t^t_t`t{t|t}t~taubucudueufuguhuiujukulumunuoupuqurusutuuuvuwuxuyuzuAuBuCuDuEuFuGuHuIuJuKuLuMuNuOuPuQuRuSuTuUuVuWuXuYuZu0u1u2u3u4u5u6u7u8u9u!u#u$u%u'u(u)u*u+u,u-u.u/u:u;u=u?u@u[u]u^u_u`u{u|u}u~uavbvcvdvevfvgvhvivjvkvlvmvnvovpvqvrvsvtvuvvvwvxvyvzvAvBvCvDvEvFvGvHvIvJvKvLvMvNvOvPvQvRvSvTvUvVvWvXvYvZv0v1v2v3v4v5v6v7v8v9v!v#v$v%v'v(v)v*v+v,v-v.v/v:v;v=v?v@v[v]v^v_v`v{v|v}v~vawbwcwdwewfwgwhwiwjwkwlwmwnwowpwqwrwswtwuwvwwwxwywzwAwBwCwDwEwFwGwHwIwJwKwLwMwNwOwPwQwRwSwTwUwVwWwXwYwZw0w1w2w3w4w5w6w7w8w9w!w#w$w%w'w(w)w*w+w,w-w.w/w:w;w=w?w@w[w]w^w_w`w{w(e|w}w~waxbxcxdxexfxgxhxixjxkxlxmxnxoxpxqxrxsxtxvguxvxwxwgxxyxzxAxBxCxDxExFxGxHxIxJxKxLxMxNxOxPxQxRxSxTxUxVxWxXxYxZx0x1x2x3x4x5x6x7x8x9x!x#x$x%x'x(x)x*x+x,x-x.x/x:x;x=x?x@x[x]x^x_x`x{x|x}x~xaybycydyeyfygyhyiyjykylymynyoypyqyrysytyuyvywyxyyyzyAyByCyDyEyFyGyHyIyJyKyLyMyNyOyPyQyRySyTyUyVyWyXyYyZy0y1y2y3y4y5y6y7y8y9y!y#y$y%y'y(y)y*y+y,y-y.y/y:y;y=y?y@y[y]y^y_y`y{y|y}y~yazbzczdzezfzgzhzizjzkzlzmznzozpzqzrzsztzuzXcvzYcwz'dxzZcyz(dzz)dAz*dBz^cCz_cDz`cEz+dFz,dGz-dHz.dIz3cJz4cKz{cLz/dMz|cNz:dOz;dPz=dQz?dRz}cSz@dTz~cUzadVzbdWzcdXz[dYz]dZz^d0z_d1z`d2z{d3z|d4z}d5z~d6zdd7zed8zfd9zc !z0c#z$z%z'z(z)z*z+z,z-z.z/z:z;z=zgd?z1c@z[z]z^z_z`z{z|z}z~zaA2cbAUgVgWgXgYgZgJhKhLhMhNhOhPhQhRhShThUhVhWhcAXhYhZh0h1h2h3h4h5h6hdAeAfAgA7h8h9haebecedeeehAfegeheiejeiAjAkAlAmAnAoApAqArAsAtAuAvAwAxAyAzAAABACADAEAFAGAHAIAJAKALAMANAOAPAQARASATAUAVAWAXAYAZA0A1A2A3A4A5A6A7A8A9A!A#A$A%A'A(A)A*A+A,A-A.A/A:A;A=A?A@A[A]A^A_A`A{A|A}A~AaBbBcBdBeBfBgBhBiBjBkBlBmBnBoBpBqBrBsBtBuBvBwBxByBzBABBBCBDBEBFBGBHBIBJBKBLBMBNBOBPBQBRBSBTBUBVBWBXBYBZB0B1B2B3B4B5B6B7B8B9B!B#B$B%B'B(B)B*B+B,B-B.B/B
215 cdef int out_len = len(result) 2a ZdTi0dUikfVi# Wi$ XilfYi% Zi' 0i( 1i) 2i* 3i+ 4i, 5i- 6i. 7i/ 8i: 9i; !i= #i? $i@ %i[ 'i] (i^ )i_ *i` +i{ ,i| -i} .i~ /iab:ibb;icb=idb?i@i[i]i^i_i`i{i|i}i~iajbjcjdjejfjgjhjijSejjkjljTemjUenjVeojWepjAbqjXerjYesjZetjujvjBbwjCbxjDbyjEbzjFbAjDcBjEcCjFcDjGcEjHcFjIcGjJcHjKcIjJj,hKj-hLj.hMjygNjzgOjAgPjQjRjSjTjUjVjWjXjYjZj0j1j2j3j4j5j6j7j8j9j!j#j$j%j'j(j)j*j+j,j-j.j/j:j;j=j?j@j[j]j^j_j`j{j|j}j~jakbkckdkekfkgkhkikjkkklkmknkokpkqkrksktkukvkwkxkykzkAkBkCkDkEkFkGkHkIkJkKkLkMkNkOkPkQkRkSkTkUkVkWkXkYkZk0k1k2k3k4k5k6k7k8k9k!k#k$k%k'k(k)k*k+k,k-k.k/k:k;k=k?k@k[k]k^k_k`k{k|k}k~kalblcldlelflglhliljlklllmlnlolplqlrlsltlulvlwlxlylzlAlXbBlClDl0eElFlGlHlIlJlKlLlMlNlOlPlQlRlSlTl1eUl2eVlWlafXlYlZl*c0l+c1l2l3lh 4li 5l6l7lGb8lHb9lIb!l#l3e$l{g%lJb'lKb(lLb)lMb*lNb+l,l-l.l/ld :l;l=l?l,c-c.c1d2d3de j f g 6 eb7 8 fbgbhb4d5d6d7d8d9d@lib[ljb]lOb^lkb_llb`lmb{lnb|lob}lpb~lambfbmcm~gah/hbhch:hdheh;hfh=hdmBgemcf?hfm5egmdfef@hffhmgf[himCgjmgh]hhhih^hhfkmiflmmmnmompmqm6erm9 ! sg!d/cYb:c+bLc;c=ctgQbRbSbTbUbVbDgsmtmumvmwmxmymzmEg7eAm8eBmCmDmEmFmGmHmImFgJmKmLmMmnhNmOmPmohGgHgQmphRmSmTmUmVmWmXmYmZm0m1m2m3m4m5m6m7m8mqhrhsh9m!mthuhvh#mwhxhyh$m%m'm(m)m*m+m,m-m.m/m:m;m=m?m@m[m]m^mqb,b-b_m`m{m|mb .b}m/b~m:ban;bbn=bcn?bdn@ben[bfn]bgn^bhn_bin`bjn{bkn|bln}bmn~bnnaconbcpnccqndcrnecsnfctngcunhcvnicwnjcxnkcynlcznmcAnncBnocCnDnpcEnFnGnHnIgqcrcJgKgLgMgNgOgPgQgzhAhke5crbtcsb6cuctbZb)e*e+ek 9bl m n o p q r s t !bu v w x 3by z 4b#ble7cubvcvb8cwcwb0b,e-e.eA $bB C D E F G H I J %bK L M N 5bO P 6b'bme9cxbxcyb!cyczb1b/e:e;eQ (bR S T U V W X Y Z )b0 1 2 3 7b4 5 8b*bRg?cInJnKn9eLnMnNnOnPnQnRnMcSn!eTn@cUn#eVnWnXnYnZn0n1n2n3n4n5n6n7n8n9n!n#n$n%n'n(n)n*n+n,n-n.n/n:n;n=n?n@n[n]n^n_n`n{n|n}n~naobocodoeofogohoiojokolomoSgTg$eno[coopoqorosotouovoscwoxoyo#dzoAoBoCoDoEoFoGoWbHoIoJoKoLoMoNoOoPoQoRoSoToUoRcVoScWoTcXoYoZo0oGhUc1oHhVc2oIhWc3o4o5o6o7o8o9o!o#o$o%o'o(o)o*o+o,o-o.o/o:o;o=o?o@o[o]o^o_o`o{o|o}o~oapbpcpdpepfpgphpipjpkplpmpnpopppqprpsptpupvpwpxpypzpApBpCpDpEpFpGpHpIpJpKpLpMpNpOpPpQpRpSpTpUpVpWpXpYpZp0p1p2p3p4p5p6p7p8p9p!p#p$p%p'p(p)p*p+p,p-p.p/p:p;p=p?p@p[p]p^p_p`p{p|p}p~paqbqcqdqeqfqgqhqiqjqkqlqmqnqoqpqqqrqsqtquqvqwqxqyqzqAqBqCqDqEqFqGqHqIqJqKqLqMqNqOqPqQqRqSqTqUqVqWqXqYqZq0q1q2q3q4q5q6q7q8q9q!q#q$q%q'q(q)q*q+q,q-q.q/q:q;q=q?q@q[q]q^q_q`q{q|q}q~qarbrcrdrerfrgrhrirjrkrlrmrnrorprqrrrsrtrurvrwrxryrzrArBrCrDrErFrGrHrIrJrKrLrMrNrOrPrQrRrSrTrUrVrWrXrYrZr0r1r2r3r4r5r6r7r8r9r!r#r$r%r'r(r)r*r+r,r-r.r/r:r;r=r?r@r[r]r^r_r`r{r|r}r~rasbscsdsesfsgshsisjskslsmsnsospsqsrssstsusvswsxsyszsAsBsCsDsEsFsGsHsIsJsKsLsMsNsOsPsQsRsSsTsUsVsWsXsYsZs0s1s2s3s4s5s6s7s8s9s!s#s$s%s's(s)s*s+s,s-s.s/s:s;s=s?s@s[s]s^s_s`s{s|s}s~satbtctdtetftgthtitjtktltmtntotptqtrtstttutvtwtxtytztAtBtCtDtEtFtGtHtItJtKtLtMtNtOtPtQtRtStTtUtVtWtXtYtZt0t1t2t3t4t5t6t7t8t9t!t#t$t%t't(t)t*t+t,t-t.t/t:t;t=t?t@t[t]t^t_t`t{t|t}t~taubucudueufuguhuiujukulumunuoupuqurusutuuuvuwuxuyuzuAuBuCuDuEuFuGuHuIuJuKuLuMuNuOuPuQuRuSuTuUuVuWuXuYuZu0u1u2u3u4u5u6u7u8u9u!u#u$u%u'u(u)u*u+u,u-u.u/u:u;u=u?u@u[u]u^u_u`u{u|u}u~uavbvcvdvevfvgvhvivjvkvlvmvnvovpvqvrvsvtvuvvvwvxvyvzvAvBvCvDvEvFvGvHvIvJvKvLvMvNvOvPvQvRvSvTvUvVvWvXvYvZv0v1v2v3v4v5v6v7v8v9v!v#v$v%v'v(v)v*v+v,v-v.v/v:v;v=v?v@v[v]v^v_v`v{v|v}v~vawbwcwdwewfwgwhwiwjwkwlwmwnwowpwqwrwswtwuwvwwwxwywzwAwBwCwDwEwFwGwHwIwJwKwLwMwNwOwPwQwRwSwTwUwVwWwXwYwZw0w1w2w3w4w5w6w7w8w9w!w#w$w%w'w(w)w*w+w,w-w.w/w:w;w=w?w@w[w]w^w_w`w{w(e|w}w~waxbxcxdxexfxgxhxixjxkxlxmxnxoxpxqxrxsxtxvguxvxwxwgxxyxzxAxBxCxDxExFxGxHxIxJxKxLxMxNxOxPxQxRxSxTxUxVxWxXxYxZx0x1x2x3x4x5x6x7x8x9x!x#x$x%x'x(x)x*x+x,x-x.x/x:x;x=x?x@x[x]x^x_x`x{x|x}x~xaybycydyeyfygyhyiyjykylymynyoypyqyrysytyuyvywyxyyyzyAyByCyDyEyFyGyHyIyJyKyLyMyNyOyPyQyRySyTyUyVyWyXyYyZy0y1y2y3y4y5y6y7y8y9y!y#y$y%y'y(y)y*y+y,y-y.y/y:y;y=y?y@y[y]y^y_y`y{y|y}y~yazbzczdzezfzgzhzizjzkzlzmznzozpzqzrzsztzuzXcvzYcwz'dxzZcyz(dzz)dAz*dBz^cCz_cDz`cEz+dFz,dGz-dHz.dIz3cJz4cKz{cLz/dMz|cNz:dOz;dPz=dQz?dRz}cSz@dTz~cUzadVzbdWzcdXz[dYz]dZz^d0z_d1z`d2z{d3z|d4z}d5z~d6zdd7zed8zfd9zc !z0c#z$z%z'z(z)z*z+z,z-z.z/z:z;z=zgd?z1c@z[z]z^z_z`z{z|z}z~zaA2cbAUgVgWgXgYgZgJhKhLhMhNhOhPhQhRhShThUhVhWhcAXhYhZh0h1h2h3h4h5h6hdAeAfAgA7h8h9haebecedeeehAfegeheiejeiAjAkAlAmAnAoApAqArAsAtAuAvAwAxAyAzAAABACADAEAFAGAHAIAJAKALAMANAOAPAQARASATAUAVAWAXAYAZA0A1A2A3A4A5A6A7A8A9A!A#A$A%A'A(A)A*A+A,A-A.A/A:A;A=A?A@A[A]A^A_A`A{A|A}A~AaBbBcBdBeBfBgBhBiBjBkBlBmBnBoBpBqBrBsBtBuBvBwBxByBzBABBBCBDBEBFBGBHBIBJBKBLBMBNBOBPBQBRBSBTBUBVBWBXBYBZB0B1B2B3B4B5B6B7B8B9B!B#B$B%B'B(B)B*B+B,B-B.B/B
216 if out_len == 1: 2a ZdTi0dUikfVi# Wi$ XilfYi% Zi' 0i( 1i) 2i* 3i+ 4i, 5i- 6i. 7i/ 8i: 9i; !i= #i? $i@ %i[ 'i] (i^ )i_ *i` +i{ ,i| -i} .i~ /iab:ibb;icb=idb?i@i[i]i^i_i`i{i|i}i~iajbjcjdjejfjgjhjijSejjkjljTemjUenjVeojWepjAbqjXerjYesjZetjujvjBbwjCbxjDbyjEbzjFbAjDcBjEcCjFcDjGcEjHcFjIcGjJcHjKcIjJj,hKj-hLj.hMjygNjzgOjAgPjQjRjSjTjUjVjWjXjYjZj0j1j2j3j4j5j6j7j8j9j!j#j$j%j'j(j)j*j+j,j-j.j/j:j;j=j?j@j[j]j^j_j`j{j|j}j~jakbkckdkekfkgkhkikjkkklkmknkokpkqkrksktkukvkwkxkykzkAkBkCkDkEkFkGkHkIkJkKkLkMkNkOkPkQkRkSkTkUkVkWkXkYkZk0k1k2k3k4k5k6k7k8k9k!k#k$k%k'k(k)k*k+k,k-k.k/k:k;k=k?k@k[k]k^k_k`k{k|k}k~kalblcldlelflglhliljlklllmlnlolplqlrlsltlulvlwlxlylzlAlXbBlClDl0eElFlGlHlIlJlKlLlMlNlOlPlQlRlSlTl1eUl2eVlWlafXlYlZl*c0l+c1l2l3lh 4li 5l6l7lGb8lHb9lIb!l#l3e$l{g%lJb'lKb(lLb)lMb*lNb+l,l-l.l/ld :l;l=l?l,c-c.c1d2d3de j f g 6 eb7 8 fbgbhb4d5d6d7d8d9d@lib[ljb]lOb^lkb_llb`lmb{lnb|lob}lpb~lambfbmcm~gah/hbhch:hdheh;hfh=hdmBgemcf?hfm5egmdfef@hffhmgf[himCgjmgh]hhhih^hhfkmiflmmmnmompmqm6erm9 ! sg!d/cYb:c+bLc;c=ctgQbRbSbTbUbVbDgsmtmumvmwmxmymzmEg7eAm8eBmCmDmEmFmGmHmImFgJmKmLmMmnhNmOmPmohGgHgQmphRmSmTmUmVmWmXmYmZm0m1m2m3m4m5m6m7m8mqhrhsh9m!mthuhvh#mwhxhyh$m%m'm(m)m*m+m,m-m.m/m:m;m=m?m@m[m]m^mqb,b-b_m`m{m|mb .b}m/b~m:ban;bbn=bcn?bdn@ben[bfn]bgn^bhn_bin`bjn{bkn|bln}bmn~bnnaconbcpnccqndcrnecsnfctngcunhcvnicwnjcxnkcynlcznmcAnncBnocCnDnpcEnFnGnHnIgqcrcJgKgLgMgNgOgPgQgzhAhke5crbtcsb6cuctbZb)e*e+ek 9bl m n o p q r s t !bu v w x 3by z 4b#ble7cubvcvb8cwcwb0b,e-e.eA $bB C D E F G H I J %bK L M N 5bO P 6b'bme9cxbxcyb!cyczb1b/e:e;eQ (bR S T U V W X Y Z )b0 1 2 3 7b4 5 8b*bRg?cInJnKn9eLnMnNnOnPnQnRnMcSn!eTn@cUn#eVnWnXnYnZn0n1n2n3n4n5n6n7n8n9n!n#n$n%n'n(n)n*n+n,n-n.n/n:n;n=n?n@n[n]n^n_n`n{n|n}n~naobocodoeofogohoiojokolomoSgTg$eno[coopoqorosotouovoscwoxoyo#dzoAoBoCoDoEoFoGoWbHoIoJoKoLoMoNoOoPoQoRoSoToUoRcVoScWoTcXoYoZo0oGhUc1oHhVc2oIhWc3o4o5o6o7o8o9o!o#o$o%o'o(o)o*o+o,o-o.o/o:o;o=o?o@o[o]o^o_o`o{o|o}o~oapbpcpdpepfpgphpipjpkplpmpnpopppqprpsptpupvpwpxpypzpApBpCpDpEpFpGpHpIpJpKpLpMpNpOpPpQpRpSpTpUpVpWpXpYpZp0p1p2p3p4p5p6p7p8p9p!p#p$p%p'p(p)p*p+p,p-p.p/p:p;p=p?p@p[p]p^p_p`p{p|p}p~paqbqcqdqeqfqgqhqiqjqkqlqmqnqoqpqqqrqsqtquqvqwqxqyqzqAqBqCqDqEqFqGqHqIqJqKqLqMqNqOqPqQqRqSqTqUqVqWqXqYqZq0q1q2q3q4q5q6q7q8q9q!q#q$q%q'q(q)q*q+q,q-q.q/q:q;q=q?q@q[q]q^q_q`q{q|q}q~qarbrcrdrerfrgrhrirjrkrlrmrnrorprqrrrsrtrurvrwrxryrzrArBrCrDrErFrGrHrIrJrKrLrMrNrOrPrQrRrSrTrUrVrWrXrYrZr0r1r2r3r4r5r6r7r8r9r!r#r$r%r'r(r)r*r+r,r-r.r/r:r;r=r?r@r[r]r^r_r`r{r|r}r~rasbscsdsesfsgshsisjskslsmsnsospsqsrssstsusvswsxsyszsAsBsCsDsEsFsGsHsIsJsKsLsMsNsOsPsQsRsSsTsUsVsWsXsYsZs0s1s2s3s4s5s6s7s8s9s!s#s$s%s's(s)s*s+s,s-s.s/s:s;s=s?s@s[s]s^s_s`s{s|s}s~satbtctdtetftgthtitjtktltmtntotptqtrtstttutvtwtxtytztAtBtCtDtEtFtGtHtItJtKtLtMtNtOtPtQtRtStTtUtVtWtXtYtZt0t1t2t3t4t5t6t7t8t9t!t#t$t%t't(t)t*t+t,t-t.t/t:t;t=t?t@t[t]t^t_t`t{t|t}t~taubucudueufuguhuiujukulumunuoupuqurusutuuuvuwuxuyuzuAuBuCuDuEuFuGuHuIuJuKuLuMuNuOuPuQuRuSuTuUuVuWuXuYuZu0u1u2u3u4u5u6u7u8u9u!u#u$u%u'u(u)u*u+u,u-u.u/u:u;u=u?u@u[u]u^u_u`u{u|u}u~uavbvcvdvevfvgvhvivjvkvlvmvnvovpvqvrvsvtvuvvvwvxvyvzvAvBvCvDvEvFvGvHvIvJvKvLvMvNvOvPvQvRvSvTvUvVvWvXvYvZv0v1v2v3v4v5v6v7v8v9v!v#v$v%v'v(v)v*v+v,v-v.v/v:v;v=v?v@v[v]v^v_v`v{v|v}v~vawbwcwdwewfwgwhwiwjwkwlwmwnwowpwqwrwswtwuwvwwwxwywzwAwBwCwDwEwFwGwHwIwJwKwLwMwNwOwPwQwRwSwTwUwVwWwXwYwZw0w1w2w3w4w5w6w7w8w9w!w#w$w%w'w(w)w*w+w,w-w.w/w:w;w=w?w@w[w]w^w_w`w{w(e|w}w~waxbxcxdxexfxgxhxixjxkxlxmxnxoxpxqxrxsxtxvguxvxwxwgxxyxzxAxBxCxDxExFxGxHxIxJxKxLxMxNxOxPxQxRxSxTxUxVxWxXxYxZx0x1x2x3x4x5x6x7x8x9x!x#x$x%x'x(x)x*x+x,x-x.x/x:x;x=x?x@x[x]x^x_x`x{x|x}x~xaybycydyeyfygyhyiyjykylymynyoypyqyrysytyuyvywyxyyyzyAyByCyDyEyFyGyHyIyJyKyLyMyNyOyPyQyRySyTyUyVyWyXyYyZy0y1y2y3y4y5y6y7y8y9y!y#y$y%y'y(y)y*y+y,y-y.y/y:y;y=y?y@y[y]y^y_y`y{y|y}y~yazbzczdzezfzgzhzizjzkzlzmznzozpzqzrzsztzuzXcvzYcwz'dxzZcyz(dzz)dAz*dBz^cCz_cDz`cEz+dFz,dGz-dHz.dIz3cJz4cKz{cLz/dMz|cNz:dOz;dPz=dQz?dRz}cSz@dTz~cUzadVzbdWzcdXz[dYz]dZz^d0z_d1z`d2z{d3z|d4z}d5z~d6zdd7zed8zfd9zc !z0c#z$z%z'z(z)z*z+z,z-z.z/z:z;z=zgd?z1c@z[z]z^z_z`z{z|z}z~zaA2cbAUgVgWgXgYgZgJhKhLhMhNhOhPhQhRhShThUhVhWhcAXhYhZh0h1h2h3h4h5h6hdAeAfAgA7h8h9haebecedeeehAfegeheiejeiAjAkAlAmAnAoApAqArAsAtAuAvAwAxAyAzAAABACADAEAFAGAHAIAJAKALAMANAOAPAQARASATAUAVAWAXAYAZA0A1A2A3A4A5A6A7A8A9A!A#A$A%A'A(A)A*A+A,A-A.A/A:A;A=A?A@A[A]A^A_A`A{A|A}A~AaBbBcBdBeBfBgBhBiBjBkBlBmBnBoBpBqBrBsBtBuBvBwBxByBzBABBBCBDBEBFBGBHBIBJBKBLBMBNBOBPBQBRBSBTBUBVBWBXBYBZB0B1B2B3B4B5B6B7B8B9B!B#B$B%B'B(B)B*B+B,B-B.B/B
217 return 2Zd0dkf# $ lf% ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdbSeTeUeVeWeAbXeYeZeBbCbDbEbFbDcEcFcGcHcIcJcKc,h-h.hygzgAgXb0e1e2eaf*c+ch i GbHbIb3eJbKbLbMbNbd ,c-c.c1d2d3de j f g 6 eb7 8 fbgbhb4d5d6d7d8d9dibjbObkblbmbnbobpb/h:h;h=hBg?h@h[hCg]h^h6esg/cYb:c+b;c=ctgQbRbSbTbUbVb7e8eqb,b-bIgqcrcJgKgLgMgNgOgPgQgzhAhke5crbtcsb6cuctbZb)e*e+ek 9bl m n o p q r s t !bu v w x 3by z 4b#ble7cubvcvb8cwcwb0b,e-e.eA $bB C D E F G H I J %bK L M N 5bO P 6b'bme9cxbxcyb!cyczb1b/e:e;eQ (bR S T U V W X Y Z )b0 1 2 3 7b4 5 8b*bRg?c9eMc!e@c#eSgTg$e[c(ekz
218 elif out_len == 2:
219 return result[1] 2a ZdTi0dUikfVi# Wi$ XilfYi% Zi' 0i( 1i) 2i* 3i+ 4i, 5i- 6i. 7i/ 8i: 9i; !i= #i? $i@ %i[ 'i] (i^ )i_ *i` +i{ ,i| -i} .i~ /iab:ibb;icb=idb?i@i[i]i^i_i`i{i|i}i~iajbjcjdjejfjgjhjijSejjkjljTemjUenjVeojWepjAbqjXerjYesjZetjujvjBbwjCbxjDbyjEbzjFbAjDcBjEcCjFcDjGcEjHcFjIcGjJcHjKcIjJj,hKj-hLj.hMjygNjzgOjAgPjQjRjSjTjUjVjWjXjYjZj0j1j2j3j4j5j6j7j8j9j!j#j$j%j'j(j)j*j+j,j-j.j/j:j;j=j?j@j[j]j^j_j`j{j|j}j~jakbkckdkekfkgkhkikjkkklkmknkokpkqkrksktkukvkwkxkykzkAkBkCkDkEkFkGkHkIkJkKkLkMkNkOkPkQkRkSkTkUkVkWkXkYkZk0k1k2k3k4k5k6k7k8k9k!k#k$k%k'k(k)k*k+k,k-k.k/k:k;k=k?k@k[k]k^k_k`k{k|k}k~kalblcldlelflglhliljlklllmlnlolplqlrlsltlulvlwlxlylzlAlXbBlClDl0eElFlGlHlIlJlKlLlMlNlOlPlQlRlSlTl1eUl2eVlWlXlYlZl*c0l+c1l2l3lh 4li 5l6l7lGb8lHb9lIb!l#l3e$l{g%lJb'lKb(lLb)lMb*lNb+l,l-l.l/ld :l;l=l?l,c-c.c1d2d3de j f g 6 eb7 8 fbgbhb4d5d6d7d8d9d@lib[ljb]lOb^lkb_llb`lmb{lnb|lob}lpb~lambfbmcm~gah/hbhch:hdheh;hfh=hdmBgemcf?hfm5egmdfef@hffhmgf[himCgjmgh]hhhih^hhfkmiflmmmnmompmqm6erm9 ! !d/cYb:c+bLc;c=cQbRbSbTbUbVbDgsmtmumvmwmxmymzmEgAmBmCmDmEmFmGmHmImFgJmKmLmMmnhNmOmPmohGgHgQmphRmSmTmUmVmWmXmYmZm0m1m2m3m4m5m6m7m8mqhrhsh9m!mthuhvh#mwhxhyh$m%m'm(m)m*m+m,m-m.m/m:m;m=m?m@m[m]m^mqb,b-b_m`m{m|m}m~manbncndnenfngnhninjnknlnmnnnonpnqnrnsntnunvnwnxnynznAnBnCnDnEnFnGnHnIgqcrcJgKgLgMgNgOgPgQgke5crbtcsb6cuctbZb)e*e+ek 9bl m n o p q r s t !bu v w x 3by z 4b#ble7cubvcvb8cwcwb0b,e-e.eA $bB C D E F G H I J %bK L M N 5bO P 6b'bme9cxbxcyb!cyczb1b/e:e;eQ (bR S T U V W X Y Z )b0 1 2 3 7b4 5 8b*bRgInJnKn9eLnMnNnOnPnQnRnSnTnUnVnWnXnYnZn0n1n2n3n4n5n6n7n8n9n!n#n$n%n'n(n)n*n+n,n-n.n/n:n;n=n?n@n[n]n^n_n`n{n|n}n~naobocodoeofogohoiojokolomoSgTg$enooopoqorosotouovowoxoyo#dzoAoBoCoDoEoFoGoWbHoIoJoKoLoMoNoOoPoQoRoSoToUoVoWoXoYoZo0o1o2o3o4o5o6o7o8o9o!o#o$o%o'o(o)o*o+o,o-o.o/o:o;o=o?o@o[o]o^o_o`o{o|o}o~oapbpcpdpepfpgphpipjpkplpmpnpopppqprpsptpupvpwpxpypzpApBpCpDpEpFpGpHpIpJpKpLpMpNpOpPpQpRpSpTpUpVpWpXpYpZp0p1p2p3p4p5p6p7p8p9p!p#p$p%p'p(p)p*p+p,p-p.p/p:p;p=p?p@p[p]p^p_p`p{p|p}p~paqbqcqdqeqfqgqhqiqjqkqlqmqnqoqpqqqrqsqtquqvqwqxqyqzqAqBqCqDqEqFqGqHqIqJqKqLqMqNqOqPqQqRqSqTqUqVqWqXqYqZq0q1q2q3q4q5q6q7q8q9q!q#q$q%q'q(q)q*q+q,q-q.q/q:q;q=q?q@q[q]q^q_q`q{q|q}q~qarbrcrdrerfrgrhrirjrkrlrmrnrorprqrrrsrtrurvrwrxryrzrArBrCrDrErFrGrHrIrJrKrLrMrNrOrPrQrRrSrTrUrVrWrXrYrZr0r1r2r3r4r5r6r7r8r9r!r#r$r%r'r(r)r*r+r,r-r.r/r:r;r=r?r@r[r]r^r_r`r{r|r}r~rasbscsdsesfsgshsisjskslsmsnsospsqsrssstsusvswsxsyszsAsBsCsDsEsFsGsHsIsJsKsLsMsNsOsPsQsRsSsTsUsVsWsXsYsZs0s1s2s3s4s5s6s7s8s9s!s#s$s%s's(s)s*s+s,s-s.s/s:s;s=s?s@s[s]s^s_s`s{s|s}s~satbtctdtetftgthtitjtktltmtntotptqtrtstttutvtwtxtytztAtBtCtDtEtFtGtHtItJtKtLtMtNtOtPtQtRtStTtUtVtWtXtYtZt0t1t2t3t4t5t6t7t8t9t!t#t$t%t't(t)t*t+t,t-t.t/t:t;t=t?t@t[t]t^t_t`t{t|t}t~taubucudueufuguhuiujukulumunuoupuqurusutuuuvuwuxuyuzuAuBuCuDuEuFuGuHuIuJuKuLuMuNuOuPuQuRuSuTuUuVuWuXuYuZu0u1u2u3u4u5u6u7u8u9u!u#u$u%u'u(u)u*u+u,u-u.u/u:u;u=u?u@u[u]u^u_u`u{u|u}u~uavbvcvdvevfvgvhvivjvkvlvmvnvovpvqvrvsvtvuvvvwvxvyvzvAvBvCvDvEvFvGvHvIvJvKvLvMvNvOvPvQvRvSvTvUvVvWvXvYvZv0v1v2v3v4v5v6v7v8v9v!v#v$v%v'v(v)v*v+v,v-v.v/v:v;v=v?v@v[v]v^v_v`v{v|v}v~vawbwcwdwewfwgwhwiwjwkwlwmwnwowpwqwrwswtwuwvwwwxwywzwAwBwCwDwEwFwGwHwIwJwKwLwMwNwOwPwQwRwSwTwUwVwWwXwYwZw0w1w2w3w4w5w6w7w8w9w!w#w$w%w'w(w)w*w+w,w-w.w/w:w;w=w?w@w[w]w^w_w`w{w(e|w}w~waxbxcxdxexfxgxhxixjxkxlxmxnxoxpxqxrxsxtxuxvxwxxxyxzxAxBxCxDxExFxGxHxIxJxKxLxMxNxOxPxQxRxSxTxUxVxWxXxYxZx0x1x2x3x4x5x6x7x8x9x!x#x$x%x'x(x)x*x+x,x-x.x/x:x;x=x?x@x[x]x^x_x`x{x|x}x~xaybycydyeyfygyhyiyjykylymynyoypyqyrysytyuyvywyxyyyzyAyByCyDyEyFyGyHyIyJyKyLyMyNyOyPyQyRySyTyUyVyWyXyYyZy0y1y2y3y4y5y6y7y8y9y!y#y$y%y'y(y)y*y+y,y-y.y/y:y;y=y?y@y[y]y^y_y`y{y|y}y~yazbzczdzezfzgzhzizjzlzmznzozpzqzrzsztzuzvzwzxzyzzzAzBzCzDzEzFzGzHzIzJzKzLzMzNzOzPzQzRzSzTzUzVzWzXzYzZz0z1z2z3z4z5z6z7z8z9z!z#z$z%z'z(z)z*z+z,z-z.z/z:z;z=z?z@z[z]z^z_z`z{z|z}z~zaAbAcAdAeAfAgAhAiAjAkAlAmAnAoApAqArAsAtAuAvAwAxAyAzAAABACADAEAFAGAHAIAJAKALAMANAOAPAQARASATAUAVAWAXAYAZA0A1A2A3A4A5A6A7A8A9A!A#A$A%A'A(A)A*A+A,A-A.A/A:A;A=A?A@A[A]A^A_A`A{A|A}A~AaBbBcBdBeBfBgBhBiBjBkBlBmBnBoBpBqBrBsBtBuBvBwBxByBzBABBBCBDBEBFBGBHBIBJBKBLBMBNBOBPBQBRBSBTBUBVBWBXBYBZB0B1B2B3B4B5B6B7B8B9B!B#B$B%B'B(B)B*B+B,B-B.B/B
220 else:
221 return result[1:] 2a Zd0dkf# $ lf% ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb,c-c.c1d2d3de f g 6 7 8 fbgbhb4d5d6d7d8d9dibjbkblbmbnbobpbb .b/b:b;b=b?b@b[b]b^b_b`b{b|b}b~bacbcccdcecfcgchcicjckclcmcncocpcscRcScTcGhUcHhVcIhWcvgwgXcYc'dZc(d)d*d^c_c`c+d,d-d.d3c4c{c/d|c:d;d=d?d}c@d~cadbdcd[d]d^d_d`d{d|d}d~dddedfdc 0cgd1c2cUgVgWgXgYgZgJhKhLhMhNhOhPhQhRhShThUhVhWhXhYhZh0h1h2h3h4h5h6h7h8h9haebecedeeefegeheieje
224cpdef object check_or_create_options(type cls, object options, str options_description="", bint keep_none=False):
225 """
226 Create the specified options dataclass from a dictionary of options or None.
227 """
228 if options is None: 2Zd0dkf# $ lf% ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdbSeTeUeVeWeAbXeYeZeBbCbDbEbFbDcEcFcGcHcIcJcKcmfnfof}E~EhdpfcFdFidqfgFhFjdrfkFlFkdsfoFpFldtfsFtFmdXb0eufvfwfxf1e2eyfafwFxF*c+cyFzFh i [BzfGbHbIbAf3eJbKbLbMbNbd ,c-c.c1d2d3de j f g 6 eb7 8 fbgbhb4d5d6d7d8d9dibjbObkblbmbnbobpbAFBFbfCFDFEFcfFFGF5edfefffHFIFgfJFKFhfLFifMFNFOFPFQFRFSFTFUFVFWF6eXFYF]BZF^B0F1F2F_B3F`B4F9 5F! 6Fjhkh'h(hlhmh)h*hxi]E~K7F8F9F!F#F$F%F'F(F)F*F+F,F-F.F/F:F;F=F?F@F[F]F^F_F`F{F|F}F~FaGbGcGdG!deG/cfGYbgG:chG+bLc;ciG=cjGkGlGmGnGQbRbSbTbUbVbxLyLzLALBLCLTE7e8eDLEL{Bzc|BFLoGGLpGqGrGsGtGuGAcvGwGxGyG}B~BaCnezGAGBGCGDGEGFGGGHGIGJGKGLGXEYEMGZEndyi#cbC$ccC0EHL1EIL2E3E4ENG5E6E7E8E9E!E#E$EuLvLwLJLKLLL;Bqb,b-bodBcpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdoepeb .bVK/bWK:bXK;bYK=bZK?b0K@b1K[b2K]b3K^b4K_b5K`b@E{b[E|b6K}b7K~b8Kac9Kbc!Kcc#Kdc$Kec%Kfc'Kgc(Khc)Kic*Kjc+Kkcjflc,Kmc-Knc.Koc/KdC:KpcqcrcOGPGke5crbtcsb6cuctbZb)e*e+ek 9bl m n o p q r s t !bu v w x 3by z 4b#ble7cubvcvb8cwcwb0b,e-e.eA $bB C D E F G H I J %bK L M N 5bO P 6b'bme9cxbxcyb!cyczb1b/e:e;eQ (bR S T U V W X Y Z )b0 1 2 3 7b4 5 8b*bQGRGSG?cTGUG9eeCfCVGWGgChCiCjCMc!e@c#eXGYGZG0G1G2G3G4G5G6G7G8G9G!G#G$GkC%GlC'GmC(GnC)GoC*GpC+GqC,GrC-G.G/G:G;G=G?G@G[G]G^G_G`G{G|G}G~G%Ezi'E(EAi)EaH$ebHcHdH[ceHfHgHhHiHjHsCtCNcOcPbPcQcscuCvC#dwCxCyCzCACCcBCWbCCDCECFCGCHCICJCKCLCMCNCOCPCRcScTcQCRCSCUcVcWcTCUCVCWCXCYCZC0C1C2C3C4C5C6C7C8C9C!C#C$C%C'C(C)C*C+C,C-C.C/C:C;C=C?C@C[C]C^C_C`C{C|C}C~CaDbDcDdDeDfDgDhDiDjDkDlDmDnDoDpDqDrDsDtDuDug%e'e]c$dkHBflHmHnHvDCfwDpHqHxDrHsHDftHuHvHyDEfzDxHFfyHzHGfBHqeCHDH*E+EEHFHGHHHIHBire=eJHCiDiKHLHMHseteHfIfJfKfLfMfNHOHNfQHOfRHSHADPfBDUHQfVHWHXHCDRfDDZHSf0H1HTf3H,Eue4H5H6H7H8H9H!H#H$H%HEDve?e'HEiFD(H)H*H+H,H-H.H/H:H;H=H?H@H[HUf]H^HVf`H{HWfGDweHDIDJDKDLDMDNDODPDxeQDBhRDSDTDXfyezeYfZf0f1f2fRd3f4f%c@e[eSd'cTd5f}H~HUdFiGiUDAeVDWDXDYDZD0D1D2D3D4DBe5DCh6D7D8D6fgIhI7fjI-ECekIlImInIoIpIqIrIsItIuI9DDe^evIHi!DwIxIyIzI8fAIBICI#D9f$DEI!fFIGI#fIIEeJIKI.E/ELIMINIOIPIIiFe_eQIJiKiRISITIGeHe$f%f'f(f)f*fUIVI+fXI,fYIZI%D-f'D1I.f2I3I4I(D/f)D6I:f7I8I;f!I:EIe#I$I%I'I(I)I*I+I,I-I*DJe`e.ILi+D/I:I;I=I?I@I[I]I^I_I`I{I|I}I=f~IaJ?fcJdJ@f,DKe-D.D/D:D;D=D?D@D[DLe]DDh^D_D`D[fMeNe]f^f_f`f{fVd|f}f(c{e|eWd)cXd~ffJgJYdMiNi{DOe|D}D~DaEbEcEdEeEfEgEPehEEhiEjEkEagnJoJbgqJ;EQerJsJtJuJvJwJxJyJzJAJBJlERe~eCJOimEDJEJFJGJ(ePiQinEoEHJcgIJJJKJpEdgqErEsEtEMJNJegOJPJQJuEfgvESJTJggUJVJWJwEhgxEYJZJig0J1J2JyEjgzE4J5Jkg6J7J8JAElgBE!J#Jmg$J%J'JCEngDE)J*Jog+J,J-JEEpgFEGE/J:Jqg;J=J?JHErgIEJE[JFKGKXcYc=E'dZc(d)d*d^c_c`c+d,d-d.d3c4c{c/d|c:d;d=d?d}c@d~cadbdcd[d]d^d_d`d{d|d}d~dddedfdc 0cKE2bLE%dHKIKJKKKLKMKNKOKPKQKRKSKTKUKRiSi]JMENEOEgdPE^J_JQE?E1c+h;K=K?K@K[K]K^K_K`K{K2caebecedeeefegezKAKBKCKheDKiejeEKMLNLOLPLQLRL`JSLTLULVLWLXLYLZL0LRE1L{J|J2L}J3L~JaKbKcKdKeKfKgKhKiKjKkKlKmKnKoKpKqKrKsKtKuKvKwKxKyK
229 if keep_none: 2Zd0dkf# $ lf% ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdbSeTeUeVeWeAbXeYeZeBbCbDbEbFbDcEcFcGcHcIcJcKc}E~EcFdFgFhFkFlFoFpFsFtF0e1e2e*c+ch i GbHbIb3eJbKbLbMbNbd ,c-c.c1d2d3de j f g 6 eb7 8 fbgbhb4d5d6d7d8d9dibjbObkblbmbnbobpbAFBFbfCFDFEFcfFFGF5edfefffHFIFgfJFKFhfLFifMFNFOFPFQFRFSFTFUFVFWF6e9 ! jhkh'h(hlhmh)h*hxi]E~K!d/cYb:c+bLc;c=cQbRbSbTbUbVbxLyLzLALBLCL7e8eDLEL{Bzc|BoGpGqGrGsGtGuGAcvGwGxGyG}B~BaCnezGAGBGCGDGEGFGGGHGIGJGKGLGMGnd#c$cILNGqb,b-bBcqcrcOGPGke5crbtcsb6cuctbZb)e*e+ek 9bl m n o p q r s t !bu v w x 3by z 4b#ble7cubvcvb8cwcwb0b,e-e.eA $bB C D E F G H I J %bK L M N 5bO P 6b'bme9cxbxcyb!cyczb1b/e:e;eQ (bR S T U V W X Y Z )b0 1 2 3 7b4 5 8b*bQGRGSG?cTG9eeCfCVGWGgChCiCjCMc!e@c#ekC%GlC'GmC(GnC)GoC*GpC+GqC,GrC-GziAi$ebHcHdH[csCtCuCvC#dwCxCyCzCACCcBCWbCCDCECFCGCHCICJCKCLCMCNCOCPCUcVcWcTCUCVCWCXCYCZC0C1C2C3C4C5C6C7C8C9C!C#C$C%C'C(C)C*C+C,C-C.C/C:C;C=C?C@C[C]C^C_C`C{C|C}C~CaDbDcDdDeDfDgDhDiDjDkDlDmDnDoDpDqDrDsDtDuD]ckHlHmHnHvDwDpHqHxDrHsHtHuHvHyDzDxHyHzHBHqeCHDH*E+EEHFHGHHHIHBire=eJHCiDiKHLHMHseteNHOHQHRHSHADBDUHVHWHXHCDDDZH0H1H3H,Eue4H5H6H7H8H9H!H#H$H%HEDve?e'HEiFD(H)H*H+H,H-H.H/H:H;H=H?H@H[H]H^H`H{HGDweHDIDJDKDLDMDNDODPDxeQDBhRDSDTDyeze@e[eSdTd}H~HFiGiUDAeVDWDXDYDZD0D1D2D3D4DBe5DCh6D7D8DgIhIjI-ECekIlImInIoIpIqIrIsItIuI9DDe^evIHi!DwIxIyIzIAIBICI#D$DEIFIGIIIEeJIKI.E/ELIMINIOIPIIiFe_eQIJiKiRISITIGeHeUIVIXIYIZI%D'D1I2I3I4I(D)D6I7I8I!I:EIe#I$I%I'I(I)I*I+I,I-I*DJe`e.ILi+D/I:I;I=I?I@I[I]I^I_I`I{I|I}I~IaJcJdJ,DKe-D.D/D:D;D=D?D@D[DLe]DDh^D_D`DMeNe{e|eWdXdfJgJMiNi{DOe|D}D~DaEbEcEdEeEfEgEPehEEhiEjEkEnJoJqJ;EQerJsJtJuJvJwJxJyJzJAJBJlERe~eCJOimEDJEJFJPiQinEoEHJIJJJKJpEqErEsEtEMJNJOJPJQJuEvESJTJUJVJWJwExEYJZJ0J1J2JyEzE4J5J6J7J8JAEBE!J#J$J%J'JCEDE)J*J+J,J-JEEFEGE/J:J;J=J?JHEIEJE[JYcc 0cKE2bLE%d]JMENEOEPE^J_JQE1c+hOLPLQLRLTLULVLWLZLRE1L2L}J~JaKbKcKdKeKfKgKhKiKjKkKlKmKnKoKpKqKrKsKtKuKvKwKxKyK
230 return options 2h i ,c-c.ce j f g 6 eb7 8 BFCFDFcfGF5eefHFgfKFLFMFOFQFSFUFWF7e8end#c$cOGPGQGRGSGTG9eeCfCVGWGgChCiCjCMc!e@c#ekC%GlC'GmC(GnC)GoC*GpC+GqC,GrC-GziAibHcHdHCckHqHsHBHqeCHDH*E+EEHFHGHHHIHBire=eJHCiDiKHLHMHQHUH,E-EzIIIEeJIKI.E/ELIMINIOIPIIiFe_eQIJiKiRISITIXI1I:E;EHJNJTJZJ5J#J*J:J~JaKbKcKdKeKfKgKhKiKjKkKlKmKnKoKpKqKrKsKtKuKvKwKxKyK
231 return cls() 2Zd0dkf# $ lf% ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdbSeTeUeVeWeAbXeYeZeBbCbDbEbFbDcEcFcGcHcIcJcKc}E~EcFdFgFhFkFlFoFpFsFtF0e1e2e*c+ch i GbHbIb3eJbKbLbMbNbd ,c-c.c1d2d3de j f g 6 eb7 8 fbgbhb4d5d6d7d8d9dibjbObkblbmbnbobpbAFbfEFFF5edfffIFJFhfifNFPFRFTFVF6e9 ! jhkh'h(hlhmh)h*hxi]E~K!d/cYb:c+bLc;c=cQbRbSbTbUbVbxLyLzLALBLCLDLEL{Bzc|BoGpGqGrGsGtGuGAcvGwGxGyG}B~BaCnezGAGBGCGDGEGFGGGHGIGJGKGLGMGnd#c$cILNGqb,b-bBcqcrcke5crbtcsb6cuctbZb)e*e+ek 9bl m n o p q r s t !bu v w x 3by z 4b#ble7cubvcvb8cwcwb0b,e-e.eA $bB C D E F G H I J %bK L M N 5bO P 6b'bme9cxbxcyb!cyczb1b/e:e;eQ (bR S T U V W X Y Z )b0 1 2 3 7b4 5 8b*b?cMc$e[csCtCuCvC#dwCxCyCzCACCcBCWbCCDCECFCGCHCICJCKCLCMCNCOCPCUcVcWcTCUCVCWCXCYCZC0C1C2C3C4C5C6C7C8C9C!C#C$C%C'C(C)C*C+C,C-C.C/C:C;C=C?C@C[C]C^C_C`C{C|C}C~CaDbDcDdDeDfDgDhDiDjDkDlDmDnDoDpDqDrDsDtDuD]clHmHnHvDwDpHxDrHtHuHvHyDzDxHyHzH*E+EBiDiseteNHOHRHSHADBDVHWHXHCDDDZH0H1H3H,Eue4H5H6H7H8H9H!H#H$H%HEDve?e'HEiFD(H)H*H+H,H-H.H/H:H;H=H?H@H[H]H^H`H{HGDweHDIDJDKDLDMDNDODPDxeQDBhRDSDTDyeze@e[eSdTd}H~HFiGiUDAeVDWDXDYDZD0D1D2D3D4DBe5DCh6D7D8DgIhIjI-ECekIlImInIoIpIqIrIsItIuI9DDe^evIHi!DwIxIyIAIBICI#D$DEIFIGI.E/EIiKiGeHeUIVIYIZI%D'D2I3I4I(D)D6I7I8I!I:EIe#I$I%I'I(I)I*I+I,I-I*DJe`e.ILi+D/I:I;I=I?I@I[I]I^I_I`I{I|I}I~IaJcJdJ,DKe-D.D/D:D;D=D?D@D[DLe]DDh^D_D`DMeNe{e|eWdXdfJgJMiNi{DOe|D}D~DaEbEcEdEeEfEgEPehEEhiEjEkEnJoJqJ;EQerJsJtJuJvJwJxJyJzJAJBJlERe~eCJOimEDJEJFJPiQinEoEIJJJKJpEqErEsEtEMJOJPJQJuEvESJUJVJWJwExEYJ0J1J2JyEzE4J6J7J8JAEBE!J$J%J'JCEDE)J+J,J-JEEFEGE/J;J=J?JHEIEJE[JYcc 0cKE2bLE%d]JMENEOEPE^J_JQE1c+hOLPLQLRLTLULVLWLZLRE1L2L}J
232 elif isinstance(options, cls): 2# $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdbAbBbCbDbEbFbmfnfofhdpfidqfjdrfkdsfldtfmdXbufvfwfxfyfafwFxF*c+cyFzFh i [BzfGbHbIbAfJbKbLbMbNbd e j f g 6 eb7 8 fbgbhbibjbObkblbmbnbobpbXFYF]BZF^B0F1F2F_B3F`B4F9 5F! 6Fjhkh'h(hlhmh)h*hxi]E~K7F8F9F!F#F$F%F'F(F)F*F+F,F-F.F/F:F;F=F?F@F[F]F^F_F`F{F|F}F~FaGbGcGdGeGfGgGhG+bLciGjGkGlGmGnGQbRbSbTbUbVbTE{Bzc|BFLGLAc}B~BaCneXEYEZEndyi#cbC$ccC0EHL1E2E3E4E5E6E7E8E9E!E#E$EuLvLwLJLKLLL;BqbodBcpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdoepeb .bVK/bWK:bXK;bYK=bZK?b0K@b1K[b2K]b3K^b4K_b5K`b@E{b[E|b6K}b7K~b8Kac9Kbc!Kcc#Kdc$Kec%Kfc'Kgc(Khc)Kic*Kjc+Kkcjflc,Kmc-Knc.Koc/KdC:Kpc?cUGeCfCgChCiCjCMc@cXGYGZG0G1G2G3G4G5G6G7G8G9G!G#G$GkClCmCnCoCpCqCrC.G/G:G;G=G?G@G[G]G^G_G`G{G|G}G~G%Ezi'E(EAi)EaH[ceHfHgHhHiHjHNcOcPbPcQcscCcRcScTcQCRCSCug%e'e]c$dBfCfDfEfFfGfqereseteHfIfJfKfLfMfNfOfPfQfRfSfTfueveUfVfWfwexeXfyezeYfZf0f1f2fRd3f4f%c'c5fUdAeBe6f7fCeDe8f9f!f#fEeFeGeHe$f%f'f(f)f*f+f,f-f.f/f:f;fIeJe=f?f@fKeLe[fMeNe]f^f_f`f{fVd|f}f(c)c~fYdOePeagbgQeReGJ(ePiQicgdgegfggghgigjgkglgmgngogpgqgrgFKGKXc=E'dZc(d)d*d^c_c`c+d,d-d.d3c4c{c/d|c:d;d=d?d}c@d~cadbdcd[d]d^d_d`d{d|d}d~dddedfdHKIKJKKKLKMKNKOKPKQKRKSKTKUKRiSigd?E+h;K=K?K@K[K]K^K_K`K{K2caebecedeeefegezKAKBKCKheDKiejeEKMLNL`JSLXLYL0LRE{J|J3L
233 return options 2# $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdbAbBbCbDbEbFbmfnfofhdpfidqfjdrfkdsfldtfmdXbufvfwfxfyfafwFxF*c+cyFzFh i [BzfGbHbIbAfJbKbLbMbNbd e j f g 6 eb7 8 fbgbhbibjbObkblbmbnbobpbXFYF]BZF^B0F1F2F_B3F`B4F9 5F! 6F'h(h)h*h~K7F8F9F!F#F$F%F'F(F)F*F+F,F-F.F/F:F;F=F?F@F[F]F^F_F`F{F|F}F~FaGbGcGdGeGfGgGhG+bLciGjGkGlGmGnGQbRbSbTbUbVb{Bzc|BFLGLAc}B~BaCneXEYEZEndyi#cbC$ccC0EHL1E2E3E4E5E6E7E8E9E!E#E$EuLvLwLJLKLLL;BqbodBcpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdoepeb .bVK/bWK:bXK;bYK=bZK?b0K@b1K[b2K]b3K^b4K_b5K`b@E{b[E|b6K}b7K~b8Kac9Kbc!Kcc#Kdc$Kec%Kfc'Kgc(Khc)Kic*Kjc+Kkcjflc,Kmc-Knc.Koc/KdC:Kpc?cUGeCfCgChCiCjCMc@cXGYGZG0G1G2G3G4G5G6G7G8G9G!G#G$GkClCmCnCoCpCqCrC.G/G:G;G=G?G@G[G]G^G_G`G{G|G}G~GaH[ceHfHgHhHiHjHNcOcPbPcQcscCcRcScTcQCRCSCug%e'e]c$dBfCfDfEfFfGfqereseteHfIfJfKfLfMfNfOfPfQfRfSfTfueveUfVfWfwexeXfyezeYfZf0f1f2fRd3f4f%c'c5fUdAeBe6f7fCeDe8f9f!f#fEeFeGeHe$f%f'f(f)f*f+f,f-f.f/f:f;fIeJe=f?f@fKeLe[fMeNe]f^f_f`f{fVd|f}f(c)c~fYdOePeagbgQeReGJ(ecgdgegfggghgigjgkglgmgngogpgqgrgFKGKXc=E'dZc(d)d*d^c_c`c+d,d-d.d3c4c{c/d|c:d;d=d?d}c@d~cadbdcd[d]d^d_d`d{d|d}d~dddedfdHKIKJKKKLKMKNKOKPKQKRKSKTKUKRiSi?E+h;K=K?K@K[K]K^K_K`K{K2caebecedeeefegezKAKBKCKheDKiejeEKMLNL`JSLXLYL0LRE{J|J3L
234 elif isinstance(options, dict): 2jhkhlhmhxi]ETE;Bod%Ezi'E(EAi)EPiQigd
235 return cls(**options) 2jhkhlhmhxi]Eod%Ezi'E(EAi)EPiQigd
236 else:
237 raise TypeError( 2TE;B
238 f"The {options_description} must be provided as an object " 2TE;B
239 f"of type {cls.__name__} or as a dict with valid {options_description}. " 2TE;B
240 f"The provided object is '{options}'." 2TE;B
241 )
244def _handle_boolean_option(option: bool) -> str:
245 """
246 Convert a boolean option to a string representation.
247 """
248 return "true" if bool(option) else "false" 2b .b/b:b;b=b?b@b[b]b^b_b`b{b|b}b~bacbcccdcecfcgchcicjckclcmcncocpcsc^c_c`c|cbdcdddedfdf3
251def precondition(checker: Callable[..., None], str what="") -> Callable[..., Any]:
252 """
253 A decorator that adds checks to ensure any preconditions are met.
255 Args:
256 checker: The function to call to check whether the preconditions are met. It has
257 the same signature as the wrapped function with the addition of the keyword argument `what`.
258 what: A string that is passed in to `checker` to provide context information.
260 Returns:
261 Callable: A decorator that creates the wrapping.
262 """
264 def outer(wrapped_function: Callable) -> Callable: 2UE
265 """
266 A decorator that actually wraps the function for checking preconditions.
267 """
269 @functools.wraps(wrapped_function) 2UE
270 def inner(*args, **kwargs) -> Any:
271 """
272 Check preconditions and if they are met, call the wrapped function.
273 """
274 checker(*args, **kwargs, what=what) 2UE
275 result = wrapped_function(*args, **kwargs) 2UE
277 return result 2UE
279 return inner 2UE
281 return outer 2UE
284def is_sequence(obj: object) -> bool:
285 """
286 Check if the given object is a sequence (list or tuple).
287 """
288 return isinstance(obj, Sequence) 29 ! zcAcXEYEZEyi#c$c0E1E2E3E4E5E6E7E8E9E!E#E$EuLvLwLqbodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQd@E[E=E3c4c{c}c~cadg3h3i3j3k3RiSil3m3n3[2]2?EUgVgWgXgYgZgo3p3
291def is_nested_sequence(obj: object) -> bool:
292 """
293 Check if the given object is a nested sequence (list or tuple with atleast one list or tuple element).
294 """
295 return is_sequence(obj) and any(is_sequence(elem) for elem in obj) 23c4c[2]2
299class Transaction:
300 """
301 A context manager for transactional operations with undo capability.
303 The Transaction class allows you to register undo actions (callbacks) that will be executed
304 if the transaction is not committed before exiting the context. This is useful for managing
305 resources or operations that need to be rolled back in case of errors or early exits.
307 Usage:
308 with Transaction() as txn:
309 txn.append(some_cleanup_function, arg1, arg2)
310 # ... perform operations ...
311 txn.commit() # Disarm undo actions; nothing will be rolled back on exit
313 Methods:
314 append(fn, *args, **kwargs): Register an undo action to be called on rollback.
315 commit(): Disarm all undo actions; nothing will be rolled back on exit.
316 """
317 def __init__(self) -> None:
318 self._stack = ExitStack() 2NcOcPbPcQc
319 self._entered = False 2NcOcPbPcQc
321 def __enter__(self):
322 self._stack.__enter__() 2NcOcPbPcQc
323 self._entered = True 2NcOcPbPcQc
324 return self 2NcOcPbPcQc
326 def __exit__(self, exc_type, exc, tb):
327 # If exit callbacks remain, they'll run in LIFO order.
328 self._entered = False 2NcOcPbPcQc
329 return self._stack.__exit__(exc_type, exc, tb) 2NcOcPbPcQc
331 def append(self, fn: Callable[..., Any], /, *args: Any, **kwargs: Any) -> None:
332 """
333 Register an undo action (runs if the with-block exits without commit()).
334 Values are bound now via partial so late mutations don't bite you.
335 """
336 if not self._entered: 2NcOcPbPcQc
337 raise RuntimeError("Transaction must be entered before append()")
338 self._stack.callback(partial(fn, *args, **kwargs)) 2NcOcPbPcQc
340 def commit(self) -> None:
341 """
342 Disarm all undo actions. After this, exiting the with-block does nothing.
343 """
344 # pop_all() empties this stack so no callbacks are triggered on exit.
345 self._stack.pop_all() 2NcOcPbPcQc
348# Track whether we've already warned about fork method
349_fork_warning_checked = False
352def reset_fork_warning() -> None:
353 """Reset the fork warning check flag for testing purposes.
355 This function is intended for use in tests to allow multiple test runs
356 to check the warning behavior.
357 """
358 global _fork_warning_checked
359 _fork_warning_checked = False 2ug%e'e]c!h$d
362cdef inline tuple _read_fill_ptr(const char* ptr, Py_ssize_t width):
363 """Extract (value, element_size) from a raw pointer of known width."""
364 cdef unsigned int val
365 if width == 1: 2liminioipiqirisitiuiviwi@Bd 5crbtcsb6cuctbk 9bl m n o p q r s t !bu v w x 3by z 4b#b7cubvcvb8cwcwbA $bB C D E F G H I J %bK L M N 5bO P 6b'b9cxbxcyb!cyczbQ (bR S T U V W X Y Z )b0 1 2 3 7b4 5 8b*b
366 val = (<uint8_t*>ptr)[0] 2rbl m n u v ubB C D K L xbR S T 0 1
367 elif width == 2:
368 val = (<uint16_t*>ptr)[0] 2linipiritivitcsbo p q w x 3bvcvbE F G M N 5bxcybU V W 2 3 7b
369 elif width == 4:
370 val = (<uint32_t*>ptr)[0] 2mioiqisiuiwid uctbk r s t y z 4bwcwbA H I J O P 6byczbQ X Y Z 4 5 8b
371 else:
372 raise ValueError(f"value must be 1, 2, or 4 bytes, got {width}") 2@B5c6c9b!b#b7c8c$b%b'b9c!c(b)b*b
373 return (val, <unsigned int>width) 2liminioipiqirisitiuiviwid rbtcsbuctbk l m n o p q r s t u v w x 3by z 4bubvcvbwcwbA B C D E F G H I J K L M N 5bO P 6bxbxcybyczbQ R S T U V W X Y Z 0 1 2 3 7b4 5 8b
376cpdef tuple _parse_fill_value(value):
377 """Parse a fill/memset value into (raw_value, element_size).
379 Parameters
380 ----------
381 value : int or buffer-protocol object
382 - int: Must be in range [0, 256). Treated as 1-byte fill.
383 - bytes or buffer-protocol: Must be 1, 2, or 4 bytes.
385 Returns
386 -------
387 tuple of (int, int)
388 (raw_value, element_size) where element_size is 1, 2, or 4.
390 Raises
391 ------
392 OverflowError
393 If int value is outside [0, 256).
394 TypeError
395 If value is not an int and does not support the buffer protocol.
396 ValueError
397 If value byte length is not 1, 2, or 4.
398 """
399 cdef uint8_t byte_val
400 cdef Py_buffer buf
402 if isinstance(value, int): 2DcEcFcGcHcIcJcKcaFbFlimieFfFnioiiFjFpiqimFnFrisiqFrFtiuiuFvFviwi@Bh i d 4ee j f g ke5crbtcsb6cuctbZb)e*e+ek 9bl m n o p q r s t !bu v w x 3by z 4b#ble7cubvcvb8cwcwb0b,e-e.eA $bB C D E F G H I J %bK L M N 5bO P 6b'bme9cxbxcyb!cyczb1b/e:e;eQ (bR S T U V W X Y Z )b0 1 2 3 7b4 5 8b*boHwHAHCiPHTHYH2HEi_H|HBhRd]eUdFiGiaIbIcIdIeIfIChiIHiDIHIJiWI0I5I9ILibJeJDhVd}eYdMiNihJiJjJkJlJmJEhpJOiLJRJXJ3J9J(J.J@J
403 byte_val = value 2DcEcFcGcHcIcJcKcaFbFeFfFiFjFmFnFqFrFuFvFh i 4ee j f g Zb)e*e+e0b,e-e.e1b/e:e;eoHwHAHCiPHTHYH2HEi_H|HBhRd]eUdFiGiaIbIcIdIeIfIChiIHiDIHIJiWI0I5I9ILibJeJDhVd}eYdMiNihJiJjJkJlJmJEhpJOiLJRJXJ3J9J(J.J@J
404 return (<unsigned int>byte_val, <unsigned int>1) 2DcEcFcGcHcIcJcKcaFbFeFfFiFjFmFnFqFrFuFvFh i 4ee j f g Zb0b1boHwHAHCiPHTHYH2HEi_H|HBhRd]eUdFiGiaIbIcIdIeIfIChiIHiDIHIJiWI0I5I9ILibJeJDhVd}eYdMiNihJiJjJkJlJmJEhpJOiLJRJXJ3J9J(J.J@J
406 if isinstance(value, bytes): 2liminioipiqirisitiuiviwi@Bd ke5crbtcsb6cuctbk 9bl m n o p q r s t !bu v w x 3by z 4b#ble7cubvcvb8cwcwbA $bB C D E F G H I J %bK L M N 5bO P 6b'bme9cxbxcyb!cyczbQ (bR S T U V W X Y Z )b0 1 2 3 7b4 5 8b*b
407 return _read_fill_ptr(<const char*><bytes>value, <Py_ssize_t>len(value)) 2liminioipiqirisitiuiviwi@B5crbtcsb6cuctb7cubvcvb8cwcwb9cxbxcyb!cyczb
409 if PyObject_GetBuffer(value, &buf, PyBUF_SIMPLE) != 0: 2d kek 9bl m n o p q r s t !bu v w x 3by z 4b#bleA $bB C D E F G H I J %bK L M N 5bO P 6b'bmeQ (bR S T U V W X Y Z )b0 1 2 3 7b4 5 8b*b
410 raise TypeError(
411 f"value must be an int or support the buffer protocol, "
412 f"got {type(value).__name__}"
413 )
414 try: 2d k 9bl m n o p q r s t !bu v w x 3by z 4b#bA $bB C D E F G H I J %bK L M N 5bO P 6b'bQ (bR S T U V W X Y Z )b0 1 2 3 7b4 5 8b*b
415 return _read_fill_ptr(<const char*>buf.buf, buf.len) 2d k 9bl m n o p q r s t !bu v w x 3by z 4b#bA $bB C D E F G H I J %bK L M N 5bO P 6b'bQ (bR S T U V W X Y Z )b0 1 2 3 7b4 5 8b*b
416 finally:
417 PyBuffer_Release(&buf)
420def check_multiprocessing_start_method() -> None:
421 """Check if multiprocessing start method is 'fork' and warn if so."""
422 global _fork_warning_checked
423 if _fork_warning_checked: 2aLbL]B^BcLdL_B`B9 ! jhkh'h(hlhmh)h*heLfL^2gL_2`2hL{2|2iL}2~2jLa3b3kLc3d3lLe3mLsgnLoLpLqL!d/cYb:c+bLc;c=ctgrLsLtLQbRbSbTbUbVbug%e'e]c!h$d
424 return 2aLbL]B^BcLdL_B`B9 ! jhkh'h(hlhmh)h*heLfL^2gL_2`2hL{2|2iL}2~2jLa3b3kLc3d3lLe3mLsgnLoLpLqL!d/cYb:c+bLc;c=ctgrLsLtLQbRbSbTbUbVb$d
425 _fork_warning_checked = True 2Ybug%e'e]c!h$d
427 # Common warning message parts
428 common_message = (
429 "CUDA does not support. Forked subprocesses exhibit undefined behavior, " 2Ybug%e'e]c!h$d
430 "including failure to initialize CUDA contexts and devices. Set the start method "
431 "to 'spawn' before creating processes that use CUDA. "
432 "Use: multiprocessing.set_start_method('spawn')"
433 )
435 try: 2Ybug%e'e]c!h$d
436 start_method = multiprocessing.get_start_method() 2Ybug%e'e]c!h$d
437 if start_method == "fork": 2Ybug%e'e]c$d
438 message = f"multiprocessing start method is 'fork', which {common_message}" 2%e'e]c$d
439 warnings.warn(message, UserWarning, stacklevel=3) 2%e'e]c$d
440 except RuntimeError: 2!h
441 # get_start_method() can raise RuntimeError if start method hasn't been set
442 # In this case, default is 'fork' on Linux, so we should warn
443 if platform.system() == "Linux": 2!h
444 message = (
445 f"multiprocessing start method is not set and defaults to 'fork' on Linux, " 2!h
446 f"which {common_message}"
447 )
448 warnings.warn(message, UserWarning, stacklevel=3) 2!h