Coverage for cuda / core / _utils / cuda_utils.pyx: 93.59%
234 statements
« prev ^ index » next coverage.py v7.14.0, created at 2026-05-22 01:37 +0000
« prev ^ index » next coverage.py v7.14.0, created at 2026-05-22 01:37 +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 Callable
15try:
16 from cuda.bindings import driver, nvrtc, runtime
17except ImportError:
18 from cuda import cuda as driver
19 from cuda import cudart as runtime
20 from cuda import nvrtc
22from cuda.bindings.nvvm import nvvmError
23from cuda.bindings.nvjitlink import nvJitLinkError
25from cpython.buffer cimport PyObject_GetBuffer, PyBuffer_Release, Py_buffer, PyBUF_SIMPLE
27from cuda.bindings cimport cynvrtc, cynvvm, cynvjitlink
29from cuda.core._utils.driver_cu_result_explanations import DRIVER_CU_RESULT_EXPLANATIONS
30from cuda.core._utils.runtime_cuda_error_explanations import RUNTIME_CUDA_ERROR_EXPLANATIONS
33class CUDAError(Exception):
34 pass
37class NVRTCError(CUDAError):
38 pass
42ComputeCapability = namedtuple("ComputeCapability", ("major", "minor"))
45def cast_to_3_tuple(label, cfg):
46 cfg_orig = cfg 2$ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdbebBbCbDbEbFbGbkflfmf?hidnfZgjdofkdpfldqfmdrfndZbsftfufvfwfi j xfHbIbJbyfKbLbMbNbObc f k g h 8 fb9 ! gbhbibjbkbPblbmbnbobpbqb# l CcDcneod$c%crbpd!E@hvg#E[hEcqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdoepeFczfAf]hBfCf^hDfEfqere;eseteFfGfHfIfJfKfLfMfNf0gOfPf1gQfRfueve=eSfTfUfwexeVfyezeWfXfYfZf0fSd1f2f2g'c3g4g?e@e5g6g7g8g9gTd(c!g[eUd#g$g%g3fVdAeBe4f5fCeDe]e6f7f_h8f9fEeFe^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_fWd`f{f)g)c*g+g`e{e,g-g.g/g:gXd*c;g|eYd=g?g@g|fZdOePe}f~fQeRe}eagbg`hcgdg{hegfg|hgghg}higjg~hkglgaimgngbiogpgci9hBECE8B$E!h
47 if isinstance(cfg, int): 2a $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdbebBbCbDbEbFbGbkflfmf?hidnfZgjdofkdpfldqfmdrfndZbsftfufvfwfi j xfHbIbJbyfKbLbMbNbObc f k g h 8 fb9 ! gbhbibjbkbPblbmbnbobpbqb# l CcDcneod$c%crbpd!E@hvg#E[hEcqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdoepeFczfAf]hBfCf^hDfEfqere;eseteFfGfHfIfJfKfLfMfNf0gOfPf1gQfRfueve=eSfTfUfwexeVfyezeWfXfYfZf0fSd1f2f2g'c3g4g?e@e5g6g7g8g9gTd(c!g[eUd#g$g%g3fVdAeBe4f5fCeDe]e6f7f_h8f9fEeFe^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_fWd`f{f)g)c*g+g`e{e,g-g.g/g:gXd*c;g|eYd=g?g@g|fZdOePe}f~fQeRe}eagbg`hcgdg{hegfg|hgghg}higjg~hkglgaimgngbiogpgci9hBECE8B$E!h
48 cfg = (cfg,) 2$ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdbebBbCbDbEbFbGbkflfmfnfofpfqfrfZbsftfufvfwfi j xfHbIbJbyfKbLbMbNbObc f k g h 8 fb9 ! gbhbibjbkbPblbmbnbobpbqb# l CcDcneod$c%crbpd!Evg#EqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdoepeFczfAfBfCfDfEfqereseteFfGfHfIfJfKfLfMfNfOfPfQfRfueveSfTfUfwexeVfyezeWfXfYfZf0fSd1f2f'c(c3fVdAeBe4f5fCeDe6f7f8f9fEeFeGeHe!f#f$f%f'f(f)f*f+f,f-f.f/fIeJe:f;f=fKeLe?fMeNe@f[f]f^f_fWd`f{f)c*c|fZdOePe}f~fQeReagbgcgdgegfggghgigjgkglgmgngogpg9hBE
49 else:
50 common = "must be an int, or a tuple with up to 3 ints" 2?hidZgjdkdldmdnd@hvg[hEc]h^h;e0g1g=e2g'c3g4g?e@e5g6g7g8g9gTd(c!g[eUd#g$g%g]e_h^e'g(g_e)g)c*g+g`e{e,g-g.g/g:gXd*c;g|eYd=g?g@g}e`h{h|h}h~haibici9hCE8B$E!h
51 if not isinstance(cfg, tuple): 2?hidZgjdkdldmdnd@hvg[hEc]h^h;e0g1g=e2g'c3g4g?e@e5g6g7g8g9gTd(c!g[eUd#g$g%g]e_h^e'g(g_e)g)c*g+g`e{e,g-g.g/g:gXd*c;g|eYd=g?g@g}e`h{h|h}h~haibici9hCE8B$E!h
52 raise ValueError(f"{label} {common} (got {type(cfg)})") 2$E
53 if len(cfg) > 3: 2?hidZgjdkdldmdnd@hvg[hEc]h^h;e0g1g=e2g'c3g4g?e@e5g6g7g8g9gTd(c!g[eUd#g$g%g]e_h^e'g(g_e)g)c*g+g`e{e,g-g.g/g:gXd*c;g|eYd=g?g@g}e`h{h|h}h~haibici9hCE8B!h
54 raise ValueError(f"{label} {common} (got tuple with length {len(cfg)})") 2CE
55 if any(not isinstance(val, int) for val in cfg): 2?hidZgjdkdldmdnd@hvg[hEc]h^h;e0g1g=e2g'c3g4g?e@e5g6g7g8g9gTd(c!g[eUd#g$g%g]e_h^e'g(g_e)g)c*g+g`e{e,g-g.g/g:gXd*c;g|eYd=g?g@g}e`h{h|h}h~haibici9h8B!h
56 raise ValueError(f"{label} {common} (got {cfg})") 28B
57 if any(val < 1 for val in cfg): 2$ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdbebBbCbDbEbFbGbkflfmf?hidnfZgjdofkdpfldqfmdrfndZbsftfufvfwfi j xfHbIbJbyfKbLbMbNbObc f k g h 8 fb9 ! gbhbibjbkbPblbmbnbobpbqb# l CcDcneod$c%crbpd!E@hvg#E[hEcqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdoepeFczfAf]hBfCf^hDfEfqere;eseteFfGfHfIfJfKfLfMfNf0gOfPf1gQfRfueve=eSfTfUfwexeVfyezeWfXfYfZf0fSd1f2f2g'c3g4g?e@e5g6g7g8g9gTd(c!g[eUd#g$g%g3fVdAeBe4f5fCeDe]e6f7f_h8f9fEeFe^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_fWd`f{f)g)c*g+g`e{e,g-g.g/g:gXd*c;g|eYd=g?g@g|fZdOePe}f~fQeRe}eagbg`hcgdg{hegfg|hgghg}higjg~hkglgaimgngbiogpgci9hBE!h
58 plural_s = "" if len(cfg) == 1 else "s" 2vgBE!h
59 raise ValueError(f"{label} value{plural_s} must be >= 1 (got {cfg_orig})") 2vgBE!h
60 return cfg + (1,) * (3 - len(cfg)) 2$ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdbebBbCbDbEbFbGbkflfmf?hidnfZgjdofkdpfldqfmdrfndZbsftfufvfwfi j xfHbIbJbyfKbLbMbNbObc f k g h 8 fb9 ! gbhbibjbkbPblbmbnbobpbqb# l CcDcneod$c%crbpd!E@hvg#E[hEcqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdoepeFczfAf]hBfCf^hDfEfqere;eseteFfGfHfIfJfKfLfMfNf0gOfPf1gQfRfueve=eSfTfUfwexeVfyezeWfXfYfZf0fSd1f2f2g'c3g4g?e@e5g6g7g8g9gTd(c!g[eUd#g$g%g3fVdAeBe4f5fCeDe]e6f7f_h8f9fEeFe^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_fWd`f{f)g)c*g+g`e{e,g-g.g/g:gXd*c;g|eYd=g?g@g|fZdOePe}f~fQeRe}eagbg`hcgdg{hegfg|hgghg}higjg~hkglgaimgngbiogpgci9h
63cdef int HANDLE_RETURN(cydriver.CUresult err) except?-1 nogil:
64 if err != cydriver.CUresult.CUDA_SUCCESS: 2a 0dAL1dBLCL$ DL% ELFL' GL( HL) IL* JL+ KL, LL- ML. NL/ OL: PL; QL= RL? SL@ TL[ UL] VL^ WL_ XL` YL{ ZL| 0L} 1L~ 2Lab3Lbb4Lcb5Ldb6Leb7L8L9L!L#L$L%L'L(L)L[g*L+L,L-L.L/L:L;L=L?L@L[L]L^L_L`L{L|L}L~LaMbMcMdMeMfMgMSehMiMjMkMTelMUemMVenMWeoMBbpMXeqMYerMZesMtMuMvMwMCbxMDbyMEbzMFbAMGbBMGcCMHcDMIcEMJcFMKcGMLcHMMcIMNcJMKMLMMMNMOMwgPMxgQMygRMSMTMUMVMWMXMkfYMlfZM0M1M2Mmf3M%E'E4M5M6M7M8Mid9M(E)Ediei!M#M$M%M'M(M)Mnf*M+M*E,M+E-M.M/M:M;M=M?M@M[MZgjd]M^M_M,E`M-E{Mfi|Mgi}M~MaNbNcNdNeNfNgNhNiNjNkNlNmNnNoNpNqNrNsNtNuNvNwNxNofyN.E/EzNANBNCNDNkdEN:E;EhiiiFNGNHNINJNKNLNpfMNNNON=EPN?EQNRNSNTNUNVNWNXNYNZN0Nld1N2N3N@E4N[E5Nji6Nki7N8N9N!N#N$N%N'N(Nqf)N*N+N]E,N^E-N.N/N:N;N=N?N@N[N]N^Nmd_N`N{N_E|N`E}Nli~NmiaObOcOdOeOfOgOhOiOrfjOkOlO{EmO|EnOoOpOqOrOsOtOuOvOwOxOndyOzOAO}EBO~ECOniDOoiEOFOGOHOIOJOKOLOMONOOOPOQOROSOTOUOVOWOXOYOZO0O1O2O3O4O5O6O7O8O9O!O#O$O%O'O(O)O*O+O,O]g-O.O/O:O;O=O?O@O[O]O^O_O`O{O|O^g}O~OaPbPcP9BdPZbePfPgP_ghP0eiPjPkPlPmPnPoPpPqPrPsPtPuPvPwPxPyPzPAPBPCPDPsfEPtfFPufGPvfHP1eIP2eJPwfKPaFLPbFMPcFNP+cOP,cPPdFQPeFRPi SPj TP!BUPxfVPHbWPIbXPJbYPyfZP3e0P`g1PKb2PLb3PMb4PNb5POb6P7P8P9P!P#P$P%P'Pc (P{g)P4e*P|g+P-c,P.c-P/c.P2d/P3d:P4d;Pf =Pk ?Pg @Ph [P8 ]Pfb^P9 _P! `Pgb{Phb|Pib}P5d~P6daQ7dbQ8dcQ9ddQ!deQjbfQkbgQPbhQlbiQmbjQnbkQoblQpbmQqbnQfFgF~ehFoQ}g~gahbhchdhpQehqQiFzgrQjFafkFlF5esQbfcfdfmFnFefoFpFAgtQuQfhghhhffqFgfrFsFtFuFvFwFxFyFzFAFBF6evQ8KCF9KDF#BEF$BFF!KGF#KHF%BIF'BJF# KFl LFihwQjhxQ#hyQ$hzQkhAQlhBQ%hCQOcDQpiEQFQGQ$KMFqgNFOFPFQF%KRFHQSFTFUF'KVFIQWFXFYF(KZFJQ0F1F2F)K3FKQ4F5F6F*K7FLQ8F9F!F+K#FMQ$F%F,K'F-K(F.K)F/K*Frg+F:K,F:c-F#d.F;c/F=c:F.bNQRbBg?c;F@c=Fsg?F;K@Ftg[F=K]FSbOQTbPQUbQQVbRQWbSQXbTQCgUQVQWQXQYQZQ0Q1QDg7e8e2Q3Q4Q5Q6Q7Q8Q9Q!Q#QEg$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~RaSbScSdSeSfSgShSiSjSkSlSmSnSoSpSqSrSsStSuSvSwSxSySzSASBSCSDSESFSGSmhHSISJSnhFgGgoh(BCc)BKS^FLS_FMS`FNS{FOS|FPS}FQS~FRSDcSSaGTSbGUScGVSdGWS*BXS+BYS,BZSne0SeGfGgGhGiG1S2S3S4S5S6SjGkGlGmG7SnGoG8SpGqG9S!S#S$SDE%S'SEEphrGFEqhodqirh$c-B%c.B(SGEshHEthIEuh)SJEvh*SKEwhsGLExh+S,S-S.S/S:SME;SNE=SOE?SPE@SQE[SRE]SSE^S_S`S{S|S}S~SaTbTrb/b:bpdcTdTeTEcfTqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdoepe;b=b?b@b[b]b^b_b`b{b|b}b~bacbcccdcecfcgchcicjckclcmcncocpcqcrc/BscgThTiTHgtcucIgjTJgKgLgtGuGkTMgNgOgPgyhzhlTmTsbnToTtbpTqTrTubsT0btTuTvTwTm xTyTn zTo ATp BTq CTr DTs ETt FTu GTv HTITw JTx KTy LTz MTNTA OTB PTQTRTSTTTvbUTVTwbWTXTYTxbZT1b0T1T2T3TC 4T5TD 6TE 7TF 8TG 9TH !TI #TJ $TK %TL 'T(TM )TN *TO +TP ,T-TQ .TR /T:T;T=T?Tyb@T[Tzb]T^T_TAb`T2b{T|T}T~TS aUbUT cUU dUV eUW fUX gUY hUZ iU0 jU1 kUlU2 mU3 nU4 oU5 pUqU6 rU7 sUtUuUQgvGwG[cvUwU9e:B;BxUxGyUyGzU=BAU?BBU@BCU[BDUPcEU!eFU]cGU#eHUzGIUAGJUBGKUCGLUDGMUEGNUFGOUGGPUHGQUIGRUJGSUKGTULGUUMGVUNGWUOGXU]BYUPGZU^B0UQG1U_B2URG3U`B4USG5U{B6UTG7U|B8UUG9U}B!UVG#U~B$UWG%UXG'UYG(UZG)U0G*U1G+U2G,U3G-U4G.U5G/U6G:U7G;U8G=U9G?U!G@U#G[U$G]UTE^Uri_UUE`UVE{Usi|UWE}U%G~UaVRgSg$e'G(G)GbV^ccV*G+GdV,GeV-GfV.GgVQcRcQbd SchViVvcjVaCkVbC$dlVcCdCmVeCnVfCoVgCFcpVhCYbiCqVjCrVkCsVlCtVmCuVnCvVoCwVpCxVqCyVrCzVsCtCuCTcUcvCwCVcWcAVxCBVyCCVzCDVACEVBCFVCCGVDCHVECIVFCJVGCKVHCLVICMVJCNVKCOVLCPVMCQVNCRVOCSVPCTVQCUVRCVVSCWVTCXVUCVCWCYVXCZVYC0VZC1V0C2V1C3V2C4V3C5V4C6V5C7V6C8V7C9V8C!V9C#V!C$V#C%V$C'V%C(V'C)V(C*V)C+V*C+C,V,C-V-C.V.C/V/C:V:C;V;C=V=C?V?C@V@C[V[C]V]C^V^C_V_C`V`C{V{C|V|C}Cug}V%e~V'eaW_cbW%dcWdW/GzfeWfWgWhW:G;G=GiWjWkWlWmW~CAfnW?GaDoW@GpWqW[GbD]GrW^GBfsWtWuWvW_G`G{GwWxWyWzWAWcDCfBW|GdDCW}GDWEWDf~GaHFWGWHWIWEfJWbHKWLWcHqedHeHXEYEfHgHhHiHjHtire;ekHuivilHmHnHseteFfGfHfIfJfMWKfNWoHpHOWPWQWRWLfSWqHTWUWVWWWrHMfXWYWsHtHZW0W1W2W3WeDNf0g4WuHfD5W6W7W8WvHOf9W!W#WwHxHyH$W%W'W(W)WgDPf1g*WzHhD+WAH,W-W.WQfBHCH/W:W;W=WRf?WDH@W[WEHZEueFHGHHHIHJHKHLHMHNHOHiDve=ePHwijDQHRHSHTHUHVHWHXHYHZH0H1H2H3HSf4H5H]W^W_W`W6H{W|W}WTf~WaX7H8HbXcXdXeXUffX9HgXhXiXjXkXlXmXnXoXpXqXrXkDwelDmDnDoDpDqDrDsDtDxeuDAhvDwDxDVfyezeWfXfYfZf0fSd1f2f2g'c3g4g?e@e5g6g7g8g9gTd(c!g[eUd#g$g%g3f!H#HsXtXuXvXwXVdxiyi$H%H'H(H)H*HyDAezDADBDCDDDEDFDGDHDIDBeJDBhKDLDMDxX4fyX+H,HzXAXBXCX5fDX-HEXFX.H0ECe/H:H;H=H?H@H[H]H^H_H`HNDDe]e{HziOD|H}H~HGXHXaI6fIXJXKXLXbIcIdIMXNXOXPXQXPD7fRXeIQDSXfITXUX8fgIhIVXWXXXYX9fZXiI0X1XjIEekIlI1E2EmInIoIpIqIAiFe^erIBiCisItIuIGeHe!f#f$f%f'f2X(f3XvIwI4X5X6X7X)f8XxI9X!X#X$XyI*f%X'XzIAI(X)X*X+X,XRD+f'g-XBISD.X/X:X;XCI,f=X?X@XDIEIFI[X]X^X_X`XTD-f(g{XGIUD|XHI}X~XaY.fIIJIbYcYdYeY/ffYKIgYhYLI3EIeMINIOIPIQIRISITIUIVIVDJe_eWIDiWDXIYIZI0I1I2I3I4I5I6I7I8I9I!I:f#I$IiYjYkYlY%ImYnYoY;fpYqY'I(IrYsYtYuY=fvY)IwYxYyYzYAYBYCYDYEYFYGYHYXDKeYDZD0D1D2D3D4D5D6DLe7DCh8D9D!D?fMeNe@f[f]f^f_fWd`f{f)g)c*g+g`e{e,g-g.g/g:gXd*c;g|eYd=g?g@g|f*I+IIYJYKYLYMYZdEiFi,I-I.I/I:I;I#DOe$D%D'D(D)D*D+D,D-D.DPe/DDh:D;D=DNY}fOY=I?IPYQYRYSY~fTY@IUYVY[I4EQe]I^I_I`I{I|I}I~IaJbJcJ?DRe}edJGi@DeJfJgJWYXYYYhJiJHiIi[D]DZYjJag0Y1Y2Y3YkJlJmJ4Y5Y6Y7Y8Y^Dbg9YnJ_D`DXc{D|D3b!YoJ#Y$Y%YpJcg'Y(Y)Y*YqJrJsJ+Y,Y-Y.Y/Y}Ddg:YtJ~D;YuJ=Y?Y@YvJeg[Y]Y^Y_YwJxJyJ`Y{Y|Y}Y~YaEfgaZzJbEbZAJcZdZeZBJggfZgZhZiZCJDJEJjZkZlZmZnZcEhgoZFJdEpZGJqZrZsZHJigtZuZvZwZIJJJKJxZyZzZAZBZeEjgCZLJfEDZMJEZFZGZNJkgHZIZJZKZOJPJQJLZMZNZOZPZgElgQZRJhERZSJSZTZUZTJmgVZWZXZYZUJVJWJZZ0Z1Z2Z3ZiEng4ZXJjEkE5ZYJ6Z7Z8ZZJog9Z!Z#Z$Z0J1J2J%Z'Z(Z)Z*ZlEpg+Z3JmEnE,Z4J-Z.ZYc/ZZc:Z5E0c;Z1c=Z?Z@Z[Z]Z^Z_Z`Z{Z|Z}Z~Za0b0c0d0e0f0g0h0i0j0k0l0m0n0o0p0q0r0s0t0u0v0w0x0y0z0A0B0oE4bpEqEC0D0E0F0G0H0I0J0K0L0M0N0O0P0JiKi5JrEsEtE5bQ0uE6J7JvE6ER0S02cT0U0V0W0X0Y0Z03c00102030405060708090!0#08J$0%0'0(0)0*0+0,0-0.0wE/0:09J;0!J=0?0#J@0[0]0^0_0`0{0|0}0~0a1b1c1d1e1f1g1h1i1j1k1l1m1n1o1p1q1r1s1t1u1v1w1x1y1z1A1B1C1D1E1F1G1H1I1J1K1L1M1N1O1P1Q1R1S1T1U1V1W1X1Y1Z101112131415161718191!1$J%J#1$1%1'J(J)J'1(1)1*1+1,1*J+J,J-J.J-1.1/1/J:J;J=J?J@J[J:1;1=1]J^J?1@1_J[1`J]1^1_1{J`1{1|1}1~1a2b2c2d2e2f2g2h2i2j2k2l2m2n2o2p2q2r2s2t2u2v2w2x2y2z2A2B2C2D2E2F2G2H2
65 return _check_driver_error(err) 2[g]g^g_g{g4e|gYb
66 return 0 2a 0dAL1dBLCL$ DL% ELFL' GL( HL) IL* JL+ KL, LL- ML. NL/ OL: PL; QL= RL? SL@ TL[ UL] VL^ WL_ XL` YL{ ZL| 0L} 1L~ 2Lab3Lbb4Lcb5Ldb6Leb7L8L9L!L#L$L%L'L(L)L[g*L+L,L-L.L/L:L;L=L?L@L[L]L^L_L`L{L|L}L~LaMbMcMdMeMfMgMSehMiMjMkMTelMUemMVenMWeoMBbpMXeqMYerMZesMtMuMvMwMCbxMDbyMEbzMFbAMGbBMGcCMHcDMIcEMJcFMKcGMLcHMMcIMNcJMKMLMMMNMOMwgPMxgQMygRMSMTMUMVMWMXMkfYMlfZM0M1M2Mmf3M%E'E4M5M6M7M8Mid9M(E)Ediei!M#M$M%M'M(M)Mnf*M+M*E,M+E-M.M/M:M;M=M?M@M[MZgjd]M^M_M,E`M-E{Mfi|Mgi}M~MaNbNcNdNeNfNgNhNiNjNkNlNmNnNoNpNqNrNsNtNuNvNwNxNofyN.E/EzNANBNCNDNkdEN:E;EhiiiFNGNHNINJNKNLNpfMNNNON=EPN?EQNRNSNTNUNVNWNXNYNZN0Nld1N2N3N@E4N[E5Nji6Nki7N8N9N!N#N$N%N'N(Nqf)N*N+N]E,N^E-N.N/N:N;N=N?N@N[N]N^Nmd_N`N{N_E|N`E}Nli~NmiaObOcOdOeOfOgOhOiOrfjOkOlO{EmO|EnOoOpOqOrOsOtOuOvOwOxOndyOzOAO}EBO~ECOniDOoiEOFOGOHOIOJOKOLOMONOOOPOQOROSOTOUOVOWOXOYOZO0O1O2O3O4O5O6O7O8O9O!O#O$O%O'O(O)O*O+O,O]g-O.O/O:O;O=O?O@O[O]O^O_O`O{O|O^g}O~OaPbPcP9BdPZbePfPgP_ghP0eiPjPkPlPmPnPoPpPqPrPsPtPuPvPwPxPyPzPAPBPCPDPsfEPtfFPufGPvfHP1eIP2eJPwfKPaFLPbFMPcFNP+cOP,cPPdFQPeFRPi SPj TP!BUPxfVPHbWPIbXPJbYPyfZP3e0P`g1PKb2PLb3PMb4PNb5POb6P7P8P9P!P#P$P%P'Pc (P{g)P4e*P|g+P-c,P.c-P/c.P2d/P3d:P4d;Pf =Pk ?Pg @Ph [P8 ]Pfb^P9 _P! `Pgb{Phb|Pib}P5d~P6daQ7dbQ8dcQ9ddQ!deQjbfQkbgQPbhQlbiQmbjQnbkQoblQpbmQqbnQfFgF~ehFoQ}g~gahbhchdhpQehqQiFzgrQjFafkFlF5esQbfcfdfmFnFefoFpFAgtQuQfhghhhffqFgfrFsFtFuFvFwFxFyFzFAFBF6evQ8KCF9KDF#BEF$BFF!KGF#KHF%BIF'BJF# KFl LFihwQjhxQ#hyQ$hzQkhAQlhBQ%hCQOcDQpiEQFQGQ$KMFqgNFOFPFQF%KRFHQSFTFUF'KVFIQWFXFYF(KZFJQ0F1F2F)K3FKQ4F5F6F*K7FLQ8F9F!F+K#FMQ$F%F,K'F-K(F.K)F/K*Frg+F:K,F:c-F#d.F;c/F=c:F.bNQRbBg?c;F@c=Fsg?F;K@Ftg[F=K]FSbOQTbPQUbQQVbRQWbSQXbTQCgUQVQWQXQYQZQ0Q1QDg7e8e2Q3Q4Q5Q6Q7Q8Q9Q!Q#QEg$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~RaSbScSdSeSfSgShSiSjSkSlSmSnSoSpSqSrSsStSuSvSwSxSySzSASBSCSDSESFSGSmhHSISJSnhFgGgoh(BCc)BKS^FLS_FMS`FNS{FOS|FPS}FQS~FRSDcSSaGTSbGUScGVSdGWS*BXS+BYS,BZSne0SeGfGgGhGiG1S2S3S4S5S6SjGkGlGmG7SnGoG8SpGqG9S!S#S$SDE%S'SEEphrGFEqhodqirh$c-B%c.B(SGEshHEthIEuh)SJEvh*SKEwhsGLExh+S,S-S.S/S:SME;SNE=SOE?SPE@SQE[SRE]SSE^S_S`S{S|S}S~SaTbTrb/b:bpdcTdTeTEcfTqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdoepe;b=b?b@b[b]b^b_b`b{b|b}b~bacbcccdcecfcgchcicjckclcmcncocpcqcrc/BscgThTiTHgtcucIgjTJgKgLgtGuGkTMgNgOgPgyhzhlTmTsbnToTtbpTqTrTubsT0btTuTvTwTm xTyTn zTo ATp BTq CTr DTs ETt FTu GTv HTITw JTx KTy LTz MTNTA OTB PTQTRTSTTTvbUTVTwbWTXTYTxbZT1b0T1T2T3TC 4T5TD 6TE 7TF 8TG 9TH !TI #TJ $TK %TL 'T(TM )TN *TO +TP ,T-TQ .TR /T:T;T=T?Tyb@T[Tzb]T^T_TAb`T2b{T|T}T~TS aUbUT cUU dUV eUW fUX gUY hUZ iU0 jU1 kUlU2 mU3 nU4 oU5 pUqU6 rU7 sUtUuUQgvGwG[cvUwU9e:B;BxUxGyUyGzU=BAU?BBU@BCU[BDUPcEU!eFU]cGU#eHUzGIUAGJUBGKUCGLUDGMUEGNUFGOUGGPUHGQUIGRUJGSUKGTULGUUMGVUNGWUOGXU]BYUPGZU^B0UQG1U_B2URG3U`B4USG5U{B6UTG7U|B8UUG9U}B!UVG#U~B$UWG%UXG'UYG(UZG)U0G*U1G+U2G,U3G-U4G.U5G/U6G:U7G;U8G=U9G?U!G@U#G[U$G]UTE^Uri_UUE`UVE{Usi|UWE}U%G~UaVRgSg$e'G(G)GbV^ccV*G+GdV,GeV-GfV.GgVQcRcQbd SchViVvcjVaCkVbC$dlVcCdCmVeCnVfCoVgCFcpVhCYbiCqVjCrVkCsVlCtVmCuVnCvVoCwVpCxVqCyVrCzVsCtCuCTcUcvCwCVcWcAVxCBVyCCVzCDVACEVBCFVCCGVDCHVECIVFCJVGCKVHCLVICMVJCNVKCOVLCPVMCQVNCRVOCSVPCTVQCUVRCVVSCWVTCXVUCVCWCYVXCZVYC0VZC1V0C2V1C3V2C4V3C5V4C6V5C7V6C8V7C9V8C!V9C#V!C$V#C%V$C'V%C(V'C)V(C*V)C+V*C+C,V,C-V-C.V.C/V/C:V:C;V;C=V=C?V?C@V@C[V[C]V]C^V^C_V_C`V`C{V{C|V|C}Cug}V%e~V'eaW_cbW%dcWdW/GzfeWfWgWhW:G;G=GiWjWkWlWmW~CAfnW?GaDoW@GpWqW[GbD]GrW^GBfsWtWuWvW_G`G{GwWxWyWzWAWcDCfBW|GdDCW}GDWEWDf~GaHFWGWHWIWEfJWbHKWLWcHqedHeHXEYEfHgHhHiHjHtire;ekHuivilHmHnHseteFfGfHfIfJfMWKfNWoHpHOWPWQWRWLfSWqHTWUWVWWWrHMfXWYWsHtHZW0W1W2W3WeDNf0g4WuHfD5W6W7W8WvHOf9W!W#WwHxHyH$W%W'W(W)WgDPf1g*WzHhD+WAH,W-W.WQfBHCH/W:W;W=WRf?WDH@W[WEHZEueFHGHHHIHJHKHLHMHNHOHiDve=ePHwijDQHRHSHTHUHVHWHXHYHZH0H1H2H3HSf4H5H]W^W_W`W6H{W|W}WTf~WaX7H8HbXcXdXeXUffX9HgXhXiXjXkXlXmXnXoXpXqXrXkDwelDmDnDoDpDqDrDsDtDxeuDAhvDwDxDVfyezeWfXfYfZf0fSd1f2f2g'c3g4g?e@e5g6g7g8g9gTd(c!g[eUd#g$g%g3f!H#HsXtXuXvXwXVdxiyi$H%H'H(H)H*HyDAezDADBDCDDDEDFDGDHDIDBeJDBhKDLDMDxX4fyX+H,HzXAXBXCX5fDX-HEXFX.H0ECe/H:H;H=H?H@H[H]H^H_H`HNDDe]e{HziOD|H}H~HGXHXaI6fIXJXKXLXbIcIdIMXNXOXPXQXPD7fRXeIQDSXfITXUX8fgIhIVXWXXXYX9fZXiI0X1XjIEekIlI1E2EmInIoIpIqIAiFe^erIBiCisItIuIGeHe!f#f$f%f'f2X(f3XvIwI4X5X6X7X)f8XxI9X!X#X$XyI*f%X'XzIAI(X)X*X+X,XRD+f'g-XBISD.X/X:X;XCI,f=X?X@XDIEIFI[X]X^X_X`XTD-f(g{XGIUD|XHI}X~XaY.fIIJIbYcYdYeY/ffYKIgYhYLI3EIeMINIOIPIQIRISITIUIVIVDJe_eWIDiWDXIYIZI0I1I2I3I4I5I6I7I8I9I!I:f#I$IiYjYkYlY%ImYnYoY;fpYqY'I(IrYsYtYuY=fvY)IwYxYyYzYAYBYCYDYEYFYGYHYXDKeYDZD0D1D2D3D4D5D6DLe7DCh8D9D!D?fMeNe@f[f]f^f_fWd`f{f)g)c*g+g`e{e,g-g.g/g:gXd*c;g|eYd=g?g@g|f*I+IIYJYKYLYMYZdEiFi,I-I.I/I:I;I#DOe$D%D'D(D)D*D+D,D-D.DPe/DDh:D;D=DNY}fOY=I?IPYQYRYSY~fTY@IUYVY[I4EQe]I^I_I`I{I|I}I~IaJbJcJ?DRe}edJGi@DeJfJgJWYXYYYhJiJHiIi[D]DZYjJag0Y1Y2Y3YkJlJmJ4Y5Y6Y7Y8Y^Dbg9YnJ_D`DXc{D|D3b!YoJ#Y$Y%YpJcg'Y(Y)Y*YqJrJsJ+Y,Y-Y.Y/Y}Ddg:YtJ~D;YuJ=Y?Y@YvJeg[Y]Y^Y_YwJxJyJ`Y{Y|Y}Y~YaEfgaZzJbEbZAJcZdZeZBJggfZgZhZiZCJDJEJjZkZlZmZnZcEhgoZFJdEpZGJqZrZsZHJigtZuZvZwZIJJJKJxZyZzZAZBZeEjgCZLJfEDZMJEZFZGZNJkgHZIZJZKZOJPJQJLZMZNZOZPZgElgQZRJhERZSJSZTZUZTJmgVZWZXZYZUJVJWJZZ0Z1Z2Z3ZiEng4ZXJjEkE5ZYJ6Z7Z8ZZJog9Z!Z#Z$Z0J1J2J%Z'Z(Z)Z*ZlEpg+Z3JmEnE,Z4J-Z.ZYc/ZZc:Z5E0c;Z1c=Z?Z@Z[Z]Z^Z_Z`Z{Z|Z}Z~Za0b0c0d0e0f0g0h0i0j0k0l0m0n0o0p0q0r0s0t0u0v0w0x0y0z0A0B0oE4bpEqEC0D0E0F0G0H0I0J0K0L0M0N0O0P0JiKi5JrEsEtE5bQ0uE6J7JvE6ER0S02cT0U0V0W0X0Y0Z03c00102030405060708090!0#08J$0%0'0(0)0*0+0,0-0.0wE/0:09J;0!J=0?0#J@0[0]0^0_0`0{0|0}0~0a1b1c1d1e1f1g1h1i1j1k1l1m1n1o1p1q1r1s1t1u1v1w1x1y1z1A1B1C1D1E1F1G1H1I1J1K1L1M1N1O1P1Q1R1S1T1U1V1W1X1Y1Z101112131415161718191!1$J%J#1$1%1'J(J)J'1(1)1*1+1,1*J+J,J-J.J-1.1/1/J:J;J=J?J@J[J:1;1=1]J^J?1@1_J[1`J]1^1_1{J`1{1|1}1~1a2b2c2d2e2f2g2h2i2j2k2l2m2n2o2p2q2r2s2t2u2v2w2x2y2z2A2B2C2D2E2F2G2H2
69cdef int HANDLE_RETURN_NVRTC(cynvrtc.nvrtcProgram prog, cynvrtc.nvrtcResult err) except?-1 nogil:
70 """Handle NVRTC result codes, raising NVRTCError with program log on failure."""
71 if err == cynvrtc.nvrtcResult.NVRTC_SUCCESS: 2$ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdbebBbCbDbEbFbGbkflfmfidnfjdofkdpfldqfmdrfndZbsftfufvfwfi j !BxfHbIbJbyfKbLbMbNbObc f k g h 8 fb9 ! gbhbibjbkbPblbmbnbobpbqb# l CcDcneqi-B.BrbpdEcqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdoepeb ;b=b?b@b[b]b^b_b`b{b|b}b~bacbcccdcecfcgchcicjckclcmcncocpcqcrc/BscvcaCbC$dcCdCeCfCgCFchCYbiCjCkClCmCnCoCpCqCrCsCtCuCTcUcvCwCVcWcxCyCzCACBCCCDCECFCGCHCICJCKCLCMCNCOCPCQCRCSCTCUCVCWCXCYCZC0C1C2C3C4C5C6C7C8C9C!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}Czf~CAfaDbDBfcDCfdDDfEfqetireviseteFfGfHfIfJfKfLfMfeDNffDOfgDPfhDQfRfueiDvejDSfTfUfkDwelDmDnDoDpDqDrDsDtDxeuDAhvDwDxDVfyezeWfXfYfZf0fSd1f2f'cTd(cUd3fVdyDAezDADBDCDDDEDFDGDHDIDBeJDBhKDLDMD4f5fCeNDDeOD6fPD7fQD8f9fEeAiFeCiGeHe!f#f$f%f'f(f)f*fRD+fSD,fTD-fUD.f/fIeVDJeWD:f;f=fXDKeYDZD0D1D2D3D4D5D6DLe7DCh8D9D!D?fMeNe@f[f]f^f_fWd`f{f)cXd*cYd|fZd#DOe$D%D'D(D)D*D+D,D-D.DPe/DDh:D;D=D}f~fQe?DRe@D[D]Dag^Dbg_D`DXc{D3bcg}Ddg~DegaEfgbEggcEhgdEigeEjgfEkggElghEmgiEngjEkEoglEpgmEnEYcZc0c1c'd(d)d`c{c|c*d+d,d-d4c5c}c.d~c/d:d;d=dad?dbdcddded@d[d]d^d_d`d{d|d}dfdgdhde ~drEsEtE5buEvE2c3caebecedeeefege|J}J~JaKhebKiejecK
72 return 0 2$ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdbebBbCbDbEbFbGbkflfmfidnfjdofkdpfldqfmdrfndZbsftfufvfwfi j !BxfHbIbJbyfKbLbMbNbObc f k g h 8 fb9 ! gbhbibjbkbPblbmbnbobpbqb# l CcDcneqi-B.BrbpdEcqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdoepeb ;b=b?b@b[b]b^b_b`b{b|b}b~bacbcccdcecfcgchcicjckclcmcncocpcqcrc/BscvcaCbC$dcCdCeCfCgCFchCYbiCjCkClCmCnCoCpCqCrCsCtCuCTcUcvCwCVcWcxCyCzCACBCCCDCECFCGCHCICJCKCLCMCNCOCPCQCRCSCTCUCVCWCXCYCZC0C1C2C3C4C5C6C7C8C9C!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}Czf~CAfaDbDBfcDCfdDDfEfqetireviseteFfGfHfIfJfKfLfMfeDNffDOfgDPfhDQfRfueiDvejDSfTfUfkDwelDmDnDoDpDqDrDsDtDxeuDAhvDwDxDVfyezeWfXfYfZf0fSd1f2f'cTd(cUd3fVdyDAezDADBDCDDDEDFDGDHDIDBeJDBhKDLDMD4f5fCeNDDeOD6fPD7fQD8f9fEeAiFeCiGeHe!f#f$f%f'f(f)f*fRD+fSD,fTD-fUD.f/fIeVDJeWD:f;f=fXDKeYDZD0D1D2D3D4D5D6DLe7DCh8D9D!D?fMeNe@f[f]f^f_fWd`f{f)cXd*cYd|fZd#DOe$D%D'D(D)D*D+D,D-D.DPe/DDh:D;D=D}f~fQe?DRe@D[D]Dag^Dbg_D`DXc{D3bcg}Ddg~DegaEfgbEggcEhgdEigeEjgfEkggElghEmgiEngjEkEoglEpgmEnEYcZc0c1c'd(d)d`c{c|c*d+d,d-d4c5c}c.d~c/d:d;d=dad?dbdcddded@d[d]d^d_d`d{d|d}dfdgdhde ~drEsEtE5buEvE2c3caebecedeeefege|J}J~JaKhebKiejecK
73 with gil: 1e
74 _raise_nvrtc_error(prog, err) 1e
77cdef int _raise_nvrtc_error(cynvrtc.nvrtcProgram prog, cynvrtc.nvrtcResult err) except -1:
78 """Build error message with program log and raise NVRTCError."""
79 cdef const char* err_str = cynvrtc.nvrtcGetErrorString(err) 1e
80 cdef size_t logsize = 0 1e
81 if prog != NULL: 1e
82 cynvrtc.nvrtcGetProgramLogSize(prog, &logsize) 1e
83 cdef bytes log_bytes
84 cdef str log_str = "" 1e
85 if logsize > 1 and prog != NULL: 1e
86 log_bytes = b" " * logsize 1e
87 if cynvrtc.nvrtcGetProgramLog(prog, <char*>log_bytes) == cynvrtc.nvrtcResult.NVRTC_SUCCESS: 1e
88 log_str = log_bytes.decode("utf-8", errors="backslashreplace") 1e
89 err_msg = f"{err}: {err_str.decode()}" if err_str != NULL else f"NVRTC error {err}" 1e
90 if log_str: 1e
91 err_msg += f", compilation log:\n\n{log_str}" 1e
92 raise NVRTCError(err_msg) 1e
95cdef int HANDLE_RETURN_NVVM(cynvvm.nvvmProgram prog, cynvvm.nvvmResult err) except?-1 nogil:
96 """Handle NVVM result codes, raising nvvmError with program log on failure."""
97 if err == cynvvm.nvvmResult.NVVM_SUCCESS: 2|DdKeKoE4bpEqEfKgKhKiKjKkKlKmKnKoKpKqKrKsKJiKi
98 return 0 2|DdKeKoE4bpEqEfKgKhKiKjKkKlKmKnKoKpKqKrKsKJiKi
99 with gil: 24b
100 _raise_nvvm_error(prog, err) 24b
103cdef int _raise_nvvm_error(cynvvm.nvvmProgram prog, cynvvm.nvvmResult err) except -1:
104 """Raise nvvmError annotated with the program log."""
105 cdef size_t logsize = 0 24b
106 if prog != NULL: 24b
107 cynvvm.nvvmGetProgramLogSize(prog, &logsize) 24b
108 cdef bytes log_bytes
109 cdef str log_str = "" 24b
110 if logsize > 1 and prog != NULL: 24b
111 log_bytes = b" " * logsize 24b
112 if cynvvm.nvvmGetProgramLog(prog, <char*>log_bytes) == cynvvm.nvvmResult.NVVM_SUCCESS: 24b
113 log_str = log_bytes.decode("utf-8", errors="backslashreplace") 24b
114 cdef object exc = nvvmError(err) 24b
115 if log_str: 2a 4b
116 exc.args = (exc.args[0] + f"\nNVVM program log: {log_str}", *exc.args[1:]) 24b
117 raise exc 24b
120cdef int HANDLE_RETURN_NVJITLINK(
121 cynvjitlink.nvJitLinkHandle handle, cynvjitlink.nvJitLinkResult err) except?-1 nogil:
122 """Handle nvJitLink result codes, raising nvJitLinkError with error log on failure."""
123 if err == cynvjitlink.nvJitLinkResult.NVJITLINK_SUCCESS: 2b tKuKvKwKxKyKzKAKBKCKDK7E8EEKFKGKHKIKJKKKLKMKNKOKPKQKhfRKSKTKUKVK3b5bxEWKXKYKZK0K1K2K3K4K
124 return 0 2b tKuKvKwKxKyKzKAKBKCKDK7E8EEKFKGKHKIKJKKKLKMKNKOKPKQKRKSKTKUKVK3b5bxEWKXKYKZK0K1K2K3K4K
125 with gil: 2b hf
126 _raise_nvjitlink_error(handle, err) 2b hf
129cdef int _raise_nvjitlink_error(
130 cynvjitlink.nvJitLinkHandle handle, cynvjitlink.nvJitLinkResult err) except -1:
131 """Raise nvJitLinkError annotated with the error log."""
132 cdef size_t logsize = 0 2b hf
133 if handle != NULL: 2a b hf
134 cynvjitlink.nvJitLinkGetErrorLogSize(handle, &logsize) 1b
135 cdef bytes log_bytes
136 cdef str log_str = "" 2b hf
137 if logsize > 1 and handle != NULL: 2b hf
138 log_bytes = b" " * logsize 1b
139 if cynvjitlink.nvJitLinkGetErrorLog(handle, <char*>log_bytes) == \ 1b
140 cynvjitlink.nvJitLinkResult.NVJITLINK_SUCCESS: 1b
141 log_str = log_bytes.decode("utf-8", errors="backslashreplace") 1b
142 cdef object exc = nvJitLinkError(err) 2b hf
143 if log_str: 2b hf
144 exc.args = (exc.args[0] + f"\nnvJitLink error log: {log_str}", *exc.args[1:]) 1b
145 raise exc 2b hf
148cdef object _RUNTIME_SUCCESS = runtime.cudaError_t.cudaSuccess
149cdef object _NVRTC_SUCCESS = nvrtc.nvrtcResult.NVRTC_SUCCESS
152cpdef inline int _check_driver_error(cydriver.CUresult error) except?-1 nogil:
153 if error == cydriver.CUresult.CUDA_SUCCESS: 2a 0dLi1dMiifNi$ Oi% PijfQi' Ri( Si) Ti* Ui+ Vi, Wi- Xi. Yi/ Zi: 0i; 1i= 2i? 3i@ 4i[ 5i] 6i^ 7i_ 8i` 9i{ !i| #i} $i~ %iab'ibb(icb)idb*ieb+i,i-i.i/i[g:i;i=i?i@i[i]i^i_i`i{i|i}i~iajSebjcjdjTeejUefjVegjWehjBbijXejjYekjZeljmjnjCbojDbpjEbqjFbrjGbsjGctjHcujIcvjJcwjKcxjLcyjMczjNcAjBj'hCj(hDj)hEjwgFjxgGjygHjIjJjKjLjMjNjOjPjQjRjSjTjUjVjWjXjYjZj0j1j2j3j4j5j6j7j8j9j!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~kalblcldlelflgl]ghliljlklllmlnlol^gplqlrlslZbtlul_gvl0ewlxlylzlAlBlClDlElFlGlHlIlJlKlLl1eMl2eNlOlPlQlRl+cSl,cTlUlVli Wlj XlYlZlHb0lIb1lJb2l3l3e4l`g5lKb6lLb7lMb8lNb9lOb!l#l$l%l'lc (l{g)l4e*l|g+l-c.c/c2d3d4df k g h 8 fb9 ! gbhbib5d6d7d8d9d!djb,lkb-lPb.llb/lmb:lnb;lob=lpb?lqb@l[l~e]l^l}g~g*hahbh+hchdh,heh-h_lzg`laf.h{l5e|lbfcf/hdf}lef:h~lAgamfh;hghhh=hffbmgfcmdmemfmgmhm6eim# l Ocqgrg:c#d;c=c.bRbBg?c@csgtgSbTbUbVbWbXbjmkmlmmmnmompmyEqmrmsmtmumvmwmxmymzmAmBmCmDmEmmhFmGmHmnhImohCcJmKmLmMmNmOmPmQmDcRmSmTmUmVmWmXmYmZm0mphqhodrh$c1m%c2mshthuh3mvhwhxh4m5m6m7m8m9m!m#m$m%m'm(m)m*m+m,m-m.m/mrb/b:b:m;m=m?mqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdoepe@m[m]m^m_m`m{m|m}m~manbncndnenfngnhninjnknlnmnnnonpnqnrnsntnunvnwnxnynznHgtcucIgJgKgLgMgNgOgPgyhzhke6csbwctb7cxcub0b(e)e*em $bn o p q r s t u v %bw x y z 6bA B 7b'ble8cvbycwb9czcxb1b+e,e-eC (bD E F G H I J K L )bM N O P 8bQ R 9b*bme!cybAczb#cBcAb2b.e/e:eS +bT U V W X Y Z 0 1 ,b2 3 4 5 !b6 7 #b-bQgAnBn9eCnDnEnFnGnHnInJnKnLnMnNnOnPnQnRnSnTnUnVnWnXnYnZn0n1n2n3n4n5n6n7n8n9n!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~naobocodoRgSg$eeofogohoiojoQcRcQbd Sckolomo$dnooopoqorosoFctouoYbvowoxoyozoAoBoCoDoEoFoGoHoIoJoKoLoMoNoOoPoQoRoSoToUoVoWoXoYoZo0o1o2o3o4o5o6o7o8o9o!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|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~AaBbBcBdBeBfBgBhBiBjBkBlBmBnBoBpBqBrBsBtBuBvBwBxByBzBABBBCBDBEBFBGBHBIBJBKBLBMBNBOBPBQBRBSBTBUBVBWBXBYBZB0B1B2B3B4B5B
154 return 0 2a 0dLi1dMiifNi$ Oi% PijfQi' Ri( Si) Ti* Ui+ Vi, Wi- Xi. Yi/ Zi: 0i; 1i= 2i? 3i@ 4i[ 5i] 6i^ 7i_ 8i` 9i{ !i| #i} $i~ %iab'ibb(icb)idb*ieb+i,i-i.i/i:i;i=i?i@i[i]i^i_i`i{i|i}i~iajSebjcjdjTeejUefjVegjWehjBbijXejjYekjZeljmjnjCbojDbpjEbqjFbrjGbsjGctjHcujIcvjJcwjKcxjLcyjMczjNcAjBj'hCj(hDj)hEjwgFjxgGjygHjIjJjKjLjMjNjOjPjQjRjSjTjUjVjWjXjYjZj0j1j2j3j4j5j6j7j8j9j!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~kalblcldlelflglhliljlklllmlnlolplqlrlslZbtlulvl0ewlxlylzlAlBlClDlElFlGlHlIlJlKlLl1eMl2eNlOlPlQlRl+cSl,cTlUlVli Wlj XlYlZlHb0lIb1lJb2l3l3e4l`g5lKb6lLb7lMb8lNb9lOb!l#l$l%l'lc (l)l4e*l+l-c.c/c2d3d4df k g h 8 fb9 ! gbhbib5d6d7d8d9d!djb,lkb-lPb.llb/lmb:lnb;lob=lpb?lqb@l[l~e]l^l}g~g*hahbh+hchdh,heh-h_lzg`laf.h{l5e|lbfcf/hdf}lef:h~lAgamfh;hghhh=hffbmgfcmdmemfmgmhm6eim# l Ocqgrg:c#d;c=c.bRbBg?c@csgtgSbTbUbVbWbXbjmkmlmmmnmompmqmrmsmtmumvmwmxmymzmAmBmCmDmEmmhFmGmHmnhImohCcJmKmLmMmNmOmPmQmDcRmSmTmUmVmWmXmYmZm0mphqhodrh$c1m%c2mshthuh3mvhwhxh4m5m6m7m8m9m!m#m$m%m'm(m)m*m+m,m-m.m/mrb/b:b:m;m=m?mqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdoepe@m[m]m^m_m`m{m|m}m~manbncndnenfngnhninjnknlnmnnnonpnqnrnsntnunvnwnxnynznHgtcucIgJgKgLgMgNgOgPgyhzhke6csbwctb7cxcub0b(e)e*em $bn o p q r s t u v %bw x y z 6bA B 7b'ble8cvbycwb9czcxb1b+e,e-eC (bD E F G H I J K L )bM N O P 8bQ R 9b*bme!cybAczb#cBcAb2b.e/e:eS +bT U V W X Y Z 0 1 ,b2 3 4 5 !b6 7 #b-bQgAnBn9eCnDnEnFnGnHnInJnKnLnMnNnOnPnQnRnSnTnUnVnWnXnYnZn0n1n2n3n4n5n6n7n8n9n!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~naobocodoRgSg$eeofogohoiojoQcRcQbd Sckolomo$dnooopoqorosoFctouoYbvowoxoyozoAoBoCoDoEoFoGoHoIoJoKoLoMoNoOoPoQoRoSoToUoVoWoXoYoZo0o1o2o3o4o5o6o7o8o9o!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|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~AaBbBcBdBeBfBgBhBiBjBkBlBmBnBoBpBqBrBsBtBuBvBwBxByBzBABBBCBDBEBFBGBHBIBJBKBLBMBNBOBPBQBRBSBTBUBVBWBXBYBZB0B1B2B3B4B5B
155 cdef const char* name
156 name_err = cydriver.cuGetErrorName(error, &name) 2[g]g^g_g{g4e|gyEqmQbd Yb
157 if name_err != cydriver.CUresult.CUDA_SUCCESS: 2[g]g^g_g{g4e|gyEqmQbd Yb
158 raise CUDAError(f"UNEXPECTED ERROR CODE: {error}")
159 with gil: 2[g]g^g_g{g4e|gyEqmQbd Yb
160 # TODO: consider lower this to Cython
161 expl = DRIVER_CU_RESULT_EXPLANATIONS.get(int(error)) 2[g]g^g_g{g4e|gyEqmQbd Yb
162 if expl is not None: 2[g]g^g_g{g4e|gyEqmQbd Yb
163 raise CUDAError(f"{name.decode()}: {expl}") 2[g]g^g_g{g4e|gyEqmQbd Yb
164 cdef const char* desc
165 desc_err = cydriver.cuGetErrorString(error, &desc)
166 if desc_err != cydriver.CUresult.CUDA_SUCCESS:
167 raise CUDAError(f"{name.decode()}")
168 raise CUDAError(f"{name.decode()}: {desc.decode()}")
171cpdef inline int _check_runtime_error(error) except?-1:
172 if error == _RUNTIME_SUCCESS: 2SbTbUbVbWbXbCg5KEhDg7e8eEgFgGg/b:btcucsbtbub0bm n o p q r s t u v w x y z A B vbwbxb1bC D E F G H I J K L M N O P Q R ybzbAb2bS T U V W X Y Z 0 1 2 3 4 5 6 7 [cPc!e]c#e^c
173 return 0 2SbTbUbVbWbXbCgEhDg7e8eEgFgGg/b:btcucsbtbub0bm n o p q r s t u v w x y z A B vbwbxb1bC D E F G H I J K L M N O P Q R ybzbAb2bS T U V W X Y Z 0 1 2 3 4 5 6 7 [cPc!e]c#e^c
174 # `_check_error()` reaches this path only for `runtime.cudaError_t` values.
175 # Use the enum name directly because Windows hybrid cudart can lag that table.
176 name = error.name 25KEh
177 expl = RUNTIME_CUDA_ERROR_EXPLANATIONS.get(int(error)) 25KEh
178 if expl is not None: 25KEh
179 raise CUDAError(f"{name}: {expl}") 25KEh
180 desc_err, desc = runtime.cudaGetErrorString(error) 2Eh
181 if desc_err != _RUNTIME_SUCCESS: 2Eh
182 raise CUDAError(f"{name}")
183 desc = desc.decode() 2Eh
184 raise CUDAError(f"{name}: {desc}") 2Eh
187cpdef inline int _check_nvrtc_error(error, handle=None) except?-1:
188 if error == _NVRTC_SUCCESS: 2a 6B6Kb ;b=b?b@b[b]b^b_b`b{b|b}b~bacbcccdcecfcgchcicjckclcmcncocpcqcrcscvcTcUcFhVcGhWcXc3bYcZc0c1c'd(d)d`c{c|c*d+d,d-d4c5c}c.d~c/d:d;d=dad?dbdcddded@d[d]d^d_d`d{d|d}dfdgdhde ~d5b2c3cTgUgVgWgXgYgHhIhJhKhLhMhNhOhPhQhRhShThUhVhWhXhYhZh0h1h2h3h4h5h6h7haebecedeeefegeheieje
189 return 0 2a 6Kb ;b=b?b@b[b]b^b_b`b{b|b}b~bacbcccdcecfcgchcicjckclcmcncocpcqcrcscvcTcUcFhVcGhWcXc3bYcZc0c1c'd(d)d`c{c|c*d+d,d-d4c5c}c.d~c/d:d;d=dad?dbdcddded@d[d]d^d_d`d{d|d}dfdgdhde ~d5b2c3cTgUgVgWgXgYgHhIhJhKhLhMhNhOhPhQhRhShThUhVhWhXhYhZh0h1h2h3h4h5h6h7haebecedeeefegeheieje
190 err = f"{error}: {nvrtc.nvrtcGetErrorString(error)[1].decode()}" 26B6K
191 if handle is not None: 26B6K
192 _, logsize = nvrtc.nvrtcGetProgramLogSize(handle) 26B
193 log = b" " * logsize 26B
194 _ = nvrtc.nvrtcGetProgramLog(handle, log) 26B
195 err += f", compilation log:\n\n{log.decode('utf-8', errors='backslashreplace')}" 26B
196 raise NVRTCError(err) 26B6K
199cdef inline int _check_error(error, handle=None) except?-1:
200 if isinstance(error, driver.CUresult): 2a 0dLi1dMiifNi$ Oi% PijfQi' Ri( Si) Ti* Ui+ Vi, Wi- Xi. Yi/ Zi: 0i; 1i= 2i? 3i@ 4i[ 5i] 6i^ 7i_ 8i` 9i{ !i| #i} $i~ %iab'ibb(icb)idb*ieb+i,i-i.i/i:i;i=i?i@i[i]i^i_i`i{i|i}i~iajSebjcjdjTeejUefjVegjWehjBbijXejjYekjZeljmjnjCbojDbpjEbqjFbrjGbsjGctjHcujIcvjJcwjKcxjLcyjMczjNcAjBj'hCj(hDj)hEjwgFjxgGjygHjIjJjKjLjMjNjOjPjQjRjSjTjUjVjWjXjYjZj0j1j2j3j4j5j6j7j8j9j!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~kalblcldlelflglhliljlklllmlnlolplqlrlslZbtlulvl0ewlxlylzlAlBlClDlElFlGlHlIlJlKlLl1eMl2eNlOlPlQlRl+cSl,cTlUlVli Wlj XlYlZlHb0lIb1lJb2l3l3e4l`g5lKb6lLb7lMb8lNb9lOb!l#l$l%l'lc (l)l*l+l-c.c/c2d3d4df k g h 8 fb9 ! gbhbib5d6d7d8d9d!djb,lkb-lPb.llb/lmb:lnb;lob=lpb?lqb@l[l~e]l^l}g~g*hahbh+hchdh,heh-h_lzg`laf.h{l5e|lbfcf/hdf}lef:h~lAgamfh;hghhh=hffbmgfcmdmemfmgmhm6eim# l Ocqgrg:c#d;c=c.bRbBg?c@csgtgSbTbUbVbWbXbCgjmkmlmmmnmompmrmDg7esm8etmumvmwmxmymzmAmEgBmCmDmEmmhFmGmHmnhFgGgImohJmKmLmMmNmOmPmQmRmSmTmUmVmWmXmYmZm0mphqhrh1m2mshthuh3mvhwhxh4m5m6m7m8m9m!m#m$m%m'm(m)m*m+m,m-m.m/mrb/b:b:m;m=m?mb ;b@m=b[m?b]m@b^m[b_m]b`m^b{m_b|m`b}m{b~m|ban}bbn~bcnacdnbcenccfndcgnechnfcingcjnhckniclnjcmnkcnnlconmcpnncqnocrnpcsnqctnrcunvnscwnxnynznHgtcucIgJgKgLgMgNgOgPgyhzhke6csbwctb7cxcub0b(e)e*em $bn o p q r s t u v %bw x y z 6bA B 7b'ble8cvbycwb9czcxb1b+e,e-eC (bD E F G H I J K L )bM N O P 8bQ R 9b*bme!cybAczb#cBcAb2b.e/e:eS +bT U V W X Y Z 0 1 ,b2 3 4 5 !b6 7 #b-bQg[cAnBn9eCnDnEnFnGnHnInPcJn!eKn]cLn#eMnNnOnPnQnRnSnTnUnVnWnXnYnZn0n1n2n3n4n5n6n7n8n9n!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~naobocodoRgSg$eeo^cfogohoiojod vckolomo$dnooopoqorosotouoYbvowoxoyozoAoBoCoDoEoFoGoHoTcIoUcJoKoLoFhVcMoGhWcNoOoPoQoRoSoToUoVoWoXoYoZo0o1o2o3o4o5o6o7o8o9o!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|w}w~waxbxcxdxXcexfxgx3bhxixjxkxlxmxnxoxpxqxrxsxtxuxvxwxxxyxzxAxBxCxDxExFxGxHxIxJxKxLxMxNxOxPxQxRxSxTxUxVxWxXxYxZx0x1x2x3x4x5x6x7x8x9x!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~yazbzczdzYcezZcfz0cgz1chz'diz(djz)dkz`clz{cmz|cnz*doz+dpz,dqz-drz4csz5ctz}cuz.dvz~cwz/dxz:dyz;dzz=dAzadBz?dCzbdDzcdEzddFzedGz@dHz[dIz]dJz^dKz_dLz`dMz{dNz|dOz}dPzfdQzgdRzhdSze Tz~dUzVzWzXzYzZz0z1z2z3z4z5z6z7z8z5b9z!z#z2c$z%z'z(z)z*z+z,z3c-zTgUgVgWgXgYgHhIhJhKhLhMhNhOhPhQhRhShThUh.zVhWhXhYhZh0h1h2h3h4h/z:z;z=z5h6h7haebecedeee?zfegeheieje@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~AaBbBcBdBeBfBgBhBiBjBkBlBmBnBoBpBqBrBsBtBuBvBwBxByBzBABBBCBDBEBFBGBHBIBJBKBLBMBNBOBPBQBRBSBTBUBVBWBXBYBZB0B1B2B3B4B5B
201 return _check_driver_error(error) 2a 0dLi1dMiifNi$ Oi% PijfQi' Ri( Si) Ti* Ui+ Vi, Wi- Xi. Yi/ Zi: 0i; 1i= 2i? 3i@ 4i[ 5i] 6i^ 7i_ 8i` 9i{ !i| #i} $i~ %iab'ibb(icb)idb*ieb+i,i-i.i/i:i;i=i?i@i[i]i^i_i`i{i|i}i~iajSebjcjdjTeejUefjVegjWehjBbijXejjYekjZeljmjnjCbojDbpjEbqjFbrjGbsjGctjHcujIcvjJcwjKcxjLcyjMczjNcAjBj'hCj(hDj)hEjwgFjxgGjygHjIjJjKjLjMjNjOjPjQjRjSjTjUjVjWjXjYjZj0j1j2j3j4j5j6j7j8j9j!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~kalblcldlelflglhliljlklllmlnlolplqlrlslZbtlulvl0ewlxlylzlAlBlClDlElFlGlHlIlJlKlLl1eMl2eNlOlPlQlRl+cSl,cTlUlVli Wlj XlYlZlHb0lIb1lJb2l3l3e4l`g5lKb6lLb7lMb8lNb9lOb!l#l$l%l'lc (l)l*l+l-c.c/c2d3d4df k g h 8 fb9 ! gbhbib5d6d7d8d9d!djb,lkb-lPb.llb/lmb:lnb;lob=lpb?lqb@l[l~e]l^l}g~g*hahbh+hchdh,heh-h_lzg`laf.h{l5e|lbfcf/hdf}lef:h~lAgamfh;hghhh=hffbmgfcmdmemfmgmhm6eim# l Ocqgrg:c#d;c=c.bRbBg?c@csgtgSbTbUbVbWbXbjmkmlmmmnmompmrmsmtmumvmwmxmymzmAmBmCmDmEmmhFmGmHmnhImohJmKmLmMmNmOmPmQmRmSmTmUmVmWmXmYmZm0mphqhrh1m2mshthuh3mvhwhxh4m5m6m7m8m9m!m#m$m%m'm(m)m*m+m,m-m.m/mrb/b:b:m;m=m?m@m[m]m^m_m`m{m|m}m~manbncndnenfngnhninjnknlnmnnnonpnqnrnsntnunvnwnxnynznHgtcucIgJgKgLgMgNgOgPgyhzhke6csbwctb7cxcub0b(e)e*em $bn o p q r s t u v %bw x y z 6bA B 7b'ble8cvbycwb9czcxb1b+e,e-eC (bD E F G H I J K L )bM N O P 8bQ R 9b*bme!cybAczb#cBcAb2b.e/e:eS +bT U V W X Y Z 0 1 ,b2 3 4 5 !b6 7 #b-bQgAnBn9eCnDnEnFnGnHnInJnKnLnMnNnOnPnQnRnSnTnUnVnWnXnYnZn0n1n2n3n4n5n6n7n8n9n!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~naobocodoRgSg$eeofogohoiojod kolomo$dnooopoqorosotouoYbvowoxoyozoAoBoCoDoEoFoGoHoIoJoKoLoMoNoOoPoQoRoSoToUoVoWoXoYoZo0o1o2o3o4o5o6o7o8o9o!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|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~AaBbBcBdBeBfBgBhBiBjBkBlBmBnBoBpBqBrBsBtBuBvBwBxByBzBABBBCBDBEBFBGBHBIBJBKBLBMBNBOBPBQBRBSBTBUBVBWBXBYBZB0B1B2B3B4B5B
202 elif isinstance(error, runtime.cudaError_t): 2a SbTbUbVbWbXbCgDg7e8eEgFgGg/b:bb ;b=b?b@b[b]b^b_b`b{b|b}b~bacbcccdcecfcgchcicjckclcmcncocpcqcrcsctcucsbtbub0bm n o p q r s t u v w x y z A B vbwbxb1bC D E F G H I J K L M N O P Q R ybzbAb2bS T U V W X Y Z 0 1 2 3 4 5 6 7 [cPc!e]c#e^cvcTcUcFhVcGhWcXc3bYcZc0c1c'd(d)d`c{c|c*d+d,d-d4c5c}c.d~c/d:d;d=dad?dbdcddded@d[d]d^d_d`d{d|d}dfdgdhde ~d5b2c3cTgUgVgWgXgYgHhIhJhKhLhMhNhOhPhQhRhShThUhVhWhXhYhZh0h1h2h3h4h5h6h7haebecedeeefegeheieje
203 return _check_runtime_error(error) 2SbTbUbVbWbXbCgDg7e8eEgFgGg/b:btcucsbtbub0bm n o p q r s t u v w x y z A B vbwbxb1bC D E F G H I J K L M N O P Q R ybzbAb2bS T U V W X Y Z 0 1 2 3 4 5 6 7 [cPc!e]c#e^c
204 elif isinstance(error, nvrtc.nvrtcResult): 2a b ;b=b?b@b[b]b^b_b`b{b|b}b~bacbcccdcecfcgchcicjckclcmcncocpcqcrcscvcTcUcFhVcGhWcXc3bYcZc0c1c'd(d)d`c{c|c*d+d,d-d4c5c}c.d~c/d:d;d=dad?dbdcddded@d[d]d^d_d`d{d|d}dfdgdhde ~d5b2c3cTgUgVgWgXgYgHhIhJhKhLhMhNhOhPhQhRhShThUhVhWhXhYhZh0h1h2h3h4h5h6h7haebecedeeefegeheieje
205 return _check_nvrtc_error(error, handle=handle) 2a b ;b=b?b@b[b]b^b_b`b{b|b}b~bacbcccdcecfcgchcicjckclcmcncocpcqcrcscvcTcUcFhVcGhWcXc3bYcZc0c1c'd(d)d`c{c|c*d+d,d-d4c5c}c.d~c/d:d;d=dad?dbdcddded@d[d]d^d_d`d{d|d}dfdgdhde ~d5b2c3cTgUgVgWgXgYgHhIhJhKhLhMhNhOhPhQhRhShThUhVhWhXhYhZh0h1h2h3h4h5h6h7haebecedeeefegeheieje
206 else:
207 raise RuntimeError(f"Unknown error type: {error}")
210def handle_return(tuple result, handle=None):
211 _check_error(result[0], handle=handle) 2a 0dLi1dMiifNi$ Oi% PijfQi' Ri( Si) Ti* Ui+ Vi, Wi- Xi. Yi/ Zi: 0i; 1i= 2i? 3i@ 4i[ 5i] 6i^ 7i_ 8i` 9i{ !i| #i} $i~ %iab'ibb(icb)idb*ieb+i,i-i.i/i:i;i=i?i@i[i]i^i_i`i{i|i}i~iajSebjcjdjTeejUefjVegjWehjBbijXejjYekjZeljmjnjCbojDbpjEbqjFbrjGbsjGctjHcujIcvjJcwjKcxjLcyjMczjNcAjBj'hCj(hDj)hEjwgFjxgGjygHjIjJjKjLjMjNjOjPjQjRjSjTjUjVjWjXjYjZj0j1j2j3j4j5j6j7j8j9j!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~kalblcldlelflglhliljlklllmlnlolplqlrlslZbtlulvl0ewlxlylzlAlBlClDlElFlGlHlIlJlKlLl1eMl2eNlOlPlQlRl+cSl,cTlUlVli Wlj XlYlZlHb0lIb1lJb2l3l3e4l`g5lKb6lLb7lMb8lNb9lOb!l#l$l%l'lc (l)l*l+l-c.c/c2d3d4df k g h 8 fb9 ! gbhbib5d6d7d8d9d!djb,lkb-lPb.llb/lmb:lnb;lob=lpb?lqb@l[l~e]l^l}g~g*hahbh+hchdh,heh-h_lzg`laf.h{l5e|lbfcf/hdf}lef:h~lAgamfh;hghhh=hffbmgfcmdmemfmgmhm6eim# l Ocqgrg:c#d;c=c.bRbBg?c@csgtgSbTbUbVbWbXbCgjmkmlmmmnmompmrmDg7esm8etmumvmwmxmymzmAmEgBmCmDmEmmhFmGmHmnhFgGgImohJmKmLmMmNmOmPmQmRmSmTmUmVmWmXmYmZm0mphqhrh1m2mshthuh3mvhwhxh4m5m6m7m8m9m!m#m$m%m'm(m)m*m+m,m-m.m/mrb/b:b:m;m=m?mb ;b@m=b[m?b]m@b^m[b_m]b`m^b{m_b|m`b}m{b~m|ban}bbn~bcnacdnbcenccfndcgnechnfcingcjnhckniclnjcmnkcnnlconmcpnncqnocrnpcsnqctnrcunvnscwnxnynznHgtcucIgJgKgLgMgNgOgPgyhzhke6csbwctb7cxcub0b(e)e*em $bn o p q r s t u v %bw x y z 6bA B 7b'ble8cvbycwb9czcxb1b+e,e-eC (bD E F G H I J K L )bM N O P 8bQ R 9b*bme!cybAczb#cBcAb2b.e/e:eS +bT U V W X Y Z 0 1 ,b2 3 4 5 !b6 7 #b-bQg[cAnBn9eCnDnEnFnGnHnInPcJn!eKn]cLn#eMnNnOnPnQnRnSnTnUnVnWnXnYnZn0n1n2n3n4n5n6n7n8n9n!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~naobocodoRgSg$eeo^cfogohoiojod vckolomo$dnooopoqorosotouoYbvowoxoyozoAoBoCoDoEoFoGoHoTcIoUcJoKoLoFhVcMoGhWcNoOoPoQoRoSoToUoVoWoXoYoZo0o1o2o3o4o5o6o7o8o9o!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|w}w~waxbxcxdxXcexfxgx3bhxixjxkxlxmxnxoxpxqxrxsxtxuxvxwxxxyxzxAxBxCxDxExFxGxHxIxJxKxLxMxNxOxPxQxRxSxTxUxVxWxXxYxZx0x1x2x3x4x5x6x7x8x9x!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~yazbzczdzYcezZcfz0cgz1chz'diz(djz)dkz`clz{cmz|cnz*doz+dpz,dqz-drz4csz5ctz}cuz.dvz~cwz/dxz:dyz;dzz=dAzadBz?dCzbdDzcdEzddFzedGz@dHz[dIz]dJz^dKz_dLz`dMz{dNz|dOz}dPzfdQzgdRzhdSze Tz~dUzVzWzXzYzZz0z1z2z3z4z5z6z7z8z5b9z!z#z2c$z%z'z(z)z*z+z,z3c-zTgUgVgWgXgYgHhIhJhKhLhMhNhOhPhQhRhShThUh.zVhWhXhYhZh0h1h2h3h4h/z:z;z=z5h6h7haebecedeee?zfegeheieje@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~AaBbBcBdBeBfBgBhBiBjBkBlBmBnBoBpBqBrBsBtBuBvBwBxByBzBABBBCBDBEBFBGBHBIBJBKBLBMBNBOBPBQBRBSBTBUBVBWBXBYBZB0B1B2B3B4B5B
212 cdef int out_len = len(result) 2a 0dLi1dMiifNi$ Oi% PijfQi' Ri( Si) Ti* Ui+ Vi, Wi- Xi. Yi/ Zi: 0i; 1i= 2i? 3i@ 4i[ 5i] 6i^ 7i_ 8i` 9i{ !i| #i} $i~ %iab'ibb(icb)idb*ieb+i,i-i.i/i:i;i=i?i@i[i]i^i_i`i{i|i}i~iajSebjcjdjTeejUefjVegjWehjBbijXejjYekjZeljmjnjCbojDbpjEbqjFbrjGbsjGctjHcujIcvjJcwjKcxjLcyjMczjNcAjBj'hCj(hDj)hEjwgFjxgGjygHjIjJjKjLjMjNjOjPjQjRjSjTjUjVjWjXjYjZj0j1j2j3j4j5j6j7j8j9j!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~kalblcldlelflglhliljlklllmlnlolplqlrlslZbtlulvl0ewlxlylzlAlBlClDlElFlGlHlIlJlKlLl1eMl2eNlOlPlQlRl+cSl,cTlUlVli Wlj XlYlZlHb0lIb1lJb2l3l3e4l`g5lKb6lLb7lMb8lNb9lOb!l#l$l%l'lc (l)l*l+l-c.c/c2d3d4df k g h 8 fb9 ! gbhbib5d6d7d8d9d!djb,lkb-lPb.llb/lmb:lnb;lob=lpb?lqb@l[l~e]l^l}g~g*hahbh+hchdh,heh-h_lzg`laf.h{l5e|lbfcf/hdf}lef:h~lAgamfh;hghhh=hffbmgfcmdmemfmgmhm6eim# l Ocqgrg:c#d;c=c.bRbBg?c@csgtgSbTbUbVbWbXbCgjmkmlmmmnmompmrmDg7esm8etmumvmwmxmymzmAmEgBmCmDmEmmhFmGmHmnhFgGgImohJmKmLmMmNmOmPmQmRmSmTmUmVmWmXmYmZm0mphqhrh1m2mshthuh3mvhwhxh4m5m6m7m8m9m!m#m$m%m'm(m)m*m+m,m-m.m/mrb/b:b:m;m=m?mb ;b@m=b[m?b]m@b^m[b_m]b`m^b{m_b|m`b}m{b~m|ban}bbn~bcnacdnbcenccfndcgnechnfcingcjnhckniclnjcmnkcnnlconmcpnncqnocrnpcsnqctnrcunvnscwnxnynznHgtcucIgJgKgLgMgNgOgPgyhzhke6csbwctb7cxcub0b(e)e*em $bn o p q r s t u v %bw x y z 6bA B 7b'ble8cvbycwb9czcxb1b+e,e-eC (bD E F G H I J K L )bM N O P 8bQ R 9b*bme!cybAczb#cBcAb2b.e/e:eS +bT U V W X Y Z 0 1 ,b2 3 4 5 !b6 7 #b-bQg[cAnBn9eCnDnEnFnGnHnInPcJn!eKn]cLn#eMnNnOnPnQnRnSnTnUnVnWnXnYnZn0n1n2n3n4n5n6n7n8n9n!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~naobocodoRgSg$eeo^cfogohoiojod vckolomo$dnooopoqorosotouoYbvowoxoyozoAoBoCoDoEoFoGoHoTcIoUcJoKoLoFhVcMoGhWcNoOoPoQoRoSoToUoVoWoXoYoZo0o1o2o3o4o5o6o7o8o9o!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|w}w~waxbxcxdxXcexfxgx3bhxixjxkxlxmxnxoxpxqxrxsxtxuxvxwxxxyxzxAxBxCxDxExFxGxHxIxJxKxLxMxNxOxPxQxRxSxTxUxVxWxXxYxZx0x1x2x3x4x5x6x7x8x9x!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~yazbzczdzYcezZcfz0cgz1chz'diz(djz)dkz`clz{cmz|cnz*doz+dpz,dqz-drz4csz5ctz}cuz.dvz~cwz/dxz:dyz;dzz=dAzadBz?dCzbdDzcdEzddFzedGz@dHz[dIz]dJz^dKz_dLz`dMz{dNz|dOz}dPzfdQzgdRzhdSze Tz~dUzVzWzXzYzZz0z1z2z3z4z5z6z7z8z5b9z!z#z2c$z%z'z(z)z*z+z,z3c-zTgUgVgWgXgYgHhIhJhKhLhMhNhOhPhQhRhShThUh.zVhWhXhYhZh0h1h2h3h4h/z:z;z=z5h6h7haebecedeee?zfegeheieje@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~AaBbBcBdBeBfBgBhBiBjBkBlBmBnBoBpBqBrBsBtBuBvBwBxByBzBABBBCBDBEBFBGBHBIBJBKBLBMBNBOBPBQBRBSBTBUBVBWBXBYBZB0B1B2B3B4B5B
213 if out_len == 1: 2a 0dLi1dMiifNi$ Oi% PijfQi' Ri( Si) Ti* Ui+ Vi, Wi- Xi. Yi/ Zi: 0i; 1i= 2i? 3i@ 4i[ 5i] 6i^ 7i_ 8i` 9i{ !i| #i} $i~ %iab'ibb(icb)idb*ieb+i,i-i.i/i:i;i=i?i@i[i]i^i_i`i{i|i}i~iajSebjcjdjTeejUefjVegjWehjBbijXejjYekjZeljmjnjCbojDbpjEbqjFbrjGbsjGctjHcujIcvjJcwjKcxjLcyjMczjNcAjBj'hCj(hDj)hEjwgFjxgGjygHjIjJjKjLjMjNjOjPjQjRjSjTjUjVjWjXjYjZj0j1j2j3j4j5j6j7j8j9j!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~kalblcldlelflglhliljlklllmlnlolplqlrlslZbtlulvl0ewlxlylzlAlBlClDlElFlGlHlIlJlKlLl1eMl2eNlOlPlQlRl+cSl,cTlUlVli Wlj XlYlZlHb0lIb1lJb2l3l3e4l`g5lKb6lLb7lMb8lNb9lOb!l#l$l%l'lc (l)l*l+l-c.c/c2d3d4df k g h 8 fb9 ! gbhbib5d6d7d8d9d!djb,lkb-lPb.llb/lmb:lnb;lob=lpb?lqb@l[l~e]l^l}g~g*hahbh+hchdh,heh-h_lzg`laf.h{l5e|lbfcf/hdf}lef:h~lAgamfh;hghhh=hffbmgfcmdmemfmgmhm6eim# l Ocqgrg:c#d;c=c.bRbBg?c@csgtgSbTbUbVbWbXbCgjmkmlmmmnmompmrmDg7esm8etmumvmwmxmymzmAmEgBmCmDmEmmhFmGmHmnhFgGgImohJmKmLmMmNmOmPmQmRmSmTmUmVmWmXmYmZm0mphqhrh1m2mshthuh3mvhwhxh4m5m6m7m8m9m!m#m$m%m'm(m)m*m+m,m-m.m/mrb/b:b:m;m=m?mb ;b@m=b[m?b]m@b^m[b_m]b`m^b{m_b|m`b}m{b~m|ban}bbn~bcnacdnbcenccfndcgnechnfcingcjnhckniclnjcmnkcnnlconmcpnncqnocrnpcsnqctnrcunvnscwnxnynznHgtcucIgJgKgLgMgNgOgPgyhzhke6csbwctb7cxcub0b(e)e*em $bn o p q r s t u v %bw x y z 6bA B 7b'ble8cvbycwb9czcxb1b+e,e-eC (bD E F G H I J K L )bM N O P 8bQ R 9b*bme!cybAczb#cBcAb2b.e/e:eS +bT U V W X Y Z 0 1 ,b2 3 4 5 !b6 7 #b-bQg[cAnBn9eCnDnEnFnGnHnInPcJn!eKn]cLn#eMnNnOnPnQnRnSnTnUnVnWnXnYnZn0n1n2n3n4n5n6n7n8n9n!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~naobocodoRgSg$eeo^cfogohoiojod vckolomo$dnooopoqorosotouoYbvowoxoyozoAoBoCoDoEoFoGoHoTcIoUcJoKoLoFhVcMoGhWcNoOoPoQoRoSoToUoVoWoXoYoZo0o1o2o3o4o5o6o7o8o9o!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|w}w~waxbxcxdxXcexfxgx3bhxixjxkxlxmxnxoxpxqxrxsxtxuxvxwxxxyxzxAxBxCxDxExFxGxHxIxJxKxLxMxNxOxPxQxRxSxTxUxVxWxXxYxZx0x1x2x3x4x5x6x7x8x9x!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~yazbzczdzYcezZcfz0cgz1chz'diz(djz)dkz`clz{cmz|cnz*doz+dpz,dqz-drz4csz5ctz}cuz.dvz~cwz/dxz:dyz;dzz=dAzadBz?dCzbdDzcdEzddFzedGz@dHz[dIz]dJz^dKz_dLz`dMz{dNz|dOz}dPzfdQzgdRzhdSze Tz~dUzVzWzXzYzZz0z1z2z3z4z5z6z7z8z5b9z!z#z2c$z%z'z(z)z*z+z,z3c-zTgUgVgWgXgYgHhIhJhKhLhMhNhOhPhQhRhShThUh.zVhWhXhYhZh0h1h2h3h4h/z:z;z=z5h6h7haebecedeee?zfegeheieje@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~AaBbBcBdBeBfBgBhBiBjBkBlBmBnBoBpBqBrBsBtBuBvBwBxByBzBABBBCBDBEBFBGBHBIBJBKBLBMBNBOBPBQBRBSBTBUBVBWBXBYBZB0B1B2B3B4B5B
214 return 20d1dif$ % jf' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdbebSeTeUeVeWeBbXeYeZeCbDbEbFbGbGcHcIcJcKcLcMcNc'h(h)hwgxgygZb0e1e2e+c,ci j HbIbJb3eKbLbMbNbObc -c.c/c2d3d4df k g h 8 fb9 ! gbhbib5d6d7d8d9d!djbkbPblbmbnbobpbqb*h+h,h-hzg.h/h:hAg;h=h6el Ocqgrg:c;c=c.bBg?c@csgtgSbTbUbVbWbXb7e8erb/b:bHgtcucIgJgKgLgMgNgOgPgyhzhke6csbwctb7cxcub0b(e)e*em $bn o p q r s t u v %bw x y z 6bA B 7b'ble8cvbycwb9czcxb1b+e,e-eC (bD E F G H I J K L )bM N O P 8bQ R 9b*bme!cybAczb#cBcAb2b.e/e:eS +bT U V W X Y Z 0 1 ,b2 3 4 5 !b6 7 #b-bQg[c9ePc!e]c#eRgSg$e^cd 5B
215 elif out_len == 2:
216 return result[1] 2a 0dLi1dMiifNi$ Oi% PijfQi' Ri( Si) Ti* Ui+ Vi, Wi- Xi. Yi/ Zi: 0i; 1i= 2i? 3i@ 4i[ 5i] 6i^ 7i_ 8i` 9i{ !i| #i} $i~ %iab'ibb(icb)idb*ieb+i,i-i.i/i:i;i=i?i@i[i]i^i_i`i{i|i}i~iajSebjcjdjTeejUefjVegjWehjBbijXejjYekjZeljmjnjCbojDbpjEbqjFbrjGbsjGctjHcujIcvjJcwjKcxjLcyjMczjNcAjBj'hCj(hDj)hEjwgFjxgGjygHjIjJjKjLjMjNjOjPjQjRjSjTjUjVjWjXjYjZj0j1j2j3j4j5j6j7j8j9j!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~kalblcldlelflglhliljlklllmlnlolplqlrlslZbtlulvl0ewlxlylzlAlBlClDlElFlGlHlIlJlKlLl1eMl2eNlOlPlQlRl+cSl,cTlUlVli Wlj XlYlZlHb0lIb1lJb2l3l3e4l`g5lKb6lLb7lMb8lNb9lOb!l#l$l%l'lc (l)l*l+l-c.c/c2d3d4df k g h 8 fb9 ! gbhbib5d6d7d8d9d!djb,lkb-lPb.llb/lmb:lnb;lob=lpb?lqb@l[l~e]l^l}g~g*hahbh+hchdh,heh-h_lzg`laf.h{l5e|lbfcf/hdf}lef:h~lAgamfh;hghhh=hffbmgfcmdmemfmgmhm6eim# l :c#d;c=c.bRbBg?c@cSbTbUbVbWbXbCgjmkmlmmmnmompmrmDgsmtmumvmwmxmymzmAmEgBmCmDmEmmhFmGmHmnhFgGgImohJmKmLmMmNmOmPmQmRmSmTmUmVmWmXmYmZm0mphqhrh1m2mshthuh3mvhwhxh4m5m6m7m8m9m!m#m$m%m'm(m)m*m+m,m-m.m/mrb/b:b:m;m=m?m@m[m]m^m_m`m{m|m}m~manbncndnenfngnhninjnknlnmnnnonpnqnrnsntnunvnwnxnynznHgtcucIgJgKgLgMgNgOgPgke6csbwctb7cxcub0b(e)e*em $bn o p q r s t u v %bw x y z 6bA B 7b'ble8cvbycwb9czcxb1b+e,e-eC (bD E F G H I J K L )bM N O P 8bQ R 9b*bme!cybAczb#cBcAb2b.e/e:eS +bT U V W X Y Z 0 1 ,b2 3 4 5 !b6 7 #b-bQgAnBn9eCnDnEnFnGnHnInJnKnLnMnNnOnPnQnRnSnTnUnVnWnXnYnZn0n1n2n3n4n5n6n7n8n9n!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~naobocodoRgSg$eeofogohoiojod kolomo$dnooopoqorosotouoYbvowoxoyozoAoBoCoDoEoFoGoHoIoJoKoLoMoNoOoPoQoRoSoToUoVoWoXoYoZo0o1o2o3o4o5o6o7o8o9o!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|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~AaBbBcBdBeBfBgBhBiBjBkBlBmBnBoBpBqBrBsBtBuBvBwBxByBzBABBBCBDBEBFBGBHBIBJBKBLBMBNBOBPBQBRBSBTBUBVBWBXBYBZB0B1B2B3B4B
217 else:
218 return result[1:] 2a 0d1dif$ % jf' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdbeb-c.c/c2d3d4df g h 8 9 ! gbhbib5d6d7d8d9d!djbkblbmbnbobpbqbb ;b=b?b@b[b]b^b_b`b{b|b}b~bacbcccdcecfcgchcicjckclcmcncocpcqcrcscvcTcUcFhVcGhWcXc3bYcZc0c1c'd(d)d`c{c|c*d+d,d-d4c5c}c.d~c/d:d;d=dad?dbdcddded@d[d]d^d_d`d{d|d}dfdgdhde ~d5b2c3cTgUgVgWgXgYgHhIhJhKhLhMhNhOhPhQhRhShThUhVhWhXhYhZh0h1h2h3h4h5h6h7haebecedeeefegeheieje
221cpdef check_or_create_options(type cls, options, str options_description="", bint keep_none=False):
222 """
223 Create the specified options dataclass from a dictionary of options or None.
224 """
225 if options is None: 20d1dif$ % jf' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdbebSeTeUeVeWeBbXeYeZeCbDbEbFbGbGcHcIcJcKcLcMcNckflfmf%E'Eidnf*E+Ejdof.E/Ekdpf=E?Eldqf]E^Emdrf{E|EndZb0esftfufvf1e2ewfaFbFcF+c,cdFeFi j !BxfHbIbJbyf3eKbLbMbNbObc -c.c/c2d3d4df k g h 8 fb9 ! gbhbib5d6d7d8d9d!djbkbPblbmbnbobpbqbfFgF~ehFiFjFafkFlF5ebfcfdfmFnFefoFpFffqFgfrFsFtFuFvFwFxFyFzFAFBF6eCFDF#BEF$BFFGFHF%BIF'BJF# KFl LFihjh#h$hkhlh%hOcpi9E7KMFNFOFPFQFRFSFTFUFVFWFXFYFZF0F1F2F3F4F5F6F7F8F9F!F#F$F%F'F(F)F*F+F,F:c-F#d.F;c/F=c:F.bRb?c;F@c=F?F@F[F]FSbTbUbVbWbXb]K^K_K`K{K|KzE7e8e}K~K(BCc)BaL^FbL_F`F{F|F}F~FDcaGbGcGdG*B+B,BneeGfGgGhGiGjGkGlGmGnGoGpGqGDEEErGFEodqi$c-B%c.BGEcLHEdLIEJEKEsGLEMENEOEPEQERESE?K@K[KeLfLgL7Brb/b:bpdEcqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdoepeb ;btK=buK?bvK@bwK[bxK]byK^bzK_bAK`bBK{bCK|bDK}b7E~b8EacEKbcFKccGKdcHKecIKfcJKgcKKhcLKicMKjcNKkcOKlcPKmcQKnchfocRKpcSKqcTKrcUK/BVKsctcuctGuGke6csbwctb7cxcub0b(e)e*em $bn o p q r s t u v %bw x y z 6bA B 7b'ble8cvbycwb9czcxb1b+e,e-eC (bD E F G H I J K L )bM N O P 8bQ R 9b*bme!cybAczb#cBcAb2b.e/e:eS +bT U V W X Y Z 0 1 ,b2 3 4 5 !b6 7 #b-bvGwG[c9e:B;BxGyG=B?B@B[BPc!e]c#ezGAGBGCGDGEGFGGGHGIGJGKGLGMGNGOG]BPG^BQG_BRG`BSG{BTG|BUG}BVG~BWGXGYGZG0G1G2G3G4G5G6G7G8G9G!G#G$GTEriUEVEsiWE%G$e'G(G)G^c*G+GhL,G-G.GQcRcQbd ScvcaCbC$dcCdCeCfCgCFchCYbiCjCkClCmCnCoCpCqCrCsCtCuCTcUcvCwCVcWcxCyCzCACBCCCDCECFCGCHCICJCKCLCMCNCOCPCQCRCSCTCUCVCWCXCYCZC0C1C2C3C4C5C6C7C8C9C!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}Cug%e'e_c%d/Gzf:G;G=G~CAfaD@G[GbD]G^GBf_G`G{GcDCfdD}GDf~GaHEfcHqedHeHXEYEfHgHhHiHjHtire;ekHuivilHmHnHseteFfGfHfIfJfKfoHpHLfrHMfsHtHeDNffDvHOfwHxHyHgDPfhDAHQfBHCHRfEHZEueFHGHHHIHJHKHLHMHNHOHiDve=ePHwijDQHRHSHTHUHVHWHXHYHZH0H1H2H3HSf4H5HTf7H8HUfkDwelDmDnDoDpDqDrDsDtDxeuDAhvDwDxDVfyezeWfXfYfZf0fSd1f2f'c?e@eTd(cUd3f!H#HVdxiyiyDAezDADBDCDDDEDFDGDHDIDBeJDBhKDLDMD4f+H,H5f.H0ECe/H:H;H=H?H@H[H]H^H_H`HNDDe]e{HziOD|H}H~HaI6fbIcIdIPD7fQDfI8fgIhI9fjIEekIlI1E2EmInIoIpIqIAiFe^erIBiCisItIuIGeHe!f#f$f%f'f(fvIwI)fyI*fzIAIRD+fSDCI,fDIEIFITD-fUDHI.fIIJI/fLI3EIeMINIOIPIQIRISITIUIVIVDJe_eWIDiWDXIYIZI0I1I2I3I4I5I6I7I8I9I!I:f#I$I;f'I(I=fXDKeYDZD0D1D2D3D4D5D6DLe7DCh8D9D!D?fMeNe@f[f]f^f_fWd`f{f)c`e{eXd*cYd|f*I+IZdEiFi#DOe$D%D'D(D)D*D+D,D-D.DPe/DDh:D;D=D}f=I?I~f[I4EQe]I^I_I`I{I|I}I~IaJbJcJ?DRe}edJGi@DeJfJgJhJiJHiIi[D]DjJagkJlJmJ^Dbg_D`DXc{D|D3boJpJcgqJrJsJ}Ddg~DuJvJegwJxJyJaEfgbEAJBJggCJDJEJcEhgdEGJHJigIJJJKJeEjgfEMJNJkgOJPJQJgElghESJTJmgUJVJWJiEngjEkEYJZJog0J1J2JlEpgmEnE4JdKeKYcZc5E0c1c'd(d)d`c{c|c*d+d,d-d4c5c}c.d~c/d:d;d=dad?dbdcddded@d[d]d^d_d`d{d|d}dfdgdhde ~doE4bpEqEfKgKhKiKjKkKlKmKnKoKpKqKrKsKJiKi5JrEsEtE5buE6J7JvE6ExEWKXK2cYKZK0K1K2K3K4K3caebecedeeefege|J}J~JaKhebKiejecKiLjLkLlLmLnL8JoLpLqLrLsLtLuLvLwLwExL9J!JyL#JzL$J%J'J(J)J*J+J,J-J.J/J:J;J=J?J@J[J]J^J_J`J{J
226 if keep_none: 20d1dif$ % jf' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdbebSeTeUeVeWeBbXeYeZeCbDbEbFbGbGcHcIcJcKcLcMcNc%E'E*E+E.E/E=E?E]E^E{E|E0e1e2e+c,ci j HbIbJb3eKbLbMbNbObc -c.c/c2d3d4df k g h 8 fb9 ! gbhbib5d6d7d8d9d!djbkbPblbmbnbobpbqbfFgF~ehFiFjFafkFlF5ebfcfdfmFnFefoFpFffqFgfrFsFtFuFvFwFxFyFzFAFBF6e# l ihjh#h$hkhlh%hOcpi9E7K:c#d;c=c.bRb?c@cSbTbUbVbWbXb]K^K_K`K{K|K7e8e}K~K(BCc)B^F_F`F{F|F}F~FDcaGbGcGdG*B+B,BneeGfGgGhGiGjGkGlGmGnGoGpGqGrGod$c%cdLsGrb/b:bEctcuctGuGke6csbwctb7cxcub0b(e)e*em $bn o p q r s t u v %bw x y z 6bA B 7b'ble8cvbycwb9czcxb1b+e,e-eC (bD E F G H I J K L )bM N O P 8bQ R 9b*bme!cybAczb#cBcAb2b.e/e:eS +bT U V W X Y Z 0 1 ,b2 3 4 5 !b6 7 #b-bvGwG[c9e:B;BxGyG=B?B@B[BPc!e]c#e]BPG^BQG_BRG`BSG{BTG|BUG}BVG~BWGrisi$e'G(G)G^caCbC$dcCdCeCfCgCFchCYbiCjCkClCmCnCoCpCqCrCsCtCuCVcWcxCyCzCACBCCCDCECFCGCHCICJCKCLCMCNCOCPCQCRCSCTCUCVCWCXCYCZC0C1C2C3C4C5C6C7C8C9C!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_c/G:G;G=G~CaD@G[GbD]G^G_G`G{GcDdD}G~GaHcHqedHeHXEYEfHgHhHiHjHtire;ekHuivilHmHnHseteoHpHrHsHtHeDfDvHwHxHyHgDhDAHBHCHEHZEueFHGHHHIHJHKHLHMHNHOHiDve=ePHwijDQHRHSHTHUHVHWHXHYHZH0H1H2H3H4H5H7H8HkDwelDmDnDoDpDqDrDsDtDxeuDAhvDwDxDyeze?e@eTdUd!H#HxiyiyDAezDADBDCDDDEDFDGDHDIDBeJDBhKDLDMD+H,H.H0ECe/H:H;H=H?H@H[H]H^H_H`HNDDe]e{HziOD|H}H~HaIbIcIdIPDQDfIgIhIjIEekIlI1E2EmInIoIpIqIAiFe^erIBiCisItIuIGeHevIwIyIzIAIRDSDCIDIEIFITDUDHIIIJILI3EIeMINIOIPIQIRISITIUIVIVDJe_eWIDiWDXIYIZI0I1I2I3I4I5I6I7I8I9I!I#I$I'I(IXDKeYDZD0D1D2D3D4D5D6DLe7DCh8D9D!DMeNe`e{eXdYd*I+IEiFi#DOe$D%D'D(D)D*D+D,D-D.DPe/DDh:D;D=D=I?I[I4EQe]I^I_I`I{I|I}I~IaJbJcJ?DRe}edJGi@DeJfJgJHiIi[D]DjJkJlJmJ^D_D`DXc{D|D3boJpJqJrJsJ}D~DuJvJwJxJyJaEbEAJBJCJDJEJcEdEGJHJIJJJKJeEfEMJNJOJPJQJgEhESJTJUJVJWJiEjEkEYJZJ0J1J2JlEmEnE4JZce ~doE4bpEqE5JrEsEtEuE6J7JvExE2ckLlLmLnLpLqLrLsLvLwExLyL#J$J%J'J(J)J*J+J,J-J.J/J:J;J=J?J@J[J]J^J_J`J{J
227 return options 2i j -c.c/cf k g h 8 fb9 ! gFhFiFaflF5ecfmFefpFqFrFtFvFxFzFBF7e8eod$c%ctGuGvGwG9e:B;BxGyG=B?B@B[BPc!e]c#e]BPG^BQG_BRG`BSG{BTG|BUG}BVG~BWGrisi'G(G)GFc/G[G^GcHqedHeHXEYEfHgHhHiHjHtire;ekHuivilHmHnHrHvHZE0EaIjIEekIlI1E2EmInIoIpIqIAiFe^erIBiCisItIuIyICI3E4EjJpJvJBJHJNJTJZJ$J%J'J(J)J*J+J,J-J.J/J:J;J=J?J@J[J]J^J_J`J{J
228 return cls() 20d1dif$ % jf' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdbebSeTeUeVeWeBbXeYeZeCbDbEbFbGbGcHcIcJcKcLcMcNc%E'E*E+E.E/E=E?E]E^E{E|E0e1e2e+c,ci j HbIbJb3eKbLbMbNbObc -c.c/c2d3d4df k g h 8 fb9 ! gbhbib5d6d7d8d9d!djbkbPblbmbnbobpbqbfF~ejFkF5ebfdfnFoFffgfsFuFwFyFAF6e# l ihjh#h$hkhlh%hOcpi9E7K:c#d;c=c.bRb?c@cSbTbUbVbWbXb]K^K_K`K{K|K}K~K(BCc)B^F_F`F{F|F}F~FDcaGbGcGdG*B+B,BneeGfGgGhGiGjGkGlGmGnGoGpGqGrGod$c%cdLsGrb/b:bEctcucke6csbwctb7cxcub0b(e)e*em $bn o p q r s t u v %bw x y z 6bA B 7b'ble8cvbycwb9czcxb1b+e,e-eC (bD E F G H I J K L )bM N O P 8bQ R 9b*bme!cybAczb#cBcAb2b.e/e:eS +bT U V W X Y Z 0 1 ,b2 3 4 5 !b6 7 #b-b[cPc$e^caCbC$dcCdCeCfCgCFchCYbiCjCkClCmCnCoCpCqCrCsCtCuCVcWcxCyCzCACBCCCDCECFCGCHCICJCKCLCMCNCOCPCQCRCSCTCUCVCWCXCYCZC0C1C2C3C4C5C6C7C8C9C!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_c:G;G=G~CaD@GbD]G_G`G{GcDdD}G~GaHXEYEtiviseteoHpHsHtHeDfDwHxHyHgDhDAHBHCHEHZEueFHGHHHIHJHKHLHMHNHOHiDve=ePHwijDQHRHSHTHUHVHWHXHYHZH0H1H2H3H4H5H7H8HkDwelDmDnDoDpDqDrDsDtDxeuDAhvDwDxDyeze?e@eTdUd!H#HxiyiyDAezDADBDCDDDEDFDGDHDIDBeJDBhKDLDMD+H,H.H0ECe/H:H;H=H?H@H[H]H^H_H`HNDDe]e{HziOD|H}H~HbIcIdIPDQDfIgIhI1E2EAiCiGeHevIwIzIAIRDSDDIEIFITDUDHIIIJILI3EIeMINIOIPIQIRISITIUIVIVDJe_eWIDiWDXIYIZI0I1I2I3I4I5I6I7I8I9I!I#I$I'I(IXDKeYDZD0D1D2D3D4D5D6DLe7DCh8D9D!DMeNe`e{eXdYd*I+IEiFi#DOe$D%D'D(D)D*D+D,D-D.DPe/DDh:D;D=D=I?I[I4EQe]I^I_I`I{I|I}I~IaJbJcJ?DRe}edJGi@DeJfJgJHiIi[D]DkJlJmJ^D_D`DXc{D|D3boJqJrJsJ}D~DuJwJxJyJaEbEAJCJDJEJcEdEGJIJJJKJeEfEMJOJPJQJgEhESJUJVJWJiEjEkEYJ0J1J2JlEmEnE4JZce ~doE4bpEqE5JrEsEtEuE6J7JvExE2ckLlLmLnLpLqLrLsLvLwExLyL#J
229 elif isinstance(options, cls): 2$ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdbebBbCbDbEbFbGbkflfmfidnfjdofkdpfldqfmdrfndZbsftfufvfwfaFbFcF+c,cdFeFi j !BxfHbIbJbyfKbLbMbNbObc f k g h 8 fb9 ! gbhbibjbkbPblbmbnbobpbqbCFDF#BEF$BFFGFHF%BIF'BJF# KFl LFihjh#h$hkhlh%hOcpi9E7KMFNFOFPFQFRFSFTFUFVFWFXFYFZF0F1F2F3F4F5F6F7F8F9F!F#F$F%F'F(F)F*F+F,F-F.F/F:F.bRb;F=F?F@F[F]FSbTbUbVbWbXbzE(BCc)BaLbLDc*B+B,BneDEEEFEodqi$c-B%c.BGEcLHEIEJEKELEMENEOEPEQERESE?K@K[KeLfLgL7BrbpdEcqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdoepeb ;btK=buK?bvK@bwK[bxK]byK^bzK_bAK`bBK{bCK|bDK}b7E~b8EacEKbcFKccGKdcHKecIKfcJKgcKKhcLKicMKjcNKkcOKlcPKmcQKnchfocRKpcSKqcTKrcUK/BVKsc[c:B;B=B?B@B[BPc]czGAGBGCGDGEGFGGGHGIGJGKGLGMGNGOG]B^B_B`B{B|B}B~BXGYGZG0G1G2G3G4G5G6G7G8G9G!G#G$GTEriUEVEsiWE%G^c*G+GhL,G-G.GQcRcQbd ScvcFcTcUcvCwCug%e'e_c%dzfAfBfCfDfEfqereseteFfGfHfIfJfKfLfMfNfOfPfQfRfueveSfTfUfwexeVfyezeWfXfYfZf0fSd1f2f'c(c3fVdAeBe4f5fCeDe6f7f8f9fEeFeGeHe!f#f$f%f'f(f)f*f+f,f-f.f/fIeJe:f;f=fKeLe?fMeNe@f[f]f^f_fWd`f{f)c*c|fZdOePe}f~fQeRehJiJHiIiagbg3bcgdgegfggghgigjgkglgmgngogpgdKeKYc5E0c1c'd(d)d`c{c|c*d+d,d-d4c5c}c.d~c/d:d;d=dad?dbdcddded@d[d]d^d_d`d{d|d}dfdgdhdfKgKhKiKjKkKlKmKnKoKpKqKrKsKJiKi5b6ExEWKXKYKZK0K1K2K3K4K3caebecedeeefege|J}J~JaKhebKiejecKiLjL8JoLtLuLwLwE9J!JzL
230 return options 2$ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdbebBbCbDbEbFbGbkflfmfidnfjdofkdpfldqfmdrfndZbsftfufvfwfaFbFcF+c,cdFeFi j !BxfHbIbJbyfKbLbMbNbObc f k g h 8 fb9 ! gbhbibjbkbPblbmbnbobpbqbCFDF#BEF$BFFGFHF%BIF'BJF# KFl LF#h$h%hOc7KMFNFOFPFQFRFSFTFUFVFWFXFYFZF0F1F2F3F4F5F6F7F8F9F!F#F$F%F'F(F)F*F+F,F-F.F/F:F.bRb;F=F?F@F[F]FSbTbUbVbWbXb(BCc)BaLbLDc*B+B,BneDEEEFEodqi$c-B%c.BGEcLHEIEJEKELEMENEOEPEQERESE?K@K[KeLfLgL7BrbpdEcqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdoepeb ;btK=buK?bvK@bwK[bxK]byK^bzK_bAK`bBK{bCK|bDK}b7E~b8EacEKbcFKccGKdcHKecIKfcJKgcKKhcLKicMKjcNKkcOKlcPKmcQKnchfocRKpcSKqcTKrcUK/BVKsc[c:B;B=B?B@B[BPc]czGAGBGCGDGEGFGGGHGIGJGKGLGMGNGOG]B^B_B`B{B|B}B~BXGYGZG0G1G2G3G4G5G6G7G8G9G!G#G$G%G^c*G+GhL,G-G.GQcRcQbd ScvcFcTcUcvCwCug%e'e_c%dzfAfBfCfDfEfqereseteFfGfHfIfJfKfLfMfNfOfPfQfRfueveSfTfUfwexeVfyezeWfXfYfZf0fSd1f2f'c(c3fVdAeBe4f5fCeDe6f7f8f9fEeFeGeHe!f#f$f%f'f(f)f*f+f,f-f.f/fIeJe:f;f=fKeLe?fMeNe@f[f]f^f_fWd`f{f)c*c|fZdOePe}f~fQeRehJiJagbg3bcgdgegfggghgigjgkglgmgngogpgdKeKYc5E0c1c'd(d)d`c{c|c*d+d,d-d4c5c}c.d~c/d:d;d=dad?dbdcddded@d[d]d^d_d`d{d|d}dfdgdhdfKgKhKiKjKkKlKmKnKoKpKqKrKsKJiKi5b6ExEWKXKYKZK0K1K2K3K4K3caebecedeeefege|J}J~JaKhebKiejecKiLjL8JoLtLuLwLwE9J!JzL
231 elif isinstance(options, dict): 2ihjhkhlhpi9EzE7BpdTEriUEVEsiWEHiIi5b
232 return cls(**options) 2ihjhkhlhpi9EpdTEriUEVEsiWEHiIi5b
233 else:
234 raise TypeError( 2zE7B
235 f"The {options_description} must be provided as an object " 2zE7B
236 f"of type {cls.__name__} or as a dict with valid {options_description}. " 2zE7B
237 f"The provided object is '{options}'." 2zE7B
238 )
241def _handle_boolean_option(option: bool) -> str:
242 """
243 Convert a boolean option to a string representation.
244 """
245 return "true" if bool(option) else "false" 2b ;b=b?b@b[b]b^b_b`b{b|b}b~bacbcccdcecfcgchcicjckclcmcncocpcqcrcscvc`c{c|c~cddedfdgdhdW2
248def precondition(checker: Callable[..., None], str what="") -> Callable:
249 """
250 A decorator that adds checks to ensure any preconditions are met.
252 Args:
253 checker: The function to call to check whether the preconditions are met. It has
254 the same signature as the wrapped function with the addition of the keyword argument `what`.
255 what: A string that is passed in to `checker` to provide context information.
257 Returns:
258 Callable: A decorator that creates the wrapping.
259 """
261 def outer(wrapped_function): 2AE
262 """
263 A decorator that actually wraps the function for checking preconditions.
264 """
266 @functools.wraps(wrapped_function) 2AE
267 def inner(*args, **kwargs):
268 """
269 Check preconditions and if they are met, call the wrapped function.
270 """
271 checker(*args, **kwargs, what=what) 2AE
272 result = wrapped_function(*args, **kwargs) 2AE
274 return result 2AE
276 return inner 2AE
278 return outer 2AE
281def is_sequence(obj):
282 """
283 Check if the given object is a sequence (list or tuple).
284 """
285 return isinstance(obj, Sequence) 2# l CcDcDEEEFEqi$c%cGEHEIEJEKELEMENEOEPEQERESE?K@K[KrbpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRd7E8E5E4c5c}cadbdcdX2Y2Z20212JiKiI2J26ETgUgVgWgXgYg2232
288def is_nested_sequence(obj):
289 """
290 Check if the given object is a nested sequence (list or tuple with atleast one list or tuple element).
291 """
292 return is_sequence(obj) and any(is_sequence(elem) for elem in obj) 24c5cI2J2
296class Transaction:
297 """
298 A context manager for transactional operations with undo capability.
300 The Transaction class allows you to register undo actions (callbacks) that will be executed
301 if the transaction is not committed before exiting the context. This is useful for managing
302 resources or operations that need to be rolled back in case of errors or early exits.
304 Usage:
305 with Transaction() as txn:
306 txn.append(some_cleanup_function, arg1, arg2)
307 # ... perform operations ...
308 txn.commit() # Disarm undo actions; nothing will be rolled back on exit
310 Methods:
311 append(fn, *args, **kwargs): Register an undo action to be called on rollback.
312 commit(): Disarm all undo actions; nothing will be rolled back on exit.
313 """
314 def __init__(self):
315 self._stack = ExitStack() 2QcRcQbd Sc
316 self._entered = False 2QcRcQbd Sc
318 def __enter__(self):
319 self._stack.__enter__() 2QcRcQbd Sc
320 self._entered = True 2QcRcQbd Sc
321 return self 2QcRcQbd Sc
323 def __exit__(self, exc_type, exc, tb):
324 # If exit callbacks remain, they'll run in LIFO order.
325 self._entered = False 2QcRcQbd Sc
326 return self._stack.__exit__(exc_type, exc, tb) 2QcRcQbd Sc
328 def append(self, fn, /, *args, **kwargs):
329 """
330 Register an undo action (runs if the with-block exits without commit()).
331 Values are bound now via partial so late mutations don't bite you.
332 """
333 if not self._entered: 2QcRcQbd Sc
334 raise RuntimeError("Transaction must be entered before append()")
335 self._stack.callback(partial(fn, *args, **kwargs)) 2QcRcQbd Sc
337 def commit(self):
338 """
339 Disarm all undo actions. After this, exiting the with-block does nothing.
340 """
341 # pop_all() empties this stack so no callbacks are triggered on exit.
342 self._stack.pop_all() 2QcRcQbd Sc
345# Track whether we've already warned about fork method
346_fork_warning_checked = False
349def reset_fork_warning():
350 """Reset the fork warning check flag for testing purposes.
352 This function is intended for use in tests to allow multiple test runs
353 to check the warning behavior.
354 """
355 global _fork_warning_checked
356 _fork_warning_checked = False 2ug%e'e_c8h%d
359cdef inline tuple _read_fill_ptr(const char* ptr, Py_ssize_t width):
360 """Extract (value, element_size) from a raw pointer of known width."""
361 cdef unsigned int val
362 if width == 1: 2dieifigihiiijikiliminioi9Bc 6csbwctb7cxcubm $bn o p q r s t u v %bw x y z 6bA B 7b'b8cvbycwb9czcxbC (bD E F G H I J K L )bM N O P 8bQ R 9b*b!cybAczb#cBcAbS +bT U V W X Y Z 0 1 ,b2 3 4 5 !b6 7 #b-b
363 val = (<uint8_t*>ptr)[0] 2sbn o p w x vbD E F M N ybT U V 2 3
364 elif width == 2:
365 val = (<uint16_t*>ptr)[0] 2difihijiliniwctbq r s y z 6bycwbG H I O P 8bAczbW X Y 4 5 !b
366 elif width == 4:
367 val = (<uint32_t*>ptr)[0] 2eigiiikimioic xcubm t u v A B 7bzcxbC J K L Q R 9bBcAbS Z 0 1 6 7 #b
368 else:
369 raise ValueError(f"value must be 1, 2, or 4 bytes, got {width}") 29B6c7c$b%b'b8c9c(b)b*b!c#c+b,b-b
370 return (val, <unsigned int>width) 2dieifigihiiijikiliminioic sbwctbxcubm n o p q r s t u v w x y z 6bA B 7bvbycwbzcxbC D E F G H I J K L M N O P 8bQ R 9bybAczbBcAbS T U V W X Y Z 0 1 2 3 4 5 !b6 7 #b
373cpdef tuple _parse_fill_value(value):
374 """Parse a fill/memset value into (raw_value, element_size).
376 Parameters
377 ----------
378 value : int or buffer-protocol object
379 - int: Must be in range [0, 256). Treated as 1-byte fill.
380 - bytes or buffer-protocol: Must be 1, 2, or 4 bytes.
382 Returns
383 -------
384 tuple of (int, int)
385 (raw_value, element_size) where element_size is 1, 2, or 4.
387 Raises
388 ------
389 OverflowError
390 If int value is outside [0, 256).
391 TypeError
392 If value is not an int and does not support the buffer protocol.
393 ValueError
394 If value byte length is not 1, 2, or 4.
395 """
396 cdef uint8_t byte_val
397 cdef Py_buffer buf
399 if isinstance(value, int): 2GcHcIcJcKcLcMcNc(E)Ediei,E-Efigi:E;Ehiii@E[Ejiki_E`Elimi}E~Enioi9Bi j c 4ef k g h ke6csbwctb7cxcub0b(e)e*em $bn o p q r s t u v %bw x y z 6bA B 7b'ble8cvbycwb9czcxb1b+e,e-eC (bD E F G H I J K L )bM N O P 8bQ R 9b*bme!cybAczb#cBcAb2b.e/e:eS +bT U V W X Y Z 0 1 ,b2 3 4 5 !b6 7 #b-b?G|GbHuiqHuHzHDHwi6H9HAhSd[eVdxiyi$H%H'H(H)H*HBh-HzieIiIBixIBIGIKIDi%I)IChWd|eZdEiFi,I-I.I/I:I;IDh@IGinJtJzJFJLJRJXJ3J
400 byte_val = value 2GcHcIcJcKcLcMcNc(E)E,E-E:E;E@E[E_E`E}E~Ei j 4ef k g h 0b(e)e*e1b+e,e-e2b.e/e:e?G|GbHuiqHuHzHDHwi6H9HAhSd[eVdxiyi$H%H'H(H)H*HBh-HzieIiIBixIBIGIKIDi%I)IChWd|eZdEiFi,I-I.I/I:I;IDh@IGinJtJzJFJLJRJXJ3J
401 return (<unsigned int>byte_val, <unsigned int>1) 2GcHcIcJcKcLcMcNc(E)E,E-E:E;E@E[E_E`E}E~Ei j 4ef k g h 0b1b2b?G|GbHuiqHuHzHDHwi6H9HAhSd[eVdxiyi$H%H'H(H)H*HBh-HzieIiIBixIBIGIKIDi%I)IChWd|eZdEiFi,I-I.I/I:I;IDh@IGinJtJzJFJLJRJXJ3J
403 if isinstance(value, bytes): 2dieifigihiiijikiliminioi9Bc ke6csbwctb7cxcubm $bn o p q r s t u v %bw x y z 6bA B 7b'ble8cvbycwb9czcxbC (bD E F G H I J K L )bM N O P 8bQ R 9b*bme!cybAczb#cBcAbS +bT U V W X Y Z 0 1 ,b2 3 4 5 !b6 7 #b-b
404 return _read_fill_ptr(<const char*><bytes>value, <Py_ssize_t>len(value)) 2dieifigihiiijikiliminioi9B6csbwctb7cxcub8cvbycwb9czcxb!cybAczb#cBcAb
406 if PyObject_GetBuffer(value, &buf, PyBUF_SIMPLE) != 0: 2c kem $bn o p q r s t u v %bw x y z 6bA B 7b'bleC (bD E F G H I J K L )bM N O P 8bQ R 9b*bmeS +bT U V W X Y Z 0 1 ,b2 3 4 5 !b6 7 #b-b
407 raise TypeError(
408 f"value must be an int or support the buffer protocol, "
409 f"got {type(value).__name__}"
410 )
411 try: 2c m $bn o p q r s t u v %bw x y z 6bA B 7b'bC (bD E F G H I J K L )bM N O P 8bQ R 9b*bS +bT U V W X Y Z 0 1 ,b2 3 4 5 !b6 7 #b-b
412 return _read_fill_ptr(<const char*>buf.buf, buf.len) 2c m $bn o p q r s t u v %bw x y z 6bA B 7b'bC (bD E F G H I J K L )bM N O P 8bQ R 9b*bS +bT U V W X Y Z 0 1 ,b2 3 4 5 !b6 7 #b-b
413 finally:
414 PyBuffer_Release(&buf)
417def check_multiprocessing_start_method():
418 """Check if multiprocessing start method is 'fork' and warn if so."""
419 global _fork_warning_checked
420 if _fork_warning_checked: 28K9K#B$B!K#K%B'B# l ihjh#h$hkhlh%hOc$KqgK2%KL2M2'KN2O2(KP2Q2)KR2S2*KT2U2+KV2,K-K.K/Krg:K:c#d;c=c.bRb?c@csg;Ktg=KSbTbUbVbWbXbug%e'e_c8h%d
421 return 28K9K#B$B!K#K%B'B# l ihjh#h$hkhlh%hOc$KqgK2%KL2M2'KN2O2(KP2Q2)KR2S2*KT2U2+KV2,K-K.K/Krg:K:c#d;c=c.bRb?c@csg;Ktg=KSbTbUbVbWbXb%d
422 _fork_warning_checked = True 2Rbug%e'e_c8h%d
424 # Common warning message parts
425 common_message = (
426 "CUDA does not support. Forked subprocesses exhibit undefined behavior, " 2Rbug%e'e_c8h%d
427 "including failure to initialize CUDA contexts and devices. Set the start method "
428 "to 'spawn' before creating processes that use CUDA. "
429 "Use: multiprocessing.set_start_method('spawn')"
430 )
432 try: 2Rbug%e'e_c8h%d
433 start_method = multiprocessing.get_start_method() 2Rbug%e'e_c8h%d
434 if start_method == "fork": 2Rbug%e'e_c%d
435 message = f"multiprocessing start method is 'fork', which {common_message}" 2%e'e_c%d
436 warnings.warn(message, UserWarning, stacklevel=3) 2%e'e_c%d
437 except RuntimeError: 28h
438 # get_start_method() can raise RuntimeError if start method hasn't been set
439 # In this case, default is 'fork' on Linux, so we should warn
440 if platform.system() == "Linux": 28h
441 message = (
442 f"multiprocessing start method is not set and defaults to 'fork' on Linux, " 28h
443 f"which {common_message}"
444 )
445 warnings.warn(message, UserWarning, stacklevel=3) 28h