Coverage for cuda / core / _utils / cuda_utils.pyx: 93.67%

237 statements  

« prev     ^ index     » next       coverage.py v7.13.5, created at 2026-04-29 01:27 +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 multiprocessing 

8import platform 

9import warnings 

10from collections import namedtuple 

11from collections.abc import Sequence 

12from contextlib import ExitStack 

13from typing import Callable 

14  

15try: 

16 from cuda.bindings import driver, nvrtc, runtime 

17except ImportError: 

18 from cuda import cuda as driver 

19 from cuda import cudart as runtime 

20 from cuda import nvrtc 

21  

22from cuda.bindings.nvvm import nvvmError 

23from cuda.bindings.nvjitlink import nvJitLinkError 

24  

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

26  

27from cuda.bindings cimport cynvrtc, cynvvm, cynvjitlink 

28  

29from cuda.core._utils.driver_cu_result_explanations import DRIVER_CU_RESULT_EXPLANATIONS 

30from cuda.core._utils.runtime_cuda_error_explanations import RUNTIME_CUDA_ERROR_EXPLANATIONS 

31  

32  

33class CUDAError(Exception): 

34 pass 

35  

36  

37class NVRTCError(CUDAError): 

38 pass 

39  

40  

41  

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

43  

44  

45def cast_to_3_tuple(label, cfg): 

46 cfg_orig = cfg 29 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbybzbAbBbCbDb,e-e.e?gid/eogjd:ekd;eld=emd?end2b@e[e]e^e_e`eEbFbGb{eHbIbJbKbLbd e h f g 4 cb5 6 dbebfbgbhbMbibjbkblbmbnb7 8 DcEcfeobodeA@g^f:C[gFcpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdgeheNb|e}e]g~eaf^gbfcfieje6ekeledfefffgfhfifjfkflfpgmfnfqgofpfmene7eqfrfsfoepetfqereufvfwfxfyfRdzfAfrg9csgtg8e9eugvgwgxgygSd!czg!eTdAgBgCgBfUdseteCfDfueve#eEfFf_gGfHfwexe$eyezeIfJfKfLfMfNfOfPfQfDgRfSfEgTfUfAeBe%eVfWfXfCeDeYfEeFeZf0f1f2f3fVd4f5fFg#cGgHg'e(eIgJgKgLgMgWd$cNg)eXdOgPgQg6fYdGeHe7f8fIeJe*e9f!f`g#f$f{g%f'f|g(f)f}g*f+f~g,f-fah.f/fbh:f;fch'g5C6CfA;C(g

47 if isinstance(cfg, int): 2a 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbybzbAbBbCbDb,e-e.e?gid/eogjd:ekd;eld=emd?end2b@e[e]e^e_e`eEbFbGb{eHbIbJbKbLbd e h f g 4 cb5 6 dbebfbgbhbMbibjbkblbmbnb7 8 DcEcfeobodeA@g^f:C[gFcpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdgeheNb|e}e]g~eaf^gbfcfieje6ekeledfefffgfhfifjfkflfpgmfnfqgofpfmene7eqfrfsfoepetfqereufvfwfxfyfRdzfAfrg9csgtg8e9eugvgwgxgygSd!czg!eTdAgBgCgBfUdseteCfDfueve#eEfFf_gGfHfwexe$eyezeIfJfKfLfMfNfOfPfQfDgRfSfEgTfUfAeBe%eVfWfXfCeDeYfEeFeZf0f1f2f3fVd4f5fFg#cGgHg'e(eIgJgKgLgMgWd$cNg)eXdOgPgQg6fYdGeHe7f8fIeJe*e9f!f`g#f$f{g%f'f|g(f)f}g*f+f~g,f-fah.f/fbh:f;fch'g5C6CfA;C(g

48 cfg = (cfg,) 29 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbybzbAbBbCbDb,e-e.e/e:e;e=e?e2b@e[e]e^e_e`eEbFbGb{eHbIbJbKbLbd e h f g 4 cb5 6 dbebfbgbhbMbibjbkblbmbnb7 8 DcEcfeobodeA^f:CpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdgeheNb|e}e~eafbfcfiejekeledfefffgfhfifjfkflfmfnfofpfmeneqfrfsfoepetfqereufvfwfxfyfRdzfAf9c!cBfUdseteCfDfueveEfFfGfHfwexeyezeIfJfKfLfMfNfOfPfQfRfSfTfUfAeBeVfWfXfCeDeYfEeFeZf0f1f2f3fVd4f5f#c$c6fYdGeHe7f8fIeJe9f!f#f$f%f'f(f)f*f+f,f-f.f/f:f;f'g5C

49 else: 

50 common = "must be an int, or a tuple with up to 3 ints" 2?gidogjdkdldmdnd@g^f[gFc]g^g6epgqg7erg9csgtg8e9eugvgwgxgygSd!czg!eTdAgBgCg#e_g$eDgEg%eFg#cGgHg'e(eIgJgKgLgMgWd$cNg)eXdOgPgQg*e`g{g|g}g~gahbhch'g6CfA;C(g

51 if not isinstance(cfg, tuple): 2?gidogjdkdldmdnd@g^f[gFc]g^g6epgqg7erg9csgtg8e9eugvgwgxgygSd!czg!eTdAgBgCg#e_g$eDgEg%eFg#cGgHg'e(eIgJgKgLgMgWd$cNg)eXdOgPgQg*e`g{g|g}g~gahbhch'g6CfA;C(g

52 raise ValueError(f"{label} {common} (got {type(cfg)})") 2;C

53 if len(cfg) > 3: 2?gidogjdkdldmdnd@g^f[gFc]g^g6epgqg7erg9csgtg8e9eugvgwgxgygSd!czg!eTdAgBgCg#e_g$eDgEg%eFg#cGgHg'e(eIgJgKgLgMgWd$cNg)eXdOgPgQg*e`g{g|g}g~gahbhch'g6CfA(g

54 raise ValueError(f"{label} {common} (got tuple with length {len(cfg)})") 26C

55 if any(not isinstance(val, int) for val in cfg): 2?gidogjdkdldmdnd@g^f[gFc]g^g6epgqg7erg9csgtg8e9eugvgwgxgygSd!czg!eTdAgBgCg#e_g$eDgEg%eFg#cGgHg'e(eIgJgKgLgMgWd$cNg)eXdOgPgQg*e`g{g|g}g~gahbhch'gfA(g

56 raise ValueError(f"{label} {common} (got {cfg})") 2fA

57 if any(val < 1 for val in cfg): 29 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbybzbAbBbCbDb,e-e.e?gid/eogjd:ekd;eld=emd?end2b@e[e]e^e_e`eEbFbGb{eHbIbJbKbLbd e h f g 4 cb5 6 dbebfbgbhbMbibjbkblbmbnb7 8 DcEcfeobodeA@g^f:C[gFcpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdgeheNb|e}e]g~eaf^gbfcfieje6ekeledfefffgfhfifjfkflfpgmfnfqgofpfmene7eqfrfsfoepetfqereufvfwfxfyfRdzfAfrg9csgtg8e9eugvgwgxgygSd!czg!eTdAgBgCgBfUdseteCfDfueve#eEfFf_gGfHfwexe$eyezeIfJfKfLfMfNfOfPfQfDgRfSfEgTfUfAeBe%eVfWfXfCeDeYfEeFeZf0f1f2f3fVd4f5fFg#cGgHg'e(eIgJgKgLgMgWd$cNg)eXdOgPgQg6fYdGeHe7f8fIeJe*e9f!f`g#f$f{g%f'f|g(f)f}g*f+f~g,f-fah.f/fbh:f;fch'g5C(g

