Coverage for cuda / core / _utils / cuda_utils.pyx: 91.46%
246 statements
« prev ^ index » next coverage.py v7.13.5, created at 2026-03-25 01:07 +0000
« prev ^ index » next coverage.py v7.13.5, created at 2026-03-25 01:07 +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 importlib.metadata
8import multiprocessing
9import platform
10import warnings
11from collections import namedtuple
12from collections.abc import Sequence
13from contextlib import ExitStack
14from typing import Callable
16try:
17 from cuda.bindings import driver, nvrtc, runtime
18except ImportError:
19 from cuda import cuda as driver
20 from cuda import cudart as runtime
21 from cuda import nvrtc
23from cuda.bindings.nvvm import nvvmError
24from cuda.bindings.nvjitlink import nvJitLinkError
26from cpython.buffer cimport PyObject_GetBuffer, PyBuffer_Release, Py_buffer, PyBUF_SIMPLE
28from cuda.bindings cimport cynvrtc, cynvvm, cynvjitlink
30from cuda.core._utils.driver_cu_result_explanations import DRIVER_CU_RESULT_EXPLANATIONS
31from cuda.core._utils.runtime_cuda_error_explanations import RUNTIME_CUDA_ERROR_EXPLANATIONS
34class CUDAError(Exception):
35 pass
38class NVRTCError(CUDAError):
39 pass
43ComputeCapability = namedtuple("ComputeCapability", ("major", "minor"))
46def cast_to_3_tuple(label, cfg):
47 cfg_orig = cfg 2a 8 9 ! # $ % ' d g e f 3 ( 4 5 ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdbebfbtbubvbwbxbyb2e3e4e4grd5e^fsd6etd7eud8evd9ewdic!e#e$e%e'ejczbAbBb(egbhbib6 7 !c[djb]cCC5g8fDC6g#cxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdUdVdWdXdYdCb)e*e7g+e,e8g-e.e]d^dPe_d`d/e:e;e=e?e@e[e]e^e_f_e`e`f{e|e{d|dQe}e~eaf}d~dbfaebecfdfefffgfZdhfif{f^c|f}fReSe~fagbgcgdg0d_cegTe1dfggghgjf2dcedekflfeefeUemfnf9gofpfgeheVeiejeqfrfsftfufvfwfxfyfigzfAfjgBfCfkeleWeDfEfFfmeneGfoepeHfIfJfKfLf3dMfNfkg`clgmgXeYengogpgqgrg4d{csgZe5dtgugvgOf6dqerePfQfsete0eRfSf!gTfUf#gVfWf$gXfYf%gZf0f'g1f2f(g3f4f)g5f6f*gTglCmC1zECUg
48 if isinstance(cfg, int): 28 9 ! # $ % ' d g e f 3 ( 4 5 ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdbebfbtbubvbwbxbyb2e3e4e4grd5e^fsd6etd7eud8evd9ewdic!e#e$e%e'ejczbAbBb(egbhbib6 7 !c[djb]cCC5g8fDC6g#cxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdUdVdWdXdYdCb)e*e7g+e,e8g-e.e]d^dPe_d`d/e:e;e=e?e@e[e]e^e_f_e`e`f{e|e{d|dQe}e~eaf}d~dbfaebecfdfefffgfZdhfif{f^c|f}fReSe~fagbgcgdg0d_cegTe1dfggghgjf2dcedekflfeefeUemfnf9gofpfgeheVeiejeqfrfsftfufvfwfxfyfigzfAfjgBfCfkeleWeDfEfFfmeneGfoepeHfIfJfKfLf3dMfNfkg`clgmgXeYengogpgqgrg4d{csgZe5dtgugvgOf6dqerePfQfsete0eRfSf!gTfUf#gVfWf$gXfYf%gZf0f'g1f2f(g3f4f)g5f6f*gTglCmC1zECUg
49 cfg = (cfg,) 28 9 ! # $ % ' d g e f 3 ( 4 5 ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdbebfbtbubvbwbxbyb2e3e4e5e6e7e8e9eic!e#e$e%e'ejczbAbBb(egbhbib6 7 !c[djb]cCC8fDCxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdUdVdWdXdYdCb)e*e+e,e-e.e]d^d_d`d/e:e;e=e?e@e[e]e^e_e`e{e|e{d|d}e~eaf}d~dbfaebecfdfefffgfZdhfif^c_cjf2dcedekflfeefemfnfofpfgeheiejeqfrfsftfufvfwfxfyfzfAfBfCfkeleDfEfFfmeneGfoepeHfIfJfKfLf3dMfNf`c{cOf6dqerePfQfseteRfSfTfUfVfWfXfYfZf0f1f2f3f4f5f6fTglC
50 else:
51 common = "must be an int, or a tuple with up to 3 ints" 24grd^fsdtdudvdwd5g8f6g#c7g8gPe_f`fQe{f^c|f}fReSe~fagbgcgdg0d_cegTe1dfggghgUe9gVeigjgWekg`clgmgXeYengogpgqgrg4d{csgZe5dtgugvg0e!g#g$g%g'g(g)g*gTgmC1zECUg
52 if not isinstance(cfg, tuple): 24grd^fsdtdudvdwd5g8f6g#c7g8gPe_f`fQe{f^c|f}fReSe~fagbgcgdg0d_cegTe1dfggghgUe9gVeigjgWekg`clgmgXeYengogpgqgrg4d{csgZe5dtgugvg0e!g#g$g%g'g(g)g*gTgmC1zECUg
53 raise ValueError(f"{label} {common} (got {type(cfg)})") 2EC
54 if len(cfg) > 3: 24grd^fsdtdudvdwd5g8f6g#c7g8gPe_f`fQe{f^c|f}fReSe~fagbgcgdg0d_cegTe1dfggghgUe9gVeigjgWekg`clgmgXeYengogpgqgrg4d{csgZe5dtgugvg0e!g#g$g%g'g(g)g*gTgmC1zUg
55 raise ValueError(f"{label} {common} (got tuple with length {len(cfg)})") 2mC
56 if any(not isinstance(val, int) for val in cfg): 24grd^fsdtdudvdwd5g8f6g#c7g8gPe_f`fQe{f^c|f}fReSe~fagbgcgdg0d_cegTe1dfggghgUe9gVeigjgWekg`clgmgXeYengogpgqgrg4d{csgZe5dtgugvg0e!g#g$g%g'g(g)g*gTg1zUg
57 raise ValueError(f"{label} {common} (got {cfg})") 21z
58 if any(val < 1 for val in cfg): 28 9 ! # $ % ' d g e f 3 ( 4 5 ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdbebfbtbubvbwbxbyb2e3e4e4grd5e^fsd6etd7eud8evd9ewdic!e#e$e%e'ejczbAbBb(egbhbib6 7 !c[djb]cCC5g8fDC6g#cxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdUdVdWdXdYdCb)e*e7g+e,e8g-e.e]d^dPe_d`d/e:e;e=e?e@e[e]e^e_f_e`e`f{e|e{d|dQe}e~eaf}d~dbfaebecfdfefffgfZdhfif{f^c|f}fReSe~fagbgcgdg0d_cegTe1dfggghgjf2dcedekflfeefeUemfnf9gofpfgeheVeiejeqfrfsftfufvfwfxfyfigzfAfjgBfCfkeleWeDfEfFfmeneGfoepeHfIfJfKfLf3dMfNfkg`clgmgXeYengogpgqgrg4d{csgZe5dtgugvgOf6dqerePfQfsete0eRfSf!gTfUf#gVfWf$gXfYf%gZf0f'g1f2f(g3f4f)g5f6f*gTglCUg
59 plural_s = "" if len(cfg) == 1 else "s" 28flCUg
60 raise ValueError(f"{label} value{plural_s} must be >= 1 (got {cfg_orig})") 28flCUg
61 return cfg + (1,) * (3 - len(cfg)) 28 9 ! # $ % ' d g e f 3 ( 4 5 ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdbebfbtbubvbwbxbyb2e3e4e4grd5e^fsd6etd7eud8evd9ewdic!e#e$e%e'ejczbAbBb(egbhbib6 7 !c[djb]cCC5g8fDC6g#cxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdUdVdWdXdYdCb)e*e7g+e,e8g-e.e]d^dPe_d`d/e:e;e=e?e@e[e]e^e_f_e`e`f{e|e{d|dQe}e~eaf}d~dbfaebecfdfefffgfZdhfif{f^c|f}fReSe~fagbgcgdg0d_cegTe1dfggghgjf2dcedekflfeefeUemfnf9gofpfgeheVeiejeqfrfsftfufvfwfxfyfigzfAfjgBfCfkeleWeDfEfFfmeneGfoepeHfIfJfKfLf3dMfNfkg`clgmgXeYengogpgqgrg4d{csgZe5dtgugvgOf6dqerePfQfsete0eRfSf!gTfUf#gVfWf$gXfYf%gZf0f'g1f2f(g3f4f)g5f6f*gTg
64def _reduce_3_tuple(t: tuple):
65 return t[0] * t[1] * t[2] 2]c
68cdef int HANDLE_RETURN(cydriver.CUresult err) except?-1 nogil:
69 if err != cydriver.CUresult.CUDA_SUCCESS: 2a NIOIPI8 QI9 RI! SI7dTI# UI8dVI$ WI% XI' YI|cZI}c0I~c1I9d2I!d3I#d4Id 5Ig 6Ie 7If 8I3 9I( !I4 #I5 $I) %I* 'I+ (I$d)I%d*I'd+I(d,I)d-I*d.I, /I- :I. ;I/ =I: ?I; @I= [I? ]I@ ^I[ _I] `I^ {I_ |I` }I{ ~I| aJ} bJ~ cJabdJbbeJcbfJdbgJebhJfbiJjJkJlJmJnJoJpJqJrJwgsJtJuJvJwJxJyJzJAJBJCJDJEJFJGJHJIJJJKJLJMJNJOJPJQJRJSJueTJUJveVJweWJxeXJyeYJtbZJze0JAe1JBe2J3J4J5J6Jub7Jvb8Jwb9Jxb!Jyb#J$c$J%c%J'c'J(c(J)c)J*c*J+c+J,c,J-J.J/J:J;J9f=J!f?J#f@J[J]J^J_J`J{J2e|J3e}J~JaKbK4ecKFCGCdKeKfKgKhKrdiKHCIC+g,gjKkKlKmKnKoKpK5eqKrKJCsKKCtKuKvKwKxKyKzK^fsdAKBKCKLCDKMCEK-gFK.gGKHKIKJKKKLKMKNKOKPKQKRKSKTKUKVKWKXKYKZK0K1K2K3K6e4KNCOC5K6K7K8K9Ktd!KPCQC/g:g#K$K%K'K(K)K*K7e+K,K-KRC.KSC/K:K;K=K?K@K[K]K^K_K`Kud{K|K}KTC~KUCaL;gbL=gcLdLeLfLgLhLiLjLkL8elLmLnLVCoLWCpLqLrLsLtLuLvLwLxLyLzLvdALBLCLXCDLYCEL?gFL@gGLHLILJLKLLLMLNLOL9ePLQLRLZCSL0CTLULVLWLXLYLZL0L1L2L3Lwd4L5L6L1C7L2C8L[g9L]g!L#L$L%L'L(L)L*L+L,L-L.L/L:L;L=L?L@L[L]L^L_L`L{L|L}L~LaMbMcMdMeMfMgMhMiMjMkMlMmMxgnMoMpMqMygrMsMtMuMvMwMxMyMzgzMAMBMCMDMEMFM2zGMicHMIMJMAgKMBgLMMMNMOMPMQMRMSMTMUMVMWMXMYMZM0M1M2M3M4M5M6M7M!e8M#e9M$e!M%e#MCe$MDe%M'e'M3C(M4C)M5C*Mad+Mbd,M6C-M7C.M3z/Mjc:Mzb;MAb=MBb?M(e@MEe[MCg]Mgb^Mhb_Mib`MoI8CpI9C4z!C5z#CqI$CrI%C6z'C7z(C6 )C7 *CDg{MEg|MVg}MWg~MFgaNGgbNXgcNYgdN^geN_gfN8zgNsI+CtI,C-C.C/CuI:ChN;CFe=CvI?CiN@C[C]CwI^CjN_C`C{CxI|CkN}C~CaDyIbDlNcDdDeDzIfDmNgDhDAIiDBIjDCIkDDIlDEImDFInD+doD,dpD-dqD.drD-cnN.coN/dsD:dtDGIuDHIvDIIwDJIxDEbpNFbqNGbrNHbsNIbtNJbuN$f%fyDvNzDwNADxNBDyNCDzNDDANBNCN'fGeHeDNEDENFNGNHNINJNKNLNMN(fNNONPNQNRNSNTNUNVNWNXNYNZN0N1N2N3N4N5N6N7N8N9N!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~NaObOcOdOeOfOgOhOiOjOkOlOmOnOoOpOqOrOsOtOuOvOwOxOyOzOAOBOCODOEOFOGOHOIOJOKOLOMONOOOPOQOROSOTOUOVOWOXOYOZO0O1O2O3O4O5O6O7O8O9O!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~OaPbPcPdPePfPgPhPHgiPjPkPIg)f*fJg9z!c!zlPFDmPGDnPHDoPIDpPJDqPKDrPLDsPMDtP#zuP$zvP%zwP[dxPyPNDODPDQDRDzPAPBPCPDPEPSDTDUDVDWDFPXDGPYDZDKgHPIPJPjb6c7c]cKPLPMP#cNPxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdUdVdWdXdYdKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'z'bOPPP+f8c9c,fQP-f.f/f0D1DRP:f;f=f?fLgMg/cSPDcTPkbUPkcVPlbWPEcXPlcYPmbZP;b0Pcd1Pdd2Ped3Ph 4P=b5Pi 6Pj 7Pk 8Pl 9Pm !Pn #Po $Pp %Pq 'P?b(Pr )Ps *Pt +Pu ,P(b-Pv .Pw /P)b:P@b;P:c=PFc?Pnb@Pmc[Pob]PGc^Pnc_Ppb`P[b{Pfd|Pgd}Phd~Px aQ]bbQy cQz dQA eQB fQC gQD hQE iQF jQG kQ^blQH mQI nQJ oQK pQ*bqQL rQM sQ+btQ_buQ;cvQHcwQqbxQocyQrbzQIcAQpcBQsbCQ`bDQidEQjdFQkdGQN HQ{bIQO JQP KQQ LQR MQS NQT OQU PQV QQW RQ|bSQX TQY UQZ VQ0 WQ,bXQ1 YQ2 ZQ-b0Q}b1Q@f2D3Dld2Q3Q(z4Q4D5Q5D6Q)z7Q*z8Q+z9Q,z!Q=c#QIe$Qmd%QJe'Q6D(Q7D)Q8D*Q9D+Q!D,Q#D-Q$D.Q%D/Q'D:Q(D;Q)D=Q*D?Q+D@Q,D[Q-D]Q.D^Q-z_Q/D`Q.z{Q:D|Q/z}Q;D~Q:zaR=DbR;zcR?DdR=zeR@DfR?zgR[DhR@ziR]DjR^DkR_DlR`DmR{DnR|DoR}DpR~DqRaErRbEsRcEtRdEuReEvRfEwRgExRhEyRiEzRnCAR`gBRoCCRpCDR{gERqCFRjEGR[f]fKekEHRndIRlEmEnEJRoEKRpELRqEMR?c@cDb[cNROR.bPR;dQR[z]zRR^zSR_zTR`zCbUR{z/b|zVR}zWR~zXRaAYRbAZRcA0RdA1ReA2RNgfA3RgAhAiAjAkA~bac4RlA5RmA6RnA7RoA8RpA9RqA!RrA#RsA$RtA%RuA'RvA(RwA)RxA*RyA+RzA,RAA-RBA.RCA/RDA:REA;RFA=RGA?RHA@RIAJAKA[RLA]RMA^RNA_ROA`RPA{RQA|RRA}RSA~RTAaSUAbSVAcSWAdSXAeSYAfSZAgS0AhS1AiS2AjS3AkS4AlS5A6AmS7AnS8AoS9ApS!AqS#ArS$AsS%AtS'AuS(AvS)AwS*AxS+AyS,AzS-AAS.ABS/A:A7fCSLeDSMeESodFS=dGSHSrE)eISJSKSLSsEtEuEMSNSOSPSQS;A*eRSvE=ASSwETSUSxE?AyEVSzE+eWSXSYSZSAEBECE0S1S2S3S4S@A,e5SDE[A6SEE7S8S-eFEGE9S!S#S$S.e%SHE'S(SIE]dJEKErCsCLEMENEOEPE|g^dPeQE}g~gRESETE_d`d/e:e;e=e?e)S@e*SUEVE+S,S-S.S[e/SWE:S;S=S?SXE]e@S[SYEZE]S^S_S`S{S]A^e_f|S0E^A}S~SaTbT1E_ecTdTeT2E3E4EfTgThTiTjT_A`e`fkT5E`AlT6EmTnToT{e7E8EpTqTrTsT|etT9EuTvT!EtC{d#E$E%E'E(E)E*E+E,E-E{A|dQe.Eah|A/E:E;E=E?E@E[E]E^E_E`E{E|E}E}e~EaFwTxTyTzTbFATBTCT~eDTETcFdFFTGTHTITafJTeFKTLTMTNTOTPTQTRTSTTTUTVT}A}d~AaBbBcBdBeBfBgBhB~diBOgjBkBlBbfaebecfdfefffgfZdhfif{f^c|f}fReSe~fagbgcgdg0d_cegTe1dfggghgjffFgFWTXTYTZT0T2dbhchhFiFjFkFlFmFmBcenBoBpBqBrBsBtBuBvBwBdexBPgyBzBAB1Tkf2TnFoF3T4T5T6Tlf7TpF8T9TqFuCeerFsFtFuFvFwFxFyFzFAFBFBBfeUeCFdhCBDFEFFF!T#TGFmf$T%T'T(THFIFJF)T*T+T,T-TDBnf.TKFEB/TLF:T;TofMFNF=T?T@T[Tpf]TOF^T_TPFgeQFRFvCwCSFTFUFVFWFehheVeXFfhghYFZF0Fiejeqfrfsftfuf`Tvf{T1F2F|T}T~TaUwfbU3FcUdUeUfU4FxfgUhU5F6FiUjUkUlUmUFByfignU7FGBoUpUqUrU8FzfsUtUuU9F!F#FvUwUxUyUzUHBAfjgAU$FIBBU%FCUDUEUBf'F(FFUGUHUIUCfJU)FKULU*FxCke+F,F-F.F/F:F;F=F?F@FJBleWe[FhhKB]F^F_F`F{F|F}F~FaGbGcGdGeGfGDfgGhGMUNUOUPUiGQURUSUEfTUUUjGkGVUWUXUYUFfZUlG0U1U2U3U4U5U6U7U8U9U!U#ULBmeMBNBOBPBQBRBSBTBUBneVBQgWBXBYBGfoepeHfIfJfKfLf3dMfNfkg`clgmgXeYengogpgqgrg4d{csgZe5dtgugvgOfmGnG$U%U'U(U)U6dihjhoGpGqGrGsGtGZBqe0B1B2B3B4B5B6B7B8B9Bre!BRg#B$B%B*UPf+UuGvG,U-U.U/UQf:UwG;U=UxGyCseyGzGAGBGCGDGEGFGGGHGIG'Bte0eJGkh(BKGLGMG?U@U[UNGOGlhmh)B*B]UPGRf^U_U`U{UQGRGSG|U}U~UaVbV+BSfcVTG,B-Bpd.B/BqddVUGeVfVgVVGTfhViVjVkVWGXGYGlVmVnVoVpV:BUfqVZG;BrV0GsVtVuV1GVfvVwVxVyV2G3G4GzVAVBVCVDV=BWfEV5G?BFV6GGVHVIV7GXfJVKVLVMV8G9G!GNVOVPVQVRV@BYfSV#G[BTV$GUVVVWV%GZfXVYVZV0V'G(G)G1V2V3V4V5V]B0f6V*G^B7V+G8V9V!V,G1f#V$V%V'V-G.G/G(V)V*V+V,V_B2f-V:G`B.V;G/V:V;V=G3f=V?V@V[V?G@G[G]V^V_V`V{V{B4f|V]G|B}B}V^G~VaWbW_G5fcWdWeWfW`G{G|GgWhWiWjWkW~B6flW}GaCbCmW~GnWoWbcpWqWzCccrWdcsWtWuWvWwWxWyWzWAWBWCWDWEWFWGWHWIWJWKWLWMWNWOWPWQWRWSWTWUWVWWWXWYWZW0W1W2W3Wc 4WeccC?d5W6W7W8W9W!W#W$W%W'W(W)W*W+WnhohdCeCfC:b,WgCaHbHhC-W.W/W:W;W=W?Wfc@W[W]W^WcH_WdH`WeH{WfH|WgH}WhH~WiHaXjHbXkHcXlHdXmHeXnHfXoHgXpHhXqHiXrHjXiCkXsHlXtHmXuHnXvHoXwHpXxHqXrXsXtXuXvXwXxXyXzXAXBXCXDXEXFXGXHXIXJXKXLXMXNXOXPXQXRXSXTXUXVXWXXXYXZX0X1X2X3X4X5X6X7X8X9X!X#X$X%X'X(X)X*X+X,X-X.X/X:X;X=X?X@X[X]X^X_X`X{XyH|XzHAHBHCH}XDH~XEHaYbYFHcYGHHHdYIHJHKHLHMHeYfYgYhYNHOHPHQHiYjYkYRHSHlYmYTHUHnYoYpYVHqYrYsYtYuYvYwYxYyYzYAYBYCYDYEYFYGYHYIYJYKYLYMYNYOYPYQYRYSYTYUYVYWYXYYYZY0Y1Y
70 return _check_driver_error(err) 2wgxgzgAg/b?d
71 return 0 2a NIOIPI8 QI9 RI! SI7dTI# UI8dVI$ WI% XI' YI|cZI}c0I~c1I9d2I!d3I#d4Id 5Ig 6Ie 7If 8I3 9I( !I4 #I5 $I) %I* 'I+ (I$d)I%d*I'd+I(d,I)d-I*d.I, /I- :I. ;I/ =I: ?I; @I= [I? ]I@ ^I[ _I] `I^ {I_ |I` }I{ ~I| aJ} bJ~ cJabdJbbeJcbfJdbgJebhJfbiJjJkJlJmJnJoJpJqJrJwgsJtJuJvJwJxJyJzJAJBJCJDJEJFJGJHJIJJJKJLJMJNJOJPJQJRJSJueTJUJveVJweWJxeXJyeYJtbZJze0JAe1JBe2J3J4J5J6Jub7Jvb8Jwb9Jxb!Jyb#J$c$J%c%J'c'J(c(J)c)J*c*J+c+J,c,J-J.J/J:J;J9f=J!f?J#f@J[J]J^J_J`J{J2e|J3e}J~JaKbK4ecKFCGCdKeKfKgKhKrdiKHCIC+g,gjKkKlKmKnKoKpK5eqKrKJCsKKCtKuKvKwKxKyKzK^fsdAKBKCKLCDKMCEK-gFK.gGKHKIKJKKKLKMKNKOKPKQKRKSKTKUKVKWKXKYKZK0K1K2K3K6e4KNCOC5K6K7K8K9Ktd!KPCQC/g:g#K$K%K'K(K)K*K7e+K,K-KRC.KSC/K:K;K=K?K@K[K]K^K_K`Kud{K|K}KTC~KUCaL;gbL=gcLdLeLfLgLhLiLjLkL8elLmLnLVCoLWCpLqLrLsLtLuLvLwLxLyLzLvdALBLCLXCDLYCEL?gFL@gGLHLILJLKLLLMLNLOL9ePLQLRLZCSL0CTLULVLWLXLYLZL0L1L2L3Lwd4L5L6L1C7L2C8L[g9L]g!L#L$L%L'L(L)L*L+L,L-L.L/L:L;L=L?L@L[L]L^L_L`L{L|L}L~LaMbMcMdMeMfMgMhMiMjMkMlMmMxgnMoMpMqMygrMsMtMuMvMwMxMyMzgzMAMBMCMDMEMFM2zGMicHMIMJMAgKMBgLMMMNMOMPMQMRMSMTMUMVMWMXMYMZM0M1M2M3M4M5M6M7M!e8M#e9M$e!M%e#MCe$MDe%M'e'M3C(M4C)M5C*Mad+Mbd,M6C-M7C.M3z/Mjc:Mzb;MAb=MBb?M(e@MEe[MCg]Mgb^Mhb_Mib`MoI8CpI9C4z!C5z#CqI$CrI%C6z'C7z(C6 )C7 *CDg{MEg|MVg}MWg~MFgaNGgbNXgcNYgdN^geN_gfN8zgNsI+CtI,C-C.C/CuI:ChN;CFe=CvI?CiN@C[C]CwI^CjN_C`C{CxI|CkN}C~CaDyIbDlNcDdDeDzIfDmNgDhDAIiDBIjDCIkDDIlDEImDFInD+doD,dpD-dqD.drD-cnN.coN/dsD:dtDGIuDHIvDIIwDJIxDEbpNFbqNGbrNHbsNIbtNJbuN$f%fyDvNzDwNADxNBDyNCDzNDDANBNCN'fGeHeDNEDENFNGNHNINJNKNLNMN(fNNONPNQNRNSNTNUNVNWNXNYNZN0N1N2N3N4N5N6N7N8N9N!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~NaObOcOdOeOfOgOhOiOjOkOlOmOnOoOpOqOrOsOtOuOvOwOxOyOzOAOBOCODOEOFOGOHOIOJOKOLOMONOOOPOQOROSOTOUOVOWOXOYOZO0O1O2O3O4O5O6O7O8O9O!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~OaPbPcPdPePfPgPhPHgiPjPkPIg)f*fJg9z!c!zlPFDmPGDnPHDoPIDpPJDqPKDrPLDsPMDtP#zuP$zvP%zwP[dxPyPNDODPDQDRDzPAPBPCPDPEPSDTDUDVDWDFPXDGPYDZDKgHPIPJPjb6c7c]cKPLPMP#cNPxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdUdVdWdXdYdKbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'z'bOPPP+f8c9c,fQP-f.f/f0D1DRP:f;f=f?fLgMg/cSPDcTPkbUPkcVPlbWPEcXPlcYPmbZP;b0Pcd1Pdd2Ped3Ph 4P=b5Pi 6Pj 7Pk 8Pl 9Pm !Pn #Po $Pp %Pq 'P?b(Pr )Ps *Pt +Pu ,P(b-Pv .Pw /P)b:P@b;P:c=PFc?Pnb@Pmc[Pob]PGc^Pnc_Ppb`P[b{Pfd|Pgd}Phd~Px aQ]bbQy cQz dQA eQB fQC gQD hQE iQF jQG kQ^blQH mQI nQJ oQK pQ*bqQL rQM sQ+btQ_buQ;cvQHcwQqbxQocyQrbzQIcAQpcBQsbCQ`bDQidEQjdFQkdGQN HQ{bIQO JQP KQQ LQR MQS NQT OQU PQV QQW RQ|bSQX TQY UQZ VQ0 WQ,bXQ1 YQ2 ZQ-b0Q}b1Q@f2D3Dld2Q3Q(z4Q4D5Q5D6Q)z7Q*z8Q+z9Q,z!Q=c#QIe$Qmd%QJe'Q6D(Q7D)Q8D*Q9D+Q!D,Q#D-Q$D.Q%D/Q'D:Q(D;Q)D=Q*D?Q+D@Q,D[Q-D]Q.D^Q-z_Q/D`Q.z{Q:D|Q/z}Q;D~Q:zaR=DbR;zcR?DdR=zeR@DfR?zgR[DhR@ziR]DjR^DkR_DlR`DmR{DnR|DoR}DpR~DqRaErRbEsRcEtRdEuReEvRfEwRgExRhEyRiEzRnCAR`gBRoCCRpCDR{gERqCFRjEGR[f]fKekEHRndIRlEmEnEJRoEKRpELRqEMR?c@cDb[cNROR.bPR;dQR[z]zRR^zSR_zTR`zCbUR{z/b|zVR}zWR~zXRaAYRbAZRcA0RdA1ReA2RNgfA3RgAhAiAjAkA~bac4RlA5RmA6RnA7RoA8RpA9RqA!RrA#RsA$RtA%RuA'RvA(RwA)RxA*RyA+RzA,RAA-RBA.RCA/RDA:REA;RFA=RGA?RHA@RIAJAKA[RLA]RMA^RNA_ROA`RPA{RQA|RRA}RSA~RTAaSUAbSVAcSWAdSXAeSYAfSZAgS0AhS1AiS2AjS3AkS4AlS5A6AmS7AnS8AoS9ApS!AqS#ArS$AsS%AtS'AuS(AvS)AwS*AxS+AyS,AzS-AAS.ABS/A:A7fCSLeDSMeESodFS=dGSHSrE)eISJSKSLSsEtEuEMSNSOSPSQS;A*eRSvE=ASSwETSUSxE?AyEVSzE+eWSXSYSZSAEBECE0S1S2S3S4S@A,e5SDE[A6SEE7S8S-eFEGE9S!S#S$S.e%SHE'S(SIE]dJEKErCsCLEMENEOEPE|g^dPeQE}g~gRESETE_d`d/e:e;e=e?e)S@e*SUEVE+S,S-S.S[e/SWE:S;S=S?SXE]e@S[SYEZE]S^S_S`S{S]A^e_f|S0E^A}S~SaTbT1E_ecTdTeT2E3E4EfTgThTiTjT_A`e`fkT5E`AlT6EmTnToT{e7E8EpTqTrTsT|etT9EuTvT!EtC{d#E$E%E'E(E)E*E+E,E-E{A|dQe.Eah|A/E:E;E=E?E@E[E]E^E_E`E{E|E}E}e~EaFwTxTyTzTbFATBTCT~eDTETcFdFFTGTHTITafJTeFKTLTMTNTOTPTQTRTSTTTUTVT}A}d~AaBbBcBdBeBfBgBhB~diBOgjBkBlBbfaebecfdfefffgfZdhfif{f^c|f}fReSe~fagbgcgdg0d_cegTe1dfggghgjffFgFWTXTYTZT0T2dbhchhFiFjFkFlFmFmBcenBoBpBqBrBsBtBuBvBwBdexBPgyBzBAB1Tkf2TnFoF3T4T5T6Tlf7TpF8T9TqFuCeerFsFtFuFvFwFxFyFzFAFBFBBfeUeCFdhCBDFEFFF!T#TGFmf$T%T'T(THFIFJF)T*T+T,T-TDBnf.TKFEB/TLF:T;TofMFNF=T?T@T[Tpf]TOF^T_TPFgeQFRFvCwCSFTFUFVFWFehheVeXFfhghYFZF0Fiejeqfrfsftfuf`Tvf{T1F2F|T}T~TaUwfbU3FcUdUeUfU4FxfgUhU5F6FiUjUkUlUmUFByfignU7FGBoUpUqUrU8FzfsUtUuU9F!F#FvUwUxUyUzUHBAfjgAU$FIBBU%FCUDUEUBf'F(FFUGUHUIUCfJU)FKULU*FxCke+F,F-F.F/F:F;F=F?F@FJBleWe[FhhKB]F^F_F`F{F|F}F~FaGbGcGdGeGfGDfgGhGMUNUOUPUiGQURUSUEfTUUUjGkGVUWUXUYUFfZUlG0U1U2U3U4U5U6U7U8U9U!U#ULBmeMBNBOBPBQBRBSBTBUBneVBQgWBXBYBGfoepeHfIfJfKfLf3dMfNfkg`clgmgXeYengogpgqgrg4d{csgZe5dtgugvgOfmGnG$U%U'U(U)U6dihjhoGpGqGrGsGtGZBqe0B1B2B3B4B5B6B7B8B9Bre!BRg#B$B%B*UPf+UuGvG,U-U.U/UQf:UwG;U=UxGyCseyGzGAGBGCGDGEGFGGGHGIG'Bte0eJGkh(BKGLGMG?U@U[UNGOGlhmh)B*B]UPGRf^U_U`U{UQGRGSG|U}U~UaVbV+BSfcVTG,B-Bpd.B/BqddVUGeVfVgVVGTfhViVjVkVWGXGYGlVmVnVoVpV:BUfqVZG;BrV0GsVtVuV1GVfvVwVxVyV2G3G4GzVAVBVCVDV=BWfEV5G?BFV6GGVHVIV7GXfJVKVLVMV8G9G!GNVOVPVQVRV@BYfSV#G[BTV$GUVVVWV%GZfXVYVZV0V'G(G)G1V2V3V4V5V]B0f6V*G^B7V+G8V9V!V,G1f#V$V%V'V-G.G/G(V)V*V+V,V_B2f-V:G`B.V;G/V:V;V=G3f=V?V@V[V?G@G[G]V^V_V`V{V{B4f|V]G|B}B}V^G~VaWbW_G5fcWdWeWfW`G{G|GgWhWiWjWkW~B6flW}GaCbCmW~GnWoWbcpWqWzCccrWdcsWtWuWvWwWxWyWzWAWBWCWDWEWFWGWHWIWJWKWLWMWNWOWPWQWRWSWTWUWVWWWXWYWZW0W1W2W3Wc 4WeccC?d5W6W7W8W9W!W#W$W%W'W(W)W*W+WnhohdCeCfC:b,WgCaHbHhC-W.W/W:W;W=W?Wfc@W[W]W^WcH_WdH`WeH{WfH|WgH}WhH~WiHaXjHbXkHcXlHdXmHeXnHfXoHgXpHhXqHiXrHjXiCkXsHlXtHmXuHnXvHoXwHpXxHqXrXsXtXuXvXwXxXyXzXAXBXCXDXEXFXGXHXIXJXKXLXMXNXOXPXQXRXSXTXUXVXWXXXYXZX0X1X2X3X4X5X6X7X8X9X!X#X$X%X'X(X)X*X+X,X-X.X/X:X;X=X?X@X[X]X^X_X`X{XyH|XzHAHBHCH}XDH~XEHaYbYFHcYGHHHdYIHJHKHLHMHeYfYgYhYNHOHPHQHiYjYkYRHSHlYmYTHUHnYoYpYVHqYrYsYtYuYvYwYxYyYzYAYBYCYDYEYFYGYHYIYJYKYLYMYNYOYPYQYRYSYTYUYVYWYXYYYZY0Y1Y
74cdef int HANDLE_RETURN_NVRTC(cynvrtc.nvrtcProgram prog, cynvrtc.nvrtcResult err) except?-1 nogil:
75 """Handle NVRTC result codes, raising NVRTCError with program log on failure."""
76 if err == cynvrtc.nvrtcResult.NVRTC_SUCCESS: 28 9 ! # $ % ' d g e f 3 ( 4 5 ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdbebfbtbubvbwbxbyb2e3e4erd5esd6etd7eud8evd9ewdic!e#e$e%e'e3zjczbAbBb(egbhbib6 7 !c[djb]c#cxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdUdVdWdXdYdb KbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'z'b.b;d[z]z^z_z`zCb{z/b|z}z~zaAbAcAdAeAWHfAgAhAiAjAkA~baclAmAnAoApAqArAsAtAuAvAwAxAyAzAAABACADAEAFAGAHAIAJAKALAMANAOAPAQARASATAUAVAWAXAYAZA0A1A2A3A4A5A6A7A8A9A!A#A$A%A'A(A)A*A+A,A-A.A/A:A)e;A*e=A?A+e@A,e[A-e.e]d|g^d~g_d`d/e:e;e=e?e@e[e]e]A^e^A_e_A`e`A{e|e{d{A|d|A}e~eaf}A}d~AaBbBcBdBeBfBgBhB~diBOgjBkBlBbfaebecfdfefffgfZdhfif^c0d_c1djf2dmBcenBoBpBqBrBsBtBuBvBwBdexBPgyBzBABkflfeeBBfeCBmfDBnfEBofpfgeehheghiejeqfrfsftfufvfwfxfFByfGBzfHBAfIBBfCfkeJBleKBDfEfFfLBmeMBNBOBPBQBRBSBTBUBneVBQgWBXBYBGfoepeHfIfJfKfLf3dMfNf`c4d{c5dOf6dZBqe0B1B2B3B4B5B6B7B8B9Bre!BRg#B$B%BPfQfse'Bte(B)B*BRf+BSf,B-B.BTf:BUf;BVf=BWf?BXf@BYf[BZf]B0f^B1f_B2f`B3f{B4f|B}B5f~B6faCbCbcJcccdcKcLcMcqcrcscNcOcPcQcgchctcRcucScTcUcVcvcWcwcxcyczcXcYcZc0c1c2c3c4c5cAcBcCcc dCeCfC:bgChCfc
77 return 0 28 9 ! # $ % ' d g e f 3 ( 4 5 ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdbebfbtbubvbwbxbyb2e3e4erd5esd6etd7eud8evd9ewdic!e#e$e%e'e3zjczbAbBb(egbhbib6 7 !c[djb]c#cxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdUdVdWdXdYdb KbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'z'b.b;d[z]z^z_z`zCb{z/b|z}z~zaAbAcAdAeAWHfAgAhAiAjAkA~baclAmAnAoApAqArAsAtAuAvAwAxAyAzAAABACADAEAFAGAHAIAJAKALAMANAOAPAQARASATAUAVAWAXAYAZA0A1A2A3A4A5A6A7A8A9A!A#A$A%A'A(A)A*A+A,A-A.A/A:A)e;A*e=A?A+e@A,e[A-e.e]d|g^d~g_d`d/e:e;e=e?e@e[e]e]A^e^A_e_A`e`A{e|e{d{A|d|A}e~eaf}A}d~AaBbBcBdBeBfBgBhB~diBOgjBkBlBbfaebecfdfefffgfZdhfif^c0d_c1djf2dmBcenBoBpBqBrBsBtBuBvBwBdexBPgyBzBABkflfeeBBfeCBmfDBnfEBofpfgeehheghiejeqfrfsftfufvfwfxfFByfGBzfHBAfIBBfCfkeJBleKBDfEfFfLBmeMBNBOBPBQBRBSBTBUBneVBQgWBXBYBGfoepeHfIfJfKfLf3dMfNf`c4d{c5dOf6dZBqe0B1B2B3B4B5B6B7B8B9Bre!BRg#B$B%BPfQfse'Bte(B)B*BRf+BSf,B-B.BTf:BUf;BVf=BWf?BXf@BYf[BZf]B0f^B1f_B2f`B3f{B4f|B}B5f~B6faCbCbcJcccdcKcLcMcqcrcscNcOcPcQcgchctcRcucScTcUcVcvcWcwcxcyczcXcYcZc0c1c2c3c4c5cAcBcCcc dCeCfC:bgChCfc
78 with gil: 1c
79 _raise_nvrtc_error(prog, err) 1c
82cdef int _raise_nvrtc_error(cynvrtc.nvrtcProgram prog, cynvrtc.nvrtcResult err) except -1:
83 """Build error message with program log and raise NVRTCError."""
84 cdef const char* err_str = cynvrtc.nvrtcGetErrorString(err) 1c
85 cdef size_t logsize = 0 1c
86 if prog != NULL: 1c
87 cynvrtc.nvrtcGetProgramLogSize(prog, &logsize) 1c
88 cdef bytes log_bytes
89 cdef str log_str = "" 1c
90 if logsize > 1 and prog != NULL: 1c
91 log_bytes = b" " * logsize 1c
92 if cynvrtc.nvrtcGetProgramLog(prog, <char*>log_bytes) == cynvrtc.nvrtcResult.NVRTC_SUCCESS: 1c
93 log_str = log_bytes.decode("utf-8", errors="backslashreplace") 1c
94 err_msg = f"{err}: {err_str.decode()}" if err_str != NULL else f"NVRTC error {err}" 1c
95 if log_str: 1c
96 err_msg += f", compilation log:\n\n{log_str}" 1c
97 raise NVRTCError(err_msg) 1c
100cdef int HANDLE_RETURN_NVVM(cynvvm.nvvmProgram prog, cynvvm.nvvmResult err) except?-1 nogil:
101 """Handle NVVM result codes, raising nvvmError with program log on failure."""
102 if err == cynvvm.nvvmResult.NVVM_SUCCESS: 2/BXHYHeccC?dZH0H1H2H3H4H5H6H7H8H9H!H#H$Hnhoh
103 return 0 2/BXHYHeccC?dZH0H1H2H3H4H5H6H7H8H9H!H#H$Hnhoh
104 with gil: 2ec
105 _raise_nvvm_error(prog, err) 2ec
108cdef int _raise_nvvm_error(cynvvm.nvvmProgram prog, cynvvm.nvvmResult err) except -1:
109 """Raise nvvmError annotated with the program log."""
110 cdef size_t logsize = 0 2ec
111 if prog != NULL: 2ec
112 cynvvm.nvvmGetProgramLogSize(prog, &logsize) 2ec
113 cdef bytes log_bytes
114 cdef str log_str = "" 2ec
115 if logsize > 1 and prog != NULL: 2a ec
116 log_bytes = b" " * logsize 2ec
117 if cynvvm.nvvmGetProgramLog(prog, <char*>log_bytes) == cynvvm.nvvmResult.NVVM_SUCCESS: 2ec
118 log_str = log_bytes.decode("utf-8", errors="backslashreplace") 2ec
119 cdef object exc = nvvmError(err) 2ec
120 if log_str: 2ec
121 exc.args = (exc.args[0] + f"\nNVVM program log: {log_str}", *exc.args[1:]) 2ec
122 raise exc 2ec
125cdef int HANDLE_RETURN_NVJITLINK(
126 cynvjitlink.nvJitLinkHandle handle, cynvjitlink.nvJitLinkResult err) except?-1 nogil:
127 """Handle nvJitLink result codes, raising nvJitLinkError with error log on failure."""
128 if err == cynvjitlink.nvJitLinkResult.NVJITLINK_SUCCESS: 2b %H'H(H)H*H+H,H-H.H/HACBC:H;H=H?H@H[H]H^H_H`H{H|H}H1e~HaIbIcIdIeIfIgIhIiIjIkIlImI
129 return 0 2b %H'H(H)H*H+H,H-H.H/HACBC:H;H=H?H@H[H]H^H_H`H{H|H}H~HaIbIcIdIeIfIgIhIiIjIkIlImI
130 with gil: 2b 1e
131 _raise_nvjitlink_error(handle, err) 2b 1e
134cdef int _raise_nvjitlink_error(
135 cynvjitlink.nvJitLinkHandle handle, cynvjitlink.nvJitLinkResult err) except -1:
136 """Raise nvJitLinkError annotated with the error log."""
137 cdef size_t logsize = 0 2b 1e
138 if handle != NULL: 2b 1e
139 cynvjitlink.nvJitLinkGetErrorLogSize(handle, &logsize) 1b
140 cdef bytes log_bytes
141 cdef str log_str = "" 2b 1e
142 if logsize > 1 and handle != NULL: 2b 1e
143 log_bytes = b" " * logsize 1b
144 if cynvjitlink.nvJitLinkGetErrorLog(handle, <char*>log_bytes) == \ 1b
145 cynvjitlink.nvJitLinkResult.NVJITLINK_SUCCESS: 1b
146 log_str = log_bytes.decode("utf-8", errors="backslashreplace") 1b
147 cdef object exc = nvJitLinkError(err) 2b 1e
148 if log_str: 2b 1e
149 exc.args = (exc.args[0] + f"\nnvJitLink error log: {log_str}", *exc.args[1:]) 1b
150 raise exc 2b 1e
153cdef object _RUNTIME_SUCCESS = runtime.cudaError_t.cudaSuccess
154cdef object _NVRTC_SUCCESS = nvrtc.nvrtcResult.NVRTC_SUCCESS
157cpdef inline int _check_driver_error(cydriver.CUresult error) except?-1 nogil:
158 if error == cydriver.CUresult.CUDA_SUCCESS: 2a phqhrhshthuhvhwhxhyhzhAhBh8 Ch9 Dh! Eh7dFh# Gh8dHh$ Ih% Jh' Kh|c}c~c9d!d#dd g e f 3 ( 4 5 ) * + $d%d'd(d)d*d, Lh- Mh. Nh/ Oh: Ph; Qh= Rh? Sh@ Th[ Uh] Vh^ Wh_ Xh` Yh{ Zh| 0h} 1h~ 2hab3hbb4hcb5hdb6heb7hfb8h9h!h#h$hwg%h'h(h)h*h+h,h-h.h/h:h;h=h?h@hue[h]hve^hwe_hxe`hye{htb|hze}hAe~hBeaibiciubdivbeiwbfixbgiybhi$cii%cji'cki(cli)cmi*cni+coi,cpiqiZgri0gsi1gti9fui!fvi#fwixiyiziAiBiCiDiEiFiGiHiIiJiKiLiMiNiOiPiQiRiSiTiUiViWiXiYiZi0i1i2i3i4i5i6i7i8i9i!i#i$i%i'i(i)i*i+i,i-i.i/i:i;i=i?i@i[i]i^i_i`i{i|i}i~iajbjcjdjejfjgjhjijjjkjljmjnjojpjqjrjsjtjujvjwjxjyjzjAjBjCjDjEjFjGjHjIjJjKjLjMjNjOjPjQjRjSjTjUjVjWjXjYjZj0j1j2j3j4j5j6j7j8j9j!j#j$j%j'j(j)j*j+j,j-j.j/j:j;j=j?j@j[j]j^j_jxg`j{jyg|j}j~jakbkzgckdkekfkgkichkikAgjkBgkklkmknkokpkqkrksktkukvkwkxkykzkCeAkDeBkCkDkEkFkadGkbdHkIkJkKkjcLkzbMkAbNkBbOkPkEeQkCgRkgbSkhbTkibUk6 7 @dFe+d,d-d.d-c.c/d:dEbFbGbHbIbJbVkWkXkYkZk0k1k2g2k3k4k5k6k7k8k9k!k#k$k%k'k(kHg)k*k+kIg,kJg!c-k.k/k:k;k=k?k@k[k]k^k_k`kKgjb6c7c{k|k}k~kxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdUdVdWdXdYdb KbalLbblMbclNbdlObelPbflQbglRbhlSbilTbjlUbklVbllWbmlXbnlYbolZbpl0bql1brl2bsl3btl4bul5bvl6bwl7bxl8byl9bzl!bAl#bBl$bCl%bDlEl'bFlGlHl+f8c9c,f-f.f/f:f;f=f?fLgMg/cDckbkclbEclcmb;bcdddedh =bi j k l m n o p q ?br s t u (bv w )b@b:cFcnbmcobGcncpb[bfdgdhdx ]by z A B C D E F G ^bH I J K *bL M +b_b;cHcqbocrbIcpcsb`bidjdkdN {bO P Q R S T U V W |bX Y Z 0 ,b1 2 -b}b@fIlJlKlLlMlNlOlPlQlRlSlTlUlVlWlXlYlZl0l1l2l3l4l5l6l7l8l9l!l#l$l%l'l(l)l*l+l,l-l.l/l:l;l=l?l@l[l]l^l_l`l{l|l}l~lambmcmdmemfmgmhmimjmkmlm[f]fKemmnmompmqmrm?c@cDb[c.bsm;dtmumvmwmxmymCbzmAm/bBmCmDmEmFmGmHmImNgJmKmLmMmNmOmPmNe~bQmOeacRmSmTmUmVmWmXmYmZm0m1m2m3m4m5m6m7m8m9m!m#m$m%m'm(m)m*m+m,m-m.m/m:m;m=m?m@m[m]m^m_m`m{m|m}m~manbncndnenfngnhninjnknlnmnnnonpnqnrnsntnunvnwnxnynznAnBnCnDnEnFnGnHnInJnKnLnMnNnOnPnQnRnSnTnUnVnWnXnYnZn0n1n2n3n4n5n6n7n8n9n!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~naobocodoeofogohoiojokolomonooopoqorosotouovowoxoyozoAoBoCoDoEoFoGoHoIoJoKoLoMoNoOoPoQoRoSoToUoVoWoXoYoZo0o1o2o3o4o5o6o7o8o9o!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~uavbvcvdvevfvgvhvpdivjvkvqdlvmvnvovpvqvrvsvtvuvvvwvxvyvzvAvBvCvDvEvFvGvHvIvJvKvLvMvNvOvPvQvRvSvTvUvVvWvXvYvZv0v1v2v3v4v5v6v7v8v9v!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~waxbxcxdxexfxgxhxbcixJcjxcckxdclxKcmxLcnxMcoxqcpxrcqxscrxNcsxOctxPcuxQcvxgcwxhcxxtcyxRczxucAxScBxTcCxUcDxVcExvcFxWcGxwcHxxcIxycJxzcKxXcLxYcMxZcNx0cOx1cPx2cQx3cRx4cSx5cTxAcUxBcVxCcWxc Xx?dYxZx0x1x2x3x4x5x6x7x8x9x!x#x:b$x%x'x(x)x*x+x,xfc-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~yazbzczdzezfzgzhzizjzkzlzmznzozpzqzrzsztzuzvzwzxzyzzzAzBzCzDzEzFzGzHzIzJzKzLzMzNzOzPzQzRzSzTzUzVzWzXzYz
159 return 0 2a phqhrhshthuhvhwhxhyhzhAhBh8 Ch9 Dh! Eh7dFh# Gh8dHh$ Ih% Jh' Kh|c}c~c9d!d#dd g e f 3 ( 4 5 ) * + $d%d'd(d)d*d, Lh- Mh. Nh/ Oh: Ph; Qh= Rh? Sh@ Th[ Uh] Vh^ Wh_ Xh` Yh{ Zh| 0h} 1h~ 2hab3hbb4hcb5hdb6heb7hfb8h9h!h#h$h%h'h(h)h*h+h,h-h.h/h:h;h=h?h@hue[h]hve^hwe_hxe`hye{htb|hze}hAe~hBeaibiciubdivbeiwbfixbgiybhi$cii%cji'cki(cli)cmi*cni+coi,cpiqiZgri0gsi1gti9fui!fvi#fwixiyiziAiBiCiDiEiFiGiHiIiJiKiLiMiNiOiPiQiRiSiTiUiViWiXiYiZi0i1i2i3i4i5i6i7i8i9i!i#i$i%i'i(i)i*i+i,i-i.i/i:i;i=i?i@i[i]i^i_i`i{i|i}i~iajbjcjdjejfjgjhjijjjkjljmjnjojpjqjrjsjtjujvjwjxjyjzjAjBjCjDjEjFjGjHjIjJjKjLjMjNjOjPjQjRjSjTjUjVjWjXjYjZj0j1j2j3j4j5j6j7j8j9j!j#j$j%j'j(j)j*j+j,j-j.j/j:j;j=j?j@j[j]j^j_j`j{jyg|j}j~jakbkckdkekfkgkichkikjkBgkklkmknkokpkqkrksktkukvkwkxkykzkCeAkDeBkCkDkEkFkadGkbdHkIkJkKkjcLkzbMkAbNkBbOkPkEeQkCgRkgbSkhbTkibUk6 7 @dFe+d,d-d.d-c.c/d:dEbFbGbHbIbJbVkWkXkYkZk0k1k2g2k3k4k5k6k7k8k9k!k#k$k%k'k(kHg)k*k+kIg,kJg!c-k.k/k:k;k=k?k@k[k]k^k_k`kKgjb6c7c{k|k}k~kxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdUdVdWdXdYdb KbalLbblMbclNbdlObelPbflQbglRbhlSbilTbjlUbklVbllWbmlXbnlYbolZbpl0bql1brl2bsl3btl4bul5bvl6bwl7bxl8byl9bzl!bAl#bBl$bCl%bDlEl'bFlGlHl+f8c9c,f-f.f/f:f;f=f?fLgMg/cDckbkclbEclcmb;bcdddedh =bi j k l m n o p q ?br s t u (bv w )b@b:cFcnbmcobGcncpb[bfdgdhdx ]by z A B C D E F G ^bH I J K *bL M +b_b;cHcqbocrbIcpcsb`bidjdkdN {bO P Q R S T U V W |bX Y Z 0 ,b1 2 -b}b@fIlJlKlLlMlNlOlPlQlRlSlTlUlVlWlXlYlZl0l1l2l3l4l5l6l7l8l9l!l#l$l%l'l(l)l*l+l,l-l.l/l:l;l=l?l@l[l]l^l_l`l{l|l}l~lambmcmdmemfmgmhmimjmkmlm[f]fKemmnmompmqmrm?c@cDb[c.bsm;dtmumvmwmxmymCbzmAm/bBmCmDmEmFmGmHmImNgJmKmLmMmNmOmPmNe~bQmOeacRmSmTmUmVmWmXmYmZm0m1m2m3m4m5m6m7m8m9m!m#m$m%m'm(m)m*m+m,m-m.m/m:m;m=m?m@m[m]m^m_m`m{m|m}m~manbncndnenfngnhninjnknlnmnnnonpnqnrnsntnunvnwnxnynznAnBnCnDnEnFnGnHnInJnKnLnMnNnOnPnQnRnSnTnUnVnWnXnYnZn0n1n2n3n4n5n6n7n8n9n!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~naobocodoeofogohoiojokolomonooopoqorosotouovowoxoyozoAoBoCoDoEoFoGoHoIoJoKoLoMoNoOoPoQoRoSoToUoVoWoXoYoZo0o1o2o3o4o5o6o7o8o9o!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~uavbvcvdvevfvgvhvpdivjvkvqdlvmvnvovpvqvrvsvtvuvvvwvxvyvzvAvBvCvDvEvFvGvHvIvJvKvLvMvNvOvPvQvRvSvTvUvVvWvXvYvZv0v1v2v3v4v5v6v7v8v9v!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~waxbxcxdxexfxgxhxbcixJcjxcckxdclxKcmxLcnxMcoxqcpxrcqxscrxNcsxOctxPcuxQcvxgcwxhcxxtcyxRczxucAxScBxTcCxUcDxVcExvcFxWcGxwcHxxcIxycJxzcKxXcLxYcMxZcNx0cOx1cPx2cQx3cRx4cSx5cTxAcUxBcVxCcWxc XxYxZx0x1x2x3x4x5x6x7x8x9x!x#x:b$x%x'x(x)x*x+x,xfc-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~yazbzczdzezfzgzhzizjzkzlzmznzozpzqzrzsztzuzvzwzxzyzzzAzBzCzDzEzFzGzHzIzJzKzLzMzNzOzPzQzRzSzTzUzVzWzXzYz
160 cdef const char* name
161 name_err = cydriver.cuGetErrorName(error, &name) 2wgxgzgAg2gDb/b?d
162 if name_err != cydriver.CUresult.CUDA_SUCCESS: 2wgxgzgAg2gDb/b?d
163 raise CUDAError(f"UNEXPECTED ERROR CODE: {error}") 22g
164 with gil: 2wgxgzgAg2gDb/b?d
165 # TODO: consider lower this to Cython
166 expl = DRIVER_CU_RESULT_EXPLANATIONS.get(int(error)) 2wgxgzgAg2gDb/b?d
167 if expl is not None: 2wgxgzgAg2gDb/b?d
168 raise CUDAError(f"{name.decode()}: {expl}") 2wgxgzgAg2gDb/b?d
169 cdef const char* desc
170 desc_err = cydriver.cuGetErrorString(error, &desc)
171 if desc_err != cydriver.CUresult.CUDA_SUCCESS:
172 raise CUDAError(f"{name.decode()}")
173 raise CUDAError(f"{name.decode()}: {desc.decode()}")
176cpdef inline int _check_runtime_error(error) except?-1:
177 if error == _RUNTIME_SUCCESS: 2EbFbGbHbIbJb$f%f3gZz'fGeHe(f)f*f6c7c8c9ckblbmb;bh i j k l m n o p q r s t u v w nbobpb[bx y z A B C D E F G H I J K L M qbrbsb`bN O P Q R S T U V W X Y Z 0 1 2 ld=cIemdJend
178 return 0 2EbFbGbHbIbJb$f%f3gZz'fGeHe(f)f*f6c7c8c9ckblbmb;bh i j k l m n o p q r s t u v w nbobpb[bx y z A B C D E F G H I J K L M qbrbsb`bN O P Q R S T U V W X Y Z 0 1 2 ld=cIemdJend
179 name_err, name = runtime.cudaGetErrorName(error) 2Zz
180 if name_err != _RUNTIME_SUCCESS: 2Zz
181 raise CUDAError(f"UNEXPECTED ERROR CODE: {error}")
182 name = name.decode() 2Zz
183 expl = RUNTIME_CUDA_ERROR_EXPLANATIONS.get(int(error)) 2Zz
184 if expl is not None: 2Zz
185 raise CUDAError(f"{name}: {expl}") 2Zz
186 desc_err, desc = runtime.cudaGetErrorString(error)
187 if desc_err != _RUNTIME_SUCCESS:
188 raise CUDAError(f"{name}")
189 desc = desc.decode()
190 raise CUDAError(f"{name}: {desc}")
193cpdef inline int _check_nvrtc_error(error, handle=None) except?-1:
194 if error == _NVRTC_SUCCESS: 2a 0znIb KbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b.bNe~bOeacpdqdbcJcccdcKcLcMcqcrcscNcOcPcQcgchctcRcucScTcUcVcvcWcwcxcyczcXcYcZc0c1c2c3c4c5cAcBcCcc :bfc
195 return 0 2a nIb KbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b.bNe~bOeacpdqdbcJcccdcKcLcMcqcrcscNcOcPcQcgchctcRcucScTcUcVcvcWcwcxcyczcXcYcZc0c1c2c3c4c5cAcBcCcc :bfc
196 err = f"{error}: {nvrtc.nvrtcGetErrorString(error)[1].decode()}" 20znI
197 if handle is not None: 20znI
198 _, logsize = nvrtc.nvrtcGetProgramLogSize(handle) 20z
199 log = b" " * logsize 20z
200 _ = nvrtc.nvrtcGetProgramLog(handle, log) 20z
201 err += f", compilation log:\n\n{log.decode('utf-8', errors='backslashreplace')}" 20z
202 raise NVRTCError(err) 20znI
205cdef inline int _check_error(error, handle=None) except?-1:
206 if isinstance(error, driver.CUresult): 2a phqhrhshthuhvhwhxhyhzhAhBh8 Ch9 Dh! Eh7dFh# Gh8dHh$ Ih% Jh' Kh|c}c~c9d!d#dd g e f 3 ( 4 5 ) * + $d%d'd(d)d*d, Lh- Mh. Nh/ Oh: Ph; Qh= Rh? Sh@ Th[ Uh] Vh^ Wh_ Xh` Yh{ Zh| 0h} 1h~ 2hab3hbb4hcb5hdb6heb7hfb8h9h!h#h$h%h'h(h)h*h+h,h-h.h/h:h;h=h?h@hue[h]hve^hwe_hxe`hye{htb|hze}hAe~hBeaibiciubdivbeiwbfixbgiybhi$cii%cji'cki(cli)cmi*cni+coi,cpiqiZgri0gsi1gti9fui!fvi#fwixiyiziAiBiCiDiEiFiGiHiIiJiKiLiMiNiOiPiQiRiSiTiUiViWiXiYiZi0i1i2i3i4i5i6i7i8i9i!i#i$i%i'i(i)i*i+i,i-i.i/i:i;i=i?i@i[i]i^i_i`i{i|i}i~iajbjcjdjejfjgjhjijjjkjljmjnjojpjqjrjsjtjujvjwjxjyjzjAjBjCjDjEjFjGjHjIjJjKjLjMjNjOjPjQjRjSjTjUjVjWjXjYjZj0j1j2j3j4j5j6j7j8j9j!j#j$j%j'j(j)j*j+j,j-j.j/j:j;j=j?j@j[j]j^j_j`j{jyg|j}j~jakbkckdkekfkgkichkikjkBgkklkmknkokpkqkrksktkukvkwkxkykzkCeAkDeBkCkDkEkFkadGkbdHkIkJkKkjcLkzbMkAbNkBbOkPkEeQkCgRkgbSkhbTkibUk6 7 @dFe+d,d-d.d-c.c/d:dEbFbGbHbIbJb$f%fVk3gWkXkYkZk0k1k2k'fGe3kHe4k5k6k7k8k9k!k#k(f$k%k'k(kHg)k*k+kIg)f*f,kJg-k.k/k:k;k=k?k@k[k]k^k_k`kKgjb6c7c{k|k}k~kb KbalLbblMbclNbdlObelPbflQbglRbhlSbilTbjlUbklVbllWbmlXbnlYbolZbpl0bql1brl2bsl3btl4bul5bvl6bwl7bxl8byl9bzl!bAl#bBl$bCl%bDlEl'bFlGlHl+f8c9c,f-f.f/f:f;f=f?fLgMg/cDckbkclbEclcmb;bcdddedh =bi j k l m n o p q ?br s t u (bv w )b@b:cFcnbmcobGcncpb[bfdgdhdx ]by z A B C D E F G ^bH I J K *bL M +b_b;cHcqbocrbIcpcsb`bidjdkdN {bO P Q R S T U V W |bX Y Z 0 ,b1 2 -b}b@fldIlJlKlLlMlNlOlPlQl=cRlIeSlmdTlJeUlVlWlXlYlZl0l1l2l3l4l5l6l7l8l9l!l#l$l%l'l(l)l*l+l,l-l.l/l:l;l=l?l@l[l]l^l_l`l{l|l}l~lambmcmdmemfmgmhmimjmkmlm[f]fKemmndnmompmqmrm.bsm;dtmumvmwmxmymCbzmAm/bBmCmDmEmFmGmHmImNgJmKmLmMmNmOmPmNe~bQmOeacRmSmTmUmVmWmXmYmZm0m1m2m3m4m5m6m7m8m9m!m#m$m%m'm(m)m*m+m,m-m.m/m:m;m=m?m@m[m]m^m_m`m{m|m}m~manbncndnenfngnhninjnknlnmnnnonpnqnrnsntnunvnwnxnynznAnBnCnDnEnFnGnHnInJnKnLnMnNnOnPnQnRnSnTnUnVnWnXnYnZn0n1n2n3n4n5n6n7n8n9n!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~naobocodoeofogohoiojokolomonooopoqorosotouovowoxoyozoAoBoCoDoEoFoGoHoIoJoKoLoMoNoOoPoQoRoSoToUoVoWoXoYoZo0o1o2o3o4o5o6o7o8o9o!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~uavbvcvdvevfvgvhvpdivjvkvqdlvmvnvovpvqvrvsvtvuvvvwvxvyvzvAvBvCvDvEvFvGvHvIvJvKvLvMvNvOvPvQvRvSvTvUvVvWvXvYvZv0v1v2v3v4v5v6v7v8v9v!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~waxbxcxdxexfxgxhxbcixJcjxcckxdclxKcmxLcnxMcoxqcpxrcqxscrxNcsxOctxPcuxQcvxgcwxhcxxtcyxRczxucAxScBxTcCxUcDxVcExvcFxWcGxwcHxxcIxycJxzcKxXcLxYcMxZcNx0cOx1cPx2cQx3cRx4cSx5cTxAcUxBcVxCcWxc XxYxZx0x1x2x3x4x5x6x7x8x9x!x#x:b$x%x'x(x)x*x+x,xfc-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~yazbzczdzezfzgzhzizjzkzlzmznzozpzqzrzsztzuzvzwzxzyzzzAzBzCzDzEzFzGzHzIzJzKzLzMzNzOzPzQzRzSzTzUzVzWzXzYz
207 return _check_driver_error(error) 2a phqhrhshthuhvhwhxhyhzhAhBh8 Ch9 Dh! Eh7dFh# Gh8dHh$ Ih% Jh' Kh|c}c~c9d!d#dd g e f 3 ( 4 5 ) * + $d%d'd(d)d*d, Lh- Mh. Nh/ Oh: Ph; Qh= Rh? Sh@ Th[ Uh] Vh^ Wh_ Xh` Yh{ Zh| 0h} 1h~ 2hab3hbb4hcb5hdb6heb7hfb8h9h!h#h$h%h'h(h)h*h+h,h-h.h/h:h;h=h?h@hue[h]hve^hwe_hxe`hye{htb|hze}hAe~hBeaibiciubdivbeiwbfixbgiybhi$cii%cji'cki(cli)cmi*cni+coi,cpiqiZgri0gsi1gti9fui!fvi#fwixiyiziAiBiCiDiEiFiGiHiIiJiKiLiMiNiOiPiQiRiSiTiUiViWiXiYiZi0i1i2i3i4i5i6i7i8i9i!i#i$i%i'i(i)i*i+i,i-i.i/i:i;i=i?i@i[i]i^i_i`i{i|i}i~iajbjcjdjejfjgjhjijjjkjljmjnjojpjqjrjsjtjujvjwjxjyjzjAjBjCjDjEjFjGjHjIjJjKjLjMjNjOjPjQjRjSjTjUjVjWjXjYjZj0j1j2j3j4j5j6j7j8j9j!j#j$j%j'j(j)j*j+j,j-j.j/j:j;j=j?j@j[j]j^j_j`j{jyg|j}j~jakbkckdkekfkgkichkikjkBgkklkmknkokpkqkrksktkukvkwkxkykzkCeAkDeBkCkDkEkFkadGkbdHkIkJkKkjcLkzbMkAbNkBbOkPkEeQkCgRkgbSkhbTkibUk6 7 @dFe+d,d-d.d-c.c/d:dEbFbGbHbIbJbVkWkXkYkZk0k1k2k3k4k5k6k7k8k9k!k#k$k%k'k(kHg)k*k+kIg,kJg-k.k/k:k;k=k?k@k[k]k^k_k`kKgjb6c7c{k|k}k~kb KbalLbblMbclNbdlObelPbflQbglRbhlSbilTbjlUbklVbllWbmlXbnlYbolZbpl0bql1brl2bsl3btl4bul5bvl6bwl7bxl8byl9bzl!bAl#bBl$bCl%bDlEl'bFlGlHl+f8c9c,f-f.f/f:f;f=f?fLgMg/cDckbkclbEclcmb;bcdddedh =bi j k l m n o p q ?br s t u (bv w )b@b:cFcnbmcobGcncpb[bfdgdhdx ]by z A B C D E F G ^bH I J K *bL M +b_b;cHcqbocrbIcpcsb`bidjdkdN {bO P Q R S T U V W |bX Y Z 0 ,b1 2 -b}b@fIlJlKlLlMlNlOlPlQlRlSlTlUlVlWlXlYlZl0l1l2l3l4l5l6l7l8l9l!l#l$l%l'l(l)l*l+l,l-l.l/l:l;l=l?l@l[l]l^l_l`l{l|l}l~lambmcmdmemfmgmhmimjmkmlm[f]fKemmnmompmqmrm.bsm;dtmumvmwmxmymCbzmAm/bBmCmDmEmFmGmHmImNgJmKmLmMmNmOmPmNe~bQmOeacRmSmTmUmVmWmXmYmZm0m1m2m3m4m5m6m7m8m9m!m#m$m%m'm(m)m*m+m,m-m.m/m:m;m=m?m@m[m]m^m_m`m{m|m}m~manbncndnenfngnhninjnknlnmnnnonpnqnrnsntnunvnwnxnynznAnBnCnDnEnFnGnHnInJnKnLnMnNnOnPnQnRnSnTnUnVnWnXnYnZn0n1n2n3n4n5n6n7n8n9n!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~naobocodoeofogohoiojokolomonooopoqorosotouovowoxoyozoAoBoCoDoEoFoGoHoIoJoKoLoMoNoOoPoQoRoSoToUoVoWoXoYoZo0o1o2o3o4o5o6o7o8o9o!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~uavbvcvdvevfvgvhvpdivjvkvqdlvmvnvovpvqvrvsvtvuvvvwvxvyvzvAvBvCvDvEvFvGvHvIvJvKvLvMvNvOvPvQvRvSvTvUvVvWvXvYvZv0v1v2v3v4v5v6v7v8v9v!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~waxbxcxdxexfxgxhxbcixJcjxcckxdclxKcmxLcnxMcoxqcpxrcqxscrxNcsxOctxPcuxQcvxgcwxhcxxtcyxRczxucAxScBxTcCxUcDxVcExvcFxWcGxwcHxxcIxycJxzcKxXcLxYcMxZcNx0cOx1cPx2cQx3cRx4cSx5cTxAcUxBcVxCcWxc XxYxZx0x1x2x3x4x5x6x7x8x9x!x#x:b$x%x'x(x)x*x+x,xfc-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~yazbzczdzezfzgzhzizjzkzlzmznzozpzqzrzsztzuzvzwzxzyzzzAzBzCzDzEzFzGzHzIzJzKzLzMzNzOzPzQzRzSzTzUzVzWzXzYz
208 elif isinstance(error, runtime.cudaError_t): 2a EbFbGbHbIbJb$f%f3g'fGeHe(f)f*f6c7cb KbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b8c9ckblbmb;bh i j k l m n o p q r s t u v w nbobpb[bx y z A B C D E F G H I J K L M qbrbsb`bN O P Q R S T U V W X Y Z 0 1 2 ld=cIemdJend.bNe~bOeacpdqdbcJcccdcKcLcMcqcrcscNcOcPcQcgchctcRcucScTcUcVcvcWcwcxcyczcXcYcZc0c1c2c3c4c5cAcBcCcc :bfc
209 return _check_runtime_error(error) 2EbFbGbHbIbJb$f%f3g'fGeHe(f)f*f6c7c8c9ckblbmb;bh i j k l m n o p q r s t u v w nbobpb[bx y z A B C D E F G H I J K L M qbrbsb`bN O P Q R S T U V W X Y Z 0 1 2 ld=cIemdJend
210 elif isinstance(error, nvrtc.nvrtcResult): 2a b KbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b.bNe~bOeacpdqdbcJcccdcKcLcMcqcrcscNcOcPcQcgchctcRcucScTcUcVcvcWcwcxcyczcXcYcZc0c1c2c3c4c5cAcBcCcc :bfc
211 return _check_nvrtc_error(error, handle=handle) 2a b KbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b.bNe~bOeacpdqdbcJcccdcKcLcMcqcrcscNcOcPcQcgchctcRcucScTcUcVcvcWcwcxcyczcXcYcZc0c1c2c3c4c5cAcBcCcc :bfc
212 else:
213 raise RuntimeError(f"Unknown error type: {error}")
216def handle_return(tuple result, handle=None):
217 _check_error(result[0], handle=handle) 2a phqhrhshthuhvhwhxhyhzhAhBh8 Ch9 Dh! Eh7dFh# Gh8dHh$ Ih% Jh' Kh|c}c~c9d!d#dd g e f 3 ( 4 5 ) * + $d%d'd(d)d*d, Lh- Mh. Nh/ Oh: Ph; Qh= Rh? Sh@ Th[ Uh] Vh^ Wh_ Xh` Yh{ Zh| 0h} 1h~ 2hab3hbb4hcb5hdb6heb7hfb8h9h!h#h$h%h'h(h)h*h+h,h-h.h/h:h;h=h?h@hue[h]hve^hwe_hxe`hye{htb|hze}hAe~hBeaibiciubdivbeiwbfixbgiybhi$cii%cji'cki(cli)cmi*cni+coi,cpiqiZgri0gsi1gti9fui!fvi#fwixiyiziAiBiCiDiEiFiGiHiIiJiKiLiMiNiOiPiQiRiSiTiUiViWiXiYiZi0i1i2i3i4i5i6i7i8i9i!i#i$i%i'i(i)i*i+i,i-i.i/i:i;i=i?i@i[i]i^i_i`i{i|i}i~iajbjcjdjejfjgjhjijjjkjljmjnjojpjqjrjsjtjujvjwjxjyjzjAjBjCjDjEjFjGjHjIjJjKjLjMjNjOjPjQjRjSjTjUjVjWjXjYjZj0j1j2j3j4j5j6j7j8j9j!j#j$j%j'j(j)j*j+j,j-j.j/j:j;j=j?j@j[j]j^j_j`j{jyg|j}j~jakbkckdkekfkgkichkikjkBgkklkmknkokpkqkrksktkukvkwkxkykzkCeAkDeBkCkDkEkFkadGkbdHkIkJkKkjcLkzbMkAbNkBbOkPkEeQkCgRkgbSkhbTkibUk6 7 @dFe+d,d-d.d-c.c/d:dEbFbGbHbIbJb$f%fVk3gWkXkYkZk0k1k2k'fGe3kHe4k5k6k7k8k9k!k#k(f$k%k'k(kHg)k*k+kIg)f*f,kJg-k.k/k:k;k=k?k@k[k]k^k_k`kKgjb6c7c{k|k}k~kb KbalLbblMbclNbdlObelPbflQbglRbhlSbilTbjlUbklVbllWbmlXbnlYbolZbpl0bql1brl2bsl3btl4bul5bvl6bwl7bxl8byl9bzl!bAl#bBl$bCl%bDlEl'bFlGlHl+f8c9c,f-f.f/f:f;f=f?fLgMg/cDckbkclbEclcmb;bcdddedh =bi j k l m n o p q ?br s t u (bv w )b@b:cFcnbmcobGcncpb[bfdgdhdx ]by z A B C D E F G ^bH I J K *bL M +b_b;cHcqbocrbIcpcsb`bidjdkdN {bO P Q R S T U V W |bX Y Z 0 ,b1 2 -b}b@fldIlJlKlLlMlNlOlPlQl=cRlIeSlmdTlJeUlVlWlXlYlZl0l1l2l3l4l5l6l7l8l9l!l#l$l%l'l(l)l*l+l,l-l.l/l:l;l=l?l@l[l]l^l_l`l{l|l}l~lambmcmdmemfmgmhmimjmkmlm[f]fKemmndnmompmqmrm.bsm;dtmumvmwmxmymCbzmAm/bBmCmDmEmFmGmHmImNgJmKmLmMmNmOmPmNe~bQmOeacRmSmTmUmVmWmXmYmZm0m1m2m3m4m5m6m7m8m9m!m#m$m%m'm(m)m*m+m,m-m.m/m:m;m=m?m@m[m]m^m_m`m{m|m}m~manbncndnenfngnhninjnknlnmnnnonpnqnrnsntnunvnwnxnynznAnBnCnDnEnFnGnHnInJnKnLnMnNnOnPnQnRnSnTnUnVnWnXnYnZn0n1n2n3n4n5n6n7n8n9n!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~naobocodoeofogohoiojokolomonooopoqorosotouovowoxoyozoAoBoCoDoEoFoGoHoIoJoKoLoMoNoOoPoQoRoSoToUoVoWoXoYoZo0o1o2o3o4o5o6o7o8o9o!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~uavbvcvdvevfvgvhvpdivjvkvqdlvmvnvovpvqvrvsvtvuvvvwvxvyvzvAvBvCvDvEvFvGvHvIvJvKvLvMvNvOvPvQvRvSvTvUvVvWvXvYvZv0v1v2v3v4v5v6v7v8v9v!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~waxbxcxdxexfxgxhxbcixJcjxcckxdclxKcmxLcnxMcoxqcpxrcqxscrxNcsxOctxPcuxQcvxgcwxhcxxtcyxRczxucAxScBxTcCxUcDxVcExvcFxWcGxwcHxxcIxycJxzcKxXcLxYcMxZcNx0cOx1cPx2cQx3cRx4cSx5cTxAcUxBcVxCcWxc XxYxZx0x1x2x3x4x5x6x7x8x9x!x#x:b$x%x'x(x)x*x+x,xfc-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~yazbzczdzezfzgzhzizjzkzlzmznzozpzqzrzsztzuzvzwzxzyzzzAzBzCzDzEzFzGzHzIzJzKzLzMzNzOzPzQzRzSzTzUzVzWzXzYz
218 cdef int out_len = len(result) 2a phqhrhshthuhvhwhxhyhzhAhBh8 Ch9 Dh! Eh7dFh# Gh8dHh$ Ih% Jh' Kh|c}c~c9d!d#dd g e f 3 ( 4 5 ) * + $d%d'd(d)d*d, Lh- Mh. Nh/ Oh: Ph; Qh= Rh? Sh@ Th[ Uh] Vh^ Wh_ Xh` Yh{ Zh| 0h} 1h~ 2hab3hbb4hcb5hdb6heb7hfb8h9h!h#h$h%h'h(h)h*h+h,h-h.h/h:h;h=h?h@hue[h]hve^hwe_hxe`hye{htb|hze}hAe~hBeaibiciubdivbeiwbfixbgiybhi$cii%cji'cki(cli)cmi*cni+coi,cpiqiZgri0gsi1gti9fui!fvi#fwixiyiziAiBiCiDiEiFiGiHiIiJiKiLiMiNiOiPiQiRiSiTiUiViWiXiYiZi0i1i2i3i4i5i6i7i8i9i!i#i$i%i'i(i)i*i+i,i-i.i/i:i;i=i?i@i[i]i^i_i`i{i|i}i~iajbjcjdjejfjgjhjijjjkjljmjnjojpjqjrjsjtjujvjwjxjyjzjAjBjCjDjEjFjGjHjIjJjKjLjMjNjOjPjQjRjSjTjUjVjWjXjYjZj0j1j2j3j4j5j6j7j8j9j!j#j$j%j'j(j)j*j+j,j-j.j/j:j;j=j?j@j[j]j^j_j`j{jyg|j}j~jakbkckdkekfkgkichkikjkBgkklkmknkokpkqkrksktkukvkwkxkykzkCeAkDeBkCkDkEkFkadGkbdHkIkJkKkjcLkzbMkAbNkBbOkPkEeQkCgRkgbSkhbTkibUk6 7 @dFe+d,d-d.d-c.c/d:dEbFbGbHbIbJb$f%fVk3gWkXkYkZk0k1k2k'fGe3kHe4k5k6k7k8k9k!k#k(f$k%k'k(kHg)k*k+kIg)f*f,kJg-k.k/k:k;k=k?k@k[k]k^k_k`kKgjb6c7c{k|k}k~kb KbalLbblMbclNbdlObelPbflQbglRbhlSbilTbjlUbklVbllWbmlXbnlYbolZbpl0bql1brl2bsl3btl4bul5bvl6bwl7bxl8byl9bzl!bAl#bBl$bCl%bDlEl'bFlGlHl+f8c9c,f-f.f/f:f;f=f?fLgMg/cDckbkclbEclcmb;bcdddedh =bi j k l m n o p q ?br s t u (bv w )b@b:cFcnbmcobGcncpb[bfdgdhdx ]by z A B C D E F G ^bH I J K *bL M +b_b;cHcqbocrbIcpcsb`bidjdkdN {bO P Q R S T U V W |bX Y Z 0 ,b1 2 -b}b@fldIlJlKlLlMlNlOlPlQl=cRlIeSlmdTlJeUlVlWlXlYlZl0l1l2l3l4l5l6l7l8l9l!l#l$l%l'l(l)l*l+l,l-l.l/l:l;l=l?l@l[l]l^l_l`l{l|l}l~lambmcmdmemfmgmhmimjmkmlm[f]fKemmndnmompmqmrm.bsm;dtmumvmwmxmymCbzmAm/bBmCmDmEmFmGmHmImNgJmKmLmMmNmOmPmNe~bQmOeacRmSmTmUmVmWmXmYmZm0m1m2m3m4m5m6m7m8m9m!m#m$m%m'm(m)m*m+m,m-m.m/m:m;m=m?m@m[m]m^m_m`m{m|m}m~manbncndnenfngnhninjnknlnmnnnonpnqnrnsntnunvnwnxnynznAnBnCnDnEnFnGnHnInJnKnLnMnNnOnPnQnRnSnTnUnVnWnXnYnZn0n1n2n3n4n5n6n7n8n9n!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~naobocodoeofogohoiojokolomonooopoqorosotouovowoxoyozoAoBoCoDoEoFoGoHoIoJoKoLoMoNoOoPoQoRoSoToUoVoWoXoYoZo0o1o2o3o4o5o6o7o8o9o!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~uavbvcvdvevfvgvhvpdivjvkvqdlvmvnvovpvqvrvsvtvuvvvwvxvyvzvAvBvCvDvEvFvGvHvIvJvKvLvMvNvOvPvQvRvSvTvUvVvWvXvYvZv0v1v2v3v4v5v6v7v8v9v!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~waxbxcxdxexfxgxhxbcixJcjxcckxdclxKcmxLcnxMcoxqcpxrcqxscrxNcsxOctxPcuxQcvxgcwxhcxxtcyxRczxucAxScBxTcCxUcDxVcExvcFxWcGxwcHxxcIxycJxzcKxXcLxYcMxZcNx0cOx1cPx2cQx3cRx4cSx5cTxAcUxBcVxCcWxc XxYxZx0x1x2x3x4x5x6x7x8x9x!x#x:b$x%x'x(x)x*x+x,xfc-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~yazbzczdzezfzgzhzizjzkzlzmznzozpzqzrzsztzuzvzwzxzyzzzAzBzCzDzEzFzGzHzIzJzKzLzMzNzOzPzQzRzSzTzUzVzWzXzYz
219 if out_len == 1: 2a phqhrhshthuhvhwhxhyhzhAhBh8 Ch9 Dh! Eh7dFh# Gh8dHh$ Ih% Jh' Kh|c}c~c9d!d#dd g e f 3 ( 4 5 ) * + $d%d'd(d)d*d, Lh- Mh. Nh/ Oh: Ph; Qh= Rh? Sh@ Th[ Uh] Vh^ Wh_ Xh` Yh{ Zh| 0h} 1h~ 2hab3hbb4hcb5hdb6heb7hfb8h9h!h#h$h%h'h(h)h*h+h,h-h.h/h:h;h=h?h@hue[h]hve^hwe_hxe`hye{htb|hze}hAe~hBeaibiciubdivbeiwbfixbgiybhi$cii%cji'cki(cli)cmi*cni+coi,cpiqiZgri0gsi1gti9fui!fvi#fwixiyiziAiBiCiDiEiFiGiHiIiJiKiLiMiNiOiPiQiRiSiTiUiViWiXiYiZi0i1i2i3i4i5i6i7i8i9i!i#i$i%i'i(i)i*i+i,i-i.i/i:i;i=i?i@i[i]i^i_i`i{i|i}i~iajbjcjdjejfjgjhjijjjkjljmjnjojpjqjrjsjtjujvjwjxjyjzjAjBjCjDjEjFjGjHjIjJjKjLjMjNjOjPjQjRjSjTjUjVjWjXjYjZj0j1j2j3j4j5j6j7j8j9j!j#j$j%j'j(j)j*j+j,j-j.j/j:j;j=j?j@j[j]j^j_j`j{jyg|j}j~jakbkckdkekfkgkichkikjkBgkklkmknkokpkqkrksktkukvkwkxkykzkCeAkDeBkCkDkEkFkadGkbdHkIkJkKkjcLkzbMkAbNkBbOkPkEeQkCgRkgbSkhbTkibUk6 7 @dFe+d,d-d.d-c.c/d:dEbFbGbHbIbJb$f%fVk3gWkXkYkZk0k1k2k'fGe3kHe4k5k6k7k8k9k!k#k(f$k%k'k(kHg)k*k+kIg)f*f,kJg-k.k/k:k;k=k?k@k[k]k^k_k`kKgjb6c7c{k|k}k~kb KbalLbblMbclNbdlObelPbflQbglRbhlSbilTbjlUbklVbllWbmlXbnlYbolZbpl0bql1brl2bsl3btl4bul5bvl6bwl7bxl8byl9bzl!bAl#bBl$bCl%bDlEl'bFlGlHl+f8c9c,f-f.f/f:f;f=f?fLgMg/cDckbkclbEclcmb;bcdddedh =bi j k l m n o p q ?br s t u (bv w )b@b:cFcnbmcobGcncpb[bfdgdhdx ]by z A B C D E F G ^bH I J K *bL M +b_b;cHcqbocrbIcpcsb`bidjdkdN {bO P Q R S T U V W |bX Y Z 0 ,b1 2 -b}b@fldIlJlKlLlMlNlOlPlQl=cRlIeSlmdTlJeUlVlWlXlYlZl0l1l2l3l4l5l6l7l8l9l!l#l$l%l'l(l)l*l+l,l-l.l/l:l;l=l?l@l[l]l^l_l`l{l|l}l~lambmcmdmemfmgmhmimjmkmlm[f]fKemmndnmompmqmrm.bsm;dtmumvmwmxmymCbzmAm/bBmCmDmEmFmGmHmImNgJmKmLmMmNmOmPmNe~bQmOeacRmSmTmUmVmWmXmYmZm0m1m2m3m4m5m6m7m8m9m!m#m$m%m'm(m)m*m+m,m-m.m/m:m;m=m?m@m[m]m^m_m`m{m|m}m~manbncndnenfngnhninjnknlnmnnnonpnqnrnsntnunvnwnxnynznAnBnCnDnEnFnGnHnInJnKnLnMnNnOnPnQnRnSnTnUnVnWnXnYnZn0n1n2n3n4n5n6n7n8n9n!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~naobocodoeofogohoiojokolomonooopoqorosotouovowoxoyozoAoBoCoDoEoFoGoHoIoJoKoLoMoNoOoPoQoRoSoToUoVoWoXoYoZo0o1o2o3o4o5o6o7o8o9o!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~uavbvcvdvevfvgvhvpdivjvkvqdlvmvnvovpvqvrvsvtvuvvvwvxvyvzvAvBvCvDvEvFvGvHvIvJvKvLvMvNvOvPvQvRvSvTvUvVvWvXvYvZv0v1v2v3v4v5v6v7v8v9v!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~waxbxcxdxexfxgxhxbcixJcjxcckxdclxKcmxLcnxMcoxqcpxrcqxscrxNcsxOctxPcuxQcvxgcwxhcxxtcyxRczxucAxScBxTcCxUcDxVcExvcFxWcGxwcHxxcIxycJxzcKxXcLxYcMxZcNx0cOx1cPx2cQx3cRx4cSx5cTxAcUxBcVxCcWxc XxYxZx0x1x2x3x4x5x6x7x8x9x!x#x:b$x%x'x(x)x*x+x,xfc-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~yazbzczdzezfzgzhzizjzkzlzmznzozpzqzrzsztzuzvzwzxzyzzzAzBzCzDzEzFzGzHzIzJzKzLzMzNzOzPzQzRzSzTzUzVzWzXzYz
220 return 28 9 ! 7d# 8d$ % ' |c}c~c9d!d#dd g e f 3 ( 4 5 ) * + $d%d'd(d)d*d, - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdbebfbuevewexeyetbzeAeBeubvbwbxbyb$c%c'c(c)c*c+c,cZg0g1g9f!f#fygCeDeadbdjczbAbBbEegbhbib@dFeEbFbGbHbIbJbGeHeKgjb6c7c+f8c9c,f-f.f/f:f;f=f?fLgMg/cDckbkclbEclcmb;bcdddedh =bi j k l m n o p q ?br s t u (bv w )b@b:cFcnbmcobGcncpb[bfdgdhdx ]by z A B C D E F G ^bH I J K *bL M +b_b;cHcqbocrbIcpcsb`bidjdkdN {bO P Q R S T U V W |bX Y Z 0 ,b1 2 -b}b@fld=cIemdJe[f]fKendCb
221 elif out_len == 2:
222 return result[1] 2a phqhrhshthuhvhwhxhyhzhAhBh8 Ch9 Dh! Eh7dFh# Gh8dHh$ Ih% Jh' Kh|c}c~c9d!d#dd g e f 3 ( 4 5 ) * + $d%d'd(d)d*d, Lh- Mh. Nh/ Oh: Ph; Qh= Rh? Sh@ Th[ Uh] Vh^ Wh_ Xh` Yh{ Zh| 0h} 1h~ 2hab3hbb4hcb5hdb6heb7hfb8h9h!h#h$h%h'h(h)h*h+h,h-h.h/h:h;h=h?h@hue[h]hve^hwe_hxe`hye{htb|hze}hAe~hBeaibiciubdivbeiwbfixbgiybhi$cii%cji'cki(cli)cmi*cni+coi,cpiqiZgri0gsi1gti9fui!fvi#fwixiyiziAiBiCiDiEiFiGiHiIiJiKiLiMiNiOiPiQiRiSiTiUiViWiXiYiZi0i1i2i3i4i5i6i7i8i9i!i#i$i%i'i(i)i*i+i,i-i.i/i:i;i=i?i@i[i]i^i_i`i{i|i}i~iajbjcjdjejfjgjhjijjjkjljmjnjojpjqjrjsjtjujvjwjxjyjzjAjBjCjDjEjFjGjHjIjJjKjLjMjNjOjPjQjRjSjTjUjVjWjXjYjZj0j1j2j3j4j5j6j7j8j9j!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~jakbkckdkekfkgkichkikjkBgkklkmknkokpkqkrksktkukvkwkxkykzkCeAkDeBkCkDkEkFkadGkbdHkIkJkKkLkzbMkAbNkBbOkPkEeQkCgRkgbSkhbTkibUk6 7 Fe+d,d-d.d-c.c/d:dEbFbGbHbIbJb$f%fVk3gWkXkYkZk0k1k2k'f3k4k5k6k7k8k9k!k#k(f$k%k'k(kHg)k*k+kIg)f*f,kJg-k.k/k:k;k=k?k@k[k]k^k_k`kjb6c7c{k|k}k~kb KbalLbblMbclNbdlObelPbflQbglRbhlSbilTbjlUbklVbllWbmlXbnlYbolZbpl0bql1brl2bsl3btl4bul5bvl6bwl7bxl8byl9bzl!bAl#bBl$bCl%bDlEl'bFlGlHl+f8c9c,f-f.f/f:f;f=f?f/cDckbkclbEclcmb;bcdddedh =bi j k l m n o p q ?br s t u (bv w )b@b:cFcnbmcobGcncpb[bfdgdhdx ]by z A B C D E F G ^bH I J K *bL M +b_b;cHcqbocrbIcpcsb`bidjdkdN {bO P Q R S T U V W |bX Y Z 0 ,b1 2 -b}b@fIlJlKlLlMlNlOlPlQlRlSlTlUlVlWlXlYlZl0l1l2l3l4l5l6l7l8l9l!l#l$l%l'l(l)l*l+l,l-l.l/l:l;l=l?l@l[l]l^l_l`l{l|l}l~lambmcmdmemfmgmhmimjmkmlm[f]fKemmnmompmqmrm.bsm;dtmumvmwmxmymzmAm/bBmCmDmEmFmGmHmImNgJmKmLmMmNmOmPmNe~bQmOeacRmSmTmUmVmWmXmYmZm0m1m2m3m4m5m6m7m8m9m!m#m$m%m'm(m)m*m+m,m-m.m/m:m;m=m?m@m[m]m^m_m`m{m|m}m~manbncndnenfngnhninjnknlnmnnnonpnqnrnsntnunvnwnxnynznAnBnCnDnEnFnGnHnInJnKnLnMnNnOnPnQnRnSnTnUnVnWnXnYnZn0n1n2n3n4n5n6n7n8n9n!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~naobocodoeofogohoiojokolomonooopoqorosotouovowoxoyozoAoBoCoDoEoFoGoHoIoJoKoLoMoNoOoPoQoRoSoToUoVoWoXoYoZo0o1o2o3o4o5o6o7o8o9o!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~uavbvcvdvevfvgvhvpdivjvkvqdlvmvnvovpvqvrvsvtvuvvvwvxvyvzvAvBvCvDvEvFvGvHvIvJvKvLvMvNvOvPvQvRvSvTvUvVvWvXvYvZv0v1v2v3v4v5v6v7v8v9v!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~waxbxcxdxexfxgxhxbcixJcjxcckxdclxKcmxLcnxMcoxqcpxrcqxscrxNcsxOctxPcuxQcvxgcwxhcxxtcyxRczxucAxScBxTcCxUcDxVcExvcFxWcGxwcHxxcIxycJxzcKxXcLxYcMxZcNx0cOx1cPx2cQx3cRx4cSx5cTxAcUxBcVxCcWxc XxYxZx0x1x2x3x4x5x6x7x8x9x!x#x:b$x%x'x(x)x*x+x,xfc-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~yazbzczdzezfzgzhzizjzkzlzmznzozpzqzrzsztzuzvzwzxzyzzzAzBzCzDzEzFzGzHzIzJzKzLzMzNzOzPzQzRzSzTzUzVzWzXzYz
223 else:
224 return result[1:] 2a 8 9 ! 7d# 8d$ % ' |c}c~c9d!d#dd e f 3 4 5 ) * + $d%d'd(d)d*d, - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdbebfbgbhbibb KbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b.bNe~bOeacpdqdbcJcccdcKcLcMcqcrcscNcOcPcQcgchctcRcucScTcUcVcvcWcwcxcyczcXcYcZc0c1c2c3c4c5cAcBcCcc :bfc
227cpdef check_or_create_options(type cls, options, str options_description="", bint keep_none=False):
228 """
229 Create the specified options dataclass from a dictionary of options or None.
230 """
231 if options is None: 28 9 ! 7d# 8d$ % ' |c}c~c9d!d#dd g e f 3 ( 4 5 ) * + $d%d'd(d)d*d, - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdbebfbuevewexeyetbzeAeBeubvbwbxbyb$c%c'c(c)c*c+c,c2e3e4eFCGCrd5eJCKCsd6eNCOCtd7eRCSCud8eVCWCvd9eZC0Cwdic!e#e$e%eCeDe'e3C4C5Cadbd6C7C3zjczbAbBb(eEegbhbib8C9C4z!C5z#C$C%C6z'C7z(C6 )C7 *CDgEgVgWgFgGgXgYg^g_g8z+C,C-C.C/C:C;CFe=C?C@C[C]C^C_C`C{C|C}C~CaDbDcDdDeDfDgDhDiDjDkDlDmDnD+doD,dpD-dqD.drD-c.c/dsD:dtDuDvDwDxDEbFbGbHbIbJbyDzDADBDCDDDjCGeHeKIED9z!c!zLIFDMIGDHDIDJDKDLDMD#z$z%z[dNDODPDQDRDSDTDUDVDWDXDYDZDjb6c7c]c#cxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdUdVdWdXdYdb Kb%HLb'HMb(HNb)HOb*HPb+HQb,HRb-HSb.HTb/HUbACVbBCWb:HXb;HYb=HZb?H0b@H1b[H2b]H3b^H4b_H5b`H6b{H7b|H8b}H9b1e!b~H#baI$bbI%bcI'zdI'b8c9c0D1D/cDckbkclbEclcmb;bcdddedh =bi j k l m n o p q ?br s t u (bv w )b@b:cFcnbmcobGcncpb[bfdgdhdx ]by z A B C D E F G ^bH I J K *bL M +b_b;cHcqbocrbIcpcsb`bidjdkdN {bO P Q R S T U V W |bX Y Z 0 ,b1 2 -b}b2D3Dld(z4D5D)z*z+z,z=cIemdJe6D7D8D9D!D#D$D%D'D(D)D*D+D,D-D.D-z/D.z:D/z;D:z=D;z?D=z@D?z[D@z]D^D_D`D{D|D}D~DaEbEcEdEeEfEgEhEiEnC`goCpC{gqCjEKekEndlEmEnEoEpEqE?c@cDb[c.b;d[z]z^z_z`zCb{z/b|z}z~zaAbAcAdAeAWHfAgAhAiAjAkA~baclAmAnAoApAqArAsAtAuAvAwAxAyAzAAABACADAEAFAGAHAIAJAKALAMANAOAPAQARASATAUAVAWAXAYAZA0A1A2A3A4A5A6A7A8A9A!A#A$A%A'A(A)A*A+A,A-A.A/A:A7fLeMeod=drE)esEtEuE;A*e=AwExE?AyEzE+eAEBECE@A,e[AEE-eFEGE.eIE]dJEKErCsCLEMENEOEPE|g^dPeQE}g~gRESETE_d`d/e:e;e=e?e@eUEVE[eXE]eYEZE]A^e^A1E_e2E3E4E_A`e`A6E{e7E8E|e!EtC{d#E$E%E'E(E)E*E+E,E-E{A|dQe.Eah|A/E:E;E=E?E@E[E]E^E_E`E{E|E}E}e~EaF~ecFdFaf}A}d~AaBbBcBdBeBfBgBhB~diBOgjBkBlBbfaebecfdfefffgfZdhfif^cReSe0d_c1djffFgF2dbhchmBcenBoBpBqBrBsBtBuBvBwBdexBPgyBzBABkfnFoFlfqFuCeerFsFtFuFvFwFxFyFzFAFBFBBfeUeCFdhCBDFEFFFGFmfHFIFJFDBnfEBLFofMFNFpfPFgeQFRFvCwCSFTFUFVFWFehheVeXFfhghYFZF0Fiejeqfrfsftfufvf1F2Fwf4Fxf5F6FFByfGB8Fzf9F!F#FHBAfIB%FBf'F(FCf*FxCke+F,F-F.F/F:F;F=F?F@FJBleWe[FhhKB]F^F_F`F{F|F}F~FaGbGcGdGeGfGDfgGhGEfjGkGFfLBmeMBNBOBPBQBRBSBTBUBneVBQgWBXBYBGfoepeHfIfJfKfLf3dMfNf`cXeYe4d{c5dOfmGnG6dihjhZBqe0B1B2B3B4B5B6B7B8B9Bre!BRg#B$B%BPfuGvGQfxGyCseyGzGAGBGCGDGEGFGGGHGIG'Bte0eJGkh(BKGLGMGNGOGlhmh)B*BPGRfQGRGSG+BSf,B-B.B/BUGVGTfWGXGYG:BUf;B0G1GVf2G3G4G=BWf?B6G7GXf8G9G!G@BYf[B$G%GZf'G(G)G]B0f^B+G,G1f-G.G/G_B2f`B;G=G3f?G@G[G{B4f|B}B^G_G5f`G{G|G~B6faCbC~GXHYHbcJczCccdcKcLcMcqcrcscNcOcPcQcgchctcRcucScTcUcVcvcWcwcxcyczcXcYcZc0c1c2c3c4c5cAcBcCcc eccC?dZH0H1H2H3H4H5H6H7H8H9H!H#H$HnhohdCeCfC:bgCaHbHhCeIfIgIhIiIjIkIlIfcmIcHdHeHfHgHhHiHjHkHlHmHnHoHpHqHrHiCsHtHuHvHwHxHyHzHAHBHCHDHEHFHGHHHIHJHKHLHMHNHOHPHQHRHSHTHUHVH
232 if keep_none: 28 9 ! 7d# 8d$ % ' |c}c~c9d!d#dd g e f 3 ( 4 5 ) * + $d%d'd(d)d*d, - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdbebfbuevewexeyetbzeAeBeubvbwbxbyb$c%c'c(c)c*c+c,cFCGCJCKCNCOCRCSCVCWCZC0CCeDeadbdzbAbBbEegbhbib6 7 DgEgVgWgFgGgXgYg^g_g8z+d,d-d.d-c.c/d:dEbFbGbHbIbJbyDzDADBDCDDDGeHeKIED9z!c!zFDGDHDIDJDKDLDMD#z$z%z[dNDODPDQDRDSDTDUDVDWDXDYDZDjb6c7c#c8c9c0D1D/cDckbkclbEclcmb;bcdddedh =bi j k l m n o p q ?br s t u (bv w )b@b:cFcnbmcobGcncpb[bfdgdhdx ]by z A B C D E F G ^bH I J K *bL M +b_b;cHcqbocrbIcpcsb`bidjdkdN {bO P Q R S T U V W |bX Y Z 0 ,b1 2 -b}b2D3Dld(z4D5D)z*z+z,z=cIemdJe-z/D.z:D/z;D:z=D;z?D=z@D?z[D@z]D`g{gKekEnd;d[z]z^z_z`zCb{z/b|z}z~zaAbAcAdAeAWHfAgAhAiA~baclAmAnAoApAqArAsAtAuAvAwAxAyAzAAABACADAEAFAGAHAIAJAKALAMANAOAPAQARASATAUAVAWAXAYAZA0A1A2A3A4A5A6A7A8A9A!A#A$A%A'A(A)A*A+A,A-A.A/A:AodrEsEtEuE;A=AwExE?AyEzEAEBECE@A[AEEFEGEIE]dJEKErCsCLEMENEOEPE|g^dPeQE}g~gRESETE_d`dUEVEXEYEZE]A^A1E2E3E4E_A`A6E7E8E!EtC{d#E$E%E'E(E)E*E+E,E-E{A|dQe.Eah|A/E:E;E=E?E@E[E]E^E_E`E{E|E}E~EaFcFdF}A}d~AaBbBcBdBeBfBgBhB~diBOgjBkBlBaebeReSe0d1dfFgFbhchmBcenBoBpBqBrBsBtBuBvBwBdexBPgyBzBABnFoFqFuCeerFsFtFuFvFwFxFyFzFAFBFBBfeUeCFdhCBDFEFFFGFHFIFJFDBEBLFMFNFPFgeQFRFvCwCSFTFUFVFWFehheVeXFfhghYFZF0Fieje1F2F4F5F6FFBGB8F9F!F#FHBIB%F'F(F*FxCke+F,F-F.F/F:F;F=F?F@FJBleWe[FhhKB]F^F_F`F{F|F}F~FaGbGcGdGeGfGgGhGjGkGLBmeMBNBOBPBQBRBSBTBUBneVBQgWBXBYBoepeXeYe4d5dmGnGihjhZBqe0B1B2B3B4B5B6B7B8B9Bre!BRg#B$B%BuGvGxGyCseyGzGAGBGCGDGEGFGGGHGIG'Bte0eJGkh(BKGLGMGlhmh)B*BPGQGRGSG+B,B-B.B/BUGVGWGXGYG:B;B0G1G2G3G4G=B?B6G7G8G9G!G@B[B$G%G'G(G)G]B^B+G,G-G.G/G_B`B;G=G?G@G[G{B|B}B^G_G`G{G|G~BaCbC~GJcc eccC?ddCeCfCgCaHbHhCfceHfHgHhHkHlHmHnHqHiCsHvHwHyHzHAHBHCHDHEHFHGHHHIHJHKHLHMHNHOHPHQHRHSHTHUHVH
233 return options 2|c}c~cd g e f 3 ( 4 5 GeHe0D1D2D3D(z4D5D)z*z+z,z=cIemdJe-z/D.z:D/z;D:z=D;z?D=z@D?z[D@z]D`g{gkECbrExEzEIE]dJEKErCsCLEMENEOEPE|g^dPeQE}g~gRESETEXE1EtCuCGFPFgeQFRFvCwCSFTFUFVFWFehheVeXFfhghYFZF0F4F8FxCyCPGVG1G7G%G,G=G_GyHzHAHBHCHDHEHFHGHHHIHJHKHLHMHNHOHPHQHRHSHTHUHVH
234 return cls() 28 9 ! 7d# 8d$ % ' |c}c~c9d!d#dd g e f 3 ( 4 5 ) * + $d%d'd(d)d*d, - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdbebfbuevewexeyetbzeAeBeubvbwbxbyb$c%c'c(c)c*c+c,cFCGCJCKCNCOCRCSCVCWCZC0CCeDeadbdzbAbBbEegbhbib6 7 DgEgVgWgFgGgXgYg^g_g8z+d,d-d.d-c.c/d:dEbFbGbHbIbJbyDzDADBDCDDDKIED9z!c!zFDGDHDIDJDKDLDMD#z$z%z[dNDODPDQDRDSDTDUDVDWDXDYDZDjb6c7c#c8c9c/cDckbkclbEclcmb;bcdddedh =bi j k l m n o p q ?br s t u (bv w )b@b:cFcnbmcobGcncpb[bfdgdhdx ]by z A B C D E F G ^bH I J K *bL M +b_b;cHcqbocrbIcpcsb`bidjdkdN {bO P Q R S T U V W |bX Y Z 0 ,b1 2 -b}bld=cKend;d[z]z^z_z`zCb{z/b|z}z~zaAbAcAdAeAWHfAgAhAiA~baclAmAnAoApAqArAsAtAuAvAwAxAyAzAAABACADAEAFAGAHAIAJAKALAMANAOAPAQARASATAUAVAWAXAYAZA0A1A2A3A4A5A6A7A8A9A!A#A$A%A'A(A)A*A+A,A-A.A/A:AodsEtEuE;A=AwE?AyEAEBECE@A[AEEFEGErCsC|g~g_d`dUEVEYEZE]A^A2E3E4E_A`A6E7E8E!EtC{d#E$E%E'E(E)E*E+E,E-E{A|dQe.Eah|A/E:E;E=E?E@E[E]E^E_E`E{E|E}E~EaFcFdF}A}d~AaBbBcBdBeBfBgBhB~diBOgjBkBlBaebeReSe0d1dfFgFbhchmBcenBoBpBqBrBsBtBuBvBwBdexBPgyBzBABnFoFqFuCeerFsFtFuFvFwFxFyFzFAFBFBBfeUeCFdhCBDFEFFFHFIFJFDBEBLFMFNFvCwCehghieje1F2F5F6FFBGB9F!F#FHBIB%F'F(F*FxCke+F,F-F.F/F:F;F=F?F@FJBleWe[FhhKB]F^F_F`F{F|F}F~FaGbGcGdGeGfGgGhGjGkGLBmeMBNBOBPBQBRBSBTBUBneVBQgWBXBYBoepeXeYe4d5dmGnGihjhZBqe0B1B2B3B4B5B6B7B8B9Bre!BRg#B$B%BuGvGxGyCseyGzGAGBGCGDGEGFGGGHGIG'Bte0eJGkh(BKGLGMGlhmh)B*BQGRGSG+B,B-B.B/BUGWGXGYG:B;B0G2G3G4G=B?B6G8G9G!G@B[B$G'G(G)G]B^B+G-G.G/G_B`B;G?G@G[G{B|B}B^G`G{G|G~BaCbC~GJcc eccC?ddCeCfCgCaHbHhCfceHfHgHhHkHlHmHnHqHiCsHvHwH
235 elif isinstance(options, cls): 28 9 ! # $ % ' d g e f 3 ( 4 5 ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdbebfbtbubvbwbxbyb2e3e4erd5esd6etd7eud8evd9ewdic!e#e$e%e'e3C4C5Cadbd6C7C3zjczbAbBb(egbhbib8C9C4z!C5z#C$C%C6z'C7z(C6 )C7 *CDgEgVgWgFgGgXgYg^g_g8z+C,C-C.C/C:C;CFe=C?C@C[C]C^C_C`C{C|C}C~CaDbDcDdDeDfDgDhDiDjDkDlDmDnDoDpDqDrD-c.csDtDuDvDwDxDEbFbGbHbIbJbjC9z!c!zLIMI#z$z%z[djb]c#cxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdUdVdWdXdYdb Kb%HLb'HMb(HNb)HOb*HPb+HQb,HRb-HSb.HTb/HUbACVbBCWb:HXb;HYb=HZb?H0b@H1b[H2b]H3b^H4b_H5b`H6b{H7b|H8b}H9b1e!b~H#baI$bbI%bcI'zdI'bld(z)z*z+z,z=cmd6D7D8D9D!D#D$D%D'D(D)D*D+D,D-D.D-z.z/z:z;z=z?z@z^D_D`D{D|D}D~DaEbEcEdEeEfEgEhEiEnC`goCpC{gqCjEndlEmEnEoEpEqE?c@cDb[c.bCbjAkA7fLeMeod=d)e*e+e,e-e.e]d^d_d`d/e:e;e=e?e@e[e]e^e_e`e{e|e{d|d}e~eaf}d~dbfaebecfdfefffgfZdhfif^c_cjf2dcedekflfeefemfnfofpfgeheiejeqfrfsftfufvfwfxfyfzfAfBfCfkeleDfEfFfmeneGfoepeHfIfJfKfLf3dMfNf`c{cOf6dqerePfQfseteNGOGlhmhRfSfTfUfVfWfXfYfZf0f1f2f3f4f5f6fXHYHbczCccdcKcLcMcqcrcscNcOcPcQcgchctcRcucScTcUcVcvcWcwcxcyczcXcYcZc0c1c2c3c4c5cAcBcCcZH0H1H2H3H4H5H6H7H8H9H!H#H$Hnhoh:beIfIgIhIiIjIkIlImIcHdHiHjHoHpHrHiCtHuHxH
236 return options 28 9 ! # $ % ' d g e f 3 ( 4 5 ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdbebfbtbubvbwbxbyb2e3e4erd5esd6etd7eud8evd9ewdic!e#e$e%e'e3C4C5Cadbd6C7C3zjczbAbBb(egbhbib8C9C4z!C5z#C$C%C6z'C7z(C6 )C7 *CVgWgXgYg8z+C,C-C.C/C:C;CFe=C?C@C[C]C^C_C`C{C|C}C~CaDbDcDdDeDfDgDhDiDjDkDlDmDnDoDpDqDrD-c.csDtDuDvDwDxDEbFbGbHbIbJb9z!c!zLIMI#z$z%z[djb]c#cxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdUdVdWdXdYdb Kb%HLb'HMb(HNb)HOb*HPb+HQb,HRb-HSb.HTb/HUbACVbBCWb:HXb;HYb=HZb?H0b@H1b[H2b]H3b^H4b_H5b`H6b{H7b|H8b}H9b1e!b~H#baI$bbI%bcI'zdI'bld(z)z*z+z,z=cmd6D7D8D9D!D#D$D%D'D(D)D*D+D,D-D.D-z.z/z:z;z=z?z@z^D_D`D{D|D}D~DaEbEcEdEeEfEgEhEiEjEndlEmEnEoEpEqE?c@cDb[c.bCbjAkA7fLeMeod=d)e*e+e,e-e.e]d^d_d`d/e:e;e=e?e@e[e]e^e_e`e{e|e{d|d}e~eaf}d~dbfaebecfdfefffgfZdhfif^c_cjf2dcedekflfeefemfnfofpfgeheiejeqfrfsftfufvfwfxfyfzfAfBfCfkeleDfEfFfmeneGfoepeHfIfJfKfLf3dMfNf`c{cOf6dqerePfQfseteNGOGRfSfTfUfVfWfXfYfZf0f1f2f3f4f5f6fXHYHbczCccdcKcLcMcqcrcscNcOcPcQcgchctcRcucScTcUcVcvcWcwcxcyczcXcYcZc0c1c2c3c4c5cAcBcCcZH0H1H2H3H4H5H6H7H8H9H!H#H$HnhoheIfIgIhIiIjIkIlImIcHdHiHjHoHpHrHiCtHuHxH
237 elif isinstance(options, dict): 2DgEgFgGg^g_gjC]cnC`goCpC{gqClhmh:b
238 return cls(**options) 2DgEgFgGg^g_g]cnC`goCpC{gqClhmh:b
239 else:
240 raise TypeError( 2jC
241 f"The {options_description} must be provided as an object " 2jC
242 f"of type {cls.__name__} or as a dict with valid {options_description}. " 2jC
243 f"The provided object is '{options}'." 2jC
244 )
247def _handle_boolean_option(option: bool) -> str:
248 """
249 Convert a boolean option to a string representation.
250 """
251 return "true" if bool(option) else "false" 2b KbLbMbNbObPbQbRbSbTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!b#b$b%b'b.bqcrcscucyczcAcBcCc%Y
254def precondition(checker: Callable[..., None], str what="") -> Callable:
255 """
256 A decorator that adds checks to ensure any preconditions are met.
258 Args:
259 checker: The function to call to check whether the preconditions are met. It has
260 the same signature as the wrapped function with the addition of the keyword argument `what`.
261 what: A string that is passed in to `checker` to provide context information.
263 Returns:
264 Callable: A decorator that creates the wrapping.
265 """
267 def outer(wrapped_function): 2kC
268 """
269 A decorator that actually wraps the function for checking preconditions.
270 """
272 @functools.wraps(wrapped_function) 2kC
273 def inner(*args, **kwargs):
274 """
275 Check preconditions and if they are met, call the wrapped function.
276 """
277 checker(*args, **kwargs, what=what) 2kC
278 result = wrapped_function(*args, **kwargs) 2kC
280 return result 2kC
282 return inner 2kC
284 return outer 2kC
287def is_sequence(obj):
288 """
289 Check if the given object is a sequence (list or tuple).
290 """
291 return isinstance(obj, Sequence) 26 7 !cjb]cxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdUdVdWdXdYdACBCzCgchctcvcwcxcnhoh
294def is_nested_sequence(obj):
295 """
296 Check if the given object is a nested sequence (list or tuple with atleast one list or tuple element).
297 """
298 return is_sequence(obj) and any(is_sequence(elem) for elem in obj) 2gchc
301@functools.lru_cache
302def get_binding_version():
303 try:
304 major_minor = importlib.metadata.version("cuda-bindings").split(".")[:2]
305 except importlib.metadata.PackageNotFoundError:
306 major_minor = importlib.metadata.version("cuda-python").split(".")[:2]
307 return tuple(int(v) for v in major_minor)
310class Transaction:
311 """
312 A context manager for transactional operations with undo capability.
314 The Transaction class allows you to register undo actions (callbacks) that will be executed
315 if the transaction is not committed before exiting the context. This is useful for managing
316 resources or operations that need to be rolled back in case of errors or early exits.
318 Usage:
319 with Transaction() as txn:
320 txn.append(some_cleanup_function, arg1, arg2)
321 # ... perform operations ...
322 txn.commit() # Disarm undo actions; nothing will be rolled back on exit
324 Methods:
325 append(fn, *args, **kwargs): Register an undo action to be called on rollback.
326 commit(): Disarm all undo actions; nothing will be rolled back on exit.
327 """
328 def __init__(self):
329 self._stack = ExitStack() 2?c@cDb[c
330 self._entered = False 2?c@cDb[c
332 def __enter__(self):
333 self._stack.__enter__() 2?c@cDb[c
334 self._entered = True 2?c@cDb[c
335 return self 2?c@cDb[c
337 def __exit__(self, exc_type, exc, tb):
338 # If exit callbacks remain, they'll run in LIFO order.
339 self._entered = False 2?c@cDb[c
340 return self._stack.__exit__(exc_type, exc, tb) 2?c@cDb[c
342 def append(self, fn, /, *args, **kwargs):
343 """
344 Register an undo action (runs if the with-block exits without commit()).
345 Values are bound now via partial so late mutations don't bite you.
346 """
347 if not self._entered: 2?c@cDb[c
348 raise RuntimeError("Transaction must be entered before append()")
349 self._stack.callback(partial(fn, *args, **kwargs)) 2?c@cDb[c
351 def commit(self):
352 """
353 Disarm all undo actions. After this, exiting the with-block does nothing.
354 """
355 # pop_all() empties this stack so no callbacks are triggered on exit.
356 self._stack.pop_all() 2?c@cDb[c
359# Track whether we've already warned about fork method
360_fork_warning_checked = False
363def reset_fork_warning():
364 """Reset the fork warning check flag for testing purposes.
366 This function is intended for use in tests to allow multiple test runs
367 to check the warning behavior.
368 """
369 global _fork_warning_checked
370 _fork_warning_checked = False 27fLeMeodSg=d
373cdef inline tuple _read_fill_ptr(const char* ptr, Py_ssize_t width):
374 """Extract (value, element_size) from a raw pointer of known width."""
375 cdef unsigned int val
376 if width == 1: 2+g,g-g.g/g:g;g=g?g@g[g]g2zDckbkclbEclcmbh =bi j k l m n o p q ?br s t u (bv w )b@bFcnbmcobGcncpbx ]by z A B C D E F G ^bH I J K *bL M +b_bHcqbocrbIcpcsbN {bO P Q R S T U V W |bX Y Z 0 ,b1 2 -b}b
377 val = (<uint8_t*>ptr)[0] 2kbi j k r s nby z A H I qbO P Q X Y
378 elif width == 2:
379 val = (<uint16_t*>ptr)[0] 2+g-g/g;g?g[gkclbl m n t u (bmcobB C D J K *bocrbR S T Z 0 ,b
380 elif width == 4:
381 val = (<uint32_t*>ptr)[0] 2,g.g:g=g@g]glcmbh o p q v w )bncpbx E F G L M +bpcsbN U V W 1 2 -b
382 else:
383 raise ValueError(f"value must be 1, 2, or 4 bytes, got {width}") 22zDcEc=b?b@bFcGc]b^b_bHcIc{b|b}b
384 return (val, <unsigned int>width) 2+g,g-g.g/g:g;g=g?g@g[g]gkbkclblcmbh i j k l m n o p q r s t u (bv w )bnbmcobncpbx y z A B C D E F G H I J K *bL M +bqbocrbpcsbN O P Q R S T U V W X Y Z 0 ,b1 2 -b
387cpdef tuple _parse_fill_value(value):
388 """Parse a fill/memset value into (raw_value, element_size).
390 Parameters
391 ----------
392 value : int or buffer-protocol object
393 - int: Must be in range [0, 256). Treated as 1-byte fill.
394 - bytes or buffer-protocol: Must be 1, 2, or 4 bytes.
396 Returns
397 -------
398 tuple of (int, int)
399 (raw_value, element_size) where element_size is 1, 2, or 4.
401 Raises
402 ------
403 OverflowError
404 If int value is outside [0, 256).
405 TypeError
406 If value is not an int and does not support the buffer protocol.
407 ValueError
408 If value byte length is not 1, 2, or 4.
409 """
410 cdef uint8_t byte_val
411 cdef Py_buffer buf
413 if isinstance(value, int): 2d g e f $c%c'c(c)c*c+c,cHCIC+g,gLCMC-g.gPCQC/g:gTCUC;g=gXCYC?g@g1C2C[g]g2z/cDckbkclbEclcmb;bcdddedh =bi j k l m n o p q ?br s t u (bv w )b@b:cFcnbmcobGcncpb[bfdgdhdx ]by z A B C D E F G ^bH I J K *bL M +b_b;cHcqbocrbIcpcsb`bidjdkdN {bO P Q R S T U V W |bX Y Z 0 ,b1 2 -b}bvEDEHE}gWE0E5E9EahbFeFOgZdTe2dbhchhFiFjFkFlFmFPgpFdhKFOFfh3F7F$F)FhhiGlGQg3dZe6dihjhoGpGqGrGsGtGRgwGkhTGZG5G#G*G:G]G}G
414 byte_val = value 2d g e f $c%c'c(c)c*c+c,cHCICLCMCPCQCTCUCXCYC1C2C;bcddded[bfdgdhd`bidjdkdvEDEHE}gWE0E5E9EahbFeFOgZdTe2dbhchhFiFjFkFlFmFPgpFdhKFOFfh3F7F$F)FhhiGlGQg3dZe6dihjhoGpGqGrGsGtGRgwGkhTGZG5G#G*G:G]G}G
415 return (<unsigned int>byte_val, <unsigned int>1) 2d g e f $c%c'c(c)c*c+c,cHCICLCMCPCQCTCUCXCYC1C2C;b[b`bvEDEHE}gWE0E5E9EahbFeFOgZdTe2dbhchhFiFjFkFlFmFPgpFdhKFOFfh3F7F$F)FhhiGlGQg3dZe6dihjhoGpGqGrGsGtGRgwGkhTGZG5G#G*G:G]G}G
417 if isinstance(value, bytes): 2+g,g-g.g/g:g;g=g?g@g[g]g2z/cDckbkclbEclcmbh =bi j k l m n o p q ?br s t u (bv w )b@b:cFcnbmcobGcncpbx ]by z A B C D E F G ^bH I J K *bL M +b_b;cHcqbocrbIcpcsbN {bO P Q R S T U V W |bX Y Z 0 ,b1 2 -b}b
418 return _read_fill_ptr(<const char*><bytes>value, <Py_ssize_t>len(value)) 2+g,g-g.g/g:g;g=g?g@g[g]g2zDckbkclbEclcmbFcnbmcobGcncpbHcqbocrbIcpcsb
420 if PyObject_GetBuffer(value, &buf, PyBUF_SIMPLE) != 0: 2/ch =bi j k l m n o p q ?br s t u (bv w )b@b:cx ]by z A B C D E F G ^bH I J K *bL M +b_b;cN {bO P Q R S T U V W |bX Y Z 0 ,b1 2 -b}b
421 raise TypeError(
422 f"value must be an int or support the buffer protocol, "
423 f"got {type(value).__name__}"
424 )
425 try: 2h =bi j k l m n o p q ?br s t u (bv w )b@bx ]by z A B C D E F G ^bH I J K *bL M +b_bN {bO P Q R S T U V W |bX Y Z 0 ,b1 2 -b}b
426 return _read_fill_ptr(<const char*>buf.buf, buf.len) 2h =bi j k l m n o p q ?br s t u (bv w )b@bx ]by z A B C D E F G ^bH I J K *bL M +b_bN {bO P Q R S T U V W |bX Y Z 0 ,b1 2 -b}b
427 finally:
428 PyBuffer_Release(&buf)
431def check_multiprocessing_start_method():
432 """Check if multiprocessing start method is 'fork' and warn if so."""
433 global _fork_warning_checked
434 if _fork_warning_checked: 2oIpI4z5zqIrI6z7z6 7 DgEgVgWgFgGgXgYgsItI2YuI@d3YvI4Y5YwI6Y7YxI8Y9YyI!Y#YzI$YAIBICIDIEIFI+d,d-d.d-c.c/d:dGIHIIIJIEbFbGbHbIbJb7fLeMeodSg=d
435 return 2oIpI4z5zqIrI6z7z6 7 DgEgVgWgFgGgXgYgsItI2YuI@d3YvI4Y5YwI6Y7YxI8Y9YyI!Y#YzI$YAIBICIDIEIFI+d,d-d.d-c.c/d:dGIHIIIJIEbFbGbHbIbJb=d
436 _fork_warning_checked = True 2@d7fLeMeodSg=d
438 # Common warning message parts
439 common_message = (
440 "CUDA does not support. Forked subprocesses exhibit undefined behavior, " 2@d7fLeMeodSg=d
441 "including failure to initialize CUDA contexts and devices. Set the start method "
442 "to 'spawn' before creating processes that use CUDA. "
443 "Use: multiprocessing.set_start_method('spawn')"
444 )
446 try: 2@d7fLeMeodSg=d
447 start_method = multiprocessing.get_start_method() 2@d7fLeMeodSg=d
448 if start_method == "fork": 2@d7fLeMeod=d
449 message = f"multiprocessing start method is 'fork', which {common_message}" 2LeMeod=d
450 warnings.warn(message, UserWarning, stacklevel=3) 2LeMeod=d
451 except RuntimeError: 2Sg
452 # get_start_method() can raise RuntimeError if start method hasn't been set
453 # In this case, default is 'fork' on Linux, so we should warn
454 if platform.system() == "Linux": 2Sg
455 message = (
456 f"multiprocessing start method is not set and defaults to 'fork' on Linux, " 2Sg
457 f"which {common_message}"
458 )
459 warnings.warn(message, UserWarning, stacklevel=3) 2Sg