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

229 statements  

« prev     ^ index     » next       coverage.py v7.14.1, created at 2026-06-13 01:38 +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 Any, Callable 

14  

15from cuda.bindings import driver as driver, nvrtc as nvrtc, runtime as runtime 

16  

17# Module-level annotations that reference `driver`, `nvrtc`, and `runtime` so 

18# that stubgen-pyx keeps these imports in the generated `.pyi` (it would 

19# otherwise trim them as unused). These names are not assigned, so they only 

20# affect __annotations__ and have no runtime cost. 

21_keep_driver_in_stub: 'driver.CUresult' 

22_keep_nvrtc_in_stub: 'nvrtc.nvrtcResult' 

23_keep_runtime_in_stub: 'runtime.cudaError_t' 

24  

25from cuda.bindings.nvvm import nvvmError 

26from cuda.bindings.nvjitlink import nvJitLinkError 

27  

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

29  

30from cuda.bindings cimport cynvrtc, cynvvm, cynvjitlink 

31  

32from cuda.core._utils.driver_cu_result_explanations import DRIVER_CU_RESULT_EXPLANATIONS 

33from cuda.core._utils.runtime_cuda_error_explanations import RUNTIME_CUDA_ERROR_EXPLANATIONS 

34  

35  

36class CUDAError(Exception): 

37 pass 

38  

39  

40class NVRTCError(CUDAError): 

41 pass 

42  

43  

44  

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

46  

47  

48def cast_to_3_tuple(label: str, cfg: int | tuple[int, ...]) -> tuple[int, int, int]: 

49 cfg_orig = cfg 2# $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdbAbBbCbDbEbFbmfnfof`hhdpf0gidqfjdrfkdsfldtfmdXbufvfwfxfyfh i zfGbHbIbAfJbKbLbMbNbd e j f g 6 eb7 8 fbgbhbibjbObkblbmbnbobpb9 ! zcAcnend#c$cqbod#h^E_E`E{hxg=B|hBcpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdoepe}h{E~hCcBfCfaiDfEfbiFfGfqere=eseteHfIfJfKfLfMfNfOfPf1gQfRf2gSfTfueve?eUfVfWfwexeXfyezeYfZf0f1f2fRd3f4f3g%c4g5g@e[e6g7g8g9g!gSd'c#g]eTd$g%g'g5fUdAeBe6f7fCeDe^e8f9fci!f#fEeFe_eGeHe$f%f'f(f)f*f+f,f-f(g.f/f)g:f;fIeJe`e=f?f@fKeLe[fMeNe]f^f_f`f{fVd|f}f*g(c+g,g{e|e-g.g/g:g;gWd)c=g}eXd?g@g[g~fYdOePeagbgQeRe~ecgdgdiegfgeigghgfiigjggikglghimgngiiogpgjiqgrgki$hVEWE?B|E%h

50 if isinstance(cfg, int): 2# $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdbAbBbCbDbEbFbmfnfof`hhdpf0gidqfjdrfkdsfldtfmdXbufvfwfxfyfh i zfGbHbIbAfJbKbLbMbNbd e j f g 6 eb7 8 fbgbhbibjbObkblbmbnbobpb9 ! zcAcnend#c$cqbod#h^E_E`E{hxg=B|hBcpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdoepe}h{E~hCcBfCfaiDfEfbiFfGfqere=eseteHfIfJfKfLfMfNfOfPf1gQfRf2gSfTfueve?eUfVfWfwexeXfyezeYfZf0f1f2fRd3f4f3g%c4g5g@e[e6g7g8g9g!gSd'c#g]eTd$g%g'g5fUdAeBe6f7fCeDe^e8f9fci!f#fEeFe_eGeHe$f%f'f(f)f*f+f,f-f(g.f/f)g:f;fIeJe`e=f?f@fKeLe[fMeNe]f^f_f`f{fVd|f}f*g(c+g,g{e|e-g.g/g:g;gWd)c=g}eXd?g@g[g~fYdOePeagbgQeRe~ecgdgdiegfgeigghgfiigjggikglghimgngiiogpgjiqgrgki$hVEWE?B|E%h

51 cfg = (cfg,) 2# $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdbAbBbCbDbEbFbmfnfofpfqfrfsftfXbufvfwfxfyfh i zfGbHbIbAfJbKbLbMbNbd e j f g 6 eb7 8 fbgbhbibjbObkblbmbnbobpb9 ! zcAcnend#c$cqbod#h^E_E`Exg=BpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdoepe{ECcBfCfDfEfFfGfqereseteHfIfJfKfLfMfNfOfPfQfRfSfTfueveUfVfWfwexeXfyezeYfZf0f1f2fRd3f4f%c'c5fUdAeBe6f7fCeDe8f9f!f#fEeFeGeHe$f%f'f(f)f*f+f,f-f.f/f:f;fIeJe=f?f@fKeLe[fMeNe]f^f_f`f{fVd|f}f(c)c~fYdOePeagbgQeRecgdgegfggghgigjgkglgmgngogpgqgrg$hVE

52 else: 