58 plural_s = "" if len(cfg) == 1 else "s" 2^f5C(g

59 raise ValueError(f"{label} value{plural_s} must be >= 1 (got {cfg_orig})") 2^f5C(g

60 return cfg + (1,) * (3 - len(cfg)) 29 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbybzbAbBbCbDb,e-e.e?gid/eogjd:ekd;eld=emd?end2b@e[e]e^e_e`eEbFbGb{eHbIbJbKbLbd e h f g 4 cb5 6 dbebfbgbhbMbibjbkblbmbnb7 8 DcEcfeobodeA@g^f:C[gFcpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdgeheNb|e}e]g~eaf^gbfcfieje6ekeledfefffgfhfifjfkflfpgmfnfqgofpfmene7eqfrfsfoepetfqereufvfwfxfyfRdzfAfrg9csgtg8e9eugvgwgxgygSd!czg!eTdAgBgCgBfUdseteCfDfueve#eEfFf_gGfHfwexe$eyezeIfJfKfLfMfNfOfPfQfDgRfSfEgTfUfAeBe%eVfWfXfCeDeYfEeFeZf0f1f2f3fVd4f5fFg#cGgHg'e(eIgJgKgLgMgWd$cNg)eXdOgPgQg6fYdGeHe7f8fIeJe*e9f!f`g#f$f{g%f'f|g(f)f}g*f+f~g,f-fah.f/fbh:f;fch'g

61  

62  

63cdef int HANDLE_RETURN(cydriver.CUresult err) except?-1 nogil: 

64 if err != cydriver.CUresult.CUDA_SUCCESS: 2a ~IZdaJ0dbJ1dcJ9 dJ! eJ2dfJ# gJ$ hJ% iJ' jJ( kJ) lJ* mJ+ nJ, oJ- pJ. qJ/ rJ: sJ; tJ= uJ? vJ@ wJ[ xJ] yJ^ zJ_ AJ` BJ{ CJ| DJ} EJ~ FJabGJbbHJIJJJKJLJMJNJOJPJQJRgRJSJTJUJVJWJXJYJZJ0J1J2J3J4J5J6J7J8J9J!J#J$J%J'J(J)J*JKe+J,J-J.JLe/JMe:JNe;JOe=Jyb?JPe@JQe[JRe]J^J_J`J{Jzb|JAb}JBb~JCbaKDbbKGccKHcdKIceKJcfKKcgKLchKMciKNcjKkKlKmKnKoK_fpK`fqK{frKsKtKuKvKwKxK,eyK-ezKAKBKCK.eDK=C?CEKFKGKHKIKidJK@C[CdhehKKLKMKNKOKPKQK/eRKSK]CTK^CUKVKWKXKYKZK0Kogjd1K2K3K_C4K`C5Kfh6Kgh7K8K9K!K#K$K%K'K(K)K*K+K,K-K.K/K:K;K=K?K@K[K]K^K:e_K{C|C`K{K|K}K~KkdaL}C~ChhihbLcLdLeLfLgLhL;eiLjLkLaDlLbDmLnLoLpLqLrLsLtLuLvLwLldxLyLzLcDALdDBLjhCLkhDLELFLGLHLILJLKLLL=eMLNLOLeDPLfDQLRLSLTLULVLWLXLYLZL0Lmd1L2L3LgD4LhD5Llh6Lmh7L8L9L!L#L$L%L'L(L?e)L*L+LiD,LjD-L.L/L:L;L=L?L@L[L]L^Lnd_L`L{LkD|LlD}Lnh~LohaMbMcMdMeMfMgMhMiMjMkMlMmMnMoMpMqMrMsMtMuMvMwMxMyMzMAMBMCMDMEMFMGMHMIMJMKMLMMMNMOMPMSgQMRMSMTMUMVMWMXMYMZM0M1M2M3M4MTg5M6M7M8M9MgA!M2b#M$M%MUg'MSe(M)M*M+M,M-M.M/M:M;M=M?M@M[M]M^M_M`M{M|M}M~MaN@ebN[ecN]edN^eeNTefNUegN_ehNmDiNnDjNoDkN%clN'cmNpDnNqDoNhApN`eqNEbrNFbsNGbtN{euNVevNwNxNHbyNIbzNJbANKbBNLbCNDNENFNGNHNINJNKNd LNVgMNWeNNWgON(cPN)cQN*cRN3dSN4dTN5dUNe VNh WNf XNg YN4 ZNcb0N5 1N6 2Ndb3Neb4Nfb5N6d6N7d7N8d8N9d9N!d!N#d#Ngb$Nhb%NMb'Nib(Njb)Nkb*Nlb+Nmb,Nnb-N$IrD%IsDiAtDjAuD'IvD(IwDkAxDlAyD7 zD8 ADXg.NYg/N)g:N*g;NZg=N0g?N+g@N,g[Nph]Nqh^NmA_N)IBD*ICDDDEDFD=fGD`NHDIDJD+IKD{NLDMDND,IOD|NPDQDRD-ISD}NTDUDVD.IWD~NXDYDZD/I0DaO1D2D:I3D;I4D=I5D?I6D@I7D[I8D$d9D%d!D'd#D(d$DOcbOPccO)d%D)bXe]I'D^I(D_I)D`I*DPbdOQbeORbfOSbgOTbhOUbiO|f+DjO,DkO-DlO.DmO/DnO:DoOpOqO}fYeZerO;DsOtOuOvOwOxOyOzOAO~fBOCODOEOFOGOHOIOJOKOLOMONOOOPOQOROSOTOUOVOWOXOYOZO0O1O2O3O4O5O6O7O8O9O!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`P1g{P|P}P2gagbg3gnADcoA~P=DaQ?DbQ@DcQ[DdQ]DeQ^DfQ_DgQEchQ`DiQ{DjQ|DkQ}DlQpAmQqAnQrAoQfepQ~DaEbEcEdEqQrQsQtQuQvQeEfEgEhEiEwQjExQkElE4gyQzQAQob]b^bodeABQCQDQFcEQpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdgehe_b`b{b|b}b~bacbcccdcecfcgchcicjckclcmcncocpcqcrcsctcucvcwcxcycsAzcFQGQHQcgAcBcdgIQegfgggmEnEJQhgigjgkg5g6gQcKQ:bLQpbMQ*bNQqbOQ;bPQ+bQQrbRQ3bSQ+cTQ,cUQ-cVQi WQ4bXQj YQk ZQl 0Qm 1Qn 2Qo 3Qp 4Qq 5Qr 6Q5b7Qs 8Qt 9Qu !Qv #QVb$Qw %Qx 'QWb(Q6b)QRc*Q=b+Qsb,Q,b-Qtb.Q?b/Q-b:Qub;Q7b=Q.c?Q/c@Q:c[Qy ]Q8b^Qz _QA `QB {QC |QD }QE ~QF aRG bRH cR9bdRI eRJ fRK gRL hRXbiRM jRN kRYblR!bmRScnR@boRvbpR.bqRwbrR[bsR/btRxbuR#bvR;cwR=cxR?cyRO zR$bARP BRQ CRR DRS ERT FRU GRV HRW IRX JR%bKRY LRZ MR0 NR1 ORZbPR2 QR3 RR0bSR'bTRlgoEpE@cURVR0etAuAWRqEXRrEYRvAZRwA0RxA1RyA2RTc3R1e4R[c5R2e6RsE7RtE8RuE9RvE!RwE#RxE$RyE%RzE'RAE(RBE)RCE*RDE+REE,RFE-RGE.RHE/RzA:RIE;RAA=RJE?RBA@RKE[RCA]RLE^RDA_RME`REA{RNE|RFA}ROE~RGAaSPEbSQEcSREdSSEeSTEfSUEgSVEhSWEiSXEjSYEkSZElS0EmS1EnS2EoS3EpS4EqS5ErS7CsSrhtS8CuS9CvSshwS!CxS6EySmgng3e7E8E9EzS]cAS!E#E$EBS%ECS'EDS(EESUcVcObWcFSGSCcHS*dISHAIAJSJAKSKALSLANbMSMA1bNANSOAOSPAPSQAQSRARSSASSTATSUAUSVAVSWAWSXAYAZAXcYc0A1AZc0cXS2AYS3AZS4A0S5A1S6A2S7A3S8A4S9A5S!A6S#A7S$A8S%A9S'A!S(A#S)A$S*A%S+A'S,A(S-A)S.A*S/A+S:A,S;A-S=A?A@A.S[A/S]A:S^A;S_A=S`A?S{A@S|A[S}A]S~A^SaB_SbB`ScB{SdB|SeB}SfB~SgBaThBbTiBcTjBdTkBeTlBmBfTnBgToBhTpBiTqBjTrBkTsBlTtBmTuBnTvBoTwBpTxBqTyBrTzBsTABtTBBuTCBDB?fvT4ewT5exT^cyT+dzTAT)E|eBTCTDTET*E+E,EFTGTHTITJTEB}eKT-EFBLT.EMTNT/EGB:EOT;E~ePTQTRTST=E?E@ETTUTVTWTXTHBafYT[EIBZT]E0T1Tbf^E_E2T3T4T5Tcf6T`E7T8T{Eie|E}E#C$C~EaFbFcFdFthje6eeFuhvhfFgFhFkeledfefffgfhf9Tif!TiFjF#T$T%T'Tjf(TkF)T*T+T,TlFkf-T.TmFnF/T:T;T=T?TJBlfpg@ToFKB[T]T^T_TpFmf`T{T|TqFrFsF}T~TaUbUcULBnfqgdUtFMBeUuFfUgUhUofvFwFiUjUkUlUpfmUxFnUoUyF%CmezFAFBFCFDFEFFFGFHFIFNBne7eJFwhOBKFLFMFNFOFPFQFRFSFTFUFVFWFXFqfYFZFpUqUrUsU0FtUuUvUrfwUxU1F2FyUzUAUBUsfCU3FDUEUFUGUHUIUJUKULUMUNUOUPBoeQBRBSBTBUBVBWBXBYBpeZB7g0B1B2BtfqereufvfwfxfyfRdzfAfrg9csgtg8e9eugvgwgxgygSd!czg!eTdAgBgCgBf4F5FPUQURUSUTUUdxhyh6F7F8F9F!F#F3Bse4B5B6B7B8B9B!B#B$B%Bte'B8g(B)B*BUUCfVU$F%FWUXUYUZUDf0U'F1U2U(F'Cue)F*F+F,F-F.F/F:F;F=F?F+Bve#e@Fzh,B[F]F^F3U4U_FEf5U6U7U8U`F{F|F9U!U#U$U%U-BFf'U}F.B(U~F)U*UGfaGbG+U,U-U.UHf/UcG:U;UdGweeGfG(C)CgGhGiGjGkGAhxe$elGBhChmGnGoGyezeIfJfKfLfMf=UNf?UpGqG@U[U]U^UOf_UrG`U{U|U}UsGPf~UaVtGuGbVcVdVeVfV/BQfDggVvG:BhViVjVkVwGRflVmVnVxGyGzGoVpVqVrVsV;BSfEgtVAG=BuVBGvVwVxVTfCGDGyVzVAVBVUfCVEGDVEVFG*CAeGGHGIGJGKGLGMGNGOGPG?BBe%eQGDh@BRGSGTGUGVGWGXGYGZG0G1G2G3G4GVf5G6GFVGVHVIV7GJVKVLVWfMVNV8G9GOVPVQVRVXfSV!GTVUVVVWVXVYVZV0V1V2V3V4V[BCe]B^B_B`B{B|B}B~BaCDebC9gcCdCeCYfEeFeZf0f1f2f3fVd4f5fFg#cGgHg'e(eIgJgKgLgMgWd$cNg)eXdOgPgQg6f#G$G5V6V7V8V9VYdEhFh%G'G(G)G*G+GfCGegChCiCjCkClCmCnCoCpCHeqC!grCsCtC!V7f#V,G-G$V%V'V(V8f)V.G*V+V/G+CIe:G;G=G?G@G[G]G^G_G`G{GuCJe*e|GGhvC}G~GaH,V-V.VbHcHHhIhwCxC/VdH9f:V;V=V?VeHfHgH@V[V]V^V_VyC!f`VhHzCAC@fBCCC[f{ViH|V}V~VjH#faWbWcWdWkHlHmHeWfWgWhWiWDC$fjWnHECkWoHlWmWnWpH%foWpWqWrWqHrHsHsWtWuWvWwWFC'fxWtHGCyWuHzWAWBWvH(fCWDWEWFWwHxHyHGWHWIWJWKWHC)fLWzHICMWAHNWOWPWBH*fQWRWSWTWCHDHEHUWVWWWXWYWJC+fZWFHKC0WGH1W2W3WHH,f4W5W6W7WIHJHKH8W9W!W#W$WLC-f%WLHMC'WMH(W)W*WNH.f+W,W-W.WOHPHQH/W:W;W=W?WNC/f@WRHOCPC[WSH]W^W_WTH:f`W{W|W}WUHVHWH~WaXbXcXdXQC;feXXHRCSCfXYHgXhX1ciX2cjX,C3ckX4clXmXnXoXpXqXrXsXtXuXvXwXxXyXzXAXBXCXDXEXFXGXHXIXJXKXLXMXNXOXPXQXRXSXTXUXVXWXc XX5cYX(bTC,dZX0X1X2X3X4X5X6X7X8X9X!X#X$XJhKhZHUCVCWC%XXC0H1HYC-C'X(X)X*X+X,X-X.X/X6c:X;X=X2H?X3H@X4H[X5H]X6H^X7H_X8H`X9H{X!H|X#H}X$H~X%HaY'HbY(HcY)HdY*HeYZCfY+HgY,HhY-HiY.HjY/HkY:HlYmYnYoYpYqYrYsYtYuYvYwYxYyYzYAYBYCYDYEYFYGYHYIYJYKYLYMYNYOYPYQYRYSYTYUYVYWYXYYYZY0Y1Y2Y3Y4Y5Y6Y7Y8Y9Y!Y#Y$Y%Y'Y(Y)Y*Y+Y,Y-Y.Y/Y:Y;Y=Y?Y@Y[Y]Y^Y;H_Y=H`Y?H@H[H{Y|Y]H^H}Y~YaZ_HbZ`HcZ{H|H}HdZeZ~HfZgZhZiZaIjZbIcIdIeIfIgIkZlZhImZnZoZiIpZqZrZsZtZuZvZwZxZyZzZAZBZCZDZEZFZGZHZIZJZKZLZMZNZOZPZQZRZSZTZUZVZWZXZYZZZ0Z1Z

65 return _check_driver_error(err) 2RgSgTgUgVgWeWg1b,d

66 return 0 2a ~IZdaJ0dbJ1dcJ9 dJ! eJ2dfJ# gJ$ hJ% iJ' jJ( kJ) lJ* mJ+ nJ, oJ- pJ. qJ/ rJ: sJ; tJ= uJ? vJ@ wJ[ xJ] yJ^ zJ_ AJ` BJ{ CJ| DJ} EJ~ FJabGJbbHJIJJJKJLJMJNJOJPJQJRgRJSJTJUJVJWJXJYJZJ0J1J2J3J4J5J6J7J8J9J!J#J$J%J'J(J)J*JKe+J,J-J.JLe/JMe:JNe;JOe=Jyb?JPe@JQe[JRe]J^J_J`J{Jzb|JAb}JBb~JCbaKDbbKGccKHcdKIceKJcfKKcgKLchKMciKNcjKkKlKmKnKoK_fpK`fqK{frKsKtKuKvKwKxK,eyK-ezKAKBKCK.eDK=C?CEKFKGKHKIKidJK@C[CdhehKKLKMKNKOKPKQK/eRKSK]CTK^CUKVKWKXKYKZK0Kogjd1K2K3K_C4K`C5Kfh6Kgh7K8K9K!K#K$K%K'K(K)K*K+K,K-K.K/K:K;K=K?K@K[K]K^K:e_K{C|C`K{K|K}K~KkdaL}C~ChhihbLcLdLeLfLgLhL;eiLjLkLaDlLbDmLnLoLpLqLrLsLtLuLvLwLldxLyLzLcDALdDBLjhCLkhDLELFLGLHLILJLKLLL=eMLNLOLeDPLfDQLRLSLTLULVLWLXLYLZL0Lmd1L2L3LgD4LhD5Llh6Lmh7L8L9L!L#L$L%L'L(L?e)L*L+LiD,LjD-L.L/L:L;L=L?L@L[L]L^Lnd_L`L{LkD|LlD}Lnh~LohaMbMcMdMeMfMgMhMiMjMkMlMmMnMoMpMqMrMsMtMuMvMwMxMyMzMAMBMCMDMEMFMGMHMIMJMKMLMMMNMOMPMSgQMRMSMTMUMVMWMXMYMZM0M1M2M3M4MTg5M6M7M8M9MgA!M2b#M$M%MUg'MSe(M)M*M+M,M-M.M/M:M;M=M?M@M[M]M^M_M`M{M|M}M~MaN@ebN[ecN]edN^eeNTefNUegN_ehNmDiNnDjNoDkN%clN'cmNpDnNqDoNhApN`eqNEbrNFbsNGbtN{euNVevNwNxNHbyNIbzNJbANKbBNLbCNDNENFNGNHNINJNKNd LNVgMNWeNNWgON(cPN)cQN*cRN3dSN4dTN5dUNe VNh WNf XNg YN4 ZNcb0N5 1N6 2Ndb3Neb4Nfb5N6d6N7d7N8d8N9d9N!d!N#d#Ngb$Nhb%NMb'Nib(Njb)Nkb*Nlb+Nmb,Nnb-N$IrD%IsDiAtDjAuD'IvD(IwDkAxDlAyD7 zD8 ADXg.NYg/N)g:N*g;NZg=N0g?N+g@N,g[Nph]Nqh^NmA_N)IBD*ICDDDEDFD=fGD`NHDIDJD+IKD{NLDMDND,IOD|NPDQDRD-ISD}NTDUDVD.IWD~NXDYDZD/I0DaO1D2D:I3D;I4D=I5D?I6D@I7D[I8D$d9D%d!D'd#D(d$DOcbOPccO)d%D)bXe]I'D^I(D_I)D`I*DPbdOQbeORbfOSbgOTbhOUbiO|f+DjO,DkO-DlO.DmO/DnO:DoOpOqO}fYeZerO;DsOtOuOvOwOxOyOzOAO~fBOCODOEOFOGOHOIOJOKOLOMONOOOPOQOROSOTOUOVOWOXOYOZO0O1O2O3O4O5O6O7O8O9O!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`P1g{P|P}P2gagbg3gnADcoA~P=DaQ?DbQ@DcQ[DdQ]DeQ^DfQ_DgQEchQ`DiQ{DjQ|DkQ}DlQpAmQqAnQrAoQfepQ~DaEbEcEdEqQrQsQtQuQvQeEfEgEhEiEwQjExQkElE4gyQzQAQob]b^bodeABQCQDQFcEQpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdgehe_b`b{b|b}b~bacbcccdcecfcgchcicjckclcmcncocpcqcrcsctcucvcwcxcycsAzcFQGQHQcgAcBcdgIQegfgggmEnEJQhgigjgkg5g6gQcKQ:bLQpbMQ*bNQqbOQ;bPQ+bQQrbRQ3bSQ+cTQ,cUQ-cVQi WQ4bXQj YQk ZQl 0Qm 1Qn 2Qo 3Qp 4Qq 5Qr 6Q5b7Qs 8Qt 9Qu !Qv #QVb$Qw %Qx 'QWb(Q6b)QRc*Q=b+Qsb,Q,b-Qtb.Q?b/Q-b:Qub;Q7b=Q.c?Q/c@Q:c[Qy ]Q8b^Qz _QA `QB {QC |QD }QE ~QF aRG bRH cR9bdRI eRJ fRK gRL hRXbiRM jRN kRYblR!bmRScnR@boRvbpR.bqRwbrR[bsR/btRxbuR#bvR;cwR=cxR?cyRO zR$bARP BRQ CRR DRS ERT FRU GRV HRW IRX JR%bKRY LRZ MR0 NR1 ORZbPR2 QR3 RR0bSR'bTRlgoEpE@cURVR0etAuAWRqEXRrEYRvAZRwA0RxA1RyA2RTc3R1e4R[c5R2e6RsE7RtE8RuE9RvE!RwE#RxE$RyE%RzE'RAE(RBE)RCE*RDE+REE,RFE-RGE.RHE/RzA:RIE;RAA=RJE?RBA@RKE[RCA]RLE^RDA_RME`REA{RNE|RFA}ROE~RGAaSPEbSQEcSREdSSEeSTEfSUEgSVEhSWEiSXEjSYEkSZElS0EmS1EnS2EoS3EpS4EqS5ErS7CsSrhtS8CuS9CvSshwS!CxS6EySmgng3e7E8E9EzS]cAS!E#E$EBS%ECS'EDS(EESUcVcObWcFSGSCcHS*dISHAIAJSJAKSKALSLANbMSMA1bNANSOAOSPAPSQAQSRARSSASSTATSUAUSVAVSWAWSXAYAZAXcYc0A1AZc0cXS2AYS3AZS4A0S5A1S6A2S7A3S8A4S9A5S!A6S#A7S$A8S%A9S'A!S(A#S)A$S*A%S+A'S,A(S-A)S.A*S/A+S:A,S;A-S=A?A@A.S[A/S]A:S^A;S_A=S`A?S{A@S|A[S}A]S~A^SaB_SbB`ScB{SdB|SeB}SfB~SgBaThBbTiBcTjBdTkBeTlBmBfTnBgToBhTpBiTqBjTrBkTsBlTtBmTuBnTvBoTwBpTxBqTyBrTzBsTABtTBBuTCBDB?fvT4ewT5exT^cyT+dzTAT)E|eBTCTDTET*E+E,EFTGTHTITJTEB}eKT-EFBLT.EMTNT/EGB:EOT;E~ePTQTRTST=E?E@ETTUTVTWTXTHBafYT[EIBZT]E0T1Tbf^E_E2T3T4T5Tcf6T`E7T8T{Eie|E}E#C$C~EaFbFcFdFthje6eeFuhvhfFgFhFkeledfefffgfhf9Tif!TiFjF#T$T%T'Tjf(TkF)T*T+T,TlFkf-T.TmFnF/T:T;T=T?TJBlfpg@ToFKB[T]T^T_TpFmf`T{T|TqFrFsF}T~TaUbUcULBnfqgdUtFMBeUuFfUgUhUofvFwFiUjUkUlUpfmUxFnUoUyF%CmezFAFBFCFDFEFFFGFHFIFNBne7eJFwhOBKFLFMFNFOFPFQFRFSFTFUFVFWFXFqfYFZFpUqUrUsU0FtUuUvUrfwUxU1F2FyUzUAUBUsfCU3FDUEUFUGUHUIUJUKULUMUNUOUPBoeQBRBSBTBUBVBWBXBYBpeZB7g0B1B2BtfqereufvfwfxfyfRdzfAfrg9csgtg8e9eugvgwgxgygSd!czg!eTdAgBgCgBf4F5FPUQURUSUTUUdxhyh6F7F8F9F!F#F3Bse4B5B6B7B8B9B!B#B$B%Bte'B8g(B)B*BUUCfVU$F%FWUXUYUZUDf0U'F1U2U(F'Cue)F*F+F,F-F.F/F:F;F=F?F+Bve#e@Fzh,B[F]F^F3U4U_FEf5U6U7U8U`F{F|F9U!U#U$U%U-BFf'U}F.B(U~F)U*UGfaGbG+U,U-U.UHf/UcG:U;UdGweeGfG(C)CgGhGiGjGkGAhxe$elGBhChmGnGoGyezeIfJfKfLfMf=UNf?UpGqG@U[U]U^UOf_UrG`U{U|U}UsGPf~UaVtGuGbVcVdVeVfV/BQfDggVvG:BhViVjVkVwGRflVmVnVxGyGzGoVpVqVrVsV;BSfEgtVAG=BuVBGvVwVxVTfCGDGyVzVAVBVUfCVEGDVEVFG*CAeGGHGIGJGKGLGMGNGOGPG?BBe%eQGDh@BRGSGTGUGVGWGXGYGZG0G1G2G3G4GVf5G6GFVGVHVIV7GJVKVLVWfMVNV8G9GOVPVQVRVXfSV!GTVUVVVWVXVYVZV0V1V2V3V4V[BCe]B^B_B`B{B|B}B~BaCDebC9gcCdCeCYfEeFeZf0f1f2f3fVd4f5fFg#cGgHg'e(eIgJgKgLgMgWd$cNg)eXdOgPgQg6f#G$G5V6V7V8V9VYdEhFh%G'G(G)G*G+GfCGegChCiCjCkClCmCnCoCpCHeqC!grCsCtC!V7f#V,G-G$V%V'V(V8f)V.G*V+V/G+CIe:G;G=G?G@G[G]G^G_G`G{GuCJe*e|GGhvC}G~GaH,V-V.VbHcHHhIhwCxC/VdH9f:V;V=V?VeHfHgH@V[V]V^V_VyC!f`VhHzCAC@fBCCC[f{ViH|V}V~VjH#faWbWcWdWkHlHmHeWfWgWhWiWDC$fjWnHECkWoHlWmWnWpH%foWpWqWrWqHrHsHsWtWuWvWwWFC'fxWtHGCyWuHzWAWBWvH(fCWDWEWFWwHxHyHGWHWIWJWKWHC)fLWzHICMWAHNWOWPWBH*fQWRWSWTWCHDHEHUWVWWWXWYWJC+fZWFHKC0WGH1W2W3WHH,f4W5W6W7WIHJHKH8W9W!W#W$WLC-f%WLHMC'WMH(W)W*WNH.f+W,W-W.WOHPHQH/W:W;W=W?WNC/f@WRHOCPC[WSH]W^W_WTH:f`W{W|W}WUHVHWH~WaXbXcXdXQC;feXXHRCSCfXYHgXhX1ciX2cjX,C3ckX4clXmXnXoXpXqXrXsXtXuXvXwXxXyXzXAXBXCXDXEXFXGXHXIXJXKXLXMXNXOXPXQXRXSXTXUXVXWXc XX5cYX(bTC,dZX0X1X2X3X4X5X6X7X8X9X!X#X$XJhKhZHUCVCWC%XXC0H1HYC-C'X(X)X*X+X,X-X.X/X6c:X;X=X2H?X3H@X4H[X5H]X6H^X7H_X8H`X9H{X!H|X#H}X$H~X%HaY'HbY(HcY)HdY*HeYZCfY+HgY,HhY-HiY.HjY/HkY:HlYmYnYoYpYqYrYsYtYuYvYwYxYyYzYAYBYCYDYEYFYGYHYIYJYKYLYMYNYOYPYQYRYSYTYUYVYWYXYYYZY0Y1Y2Y3Y4Y5Y6Y7Y8Y9Y!Y#Y$Y%Y'Y(Y)Y*Y+Y,Y-Y.Y/Y:Y;Y=Y?Y@Y[Y]Y^Y;H_Y=H`Y?H@H[H{Y|Y]H^H}Y~YaZ_HbZ`HcZ{H|H}HdZeZ~HfZgZhZiZaIjZbIcIdIeIfIgIkZlZhImZnZoZiIpZqZrZsZtZuZvZwZxZyZzZAZBZCZDZEZFZGZHZIZJZKZLZMZNZOZPZQZRZSZTZUZVZWZXZYZZZ0Z1Z

67  

68  

69cdef int HANDLE_RETURN_NVRTC(cynvrtc.nvrtcProgram prog, cynvrtc.nvrtcResult err) except?-1 nogil: 

70 """Handle NVRTC result codes, raising NVRTCError with program log on failure.""" 

71 if err == cynvrtc.nvrtcResult.NVRTC_SUCCESS: 29 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbybzbAbBbCbDb,e-e.eid/ejd:ekd;eld=emd?end2b@e[e]e^e_ehA`eEbFbGb{eHbIbJbKbLbd e h f g 4 cb5 6 dbebfbgbhbMbibjbkblbmbnb7 8 DcEcfeobodFcpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdgeheb _b`b{b|b}b~bacbcccdcecfcgchcicjckclcmcncocpcqcrcsctcucvcwcxcycsAzcCc*dHAIAJAKALANbMA1bNAOAPAQARASATAUAVAWAXAYAZAXcYc0A1AZc0c2A3A4A5A6A7A8A9A!A#A$A%A'A(A)A*A+A,A-A.A/A:A;A=A?A@A[A]A^A_A`A{A|A}A~AaBbBcBdBeBfBgBhBiBjBkBlBmBnBoBpBqBrBsBtBuBvBwBxByBzBABBBCBDB|eEB}eFBGB~eHBafIBbfcfiethjevhkeledfefffgfhfifjfkfJBlfKBmfLBnfMBofpfmeNBneOBqfrfsfPBoeQBRBSBTBUBVBWBXBYBpeZB7g0B1B2BtfqereufvfwfxfyfRdzfAf9cSd!cTdBfUd3Bse4B5B6B7B8B9B!B#B$B%Bte'B8g(B)B*BCfDfue+Bve,BEf-BFf.BGfHfweAhxeChyezeIfJfKfLfMfNfOfPf/BQf:BRf;BSf=BTfUfAe?BBe@BVfWfXf[BCe]B^B_B`B{B|B}B~BaCDebC9gcCdCeCYfEeFeZf0f1f2f3fVd4f5f#cWd$cXd6fYdfCGegChCiCjCkClCmCnCoCpCHeqC!grCsCtC7f8fIeuCJevCwCxC9fyC!fzCACBC#fDC$fEC%fFC'fGC(fHC)fIC*fJC+fKC,fLC-fMC.fNC/fOCPC:fQC;fRCSC1c2c3c4c-d.d/d_c`c{c:d;d=d?d7c8c|c@d}c[d]d^d_d~c`dadbdcddd{d|d}d~daebecedeeeedfdgdc 5cUCVCWChdXCYC6c

72 return 0 29 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbybzbAbBbCbDb,e-e.eid/ejd:ekd;eld=emd?end2b@e[e]e^e_ehA`eEbFbGb{eHbIbJbKbLbd e h f g 4 cb5 6 dbebfbgbhbMbibjbkblbmbnb7 8 DcEcfeobodFcpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdgeheb _b`b{b|b}b~bacbcccdcecfcgchcicjckclcmcncocpcqcrcsctcucvcwcxcycsAzcCc*dHAIAJAKALANbMA1bNAOAPAQARASATAUAVAWAXAYAZAXcYc0A1AZc0c2A3A4A5A6A7A8A9A!A#A$A%A'A(A)A*A+A,A-A.A/A:A;A=A?A@A[A]A^A_A`A{A|A}A~AaBbBcBdBeBfBgBhBiBjBkBlBmBnBoBpBqBrBsBtBuBvBwBxByBzBABBBCBDB|eEB}eFBGB~eHBafIBbfcfiethjevhkeledfefffgfhfifjfkfJBlfKBmfLBnfMBofpfmeNBneOBqfrfsfPBoeQBRBSBTBUBVBWBXBYBpeZB7g0B1B2BtfqereufvfwfxfyfRdzfAf9cSd!cTdBfUd3Bse4B5B6B7B8B9B!B#B$B%Bte'B8g(B)B*BCfDfue+Bve,BEf-BFf.BGfHfweAhxeChyezeIfJfKfLfMfNfOfPf/BQf:BRf;BSf=BTfUfAe?BBe@BVfWfXf[BCe]B^B_B`B{B|B}B~BaCDebC9gcCdCeCYfEeFeZf0f1f2f3fVd4f5f#cWd$cXd6fYdfCGegChCiCjCkClCmCnCoCpCHeqC!grCsCtC7f8fIeuCJevCwCxC9fyC!fzCACBC#fDC$fEC%fFC'fGC(fHC)fIC*fJC+fKC,fLC-fMC.fNC/fOCPC:fQC;fRCSC1c2c3c4c-d.d/d_c`c{c:d;d=d?d7c8c|c@d}c[d]d^d_d~c`dadbdcddd{d|d}d~daebecedeeeedfdgdc 5cUCVCWChdXCYC6c

73 with gil: 1c

74 _raise_nvrtc_error(prog, err) 1c

75  

76  

77cdef int _raise_nvrtc_error(cynvrtc.nvrtcProgram prog, cynvrtc.nvrtcResult err) except -1: 

78 """Build error message with program log and raise NVRTCError.""" 

79 cdef const char* err_str = cynvrtc.nvrtcGetErrorString(err) 1c

80 cdef size_t logsize = 0 1c

81 if prog != NULL: 1c

82 cynvrtc.nvrtcGetProgramLogSize(prog, &logsize) 1c

83 cdef bytes log_bytes 

84 cdef str log_str = "" 1c

85 if logsize > 1 and prog != NULL: 1c

86 log_bytes = b" " * logsize 1c

87 if cynvrtc.nvrtcGetProgramLog(prog, <char*>log_bytes) == cynvrtc.nvrtcResult.NVRTC_SUCCESS: 1c

88 log_str = log_bytes.decode("utf-8", errors="backslashreplace") 1c

89 err_msg = f"{err}: {err_str.decode()}" if err_str != NULL else f"NVRTC error {err}" 1c

90 if log_str: 1c

91 err_msg += f", compilation log:\n\n{log_str}" 1c

92 raise NVRTCError(err_msg) 1c

93  

94  

95cdef int HANDLE_RETURN_NVVM(cynvvm.nvvmProgram prog, cynvvm.nvvmResult err) except?-1 nogil: 

96 """Handle NVVM result codes, raising nvvmError with program log on failure.""" 

97 if err == cynvvm.nvvmResult.NVVM_SUCCESS: 2CCjIkI(bTC,dlImInIoIpIqIrIsItIuIvIwIxIyIJhKh

98 return 0 2CCjIkI(bTC,dlImInIoIpIqIrIsItIuIvIwIxIyIJhKh

99 with gil: 2(b

100 _raise_nvvm_error(prog, err) 2(b

101  

102  

103cdef int _raise_nvvm_error(cynvvm.nvvmProgram prog, cynvvm.nvvmResult err) except -1: 

104 """Raise nvvmError annotated with the program log.""" 

105 cdef size_t logsize = 0 2(b

106 if prog != NULL: 2(b

107 cynvvm.nvvmGetProgramLogSize(prog, &logsize) 2(b

108 cdef bytes log_bytes 

109 cdef str log_str = "" 2(b

110 if logsize > 1 and prog != NULL: 2(b

111 log_bytes = b" " * logsize 2(b

112 if cynvvm.nvvmGetProgramLog(prog, <char*>log_bytes) == cynvvm.nvvmResult.NVVM_SUCCESS: 2(b

113 log_str = log_bytes.decode("utf-8", errors="backslashreplace") 2(b

114 cdef object exc = nvvmError(err) 2(b

115 if log_str: 2a (b

116 exc.args = (exc.args[0] + f"\nNVVM program log: {log_str}", *exc.args[1:]) 2(b

117 raise exc 2(b

118  

119  

120cdef int HANDLE_RETURN_NVJITLINK( 

121 cynvjitlink.nvJitLinkHandle handle, cynvjitlink.nvJitLinkResult err) except?-1 nogil: 

122 """Handle nvJitLink result codes, raising nvJitLinkError with error log on failure.""" 

123 if err == cynvjitlink.nvJitLinkResult.NVJITLINK_SUCCESS: 2b zIAIBICIDIEIFIGIHIIIJI.C/CKILIMINIOIPIQIRISITIUIVIWI+eXIYIZI0I1I0C2I3I4I5I6I7I8I9I!I

124 return 0 2b zIAIBICIDIEIFIGIHIIIJI.C/CKILIMINIOIPIQIRISITIUIVIWIXIYIZI0I1I0C2I3I4I5I6I7I8I9I!I

125 with gil: 2b +e

126 _raise_nvjitlink_error(handle, err) 2b +e

127  

128  

129cdef int _raise_nvjitlink_error( 

130 cynvjitlink.nvJitLinkHandle handle, cynvjitlink.nvJitLinkResult err) except -1: 

131 """Raise nvJitLinkError annotated with the error log.""" 

132 cdef size_t logsize = 0 2b +e

133 if handle != NULL: 2a b +e

134 cynvjitlink.nvJitLinkGetErrorLogSize(handle, &logsize) 1b

135 cdef bytes log_bytes 

136 cdef str log_str = "" 2b +e

137 if logsize > 1 and handle != NULL: 2b +e

138 log_bytes = b" " * logsize 1b

139 if cynvjitlink.nvJitLinkGetErrorLog(handle, <char*>log_bytes) == \ 1b

140 cynvjitlink.nvJitLinkResult.NVJITLINK_SUCCESS: 1b

141 log_str = log_bytes.decode("utf-8", errors="backslashreplace") 1b

142 cdef object exc = nvJitLinkError(err) 2b +e

143 if log_str: 2b +e

144 exc.args = (exc.args[0] + f"\nnvJitLink error log: {log_str}", *exc.args[1:]) 1b

145 raise exc 2b +e

146  

147  

148cdef object _RUNTIME_SUCCESS = runtime.cudaError_t.cudaSuccess 

149cdef object _NVRTC_SUCCESS = nvrtc.nvrtcResult.NVRTC_SUCCESS 

150  

151  

152cpdef inline int _check_driver_error(cydriver.CUresult error) except?-1 nogil: 

153 if error == cydriver.CUresult.CUDA_SUCCESS: 2a ZdLh0dMh1dNh9 Oh! Ph2dQh# Rh$ Sh% Th' Uh( Vh) Wh* Xh+ Yh, Zh- 0h. 1h/ 2h: 3h; 4h= 5h? 6h@ 7h[ 8h] 9h^ !h_ #h` $h{ %h| 'h} (h~ )hab*hbb+h,h-h.h/hRg:h;h=h?h@h[h]h^h_h`h{h|h}h~haiKebicidiLeeiMefiNegiOehiybiiPejiQekiReliminizboiAbpiBbqiCbriDbsiGctiHcuiIcviJcwiKcxiLcyiMcziNcAiBi-gCi.gDi/gEi_fFi`fGi{fHiIiJiKiLiMiNiOiPiQiRiSiTiUiViWiXiYiZi0i1i2i3i4i5i6i7i8i9i!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~jakbkckdkekfkgkSghkikjkkklkmknkokTgpkqkrksk2btkukUgvkSewkxkykzkAkBkCkDkEkFkGkHkIkJkKkLkTeMkUeNkOkPkQkRk%cSk'cTkUkVkWkXkEbYkFbZkGb0k1kVe2k3kHb4kIb5kJb6kKb7kLb8k9k!k#k$kd %kVg'kWe(kWg)k(c)c*c3d4d5de h f g 4 cb5 6 dbebfb6d7d8d9d!d#dgb*khb+kMb,kib-kjb.kkb/klb:kmb;knb=k7 8 =f$d%d'd(dOcPc)d)bXePbQbRbSbTbUb?k@k[k]k^k_k`k1C:g{k|k}k~kalblcldlelflglhliljl1gklllml2gnl3gDcolplqlrlsltlulvlEcwlxlylzlAlBlClDlEl4gob]b^bFlGlHlIlpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdgeheJlKlLlMlNlOlPlQlRlSlTlUlVlWlXlYlZl0l1l2l3l4l5l6l7l8l9l!l#l$l%l'l(l)l*l+lcgAcBcdgegfggghgigjgkg5g6gQc:bpb*bqb;b+brb3b+c,c-ci 4bj k l m n o p q r 5bs t u v Vbw x Wb6bRc=bsb,btb?b-bub7b.c/c:cy 8bz A B C D E F G H 9bI J K L XbM N Yb!bSc@bvb.bwb[b/bxb#b;c=c?cO $bP Q R S T U V W X %bY Z 0 1 Zb2 3 0b'blg,l-l0e.l/l:l;l=l?l@l[l]l^l_l`l{l|l}l~lambmcmdmemfmgmhmimjmkmlmmmnmompmqmrmsmtmumvmwmxmymzmAmBmCmDmEmFmGmHmImJmKmLmMmNmOmPmQmRmSmTmUmVmWmmgng3eXmYmZm0m1m2mUcVcObWc3m*d4m5m6m7m8m9mNb!m#m1b$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~uavbvcvdvevfvgvhvivjvkvlvmvnvovpvqvrvsvtvuvvvwvxvyvzvAvBvCvDvEvFvGvHvIvJvKvLvMvNvOvPvQvRvSvTvUvVvWvXvYvZv0v1v2v3v4v5v6v7v8v9v!v#v$v%v'v(v)v*v+v,v-v.v/v:v;v=v?v@v[v]v^v_v`v{v|v}v~vawbwcwdwewfwgwhwiwjwkwlwmwnwowpwqwrwswtwuwvwwwxwywzwAwBwCwDwEwFwGwHwIwJwKwLwMwNwOwPwQwRwSwTwUwVwWwXwYwZw0w1w2w3w4w5w6w7w8w9w!w#w$w%w'w(w)w*w+w,w-w.w/w:w;w=w?w@w[w]w^w_w`w{w|w}w~waxbxcxdxexfxgxhxixjxkxlxmxnxoxpxqxrxsxtxuxvxwxxxyxzxAxBxCxDxExFxGxHxIxJxKxLxMxNxOxPxQxRxSxTxUxVxWxXxYxZx0x1x2x3x4x5x6x7x8x9x!x#x$x%x'x(x)x*x+x,x-x.x/x:x;x=x?x@x[x]x^x_x`x{x|x}x~xay,dbycydyeyfygyhyiyjykylymynyoypyqyrysytyuyvywyxyyyzyAyByCyDyEyFyGyHyIyJyKyLyMyNyOyPyQyRySyTyUyVyWyXyYyZy0y1y2y3y4y5y6y7y8y9y!y#y$y%y'y(y)y*y+y,y-y.y/y:y;y=y?y@y[y]y^y_y`y{y|y}y~yazbzczdzezfzgzhzizjzkzlzmznzozpzqzrzsztzuzvzwzxzyzzzAzBzCzDzEzFzGzHzIzJzKzLzMzNzOzPzQzRzSzTzUzVzWzXzYzZz0z1z2z3z4z5z6z7z8z9z!z#z$z%z'z(z)z*z+z,z-z.z/z:z;z=z?z@z[z]z^z_z`z{z|z}z~zaAbAcA

154 return 0 2a ZdLh0dMh1dNh9 Oh! Ph2dQh# Rh$ Sh% Th' Uh( Vh) Wh* Xh+ Yh, Zh- 0h. 1h/ 2h: 3h; 4h= 5h? 6h@ 7h[ 8h] 9h^ !h_ #h` $h{ %h| 'h} (h~ )hab*hbb+h,h-h.h/h:h;h=h?h@h[h]h^h_h`h{h|h}h~haiKebicidiLeeiMefiNegiOehiybiiPejiQekiReliminizboiAbpiBbqiCbriDbsiGctiHcuiIcviJcwiKcxiLcyiMcziNcAiBi-gCi.gDi/gEi_fFi`fGi{fHiIiJiKiLiMiNiOiPiQiRiSiTiUiViWiXiYiZi0i1i2i3i4i5i6i7i8i9i!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~jakbkckdkekfkgkhkikjkkklkmknkokpkqkrksk2btkukvkSewkxkykzkAkBkCkDkEkFkGkHkIkJkKkLkTeMkUeNkOkPkQkRk%cSk'cTkUkVkWkXkEbYkFbZkGb0k1kVe2k3kHb4kIb5kJb6kKb7kLb8k9k!k#k$kd %k'kWe(k)k(c)c*c3d4d5de h f g 4 cb5 6 dbebfb6d7d8d9d!d#dgb*khb+kMb,kib-kjb.kkb/klb:kmb;knb=k7 8 =f$d%d'd(dOcPc)d)bXePbQbRbSbTbUb?k@k[k]k^k_k`k:g{k|k}k~kalblcldlelflglhliljl1gklllml2gnl3gDcolplqlrlsltlulvlEcwlxlylzlAlBlClDlEl4gob]b^bFlGlHlIlpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdgeheJlKlLlMlNlOlPlQlRlSlTlUlVlWlXlYlZl0l1l2l3l4l5l6l7l8l9l!l#l$l%l'l(l)l*l+lcgAcBcdgegfggghgigjgkg5g6gQc:bpb*bqb;b+brb3b+c,c-ci 4bj k l m n o p q r 5bs t u v Vbw x Wb6bRc=bsb,btb?b-bub7b.c/c:cy 8bz A B C D E F G H 9bI J K L XbM N Yb!bSc@bvb.bwb[b/bxb#b;c=c?cO $bP Q R S T U V W X %bY Z 0 1 Zb2 3 0b'blg,l-l0e.l/l:l;l=l?l@l[l]l^l_l`l{l|l}l~lambmcmdmemfmgmhmimjmkmlmmmnmompmqmrmsmtmumvmwmxmymzmAmBmCmDmEmFmGmHmImJmKmLmMmNmOmPmQmRmSmTmUmVmWmmgng3eXmYmZm0m1m2mUcVcObWc3m*d4m5m6m7m8m9mNb!m#m1b$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~uavbvcvdvevfvgvhvivjvkvlvmvnvovpvqvrvsvtvuvvvwvxvyvzvAvBvCvDvEvFvGvHvIvJvKvLvMvNvOvPvQvRvSvTvUvVvWvXvYvZv0v1v2v3v4v5v6v7v8v9v!v#v$v%v'v(v)v*v+v,v-v.v/v:v;v=v?v@v[v]v^v_v`v{v|v}v~vawbwcwdwewfwgwhwiwjwkwlwmwnwowpwqwrwswtwuwvwwwxwywzwAwBwCwDwEwFwGwHwIwJwKwLwMwNwOwPwQwRwSwTwUwVwWwXwYwZw0w1w2w3w4w5w6w7w8w9w!w#w$w%w'w(w)w*w+w,w-w.w/w:w;w=w?w@w[w]w^w_w`w{w|w}w~waxbxcxdxexfxgxhxixjxkxlxmxnxoxpxqxrxsxtxuxvxwxxxyxzxAxBxCxDxExFxGxHxIxJxKxLxMxNxOxPxQxRxSxTxUxVxWxXxYxZx0x1x2x3x4x5x6x7x8x9x!x#x$x%x'x(x)x*x+x,x-x.x/x:x;x=x?x@x[x]x^x_x`x{x|x}x~xaybycydyeyfygyhyiyjykylymynyoypyqyrysytyuyvywyxyyyzyAyByCyDyEyFyGyHyIyJyKyLyMyNyOyPyQyRySyTyUyVyWyXyYyZy0y1y2y3y4y5y6y7y8y9y!y#y$y%y'y(y)y*y+y,y-y.y/y:y;y=y?y@y[y]y^y_y`y{y|y}y~yazbzczdzezfzgzhzizjzkzlzmznzozpzqzrzsztzuzvzwzxzyzzzAzBzCzDzEzFzGzHzIzJzKzLzMzNzOzPzQzRzSzTzUzVzWzXzYzZz0z1z2z3z4z5z6z7z8z9z!z#z$z%z'z(z)z*z+z,z-z.z/z:z;z=z?z@z[z]z^z_z`z{z|z}z~zaAbAcA

155 cdef const char* name 

156 name_err = cydriver.cuGetErrorName(error, &name) 2RgSgTgUgVgWeWg1C:gOb1b,d

157 if name_err != cydriver.CUresult.CUDA_SUCCESS: 2RgSgTgUgVgWeWg1C:gOb1b,d

158 raise CUDAError(f"UNEXPECTED ERROR CODE: {error}") 2:g

159 with gil: 2RgSgTgUgVgWeWg1C:gOb1b,d

160 # TODO: consider lower this to Cython 

161 expl = DRIVER_CU_RESULT_EXPLANATIONS.get(int(error)) 2RgSgTgUgVgWeWg1C:gOb1b,d

162 if expl is not None: 2RgSgTgUgVgWeWg1C:gOb1b,d

163 raise CUDAError(f"{name.decode()}: {expl}") 2RgSgTgUgVgWeWg1C:gOb1b,d

164 cdef const char* desc 

165 desc_err = cydriver.cuGetErrorString(error, &desc) 

166 if desc_err != cydriver.CUresult.CUDA_SUCCESS: 

167 raise CUDAError(f"{name.decode()}") 

168 raise CUDAError(f"{name.decode()}: {desc.decode()}") 

169  

170  

171cpdef inline int _check_runtime_error(error) except?-1: 

172 if error == _RUNTIME_SUCCESS: 2PbQbRbSbTbUb;g|f=g2C]f}fYeZe~fagbg]b^bAcBcpbqbrb3bi j k l m n o p q r s t u v w x sbtbub7by z A B C D E F G H I J K L M N vbwbxb#bO P Q R S T U V W X Y Z 0 1 2 3 @cTc1e[c2e]c

173 return 0 2PbQbRbSbTbUb;g|f=g]f}fYeZe~fagbg]b^bAcBcpbqbrb3bi j k l m n o p q r s t u v w x sbtbub7by z A B C D E F G H I J K L M N vbwbxb#bO P Q R S T U V W X Y Z 0 1 2 3 @cTc1e[c2e]c

174 name_err, name = runtime.cudaGetErrorName(error) 22C]f

175 if name_err != _RUNTIME_SUCCESS: 22C]f

176 raise CUDAError(f"UNEXPECTED ERROR CODE: {error}") 

177 name = name.decode() 22C]f

178 expl = RUNTIME_CUDA_ERROR_EXPLANATIONS.get(int(error)) 22C]f

179 if expl is not None: 22C]f

180 raise CUDAError(f"{name}: {expl}") 22C]f

181 desc_err, desc = runtime.cudaGetErrorString(error) 2]f

182 if desc_err != _RUNTIME_SUCCESS: 2]f

183 raise CUDAError(f"{name}") 

184 desc = desc.decode() 2]f

185 raise CUDAError(f"{name}: {desc}") 2]f

186  

187  

188cpdef inline int _check_nvrtc_error(error, handle=None) except?-1: 

189 if error == _NVRTC_SUCCESS: 2a dA#Ib _b`b{b|b}b~bacbcccdcecfcgchcicjckclcmcncocpcqcrcsctcucvcwcxcyczcCcXcYc#gZc$g0c@f[f1c2c3c4c-d.d/d_c`c{c:d;d=d?d7c8c|c@d}c[d]d^d_d~c`dadbdcddd{d|d}d~daebecedeeeedfdgdc 5chd6c

190 return 0 2a #Ib _b`b{b|b}b~bacbcccdcecfcgchcicjckclcmcncocpcqcrcsctcucvcwcxcyczcCcXcYc#gZc$g0c@f[f1c2c3c4c-d.d/d_c`c{c:d;d=d?d7c8c|c@d}c[d]d^d_d~c`dadbdcddd{d|d}d~daebecedeeeedfdgdc 5chd6c

191 err = f"{error}: {nvrtc.nvrtcGetErrorString(error)[1].decode()}" 2dA#I

192 if handle is not None: 2dA#I

193 _, logsize = nvrtc.nvrtcGetProgramLogSize(handle) 2dA

194 log = b" " * logsize 2dA

195 _ = nvrtc.nvrtcGetProgramLog(handle, log) 2dA

196 err += f", compilation log:\n\n{log.decode('utf-8', errors='backslashreplace')}" 2dA

197 raise NVRTCError(err) 2dA#I

198  

199  

200cdef inline int _check_error(error, handle=None) except?-1: 

201 if isinstance(error, driver.CUresult): 2a ZdLh0dMh1dNh9 Oh! Ph2dQh# Rh$ Sh% Th' Uh( Vh) Wh* Xh+ Yh, Zh- 0h. 1h/ 2h: 3h; 4h= 5h? 6h@ 7h[ 8h] 9h^ !h_ #h` $h{ %h| 'h} (h~ )hab*hbb+h,h-h.h/h:h;h=h?h@h[h]h^h_h`h{h|h}h~haiKebicidiLeeiMefiNegiOehiybiiPejiQekiReliminizboiAbpiBbqiCbriDbsiGctiHcuiIcviJcwiKcxiLcyiMcziNcAiBi-gCi.gDi/gEi_fFi`fGi{fHiIiJiKiLiMiNiOiPiQiRiSiTiUiViWiXiYiZi0i1i2i3i4i5i6i7i8i9i!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~jakbkckdkekfkgkhkikjkkklkmknkokpkqkrksk2btkukvkSewkxkykzkAkBkCkDkEkFkGkHkIkJkKkLkTeMkUeNkOkPkQkRk%cSk'cTkUkVkWkXkEbYkFbZkGb0k1kVe2k3kHb4kIb5kJb6kKb7kLb8k9k!k#k$kd %k'k(k)k(c)c*c3d4d5de h f g 4 cb5 6 dbebfb6d7d8d9d!d#dgb*khb+kMb,kib-kjb.kkb/klb:kmb;knb=k7 8 =f$d%d'd(dOcPc)d)bXePbQbRbSbTbUb;g|f?k=g@k[k]k^k_k`k{k}fYe|kZe}k~kalblcldlelfl~fglhliljl1gklllml2gagbgnl3golplqlrlsltlulvlwlxlylzlAlBlClDlEl4gob]b^bFlGlHlIlb _bJl`bKl{bLl|bMl}bNl~bOlacPlbcQlccRldcSlecTlfcUlgcVlhcWlicXljcYlkcZllc0lmc1lnc2loc3lpc4lqc5lrc6lsc7ltc8luc9lvc!lwc#lxc$lyc%l'lzc(l)l*l+lcgAcBcdgegfggghgigjgkg5g6gQc:bpb*bqb;b+brb3b+c,c-ci 4bj k l m n o p q r 5bs t u v Vbw x Wb6bRc=bsb,btb?b-bub7b.c/c:cy 8bz A B C D E F G H 9bI J K L XbM N Yb!bSc@bvb.bwb[b/bxb#b;c=c?cO $bP Q R S T U V W X %bY Z 0 1 Zb2 3 0b'blg@c,l-l0e.l/l:l;l=l?l@lTc[l1e]l[c^l2e_l`l{l|l}l~lambmcmdmemfmgmhmimjmkmlmmmnmompmqmrmsmtmumvmwmxmymzmAmBmCmDmEmFmGmHmImJmKmLmMmNmOmPmQmRmSmTmUmVmWmmgng3eXm]cYmZm0m1m2mCc3m*d4m5m6m7m8m9mNb!m#m1b$m%m'm(m)m*m+m,m-m.m/m:m;mXc=mYc?m@m[m#gZc]m$g0c^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~uavbvcvdvevfvgvhvivjvkvlvmvnvovpvqvrvsvtvuvvvwvxvyvzvAvBvCvDvEvFvGvHvIvJvKvLvMvNvOvPvQvRvSvTvUv@fVvWvXv[fYvZv0v1v2v3v4v5v6v7v8v9v!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~waxbxcxdxexfxgxhxixjxkxlxmxnxoxpxqxrxsxtxuxvxwxxxyxzxAxBxCxDxExFxGxHxIxJxKxLxMxNxOxPxQxRxSxTxUx1cVx2cWx3cXx4cYx-dZx.d0x/d1x_c2x`c3x{c4x:d5x;d6x=d7x?d8x7c9x8c!x|c#x@d$x}c%x[d'x]d(x^d)x_d*x~c+x`d,xad-xbd.xcd/xdd:x{d;x|d=x}d?x~d@xae[xbe]xce^xde_xee`xed{xfd|xgd}xc ~x5caybycydyeyfygyhyiyjykylymynyoyhdpyqyrysytyuyvywyxyyy6czyAyByCyDyEyFyGyHyIyJyKyLyMyNyOyPyQyRySyTyUyVyWyXyYyZy0y1y2y3y4y5y6y7y8y9y!y#y$y%y'y(y)y*y+y,y-y.y/y:y;y=y?y@y[y]y^y_y`y{y|y}y~yazbzczdzezfzgzhzizjzkzlzmznzozpzqzrzsztzuzvzwzxzyzzzAzBzCzDzEzFzGzHzIzJzKzLzMzNzOzPzQzRzSzTzUzVzWzXzYzZz0z1z2z3z4z5z6z7z8z9z!z#z$z%z'z(z)z*z+z,z-z.z/z:z;z=z?z@z[z]z^z_z`z{z|z}z~zaAbAcA

202 return _check_driver_error(error) 2a ZdLh0dMh1dNh9 Oh! Ph2dQh# Rh$ Sh% Th' Uh( Vh) Wh* Xh+ Yh, Zh- 0h. 1h/ 2h: 3h; 4h= 5h? 6h@ 7h[ 8h] 9h^ !h_ #h` $h{ %h| 'h} (h~ )hab*hbb+h,h-h.h/h:h;h=h?h@h[h]h^h_h`h{h|h}h~haiKebicidiLeeiMefiNegiOehiybiiPejiQekiReliminizboiAbpiBbqiCbriDbsiGctiHcuiIcviJcwiKcxiLcyiMcziNcAiBi-gCi.gDi/gEi_fFi`fGi{fHiIiJiKiLiMiNiOiPiQiRiSiTiUiViWiXiYiZi0i1i2i3i4i5i6i7i8i9i!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~jakbkckdkekfkgkhkikjkkklkmknkokpkqkrksk2btkukvkSewkxkykzkAkBkCkDkEkFkGkHkIkJkKkLkTeMkUeNkOkPkQkRk%cSk'cTkUkVkWkXkEbYkFbZkGb0k1kVe2k3kHb4kIb5kJb6kKb7kLb8k9k!k#k$kd %k'k(k)k(c)c*c3d4d5de h f g 4 cb5 6 dbebfb6d7d8d9d!d#dgb*khb+kMb,kib-kjb.kkb/klb:kmb;knb=k7 8 =f$d%d'd(dOcPc)d)bXePbQbRbSbTbUb?k@k[k]k^k_k`k{k|k}k~kalblcldlelflglhliljl1gklllml2gnl3golplqlrlsltlulvlwlxlylzlAlBlClDlEl4gob]b^bFlGlHlIlJlKlLlMlNlOlPlQlRlSlTlUlVlWlXlYlZl0l1l2l3l4l5l6l7l8l9l!l#l$l%l'l(l)l*l+lcgAcBcdgegfggghgigjgkg5g6gQc:bpb*bqb;b+brb3b+c,c-ci 4bj k l m n o p q r 5bs t u v Vbw x Wb6bRc=bsb,btb?b-bub7b.c/c:cy 8bz A B C D E F G H 9bI J K L XbM N Yb!bSc@bvb.bwb[b/bxb#b;c=c?cO $bP Q R S T U V W X %bY Z 0 1 Zb2 3 0b'blg,l-l0e.l/l:l;l=l?l@l[l]l^l_l`l{l|l}l~lambmcmdmemfmgmhmimjmkmlmmmnmompmqmrmsmtmumvmwmxmymzmAmBmCmDmEmFmGmHmImJmKmLmMmNmOmPmQmRmSmTmUmVmWmmgng3eXmYmZm0m1m2m3m*d4m5m6m7m8m9mNb!m#m1b$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~uavbvcvdvevfvgvhvivjvkvlvmvnvovpvqvrvsvtvuvvvwvxvyvzvAvBvCvDvEvFvGvHvIvJvKvLvMvNvOvPvQvRvSvTvUvVvWvXvYvZv0v1v2v3v4v5v6v7v8v9v!v#v$v%v'v(v)v*v+v,v-v.v/v:v;v=v?v@v[v]v^v_v`v{v|v}v~vawbwcwdwewfwgwhwiwjwkwlwmwnwowpwqwrwswtwuwvwwwxwywzwAwBwCwDwEwFwGwHwIwJwKwLwMwNwOwPwQwRwSwTwUwVwWwXwYwZw0w1w2w3w4w5w6w7w8w9w!w#w$w%w'w(w)w*w+w,w-w.w/w:w;w=w?w@w[w]w^w_w`w{w|w}w~waxbxcxdxexfxgxhxixjxkxlxmxnxoxpxqxrxsxtxuxvxwxxxyxzxAxBxCxDxExFxGxHxIxJxKxLxMxNxOxPxQxRxSxTxUxVxWxXxYxZx0x1x2x3x4x5x6x7x8x9x!x#x$x%x'x(x)x*x+x,x-x.x/x:x;x=x?x@x[x]x^x_x`x{x|x}x~xaybycydyeyfygyhyiyjykylymynyoypyqyrysytyuyvywyxyyyzyAyByCyDyEyFyGyHyIyJyKyLyMyNyOyPyQyRySyTyUyVyWyXyYyZy0y1y2y3y4y5y6y7y8y9y!y#y$y%y'y(y)y*y+y,y-y.y/y:y;y=y?y@y[y]y^y_y`y{y|y}y~yazbzczdzezfzgzhzizjzkzlzmznzozpzqzrzsztzuzvzwzxzyzzzAzBzCzDzEzFzGzHzIzJzKzLzMzNzOzPzQzRzSzTzUzVzWzXzYzZz0z1z2z3z4z5z6z7z8z9z!z#z$z%z'z(z)z*z+z,z-z.z/z:z;z=z?z@z[z]z^z_z`z{z|z}z~zaAbAcA

203 elif isinstance(error, runtime.cudaError_t): 2a PbQbRbSbTbUb;g|f=g}fYeZe~fagbg]b^bb _b`b{b|b}b~bacbcccdcecfcgchcicjckclcmcncocpcqcrcsctcucvcwcxcyczcAcBcpbqbrb3bi j k l m n o p q r s t u v w x sbtbub7by z A B C D E F G H I J K L M N vbwbxb#bO P Q R S T U V W X Y Z 0 1 2 3 @cTc1e[c2e]cCcXcYc#gZc$g0c@f[f1c2c3c4c-d.d/d_c`c{c:d;d=d?d7c8c|c@d}c[d]d^d_d~c`dadbdcddd{d|d}d~daebecedeeeedfdgdc 5chd6c

204 return _check_runtime_error(error) 2PbQbRbSbTbUb;g|f=g}fYeZe~fagbg]b^bAcBcpbqbrb3bi j k l m n o p q r s t u v w x sbtbub7by z A B C D E F G H I J K L M N vbwbxb#bO P Q R S T U V W X Y Z 0 1 2 3 @cTc1e[c2e]c

205 elif isinstance(error, nvrtc.nvrtcResult): 2a b _b`b{b|b}b~bacbcccdcecfcgchcicjckclcmcncocpcqcrcsctcucvcwcxcyczcCcXcYc#gZc$g0c@f[f1c2c3c4c-d.d/d_c`c{c:d;d=d?d7c8c|c@d}c[d]d^d_d~c`dadbdcddd{d|d}d~daebecedeeeedfdgdc 5chd6c

