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
« 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
5import functools
6from functools import partial
7import multiprocessing
8import platform
9import warnings
10from collections import namedtuple
11from collections.abc import Sequence
12from contextlib import ExitStack
13from typing import Callable
15try:
16 from cuda.bindings import driver, nvrtc, runtime
17except ImportError:
18 from cuda import cuda as driver
19 from cuda import cudart as runtime
20 from cuda import nvrtc
22from cuda.bindings.nvvm import nvvmError
23from cuda.bindings.nvjitlink import nvJitLinkError
25from cpython.buffer cimport PyObject_GetBuffer, PyBuffer_Release, Py_buffer, PyBUF_SIMPLE
27from cuda.bindings cimport cynvrtc, cynvvm, cynvjitlink
29from cuda.core._utils.driver_cu_result_explanations import DRIVER_CU_RESULT_EXPLANATIONS
30from cuda.core._utils.runtime_cuda_error_explanations import RUNTIME_CUDA_ERROR_EXPLANATIONS
33class CUDAError(Exception):
34 pass
37class NVRTCError(CUDAError):
38 pass
42ComputeCapability = namedtuple("ComputeCapability", ("major", "minor"))
45def cast_to_3_tuple(label, cfg):
46 cfg_orig = cfg 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
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
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
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
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
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
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
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
148cdef object _RUNTIME_SUCCESS = runtime.cudaError_t.cudaSuccess
149cdef object _NVRTC_SUCCESS = nvrtc.nvrtcResult.NVRTC_SUCCESS
152cpdef inline int _check_driver_error(cydriver.CUresult error) except?-1 nogil:
153 if error == cydriver.CUresult.CUDA_SUCCESS: 2a 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()}")
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
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
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}")
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
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 )
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
249def precondition(checker: Callable[..., None], str what="") -> Callable:
250 """
251 A decorator that adds checks to ensure any preconditions are met.
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.
258 Returns:
259 Callable: A decorator that creates the wrapping.
260 """
262 def outer(wrapped_function): 24C
263 """
264 A decorator that actually wraps the function for checking preconditions.
265 """
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
275 return result 24C
277 return inner 24C
279 return outer 24C
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
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
297class Transaction:
298 """
299 A context manager for transactional operations with undo capability.
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.
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
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
319 def __enter__(self):
320 self._stack.__enter__() 2UcVcObWc
321 self._entered = True 2UcVcObWc
322 return self 2UcVcObWc
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
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
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
346# Track whether we've already warned about fork method
347_fork_warning_checked = False
350def reset_fork_warning():
351 """Reset the fork warning check flag for testing purposes.
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
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
374cpdef tuple _parse_fill_value(value):
375 """Parse a fill/memset value into (raw_value, element_size).
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.
383 Returns
384 -------
385 tuple of (int, int)
386 (raw_value, element_size) where element_size is 1, 2, or 4.
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
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
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
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)
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
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 )
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