53 common = "must be an int, or a tuple with up to 3 ints" 2`hhd0gidjdkdldmd#h{hxg|hBc}h~haibi=e1g2g?e3g%c4g5g@e[e6g7g8g9g!gSd'c#g]eTd$g%g'g^eci_e(g)g`e*g(c+g,g{e|e-g.g/g:g;gWd)c=g}eXd?g@g[g~edieifigihiiijiki$hWE?B|E%h

54 if not isinstance(cfg, tuple): 2`hhd0gidjdkdldmd#h{hxg|hBc}h~haibi=e1g2g?e3g%c4g5g@e[e6g7g8g9g!gSd'c#g]eTd$g%g'g^eci_e(g)g`e*g(c+g,g{e|e-g.g/g:g;gWd)c=g}eXd?g@g[g~edieifigihiiijiki$hWE?B|E%h

55 raise ValueError(f"{label} {common} (got {type(cfg)})") 2|E

56 if len(cfg) > 3: 2`hhd0gidjdkdldmd#h{hxg|hBc}h~haibi=e1g2g?e3g%c4g5g@e[e6g7g8g9g!gSd'c#g]eTd$g%g'g^eci_e(g)g`e*g(c+g,g{e|e-g.g/g:g;gWd)c=g}eXd?g@g[g~edieifigihiiijiki$hWE?B%h

57 raise ValueError(f"{label} {common} (got tuple with length {len(cfg)})") 2WE

58 if any(not isinstance(val, int) for val in cfg): 2`hhd0gidjdkdldmd#h{hxg|hBc}h~haibi=e1g2g?e3g%c4g5g@e[e6g7g8g9g!gSd'c#g]eTd$g%g'g^eci_e(g)g`e*g(c+g,g{e|e-g.g/g:g;gWd)c=g}eXd?g@g[g~edieifigihiiijiki$h?B%h

59 raise ValueError(f"{label} {common} (got {cfg})") 2?B

60 if any(val < 1 for val in cfg): 2# $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdbAbBbCbDbEbFbmfnfof`hhdpf0gidqfjdrfkdsfldtfmdXbufvfwfxfyfh i zfGbHbIbAfJbKbLbMbNbd e j f g 6 eb7 8 fbgbhbibjbObkblbmbnbobpb9 ! zcAcnend#c$cqbod#h^E_E`E{hxg=B|hBcpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdoepe}h{E~hCcBfCfaiDfEfbiFfGfqere=eseteHfIfJfKfLfMfNfOfPf1gQfRf2gSfTfueve?eUfVfWfwexeXfyezeYfZf0f1f2fRd3f4f3g%c4g5g@e[e6g7g8g9g!gSd'c#g]eTd$g%g'g5fUdAeBe6f7fCeDe^e8f9fci!f#fEeFe_eGeHe$f%f'f(f)f*f+f,f-f(g.f/f)g:f;fIeJe`e=f?f@fKeLe[fMeNe]f^f_f`f{fVd|f}f*g(c+g,g{e|e-g.g/g:g;gWd)c=g}eXd?g@g[g~fYdOePeagbgQeRe~ecgdgdiegfgeigghgfiigjggikglghimgngiiogpgjiqgrgki$hVE%h

61 plural_s = "" if len(cfg) == 1 else "s" 2xgVE%h

62 raise ValueError(f"{label} value{plural_s} must be >= 1 (got {cfg_orig})") 2xgVE%h

63 return cfg + (1,) * (3 - len(cfg)) 2# $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdbAbBbCbDbEbFbmfnfof`hhdpf0gidqfjdrfkdsfldtfmdXbufvfwfxfyfh i zfGbHbIbAfJbKbLbMbNbd e j f g 6 eb7 8 fbgbhbibjbObkblbmbnbobpb9 ! zcAcnend#c$cqbod#h^E_E`E{hxg=B|hBcpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdoepe}h{E~hCcBfCfaiDfEfbiFfGfqere=eseteHfIfJfKfLfMfNfOfPf1gQfRf2gSfTfueve?eUfVfWfwexeXfyezeYfZf0f1f2fRd3f4f3g%c4g5g@e[e6g7g8g9g!gSd'c#g]eTd$g%g'g5fUdAeBe6f7fCeDe^e8f9fci!f#fEeFe_eGeHe$f%f'f(f)f*f+f,f-f(g.f/f)g:f;fIeJe`e=f?f@fKeLe[fMeNe]f^f_f`f{fVd|f}f*g(c+g,g{e|e-g.g/g:g;gWd)c=g}eXd?g@g[g~fYdOePeagbgQeRe~ecgdgdiegfgeigghgfiigjggikglghimgngiiogpgjiqgrgki$h

64  

65  

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

67 if err != cydriver.CUresult.CUDA_SUCCESS: 2a 4LZd5L0d6L7L# 8L$ 9L!L% #L' $L( %L) 'L* (L+ )L, *L- +L. ,L/ -L: .L; /L= :L? ;L@ =L[ ?L] @L^ [L_ ]L` ^L{ _L| `L} {L~ |Lab}Lbb~LcbaMdbbMcMdMeMfMgMhMiMjMkM]glMmMnMoMpMqMrMsMtMuMvMwMxMyMzMAMBMCMDMEMFMGMHMIMJMKMLMSeMMNMOMPMTeQMUeRMVeSMWeTMAbUMXeVMYeWMZeXMYMZM0M1MBb2MCb3MDb4MEb5MFb6MDc7MEc8MFc9MGc!MHc#MIc$MJc%MKc'M(M)M*M+M,Myg-Mzg.MAg/M:M;M=M?M@M[Mmf]Mnf^M_M`M{Mof|M}E~E}M~MaNbNcNhddNaFbFlimieNfNgNhNiNjNkNpflNmNcFnNdFoNpNqNrNsNtNuNvNwN0gidxNyNzNeFANfFBNniCNoiDNENFNGNHNINJNKNLNMNNNONPNQNRNSNTNUNVNWNXNYNZN0N1N2Nqf3NgFhF4N5N6N7N8Njd9NiFjFpiqi!N#N$N%N'N(N)Nrf*N+N,NkF-NlF.N/N:N;N=N?N@N[N]N^N_Nkd`N{N|NmF}NnF~NriaOsibOcOdOeOfOgOhOiOjOsfkOlOmOoFnOpFoOpOqOrOsOtOuOvOwOxOyOldzOAOBOqFCOrFDOtiEOuiFOGOHOIOJOKOLOMONOtfOOPOQOsFROtFSOTOUOVOWOXOYOZO0O1O2Omd3O4O5OuF6OvF7Ovi8Owi9O!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~OaPbPcPdPePfPgPhPiPjPkPlPmPnP^goPpPqPrPsPtPuPvPwPxPyPzPAPBPCP_gDPEPFPGPHP@BIPXbJPKPLP`gMP0eNPOPPPQPRPSPTPUPVPWPXPYPZP0P1P2P3P4P5P6P7P8P9Puf!Pvf#Pwf$Pxf%P1e'P2e(Pyf)Paf*PwF+PxF,P*c-P+c.PyF/PzF:Ph ;Pi =P[B?Pzf@PGb[PHb]PIb^PAf_P3e`P{g{PJb|PKb}PLb~PMbaQNbbQcQdQeQfQgQhQiQjQd kQ|glQ4emQ}gnQ,coQ-cpQ.cqQ1drQ2dsQ3dtQe uQj vQf wQg xQ6 yQebzQ7 AQ8 BQfbCQgbDQhbEQ4dFQ5dGQ6dHQ7dIQ8dJQ9dKQLQMQNQibOQjbPQObQQkbRQlbSQmbTQnbUQobVQpbWQAFBFbfCFXQ~gahbhchdhehYQfhZQDFBg0QEFcfFFGF5e1QdfefffHFIFgfJFKFCg2Q3QghhhihhfLFifMFNFOFPFQFRFSFTFUFVFWF6e4QaLXFbLYF]BZF^B0FcL1FdL2F_B3F`B4F9 5F! 6Fjh5Qkh6Q'h7Q(h8Qlh9Qmh!Q)h#Q*h$Qxi%Q'Q(QeL7FfL8F9F!F#FgL$F)Q%F'F(FhL)F*Q*F+F,FiL-F+Q.F/F:FjL;F,Q=F?F@FkL[F-Q]F^F_FlL`F.Q{F|FmL}Fsg~FnLaGoLbGpLcGqLdG!deG/cfG/QYbgG:chG+b:QLc;Q;ciG=cjGtgkGrLlGsLmGtLnGQb=QRb?QSb@QTb[QUb]QVb^QDg_Q`Q{Q|Q}Q~QaRbREg7e8ecRdReRfRgRhRiRjRkRlRFgmRnRoRpRqRrRsRtRuRvRwRxRyRzRARBRCRDRERFRGRHRIRJRKRLRMRNRORPRQRRRSRTRURVRWRXRYRZR0R1R2R3R4R5R6R7R8R9R!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*Snh+S,S-SohGgHgph{Bzc|B.SoG/SpG:SqG;SrG=SsG?StG@SuG[SAc]SvG^SwG_SxG`SyG{S}B|S~B}SaC~SneaTzGAGBGCGDGbTcTdTeTfTgTEGFGGGHGhTIGJGiTKGLGjTkTlTmTXEnToTYEqhMGZErhndyish#cbC$ccCpT0Eth1Euh2EvhqT3EwhrT4ExhNG5EyhsTtTuTvTwTxT6EyT7EzT8EAT9EBT!ECT#EDT$EETFTGTHTITJTKTLTMTqb,b-bodNTOT=BPTBcQTpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdoepe.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b~bacbcccdcecfcgchcicjckclcmcncocdCpcRTSTTTIgqcrcJgUTKgLgMgOGPGVTNgOgPgQgzhAhWTXTrbYTZTsb0T1T2Ttb3TZb4T5T6T7Tk 8T9Tl !Tm #Tn $To %Tp 'Tq (Tr )Ts *Tt +T,Tu -Tv .Tw /Tx :T;Ty =Tz ?T@T[T]T^Tub_T`Tvb{T|T}Twb~T0baUbUcUdUA eUfUB gUC hUD iUE jUF kUG lUH mUI nUJ oUpUK qUL rUM sUN tUuUO vUP wUxUyUzUAUxbBUCUybDUEUFUzbGU1bHUIUJUKUQ LUMUR NUS OUT PUU QUV RUW SUX TUY UUZ VUWU0 XU1 YU2 ZU3 0U1U4 2U5 3U4U5URgQGRGSG6U?c7UTG8UUG9U!U#U$U9eeCfC%UVG'UWG(UgC)UhC*UiC+UjC,UMc-U!e.U@c/U#e:U;UXG=UYG?UZG@U0G[U1G]U2G^U3G_U4G`U5G{U6G|U7G}U8G~U9GaV!GbV#GcV$GdVkCeV%GfVlCgV'GhVmCiV(GjVnCkV)GlVoCmV*GnVpCoV+GpVqCqV,GrVrCsV-GtV.GuV/GvV:GwV;GxV=GyV?GzV@GAV[GBV]GCV^GDV_GEV`GFV{GGV|GHV}GIV~GJV%EKVziLV'EMV(ENVAiOV)EPVaHQVRVSgTg$ebHcHdHSV[cTVeHfHgHUVhHVViHWVjHXVsCYVtCZVNcOcPbPc0VQc1V2Vsc3VuC4VvC#d5VwCxC6VyC7VzC8VACCc9VBCWbCC!VDC#VEC$VFC%VGC'VHC(VIC)VJC*VKC+VLC,VMCNCOCPCRcScTcQCRCSCUcVcWc-VTC.VUC/VVC:VWC;VXC=VYC?VZC@V0C[V1C]V2C^V3C_V4C`V5C{V6C|V7C}V8C~V9CaW!CbW#CcW$CdW%CeW'CfW(CgW)C*C+ChW,CiW-CjW.CkW/ClW:CmW;CnW=CoW?CpW@CqW[CrW]CsW^CtW_CuW`CvW{CwW|CxW}CyW~CzWaDAWbDBWcDdDCWeDDWfDEWgDFWhDGWiDHWjDIWkDJWlDKWmDLWnDMWoDNWpDOWqDPWrDQWsDRWtDuDugSW%eTW'eUW]cVW$dWWXWkHBfYWZW0W1WlHmHnH2W3W4W5W6WvDCf7WoHwD8WpH9W!WqHxDrH#WsHDf$W%W'W(WtHuHvH)W*W+W,W-WyDEf.WwHzD/WxH:W;WFfyHzH=W?W@W[WGf]WAH^W_WBHqeCHDH*E+EEHFHGHHHIHBire=eJHCiDiKHLHMHseteHfIfJfKfLf`WMf{WNHOH|W}W~WaXNfbXPHcXdXeXfXQHOfgXhXRHSHiXjXkXlXmXADPf1gnXTHBDoXpXqXrXUHQfsXtXuXVHWHXHvXwXxXyXzXCDRf2gAXYHDDBXZHCXDXEXSf0H1HFXGXHXIXTfJX2HKXLX3H,Eue4H5H6H7H8H9H!H#H$H%HEDve?e'HEiFD(H)H*H+H,H-H.H/H:H;H=H?H@H[HUf]H^HMXNXOXPX_HQXRXSXVfTXUX`H{HVXWXXXYXWfZX|H0X1X2X3X4X5X6X7X8X9X!X#XGDweHDIDJDKDLDMDNDODPDxeQDBhRDSDTDXfyezeYfZf0f1f2fRd3f4f3g%c4g5g@e[e6g7g8g9g!gSd'c#g]eTd$g%g'g5f}H~H$X%X'X(X)XUdFiGiaIbIcIdIeIfIUDAeVDWDXDYDZD0D1D2D3D4DBe5DCh6D7D8D*X6f+XgIhI,X-X.X/X7f:XiI;X=XjI-ECekIlImInIoIpIqIrIsItIuI9DDe^evIHi!DwIxIyI?X@XzI8f[X]X^X_XAIBICI`X{X|X}X~X#D9faYDI$DbYEIcYdY!fFIGIeYfYgYhY#fiYHIjYkYIIEeJIKI.E/ELIMINIOIPIIiFe_eQIJiKiRISITIGeHe$f%f'f(f)flY*fmYUIVInYoYpYqY+frYWIsYtYuYvYXI,fwYxYYIZIyYzYAYBYCY%D-f(gDY0I'DEYFYGYHY1I.fIYJYKY2I3I4ILYMYNYOYPY(D/f)gQY5I)DRY6ISYTYUY:f7I8IVYWYXYYY;fZY9I0Y1Y!I:EIe#I$I%I'I(I)I*I+I,I-I*DJe`e.ILi+D/I:I;I=I?I@I[I]I^I_I`I{I|I}I=f~IaJ2Y3Y4Y5YbJ6Y7Y8Y?f9Y!YcJdJ#Y$Y%Y'Y@f(YeJ)Y*Y+Y,Y-Y.Y/Y:Y;Y=Y?Y@Y,DKe-D.D/D:D;D=D?D@D[DLe]DDh^D_D`D[fMeNe]f^f_f`f{fVd|f}f*g(c+g,g{e|e-g.g/g:g;gWd)c=g}eXd?g@g[g~ffJgJ[Y]Y^Y_Y`YYdMiNihJiJjJkJlJmJ{DOe|D}D~DaEbEcEdEeEfEgEPehEEhiEjEkE{Yag|YnJoJ}Y~YaZbZbgcZpJdZeZqJ;EQerJsJtJuJvJwJxJyJzJAJBJlERe~eCJOimEDJEJFJfZgZhZGJ(ePiQinEoEiZHJcgjZkZlZmZIJJJKJnZoZpZqZrZpEdgsZLJqErEvgsEtEwgtZMJuZvZwZNJegxZyZzZAZOJPJQJBZCZDZEZFZuEfgGZRJvEHZSJIZJZKZTJggLZMZNZOZUJVJWJPZQZRZSZTZwEhgUZXJxEVZYJWZXZYZZJigZZ0Z1Z2Z0J1J2J3Z4Z5Z6Z7ZyEjg8Z3JzE9Z4J!Z#Z$Z5Jkg%Z'Z(Z)Z6J7J8J*Z+Z,Z-Z.ZAElg/Z9JBE:Z!J;Z=Z?Z#Jmg@Z[Z]Z^Z$J%J'J_Z`Z{Z|Z}ZCEng~Z(JDEa0)Jb0c0d0*Joge0f0g0h0+J,J-Ji0j0k0l0m0EEpgn0.JFEGEo0/Jp0q0r0:Jqgs0t0u0v0;J=J?Jw0x0y0z0A0HErgB0@JIEJEC0[JD0E0XcF0YcG0=EH0ZcI0J0K0L0M0N0O0P0Q0R0S0T0U0V0W0X0Y0Z000102030405060708090!0#0$0%0'0(0)0*0+0,0c -00c.0KE2bLE%d/0:0;0=0?0@0[0]0^0_0`0{0|0}0RiSi]JMENEOE~0PE^J_JQE?E1c+ha1b1c1d1e1f1g1h1i1j1k12cl1m1n1o1p1q1r1s1t1u1v1w1x1`Jy1z1A1B1C1D1E1F1G1H1REI1J1{JK1|JL1M1}JN1O1P1Q1R1S1T1U1V1W1X1Y1Z101112131415161718191!1#1$1%1'1(1)1*1+1,1-1.1/1:1;1=1?1@1[1]1^1_1`1{1|1}1~1a2b2c2d2e2f2g2h2i2j2k2l2m2n2o2p2q2r2s2t2u2v2w2x2~JaKy2z2bKcKdKeKfKA2gKhKiKB2jKC2kKD2lKmKE2nKF2oKG2pKH2qKI2J2rKsKtKuKvKK2wKL2xKM2N2O2yKP2Q2R2S2T2U2V2W2X2Y2Z202122232425262728292!2#2$2%2'2(2)2*2+2,2-2.2/2:2;2=2?2@2

68 return _check_driver_error(err) 2]g^g_g`g|g4e}gWb%d

69 return 0 2a 4LZd5L0d6L7L# 8L$ 9L!L% #L' $L( %L) 'L* (L+ )L, *L- +L. ,L/ -L: .L; /L= :L? ;L@ =L[ ?L] @L^ [L_ ]L` ^L{ _L| `L} {L~ |Lab}Lbb~LcbaMdbbMcMdMeMfMgMhMiMjMkM]glMmMnMoMpMqMrMsMtMuMvMwMxMyMzMAMBMCMDMEMFMGMHMIMJMKMLMSeMMNMOMPMTeQMUeRMVeSMWeTMAbUMXeVMYeWMZeXMYMZM0M1MBb2MCb3MDb4MEb5MFb6MDc7MEc8MFc9MGc!MHc#MIc$MJc%MKc'M(M)M*M+M,Myg-Mzg.MAg/M:M;M=M?M@M[Mmf]Mnf^M_M`M{Mof|M}E~E}M~MaNbNcNhddNaFbFlimieNfNgNhNiNjNkNpflNmNcFnNdFoNpNqNrNsNtNuNvNwN0gidxNyNzNeFANfFBNniCNoiDNENFNGNHNINJNKNLNMNNNONPNQNRNSNTNUNVNWNXNYNZN0N1N2Nqf3NgFhF4N5N6N7N8Njd9NiFjFpiqi!N#N$N%N'N(N)Nrf*N+N,NkF-NlF.N/N:N;N=N?N@N[N]N^N_Nkd`N{N|NmF}NnF~NriaOsibOcOdOeOfOgOhOiOjOsfkOlOmOoFnOpFoOpOqOrOsOtOuOvOwOxOyOldzOAOBOqFCOrFDOtiEOuiFOGOHOIOJOKOLOMONOtfOOPOQOsFROtFSOTOUOVOWOXOYOZO0O1O2Omd3O4O5OuF6OvF7Ovi8Owi9O!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~OaPbPcPdPePfPgPhPiPjPkPlPmPnP^goPpPqPrPsPtPuPvPwPxPyPzPAPBPCP_gDPEPFPGPHP@BIPXbJPKPLP`gMP0eNPOPPPQPRPSPTPUPVPWPXPYPZP0P1P2P3P4P5P6P7P8P9Puf!Pvf#Pwf$Pxf%P1e'P2e(Pyf)Paf*PwF+PxF,P*c-P+c.PyF/PzF:Ph ;Pi =P[B?Pzf@PGb[PHb]PIb^PAf_P3e`P{g{PJb|PKb}PLb~PMbaQNbbQcQdQeQfQgQhQiQjQd kQ|glQ4emQ}gnQ,coQ-cpQ.cqQ1drQ2dsQ3dtQe uQj vQf wQg xQ6 yQebzQ7 AQ8 BQfbCQgbDQhbEQ4dFQ5dGQ6dHQ7dIQ8dJQ9dKQLQMQNQibOQjbPQObQQkbRQlbSQmbTQnbUQobVQpbWQAFBFbfCFXQ~gahbhchdhehYQfhZQDFBg0QEFcfFFGF5e1QdfefffHFIFgfJFKFCg2Q3QghhhihhfLFifMFNFOFPFQFRFSFTFUFVFWF6e4QaLXFbLYF]BZF^B0FcL1FdL2F_B3F`B4F9 5F! 6Fjh5Qkh6Q'h7Q(h8Qlh9Qmh!Q)h#Q*h$Qxi%Q'Q(QeL7FfL8F9F!F#FgL$F)Q%F'F(FhL)F*Q*F+F,FiL-F+Q.F/F:FjL;F,Q=F?F@FkL[F-Q]F^F_FlL`F.Q{F|FmL}Fsg~FnLaGoLbGpLcGqLdG!deG/cfG/QYbgG:chG+b:QLc;Q;ciG=cjGtgkGrLlGsLmGtLnGQb=QRb?QSb@QTb[QUb]QVb^QDg_Q`Q{Q|Q}Q~QaRbREg7e8ecRdReRfRgRhRiRjRkRlRFgmRnRoRpRqRrRsRtRuRvRwRxRyRzRARBRCRDRERFRGRHRIRJRKRLRMRNRORPRQRRRSRTRURVRWRXRYRZR0R1R2R3R4R5R6R7R8R9R!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*Snh+S,S-SohGgHgph{Bzc|B.SoG/SpG:SqG;SrG=SsG?StG@SuG[SAc]SvG^SwG_SxG`SyG{S}B|S~B}SaC~SneaTzGAGBGCGDGbTcTdTeTfTgTEGFGGGHGhTIGJGiTKGLGjTkTlTmTXEnToTYEqhMGZErhndyish#cbC$ccCpT0Eth1Euh2EvhqT3EwhrT4ExhNG5EyhsTtTuTvTwTxT6EyT7EzT8EAT9EBT!ECT#EDT$EETFTGTHTITJTKTLTMTqb,b-bodNTOT=BPTBcQTpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdoepe.b/b:b;b=b?b@b[b]b^b_b`b{b|b}b~bacbcccdcecfcgchcicjckclcmcncocdCpcRTSTTTIgqcrcJgUTKgLgMgOGPGVTNgOgPgQgzhAhWTXTrbYTZTsb0T1T2Ttb3TZb4T5T6T7Tk 8T9Tl !Tm #Tn $To %Tp 'Tq (Tr )Ts *Tt +T,Tu -Tv .Tw /Tx :T;Ty =Tz ?T@T[T]T^Tub_T`Tvb{T|T}Twb~T0baUbUcUdUA eUfUB gUC hUD iUE jUF kUG lUH mUI nUJ oUpUK qUL rUM sUN tUuUO vUP wUxUyUzUAUxbBUCUybDUEUFUzbGU1bHUIUJUKUQ LUMUR NUS OUT PUU QUV RUW SUX TUY UUZ VUWU0 XU1 YU2 ZU3 0U1U4 2U5 3U4U5URgQGRGSG6U?c7UTG8UUG9U!U#U$U9eeCfC%UVG'UWG(UgC)UhC*UiC+UjC,UMc-U!e.U@c/U#e:U;UXG=UYG?UZG@U0G[U1G]U2G^U3G_U4G`U5G{U6G|U7G}U8G~U9GaV!GbV#GcV$GdVkCeV%GfVlCgV'GhVmCiV(GjVnCkV)GlVoCmV*GnVpCoV+GpVqCqV,GrVrCsV-GtV.GuV/GvV:GwV;GxV=GyV?GzV@GAV[GBV]GCV^GDV_GEV`GFV{GGV|GHV}GIV~GJV%EKVziLV'EMV(ENVAiOV)EPVaHQVRVSgTg$ebHcHdHSV[cTVeHfHgHUVhHVViHWVjHXVsCYVtCZVNcOcPbPc0VQc1V2Vsc3VuC4VvC#d5VwCxC6VyC7VzC8VACCc9VBCWbCC!VDC#VEC$VFC%VGC'VHC(VIC)VJC*VKC+VLC,VMCNCOCPCRcScTcQCRCSCUcVcWc-VTC.VUC/VVC:VWC;VXC=VYC?VZC@V0C[V1C]V2C^V3C_V4C`V5C{V6C|V7C}V8C~V9CaW!CbW#CcW$CdW%CeW'CfW(CgW)C*C+ChW,CiW-CjW.CkW/ClW:CmW;CnW=CoW?CpW@CqW[CrW]CsW^CtW_CuW`CvW{CwW|CxW}CyW~CzWaDAWbDBWcDdDCWeDDWfDEWgDFWhDGWiDHWjDIWkDJWlDKWmDLWnDMWoDNWpDOWqDPWrDQWsDRWtDuDugSW%eTW'eUW]cVW$dWWXWkHBfYWZW0W1WlHmHnH2W3W4W5W6WvDCf7WoHwD8WpH9W!WqHxDrH#WsHDf$W%W'W(WtHuHvH)W*W+W,W-WyDEf.WwHzD/WxH:W;WFfyHzH=W?W@W[WGf]WAH^W_WBHqeCHDH*E+EEHFHGHHHIHBire=eJHCiDiKHLHMHseteHfIfJfKfLf`WMf{WNHOH|W}W~WaXNfbXPHcXdXeXfXQHOfgXhXRHSHiXjXkXlXmXADPf1gnXTHBDoXpXqXrXUHQfsXtXuXVHWHXHvXwXxXyXzXCDRf2gAXYHDDBXZHCXDXEXSf0H1HFXGXHXIXTfJX2HKXLX3H,Eue4H5H6H7H8H9H!H#H$H%HEDve?e'HEiFD(H)H*H+H,H-H.H/H:H;H=H?H@H[HUf]H^HMXNXOXPX_HQXRXSXVfTXUX`H{HVXWXXXYXWfZX|H0X1X2X3X4X5X6X7X8X9X!X#XGDweHDIDJDKDLDMDNDODPDxeQDBhRDSDTDXfyezeYfZf0f1f2fRd3f4f3g%c4g5g@e[e6g7g8g9g!gSd'c#g]eTd$g%g'g5f}H~H$X%X'X(X)XUdFiGiaIbIcIdIeIfIUDAeVDWDXDYDZD0D1D2D3D4DBe5DCh6D7D8D*X6f+XgIhI,X-X.X/X7f:XiI;X=XjI-ECekIlImInIoIpIqIrIsItIuI9DDe^evIHi!DwIxIyI?X@XzI8f[X]X^X_XAIBICI`X{X|X}X~X#D9faYDI$DbYEIcYdY!fFIGIeYfYgYhY#fiYHIjYkYIIEeJIKI.E/ELIMINIOIPIIiFe_eQIJiKiRISITIGeHe$f%f'f(f)flY*fmYUIVInYoYpYqY+frYWIsYtYuYvYXI,fwYxYYIZIyYzYAYBYCY%D-f(gDY0I'DEYFYGYHY1I.fIYJYKY2I3I4ILYMYNYOYPY(D/f)gQY5I)DRY6ISYTYUY:f7I8IVYWYXYYY;fZY9I0Y1Y!I:EIe#I$I%I'I(I)I*I+I,I-I*DJe`e.ILi+D/I:I;I=I?I@I[I]I^I_I`I{I|I}I=f~IaJ2Y3Y4Y5YbJ6Y7Y8Y?f9Y!YcJdJ#Y$Y%Y'Y@f(YeJ)Y*Y+Y,Y-Y.Y/Y:Y;Y=Y?Y@Y,DKe-D.D/D:D;D=D?D@D[DLe]DDh^D_D`D[fMeNe]f^f_f`f{fVd|f}f*g(c+g,g{e|e-g.g/g:g;gWd)c=g}eXd?g@g[g~ffJgJ[Y]Y^Y_Y`YYdMiNihJiJjJkJlJmJ{DOe|D}D~DaEbEcEdEeEfEgEPehEEhiEjEkE{Yag|YnJoJ}Y~YaZbZbgcZpJdZeZqJ;EQerJsJtJuJvJwJxJyJzJAJBJlERe~eCJOimEDJEJFJfZgZhZGJ(ePiQinEoEiZHJcgjZkZlZmZIJJJKJnZoZpZqZrZpEdgsZLJqErEvgsEtEwgtZMJuZvZwZNJegxZyZzZAZOJPJQJBZCZDZEZFZuEfgGZRJvEHZSJIZJZKZTJggLZMZNZOZUJVJWJPZQZRZSZTZwEhgUZXJxEVZYJWZXZYZZJigZZ0Z1Z2Z0J1J2J3Z4Z5Z6Z7ZyEjg8Z3JzE9Z4J!Z#Z$Z5Jkg%Z'Z(Z)Z6J7J8J*Z+Z,Z-Z.ZAElg/Z9JBE:Z!J;Z=Z?Z#Jmg@Z[Z]Z^Z$J%J'J_Z`Z{Z|Z}ZCEng~Z(JDEa0)Jb0c0d0*Joge0f0g0h0+J,J-Ji0j0k0l0m0EEpgn0.JFEGEo0/Jp0q0r0:Jqgs0t0u0v0;J=J?Jw0x0y0z0A0HErgB0@JIEJEC0[JD0E0XcF0YcG0=EH0ZcI0J0K0L0M0N0O0P0Q0R0S0T0U0V0W0X0Y0Z000102030405060708090!0#0$0%0'0(0)0*0+0,0c -00c.0KE2bLE%d/0:0;0=0?0@0[0]0^0_0`0{0|0}0RiSi]JMENEOE~0PE^J_JQE?E1c+ha1b1c1d1e1f1g1h1i1j1k12cl1m1n1o1p1q1r1s1t1u1v1w1x1`Jy1z1A1B1C1D1E1F1G1H1REI1J1{JK1|JL1M1}JN1O1P1Q1R1S1T1U1V1W1X1Y1Z101112131415161718191!1#1$1%1'1(1)1*1+1,1-1.1/1:1;1=1?1@1[1]1^1_1`1{1|1}1~1a2b2c2d2e2f2g2h2i2j2k2l2m2n2o2p2q2r2s2t2u2v2w2x2~JaKy2z2bKcKdKeKfKA2gKhKiKB2jKC2kKD2lKmKE2nKF2oKG2pKH2qKI2J2rKsKtKuKvKK2wKL2xKM2N2O2yKP2Q2R2S2T2U2V2W2X2Y2Z202122232425262728292!2#2$2%2'2(2)2*2+2,2-2.2/2:2;2=2?2@2

70  

71  

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

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

74 if err == cynvrtc.nvrtcResult.NVRTC_SUCCESS: 2# $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdbAbBbCbDbEbFbmfnfofhdpfidqfjdrfkdsfldtfmdXbufvfwfxfyfh i [BzfGbHbIbAfJbKbLbMbNbd e j f g 6 eb7 8 fbgbhbibjbObkblbmbnbobpb9 ! zcAcneyibCcCqbodBcpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdoepeb .b/b:b;b=b?b@b[b]b^b_b`b{b|b}b~bacbcccdcecfcgchcicjckclcmcncocdCpcscuCvC#dwCxCyCzCACCcBCWbCCDCECFCGCHCICJCKCLCMCNCOCPCRcScTcQCRCSCUcVcWcTCUCVCWCXCYCZC0C1C2C3C4C5C6C7C8C9C!C#C$C%C'C(C)C*C+C,C-C.C/C:C;C=C?C@C[C]C^C_C`C{C|C}C~CaDbDcDdDeDfDgDhDiDjDkDlDmDnDoDpDqDrDsDtDuDBfvDCfwDxDDfyDEfzDFfGfqeBireDiseteHfIfJfKfLfMfNfOfADPfBDQfCDRfDDSfTfueEDveFDUfVfWfGDweHDIDJDKDLDMDNDODPDxeQDBhRDSDTDXfyezeYfZf0f1f2fRd3f4f%cSd'cTd5fUdUDAeVDWDXDYDZD0D1D2D3D4DBe5DCh6D7D8D6f7fCe9DDe!D8f#D9f$D!f#fEeIiFeKiGeHe$f%f'f(f)f*f+f,f%D-f'D.f(D/f)D:f;fIe*DJe+D=f?f@f,DKe-D.D/D:D;D=D?D@D[DLe]DDh^D_D`D[fMeNe]f^f_f`f{fVd|f}f(cWd)cXd~fYd{DOe|D}D~DaEbEcEdEeEfEgEPehEEhiEjEkEagbgQelERemEnEoEcgpEdgqErEsEeguEfgvEggwEhgxEigyEjgzEkgAElgBEmgCEngDEogEEpgFEGEqgHErgIEJEXcYc'dZc(d)d*d^c_c`c+d,d-d.d3c4c{c/d|c:d;d=d?d}c@d~cadbdcd[d]d^d_d`d{d|d}d~dddedfdc 0cMENEOEgdPEQE1c2caebecedeeefegezKAKBKCKheDKiejeEK

75 return 0 2# $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdbAbBbCbDbEbFbmfnfofhdpfidqfjdrfkdsfldtfmdXbufvfwfxfyfh i [BzfGbHbIbAfJbKbLbMbNbd e j f g 6 eb7 8 fbgbhbibjbObkblbmbnbobpb9 ! zcAcneyibCcCqbodBcpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdoepeb .b/b:b;b=b?b@b[b]b^b_b`b{b|b}b~bacbcccdcecfcgchcicjckclcmcncocdCpcscuCvC#dwCxCyCzCACCcBCWbCCDCECFCGCHCICJCKCLCMCNCOCPCRcScTcQCRCSCUcVcWcTCUCVCWCXCYCZC0C1C2C3C4C5C6C7C8C9C!C#C$C%C'C(C)C*C+C,C-C.C/C:C;C=C?C@C[C]C^C_C`C{C|C}C~CaDbDcDdDeDfDgDhDiDjDkDlDmDnDoDpDqDrDsDtDuDBfvDCfwDxDDfyDEfzDFfGfqeBireDiseteHfIfJfKfLfMfNfOfADPfBDQfCDRfDDSfTfueEDveFDUfVfWfGDweHDIDJDKDLDMDNDODPDxeQDBhRDSDTDXfyezeYfZf0f1f2fRd3f4f%cSd'cTd5fUdUDAeVDWDXDYDZD0D1D2D3D4DBe5DCh6D7D8D6f7fCe9DDe!D8f#D9f$D!f#fEeIiFeKiGeHe$f%f'f(f)f*f+f,f%D-f'D.f(D/f)D:f;fIe*DJe+D=f?f@f,DKe-D.D/D:D;D=D?D@D[DLe]DDh^D_D`D[fMeNe]f^f_f`f{fVd|f}f(cWd)cXd~fYd{DOe|D}D~DaEbEcEdEeEfEgEPehEEhiEjEkEagbgQelERemEnEoEcgpEdgqErEsEeguEfgvEggwEhgxEigyEjgzEkgAElgBEmgCEngDEogEEpgFEGEqgHErgIEJEXcYc'dZc(d)d*d^c_c`c+d,d-d.d3c4c{c/d|c:d;d=d?d}c@d~cadbdcd[d]d^d_d`d{d|d}d~dddedfdc 0cMENEOEgdPEQE1c2caebecedeeefegezKAKBKCKheDKiejeEK

76 with gil: 1c

77 _raise_nvrtc_error(prog, err) 1c

78  

79  

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

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

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

83 cdef size_t logsize = 0 1c

84 if prog != NULL: 1c

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

86 cdef bytes log_bytes 

87 cdef str log_str = "" 1c

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

89 log_bytes = b" " * logsize 1c

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

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

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

93 if log_str: 1c

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

95 raise NVRTCError(err_msg) 1c

96  

97  

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

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

100 if err == cynvvm.nvvmResult.NVVM_SUCCESS: 2tEFKGKKE2bLE%dHKIKJKKKLKMKNKOKPKQKRKSKTKUKRiSi

101 return 0 2tEFKGKKE2bLE%dHKIKJKKKLKMKNKOKPKQKRKSKTKUKRiSi

102 with gil: 22b

103 _raise_nvvm_error(prog, err) 22b

104  

105  

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

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

108 cdef size_t logsize = 0 22b

109 if prog != NULL: 22b

110 cynvvm.nvvmGetProgramLogSize(prog, &logsize) 22b

111 cdef bytes log_bytes 

112 cdef str log_str = "" 22b

113 if logsize > 1 and prog != NULL: 22b

114 log_bytes = b" " * logsize 22b

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

116 log_str = log_bytes.decode("utf-8", errors="backslashreplace") 22b

117 cdef object exc = nvvmError(err) 22b

118 if log_str: 22b

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

120 raise exc 22b

121  

122  

123cdef int HANDLE_RETURN_NVJITLINK( 

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

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

126 if err == cynvjitlink.nvJitLinkResult.NVJITLINK_SUCCESS: 2b VKWKXKYKZK0K1K2K3K4K5K@E[E6K7K8K9K!K#K$K%K'K(K)K*K+Kjf,K-K.K/K:K+h;K=K?K@K[K]K^K_K`K{K

127 return 0 2b VKWKXKYKZK0K1K2K3K4K5K@E[E6K7K8K9K!K#K$K%K'K(K)K*K+K,K-K.K/K:K+h;K=K?K@K[K]K^K_K`K{K

128 with gil: 2b jf

129 _raise_nvjitlink_error(handle, err) 2b jf

130  

131  

132cdef int _raise_nvjitlink_error( 

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

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

135 cdef size_t logsize = 0 2b jf

136 if handle != NULL: 2b jf

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

138 cdef bytes log_bytes 

139 cdef str log_str = "" 2b jf

140 if logsize > 1 and handle != NULL: 2b jf

141 log_bytes = b" " * logsize 1b

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

143 cynvjitlink.nvJitLinkResult.NVJITLINK_SUCCESS: 1b

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

145 cdef object exc = nvJitLinkError(err) 2b jf

146 if log_str: 2b jf

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

148 raise exc 2b jf

149  

150  

151cdef object _RUNTIME_SUCCESS = runtime.cudaError_t.cudaSuccess 

152cdef object _NVRTC_SUCCESS = nvrtc.nvrtcResult.NVRTC_SUCCESS 

153  

154  

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

156 if error == cydriver.CUresult.CUDA_SUCCESS: 2a ZdTi0dUikfVi# Wi$ XilfYi% Zi' 0i( 1i) 2i* 3i+ 4i, 5i- 6i. 7i/ 8i: 9i; !i= #i? $i@ %i[ 'i] (i^ )i_ *i` +i{ ,i| -i} .i~ /iab:ibb;icb=idb?i@i[i]i^i]g_i`i{i|i}i~iajbjcjdjejfjgjhjijSejjkjljTemjUenjVeojWepjAbqjXerjYesjZetjujvjBbwjCbxjDbyjEbzjFbAjDcBjEcCjFcDjGcEjHcFjIcGjJcHjKcIjJj,hKj-hLj.hMjygNjzgOjAgPjQjRjSjTjUjVjWjXjYjZj0j1j2j3j4j5j6j7j8j9j!j#j$j%j'j(j)j*j+j,j-j.j/j:j;j=j?j@j[j]j^j_j`j{j|j}j~jakbkckdkekfkgkhkikjkkklkmknkokpkqkrksktkukvkwkxkykzkAkBkCkDkEkFkGkHkIkJkKkLkMkNkOkPkQkRkSkTkUkVkWkXkYkZk0k1k2k3k4k5k6k7k8k9k!k#k$k%k'k(k)k*k+k,k-k.k/k:k;k=k?k@k[k]k^k_k`k{k|k}k~kalblcldlelflglhliljlklllmlnlol^gplqlrlsltlulvlwl_gxlylzlAlXbBlCl`gDl0eElFlGlHlIlJlKlLlMlNlOlPlQlRlSlTl1eUl2eVlWlafXlYlZl*c0l+c1l2l3lh 4li 5l6l7lGb8lHb9lIb!l#l3e$l{g%lJb'lKb(lLb)lMb*lNb+l,l-l.l/ld :l|g;l4e=l}g?l,c-c.c1d2d3de j f g 6 eb7 8 fbgbhb4d5d6d7d8d9d@lib[ljb]lOb^lkb_llb`lmb{lnb|lob}lpb~lambfbmcm~gah/hbhch:hdheh;hfh=hdmBgemcf?hfm5egmdfef@hffhmgf[himCgjmgh]hhhih^hhfkmiflmmmnmompmqm6erm9 ! sg!d/cYb:c+bLc;c=ctgQbRbSbTbUbVbsmtmumvmwmxmymSE_hzmAmBmCmDmEmFmGmHmImJmKmLmMmnhNmOmPmohQmphzcRmSmTmUmVmWmXmYmAcZm0m1m2m3m4m5m6m7m8mqhrhndsh#c9m$c!mthuhvh#mwhxhyh$m%m'm(m)m*m+m,m-m.m/m:m;m=m?m@m[m]m^mqb,b-b_m`m{m|mpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdoepe}m~manbncndnenfngnhninjnknlnmnnnonpnqnrnsntnunvnwnxnynznAnBnCnDnEnFnGnHnIgqcrcJgKgLgMgNgOgPgQgzhAhke5crbtcsb6cuctbZb)e*e+ek 9bl m n o p q r s t !bu v w x 3by z 4b#ble7cubvcvb8cwcwb0b,e-e.eA $bB C D E F G H I J %bK L M N 5bO P 6b'bme9cxbxcyb!cyczb1b/e:e;eQ (bR S T U V W X Y Z )b0 1 2 3 7b4 5 8b*bRgInJnKn9eLnMnNnOnPnQnRnSnTnUnVnWnXnYnZn0n1n2n3n4n5n6n7n8n9n!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~naobocodoeofogohoiojokolomoSgTg$enooopoqorososCtotCuoNcOcPbPcvoQcwoxoyo#dzoAoBoCoDoEoCcFoGoWbHoIoJoKoLoMoNoOoPoQoRoSoToUoVoWoXoYoZo0o1o2o3o4o5o6o7o8o9o!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(e|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%d$z%z'z(z)z*z+z,z-z.z/z:z;z=z?z@z[z]z^z_z`z{z|z}z~zaAbAcAdAeAfAgAhAiAjAkAlAmAnAoApAqArAsAtAuAvAwAxAyAzAAABACADAEAFAGAHAIAJAKALAMANAOAPAQARASATAUAVAWAXAYAZA0A1A2A3A4A5A6A7A8A9A!A#A$A%A'A(A)A*A+A,A-A.A/A:A;A=A?A@A[A]A^A_A`A{A|A}A~AaBbBcBdBeBfBgBhBiBjBkBlBmBnBoBpBqBrBsBtBuBvBwBxByBzBABBBCBDBEBFBGBHBIBJBKBLBMBNBOBPBQBRBSBTBUBVBWBXBYBZB0B1B2B3B4B5B6B7B8B9B!B#B$B%B'B(B)B*B+B,B-B.B/B

157 return 0 2a ZdTi0dUikfVi# Wi$ XilfYi% Zi' 0i( 1i) 2i* 3i+ 4i, 5i- 6i. 7i/ 8i: 9i; !i= #i? $i@ %i[ 'i] (i^ )i_ *i` +i{ ,i| -i} .i~ /iab:ibb;icb=idb?i@i[i]i^i_i`i{i|i}i~iajbjcjdjejfjgjhjijSejjkjljTemjUenjVeojWepjAbqjXerjYesjZetjujvjBbwjCbxjDbyjEbzjFbAjDcBjEcCjFcDjGcEjHcFjIcGjJcHjKcIjJj,hKj-hLj.hMjygNjzgOjAgPjQjRjSjTjUjVjWjXjYjZj0j1j2j3j4j5j6j7j8j9j!j#j$j%j'j(j)j*j+j,j-j.j/j:j;j=j?j@j[j]j^j_j`j{j|j}j~jakbkckdkekfkgkhkikjkkklkmknkokpkqkrksktkukvkwkxkykzkAkBkCkDkEkFkGkHkIkJkKkLkMkNkOkPkQkRkSkTkUkVkWkXkYkZk0k1k2k3k4k5k6k7k8k9k!k#k$k%k'k(k)k*k+k,k-k.k/k:k;k=k?k@k[k]k^k_k`k{k|k}k~kalblcldlelflglhliljlklllmlnlolplqlrlsltlulvlwlxlylzlAlXbBlClDl0eElFlGlHlIlJlKlLlMlNlOlPlQlRlSlTl1eUl2eVlWlafXlYlZl*c0l+c1l2l3lh 4li 5l6l7lGb8lHb9lIb!l#l3e$l{g%lJb'lKb(lLb)lMb*lNb+l,l-l.l/ld :l;l4e=l?l,c-c.c1d2d3de j f g 6 eb7 8 fbgbhb4d5d6d7d8d9d@lib[ljb]lOb^lkb_llb`lmb{lnb|lob}lpb~lambfbmcm~gah/hbhch:hdheh;hfh=hdmBgemcf?hfm5egmdfef@hffhmgf[himCgjmgh]hhhih^hhfkmiflmmmnmompmqm6erm9 ! sg!d/cYb:c+bLc;c=ctgQbRbSbTbUbVbsmtmumvmwmxmym_hzmAmBmCmDmEmFmGmHmImJmKmLmMmnhNmOmPmohQmphzcRmSmTmUmVmWmXmYmAcZm0m1m2m3m4m5m6m7m8mqhrhndsh#c9m$c!mthuhvh#mwhxhyh$m%m'm(m)m*m+m,m-m.m/m:m;m=m?m@m[m]m^mqb,b-b_m`m{m|mpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdoepe}m~manbncndnenfngnhninjnknlnmnnnonpnqnrnsntnunvnwnxnynznAnBnCnDnEnFnGnHnIgqcrcJgKgLgMgNgOgPgQgzhAhke5crbtcsb6cuctbZb)e*e+ek 9bl m n o p q r s t !bu v w x 3by z 4b#ble7cubvcvb8cwcwb0b,e-e.eA $bB C D E F G H I J %bK L M N 5bO P 6b'bme9cxbxcyb!cyczb1b/e:e;eQ (bR S T U V W X Y Z )b0 1 2 3 7b4 5 8b*bRgInJnKn9eLnMnNnOnPnQnRnSnTnUnVnWnXnYnZn0n1n2n3n4n5n6n7n8n9n!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~naobocodoeofogohoiojokolomoSgTg$enooopoqorososCtotCuoNcOcPbPcvoQcwoxoyo#dzoAoBoCoDoEoCcFoGoWbHoIoJoKoLoMoNoOoPoQoRoSoToUoVoWoXoYoZo0o1o2o3o4o5o6o7o8o9o!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(e|w}w~waxbxcxdxexfxgxhxixjxkxlxmxnxoxpxqxrxsxtxuxvxwxxxyxzxAxBxCxDxExFxGxHxIxJxKxLxMxNxOxPxQxRxSxTxUxVxWxXxYxZx0x1x2x3x4x5x6x7x8x9x!x#x$x%x'x(x)x*x+x,x-x.x/x:x;x=x?x@x[x]x^x_x`x{x|x}x~xaybycydyeyfygyhyiyjykylymynyoypyqyrysytyuyvywyxyyyzyAyByCyDyEyFyGyHyIyJyKyLyMyNyOyPyQyRySyTyUyVyWyXyYyZy0y1y2y3y4y5y6y7y8y9y!y#y$y%y'y(y)y*y+y,y-y.y/y:y;y=y?y@y[y]y^y_y`y{y|y}y~yazbzczdzezfzgzhzizjzkzlzmznzozpzqzrzsztzuzvzwzxzyzzzAzBzCzDzEzFzGzHzIzJzKzLzMzNzOzPzQzRzSzTzUzVzWzXzYzZz0z1z2z3z4z5z6z7z8z9z!z#z$z%z'z(z)z*z+z,z-z.z/z:z;z=z?z@z[z]z^z_z`z{z|z}z~zaAbAcAdAeAfAgAhAiAjAkAlAmAnAoApAqArAsAtAuAvAwAxAyAzAAABACADAEAFAGAHAIAJAKALAMANAOAPAQARASATAUAVAWAXAYAZA0A1A2A3A4A5A6A7A8A9A!A#A$A%A'A(A)A*A+A,A-A.A/A:A;A=A?A@A[A]A^A_A`A{A|A}A~AaBbBcBdBeBfBgBhBiBjBkBlBmBnBoBpBqBrBsBtBuBvBwBxByBzBABBBCBDBEBFBGBHBIBJBKBLBMBNBOBPBQBRBSBTBUBVBWBXBYBZB0B1B2B3B4B5B6B7B8B9B!B#B$B%B'B(B)B*B+B,B-B.B/B

158 cdef const char* name 

159 name_err = cydriver.cuGetErrorName(error, &name) 2]g^g_g`g|g4e}gSE_hPbWb%d

160 if name_err != cydriver.CUresult.CUDA_SUCCESS: 2]g^g_g`g|g4e}gSE_hPbWb%d

161 raise CUDAError(f"UNEXPECTED ERROR CODE: {error}") 2_h

162 with gil: 2]g^g_g`g|g4e}gSE_hPbWb%d

163 # TODO: consider lower this to Cython 

164 expl = DRIVER_CU_RESULT_EXPLANATIONS.get(int(error)) 2]g^g_g`g|g4e}gSE_hPbWb%d

165 if expl is not None: 2]g^g_g`g|g4e}gSE_hPbWb%d

166 raise CUDAError(f"{name.decode()}: {expl}") 2]g^g_g`g|g4e}gSE_hPbWb%d

167 cdef const char* desc 

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

169 if desc_err != cydriver.CUresult.CUDA_SUCCESS: 

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

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

172  

173  

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

175 if error == _RUNTIME_SUCCESS: 2QbRbSbTbUbVbDg|KFhEg7e8eFgGgHg,b-bqcrcrbsbtbZbk l m n o p q r s t u v w x y z ubvbwb0bA B C D E F G H I J K L M N O P xbybzb1bQ R S T U V W X Y Z 0 1 2 3 4 5 ?cMc!e@c#e[c

176 return 0 2QbRbSbTbUbVbDgFhEg7e8eFgGgHg,b-bqcrcrbsbtbZbk l m n o p q r s t u v w x y z ubvbwb0bA B C D E F G H I J K L M N O P xbybzb1bQ R S T U V W X Y Z 0 1 2 3 4 5 ?cMc!e@c#e[c

177 # `_check_error()` reaches this path only for `runtime.cudaError_t` values. 

178 # Use the enum name directly because Windows hybrid cudart can lag that table. 

179 name = error.name 2|KFh

180 expl = RUNTIME_CUDA_ERROR_EXPLANATIONS.get(int(error)) 2|KFh

181 if expl is not None: 2|KFh

182 raise CUDAError(f"{name}: {expl}") 2|KFh

183 desc_err, desc = runtime.cudaGetErrorString(error) 2Fh

184 if desc_err != _RUNTIME_SUCCESS: 2Fh

185 raise CUDAError(f"{name}") 

186 desc = desc.decode() 2Fh

187 raise CUDAError(f"{name}: {desc}") 2Fh

188  

189  

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

191 if error == _NVRTC_SUCCESS: 2a :B}Kb .b/b:b;b=b?b@b[b]b^b_b`b{b|b}b~bacbcccdcecfcgchcicjckclcmcncocpcscRcScTcGhUcHhVcIhWcvgwgXcYc'dZc(d)d*d^c_c`c+d,d-d.d3c4c{c/d|c:d;d=d?d}c@d~cadbdcd[d]d^d_d`d{d|d}d~dddedfdc 0cgd1c2cUgVgWgXgYgZgJhKhLhMhNhOhPhQhRhShThUhVhWhXhYhZh0h1h2h3h4h5h6h7h8h9haebecedeeefegeheieje

192 return 0 2a }Kb .b/b:b;b=b?b@b[b]b^b_b`b{b|b}b~bacbcccdcecfcgchcicjckclcmcncocpcscRcScTcGhUcHhVcIhWcvgwgXcYc'dZc(d)d*d^c_c`c+d,d-d.d3c4c{c/d|c:d;d=d?d}c@d~cadbdcd[d]d^d_d`d{d|d}d~dddedfdc 0cgd1c2cUgVgWgXgYgZgJhKhLhMhNhOhPhQhRhShThUhVhWhXhYhZh0h1h2h3h4h5h6h7h8h9haebecedeeefegeheieje