206 return _check_nvrtc_error(error, handle=handle) 2a b _b`b{b|b}b~bacbcccdcecfcgchcicjckclcmcncocpcqcrcsctcucvcwcxcyczcCcXcYc#gZc$g0c@f[f1c2c3c4c-d.d/d_c`c{c:d;d=d?d7c8c|c@d}c[d]d^d_d~c`dadbdcddd{d|d}d~daebecedeeeedfdgdc 5chd6c

207 else: 

208 raise RuntimeError(f"Unknown error type: {error}") 

209  

210  

211def handle_return(tuple result, handle=None): 

212 _check_error(result[0], handle=handle) 2a ZdLh0dMh1dNh9 Oh! Ph2dQh# Rh$ Sh% Th' Uh( Vh) Wh* Xh+ Yh, Zh- 0h. 1h/ 2h: 3h; 4h= 5h? 6h@ 7h[ 8h] 9h^ !h_ #h` $h{ %h| 'h} (h~ )hab*hbb+h,h-h.h/h:h;h=h?h@h[h]h^h_h`h{h|h}h~haiKebicidiLeeiMefiNegiOehiybiiPejiQekiReliminizboiAbpiBbqiCbriDbsiGctiHcuiIcviJcwiKcxiLcyiMcziNcAiBi-gCi.gDi/gEi_fFi`fGi{fHiIiJiKiLiMiNiOiPiQiRiSiTiUiViWiXiYiZi0i1i2i3i4i5i6i7i8i9i!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~jakbkckdkekfkgkhkikjkkklkmknkokpkqkrksk2btkukvkSewkxkykzkAkBkCkDkEkFkGkHkIkJkKkLkTeMkUeNkOkPkQkRk%cSk'cTkUkVkWkXkEbYkFbZkGb0k1kVe2k3kHb4kIb5kJb6kKb7kLb8k9k!k#k$kd %k'k(k)k(c)c*c3d4d5de h f g 4 cb5 6 dbebfb6d7d8d9d!d#dgb*khb+kMb,kib-kjb.kkb/klb:kmb;knb=k7 8 =f$d%d'd(dOcPc)d)bXePbQbRbSbTbUb;g|f?k=g@k[k]k^k_k`k{k}fYe|kZe}k~kalblcldlelfl~fglhliljl1gklllml2gagbgnl3golplqlrlsltlulvlwlxlylzlAlBlClDlEl4gob]b^bFlGlHlIlb _bJl`bKl{bLl|bMl}bNl~bOlacPlbcQlccRldcSlecTlfcUlgcVlhcWlicXljcYlkcZllc0lmc1lnc2loc3lpc4lqc5lrc6lsc7ltc8luc9lvc!lwc#lxc$lyc%l'lzc(l)l*l+lcgAcBcdgegfggghgigjgkg5g6gQc:bpb*bqb;b+brb3b+c,c-ci 4bj k l m n o p q r 5bs t u v Vbw x Wb6bRc=bsb,btb?b-bub7b.c/c:cy 8bz A B C D E F G H 9bI J K L XbM N Yb!bSc@bvb.bwb[b/bxb#b;c=c?cO $bP Q R S T U V W X %bY Z 0 1 Zb2 3 0b'blg@c,l-l0e.l/l:l;l=l?l@lTc[l1e]l[c^l2e_l`l{l|l}l~lambmcmdmemfmgmhmimjmkmlmmmnmompmqmrmsmtmumvmwmxmymzmAmBmCmDmEmFmGmHmImJmKmLmMmNmOmPmQmRmSmTmUmVmWmmgng3eXm]cYmZm0m1m2mCc3m*d4m5m6m7m8m9mNb!m#m1b$m%m'm(m)m*m+m,m-m.m/m:m;mXc=mYc?m@m[m#gZc]m$g0c^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~uavbvcvdvevfvgvhvivjvkvlvmvnvovpvqvrvsvtvuvvvwvxvyvzvAvBvCvDvEvFvGvHvIvJvKvLvMvNvOvPvQvRvSvTvUv@fVvWvXv[fYvZv0v1v2v3v4v5v6v7v8v9v!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~waxbxcxdxexfxgxhxixjxkxlxmxnxoxpxqxrxsxtxuxvxwxxxyxzxAxBxCxDxExFxGxHxIxJxKxLxMxNxOxPxQxRxSxTxUx1cVx2cWx3cXx4cYx-dZx.d0x/d1x_c2x`c3x{c4x:d5x;d6x=d7x?d8x7c9x8c!x|c#x@d$x}c%x[d'x]d(x^d)x_d*x~c+x`d,xad-xbd.xcd/xdd:x{d;x|d=x}d?x~d@xae[xbe]xce^xde_xee`xed{xfd|xgd}xc ~x5caybycydyeyfygyhyiyjykylymynyoyhdpyqyrysytyuyvywyxyyy6czyAyByCyDyEyFyGyHyIyJyKyLyMyNyOyPyQyRySyTyUyVyWyXyYyZy0y1y2y3y4y5y6y7y8y9y!y#y$y%y'y(y)y*y+y,y-y.y/y:y;y=y?y@y[y]y^y_y`y{y|y}y~yazbzczdzezfzgzhzizjzkzlzmznzozpzqzrzsztzuzvzwzxzyzzzAzBzCzDzEzFzGzHzIzJzKzLzMzNzOzPzQzRzSzTzUzVzWzXzYzZz0z1z2z3z4z5z6z7z8z9z!z#z$z%z'z(z)z*z+z,z-z.z/z:z;z=z?z@z[z]z^z_z`z{z|z}z~zaAbAcA

213 cdef int out_len = len(result) 2a ZdLh0dMh1dNh9 Oh! Ph2dQh# Rh$ Sh% Th' Uh( Vh) Wh* Xh+ Yh, Zh- 0h. 1h/ 2h: 3h; 4h= 5h? 6h@ 7h[ 8h] 9h^ !h_ #h` $h{ %h| 'h} (h~ )hab*hbb+h,h-h.h/h:h;h=h?h@h[h]h^h_h`h{h|h}h~haiKebicidiLeeiMefiNegiOehiybiiPejiQekiReliminizboiAbpiBbqiCbriDbsiGctiHcuiIcviJcwiKcxiLcyiMcziNcAiBi-gCi.gDi/gEi_fFi`fGi{fHiIiJiKiLiMiNiOiPiQiRiSiTiUiViWiXiYiZi0i1i2i3i4i5i6i7i8i9i!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~jakbkckdkekfkgkhkikjkkklkmknkokpkqkrksk2btkukvkSewkxkykzkAkBkCkDkEkFkGkHkIkJkKkLkTeMkUeNkOkPkQkRk%cSk'cTkUkVkWkXkEbYkFbZkGb0k1kVe2k3kHb4kIb5kJb6kKb7kLb8k9k!k#k$kd %k'k(k)k(c)c*c3d4d5de h f g 4 cb5 6 dbebfb6d7d8d9d!d#dgb*khb+kMb,kib-kjb.kkb/klb:kmb;knb=k7 8 =f$d%d'd(dOcPc)d)bXePbQbRbSbTbUb;g|f?k=g@k[k]k^k_k`k{k}fYe|kZe}k~kalblcldlelfl~fglhliljl1gklllml2gagbgnl3golplqlrlsltlulvlwlxlylzlAlBlClDlEl4gob]b^bFlGlHlIlb _bJl`bKl{bLl|bMl}bNl~bOlacPlbcQlccRldcSlecTlfcUlgcVlhcWlicXljcYlkcZllc0lmc1lnc2loc3lpc4lqc5lrc6lsc7ltc8luc9lvc!lwc#lxc$lyc%l'lzc(l)l*l+lcgAcBcdgegfggghgigjgkg5g6gQc:bpb*bqb;b+brb3b+c,c-ci 4bj k l m n o p q r 5bs t u v Vbw x Wb6bRc=bsb,btb?b-bub7b.c/c:cy 8bz A B C D E F G H 9bI J K L XbM N Yb!bSc@bvb.bwb[b/bxb#b;c=c?cO $bP Q R S T U V W X %bY Z 0 1 Zb2 3 0b'blg@c,l-l0e.l/l:l;l=l?l@lTc[l1e]l[c^l2e_l`l{l|l}l~lambmcmdmemfmgmhmimjmkmlmmmnmompmqmrmsmtmumvmwmxmymzmAmBmCmDmEmFmGmHmImJmKmLmMmNmOmPmQmRmSmTmUmVmWmmgng3eXm]cYmZm0m1m2mCc3m*d4m5m6m7m8m9mNb!m#m1b$m%m'm(m)m*m+m,m-m.m/m:m;mXc=mYc?m@m[m#gZc]m$g0c^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~uavbvcvdvevfvgvhvivjvkvlvmvnvovpvqvrvsvtvuvvvwvxvyvzvAvBvCvDvEvFvGvHvIvJvKvLvMvNvOvPvQvRvSvTvUv@fVvWvXv[fYvZv0v1v2v3v4v5v6v7v8v9v!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~waxbxcxdxexfxgxhxixjxkxlxmxnxoxpxqxrxsxtxuxvxwxxxyxzxAxBxCxDxExFxGxHxIxJxKxLxMxNxOxPxQxRxSxTxUx1cVx2cWx3cXx4cYx-dZx.d0x/d1x_c2x`c3x{c4x:d5x;d6x=d7x?d8x7c9x8c!x|c#x@d$x}c%x[d'x]d(x^d)x_d*x~c+x`d,xad-xbd.xcd/xdd:x{d;x|d=x}d?x~d@xae[xbe]xce^xde_xee`xed{xfd|xgd}xc ~x5caybycydyeyfygyhyiyjykylymynyoyhdpyqyrysytyuyvywyxyyy6czyAyByCyDyEyFyGyHyIyJyKyLyMyNyOyPyQyRySyTyUyVyWyXyYyZy0y1y2y3y4y5y6y7y8y9y!y#y$y%y'y(y)y*y+y,y-y.y/y:y;y=y?y@y[y]y^y_y`y{y|y}y~yazbzczdzezfzgzhzizjzkzlzmznzozpzqzrzsztzuzvzwzxzyzzzAzBzCzDzEzFzGzHzIzJzKzLzMzNzOzPzQzRzSzTzUzVzWzXzYzZz0z1z2z3z4z5z6z7z8z9z!z#z$z%z'z(z)z*z+z,z-z.z/z:z;z=z?z@z[z]z^z_z`z{z|z}z~zaAbAcA

214 if out_len == 1: 2a ZdLh0dMh1dNh9 Oh! Ph2dQh# Rh$ Sh% Th' Uh( Vh) Wh* Xh+ Yh, Zh- 0h. 1h/ 2h: 3h; 4h= 5h? 6h@ 7h[ 8h] 9h^ !h_ #h` $h{ %h| 'h} (h~ )hab*hbb+h,h-h.h/h:h;h=h?h@h[h]h^h_h`h{h|h}h~haiKebicidiLeeiMefiNegiOehiybiiPejiQekiReliminizboiAbpiBbqiCbriDbsiGctiHcuiIcviJcwiKcxiLcyiMcziNcAiBi-gCi.gDi/gEi_fFi`fGi{fHiIiJiKiLiMiNiOiPiQiRiSiTiUiViWiXiYiZi0i1i2i3i4i5i6i7i8i9i!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~jakbkckdkekfkgkhkikjkkklkmknkokpkqkrksk2btkukvkSewkxkykzkAkBkCkDkEkFkGkHkIkJkKkLkTeMkUeNkOkPkQkRk%cSk'cTkUkVkWkXkEbYkFbZkGb0k1kVe2k3kHb4kIb5kJb6kKb7kLb8k9k!k#k$kd %k'k(k)k(c)c*c3d4d5de h f g 4 cb5 6 dbebfb6d7d8d9d!d#dgb*khb+kMb,kib-kjb.kkb/klb:kmb;knb=k7 8 =f$d%d'd(dOcPc)d)bXePbQbRbSbTbUb;g|f?k=g@k[k]k^k_k`k{k}fYe|kZe}k~kalblcldlelfl~fglhliljl1gklllml2gagbgnl3golplqlrlsltlulvlwlxlylzlAlBlClDlEl4gob]b^bFlGlHlIlb _bJl`bKl{bLl|bMl}bNl~bOlacPlbcQlccRldcSlecTlfcUlgcVlhcWlicXljcYlkcZllc0lmc1lnc2loc3lpc4lqc5lrc6lsc7ltc8luc9lvc!lwc#lxc$lyc%l'lzc(l)l*l+lcgAcBcdgegfggghgigjgkg5g6gQc:bpb*bqb;b+brb3b+c,c-ci 4bj k l m n o p q r 5bs t u v Vbw x Wb6bRc=bsb,btb?b-bub7b.c/c:cy 8bz A B C D E F G H 9bI J K L XbM N Yb!bSc@bvb.bwb[b/bxb#b;c=c?cO $bP Q R S T U V W X %bY Z 0 1 Zb2 3 0b'blg@c,l-l0e.l/l:l;l=l?l@lTc[l1e]l[c^l2e_l`l{l|l}l~lambmcmdmemfmgmhmimjmkmlmmmnmompmqmrmsmtmumvmwmxmymzmAmBmCmDmEmFmGmHmImJmKmLmMmNmOmPmQmRmSmTmUmVmWmmgng3eXm]cYmZm0m1m2mCc3m*d4m5m6m7m8m9mNb!m#m1b$m%m'm(m)m*m+m,m-m.m/m:m;mXc=mYc?m@m[m#gZc]m$g0c^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~uavbvcvdvevfvgvhvivjvkvlvmvnvovpvqvrvsvtvuvvvwvxvyvzvAvBvCvDvEvFvGvHvIvJvKvLvMvNvOvPvQvRvSvTvUv@fVvWvXv[fYvZv0v1v2v3v4v5v6v7v8v9v!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~waxbxcxdxexfxgxhxixjxkxlxmxnxoxpxqxrxsxtxuxvxwxxxyxzxAxBxCxDxExFxGxHxIxJxKxLxMxNxOxPxQxRxSxTxUx1cVx2cWx3cXx4cYx-dZx.d0x/d1x_c2x`c3x{c4x:d5x;d6x=d7x?d8x7c9x8c!x|c#x@d$x}c%x[d'x]d(x^d)x_d*x~c+x`d,xad-xbd.xcd/xdd:x{d;x|d=x}d?x~d@xae[xbe]xce^xde_xee`xed{xfd|xgd}xc ~x5caybycydyeyfygyhyiyjykylymynyoyhdpyqyrysytyuyvywyxyyy6czyAyByCyDyEyFyGyHyIyJyKyLyMyNyOyPyQyRySyTyUyVyWyXyYyZy0y1y2y3y4y5y6y7y8y9y!y#y$y%y'y(y)y*y+y,y-y.y/y:y;y=y?y@y[y]y^y_y`y{y|y}y~yazbzczdzezfzgzhzizjzkzlzmznzozpzqzrzsztzuzvzwzxzyzzzAzBzCzDzEzFzGzHzIzJzKzLzMzNzOzPzQzRzSzTzUzVzWzXzYzZz0z1z2z3z4z5z6z7z8z9z!z#z$z%z'z(z)z*z+z,z-z.z/z:z;z=z?z@z[z]z^z_z`z{z|z}z~zaAbAcA

215 return 2Zd0d1d9 ! 2d# $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbKeLeMeNeOeybPeQeRezbAbBbCbDbGcHcIcJcKcLcMcNc-g.g/g_f`f{f2bSeTeUe%c'cEbFbGbVeHbIbJbKbLbd (c)c*c3d4d5de h f g 4 cb5 6 dbebfb6d7d8d9d!d#dgbhbMbibjbkblbmbnb=fXePbQbRbSbTbUbYeZe4gob]b^bcgAcBcdgegfggghgigjgkg5g6gQc:bpb*bqb;b+brb3b+c,c-ci 4bj k l m n o p q r 5bs t u v Vbw x Wb6bRc=bsb,btb?b-bub7b.c/c:cy 8bz A B C D E F G H 9bI J K L XbM N Yb!bSc@bvb.bwb[b/bxb#b;c=c?cO $bP Q R S T U V W X %bY Z 0 1 Zb2 3 0b'blg@c0eTc1e[c2emgng3e]cNb

216 elif out_len == 2: 

217 return result[1] 2a ZdLh0dMh1dNh9 Oh! Ph2dQh# Rh$ Sh% Th' Uh( Vh) Wh* Xh+ Yh, Zh- 0h. 1h/ 2h: 3h; 4h= 5h? 6h@ 7h[ 8h] 9h^ !h_ #h` $h{ %h| 'h} (h~ )hab*hbb+h,h-h.h/h:h;h=h?h@h[h]h^h_h`h{h|h}h~haiKebicidiLeeiMefiNegiOehiybiiPejiQekiReliminizboiAbpiBbqiCbriDbsiGctiHcuiIcviJcwiKcxiLcyiMcziNcAiBi-gCi.gDi/gEi_fFi`fGi{fHiIiJiKiLiMiNiOiPiQiRiSiTiUiViWiXiYiZi0i1i2i3i4i5i6i7i8i9i!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~jakbkckdkekfkgkhkikjkkklkmknkokpkqkrksk2btkukvkSewkxkykzkAkBkCkDkEkFkGkHkIkJkKkLkTeMkUeNkOkPkQkRk%cSk'cTkUkVkWkXkEbYkFbZkGb0k1kVe2k3kHb4kIb5kJb6kKb7kLb8k9k!k#k$kd %k'k(k)k(c)c*c3d4d5de h f g 4 cb5 6 dbebfb6d7d8d9d!d#dgb*khb+kMb,kib-kjb.kkb/klb:kmb;knb=k7 8 $d%d'd(dOcPc)d)bXePbQbRbSbTbUb;g|f?k=g@k[k]k^k_k`k{k}f|k}k~kalblcldlelfl~fglhliljl1gklllml2gagbgnl3golplqlrlsltlulvlwlxlylzlAlBlClDlElob]b^bFlGlHlIlJlKlLlMlNlOlPlQlRlSlTlUlVlWlXlYlZl0l1l2l3l4l5l6l7l8l9l!l#l$l%l'l(l)l*l+lcgAcBcdgegfggghgigjgkgQc:bpb*bqb;b+brb3b+c,c-ci 4bj k l m n o p q r 5bs t u v Vbw x Wb6bRc=bsb,btb?b-bub7b.c/c:cy 8bz A B C D E F G H 9bI J K L XbM N Yb!bSc@bvb.bwb[b/bxb#b;c=c?cO $bP Q R S T U V W X %bY Z 0 1 Zb2 3 0b'blg,l-l0e.l/l:l;l=l?l@l[l]l^l_l`l{l|l}l~lambmcmdmemfmgmhmimjmkmlmmmnmompmqmrmsmtmumvmwmxmymzmAmBmCmDmEmFmGmHmImJmKmLmMmNmOmPmQmRmSmTmUmVmWmmgng3eXmYmZm0m1m2m3m*d4m5m6m7m8m9m!m#m1b$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~uavbvcvdvevfvgvhvivjvkvlvmvnvovpvqvrvsvtvuvvvwvxvyvzvAvBvCvDvEvFvGvHvIvJvKvLvMvNvOvPvQvRvSvTvUvVvWvXvYvZv0v1v2v3v4v5v6v7v8v9v!v#v$v%v'v(v)v*v+v,v-v.v/v:v;v=v?v@v[v]v^v_v`v{v|v}v~vawbwcwdwewfwgwhwiwjwkwlwmwnwowpwqwrwswtwuwvwwwxwywzwAwBwCwDwEwFwGwHwIwJwKwLwMwNwOwPwQwRwSwTwUwVwWwXwYwZw0w1w2w3w4w5w6w7w8w9w!w#w$w%w'w(w)w*w+w,w-w.w/w:w;w=w?w@w[w]w^w_w`w{w|w}w~waxbxcxdxexfxgxhxixjxkxlxmxnxoxpxqxrxsxtxuxvxwxxxyxzxAxBxCxDxExFxGxHxIxJxKxLxMxNxOxPxQxRxSxTxUxVxWxXxYxZx0x1x2x3x4x5x6x7x8x9x!x#x$x%x'x(x)x*x+x,x-x.x/x:x;x=x?x@x[x]x^x_x`x{x|x}x~xaybycydyeyfygyhyiyjykylymynyoypyqyrysytyuyvywyxyyyzyAyByCyDyEyFyGyHyIyJyKyLyMyNyOyPyQyRySyTyUyVyWyXyYyZy0y1y2y3y4y5y6y7y8y9y!y#y$y%y'y(y)y*y+y,y-y.y/y:y;y=y?y@y[y]y^y_y`y{y|y}y~yazbzczdzezfzgzhzizjzkzlzmznzozpzqzrzsztzuzvzwzxzyzzzAzBzCzDzEzFzGzHzIzJzKzLzMzNzOzPzQzRzSzTzUzVzWzXzYzZz0z1z2z3z4z5z6z7z8z9z!z#z$z%z'z(z)z*z+z,z-z.z/z:z;z=z?z@z[z]z^z_z`z{z|z}z~zaAbAcA

218 else: 

219 return result[1:] 2a Zd0d1d9 ! 2d# $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbb(c)c*c3d4d5de f g 4 5 6 dbebfb6d7d8d9d!d#dgbhbibjbkblbmbnbb _b`b{b|b}b~bacbcccdcecfcgchcicjckclcmcncocpcqcrcsctcucvcwcxcyczcCcXcYc#gZc$g0c@f[f1c2c3c4c-d.d/d_c`c{c:d;d=d?d7c8c|c@d}c[d]d^d_d~c`dadbdcddd{d|d}d~daebecedeeeedfdgdc 5chd6c

220  

221  

222cpdef check_or_create_options(type cls, options, str options_description="", bint keep_none=False): 

223 """ 

224 Create the specified options dataclass from a dictionary of options or None. 

225 """ 

226 if options is None: 2Zd0d1d9 ! 2d# $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbKeLeMeNeOeybPeQeRezbAbBbCbDbGcHcIcJcKcLcMcNc,e-e.e=C?Cid/e]C^Cjd:e{C|Ckd;eaDbDld=eeDfDmd?eiDjDnd2bSe@e[e]e^eTeUe_emDnDoD%c'cpDqDhA`eEbFbGb{eVeHbIbJbKbLbd (c)c*c3d4d5de h f g 4 cb5 6 dbebfb6d7d8d9d!d#dgbhbMbibjbkblbmbnbrDsDiAtDjAuDvDwDkAxDlAyD7 zD8 ADXgYg)g*gZg0g+g,gphqhmABDCDDDEDFDGDHDIDJDKDLDMDNDODPDQDRDSDTDUDVDWDXDYDZD0D1D2D3D4D5D6D7D8D$d9D%d!D'd#D(d$DOcPc)d%D)bXe'D(D)D*DPbQbRbSbTbUb+D,D-D.D/D:D3CYeZe{I;DnADcoA|I=D}I?D@D[D]D^D_DEc`D{D|D}DpAqArAfe~DaEbEcEdEeEfEgEhEiEjEkElEob]b^bodFcpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdgeheb _bzI`bAI{bBI|bCI}bDI~bEIacFIbcGIccHIdcIIecJIfc.Cgc/ChcKIicLIjcMIkcNIlcOImcPIncQIocRIpcSIqcTIrcUIscVItcWIuc+evcXIwcYIxcZIyc0IsA1IzcAcBcmEnEQc:bpb*bqb;b+brb3b+c,c-ci 4bj k l m n o p q r 5bs t u v Vbw x Wb6bRc=bsb,btb?b-bub7b.c/c:cy 8bz A B C D E F G H 9bI J K L XbM N Yb!bSc@bvb.bwb[b/bxb#b;c=c?cO $bP Q R S T U V W X %bY Z 0 1 Zb2 3 0b'boEpE@c0etAuAqErEvAwAxAyATc1e[c2esEtEuEvEwExEyEzEAEBECEDEEEFEGEHEzAIEAAJEBAKECALEDAMEEANEFAOEGAPEQERESETEUEVEWEXEYEZE0E1E2E3E4E5E7Crh8C9Csh!C6E3e7E8E9E]c!E#E$E%E'E(EUcVcObWcCc*dHAIAJAKALANbMA1bNAOAPAQARASATAUAVAWAXAYAZAXcYc0A1AZc0c2A3A4A5A6A7A8A9A!A#A$A%A'A(A)A*A+A,A-A.A/A:A;A=A?A@A[A]A^A_A`A{A|A}A~AaBbBcBdBeBfBgBhBiBjBkBlBmBnBoBpBqBrBsBtBuBvBwBxByBzBABBBCBDB?f4e5e^c+d)E|e*E+E,EEB}eFB.E/EGB:E;E~e=E?E@EHBafIB]Ebf^E_Ecf{Eie|E}E#C$C~EaFbFcFdFthje6eeFuhvhfFgFhFkeledfefffgfhfifiFjFjflFkfmFnFJBlfKBpFmfqFrFsFLBnfMBuFofvFwFpfyF%CmezFAFBFCFDFEFFFGFHFIFNBne7eJFwhOBKFLFMFNFOFPFQFRFSFTFUFVFWFXFqfYFZFrf1F2FsfPBoeQBRBSBTBUBVBWBXBYBpeZB7g0B1B2BtfqereufvfwfxfyfRdzfAf9c8e9eSd!cTdBf4F5FUdxhyh3Bse4B5B6B7B8B9B!B#B$B%Bte'B8g(B)B*BCf$F%FDf(F'Cue)F*F+F,F-F.F/F:F;F=F?F+Bve#e@Fzh,B[F]F^F_FEf`F{F|F-BFf.B~FGfaGbGHfdGweeGfG(C)CgGhGiGjGkGAhxe$elGBhChmGnGoGyezeIfJfKfLfMfNfpGqGOfsGPftGuG/BQf:BwGRfxGyGzG;BSf=BBGTfCGDGUfFG*CAeGGHGIGJGKGLGMGNGOGPG?BBe%eQGDh@BRGSGTGUGVGWGXGYGZG0G1G2G3G4GVf5G6GWf8G9GXf[BCe]B^B_B`B{B|B}B~BaCDebC9gcCdCeCYfEeFeZf0f1f2f3fVd4f5f#c'e(eWd$cXd6f#G$GYdEhFhfCGegChCiCjCkClCmCnCoCpCHeqC!grCsCtC7f,G-G8f/G+CIe:G;G=G?G@G[G]G^G_G`G{GuCJe*e|GGhvC}G~GaHbHcHHhIhwCxCdH9feHfHgHyC!fzCACBCCCiHjH#fkHlHmHDC$fECoHpH%fqHrHsHFC'fGCuHvH(fwHxHyHHC)fICAHBH*fCHDHEHJC+fKCGHHH,fIHJHKHLC-fMCMHNH.fOHPHQHNC/fOCPCSHTH:fUHVHWHQC;fRCSCYHjIkI1c2c,C3c4c-d.d/d_c`c{c:d;d=d?d7c8c|c@d}c[d]d^d_d~c`dadbdcddd{d|d}d~daebecedeeeedfdgdc 5c(bTC,dlImInIoIpIqIrIsItIuIvIwIxIyIJhKhZHUCVCWChdXC0H1HYC-C0C2I3I4I5I6I7I8I9I!I6c2H3H4H5H6H7H8H9H!H#H$H%H'H(H)H*HZC+H,H-H.H/H:H;H=H?H@H[H]H^H_H`H{H|H}H~HaIbIcIdIeIfIgIhIiI

227 if keep_none: 2Zd0d1d9 ! 2d# $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbKeLeMeNeOeybPeQeRezbAbBbCbDbGcHcIcJcKcLcMcNc=C?C]C^C{C|CaDbDeDfDiDjDSeTeUe%c'cEbFbGbVeHbIbJbKbLbd (c)c*c3d4d5de h f g 4 cb5 6 dbebfb6d7d8d9d!d#dgbhbMbibjbkblbmbnb7 8 XgYg)g*gZg0g+g,gphqhmA$d%d'd(dOcPc)d)bPbQbRbSbTbUb+D,D-D.D/D:DYeZe{I;DnADcoA=D?D@D[D]D^D_DEc`D{D|D}DpAqArAfe~DaEbEcEdEeEfEgEhEiEjEkElEob]b^bFcAcBcmEnEQc:bpb*bqb;b+brb3b+c,c-ci 4bj k l m n o p q r 5bs t u v Vbw x Wb6bRc=bsb,btb?b-bub7b.c/c:cy 8bz A B C D E F G H 9bI J K L XbM N Yb!bSc@bvb.bwb[b/bxb#b;c=c?cO $bP Q R S T U V W X %bY Z 0 1 Zb2 3 0b'boEpE@c0etAuAqErEvAwAxAyATc1e[c2ezAIEAAJEBAKECALEDAMEEANEFAOEGAPErhsh3e7E8E9E]c*dHAIAJAKALANbMA1bNAOAPAQARASATAUAVAWAXAYAZAZc0c2A3A4A5A6A7A8A9A!A#A$A%A'A(A)A*A+A,A-A.A/A:A;A=A?A@A[A]A^A_A`A{A|A}A~AaBbBcBdBeBfBgBhBiBjBkBlBmBnBoBpBqBrBsBtBuBvBwBxByBzBABBBCBDB^c)E*E+E,EEBFB.E/EGB:E;E=E?E@EHBIB]E^E_E{Eie|E}E#C$C~EaFbFcFdFthje6eeFuhvhfFgFhFkeleiFjFlFmFnFJBKBpFqFrFsFLBMBuFvFwFyF%CmezFAFBFCFDFEFFFGFHFIFNBne7eJFwhOBKFLFMFNFOFPFQFRFSFTFUFVFWFXFYFZF1F2FPBoeQBRBSBTBUBVBWBXBYBpeZB7g0B1B2Bqere8e9eSdTd4F5Fxhyh3Bse4B5B6B7B8B9B!B#B$B%Bte'B8g(B)B*B$F%F(F'Cue)F*F+F,F-F.F/F:F;F=F?F+Bve#e@Fzh,B[F]F^F_F`F{F|F-B.B~FaGbGdGweeGfG(C)CgGhGiGjGkGAhxe$elGBhChmGnGoGyezepGqGsGtGuG/B:BwGxGyGzG;B=BBGCGDGFG*CAeGGHGIGJGKGLGMGNGOGPG?BBe%eQGDh@BRGSGTGUGVGWGXGYGZG0G1G2G3G4G5G6G8G9G[BCe]B^B_B`B{B|B}B~BaCDebC9gcCdCeCEeFe'e(eWdXd#G$GEhFhfCGegChCiCjCkClCmCnCoCpCHeqC!grCsCtC,G-G/G+CIe:G;G=G?G@G[G]G^G_G`G{GuCJe*e|GGhvC}G~GaHHhIhwCxCdHeHfHgHyCzCACBCCCiHjHkHlHmHDCECoHpHqHrHsHFCGCuHvHwHxHyHHCICAHBHCHDHEHJCKCGHHHIHJHKHLCMCMHNHOHPHQHNCOCPCSHTHUHVHWHQCRCSCYH2cc 5c(bTC,dZHUCVCWCXC0H1HYC0C6c4H5H6H7H!H#H$H%H)HZC+H.H/H;H=H?H@H[H]H^H_H`H{H|H}H~HaIbIcIdIeIfIgIhIiI

228 return options 2(c)c*ce h f g 4 cb5 6 YeZemEnEoEpE0etAuAqErEvAwAxAyATc1e[c2ezAIEAAJEBAKECALEDAMEEANEFAOEGAPErhsh7E8E9ENb)E/E;E{Eie|E}E#C$C~EaFbFcFdFthje6eeFuhvhfFgFhFlFpF%C'C_FdGweeGfG(C)CgGhGiGjGkGAhxe$elGBhChmGnGoGsGwG*C+CdHjHpHvHBHHHNHTH;H=H?H@H[H]H^H_H`H{H|H}H~HaIbIcIdIeIfIgIhIiI

229 return cls() 2Zd0d1d9 ! 2d# $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbKeLeMeNeOeybPeQeRezbAbBbCbDbGcHcIcJcKcLcMcNc=C?C]C^C{C|CaDbDeDfDiDjDSeTeUe%c'cEbFbGbVeHbIbJbKbLbd (c)c*c3d4d5de h f g 4 cb5 6 dbebfb6d7d8d9d!d#dgbhbMbibjbkblbmbnb7 8 XgYg)g*gZg0g+g,gphqhmA$d%d'd(dOcPc)d)bPbQbRbSbTbUb+D,D-D.D/D:D{I;DnADcoA=D?D@D[D]D^D_DEc`D{D|D}DpAqArAfe~DaEbEcEdEeEfEgEhEiEjEkElEob]b^bFcAcBcQc:bpb*bqb;b+brb3b+c,c-ci 4bj k l m n o p q r 5bs t u v Vbw x Wb6bRc=bsb,btb?b-bub7b.c/c:cy 8bz A B C D E F G H 9bI J K L XbM N Yb!bSc@bvb.bwb[b/bxb#b;c=c?cO $bP Q R S T U V W X %bY Z 0 1 Zb2 3 0b'b@cTc3e]c*dHAIAJAKALANbMA1bNAOAPAQARASATAUAVAWAXAYAZAZc0c2A3A4A5A6A7A8A9A!A#A$A%A'A(A)A*A+A,A-A.A/A:A;A=A?A@A[A]A^A_A`A{A|A}A~AaBbBcBdBeBfBgBhBiBjBkBlBmBnBoBpBqBrBsBtBuBvBwBxByBzBABBBCBDB^c*E+E,EEBFB.EGB:E=E?E@EHBIB]E^E_E#C$CthvhkeleiFjFmFnFJBKBqFrFsFLBMBuFvFwFyF%CmezFAFBFCFDFEFFFGFHFIFNBne7eJFwhOBKFLFMFNFOFPFQFRFSFTFUFVFWFXFYFZF1F2FPBoeQBRBSBTBUBVBWBXBYBpeZB7g0B1B2Bqere8e9eSdTd4F5Fxhyh3Bse4B5B6B7B8B9B!B#B$B%Bte'B8g(B)B*B$F%F(F'Cue)F*F+F,F-F.F/F:F;F=F?F+Bve#e@Fzh,B[F]F^F`F{F|F-B.B~FaGbG(C)CAhChyezepGqGtGuG/B:BxGyGzG;B=BBGCGDGFG*CAeGGHGIGJGKGLGMGNGOGPG?BBe%eQGDh@BRGSGTGUGVGWGXGYGZG0G1G2G3G4G5G6G8G9G[BCe]B^B_B`B{B|B}B~BaCDebC9gcCdCeCEeFe'e(eWdXd#G$GEhFhfCGegChCiCjCkClCmCnCoCpCHeqC!grCsCtC,G-G/G+CIe:G;G=G?G@G[G]G^G_G`G{GuCJe*e|GGhvC}G~GaHHhIhwCxCeHfHgHyCzCACBCCCiHkHlHmHDCECoHqHrHsHFCGCuHwHxHyHHCICAHCHDHEHJCKCGHIHJHKHLCMCMHOHPHQHNCOCPCSHUHVHWHQCRCSCYH2cc 5c(bTC,dZHUCVCWCXC0H1HYC0C6c4H5H6H7H!H#H$H%H)HZC+H.H/H

230 elif isinstance(options, cls): 29 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbybzbAbBbCbDb,e-e.eid/ejd:ekd;eld=emd?end2b@e[e]e^e_emDnDoD%c'cpDqDhA`eEbFbGb{eHbIbJbKbLbd e h f g 4 cb5 6 dbebfbgbhbMbibjbkblbmbnbrDsDiAtDjAuDvDwDkAxDlAyD7 zD8 ADXgYg)g*gZg0g+g,gphqhmABDCDDDEDFDGDHDIDJDKDLDMDNDODPDQDRDSDTDUDVDWDXDYDZD0D1D2D3D4D5D6D7D8D9D!D#D$DOcPc%DXe'D(D)D*DPbQbRbSbTbUb3CnADcoA|I}IEcpAqArAfeobodFcpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdgeheb _bzI`bAI{bBI|bCI}bDI~bEIacFIbcGIccHIdcIIecJIfc.Cgc/ChcKIicLIjcMIkcNIlcOImcPIncQIocRIpcSIqcTIrcUIscVItcWIuc+evcXIwcYIxcZIyc0IsA1Izc@ctAuAvAwAxAyATc[csEtEuEvEwExEyEzEAEBECEDEEEFEGEHEzAAABACADAEAFAGAQERESETEUEVEWEXEYEZE0E1E2E3E4E5E7Crh8C9Csh!C6E]c!E#E$E%E'E(EUcVcObWcCcNbXcYc0A1A?f4e5e^c+d|e}e~eafbfcfiejekeledfefffgfhfifjfkflfmfnfofpfmeneqfrfsfoepetfqereufvfwfxfyfRdzfAf9c!cBfUdseteCfDfueveEfFfGfHfwexeyezeIfJfKfLfMfNfOfPfQfRfSfTfUfAeBeVfWfXfCeDeYfEeFeZf0f1f2f3fVd4f5f#c$c6fYdGeHe7f8fIeJebHcHHhIh9f!f#f$f%f'f(f)f*f+f,f-f.f/f:f;fjIkI1c,C3c4c-d.d/d_c`c{c:d;d=d?d7c8c|c@d}c[d]d^d_d~c`dadbdcddd{d|d}d~daebecedeeeedfdgdlImInIoIpIqIrIsItIuIvIwIxIyIJhKhhd-C0C2I3I4I5I6I7I8I9I!I2H3H8H9H'H(H*HZC,H-H:H

231 return options 29 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbybzbAbBbCbDb,e-e.eid/ejd:ekd;eld=emd?end2b@e[e]e^e_emDnDoD%c'cpDqDhA`eEbFbGb{eHbIbJbKbLbd e h f g 4 cb5 6 dbebfbgbhbMbibjbkblbmbnbrDsDiAtDjAuDvDwDkAxDlAyD7 zD8 AD)g*g+g,gmABDCDDDEDFDGDHDIDJDKDLDMDNDODPDQDRDSDTDUDVDWDXDYDZD0D1D2D3D4D5D6D7D8D9D!D#D$DOcPc%DXe'D(D)D*DPbQbRbSbTbUbnADcoA|I}IEcpAqArAfeobodFcpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdgeheb _bzI`bAI{bBI|bCI}bDI~bEIacFIbcGIccHIdcIIecJIfc.Cgc/ChcKIicLIjcMIkcNIlcOImcPIncQIocRIpcSIqcTIrcUIscVItcWIuc+evcXIwcYIxcZIyc0IsA1Izc@ctAuAvAwAxAyATc[csEtEuEvEwExEyEzEAEBECEDEEEFEGEHEzAAABACADAEAFAGAQERESETEUEVEWEXEYEZE0E1E2E3E4E5E6E]c!E#E$E%E'E(EUcVcObWcCcNbXcYc0A1A?f4e5e^c+d|e}e~eafbfcfiejekeledfefffgfhfifjfkflfmfnfofpfmeneqfrfsfoepetfqereufvfwfxfyfRdzfAf9c!cBfUdseteCfDfueveEfFfGfHfwexeyezeIfJfKfLfMfNfOfPfQfRfSfTfUfAeBeVfWfXfCeDeYfEeFeZf0f1f2f3fVd4f5f#c$c6fYdGeHe7f8fIeJebHcH9f!f#f$f%f'f(f)f*f+f,f-f.f/f:f;fjIkI1c,C3c4c-d.d/d_c`c{c:d;d=d?d7c8c|c@d}c[d]d^d_d~c`dadbdcddd{d|d}d~daebecedeeeedfdgdlImInIoIpIqIrIsItIuIvIwIxIyIJhKh-C0C2I3I4I5I6I7I8I9I!I2H3H8H9H'H(H*HZC,H-H:H

