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

1# SPDX-FileCopyrightText: Copyright (c) 2024-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. 

2# 

3# SPDX-License-Identifier: Apache-2.0 

4  

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 

15  

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 

22  

23from cuda.bindings.nvvm import nvvmError 

24from cuda.bindings.nvjitlink import nvJitLinkError 

25  

26from cpython.buffer cimport PyObject_GetBuffer, PyBuffer_Release, Py_buffer, PyBUF_SIMPLE 

27  

28from cuda.bindings cimport cynvrtc, cynvvm, cynvjitlink 

29  

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 

32  

33  

34class CUDAError(Exception): 

35 pass 

36  

37  

38class NVRTCError(CUDAError): 

39 pass 

40  

41  

42  

43ComputeCapability = namedtuple("ComputeCapability", ("major", "minor")) 

44  

45  

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

62  

63  

64def _reduce_3_tuple(t: tuple): 

65 return t[0] * t[1] * t[2] 2]c

66  

67  

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

72  

73  

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

80  

81  

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

98  

99  

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

106  

107  

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

123  

124  

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

132  

133  

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

151  

152  

153cdef object _RUNTIME_SUCCESS = runtime.cudaError_t.cudaSuccess 

154cdef object _NVRTC_SUCCESS = nvrtc.nvrtcResult.NVRTC_SUCCESS 

155  

156  

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()}") 

174  

175  

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}") 

191  

192  

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

203  

204  

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}") 

214  

215  

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

225  

226  

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 ) 

245  

246  

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

252  

253  

254def precondition(checker: Callable[..., None], str what="") -> Callable: 

255 """ 

256 A decorator that adds checks to ensure any preconditions are met. 

257  

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. 

262  

263 Returns: 

264 Callable: A decorator that creates the wrapping. 

265 """ 

266  

267 def outer(wrapped_function): 2kC

268 """ 

269 A decorator that actually wraps the function for checking preconditions. 

270 """ 

271  

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

279  

280 return result 2kC

281  

282 return inner 2kC

283  

284 return outer 2kC

285  

286  

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

292  

293  

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

299  

300  

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) 

308  

309  

310class Transaction: 

311 """ 

312 A context manager for transactional operations with undo capability. 

313  

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. 

317  

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 

323  

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

331  

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

336  

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

341  

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

350  

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

357  

358  

359# Track whether we've already warned about fork method 

360_fork_warning_checked = False 

361  

362  

363def reset_fork_warning(): 

364 """Reset the fork warning check flag for testing purposes. 

365  

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

371  

372  

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

385  

386  

387cpdef tuple _parse_fill_value(value): 

388 """Parse a fill/memset value into (raw_value, element_size). 

389  

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. 

395  

396 Returns 

397 ------- 

398 tuple of (int, int) 

399 (raw_value, element_size) where element_size is 1, 2, or 4. 

400  

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 

412  

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

416  

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

419  

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) 

429  

430  

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

437  

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 ) 

445  

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