193 err = f"{error}: {nvrtc.nvrtcGetErrorString(error)[1].decode()}" 2:B}K

194 if handle is not None: 2:B}K

195 _, logsize = nvrtc.nvrtcGetProgramLogSize(handle) 2:B

196 log = b" " * logsize 2:B

197 _ = nvrtc.nvrtcGetProgramLog(handle, log) 2:B

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

199 raise NVRTCError(err) 2:B}K

200  

201  

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

203 if isinstance(error, driver.CUresult): 2a ZdTi0dUikfVi# Wi$ XilfYi% Zi' 0i( 1i) 2i* 3i+ 4i, 5i- 6i. 7i/ 8i: 9i; !i= #i? $i@ %i[ 'i] (i^ )i_ *i` +i{ ,i| -i} .i~ /iab:ibb;icb=idb?i@i[i]i^i_i`i{i|i}i~iajbjcjdjejfjgjhjijSejjkjljTemjUenjVeojWepjAbqjXerjYesjZetjujvjBbwjCbxjDbyjEbzjFbAjDcBjEcCjFcDjGcEjHcFjIcGjJcHjKcIjJj,hKj-hLj.hMjygNjzgOjAgPjQjRjSjTjUjVjWjXjYjZj0j1j2j3j4j5j6j7j8j9j!j#j$j%j'j(j)j*j+j,j-j.j/j:j;j=j?j@j[j]j^j_j`j{j|j}j~jakbkckdkekfkgkhkikjkkklkmknkokpkqkrksktkukvkwkxkykzkAkBkCkDkEkFkGkHkIkJkKkLkMkNkOkPkQkRkSkTkUkVkWkXkYkZk0k1k2k3k4k5k6k7k8k9k!k#k$k%k'k(k)k*k+k,k-k.k/k:k;k=k?k@k[k]k^k_k`k{k|k}k~kalblcldlelflglhliljlklllmlnlolplqlrlsltlulvlwlxlylzlAlXbBlClDl0eElFlGlHlIlJlKlLlMlNlOlPlQlRlSlTl1eUl2eVlWlafXlYlZl*c0l+c1l2l3lh 4li 5l6l7lGb8lHb9lIb!l#l3e$l{g%lJb'lKb(lLb)lMb*lNb+l,l-l.l/ld :l;l=l?l,c-c.c1d2d3de j f g 6 eb7 8 fbgbhb4d5d6d7d8d9d@lib[ljb]lOb^lkb_llb`lmb{lnb|lob}lpb~lambfbmcm~gah/hbhch:hdheh;hfh=hdmBgemcf?hfm5egmdfef@hffhmgf[himCgjmgh]hhhih^hhfkmiflmmmnmompmqm6erm9 ! sg!d/cYb:c+bLc;c=ctgQbRbSbTbUbVbDgsmtmumvmwmxmymzmEg7eAm8eBmCmDmEmFmGmHmImFgJmKmLmMmnhNmOmPmohGgHgQmphRmSmTmUmVmWmXmYmZm0m1m2m3m4m5m6m7m8mqhrhsh9m!mthuhvh#mwhxhyh$m%m'm(m)m*m+m,m-m.m/m:m;m=m?m@m[m]m^mqb,b-b_m`m{m|mb .b}m/b~m:ban;bbn=bcn?bdn@ben[bfn]bgn^bhn_bin`bjn{bkn|bln}bmn~bnnaconbcpnccqndcrnecsnfctngcunhcvnicwnjcxnkcynlcznmcAnncBnocCnDnpcEnFnGnHnIgqcrcJgKgLgMgNgOgPgQgzhAhke5crbtcsb6cuctbZb)e*e+ek 9bl m n o p q r s t !bu v w x 3by z 4b#ble7cubvcvb8cwcwb0b,e-e.eA $bB C D E F G H I J %bK L M N 5bO P 6b'bme9cxbxcyb!cyczb1b/e:e;eQ (bR S T U V W X Y Z )b0 1 2 3 7b4 5 8b*bRg?cInJnKn9eLnMnNnOnPnQnRnMcSn!eTn@cUn#eVnWnXnYnZn0n1n2n3n4n5n6n7n8n9n!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~naobocodoeofogohoiojokolomoSgTg$eno[coopoqorosotouovoscwoxoyo#dzoAoBoCoDoEoFoGoWbHoIoJoKoLoMoNoOoPoQoRoSoToUoRcVoScWoTcXoYoZo0oGhUc1oHhVc2oIhWc3o4o5o6o7o8o9o!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(e|w}w~waxbxcxdxexfxgxhxixjxkxlxmxnxoxpxqxrxsxtxvguxvxwxwgxxyxzxAxBxCxDxExFxGxHxIxJxKxLxMxNxOxPxQxRxSxTxUxVxWxXxYxZx0x1x2x3x4x5x6x7x8x9x!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~yazbzczdzezfzgzhzizjzkzlzmznzozpzqzrzsztzuzXcvzYcwz'dxzZcyz(dzz)dAz*dBz^cCz_cDz`cEz+dFz,dGz-dHz.dIz3cJz4cKz{cLz/dMz|cNz:dOz;dPz=dQz?dRz}cSz@dTz~cUzadVzbdWzcdXz[dYz]dZz^d0z_d1z`d2z{d3z|d4z}d5z~d6zdd7zed8zfd9zc !z0c#z$z%z'z(z)z*z+z,z-z.z/z:z;z=zgd?z1c@z[z]z^z_z`z{z|z}z~zaA2cbAUgVgWgXgYgZgJhKhLhMhNhOhPhQhRhShThUhVhWhcAXhYhZh0h1h2h3h4h5h6hdAeAfAgA7h8h9haebecedeeehAfegeheiejeiAjAkAlAmAnAoApAqArAsAtAuAvAwAxAyAzAAABACADAEAFAGAHAIAJAKALAMANAOAPAQARASATAUAVAWAXAYAZA0A1A2A3A4A5A6A7A8A9A!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~AaBbBcBdBeBfBgBhBiBjBkBlBmBnBoBpBqBrBsBtBuBvBwBxByBzBABBBCBDBEBFBGBHBIBJBKBLBMBNBOBPBQBRBSBTBUBVBWBXBYBZB0B1B2B3B4B5B6B7B8B9B!B#B$B%B'B(B)B*B+B,B-B.B/B