232 elif isinstance(options, dict): 2XgYgZg0gphqh3Cod7Crh8C9Csh!CHhIhhd

233 return cls(**options) 2XgYgZg0gphqhod7Crh8C9Csh!CHhIhhd

234 else: 

235 raise TypeError( 23C

236 f"The {options_description} must be provided as an object " 23C

237 f"of type {cls.__name__} or as a dict with valid {options_description}. " 23C

238 f"The provided object is '{options}'." 23C

239 ) 

240  

241  

242def _handle_boolean_option(option: bool) -> str: 

243 """ 

244 Convert a boolean option to a string representation. 

245 """ 

246 return "true" if bool(option) else "false" 2b _b`b{b|b}b~bacbcccdcecfcgchcicjckclcmcncocpcqcrcsctcucvcwcxcyczcCc_c`c{c}ccdddedfdgd)Z

247  

248  

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

250 """ 

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

252  

253 Args: 

254 checker: The function to call to check whether the preconditions are met. It has 

255 the same signature as the wrapped function with the addition of the keyword argument `what`. 

256 what: A string that is passed in to `checker` to provide context information. 

257  

258 Returns: 

259 Callable: A decorator that creates the wrapping. 

260 """ 

261  

262 def outer(wrapped_function): 24C

263 """ 

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

265 """ 