204 return _check_driver_error(error) 2a ZdTi0dUikfVi# Wi$ XilfYi% Zi' 0i( 1i) 2i* 3i+ 4i, 5i- 6i. 7i/ 8i: 9i; !i= #i? $i@ %i[ 'i] (i^ )i_ *i` +i{ ,i| -i} .i~ /iab:ibb;icb=idb?i@i[i]i^i_i`i{i|i}i~iajbjcjdjejfjgjhjijSejjkjljTemjUenjVeojWepjAbqjXerjYesjZetjujvjBbwjCbxjDbyjEbzjFbAjDcBjEcCjFcDjGcEjHcFjIcGjJcHjKcIjJj,hKj-hLj.hMjygNjzgOjAgPjQjRjSjTjUjVjWjXjYjZj0j1j2j3j4j5j6j7j8j9j!j#j$j%j'j(j)j*j+j,j-j.j/j:j;j=j?j@j[j]j^j_j`j{j|j}j~jakbkckdkekfkgkhkikjkkklkmknkokpkqkrksktkukvkwkxkykzkAkBkCkDkEkFkGkHkIkJkKkLkMkNkOkPkQkRkSkTkUkVkWkXkYkZk0k1k2k3k4k5k6k7k8k9k!k#k$k%k'k(k)k*k+k,k-k.k/k:k;k=k?k@k[k]k^k_k`k{k|k}k~kalblcldlelflglhliljlklllmlnlolplqlrlsltlulvlwlxlylzlAlXbBlClDl0eElFlGlHlIlJlKlLlMlNlOlPlQlRlSlTl1eUl2eVlWlafXlYlZl*c0l+c1l2l3lh 4li 5l6l7lGb8lHb9lIb!l#l3e$l{g%lJb'lKb(lLb)lMb*lNb+l,l-l.l/ld :l;l=l?l,c-c.c1d2d3de j f g 6 eb7 8 fbgbhb4d5d6d7d8d9d@lib[ljb]lOb^lkb_llb`lmb{lnb|lob}lpb~lambfbmcm~gah/hbhch:hdheh;hfh=hdmBgemcf?hfm5egmdfef@hffhmgf[himCgjmgh]hhhih^hhfkmiflmmmnmompmqm6erm9 ! sg!d/cYb:c+bLc;c=ctgQbRbSbTbUbVbsmtmumvmwmxmymzmAmBmCmDmEmFmGmHmImJmKmLmMmnhNmOmPmohQmphRmSmTmUmVmWmXmYmZm0m1m2m3m4m5m6m7m8mqhrhsh9m!mthuhvh#mwhxhyh$m%m'm(m)m*m+m,m-m.m/m:m;m=m?m@m[m]m^mqb,b-b_m`m{m|m}m~manbncndnenfngnhninjnknlnmnnnonpnqnrnsntnunvnwnxnynznAnBnCnDnEnFnGnHnIgqcrcJgKgLgMgNgOgPgQgzhAhke5crbtcsb6cuctbZb)e*e+ek 9bl m n o p q r s t !bu v w x 3by z 4b#ble7cubvcvb8cwcwb0b,e-e.eA $bB C D E F G H I J %bK L M N 5bO P 6b'bme9cxbxcyb!cyczb1b/e:e;eQ (bR S T U V W X Y Z )b0 1 2 3 7b4 5 8b*bRgInJnKn9eLnMnNnOnPnQnRnSnTnUnVnWnXnYnZn0n1n2n3n4n5n6n7n8n9n!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~naobocodoeofogohoiojokolomoSgTg$enooopoqorosotouovowoxoyo#dzoAoBoCoDoEoFoGoWbHoIoJoKoLoMoNoOoPoQoRoSoToUoVoWoXoYoZo0o1o2o3o4o5o6o7o8o9o!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(e|w}w~waxbxcxdxexfxgxhxixjxkxlxmxnxoxpxqxrxsxtxuxvxwxxxyxzxAxBxCxDxExFxGxHxIxJxKxLxMxNxOxPxQxRxSxTxUxVxWxXxYxZx0x1x2x3x4x5x6x7x8x9x!x#x$x%x'x(x)x*x+x,x-x.x/x:x;x=x?x@x[x]x^x_x`x{x|x}x~xaybycydyeyfygyhyiyjykylymynyoypyqyrysytyuyvywyxyyyzyAyByCyDyEyFyGyHyIyJyKyLyMyNyOyPyQyRySyTyUyVyWyXyYyZy0y1y2y3y4y5y6y7y8y9y!y#y$y%y'y(y)y*y+y,y-y.y/y:y;y=y?y@y[y]y^y_y`y{y|y}y~yazbzczdzezfzgzhzizjzkzlzmznzozpzqzrzsztzuzvzwzxzyzzzAzBzCzDzEzFzGzHzIzJzKzLzMzNzOzPzQzRzSzTzUzVzWzXzYzZz0z1z2z3z4z5z6z7z8z9z!z#z$z%z'z(z)z*z+z,z-z.z/z:z;z=z?z@z[z]z^z_z`z{z|z}z~zaAbAcAdAeAfAgAhAiAjAkAlAmAnAoApAqArAsAtAuAvAwAxAyAzAAABACADAEAFAGAHAIAJAKALAMANAOAPAQARASATAUAVAWAXAYAZA0A1A2A3A4A5A6A7A8A9A!A#A$A%A'A(A)A*A+A,A-A.A/A:A;A=A?A@A[A]A^A_A`A{A|A}A~AaBbBcBdBeBfBgBhBiBjBkBlBmBnBoBpBqBrBsBtBuBvBwBxByBzBABBBCBDBEBFBGBHBIBJBKBLBMBNBOBPBQBRBSBTBUBVBWBXBYBZB0B1B2B3B4B5B6B7B8B9B!B#B$B%B'B(B)B*B+B,B-B.B/B