266  

267 @functools.wraps(wrapped_function) 24C

268 def inner(*args, **kwargs): 

269 """ 

270 Check preconditions and if they are met, call the wrapped function. 

271 """ 

272 checker(*args, **kwargs, what=what) 24C

273 result = wrapped_function(*args, **kwargs) 24C

274  

275 return result 24C

276  

277 return inner 24C

278  

279 return outer 24C

280  

281  

282def is_sequence(obj): 

283 """ 

284 Check if the given object is a sequence (list or tuple). 

285 """ 

286 return isinstance(obj, Sequence) 27 8 DcEcobodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQd.C/C,C7c8c|c~cadbd*Z+Z,Z-Z.ZJhKh2Z3Z-C

287  

288  

289def is_nested_sequence(obj): 

290 """ 

291 Check if the given object is a nested sequence (list or tuple with atleast one list or tuple element). 

292 """ 

293 return is_sequence(obj) and any(is_sequence(elem) for elem in obj) 27c8c2Z3Z

294  

295  

296  

297class Transaction: 

298 """ 

299 A context manager for transactional operations with undo capability. 

300  

301 The Transaction class allows you to register undo actions (callbacks) that will be executed 

302 if the transaction is not committed before exiting the context. This is useful for managing 

303 resources or operations that need to be rolled back in case of errors or early exits. 

304  

305 Usage: 

306 with Transaction() as txn: 

307 txn.append(some_cleanup_function, arg1, arg2) 

308 # ... perform operations ... 

309 txn.commit() # Disarm undo actions; nothing will be rolled back on exit 

310  

311 Methods: 

312 append(fn, *args, **kwargs): Register an undo action to be called on rollback. 

313 commit(): Disarm all undo actions; nothing will be rolled back on exit. 

314 """ 

315 def __init__(self): 

316 self._stack = ExitStack() 2UcVcObWc

317 self._entered = False 2UcVcObWc

318  

319 def __enter__(self): 

320 self._stack.__enter__() 2UcVcObWc

321 self._entered = True 2UcVcObWc

322 return self 2UcVcObWc

323  

324 def __exit__(self, exc_type, exc, tb): 

325 # If exit callbacks remain, they'll run in LIFO order. 

326 self._entered = False 2UcVcObWc

327 return self._stack.__exit__(exc_type, exc, tb) 2UcVcObWc

328  

329 def append(self, fn, /, *args, **kwargs): 

330 """ 

331 Register an undo action (runs if the with-block exits without commit()). 

332 Values are bound now via partial so late mutations don't bite you. 

333 """ 

334 if not self._entered: 2UcVcObWc

335 raise RuntimeError("Transaction must be entered before append()") 

336 self._stack.callback(partial(fn, *args, **kwargs)) 2UcVcObWc

337  

338 def commit(self): 

339 """ 

340 Disarm all undo actions. After this, exiting the with-block does nothing. 

341 """ 

342 # pop_all() empties this stack so no callbacks are triggered on exit. 

343 self._stack.pop_all() 2UcVcObWc

344  

345  

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

347_fork_warning_checked = False 