205 elif isinstance(error, runtime.cudaError_t): 2a QbRbSbTbUbVbDgEg7e8eFgGgHg,b-bb .b/b:b;b=b?b@b[b]b^b_b`b{b|b}b~bacbcccdcecfcgchcicjckclcmcncocpcqcrcrbsbtbZbk l m n o p q r s t u v w x y z ubvbwb0bA B C D E F G H I J K L M N O P xbybzb1bQ R S T U V W X Y Z 0 1 2 3 4 5 ?cMc!e@c#e[cscRcScTcGhUcHhVcIhWcvgwgXcYc'dZc(d)d*d^c_c`c+d,d-d.d3c4c{c/d|c:d;d=d?d}c@d~cadbdcd[d]d^d_d`d{d|d}d~dddedfdc 0cgd1c2cUgVgWgXgYgZgJhKhLhMhNhOhPhQhRhShThUhVhWhXhYhZh0h1h2h3h4h5h6h7h8h9haebecedeeefegeheieje

206 return _check_runtime_error(error) 2QbRbSbTbUbVbDgEg7e8eFgGgHg,b-bqcrcrbsbtbZbk l m n o p q r s t u v w x y z ubvbwb0bA B C D E F G H I J K L M N O P xbybzb1bQ R S T U V W X Y Z 0 1 2 3 4 5 ?cMc!e@c#e[c

207 elif isinstance(error, nvrtc.nvrtcResult): 2a b .b/b:b;b=b?b@b[b]b^b_b`b{b|b}b~bacbcccdcecfcgchcicjckclcmcncocpcscRcScTcGhUcHhVcIhWcvgwgXcYc'dZc(d)d*d^c_c`c+d,d-d.d3c4c{c/d|c:d;d=d?d}c@d~cadbdcd[d]d^d_d`d{d|d}d~dddedfdc 0cgd1c2cUgVgWgXgYgZgJhKhLhMhNhOhPhQhRhShThUhVhWhXhYhZh0h1h2h3h4h5h6h7h8h9haebecedeeefegeheieje

208 return _check_nvrtc_error(error, handle=handle) 2a b .b/b:b;b=b?b@b[b]b^b_b`b{b|b}b~bacbcccdcecfcgchcicjckclcmcncocpcscRcScTcGhUcHhVcIhWcvgwgXcYc'dZc(d)d*d^c_c`c+d,d-d.d3c4c{c/d|c:d;d=d?d}c@d~cadbdcd[d]d^d_d`d{d|d}d~dddedfdc 0cgd1c2cUgVgWgXgYgZgJhKhLhMhNhOhPhQhRhShThUhVhWhXhYhZh0h1h2h3h4h5h6h7h8h9haebecedeeefegeheieje

209 else: 

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

211  

212  

213def handle_return(result: tuple[Any, ...], handle: object = None) -> Any: 

214 _check_error(result[0], handle=handle) 2a ZdTi0dUikfVi# Wi$ XilfYi% Zi' 0i( 1i) 2i* 3i+ 4i, 5i- 6i. 7i/ 8i: 9i; !i= #i? $i@ %i[ 'i] (i^ )i_ *i` +i{ ,i| -i} .i~ /iab:ibb;icb=idb?i@i[i]i^i_i`i{i|i}i~iajbjcjdjejfjgjhjijSejjkjljTemjUenjVeojWepjAbqjXerjYesjZetjujvjBbwjCbxjDbyjEbzjFbAjDcBjEcCjFcDjGcEjHcFjIcGjJcHjKcIjJj,hKj-hLj.hMjygNjzgOjAgPjQjRjSjTjUjVjWjXjYjZj0j1j2j3j4j5j6j7j8j9j!j#j$j%j'j(j)j*j+j,j-j.j/j:j;j=j?j@j[j]j^j_j`j{j|j}j~jakbkckdkekfkgkhkikjkkklkmknkokpkqkrksktkukvkwkxkykzkAkBkCkDkEkFkGkHkIkJkKkLkMkNkOkPkQkRkSkTkUkVkWkXkYkZk0k1k2k3k4k5k6k7k8k9k!k#k$k%k'k(k)k*k+k,k-k.k/k:k;k=k?k@k[k]k^k_k`k{k|k}k~kalblcldlelflglhliljlklllmlnlolplqlrlsltlulvlwlxlylzlAlXbBlClDl0eElFlGlHlIlJlKlLlMlNlOlPlQlRlSlTl1eUl2eVlWlafXlYlZl*c0l+c1l2l3lh 4li 5l6l7lGb8lHb9lIb!l#l3e$l{g%lJb'lKb(lLb)lMb*lNb+l,l-l.l/ld :l;l=l?l,c-c.c1d2d3de j f g 6 eb7 8 fbgbhb4d5d6d7d8d9d@lib[ljb]lOb^lkb_llb`lmb{lnb|lob}lpb~lambfbmcm~gah/hbhch:hdheh;hfh=hdmBgemcf?hfm5egmdfef@hffhmgf[himCgjmgh]hhhih^hhfkmiflmmmnmompmqm6erm9 ! sg!d/cYb:c+bLc;c=ctgQbRbSbTbUbVbDgsmtmumvmwmxmymzmEg7eAm8eBmCmDmEmFmGmHmImFgJmKmLmMmnhNmOmPmohGgHgQmphRmSmTmUmVmWmXmYmZm0m1m2m3m4m5m6m7m8mqhrhsh9m!mthuhvh#mwhxhyh$m%m'm(m)m*m+m,m-m.m/m:m;m=m?m@m[m]m^mqb,b-b_m`m{m|mb .b}m/b~m:ban;bbn=bcn?bdn@ben[bfn]bgn^bhn_bin`bjn{bkn|bln}bmn~bnnaconbcpnccqndcrnecsnfctngcunhcvnicwnjcxnkcynlcznmcAnncBnocCnDnpcEnFnGnHnIgqcrcJgKgLgMgNgOgPgQgzhAhke5crbtcsb6cuctbZb)e*e+ek 9bl m n o p q r s t !bu v w x 3by z 4b#ble7cubvcvb8cwcwb0b,e-e.eA $bB C D E F G H I J %bK L M N 5bO P 6b'bme9cxbxcyb!cyczb1b/e:e;eQ (bR S T U V W X Y Z )b0 1 2 3 7b4 5 8b*bRg?cInJnKn9eLnMnNnOnPnQnRnMcSn!eTn@cUn#eVnWnXnYnZn0n1n2n3n4n5n6n7n8n9n!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~naobocodoeofogohoiojokolomoSgTg$eno[coopoqorosotouovoscwoxoyo#dzoAoBoCoDoEoFoGoWbHoIoJoKoLoMoNoOoPoQoRoSoToUoRcVoScWoTcXoYoZo0oGhUc1oHhVc2oIhWc3o4o5o6o7o8o9o!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(e|w}w~waxbxcxdxexfxgxhxixjxkxlxmxnxoxpxqxrxsxtxvguxvxwxwgxxyxzxAxBxCxDxExFxGxHxIxJxKxLxMxNxOxPxQxRxSxTxUxVxWxXxYxZx0x1x2x3x4x5x6x7x8x9x!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~yazbzczdzezfzgzhzizjzkzlzmznzozpzqzrzsztzuzXcvzYcwz'dxzZcyz(dzz)dAz*dBz^cCz_cDz`cEz+dFz,dGz-dHz.dIz3cJz4cKz{cLz/dMz|cNz:dOz;dPz=dQz?dRz}cSz@dTz~cUzadVzbdWzcdXz[dYz]dZz^d0z_d1z`d2z{d3z|d4z}d5z~d6zdd7zed8zfd9zc !z0c#z$z%z'z(z)z*z+z,z-z.z/z:z;z=zgd?z1c@z[z]z^z_z`z{z|z}z~zaA2cbAUgVgWgXgYgZgJhKhLhMhNhOhPhQhRhShThUhVhWhcAXhYhZh0h1h2h3h4h5h6hdAeAfAgA7h8h9haebecedeeehAfegeheiejeiAjAkAlAmAnAoApAqArAsAtAuAvAwAxAyAzAAABACADAEAFAGAHAIAJAKALAMANAOAPAQARASATAUAVAWAXAYAZA0A1A2A3A4A5A6A7A8A9A!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~AaBbBcBdBeBfBgBhBiBjBkBlBmBnBoBpBqBrBsBtBuBvBwBxByBzBABBBCBDBEBFBGBHBIBJBKBLBMBNBOBPBQBRBSBTBUBVBWBXBYBZB0B1B2B3B4B5B6B7B8B9B!B#B$B%B'B(B)B*B+B,B-B.B/B