348  

349  

350def reset_fork_warning(): 

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

352  

353 This function is intended for use in tests to allow multiple test runs 

354 to check the warning behavior. 

355 """ 

356 global _fork_warning_checked 

357 _fork_warning_checked = False 2?f4e5e^c%g+d

358  

359  

360cdef inline tuple _read_fill_ptr(const char* ptr, Py_ssize_t width): 

361 """Extract (value, element_size) from a raw pointer of known width.""" 

362 cdef unsigned int val 

363 if width == 1: 2dhehfhghhhihjhkhlhmhnhohgAd :bpb*bqb;b+brbi 4bj k l m n o p q r 5bs t u v Vbw x Wb6b=bsb,btb?b-buby 8bz A B C D E F G H 9bI J K L XbM N Yb!b@bvb.bwb[b/bxbO $bP Q R S T U V W X %bY Z 0 1 Zb2 3 0b'b

364 val = (<uint8_t*>ptr)[0] 2pbj k l s t sbz A B I J vbP Q R Y Z

365 elif width == 2: 

366 val = (<uint16_t*>ptr)[0] 2dhfhhhjhlhnh*bqbm n o u v Vb,btbC D E K L Xb.bwbS T U 0 1 Zb

367 elif width == 4: 

368 val = (<uint32_t*>ptr)[0] 2ehghihkhmhohd +brbi p q r w x Wb-buby F G H M N Yb/bxbO V W X 2 3 0b

369 else: 

370 raise ValueError(f"value must be 1, 2, or 4 bytes, got {width}") 2gA:b;b4b5b6b=b?b8b9b!b@b[b$b%b'b

371 return (val, <unsigned int>width) 2dhehfhghhhihjhkhlhmhnhohd pb*bqb+brbi j k l m n o p q r s t u v Vbw x Wbsb,btb-buby z A B C D E F G H I J K L XbM N Ybvb.bwb/bxbO P Q R S T U V W X Y Z 0 1 Zb2 3 0b

372  

373  

374cpdef tuple _parse_fill_value(value): 

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

376  

377 Parameters 

378 ---------- 

379 value : int or buffer-protocol object 

380 - int: Must be in range [0, 256). Treated as 1-byte fill. 

381 - bytes or buffer-protocol: Must be 1, 2, or 4 bytes. 

382  

383 Returns 

384 ------- 

385 tuple of (int, int) 

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

387  

388 Raises 

389 ------ 

390 OverflowError 

391 If int value is outside [0, 256). 

392 TypeError 

393 If value is not an int and does not support the buffer protocol. 

394 ValueError 

395 If value byte length is not 1, 2, or 4. 

396 """ 

397 cdef uint8_t byte_val 

398 cdef Py_buffer buf 

399  

400 if isinstance(value, int): 2GcHcIcJcKcLcMcNc@C[Cdheh_C`Cfhgh}C~ChhihcDdDjhkhgDhDlhmhkDlDnhohgAd Wee h f g Qc:bpb*bqb;b+brb3b+c,c-ci 4bj k l m n o p q r 5bs t u v Vbw x Wb6bRc=bsb,btb?b-bub7b.c/c:cy 8bz A B C D E F G H 9bI J K L XbM N Yb!bSc@bvb.bwb[b/bxb#b;c=c?cO $bP Q R S T U V W X %bY Z 0 1 Zb2 3 0b'b-E[E`EuhkFoFtFxFwh0F3F7gRd!eUdxhyh6F7F8F9F!F#F8g'Fzh}FcGBhrGvGAGEGDh7G!G9gVd)eYdEhFh%G'G(G)G*G+G!g.GGhhHnHtHzHFHLHRHXH

401 byte_val = value 2GcHcIcJcKcLcMcNc@C[C_C`C}C~CcDdDgDhDkDlDWee h f g 3b+c,c-c7b.c/c:c#b;c=c?c-E[E`EuhkFoFtFxFwh0F3F7gRd!eUdxhyh6F7F8F9F!F#F8g'Fzh}FcGBhrGvGAGEGDh7G!G9gVd)eYdEhFh%G'G(G)G*G+G!g.GGhhHnHtHzHFHLHRHXH

402 return (<unsigned int>byte_val, <unsigned int>1) 2GcHcIcJcKcLcMcNc@C[C_C`C}C~CcDdDgDhDkDlDWee h f g 3b7b#b-E[E`EuhkFoFtFxFwh0F3F7gRd!eUdxhyh6F7F8F9F!F#F8g'Fzh}FcGBhrGvGAGEGDh7G!G9gVd)eYdEhFh%G'G(G)G*G+G!g.GGhhHnHtHzHFHLHRHXH

403  

404 if isinstance(value, bytes): 2dhehfhghhhihjhkhlhmhnhohgAd Qc:bpb*bqb;b+brbi 4bj k l m n o p q r 5bs t u v Vbw x Wb6bRc=bsb,btb?b-buby 8bz A B C D E F G H 9bI J K L XbM N Yb!bSc@bvb.bwb[b/bxbO $bP Q R S T U V W X %bY Z 0 1 Zb2 3 0b'b

405 return _read_fill_ptr(<const char*><bytes>value, <Py_ssize_t>len(value)) 2dhehfhghhhihjhkhlhmhnhohgA:bpb*bqb;b+brb=bsb,btb?b-bub@bvb.bwb[b/bxb

406  

407 if PyObject_GetBuffer(value, &buf, PyBUF_SIMPLE) != 0: 2d Qci 4bj k l m n o p q r 5bs t u v Vbw x Wb6bRcy 8bz A B C D E F G H 9bI J K L XbM N Yb!bScO $bP Q R S T U V W X %bY Z 0 1 Zb2 3 0b'b

408 raise TypeError( 

409 f"value must be an int or support the buffer protocol, " 

410 f"got {type(value).__name__}" 

411 ) 

412 try: 2d i 4bj k l m n o p q r 5bs t u v Vbw x Wb6by 8bz A B C D E F G H 9bI J K L XbM N Yb!bO $bP Q R S T U V W X %bY Z 0 1 Zb2 3 0b'b

413 return _read_fill_ptr(<const char*>buf.buf, buf.len) 2d i 4bj k l m n o p q r 5bs t u v Vbw x Wb6by 8bz A B C D E F G H 9bI J K L XbM N Yb!bO $bP Q R S T U V W X %bY Z 0 1 Zb2 3 0b'b

414 finally: 

415 PyBuffer_Release(&buf) 

416  

417  

418def check_multiprocessing_start_method(): 

419 """Check if multiprocessing start method is 'fork' and warn if so.""" 

420 global _fork_warning_checked 

421 if _fork_warning_checked: 2$I%IiAjA'I(IkAlA7 8 XgYg)g*gZg0g+g,g)I*I4Z=f5Z6Z+I7Z8Z,I9Z!Z-I#Z$Z.I%Z'Z/I(Z:I;I=I?I@I[I$d%d'd(dOcPc)d)b]I^I_I`IPbQbRbSbTbUb?f4e5e^c%g+d

422 return 2$I%IiAjA'I(IkAlA7 8 XgYg)g*gZg0g+g,g)I*I4Z=f5Z6Z+I7Z8Z,I9Z!Z-I#Z$Z.I%Z'Z/I(Z:I;I=I?I@I[I$d%d'd(dOcPc)d)b]I^I_I`IPbQbRbSbTbUb+d

423 _fork_warning_checked = True 2)b?f4e5e^c%g+d

424  

425 # Common warning message parts 

426 common_message = ( 

427 "CUDA does not support. Forked subprocesses exhibit undefined behavior, " 2)b?f4e5e^c%g+d

428 "including failure to initialize CUDA contexts and devices. Set the start method " 

429 "to 'spawn' before creating processes that use CUDA. " 

430 "Use: multiprocessing.set_start_method('spawn')" 

431 ) 

432  

433 try: 2)b?f4e5e^c%g+d

434 start_method = multiprocessing.get_start_method() 2)b?f4e5e^c%g+d

435 if start_method == "fork": 2)b?f4e5e^c+d

436 message = f"multiprocessing start method is 'fork', which {common_message}" 24e5e^c+d

437 warnings.warn(message, UserWarning, stacklevel=3) 24e5e^c+d

438 except RuntimeError: 2%g

439 # get_start_method() can raise RuntimeError if start method hasn't been set 

440 # In this case, default is 'fork' on Linux, so we should warn 

441 if platform.system() == "Linux": 2%g

442 message = ( 

443 f"multiprocessing start method is not set and defaults to 'fork' on Linux, " 2%g

444 f"which {common_message}" 

445 ) 

446 warnings.warn(message, UserWarning, stacklevel=3) 2%g