215 cdef int out_len = len(result) 2a ZdTi0dUikfVi# Wi$ XilfYi% Zi' 0i( 1i) 2i* 3i+ 4i, 5i- 6i. 7i/ 8i: 9i; !i= #i? $i@ %i[ 'i] (i^ )i_ *i` +i{ ,i| -i} .i~ /iab:ibb;icb=idb?i@i[i]i^i_i`i{i|i}i~iajbjcjdjejfjgjhjijSejjkjljTemjUenjVeojWepjAbqjXerjYesjZetjujvjBbwjCbxjDbyjEbzjFbAjDcBjEcCjFcDjGcEjHcFjIcGjJcHjKcIjJj,hKj-hLj.hMjygNjzgOjAgPjQjRjSjTjUjVjWjXjYjZj0j1j2j3j4j5j6j7j8j9j!j#j$j%j'j(j)j*j+j,j-j.j/j:j;j=j?j@j[j]j^j_j`j{j|j}j~jakbkckdkekfkgkhkikjkkklkmknkokpkqkrksktkukvkwkxkykzkAkBkCkDkEkFkGkHkIkJkKkLkMkNkOkPkQkRkSkTkUkVkWkXkYkZk0k1k2k3k4k5k6k7k8k9k!k#k$k%k'k(k)k*k+k,k-k.k/k:k;k=k?k@k[k]k^k_k`k{k|k}k~kalblcldlelflglhliljlklllmlnlolplqlrlsltlulvlwlxlylzlAlXbBlClDl0eElFlGlHlIlJlKlLlMlNlOlPlQlRlSlTl1eUl2eVlWlafXlYlZl*c0l+c1l2l3lh 4li 5l6l7lGb8lHb9lIb!l#l3e$l{g%lJb'lKb(lLb)lMb*lNb+l,l-l.l/ld :l;l=l?l,c-c.c1d2d3de j f g 6 eb7 8 fbgbhb4d5d6d7d8d9d@lib[ljb]lOb^lkb_llb`lmb{lnb|lob}lpb~lambfbmcm~gah/hbhch:hdheh;hfh=hdmBgemcf?hfm5egmdfef@hffhmgf[himCgjmgh]hhhih^hhfkmiflmmmnmompmqm6erm9 ! sg!d/cYb:c+bLc;c=ctgQbRbSbTbUbVbDgsmtmumvmwmxmymzmEg7eAm8eBmCmDmEmFmGmHmImFgJmKmLmMmnhNmOmPmohGgHgQmphRmSmTmUmVmWmXmYmZm0m1m2m3m4m5m6m7m8mqhrhsh9m!mthuhvh#mwhxhyh$m%m'm(m)m*m+m,m-m.m/m:m;m=m?m@m[m]m^mqb,b-b_m`m{m|mb .b}m/b~m:ban;bbn=bcn?bdn@ben[bfn]bgn^bhn_bin`bjn{bkn|bln}bmn~bnnaconbcpnccqndcrnecsnfctngcunhcvnicwnjcxnkcynlcznmcAnncBnocCnDnpcEnFnGnHnIgqcrcJgKgLgMgNgOgPgQgzhAhke5crbtcsb6cuctbZb)e*e+ek 9bl m n o p q r s t !bu v w x 3by z 4b#ble7cubvcvb8cwcwb0b,e-e.eA $bB C D E F G H I J %bK L M N 5bO P 6b'bme9cxbxcyb!cyczb1b/e:e;eQ (bR S T U V W X Y Z )b0 1 2 3 7b4 5 8b*bRg?cInJnKn9eLnMnNnOnPnQnRnMcSn!eTn@cUn#eVnWnXnYnZn0n1n2n3n4n5n6n7n8n9n!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~naobocodoeofogohoiojokolomoSgTg$eno[coopoqorosotouovoscwoxoyo#dzoAoBoCoDoEoFoGoWbHoIoJoKoLoMoNoOoPoQoRoSoToUoRcVoScWoTcXoYoZo0oGhUc1oHhVc2oIhWc3o4o5o6o7o8o9o!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(e|w}w~waxbxcxdxexfxgxhxixjxkxlxmxnxoxpxqxrxsxtxvguxvxwxwgxxyxzxAxBxCxDxExFxGxHxIxJxKxLxMxNxOxPxQxRxSxTxUxVxWxXxYxZx0x1x2x3x4x5x6x7x8x9x!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~yazbzczdzezfzgzhzizjzkzlzmznzozpzqzrzsztzuzXcvzYcwz'dxzZcyz(dzz)dAz*dBz^cCz_cDz`cEz+dFz,dGz-dHz.dIz3cJz4cKz{cLz/dMz|cNz:dOz;dPz=dQz?dRz}cSz@dTz~cUzadVzbdWzcdXz[dYz]dZz^d0z_d1z`d2z{d3z|d4z}d5z~d6zdd7zed8zfd9zc !z0c#z$z%z'z(z)z*z+z,z-z.z/z:z;z=zgd?z1c@z[z]z^z_z`z{z|z}z~zaA2cbAUgVgWgXgYgZgJhKhLhMhNhOhPhQhRhShThUhVhWhcAXhYhZh0h1h2h3h4h5h6hdAeAfAgA7h8h9haebecedeeehAfegeheiejeiAjAkAlAmAnAoApAqArAsAtAuAvAwAxAyAzAAABACADAEAFAGAHAIAJAKALAMANAOAPAQARASATAUAVAWAXAYAZA0A1A2A3A4A5A6A7A8A9A!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~AaBbBcBdBeBfBgBhBiBjBkBlBmBnBoBpBqBrBsBtBuBvBwBxByBzBABBBCBDBEBFBGBHBIBJBKBLBMBNBOBPBQBRBSBTBUBVBWBXBYBZB0B1B2B3B4B5B6B7B8B9B!B#B$B%B'B(B)B*B+B,B-B.B/B

216 if out_len == 1: 2a ZdTi0dUikfVi# Wi$ XilfYi% Zi' 0i( 1i) 2i* 3i+ 4i, 5i- 6i. 7i/ 8i: 9i; !i= #i? $i@ %i[ 'i] (i^ )i_ *i` +i{ ,i| -i} .i~ /iab:ibb;icb=idb?i@i[i]i^i_i`i{i|i}i~iajbjcjdjejfjgjhjijSejjkjljTemjUenjVeojWepjAbqjXerjYesjZetjujvjBbwjCbxjDbyjEbzjFbAjDcBjEcCjFcDjGcEjHcFjIcGjJcHjKcIjJj,hKj-hLj.hMjygNjzgOjAgPjQjRjSjTjUjVjWjXjYjZj0j1j2j3j4j5j6j7j8j9j!j#j$j%j'j(j)j*j+j,j-j.j/j:j;j=j?j@j[j]j^j_j`j{j|j}j~jakbkckdkekfkgkhkikjkkklkmknkokpkqkrksktkukvkwkxkykzkAkBkCkDkEkFkGkHkIkJkKkLkMkNkOkPkQkRkSkTkUkVkWkXkYkZk0k1k2k3k4k5k6k7k8k9k!k#k$k%k'k(k)k*k+k,k-k.k/k:k;k=k?k@k[k]k^k_k`k{k|k}k~kalblcldlelflglhliljlklllmlnlolplqlrlsltlulvlwlxlylzlAlXbBlClDl0eElFlGlHlIlJlKlLlMlNlOlPlQlRlSlTl1eUl2eVlWlafXlYlZl*c0l+c1l2l3lh 4li 5l6l7lGb8lHb9lIb!l#l3e$l{g%lJb'lKb(lLb)lMb*lNb+l,l-l.l/ld :l;l=l?l,c-c.c1d2d3de j f g 6 eb7 8 fbgbhb4d5d6d7d8d9d@lib[ljb]lOb^lkb_llb`lmb{lnb|lob}lpb~lambfbmcm~gah/hbhch:hdheh;hfh=hdmBgemcf?hfm5egmdfef@hffhmgf[himCgjmgh]hhhih^hhfkmiflmmmnmompmqm6erm9 ! sg!d/cYb:c+bLc;c=ctgQbRbSbTbUbVbDgsmtmumvmwmxmymzmEg7eAm8eBmCmDmEmFmGmHmImFgJmKmLmMmnhNmOmPmohGgHgQmphRmSmTmUmVmWmXmYmZm0m1m2m3m4m5m6m7m8mqhrhsh9m!mthuhvh#mwhxhyh$m%m'm(m)m*m+m,m-m.m/m:m;m=m?m@m[m]m^mqb,b-b_m`m{m|mb .b}m/b~m:ban;bbn=bcn?bdn@ben[bfn]bgn^bhn_bin`bjn{bkn|bln}bmn~bnnaconbcpnccqndcrnecsnfctngcunhcvnicwnjcxnkcynlcznmcAnncBnocCnDnpcEnFnGnHnIgqcrcJgKgLgMgNgOgPgQgzhAhke5crbtcsb6cuctbZb)e*e+ek 9bl m n o p q r s t !bu v w x 3by z 4b#ble7cubvcvb8cwcwb0b,e-e.eA $bB C D E F G H I J %bK L M N 5bO P 6b'bme9cxbxcyb!cyczb1b/e:e;eQ (bR S T U V W X Y Z )b0 1 2 3 7b4 5 8b*bRg?cInJnKn9eLnMnNnOnPnQnRnMcSn!eTn@cUn#eVnWnXnYnZn0n1n2n3n4n5n6n7n8n9n!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~naobocodoeofogohoiojokolomoSgTg$eno[coopoqorosotouovoscwoxoyo#dzoAoBoCoDoEoFoGoWbHoIoJoKoLoMoNoOoPoQoRoSoToUoRcVoScWoTcXoYoZo0oGhUc1oHhVc2oIhWc3o4o5o6o7o8o9o!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(e|w}w~waxbxcxdxexfxgxhxixjxkxlxmxnxoxpxqxrxsxtxvguxvxwxwgxxyxzxAxBxCxDxExFxGxHxIxJxKxLxMxNxOxPxQxRxSxTxUxVxWxXxYxZx0x1x2x3x4x5x6x7x8x9x!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~yazbzczdzezfzgzhzizjzkzlzmznzozpzqzrzsztzuzXcvzYcwz'dxzZcyz(dzz)dAz*dBz^cCz_cDz`cEz+dFz,dGz-dHz.dIz3cJz4cKz{cLz/dMz|cNz:dOz;dPz=dQz?dRz}cSz@dTz~cUzadVzbdWzcdXz[dYz]dZz^d0z_d1z`d2z{d3z|d4z}d5z~d6zdd7zed8zfd9zc !z0c#z$z%z'z(z)z*z+z,z-z.z/z:z;z=zgd?z1c@z[z]z^z_z`z{z|z}z~zaA2cbAUgVgWgXgYgZgJhKhLhMhNhOhPhQhRhShThUhVhWhcAXhYhZh0h1h2h3h4h5h6hdAeAfAgA7h8h9haebecedeeehAfegeheiejeiAjAkAlAmAnAoApAqArAsAtAuAvAwAxAyAzAAABACADAEAFAGAHAIAJAKALAMANAOAPAQARASATAUAVAWAXAYAZA0A1A2A3A4A5A6A7A8A9A!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~AaBbBcBdBeBfBgBhBiBjBkBlBmBnBoBpBqBrBsBtBuBvBwBxByBzBABBBCBDBEBFBGBHBIBJBKBLBMBNBOBPBQBRBSBTBUBVBWBXBYBZB0B1B2B3B4B5B6B7B8B9B!B#B$B%B'B(B)B*B+B,B-B.B/B

217 return 2Zd0dkf# $ lf% ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdbSeTeUeVeWeAbXeYeZeBbCbDbEbFbDcEcFcGcHcIcJcKc,h-h.hygzgAgXb0e1e2eaf*c+ch i GbHbIb3eJbKbLbMbNbd ,c-c.c1d2d3de j f g 6 eb7 8 fbgbhb4d5d6d7d8d9dibjbObkblbmbnbobpb/h:h;h=hBg?h@h[hCg]h^h6esg/cYb:c+b;c=ctgQbRbSbTbUbVb7e8eqb,b-bIgqcrcJgKgLgMgNgOgPgQgzhAhke5crbtcsb6cuctbZb)e*e+ek 9bl m n o p q r s t !bu v w x 3by z 4b#ble7cubvcvb8cwcwb0b,e-e.eA $bB C D E F G H I J %bK L M N 5bO P 6b'bme9cxbxcyb!cyczb1b/e:e;eQ (bR S T U V W X Y Z )b0 1 2 3 7b4 5 8b*bRg?c9eMc!e@c#eSgTg$e[c(ekz

218 elif out_len == 2: 

219 return result[1] 2a ZdTi0dUikfVi# Wi$ XilfYi% Zi' 0i( 1i) 2i* 3i+ 4i, 5i- 6i. 7i/ 8i: 9i; !i= #i? $i@ %i[ 'i] (i^ )i_ *i` +i{ ,i| -i} .i~ /iab:ibb;icb=idb?i@i[i]i^i_i`i{i|i}i~iajbjcjdjejfjgjhjijSejjkjljTemjUenjVeojWepjAbqjXerjYesjZetjujvjBbwjCbxjDbyjEbzjFbAjDcBjEcCjFcDjGcEjHcFjIcGjJcHjKcIjJj,hKj-hLj.hMjygNjzgOjAgPjQjRjSjTjUjVjWjXjYjZj0j1j2j3j4j5j6j7j8j9j!j#j$j%j'j(j)j*j+j,j-j.j/j:j;j=j?j@j[j]j^j_j`j{j|j}j~jakbkckdkekfkgkhkikjkkklkmknkokpkqkrksktkukvkwkxkykzkAkBkCkDkEkFkGkHkIkJkKkLkMkNkOkPkQkRkSkTkUkVkWkXkYkZk0k1k2k3k4k5k6k7k8k9k!k#k$k%k'k(k)k*k+k,k-k.k/k:k;k=k?k@k[k]k^k_k`k{k|k}k~kalblcldlelflglhliljlklllmlnlolplqlrlsltlulvlwlxlylzlAlXbBlClDl0eElFlGlHlIlJlKlLlMlNlOlPlQlRlSlTl1eUl2eVlWlXlYlZl*c0l+c1l2l3lh 4li 5l6l7lGb8lHb9lIb!l#l3e$l{g%lJb'lKb(lLb)lMb*lNb+l,l-l.l/ld :l;l=l?l,c-c.c1d2d3de j f g 6 eb7 8 fbgbhb4d5d6d7d8d9d@lib[ljb]lOb^lkb_llb`lmb{lnb|lob}lpb~lambfbmcm~gah/hbhch:hdheh;hfh=hdmBgemcf?hfm5egmdfef@hffhmgf[himCgjmgh]hhhih^hhfkmiflmmmnmompmqm6erm9 ! !d/cYb:c+bLc;c=cQbRbSbTbUbVbDgsmtmumvmwmxmymzmEgAmBmCmDmEmFmGmHmImFgJmKmLmMmnhNmOmPmohGgHgQmphRmSmTmUmVmWmXmYmZm0m1m2m3m4m5m6m7m8mqhrhsh9m!mthuhvh#mwhxhyh$m%m'm(m)m*m+m,m-m.m/m:m;m=m?m@m[m]m^mqb,b-b_m`m{m|m}m~manbncndnenfngnhninjnknlnmnnnonpnqnrnsntnunvnwnxnynznAnBnCnDnEnFnGnHnIgqcrcJgKgLgMgNgOgPgQgke5crbtcsb6cuctbZb)e*e+ek 9bl m n o p q r s t !bu v w x 3by z 4b#ble7cubvcvb8cwcwb0b,e-e.eA $bB C D E F G H I J %bK L M N 5bO P 6b'bme9cxbxcyb!cyczb1b/e:e;eQ (bR S T U V W X Y Z )b0 1 2 3 7b4 5 8b*bRgInJnKn9eLnMnNnOnPnQnRnSnTnUnVnWnXnYnZn0n1n2n3n4n5n6n7n8n9n!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~naobocodoeofogohoiojokolomoSgTg$enooopoqorosotouovowoxoyo#dzoAoBoCoDoEoFoGoWbHoIoJoKoLoMoNoOoPoQoRoSoToUoVoWoXoYoZo0o1o2o3o4o5o6o7o8o9o!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(e|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~yazbzczdzezfzgzhzizjzlzmznzozpzqzrzsztzuzvzwzxzyzzzAzBzCzDzEzFzGzHzIzJzKzLzMzNzOzPzQzRzSzTzUzVzWzXzYzZz0z1z2z3z4z5z6z7z8z9z!z#z$z%z'z(z)z*z+z,z-z.z/z:z;z=z?z@z[z]z^z_z`z{z|z}z~zaAbAcAdAeAfAgAhAiAjAkAlAmAnAoApAqArAsAtAuAvAwAxAyAzAAABACADAEAFAGAHAIAJAKALAMANAOAPAQARASATAUAVAWAXAYAZA0A1A2A3A4A5A6A7A8A9A!A#A$A%A'A(A)A*A+A,A-A.A/A:A;A=A?A@A[A]A^A_A`A{A|A}A~AaBbBcBdBeBfBgBhBiBjBkBlBmBnBoBpBqBrBsBtBuBvBwBxByBzBABBBCBDBEBFBGBHBIBJBKBLBMBNBOBPBQBRBSBTBUBVBWBXBYBZB0B1B2B3B4B5B6B7B8B9B!B#B$B%B'B(B)B*B+B,B-B.B/B

220 else: 

221 return result[1:] 2a Zd0dkf# $ lf% ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdb,c-c.c1d2d3de f g 6 7 8 fbgbhb4d5d6d7d8d9dibjbkblbmbnbobpbb .b/b:b;b=b?b@b[b]b^b_b`b{b|b}b~bacbcccdcecfcgchcicjckclcmcncocpcscRcScTcGhUcHhVcIhWcvgwgXcYc'dZc(d)d*d^c_c`c+d,d-d.d3c4c{c/d|c:d;d=d?d}c@d~cadbdcd[d]d^d_d`d{d|d}d~dddedfdc 0cgd1c2cUgVgWgXgYgZgJhKhLhMhNhOhPhQhRhShThUhVhWhXhYhZh0h1h2h3h4h5h6h7h8h9haebecedeeefegeheieje

222  

223  

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

225 """ 

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

227 """ 

228 if options is None: 2Zd0dkf# $ lf% ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdbSeTeUeVeWeAbXeYeZeBbCbDbEbFbDcEcFcGcHcIcJcKcmfnfof}E~EhdpfcFdFidqfgFhFjdrfkFlFkdsfoFpFldtfsFtFmdXb0eufvfwfxf1e2eyfafwFxF*c+cyFzFh i [BzfGbHbIbAf3eJbKbLbMbNbd ,c-c.c1d2d3de j f g 6 eb7 8 fbgbhb4d5d6d7d8d9dibjbObkblbmbnbobpbAFBFbfCFDFEFcfFFGF5edfefffHFIFgfJFKFhfLFifMFNFOFPFQFRFSFTFUFVFWF6eXFYF]BZF^B0F1F2F_B3F`B4F9 5F! 6Fjhkh'h(hlhmh)h*hxi]E~K7F8F9F!F#F$F%F'F(F)F*F+F,F-F.F/F:F;F=F?F@F[F]F^F_F`F{F|F}F~FaGbGcGdG!deG/cfGYbgG:chG+bLc;ciG=cjGkGlGmGnGQbRbSbTbUbVbxLyLzLALBLCLTE7e8eDLEL{Bzc|BFLoGGLpGqGrGsGtGuGAcvGwGxGyG}B~BaCnezGAGBGCGDGEGFGGGHGIGJGKGLGXEYEMGZEndyi#cbC$ccC0EHL1EIL2E3E4ENG5E6E7E8E9E!E#E$EuLvLwLJLKLLL;Bqb,b-bodBcpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdoepeb .bVK/bWK:bXK;bYK=bZK?b0K@b1K[b2K]b3K^b4K_b5K`b@E{b[E|b6K}b7K~b8Kac9Kbc!Kcc#Kdc$Kec%Kfc'Kgc(Khc)Kic*Kjc+Kkcjflc,Kmc-Knc.Koc/KdC:KpcqcrcOGPGke5crbtcsb6cuctbZb)e*e+ek 9bl m n o p q r s t !bu v w x 3by z 4b#ble7cubvcvb8cwcwb0b,e-e.eA $bB C D E F G H I J %bK L M N 5bO P 6b'bme9cxbxcyb!cyczb1b/e:e;eQ (bR S T U V W X Y Z )b0 1 2 3 7b4 5 8b*bQGRGSG?cTGUG9eeCfCVGWGgChCiCjCMc!e@c#eXGYGZG0G1G2G3G4G5G6G7G8G9G!G#G$GkC%GlC'GmC(GnC)GoC*GpC+GqC,GrC-G.G/G:G;G=G?G@G[G]G^G_G`G{G|G}G~G%Ezi'E(EAi)EaH$ebHcHdH[ceHfHgHhHiHjHsCtCNcOcPbPcQcscuCvC#dwCxCyCzCACCcBCWbCCDCECFCGCHCICJCKCLCMCNCOCPCRcScTcQCRCSCUcVcWcTCUCVCWCXCYCZC0C1C2C3C4C5C6C7C8C9C!C#C$C%C'C(C)C*C+C,C-C.C/C:C;C=C?C@C[C]C^C_C`C{C|C}C~CaDbDcDdDeDfDgDhDiDjDkDlDmDnDoDpDqDrDsDtDuDug%e'e]c$dkHBflHmHnHvDCfwDpHqHxDrHsHDftHuHvHyDEfzDxHFfyHzHGfBHqeCHDH*E+EEHFHGHHHIHBire=eJHCiDiKHLHMHseteHfIfJfKfLfMfNHOHNfQHOfRHSHADPfBDUHQfVHWHXHCDRfDDZHSf0H1HTf3H,Eue4H5H6H7H8H9H!H#H$H%HEDve?e'HEiFD(H)H*H+H,H-H.H/H:H;H=H?H@H[HUf]H^HVf`H{HWfGDweHDIDJDKDLDMDNDODPDxeQDBhRDSDTDXfyezeYfZf0f1f2fRd3f4f%c@e[eSd'cTd5f}H~HUdFiGiUDAeVDWDXDYDZD0D1D2D3D4DBe5DCh6D7D8D6fgIhI7fjI-ECekIlImInIoIpIqIrIsItIuI9DDe^evIHi!DwIxIyIzI8fAIBICI#D9f$DEI!fFIGI#fIIEeJIKI.E/ELIMINIOIPIIiFe_eQIJiKiRISITIGeHe$f%f'f(f)f*fUIVI+fXI,fYIZI%D-f'D1I.f2I3I4I(D/f)D6I:f7I8I;f!I:EIe#I$I%I'I(I)I*I+I,I-I*DJe`e.ILi+D/I:I;I=I?I@I[I]I^I_I`I{I|I}I=f~IaJ?fcJdJ@f,DKe-D.D/D:D;D=D?D@D[DLe]DDh^D_D`D[fMeNe]f^f_f`f{fVd|f}f(c{e|eWd)cXd~ffJgJYdMiNi{DOe|D}D~DaEbEcEdEeEfEgEPehEEhiEjEkEagnJoJbgqJ;EQerJsJtJuJvJwJxJyJzJAJBJlERe~eCJOimEDJEJFJGJ(ePiQinEoEHJcgIJJJKJpEdgqErEsEtEMJNJegOJPJQJuEfgvESJTJggUJVJWJwEhgxEYJZJig0J1J2JyEjgzE4J5Jkg6J7J8JAElgBE!J#Jmg$J%J'JCEngDE)J*Jog+J,J-JEEpgFEGE/J:Jqg;J=J?JHErgIEJE[JFKGKXcYc=E'dZc(d)d*d^c_c`c+d,d-d.d3c4c{c/d|c:d;d=d?d}c@d~cadbdcd[d]d^d_d`d{d|d}d~dddedfdc 0cKE2bLE%dHKIKJKKKLKMKNKOKPKQKRKSKTKUKRiSi]JMENEOEgdPE^J_JQE?E1c+h;K=K?K@K[K]K^K_K`K{K2caebecedeeefegezKAKBKCKheDKiejeEKMLNLOLPLQLRL`JSLTLULVLWLXLYLZL0LRE1L{J|J2L}J3L~JaKbKcKdKeKfKgKhKiKjKkKlKmKnKoKpKqKrKsKtKuKvKwKxKyK

229 if keep_none: 2Zd0dkf# $ lf% ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdbSeTeUeVeWeAbXeYeZeBbCbDbEbFbDcEcFcGcHcIcJcKc}E~EcFdFgFhFkFlFoFpFsFtF0e1e2e*c+ch i GbHbIb3eJbKbLbMbNbd ,c-c.c1d2d3de j f g 6 eb7 8 fbgbhb4d5d6d7d8d9dibjbObkblbmbnbobpbAFBFbfCFDFEFcfFFGF5edfefffHFIFgfJFKFhfLFifMFNFOFPFQFRFSFTFUFVFWF6e9 ! jhkh'h(hlhmh)h*hxi]E~K!d/cYb:c+bLc;c=cQbRbSbTbUbVbxLyLzLALBLCL7e8eDLEL{Bzc|BoGpGqGrGsGtGuGAcvGwGxGyG}B~BaCnezGAGBGCGDGEGFGGGHGIGJGKGLGMGnd#c$cILNGqb,b-bBcqcrcOGPGke5crbtcsb6cuctbZb)e*e+ek 9bl m n o p q r s t !bu v w x 3by z 4b#ble7cubvcvb8cwcwb0b,e-e.eA $bB C D E F G H I J %bK L M N 5bO P 6b'bme9cxbxcyb!cyczb1b/e:e;eQ (bR S T U V W X Y Z )b0 1 2 3 7b4 5 8b*bQGRGSG?cTG9eeCfCVGWGgChCiCjCMc!e@c#ekC%GlC'GmC(GnC)GoC*GpC+GqC,GrC-GziAi$ebHcHdH[csCtCuCvC#dwCxCyCzCACCcBCWbCCDCECFCGCHCICJCKCLCMCNCOCPCUcVcWcTCUCVCWCXCYCZC0C1C2C3C4C5C6C7C8C9C!C#C$C%C'C(C)C*C+C,C-C.C/C:C;C=C?C@C[C]C^C_C`C{C|C}C~CaDbDcDdDeDfDgDhDiDjDkDlDmDnDoDpDqDrDsDtDuD]ckHlHmHnHvDwDpHqHxDrHsHtHuHvHyDzDxHyHzHBHqeCHDH*E+EEHFHGHHHIHBire=eJHCiDiKHLHMHseteNHOHQHRHSHADBDUHVHWHXHCDDDZH0H1H3H,Eue4H5H6H7H8H9H!H#H$H%HEDve?e'HEiFD(H)H*H+H,H-H.H/H:H;H=H?H@H[H]H^H`H{HGDweHDIDJDKDLDMDNDODPDxeQDBhRDSDTDyeze@e[eSdTd}H~HFiGiUDAeVDWDXDYDZD0D1D2D3D4DBe5DCh6D7D8DgIhIjI-ECekIlImInIoIpIqIrIsItIuI9DDe^evIHi!DwIxIyIzIAIBICI#D$DEIFIGIIIEeJIKI.E/ELIMINIOIPIIiFe_eQIJiKiRISITIGeHeUIVIXIYIZI%D'D1I2I3I4I(D)D6I7I8I!I:EIe#I$I%I'I(I)I*I+I,I-I*DJe`e.ILi+D/I:I;I=I?I@I[I]I^I_I`I{I|I}I~IaJcJdJ,DKe-D.D/D:D;D=D?D@D[DLe]DDh^D_D`DMeNe{e|eWdXdfJgJMiNi{DOe|D}D~DaEbEcEdEeEfEgEPehEEhiEjEkEnJoJqJ;EQerJsJtJuJvJwJxJyJzJAJBJlERe~eCJOimEDJEJFJPiQinEoEHJIJJJKJpEqErEsEtEMJNJOJPJQJuEvESJTJUJVJWJwExEYJZJ0J1J2JyEzE4J5J6J7J8JAEBE!J#J$J%J'JCEDE)J*J+J,J-JEEFEGE/J:J;J=J?JHEIEJE[JYcc 0cKE2bLE%d]JMENEOEPE^J_JQE1c+hOLPLQLRLTLULVLWLZLRE1L2L}J~JaKbKcKdKeKfKgKhKiKjKkKlKmKnKoKpKqKrKsKtKuKvKwKxKyK

230 return options 2h i ,c-c.ce j f g 6 eb7 8 BFCFDFcfGF5eefHFgfKFLFMFOFQFSFUFWF7e8end#c$cOGPGQGRGSGTG9eeCfCVGWGgChCiCjCMc!e@c#ekC%GlC'GmC(GnC)GoC*GpC+GqC,GrC-GziAibHcHdHCckHqHsHBHqeCHDH*E+EEHFHGHHHIHBire=eJHCiDiKHLHMHQHUH,E-EzIIIEeJIKI.E/ELIMINIOIPIIiFe_eQIJiKiRISITIXI1I:E;EHJNJTJZJ5J#J*J:J~JaKbKcKdKeKfKgKhKiKjKkKlKmKnKoKpKqKrKsKtKuKvKwKxKyK

231 return cls() 2Zd0dkf# $ lf% ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdbSeTeUeVeWeAbXeYeZeBbCbDbEbFbDcEcFcGcHcIcJcKc}E~EcFdFgFhFkFlFoFpFsFtF0e1e2e*c+ch i GbHbIb3eJbKbLbMbNbd ,c-c.c1d2d3de j f g 6 eb7 8 fbgbhb4d5d6d7d8d9dibjbObkblbmbnbobpbAFbfEFFF5edfffIFJFhfifNFPFRFTFVF6e9 ! jhkh'h(hlhmh)h*hxi]E~K!d/cYb:c+bLc;c=cQbRbSbTbUbVbxLyLzLALBLCLDLEL{Bzc|BoGpGqGrGsGtGuGAcvGwGxGyG}B~BaCnezGAGBGCGDGEGFGGGHGIGJGKGLGMGnd#c$cILNGqb,b-bBcqcrcke5crbtcsb6cuctbZb)e*e+ek 9bl m n o p q r s t !bu v w x 3by z 4b#ble7cubvcvb8cwcwb0b,e-e.eA $bB C D E F G H I J %bK L M N 5bO P 6b'bme9cxbxcyb!cyczb1b/e:e;eQ (bR S T U V W X Y Z )b0 1 2 3 7b4 5 8b*b?cMc$e[csCtCuCvC#dwCxCyCzCACCcBCWbCCDCECFCGCHCICJCKCLCMCNCOCPCUcVcWcTCUCVCWCXCYCZC0C1C2C3C4C5C6C7C8C9C!C#C$C%C'C(C)C*C+C,C-C.C/C:C;C=C?C@C[C]C^C_C`C{C|C}C~CaDbDcDdDeDfDgDhDiDjDkDlDmDnDoDpDqDrDsDtDuD]clHmHnHvDwDpHxDrHtHuHvHyDzDxHyHzH*E+EBiDiseteNHOHRHSHADBDVHWHXHCDDDZH0H1H3H,Eue4H5H6H7H8H9H!H#H$H%HEDve?e'HEiFD(H)H*H+H,H-H.H/H:H;H=H?H@H[H]H^H`H{HGDweHDIDJDKDLDMDNDODPDxeQDBhRDSDTDyeze@e[eSdTd}H~HFiGiUDAeVDWDXDYDZD0D1D2D3D4DBe5DCh6D7D8DgIhIjI-ECekIlImInIoIpIqIrIsItIuI9DDe^evIHi!DwIxIyIAIBICI#D$DEIFIGI.E/EIiKiGeHeUIVIYIZI%D'D2I3I4I(D)D6I7I8I!I:EIe#I$I%I'I(I)I*I+I,I-I*DJe`e.ILi+D/I:I;I=I?I@I[I]I^I_I`I{I|I}I~IaJcJdJ,DKe-D.D/D:D;D=D?D@D[DLe]DDh^D_D`DMeNe{e|eWdXdfJgJMiNi{DOe|D}D~DaEbEcEdEeEfEgEPehEEhiEjEkEnJoJqJ;EQerJsJtJuJvJwJxJyJzJAJBJlERe~eCJOimEDJEJFJPiQinEoEIJJJKJpEqErEsEtEMJOJPJQJuEvESJUJVJWJwExEYJ0J1J2JyEzE4J6J7J8JAEBE!J$J%J'JCEDE)J+J,J-JEEFEGE/J;J=J?JHEIEJE[JYcc 0cKE2bLE%d]JMENEOEPE^J_JQE1c+hOLPLQLRLTLULVLWLZLRE1L2L}J

232 elif isinstance(options, cls): 2# $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdbAbBbCbDbEbFbmfnfofhdpfidqfjdrfkdsfldtfmdXbufvfwfxfyfafwFxF*c+cyFzFh i [BzfGbHbIbAfJbKbLbMbNbd e j f g 6 eb7 8 fbgbhbibjbObkblbmbnbobpbXFYF]BZF^B0F1F2F_B3F`B4F9 5F! 6Fjhkh'h(hlhmh)h*hxi]E~K7F8F9F!F#F$F%F'F(F)F*F+F,F-F.F/F:F;F=F?F@F[F]F^F_F`F{F|F}F~FaGbGcGdGeGfGgGhG+bLciGjGkGlGmGnGQbRbSbTbUbVbTE{Bzc|BFLGLAc}B~BaCneXEYEZEndyi#cbC$ccC0EHL1E2E3E4E5E6E7E8E9E!E#E$EuLvLwLJLKLLL;BqbodBcpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdoepeb .bVK/bWK:bXK;bYK=bZK?b0K@b1K[b2K]b3K^b4K_b5K`b@E{b[E|b6K}b7K~b8Kac9Kbc!Kcc#Kdc$Kec%Kfc'Kgc(Khc)Kic*Kjc+Kkcjflc,Kmc-Knc.Koc/KdC:Kpc?cUGeCfCgChCiCjCMc@cXGYGZG0G1G2G3G4G5G6G7G8G9G!G#G$GkClCmCnCoCpCqCrC.G/G:G;G=G?G@G[G]G^G_G`G{G|G}G~G%Ezi'E(EAi)EaH[ceHfHgHhHiHjHNcOcPbPcQcscCcRcScTcQCRCSCug%e'e]c$dBfCfDfEfFfGfqereseteHfIfJfKfLfMfNfOfPfQfRfSfTfueveUfVfWfwexeXfyezeYfZf0f1f2fRd3f4f%c'c5fUdAeBe6f7fCeDe8f9f!f#fEeFeGeHe$f%f'f(f)f*f+f,f-f.f/f:f;fIeJe=f?f@fKeLe[fMeNe]f^f_f`f{fVd|f}f(c)c~fYdOePeagbgQeReGJ(ePiQicgdgegfggghgigjgkglgmgngogpgqgrgFKGKXc=E'dZc(d)d*d^c_c`c+d,d-d.d3c4c{c/d|c:d;d=d?d}c@d~cadbdcd[d]d^d_d`d{d|d}d~dddedfdHKIKJKKKLKMKNKOKPKQKRKSKTKUKRiSigd?E+h;K=K?K@K[K]K^K_K`K{K2caebecedeeefegezKAKBKCKheDKiejeEKMLNL`JSLXLYL0LRE{J|J3L

233 return options 2# $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ abbbcbdbAbBbCbDbEbFbmfnfofhdpfidqfjdrfkdsfldtfmdXbufvfwfxfyfafwFxF*c+cyFzFh i [BzfGbHbIbAfJbKbLbMbNbd e j f g 6 eb7 8 fbgbhbibjbObkblbmbnbobpbXFYF]BZF^B0F1F2F_B3F`B4F9 5F! 6F'h(h)h*h~K7F8F9F!F#F$F%F'F(F)F*F+F,F-F.F/F:F;F=F?F@F[F]F^F_F`F{F|F}F~FaGbGcGdGeGfGgGhG+bLciGjGkGlGmGnGQbRbSbTbUbVb{Bzc|BFLGLAc}B~BaCneXEYEZEndyi#cbC$ccC0EHL1E2E3E4E5E6E7E8E9E!E#E$EuLvLwLJLKLLL;BqbodBcpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdoepeb .bVK/bWK:bXK;bYK=bZK?b0K@b1K[b2K]b3K^b4K_b5K`b@E{b[E|b6K}b7K~b8Kac9Kbc!Kcc#Kdc$Kec%Kfc'Kgc(Khc)Kic*Kjc+Kkcjflc,Kmc-Knc.Koc/KdC:Kpc?cUGeCfCgChCiCjCMc@cXGYGZG0G1G2G3G4G5G6G7G8G9G!G#G$GkClCmCnCoCpCqCrC.G/G:G;G=G?G@G[G]G^G_G`G{G|G}G~GaH[ceHfHgHhHiHjHNcOcPbPcQcscCcRcScTcQCRCSCug%e'e]c$dBfCfDfEfFfGfqereseteHfIfJfKfLfMfNfOfPfQfRfSfTfueveUfVfWfwexeXfyezeYfZf0f1f2fRd3f4f%c'c5fUdAeBe6f7fCeDe8f9f!f#fEeFeGeHe$f%f'f(f)f*f+f,f-f.f/f:f;fIeJe=f?f@fKeLe[fMeNe]f^f_f`f{fVd|f}f(c)c~fYdOePeagbgQeReGJ(ecgdgegfggghgigjgkglgmgngogpgqgrgFKGKXc=E'dZc(d)d*d^c_c`c+d,d-d.d3c4c{c/d|c:d;d=d?d}c@d~cadbdcd[d]d^d_d`d{d|d}d~dddedfdHKIKJKKKLKMKNKOKPKQKRKSKTKUKRiSi?E+h;K=K?K@K[K]K^K_K`K{K2caebecedeeefegezKAKBKCKheDKiejeEKMLNL`JSLXLYL0LRE{J|J3L

234 elif isinstance(options, dict): 2jhkhlhmhxi]ETE;Bod%Ezi'E(EAi)EPiQigd

235 return cls(**options) 2jhkhlhmhxi]Eod%Ezi'E(EAi)EPiQigd

236 else: 

237 raise TypeError( 2TE;B

238 f"The {options_description} must be provided as an object " 2TE;B

239 f"of type {cls.__name__} or as a dict with valid {options_description}. " 2TE;B

240 f"The provided object is '{options}'." 2TE;B

241 ) 

242  

243  

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

245 """ 

246 Convert a boolean option to a string representation. 

247 """ 

248 return "true" if bool(option) else "false" 2b .b/b:b;b=b?b@b[b]b^b_b`b{b|b}b~bacbcccdcecfcgchcicjckclcmcncocpcsc^c_c`c|cbdcdddedfdf3

249  

250  

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

252 """ 

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

254  

255 Args: 

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

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

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

259  

260 Returns: 

261 Callable: A decorator that creates the wrapping. 

262 """ 

263  

264 def outer(wrapped_function: Callable) -> Callable: 2UE

265 """ 

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

267 """ 

268  

269 @functools.wraps(wrapped_function) 2UE

270 def inner(*args, **kwargs) -> Any: 

271 """ 

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

273 """ 

274 checker(*args, **kwargs, what=what) 2UE

275 result = wrapped_function(*args, **kwargs) 2UE

276  

277 return result 2UE

278  

279 return inner 2UE

280  

281 return outer 2UE

282  

283  

284def is_sequence(obj: object) -> bool: 

285 """ 

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

287 """ 

288 return isinstance(obj, Sequence) 29 ! zcAcXEYEZEyi#c$c0E1E2E3E4E5E6E7E8E9E!E#E$EuLvLwLqbodpdqdrdsdtdudvdwdxdydzdAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQd@E[E=E3c4c{c}c~cadg3h3i3j3k3RiSil3m3n3[2]2?EUgVgWgXgYgZgo3p3

289  

290  

291def is_nested_sequence(obj: object) -> bool: 

292 """ 

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

294 """ 

295 return is_sequence(obj) and any(is_sequence(elem) for elem in obj) 23c4c[2]2

296  

297  

298  

299class Transaction: 

300 """ 

301 A context manager for transactional operations with undo capability. 

302  

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

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

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

306  

307 Usage: 

308 with Transaction() as txn: 

309 txn.append(some_cleanup_function, arg1, arg2) 

310 # ... perform operations ... 

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

312  

313 Methods: 

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

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

316 """ 

317 def __init__(self) -> None: 

318 self._stack = ExitStack() 2NcOcPbPcQc

319 self._entered = False 2NcOcPbPcQc

320  

321 def __enter__(self): 

322 self._stack.__enter__() 2NcOcPbPcQc

323 self._entered = True 2NcOcPbPcQc

324 return self 2NcOcPbPcQc

325  

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

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

328 self._entered = False 2NcOcPbPcQc

329 return self._stack.__exit__(exc_type, exc, tb) 2NcOcPbPcQc

330  

331 def append(self, fn: Callable[..., Any], /, *args: Any, **kwargs: Any) -> None: 

332 """ 

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

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

335 """ 

336 if not self._entered: 2NcOcPbPcQc

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

338 self._stack.callback(partial(fn, *args, **kwargs)) 2NcOcPbPcQc

339  

340 def commit(self) -> None: 

341 """ 

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

343 """ 

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

345 self._stack.pop_all() 2NcOcPbPcQc

346  

347  

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

349_fork_warning_checked = False 

350  

351  

352def reset_fork_warning() -> None: 

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

354  

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

356 to check the warning behavior. 

357 """ 

358 global _fork_warning_checked 

359 _fork_warning_checked = False 2ug%e'e]c!h$d

360  

361  

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

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

364 cdef unsigned int val 

365 if width == 1: 2liminioipiqirisitiuiviwi@Bd 5crbtcsb6cuctbk 9bl m n o p q r s t !bu v w x 3by z 4b#b7cubvcvb8cwcwbA $bB C D E F G H I J %bK L M N 5bO P 6b'b9cxbxcyb!cyczbQ (bR S T U V W X Y Z )b0 1 2 3 7b4 5 8b*b

366 val = (<uint8_t*>ptr)[0] 2rbl m n u v ubB C D K L xbR S T 0 1

367 elif width == 2: 

368 val = (<uint16_t*>ptr)[0] 2linipiritivitcsbo p q w x 3bvcvbE F G M N 5bxcybU V W 2 3 7b

369 elif width == 4: 

370 val = (<uint32_t*>ptr)[0] 2mioiqisiuiwid uctbk r s t y z 4bwcwbA H I J O P 6byczbQ X Y Z 4 5 8b

371 else: 

372 raise ValueError(f"value must be 1, 2, or 4 bytes, got {width}") 2@B5c6c9b!b#b7c8c$b%b'b9c!c(b)b*b

373 return (val, <unsigned int>width) 2liminioipiqirisitiuiviwid rbtcsbuctbk l m n o p q r s t u v w x 3by z 4bubvcvbwcwbA B C D E F G H I J K L M N 5bO P 6bxbxcybyczbQ R S T U V W X Y Z 0 1 2 3 7b4 5 8b

374  

375  

376cpdef tuple _parse_fill_value(value): 

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

378  

379 Parameters 

380 ---------- 

381 value : int or buffer-protocol object 

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

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

384  

385 Returns 

386 ------- 

387 tuple of (int, int) 

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

389  

390 Raises 

391 ------ 

392 OverflowError 

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

394 TypeError 

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

396 ValueError 

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

398 """ 

399 cdef uint8_t byte_val 

400 cdef Py_buffer buf 

401  

402 if isinstance(value, int): 2DcEcFcGcHcIcJcKcaFbFlimieFfFnioiiFjFpiqimFnFrisiqFrFtiuiuFvFviwi@Bh i d 4ee j f g ke5crbtcsb6cuctbZb)e*e+ek 9bl m n o p q r s t !bu v w x 3by z 4b#ble7cubvcvb8cwcwb0b,e-e.eA $bB C D E F G H I J %bK L M N 5bO P 6b'bme9cxbxcyb!cyczb1b/e:e;eQ (bR S T U V W X Y Z )b0 1 2 3 7b4 5 8b*boHwHAHCiPHTHYH2HEi_H|HBhRd]eUdFiGiaIbIcIdIeIfIChiIHiDIHIJiWI0I5I9ILibJeJDhVd}eYdMiNihJiJjJkJlJmJEhpJOiLJRJXJ3J9J(J.J@J

403 byte_val = value 2DcEcFcGcHcIcJcKcaFbFeFfFiFjFmFnFqFrFuFvFh i 4ee j f g Zb)e*e+e0b,e-e.e1b/e:e;eoHwHAHCiPHTHYH2HEi_H|HBhRd]eUdFiGiaIbIcIdIeIfIChiIHiDIHIJiWI0I5I9ILibJeJDhVd}eYdMiNihJiJjJkJlJmJEhpJOiLJRJXJ3J9J(J.J@J

404 return (<unsigned int>byte_val, <unsigned int>1) 2DcEcFcGcHcIcJcKcaFbFeFfFiFjFmFnFqFrFuFvFh i 4ee j f g Zb0b1boHwHAHCiPHTHYH2HEi_H|HBhRd]eUdFiGiaIbIcIdIeIfIChiIHiDIHIJiWI0I5I9ILibJeJDhVd}eYdMiNihJiJjJkJlJmJEhpJOiLJRJXJ3J9J(J.J@J

405  

406 if isinstance(value, bytes): 2liminioipiqirisitiuiviwi@Bd ke5crbtcsb6cuctbk 9bl m n o p q r s t !bu v w x 3by z 4b#ble7cubvcvb8cwcwbA $bB C D E F G H I J %bK L M N 5bO P 6b'bme9cxbxcyb!cyczbQ (bR S T U V W X Y Z )b0 1 2 3 7b4 5 8b*b

407 return _read_fill_ptr(<const char*><bytes>value, <Py_ssize_t>len(value)) 2liminioipiqirisitiuiviwi@B5crbtcsb6cuctb7cubvcvb8cwcwb9cxbxcyb!cyczb

408  

409 if PyObject_GetBuffer(value, &buf, PyBUF_SIMPLE) != 0: 2d kek 9bl m n o p q r s t !bu v w x 3by z 4b#bleA $bB C D E F G H I J %bK L M N 5bO P 6b'bmeQ (bR S T U V W X Y Z )b0 1 2 3 7b4 5 8b*b

410 raise TypeError( 

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

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

413 ) 

414 try: 2d k 9bl m n o p q r s t !bu v w x 3by z 4b#bA $bB C D E F G H I J %bK L M N 5bO P 6b'bQ (bR S T U V W X Y Z )b0 1 2 3 7b4 5 8b*b

415 return _read_fill_ptr(<const char*>buf.buf, buf.len) 2d k 9bl m n o p q r s t !bu v w x 3by z 4b#bA $bB C D E F G H I J %bK L M N 5bO P 6b'bQ (bR S T U V W X Y Z )b0 1 2 3 7b4 5 8b*b

416 finally: 

417 PyBuffer_Release(&buf) 

418  

419  

420def check_multiprocessing_start_method() -> None: 

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

422 global _fork_warning_checked 

423 if _fork_warning_checked: 2aLbL]B^BcLdL_B`B9 ! jhkh'h(hlhmh)h*heLfL^2gL_2`2hL{2|2iL}2~2jLa3b3kLc3d3lLe3mLsgnLoLpLqL!d/cYb:c+bLc;c=ctgrLsLtLQbRbSbTbUbVbug%e'e]c!h$d

424 return 2aLbL]B^BcLdL_B`B9 ! jhkh'h(hlhmh)h*heLfL^2gL_2`2hL{2|2iL}2~2jLa3b3kLc3d3lLe3mLsgnLoLpLqL!d/cYb:c+bLc;c=ctgrLsLtLQbRbSbTbUbVb$d

425 _fork_warning_checked = True 2Ybug%e'e]c!h$d

426  

427 # Common warning message parts 

428 common_message = ( 

429 "CUDA does not support. Forked subprocesses exhibit undefined behavior, " 2Ybug%e'e]c!h$d

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

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

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

433 ) 

434  

435 try: 2Ybug%e'e]c!h$d

436 start_method = multiprocessing.get_start_method() 2Ybug%e'e]c!h$d

437 if start_method == "fork": 2Ybug%e'e]c$d

438 message = f"multiprocessing start method is 'fork', which {common_message}" 2%e'e]c$d

439 warnings.warn(message, UserWarning, stacklevel=3) 2%e'e]c$d

440 except RuntimeError: 2!h

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

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

443 if platform.system() == "Linux": 2!h

444 message = ( 

445 f"multiprocessing start method is not set and defaults to 'fork' on Linux, " 2!h

446 f"which {common_message}" 

447 ) 

448 warnings.warn(message, UserWarning, stacklevel=3) 2!h