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

233 statements  

« prev     ^ index     » next       coverage.py v7.15.2, created at 2026-07-29 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 28 vbve9 ! wexeyewbzexbAeybzbAbBbCb# $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ DbEbFbGbHbIbBgCgDgFjvdEg%hwdFgxdGgydHgzdIgAd!bJgKgLgMgNge f OgJbqbKbPgLbMbNbObPbd g h i j abbbcbdbrbsbtbebQbRbSbTbUbBeVbfbk 7 RcScCeBd-c.cgbCdijcJdJeJGjMhfJHjTcDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdUdVdWdXdYdZd0d1d2d3d4dDeEeIjgJJjUc#b$bQgRgKjSgTgLjUgVgFeGecfHeIeWgXgYgZg0g1g2g3g4g'h5g6g(h7g8gJeKedf9g!g#gLeMe$gNeOe%g'g(g)g*g5d+g,g)h/c*h+hefff,h-h.h/h:h6d:c;hgf7d=h?h@h-g8dPeQe.g/gReSehf:g;gMj=g?gTeUeifVeWe@g[g]g^g_g`g{g|g}g[h~gah]hbhchXeYejfdhehfhZe0egh1e2ehhihjhkhlh9dmhnh^h;c_h`hkflf{h|h}h~hai!d=cbimf#dcidieioh$d3e4ephqh5e6enfrhshNjthuhOjvhwhPjxhyhQjzhAhRjBhChSjDhEhTjFhGhUjjj4I5IaGhJkj

50 if isinstance(cfg, int): 28 vbve9 ! wexeyewbzexbAeybzbAbBbCb# $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ DbEbFbGbHbIbBgCgDgFjvdEg%hwdFgxdGgydHgzdIgAd!bJgKgLgMgNge f OgJbqbKbPgLbMbNbObPbd g h i j abbbcbdbrbsbtbebQbRbSbTbUbBeVbfbk 7 RcScCeBd-c.cgbCdijcJdJeJGjMhfJHjTcDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdUdVdWdXdYdZd0d1d2d3d4dDeEeIjgJJjUc#b$bQgRgKjSgTgLjUgVgFeGecfHeIeWgXgYgZg0g1g2g3g4g'h5g6g(h7g8gJeKedf9g!g#gLeMe$gNeOe%g'g(g)g*g5d+g,g)h/c*h+hefff,h-h.h/h:h6d:c;hgf7d=h?h@h-g8dPeQe.g/gReSehf:g;gMj=g?gTeUeifVeWe@g[g]g^g_g`g{g|g}g[h~gah]hbhchXeYejfdhehfhZe0egh1e2ehhihjhkhlh9dmhnh^h;c_h`hkflf{h|h}h~hai!d=cbimf#dcidieioh$d3e4ephqh5e6enfrhshNjthuhOjvhwhPjxhyhQjzhAhRjBhChSjDhEhTjFhGhUjjj4I5IaGhJkj

51 cfg = (cfg,) 28 vbve9 ! wexeyewbzexbAeybzbAbBbCb# $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ DbEbFbGbHbIbBgCgDgEgFgGgHgIg!bJgKgLgMgNge f OgJbqbKbPgLbMbNbObPbd g h i j abbbcbdbrbsbtbebQbRbSbTbUbBeVbfbk 7 RcScCeBd-c.cgbCdijcJdJeJMhfJDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdUdVdWdXdYdZd0d1d2d3d4dDeEegJUc#b$bQgRgSgTgUgVgFeGeHeIeWgXgYgZg0g1g2g3g4g5g6g7g8gJeKe9g!g#gLeMe$gNeOe%g'g(g)g*g5d+g,g/c:c-g8dPeQe.g/gReSe:g;g=g?gTeUeVeWe@g[g]g^g_g`g{g|g}g~gahbhchXeYedhehfhZe0egh1e2ehhihjhkhlh9dmhnh;c=coh$d3e4ephqh5e6erhshthuhvhwhxhyhzhAhBhChDhEhFhGhjj4I

52 else: 

53 common = "must be an int, or a tuple with up to 3 ints" 2Fjvd%hwdxdydzdAdijGjMhHjTcIjJjKjLjcf'h(hdf)h/c*h+hefff,h-h.h/h:h6d:c;hgf7d=h?h@hhfMjif[h]hjf^h;c_h`hkflf{h|h}h~hai!d=cbimf#dcidieinfNjOjPjQjRjSjTjUjjj5IaGhJkj

54 if not isinstance(cfg, tuple): 2Fjvd%hwdxdydzdAdijGjMhHjTcIjJjKjLjcf'h(hdf)h/c*h+hefff,h-h.h/h:h6d:c;hgf7d=h?h@hhfMjif[h]hjf^h;c_h`hkflf{h|h}h~hai!d=cbimf#dcidieinfNjOjPjQjRjSjTjUjjj5IaGhJkj

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

56 if len(cfg) > 3: 2Fjvd%hwdxdydzdAdijGjMhHjTcIjJjKjLjcf'h(hdf)h/c*h+hefff,h-h.h/h:h6d:c;hgf7d=h?h@hhfMjif[h]hjf^h;c_h`hkflf{h|h}h~hai!d=cbimf#dcidieinfNjOjPjQjRjSjTjUjjj5IaGkj

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

58 if any(not isinstance(val, int) for val in cfg): 2Fjvd%hwdxdydzdAdijGjMhHjTcIjJjKjLjcf'h(hdf)h/c*h+hefff,h-h.h/h:h6d:c;hgf7d=h?h@hhfMjif[h]hjf^h;c_h`hkflf{h|h}h~hai!d=cbimf#dcidieinfNjOjPjQjRjSjTjUjjjaGkj

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

60 if any(val < 1 for val in cfg): 28 vbve9 ! wexeyewbzexbAeybzbAbBbCb# $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ DbEbFbGbHbIbBgCgDgFjvdEg%hwdFgxdGgydHgzdIgAd!bJgKgLgMgNge f OgJbqbKbPgLbMbNbObPbd g h i j abbbcbdbrbsbtbebQbRbSbTbUbBeVbfbk 7 RcScCeBd-c.cgbCdijcJdJeJGjMhfJHjTcDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdUdVdWdXdYdZd0d1d2d3d4dDeEeIjgJJjUc#b$bQgRgKjSgTgLjUgVgFeGecfHeIeWgXgYgZg0g1g2g3g4g'h5g6g(h7g8gJeKedf9g!g#gLeMe$gNeOe%g'g(g)g*g5d+g,g)h/c*h+hefff,h-h.h/h:h6d:c;hgf7d=h?h@h-g8dPeQe.g/gReSehf:g;gMj=g?gTeUeifVeWe@g[g]g^g_g`g{g|g}g[h~gah]hbhchXeYejfdhehfhZe0egh1e2ehhihjhkhlh9dmhnh^h;c_h`hkflf{h|h}h~hai!d=cbimf#dcidieioh$d3e4ephqh5e6enfrhshNjthuhOjvhwhPjxhyhQjzhAhRjBhChSjDhEhTjFhGhUjjj4Ikj

61 plural_s = "" if len(cfg) == 1 else "s" 2Mh4Ikj

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

63 return cfg + (1,) * (3 - len(cfg)) 28 vbve9 ! wexeyewbzexbAeybzbAbBbCb# $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ DbEbFbGbHbIbBgCgDgFjvdEg%hwdFgxdGgydHgzdIgAd!bJgKgLgMgNge f OgJbqbKbPgLbMbNbObPbd g h i j abbbcbdbrbsbtbebQbRbSbTbUbBeVbfbk 7 RcScCeBd-c.cgbCdijcJdJeJGjMhfJHjTcDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdUdVdWdXdYdZd0d1d2d3d4dDeEeIjgJJjUc#b$bQgRgKjSgTgLjUgVgFeGecfHeIeWgXgYgZg0g1g2g3g4g'h5g6g(h7g8gJeKedf9g!g#gLeMe$gNeOe%g'g(g)g*g5d+g,g)h/c*h+hefff,h-h.h/h:h6d:c;hgf7d=h?h@h-g8dPeQe.g/gReSehf:g;gMj=g?gTeUeifVeWe@g[g]g^g_g`g{g|g}g[h~gah]hbhchXeYejfdhehfhZe0egh1e2ehhihjhkhlh9dmhnh^h;c_h`hkflf{h|h}h~hai!d=cbimf#dcidieioh$d3e4ephqh5e6enfrhshNjthuhOjvhwhPjxhyhQjzhAhRjBhChSjDhEhTjFhGhUjjj

64  

65  

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

67 if err != cydriver.CUresult.CUDA_SUCCESS: 2a kQoflQiJmQpfnQqfoQrfpQsfqQtfrQ8 sQvbtQveuQ9 vQ%dwQ! xQweyQzQjJAQxeBQyeCQwbDQzeEQkJFQxbGQAeHQlJIQybJQzbKQLQAbMQBbNQCbOQ# PQ$ QQ% RQ' SQ( TQ) UQ* VQ+ WQ, XQ- YQ. ZQ/ 0Q: 1Q; 2Q= 3Q? 4Q@ 5Q[ 6Q] 7Q^ 8Q_ 9Q` !Q{ #Q| $Q} %Q~ 'Q(Q)Q*Q+Q,Q-Q.Q/Q:Qfi;Q=Q?Q@Q[Q]Q^Q_Q`Q{Q|Q}Q~QaRbRcRdReRfRgRhRiRjRkRlRmRnRufoRpRqRrRvfsRwftRxfuRyfvRDbwRzfxRAfyRBfzRARBRCRDREbERFbFRGbGRHbHRIbIR?cJR@cKR[cLR]cMR^cNR_cOR`cPR{cQRRRSRTRURVRgiWRhiXRiiYRZR0R1R2R3R4RBg5RCg6R7R8R9RDg!RmJnJ#R$R%R'R(Rvd)RoJpJVjWj*R+R,R-R.R/R:REg;R=RqJ?RrJ@R[R]R^R_R`R{R|R}R%hwd~RaSbSsJcStJdSXjeSYjfSgShSiSjSkSlSmSnSoSpSqSrSsStSuSvSwSxSySzSASBSCSDSESFgFSuJvJGSHSISJSKSxdLSwJxJZj0jMSNSOSPSQSRSSSGgTSUSVSyJWSzJXSYSZS0S1S2S3S4S5S6S7Syd8S9S!SAJ#SBJ$S1j%S2j'S(S)S*S+S,S-S.S/SHg:S;S=SCJ?SDJ@S[S]S^S_S`S{S|S}S~SaTzdbTcTdTEJeTFJfT3jgT4jhTiTjTkTlTmTnToTpTIgqTrTsTGJtTHJuTvTwTxTyTzTATBTCTDTETAdFTGTHTIJITJJJT5jKT6jLTMTNTOTPTQTRTSTTTUTVTWTXTYTZT0T1T2T3T4T5T6T7T8T9T!T#T$Tji%T'T(T)T*T+T,T-T.T/T:T;T=Tki?T@T[T]T^T_T`T{T|T}T~TaUbUcUdUlieUfUgUhUiUbGjU!bkUlUmUminUCfoUpUqUrUsUtUuUvUwUxUyUzUAUBUCUKJDUDfEUFUGUHUIUJUKULUMUJgNUKgOULgPUMgQU'dRU7eSUEfTUFfUUGfVUNgWUHfXUYUZU0U1U2U3Uni4Uoi5Upi6Uqi7Uri8Usi9Uti!Uui#ULJ$UMJ%UNJ'U|c(U(d)UOJ*UPJ+U8e,Ue -Uf .UcG/UOg:UJb;Uqb=UKb?UQJ@URJ[USJ]U]b^U^b_U_b`UTJ{U`b|U{b}U|b~U}baVUJbVVJcV~bdVeVfVPggVIfhViVjVkVJflVvimVnVoVpVqVLbrVMbsVNbtVObuVPbvVwVxVyVzVAVBVCVDVd EVwiFV9eGV5bHVxiIV!eJV#eKV$eLVKfMVLfNVMfOVg PVh QVi RVj SVabTVbbUVcbVVdbWVrbXVsbYVtbZVNf0VOf1VPf2VQf3VRf4VSf5V6V7V8VTf9VUf!Veb#VQb$VRb%VVf'VSb(VTb)VUb*VBe+VVb,Vfb-VWJXJWfYJ.VyiziAiBiCiDiEiFi/VGi:VZJNh;V0JXf1J2J%e=VYfZf0f3J4J1f5J6JOh?V@VHiIiJi2f7J3f8J9J!J#J$J%J'J(J)J*J+J'e[VLP,JMP-JdG.J)d/JNP:JOP;JPP=JQP?JeG@JfG[J]J^Jk _J7 `JKi]VLi^Vlj_Vmj`VMi{VNi|Vnj}Voj~V7jaWbWcWRP{JSP|J}J~JaKTPbKdWcKdKeKUPfKeWgKhKiKVPjKfWkKlKmKWPnKgWoKpKqKXPrKhWsKtKuKYPvKiWwKxKZPyK0PzK1PAKpj(e2PBK3PCK}cDK*dEKjWkWlW~cFKadGKWbmWXbnWbdHKcdIKHhJK4PKKIhLK5PMKYboWZbpW0bqW1brW2bsW3btWPhuWvWwWxWyWzWAWBWQh)e*eCWDWEWFWGWHWIWJWKWRhLWMWNWOWPWQWRWSWTWUWVWWWXWYWZW0W1W2W3W4W5W6W7W8W9W!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~XaYbYcYdYeYfYOigYhYiYPiShThQigGRchGjYNKkYOKlYPKmYQKnYRKoYSKpYTKqYScrYUKsYVKtYWKuYXKvYiGwYjGxYkGyYCezYYKZK0K1K2KAYBYCYDYEYqj3K4K5K6Krj7K8Ksj9K!KFYGYHYIY6IJYKY7IRi#K8ISiBd8jTi-clG.cmGLY9IUi!IVi#IWiMY$IXiNY%IYi$K'IZiOYPYQYRYSYTY(IUY)IVY*IWY+IXY,IYY-IZY.I0Y1Y2Y3Y4Y5Y6Y7Y8Y9Y!Y#YgbacbcCd$Y%Y'YTc(YDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdUdVdWdXdYdZd0d1d2d3d4dDeEeccdcecfcgchcicjckclcmcncocpcqcrcsctcucvcwcxcyczcAcBcCcDcEcFcGcnGHc)Y*Y+YUhIcJcVh,YWhXhYh%K'K-YZh0h1h2h0i1i.Y/Yhb:Y;Yib=Y?Y@Yjb[Y6b]Y^Y_Y`Yl {Y|Ym }Yn ~Yo aZp bZq cZr dZs eZt fZu gZhZv iZw jZx kZy lZmZz nZA oZpZqZrZsZkbtZuZlbvZwZxZmbyZ7bzZAZBZCZB DZEZC FZD GZE HZF IZG JZH KZI LZJ MZK NZOZL PZM QZN RZO SZTZP UZQ VZWZXZYZZZnb0Z1Zob2Z3Z4Zpb5Z8b6Z7Z8Z9ZR !Z#ZS $ZT %ZU 'ZV (ZW )ZX *ZY +ZZ ,Z0 -Z.Z1 /Z2 :Z3 ;Z4 =Z?Z5 @Z6 [Z]Z^Z3h(K)K*K_Zdd`Z{Z+K|Z,K}Z~Za0b0+eoGpGc0-Kd0.Ke0qGf0rGg0sGh0tGi0Vcj0,ek0edl0-em0n0/Ko0:Kp0;Kq0=Kr0?Ks0@Kt0[Ku0]Kv0^Kw0_Kx0`Ky0{Kz0|KA0}KB0~KC0aLD0uGE0bLF0vGG0cLH0wGI0dLJ0xGK0eLL0yGM0fLN0zGO0gLP0AGQ0hLR0BGS0iLT0jLU0kLV0lLW0mLX0nLY0oLZ0pL00qL10rL20sL30tL40uL50vL60wL70xL80yL90/I!09j#0:I$0;I%0!j'0=I(0.e)0*04h5h/ezLALBL+0fd,0:e;eCL-0DL.0EL/0FL:0CG;0DG=0WcXcubYc?0Zc@0[0]0^0_0`0{0|0}0~0a1b1c1d1Kce1EGf1FG+dg1GGHGh1IGi1JGj1KGUck1LG4bMGl1NGm1OGn1PGo1QGp1RGq1SGr1s1TGt1UGu1,dv1VGWGXGYG0c1c2cZG0G1G3c4c5c#bw1$bx1y12Gz13GA14GB15GC16GD17GE18GF19GG1!GH1#GI1$GJ1%GK1'GL1(GM1)GN1*GO1+GP1,GQ1-GR1.GS1/GT1:GU1;GV1=G?G@GW1[GX1]GY1^GZ1_G01`G11{G21|G31}G41~G51aH61bH71cH81dH91eH!1fH#1gH$1hH%1iH'1jH(1kH)1lHmH*1nH+1oH,1pH-1qH.1rH/1sH:1tH;1uH=1vH?1wH@1xH[1yH]1zH^1AH_1BH`1CHDHJh{1=e|1?e}1gd~1-da2b2GLQgc2d2e2f2HLILJLg2h2i2j2k2EHRgl2KLFHm2LLn2o2MLGHNLp2OLSgq2r2s2t2PLQLRLu2v2w2x2y2HHTgz2SLIHA2TLB2C2UgULVLD2E2F2G2VgH2WLI2J2XLFeYLZL?I@I0L1L2L3L4L#jGecf5L$j%j6L7L8LHeIeWgXgYgZg0gK21gL29L!LM2N2O2P22gQ2#LR2S2T2U2$L3gV2W2%L'LX2Y2Z20212JH4g'h22(LKH32425262)L5g728292*L+L,L!2#2$2%2'2LH6g(h(2-LMH)2.L*2+2,27g/L:L-2.2/2:28g;2;L=2?2=L[IJe?L@L[L]L^L_L`L{L|L}LNHKedf~L'jOHaMbMcMdMeMfMgMhMiMjMkMlMmMnM9goMpM@2[2]2^2qM_2`2{2!g|2}2rMsM~2a3b3c3#gd3tMe3f3g3h3i3j3k3l3m3n3o3p3PHLeQHRHSHTHUHVHWHXHYHMeZH2i0H1H2H$gNeOe%g'g(g)g*g5d+g,g)h/c*h+hefff,h-h.h/h:h6d:c;hgf7d=h?h@h-guMvMq3r3s3t3u38d(j)jwMxMyMzMAMBM3HPe4H5H6H7H8H9H!H#H$H%HQe'H3i(H)H*Hv3.gw3CMDMx3y3z3A3/gB3EMC3D3FM]IReGMHMIMJMKMLMMMNMOMPMQM+HSehfRM*j,HSMTMUME3F3VM:gG3H3I3J3WMXMYMK3L3M3N3O3-H;gP3ZM.HQ30MR3S3=g1M2MT3U3V3W3?gX33MY3Z34MTe5M6M^I_I7M8M9M!M#M+jUeif$M,j-j%M'M(MVeWe@g[g]g^g_g03`g13)M*M23334353{g63+M738393!3,M|g#3$3-M.M%3'3(3)3*3/H}g[h+3/M:H,3-3.3/3:M~g:3;3=3;M=M?M?3@3[3]3^3;Hah]h_3@M=H`3[M{3|3}3bh]M^M~3a4b4c4chd4_Me4f4`M`IXe{M|M}M~MaNbNcNdNeNfN?HYejfgN.j@HhNiNjNkNlNmNnNoNpNqNrNsNtNuNdhvNwNg4h4i4j4xNk4l4m4ehn4o4yNzNp4q4r4s4fht4ANu4v4w4x4y4z4A4B4C4D4E4F4[HZe]H^H_H`H{H|H}H~HaI0ebI4icIdIeIgh1e2ehhihjhkhlh9dmhnh^h;c_h`hkflf{h|h}h~hai!d=cbimf#dcidieiohBNCNG4H4I4J4K4$d/j:jDNENFNGNHNINfI3egIhIiIjIkIlImInIoIpI4eqI5irIsItIL4phM4JNKNN4O4P4Q4qhR4LNS4T4MN{I5eNNONPNQNRNSNTNUNVNWNXNuI6enfYN;jvIZN0N1NU4V4W4@e[e=j?jwIxIX42NrhY4Z404143N4N5N2434445464yIsh746NzIAIKhBICILh847N94!4#48Nth$4%4'4(49N!N#N)4*4+4,4-4DIuh.4$NEI/4%N:4;4=4'Nvh?4@4[4]4(N)N*N^4_4`4{4|4FIwh}4+NGI~4,Na5b5c5-Nxhd5e5f5g5.N/N:Nh5i5j5k5l5HIyhm5;NIIn5=No5p5q5?Nzhr5s5t5u5@N[N]Nv5w5x5y5z5JIAhA5^NKIB5_NC5D5E5`NBhF5G5H5I5{N|N}NJ5K5L5M5N5LIChO5~NMIP5aOQ5R5S5bODhT5U5V5W5cOdOeOX5Y5Z50515NIEh25fOOIPI35gO455565hOFh758595!5iOjOkO#5$5%5'5(5QIGh)5lORISI*5mO+5,56c-57c.5|I/58c:5;5=5?5@5[5]5^5_5`5{5|5}5~5a6b6c6d6e6f6g6h6i6j6k6l6m6n6o6p6q6r6s6t6u6v6w6x6c y69cz6TI9bUI.dA6B6C6D6E6F6G6H6I6J6K6L6M6N6@j[jnOVIWIXIO6YIoOpOZI}ItjP6Q6R6S6T6U6V6W6!cX6Y6#cZ606162636465666768696!6#6qO$6%6'6(6)6*6+6,6-6.60I/6:6rO;6sO=6?6tO@6[6]6^6_6`6{6|6}6~6a7b7c7d7e7f7g7h7i7j7k7l7m7n7o7p7q7r7s7t7u7v7w7x7y7z7A7B7C7D7E7F7G7H7I7J7]jK7^jL7_jM7`jN7{jO7|jP7}jQ7~jR7akS7bkT7ckU7dkV7,bW7ekX76hY76iZ77i071727374757678i77fk87gk97hk!7ik#79i$7%7'7(7)7uO*7vO+7wO,7xO-7yO.7zO/7jk:7AO;7BO=7CO?7DO@7EO[7kk]7lk^7FO_7mk`7nk{7ok|7!i}7pk~7#ia8qkb8rkc8d8e84ff8g8skh8i8tkj8$ik8l8m8n8o8p8q8r8s8t8u8v8w8x8y8z8A8B8C8D8E8F8G8H8I8J8K8L8M8GON8HOO8IOJOKOLOP8MONOQ8OOR8POQOROSOTOUOS8T8VOU8WOXOV8YOW8ZO0O1O2O3O4OX85O6OY8Z808187O28388O485868]e788898!8#8$8%8'8(8)8*8+8,8-8.8/8:8;8=8?8@8[8]8^8_8`8{8|8}8~8a9b9c9d9e9f9g9

68 return _check_driver_error(err) 2fikilimiwi9e5bxiqjrjsj4b.d

69 return 0 2a kQoflQiJmQpfnQqfoQrfpQsfqQtfrQ8 sQvbtQveuQ9 vQ%dwQ! xQweyQzQjJAQxeBQyeCQwbDQzeEQkJFQxbGQAeHQlJIQybJQzbKQLQAbMQBbNQCbOQ# PQ$ QQ% RQ' SQ( TQ) UQ* VQ+ WQ, XQ- YQ. ZQ/ 0Q: 1Q; 2Q= 3Q? 4Q@ 5Q[ 6Q] 7Q^ 8Q_ 9Q` !Q{ #Q| $Q} %Q~ 'Q(Q)Q*Q+Q,Q-Q.Q/Q:Qfi;Q=Q?Q@Q[Q]Q^Q_Q`Q{Q|Q}Q~QaRbRcRdReRfRgRhRiRjRkRlRmRnRufoRpRqRrRvfsRwftRxfuRyfvRDbwRzfxRAfyRBfzRARBRCRDREbERFbFRGbGRHbHRIbIR?cJR@cKR[cLR]cMR^cNR_cOR`cPR{cQRRRSRTRURVRgiWRhiXRiiYRZR0R1R2R3R4RBg5RCg6R7R8R9RDg!RmJnJ#R$R%R'R(Rvd)RoJpJVjWj*R+R,R-R.R/R:REg;R=RqJ?RrJ@R[R]R^R_R`R{R|R}R%hwd~RaSbSsJcStJdSXjeSYjfSgShSiSjSkSlSmSnSoSpSqSrSsStSuSvSwSxSySzSASBSCSDSESFgFSuJvJGSHSISJSKSxdLSwJxJZj0jMSNSOSPSQSRSSSGgTSUSVSyJWSzJXSYSZS0S1S2S3S4S5S6S7Syd8S9S!SAJ#SBJ$S1j%S2j'S(S)S*S+S,S-S.S/SHg:S;S=SCJ?SDJ@S[S]S^S_S`S{S|S}S~SaTzdbTcTdTEJeTFJfT3jgT4jhTiTjTkTlTmTnToTpTIgqTrTsTGJtTHJuTvTwTxTyTzTATBTCTDTETAdFTGTHTIJITJJJT5jKT6jLTMTNTOTPTQTRTSTTTUTVTWTXTYTZT0T1T2T3T4T5T6T7T8T9T!T#T$Tji%T'T(T)T*T+T,T-T.T/T:T;T=Tki?T@T[T]T^T_T`T{T|T}T~TaUbUcUdUlieUfUgUhUiUbGjU!bkUlUmUminUCfoUpUqUrUsUtUuUvUwUxUyUzUAUBUCUKJDUDfEUFUGUHUIUJUKULUMUJgNUKgOULgPUMgQU'dRU7eSUEfTUFfUUGfVUNgWUHfXUYUZU0U1U2U3Uni4Uoi5Upi6Uqi7Uri8Usi9Uti!Uui#ULJ$UMJ%UNJ'U|c(U(d)UOJ*UPJ+U8e,Ue -Uf .UcG/UOg:UJb;Uqb=UKb?UQJ@URJ[USJ]U]b^U^b_U_b`UTJ{U`b|U{b}U|b~U}baVUJbVVJcV~bdVeVfVPggVIfhViVjVkVJflVvimVnVoVpVqVLbrVMbsVNbtVObuVPbvVwVxVyVzVAVBVCVDVd EVwiFV9eGV5bHVxiIV!eJV#eKV$eLVKfMVLfNVMfOVg PVh QVi RVj SVabTVbbUVcbVVdbWVrbXVsbYVtbZVNf0VOf1VPf2VQf3VRf4VSf5V6V7V8VTf9VUf!Veb#VQb$VRb%VVf'VSb(VTb)VUb*VBe+VVb,Vfb-VWJXJWfYJ.VyiziAiBiCiDiEiFi/VGi:VZJNh;V0JXf1J2J%e=VYfZf0f3J4J1f5J6JOh?V@VHiIiJi2f7J3f8J9J!J#J$J%J'J(J)J*J+J'e[VLP,JMP-JdG.J)d/JNP:JOP;JPP=JQP?JeG@JfG[J]J^Jk _J7 `JKi]VLi^Vlj_Vmj`VMi{VNi|Vnj}Voj~V7jaWbWcWRP{JSP|J}J~JaKTPbKdWcKdKeKUPfKeWgKhKiKVPjKfWkKlKmKWPnKgWoKpKqKXPrKhWsKtKuKYPvKiWwKxKZPyK0PzK1PAKpj(e2PBK3PCK}cDK*dEKjWkWlW~cFKadGKWbmWXbnWbdHKcdIKHhJK4PKKIhLK5PMKYboWZbpW0bqW1brW2bsW3btWPhuWvWwWxWyWzWAWBWQh)e*eCWDWEWFWGWHWIWJWKWRhLWMWNWOWPWQWRWSWTWUWVWWWXWYWZW0W1W2W3W4W5W6W7W8W9W!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~XaYbYcYdYeYfYOigYhYiYPiShThQigGRchGjYNKkYOKlYPKmYQKnYRKoYSKpYTKqYScrYUKsYVKtYWKuYXKvYiGwYjGxYkGyYCezYYKZK0K1K2KAYBYCYDYEY3K4K5K6K7K8K9K!KFYGYHYIY6IJYKY7IRi#K8ISiBd8jTi-clG.cmGLY9IUi!IVi#IWiMY$IXiNY%IYi$K'IZiOYPYQYRYSYTY(IUY)IVY*IWY+IXY,IYY-IZY.I0Y1Y2Y3Y4Y5Y6Y7Y8Y9Y!Y#YgbacbcCd$Y%Y'YTc(YDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdUdVdWdXdYdZd0d1d2d3d4dDeEeccdcecfcgchcicjckclcmcncocpcqcrcsctcucvcwcxcyczcAcBcCcDcEcFcGcnGHc)Y*Y+YUhIcJcVh,YWhXhYh%K'K-YZh0h1h2h0i1i.Y/Yhb:Y;Yib=Y?Y@Yjb[Y6b]Y^Y_Y`Yl {Y|Ym }Yn ~Yo aZp bZq cZr dZs eZt fZu gZhZv iZw jZx kZy lZmZz nZA oZpZqZrZsZkbtZuZlbvZwZxZmbyZ7bzZAZBZCZB DZEZC FZD GZE HZF IZG JZH KZI LZJ MZK NZOZL PZM QZN RZO SZTZP UZQ VZWZXZYZZZnb0Z1Zob2Z3Z4Zpb5Z8b6Z7Z8Z9ZR !Z#ZS $ZT %ZU 'ZV (ZW )ZX *ZY +ZZ ,Z0 -Z.Z1 /Z2 :Z3 ;Z4 =Z?Z5 @Z6 [Z]Z^Z3h(K)K*K_Zdd`Z{Z+K|Z,K}Z~Za0b0+eoGpGc0-Kd0.Ke0qGf0rGg0sGh0tGi0Vcj0,ek0edl0-em0n0/Ko0:Kp0;Kq0=Kr0?Ks0@Kt0[Ku0]Kv0^Kw0_Kx0`Ky0{Kz0|KA0}KB0~KC0aLD0uGE0bLF0vGG0cLH0wGI0dLJ0xGK0eLL0yGM0fLN0zGO0gLP0AGQ0hLR0BGS0iLT0jLU0kLV0lLW0mLX0nLY0oLZ0pL00qL10rL20sL30tL40uL50vL60wL70xL80yL90/I!09j#0:I$0;I%0!j'0=I(0.e)0*04h5h/ezLALBL+0fd,0:e;eCL-0DL.0EL/0FL:0CG;0DG=0WcXcubYc?0Zc@0[0]0^0_0`0{0|0}0~0a1b1c1d1Kce1EGf1FG+dg1GGHGh1IGi1JGj1KGUck1LG4bMGl1NGm1OGn1PGo1QGp1RGq1SGr1s1TGt1UGu1,dv1VGWGXGYG0c1c2cZG0G1G3c4c5c#bw1$bx1y12Gz13GA14GB15GC16GD17GE18GF19GG1!GH1#GI1$GJ1%GK1'GL1(GM1)GN1*GO1+GP1,GQ1-GR1.GS1/GT1:GU1;GV1=G?G@GW1[GX1]GY1^GZ1_G01`G11{G21|G31}G41~G51aH61bH71cH81dH91eH!1fH#1gH$1hH%1iH'1jH(1kH)1lHmH*1nH+1oH,1pH-1qH.1rH/1sH:1tH;1uH=1vH?1wH@1xH[1yH]1zH^1AH_1BH`1CHDHJh{1=e|1?e}1gd~1-da2b2GLQgc2d2e2f2HLILJLg2h2i2j2k2EHRgl2KLFHm2LLn2o2MLGHNLp2OLSgq2r2s2t2PLQLRLu2v2w2x2y2HHTgz2SLIHA2TLB2C2UgULVLD2E2F2G2VgH2WLI2J2XLFeYLZL?I@I0L1L2L3L4L#jGecf5L$j%j6L7L8LHeIeWgXgYgZg0gK21gL29L!LM2N2O2P22gQ2#LR2S2T2U2$L3gV2W2%L'LX2Y2Z20212JH4g'h22(LKH32425262)L5g728292*L+L,L!2#2$2%2'2LH6g(h(2-LMH)2.L*2+2,27g/L:L-2.2/2:28g;2;L=2?2=L[IJe?L@L[L]L^L_L`L{L|L}LNHKedf~L'jOHaMbMcMdMeMfMgMhMiMjMkMlMmMnM9goMpM@2[2]2^2qM_2`2{2!g|2}2rMsM~2a3b3c3#gd3tMe3f3g3h3i3j3k3l3m3n3o3p3PHLeQHRHSHTHUHVHWHXHYHMeZH2i0H1H2H$gNeOe%g'g(g)g*g5d+g,g)h/c*h+hefff,h-h.h/h:h6d:c;hgf7d=h?h@h-guMvMq3r3s3t3u38d(j)jwMxMyMzMAMBM3HPe4H5H6H7H8H9H!H#H$H%HQe'H3i(H)H*Hv3.gw3CMDMx3y3z3A3/gB3EMC3D3FM]IReGMHMIMJMKMLMMMNMOMPMQM+HSehfRM*j,HSMTMUME3F3VM:gG3H3I3J3WMXMYMK3L3M3N3O3-H;gP3ZM.HQ30MR3S3=g1M2MT3U3V3W3?gX33MY3Z34MTe5M6M^I_I7M8M9M!M#M+jUeif$M,j-j%M'M(MVeWe@g[g]g^g_g03`g13)M*M23334353{g63+M738393!3,M|g#3$3-M.M%3'3(3)3*3/H}g[h+3/M:H,3-3.3/3:M~g:3;3=3;M=M?M?3@3[3]3^3;Hah]h_3@M=H`3[M{3|3}3bh]M^M~3a4b4c4chd4_Me4f4`M`IXe{M|M}M~MaNbNcNdNeNfN?HYejfgN.j@HhNiNjNkNlNmNnNoNpNqNrNsNtNuNdhvNwNg4h4i4j4xNk4l4m4ehn4o4yNzNp4q4r4s4fht4ANu4v4w4x4y4z4A4B4C4D4E4F4[HZe]H^H_H`H{H|H}H~HaI0ebI4icIdIeIgh1e2ehhihjhkhlh9dmhnh^h;c_h`hkflf{h|h}h~hai!d=cbimf#dcidieiohBNCNG4H4I4J4K4$d/j:jDNENFNGNHNINfI3egIhIiIjIkIlImInIoIpI4eqI5irIsItIL4phM4JNKNN4O4P4Q4qhR4LNS4T4MN{I5eNNONPNQNRNSNTNUNVNWNXNuI6enfYN;jvIZN0N1NU4V4W4@e[e=j?jwIxIX42NrhY4Z404143N4N5N2434445464yIsh746NzIAIKhBICILh847N94!4#48Nth$4%4'4(49N!N#N)4*4+4,4-4DIuh.4$NEI/4%N:4;4=4'Nvh?4@4[4]4(N)N*N^4_4`4{4|4FIwh}4+NGI~4,Na5b5c5-Nxhd5e5f5g5.N/N:Nh5i5j5k5l5HIyhm5;NIIn5=No5p5q5?Nzhr5s5t5u5@N[N]Nv5w5x5y5z5JIAhA5^NKIB5_NC5D5E5`NBhF5G5H5I5{N|N}NJ5K5L5M5N5LIChO5~NMIP5aOQ5R5S5bODhT5U5V5W5cOdOeOX5Y5Z50515NIEh25fOOIPI35gO455565hOFh758595!5iOjOkO#5$5%5'5(5QIGh)5lORISI*5mO+5,56c-57c.5|I/58c:5;5=5?5@5[5]5^5_5`5{5|5}5~5a6b6c6d6e6f6g6h6i6j6k6l6m6n6o6p6q6r6s6t6u6v6w6x6c y69cz6TI9bUI.dA6B6C6D6E6F6G6H6I6J6K6L6M6N6@j[jnOVIWIXIO6YIoOpOZI}ItjP6Q6R6S6T6U6V6W6!cX6Y6#cZ606162636465666768696!6#6qO$6%6'6(6)6*6+6,6-6.60I/6:6rO;6sO=6?6tO@6[6]6^6_6`6{6|6}6~6a7b7c7d7e7f7g7h7i7j7k7l7m7n7o7p7q7r7s7t7u7v7w7x7y7z7A7B7C7D7E7F7G7H7I7J7]jK7^jL7_jM7`jN7{jO7|jP7}jQ7~jR7akS7bkT7ckU7dkV7,bW7ekX76hY76iZ77i071727374757678i77fk87gk97hk!7ik#79i$7%7'7(7)7uO*7vO+7wO,7xO-7yO.7zO/7jk:7AO;7BO=7CO?7DO@7EO[7kk]7lk^7FO_7mk`7nk{7ok|7!i}7pk~7#ia8qkb8rkc8d8e84ff8g8skh8i8tkj8$ik8l8m8n8o8p8q8r8s8t8u8v8w8x8y8z8A8B8C8D8E8F8G8H8I8J8K8L8M8GON8HOO8IOJOKOLOP8MONOQ8OOR8POQOROSOTOUOS8T8VOU8WOXOV8YOW8ZO0O1O2O3O4OX85O6OY8Z808187O28388O485868]e788898!8#8$8%8'8(8)8*8+8,8-8.8/8:8;8=8?8@8[8]8^8_8`8{8|8}8~8a9b9c9d9e9f9g9

70  

71  

72cdef int _get_current_device_id() except? -1: 

73 """Return the current thread's bound CUdevice ordinal.""" 

74 cdef cydriver.CUdevice dev 

75 with nogil: 2]j^j_j`j{j|j}j~jakbkckdk,bek6h6i7i8ifkgkhkik9ijkkklkmknkok!ipk#iqkrk4fsktk

76 HANDLE_RETURN(cydriver.cuCtxGetDevice(&dev)) 2]j^j_j`j{j|j}j~jakbkckdk,bek6h6i7i8ifkgkhkik9ijkkklkmknkok!ipk#iqkrk4fsktk

77 return <int>dev 2]j^j_j`j{j|j}j~jakbkckdk,bek6h6i7i8ifkgkhkik9ijkkklkmknkok!ipk#iqkrk4fsktk

78  

79  

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

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

82 if err == cynvrtc.nvrtcResult.NVRTC_SUCCESS: 28 vbve9 ! wexeyewbzexbAeybzbAbBbCb# $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ DbEbFbGbHbIbBgCgDgvdEgwdFgxdGgydHgzdIgAd!bJgKgLgMgNge f cGOgJbqbKbPgLbMbNbObPbd g h i j abbbcbdbrbsbtbebQbRbSbTbUbBeVbfbk 7 RcScCe8jlGmGgbCdTcDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdUdVdWdXdYdZd0d1d2d3d4dDeEeb ccdcecfcgchcicjckclcmcncocpcqcrcsctcucvcwcxcyczcAcBcCcDcEcFcGcnGHcKcEGFG+dGGHGIGJGKGUcLG4bMGNGOGPGQGRGSG9OTGUG,dVGWGXGYG0c1c2cZG0G1G3c4c5c#b$b2G3G4G5G6G7G8G9G!G#G$G%G'G(G)G*G+G,G-G.G/G:G;G=G?G@G[G]G^G_G`G{G|G}G~GaHbHcHdHeHfHgHhHiHjHkHlHmHnHoHpHqHrHsHtHuHvHwHxHyHzHAHBHCHDHQgEHRgFHGHSgHHTgIHUgVgFe#jGe%jHeIeWgXgYgZg0g1g2g3gJH4gKH5gLH6gMH7g8gJeNHKeOH9g!g#gPHLeQHRHSHTHUHVHWHXHYHMeZH2i0H1H2H$gNeOe%g'g(g)g*g5d+g,g/c6d:c7d-g8d3HPe4H5H6H7H8H9H!H#H$H%HQe'H3i(H)H*H.g/gRe+HSe,H:g-H;g.H=g?gTe+jUe-jVeWe@g[g]g^g_g`g{g|g/H}g:H~g;Hah=HbhchXe?HYe@Hdhehfh[HZe]H^H_H`H{H|H}H~HaI0ebI4icIdIeIgh1e2ehhihjhkhlh9dmhnh;c!d=c#doh$dfI3egIhIiIjIkIlImInIoIpI4eqI5irIsItIphqh5euI6evIwIxIrhyIshzIAIBIthDIuhEIvhFIwhGIxhHIyhIIzhJIAhKIBhLIChMIDhNIEhOIPIFhQIGhRISI6c7c/d8c:d;d=dhdidjd?d@d[d]d$c%ckd^dld_d`d{d|dmd}dndodpdqd~daebecedeeefegeherdsdtdc 9cVIWIXIudYIZI!c#ciejekelemeneoe!O#O$O%Ope'Oqere(O

83 return 0 28 vbve9 ! wexeyewbzexbAeybzbAbBbCb# $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ DbEbFbGbHbIbBgCgDgvdEgwdFgxdGgydHgzdIgAd!bJgKgLgMgNge f cGOgJbqbKbPgLbMbNbObPbd g h i j abbbcbdbrbsbtbebQbRbSbTbUbBeVbfbk 7 RcScCe8jlGmGgbCdTcDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdUdVdWdXdYdZd0d1d2d3d4dDeEeb ccdcecfcgchcicjckclcmcncocpcqcrcsctcucvcwcxcyczcAcBcCcDcEcFcGcnGHcKcEGFG+dGGHGIGJGKGUcLG4bMGNGOGPGQGRGSG9OTGUG,dVGWGXGYG0c1c2cZG0G1G3c4c5c#b$b2G3G4G5G6G7G8G9G!G#G$G%G'G(G)G*G+G,G-G.G/G:G;G=G?G@G[G]G^G_G`G{G|G}G~GaHbHcHdHeHfHgHhHiHjHkHlHmHnHoHpHqHrHsHtHuHvHwHxHyHzHAHBHCHDHQgEHRgFHGHSgHHTgIHUgVgFe#jGe%jHeIeWgXgYgZg0g1g2g3gJH4gKH5gLH6gMH7g8gJeNHKeOH9g!g#gPHLeQHRHSHTHUHVHWHXHYHMeZH2i0H1H2H$gNeOe%g'g(g)g*g5d+g,g/c6d:c7d-g8d3HPe4H5H6H7H8H9H!H#H$H%HQe'H3i(H)H*H.g/gRe+HSe,H:g-H;g.H=g?gTe+jUe-jVeWe@g[g]g^g_g`g{g|g/H}g:H~g;Hah=HbhchXe?HYe@Hdhehfh[HZe]H^H_H`H{H|H}H~HaI0ebI4icIdIeIgh1e2ehhihjhkhlh9dmhnh;c!d=c#doh$dfI3egIhIiIjIkIlImInIoIpI4eqI5irIsItIphqh5euI6evIwIxIrhyIshzIAIBIthDIuhEIvhFIwhGIxhHIyhIIzhJIAhKIBhLIChMIDhNIEhOIPIFhQIGhRISI6c7c/d8c:d;d=dhdidjd?d@d[d]d$c%ckd^dld_d`d{d|dmd}dndodpdqd~daebecedeeefegeherdsdtdc 9cVIWIXIudYIZI!c#ciejekelemeneoe!O#O$O%Ope'Oqere(O

84 with gil: 1c

85 _raise_nvrtc_error(prog, err) 1c

86  

87  

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

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

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

91 cdef size_t logsize = 0 1c

92 if prog != NULL: 1c

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

94 cdef bytes log_bytes 

95 cdef str log_str = "" 1c

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

97 log_bytes = b" " * logsize 1c

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

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

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

101 if log_str: 1c

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

103 raise NVRTCError(err_msg) 1c

104  

105  

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

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

108 if err == cynvvm.nvvmResult.NVVM_SUCCESS: 2CI)O*OTI9bUI.d+O,O-O.O/O:O;O=O?O@O[O]O^O_O@j[j

109 return 0 2CI)O*OTI9bUI.d+O,O-O.O/O:O;O=O?O@O[O]O^O_O@j[j

110 with gil: 29b

111 _raise_nvvm_error(prog, err) 29b

112  

113  

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

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

116 cdef size_t logsize = 0 29b

117 if prog != NULL: 29b

118 cynvvm.nvvmGetProgramLogSize(prog, &logsize) 29b

119 cdef bytes log_bytes 

120 cdef str log_str = "" 29b

121 if logsize > 1 and prog != NULL: 29b

122 log_bytes = b" " * logsize 29b

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

124 log_str = log_bytes.decode("utf-8", errors="backslashreplace") 29b

125 cdef object exc = nvvmError(err) 29b

126 if log_str: 29b

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

128 raise exc 29b

129  

130  

131cdef int HANDLE_RETURN_NVJITLINK( 

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

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

134 if err == cynvjitlink.nvJitLinkResult.NVJITLINK_SUCCESS: 2b `O{O|O}O~OaPbPcPdPePfP~IaJgPhPiPjPkPlPmPnPoPpPqPrPsP5ftPuPvPwPxP#b$btjyPzPAPBPCPDPEPFPGPHP

135 return 0 2b `O{O|O}O~OaPbPcPdPePfP~IaJgPhPiPjPkPlPmPnPoPpPqPrPsPtPuPvPwPxP#b$btjyPzPAPBPCPDPEPFPGPHP

136 with gil: 2b 5f

137 _raise_nvjitlink_error(handle, err) 2b 5f

138  

139  

140cdef int _raise_nvjitlink_error( 

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

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

143 cdef size_t logsize = 0 2b 5f

144 if handle != NULL: 2b 5f

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

146 cdef bytes log_bytes 

147 cdef str log_str = "" 2b 5f

148 if logsize > 1 and handle != NULL: 2b 5f

149 log_bytes = b" " * logsize 1b

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

151 cynvjitlink.nvJitLinkResult.NVJITLINK_SUCCESS: 1b

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

153 cdef object exc = nvJitLinkError(err) 2b 5f

154 if log_str: 2b 5f

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

156 raise exc 2b 5f

157  

158  

159cdef object _RUNTIME_SUCCESS = runtime.cudaError_t.cudaSuccess 

160cdef object _NVRTC_SUCCESS = nvrtc.nvrtcResult.NVRTC_SUCCESS 

161  

162  

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

164 if error == cydriver.CUresult.CUDA_SUCCESS: 2a ofukvkpfwkqfxkrfyksfzktfAk8 BkvbCkDk9 Ek%dFk! GkHkIkJkKkLkwbMkNkOkxbPkQkRkybSkzbTkUkAbVkBbWkCbXk# Yk$ Zk% 0k' 1k( 2k) 3k* 4k+ 5k, 6k- 7k. 8k/ 9k: !k; #k= $k? %k@ 'k[ (k] )k^ *k_ +k` ,k{ -k| .k} /k~ :k;k=k?k@k[kfi]k^k_k`k{k|k}k~kalblcldlelflglufhliljlvfklwfllxfmlyfnlDbolzfplAfqlBfrlsltlEbulFbvlGbwlHbxlIbyl?czl@cAl[cBl]cCl^cDl_cEl`cFl{cGlHlIlJlKlLlMlNlgiOlhiPliiQlRlSlTlUlVlWlXlYlZl0l1l2l3l4l5l6l7l8l9l!l#l$l%l'l(l)l*l+l,l-l.l/l:l;l=l?l@l[l]l^l_l`l{l|l}l~lambmcmdmemfmgmhmimjmkmlmmmnmompmqmrmsmtmumvmwmxmymzmAmBmCmDmEmFmGmHmImJmKmLmMmNmOmPmQmRmSmTmUmVmWmXmYmZm0m1m2m3m4m5m6m7m8m9m!m#m$m%m'm(m)m*m+m,m-m.m/m:m;m=m?m@m[m]m^m_m`m{m|m}m~manbncndnenfngnhninjijnknlnmnnnonpnkiqnrnsntnunvnwnxnliynznAnBn!bCnDnmiEnCfFnGnHnInJnKnLnMnNnDfOnPnQnRnSnTnUnVnWn'dXn7eYnEfZnFf0nGf1n2nHf3n4n5n6nni7noi8npi9nqi!nri#nsi$nti%nui'n(n)n*n|c+n(d,n-n.n8e/ne :nf ;n=n?nJb@nqb[nKb]n^n_n`n]b{n^b|n_b}n~n`bao{bbo|bco}bdoeofo~bgohoioIfjokoloJfmovinooopoLbqoMbroNbsoObtoPbuovowoxoyod zowiAo9eBo5bCoxiDo!eEo#eFo$eGoKfHoLfIoMfJog Koh Loi Moj NoabOobbPocbQodbRorbSosbTotbUoNfVoOfWoPfXoQfYoRfZoSf0o1o2oTf3oUf4oeb5oQb6oRb7oVf8oSb9oTb!oUb#o$oVb%ofb'o(oWf)o*oyiziujAiBivjCiDiwjEiFixjGiyj+oNh,oXfzj-o%e.oYfZfAj0f/o1fBj:oOh;oHiCjIiJiDj2f=o3f?o@o[o]o^o_o'e`o6f7f8f)d9f!f#f$f%f'f(f)f*fk +f7 ,f{o|o}o~oapbpcpdpepfpgp-f.f/f:f;f=f?f@f[f]f^f_f`f{f|f}f~fagbgcgdgegfggghgigjgkglgmgng(eogpgqg}crg*dsghpipjp~ctgadugWbkpXblpbdvgcdwgHhxgygIhzgAgYbmpZbnp0bop1bpp2bqp3brpsptpupvpwpxpypzp1IEjApBpCpDpEpFpGpHpIpJpKpLpMpNpOiOpPpQpPiRpQiRcSpTpUpVpWpXpYpZpSc0p1p2p3p4p5p6p7p8pqjrjsj9pRiSiBdTi-c!p.c#pUiViWi$pXiYiZi%p'p(p)p*p+p,p-p.p/p:p;p=p?p@p[p]p^p_p`p{p|pgbacbc}p~paqbqDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdUdVdWdXdYdZd0d1d2d3d4dDeEecqdqeqfqgqhqiqjqkqlqmqnqoqpqqqrqsqtquqvqwqxqyqzqAqBqCqDqEqFqGqHqIqJqKqLqUhIcJcVhWhXhYhZh0h1h2h0i1ise'chbLcib(cMcjb6b^e_e`el -bm n o p q r s t u .bv w x y %bz A 'b/bte)ckbNclb*cOcmb7b{e|e}eB :bC D E F G H I J K ;bL M N O (bP Q )b=bue+cnbPcob,cQcpb8b~eafbfR ?bS T U V W X Y Z 0 @b1 2 3 4 *b5 6 +b[b3hMqNqOqPqQqRqSq+eTqUqVqWqXqYqZq0q1q2q3q4q5q6q7q8q9q!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~qarbrcrdrerfrgrhrirjrkrlrmrnrorprqrrrsrtrurvr4h5h/ewrxryrzrArBrCGCrDGDrWcXcubYcErZcFrGrHrIrJrKrLrMrNrOrPrQrRrSrTrUr+dVrWrXrYrZr0rUc1r2r4b3r4r5r6r7r8r9r!r#r$r,d%r'r(r)r*r+r,r-r.r/r:r;r=r?r#b@r$b[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~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:B;B=B?B@B[B]B^B_B`B{B|B}B~BaCbCcCdCeCfCgChCiCjCkClCmCnCoCpCqCrCsCtCuCvCwCxCyCzCACBCCCDCECFCGCHCICJCKCLCMCNCOCPCQCRCSCTCUCVCWCXCYCZC0C1C2C3C4C5C6C7C8C9C!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~CaDbDcDdDeDfDgDhD.diDjDkDlDmDnDoDpDqDrDsDtDuDvDwDxDyDzDADBDCDDDEDFDGDHDIDJDKDLDMDNDODPDQDRDSDTDUDVDWDXDYDZD0D1D2D3D4D5D6D7D8D9D!D#D$D%D'D(D)D*D+D,D-D.D/D:D;D=D?D@D[D]D^D_D`D{D|D}D~DaEbEcEdEeEfEgEhEiEjEkElEmEnEoEpEqErEsEtEuEvEwExEyEzEAEBECEDEEEFEGEHE,bIEJEKELEMENEOEPEQERESETEUEVEWEXEYEZE0E1E2E3E4E5E6E7E8E9E!E#E$E%E'E(E)E*E+E,E-E.E/E:E;E=E?E@E[E]E^E_E`E$i{E|E}E~EaFbFcFdFeFfFgFhFiFjFkFlFmFnFoFpFqFrFsFtFuFvFwFxFyFzFAFBFCFDFEFFFGFHFIFJFKFLFMFNFOFPFQFRFSFTFUFVFWFXFYFZF0F1F2F3F]e4F5F6F7F8F9F!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

165 return 0 2a ofukvkpfwkqfxkrfyksfzktfAk8 BkvbCkDk9 Ek%dFk! GkHkIkJkKkLkwbMkNkOkxbPkQkRkybSkzbTkUkAbVkBbWkCbXk# Yk$ Zk% 0k' 1k( 2k) 3k* 4k+ 5k, 6k- 7k. 8k/ 9k: !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~kalblcldlelflglufhliljlvfklwfllxfmlyfnlDbolzfplAfqlBfrlsltlEbulFbvlGbwlHbxlIbyl?czl@cAl[cBl]cCl^cDl_cEl`cFl{cGlHlIlJlKlLlMlNlgiOlhiPliiQlRlSlTlUlVlWlXlYlZl0l1l2l3l4l5l6l7l8l9l!l#l$l%l'l(l)l*l+l,l-l.l/l:l;l=l?l@l[l]l^l_l`l{l|l}l~lambmcmdmemfmgmhmimjmkmlmmmnmompmqmrmsmtmumvmwmxmymzmAmBmCmDmEmFmGmHmImJmKmLmMmNmOmPmQmRmSmTmUmVmWmXmYmZm0m1m2m3m4m5m6m7m8m9m!m#m$m%m'm(m)m*m+m,m-m.m/m:m;m=m?m@m[m]m^m_m`m{m|m}m~manbncndnenfngnhninjijnknlnmnnnonpnqnrnsntnunvnwnxnynznAnBn!bCnDnEnCfFnGnHnInJnKnLnMnNnDfOnPnQnRnSnTnUnVnWn'dXn7eYnEfZnFf0nGf1n2nHf3n4n5n6nni7noi8npi9nqi!nri#nsi$nti%nui'n(n)n*n|c+n(d,n-n.n8e/ne :nf ;n=n?nJb@nqb[nKb]n^n_n`n]b{n^b|n_b}n~n`bao{bbo|bco}bdoeofo~bgohoioIfjokoloJfmovinooopoLbqoMbroNbsoObtoPbuovowoxoyod zoAo9eBo5bCoDo!eEo#eFo$eGoKfHoLfIoMfJog Koh Loi Moj NoabOobbPocbQodbRorbSosbTotbUoNfVoOfWoPfXoQfYoRfZoSf0o1o2oTf3oUf4oeb5oQb6oRb7oVf8oSb9oTb!oUb#o$oVb%ofb'o(oWf)o*oyiziujAiBivjCiDiwjEiFixjGiyj+oNh,oXfzj-o%e.oYfZfAj0f/o1fBj:oOh;oHiCjIiJiDj2f=o3f?o@o[o]o^o_o'e`o6f7f8f)d9f!f#f$f%f'f(f)f*fk +f7 ,f{o|o}o~oapbpcpdpepfpgp-f.f/f:f;f=f?f@f[f]f^f_f`f{f|f}f~fagbgcgdgegfggghgigjgkglgmgng(eogpgqg}crg*dsghpipjp~ctgadugWbkpXblpbdvgcdwgHhxgygIhzgAgYbmpZbnp0bop1bpp2bqp3brpsptpupvpwpxpypzpEjApBpCpDpEpFpGpHpIpJpKpLpMpNpOiOpPpQpPiRpQiRcSpTpUpVpWpXpYpZpSc0p1p2p3p4p5p6p7p8p9pRiSiBdTi-c!p.c#pUiViWi$pXiYiZi%p'p(p)p*p+p,p-p.p/p:p;p=p?p@p[p]p^p_p`p{p|pgbacbc}p~paqbqDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdUdVdWdXdYdZd0d1d2d3d4dDeEecqdqeqfqgqhqiqjqkqlqmqnqoqpqqqrqsqtquqvqwqxqyqzqAqBqCqDqEqFqGqHqIqJqKqLqUhIcJcVhWhXhYhZh0h1h2h0i1ise'chbLcib(cMcjb6b^e_e`el -bm n o p q r s t u .bv w x y %bz A 'b/bte)ckbNclb*cOcmb7b{e|e}eB :bC D E F G H I J K ;bL M N O (bP Q )b=bue+cnbPcob,cQcpb8b~eafbfR ?bS T U V W X Y Z 0 @b1 2 3 4 *b5 6 +b[b3hMqNqOqPqQqRqSq+eTqUqVqWqXqYqZq0q1q2q3q4q5q6q7q8q9q!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~qarbrcrdrerfrgrhrirjrkrlrmrnrorprqrrrsrtrurvr4h5h/ewrxryrzrArBrCGCrDGDrWcXcubYcErZcFrGrHrIrJrKrLrMrNrOrPrQrRrSrTrUr+dVrWrXrYrZr0rUc1r2r4b3r4r5r6r7r8r9r!r#r$r,d%r'r(r)r*r+r,r-r.r/r:r;r=r?r#b@r$b[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~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:B;B=B?B@B[B]B^B_B`B{B|B}B~BaCbCcCdCeCfCgChCiCjCkClCmCnCoCpCqCrCsCtCuCvCwCxCyCzCACBCCCDCECFCGCHCICJCKCLCMCNCOCPCQCRCSCTCUCVCWCXCYCZC0C1C2C3C4C5C6C7C8C9C!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~CaDbDcDdDeDfDgDhDiDjDkDlDmDnDoDpDqDrDsDtDuDvDwDxDyDzDADBDCDDDEDFDGDHDIDJDKDLDMDNDODPDQDRDSDTDUDVDWDXDYDZD0D1D2D3D4D5D6D7D8D9D!D#D$D%D'D(D)D*D+D,D-D.D/D:D;D=D?D@D[D]D^D_D`D{D|D}D~DaEbEcEdEeEfEgEhEiEjEkElEmEnEoEpEqErEsEtEuEvEwExEyEzEAEBECEDEEEFEGEHE,bIEJEKELEMENEOEPEQERESETEUEVEWEXEYEZE0E1E2E3E4E5E6E7E8E9E!E#E$E%E'E(E)E*E+E,E-E.E/E:E;E=E?E@E[E]E^E_E`E$i{E|E}E~EaFbFcFdFeFfFgFhFiFjFkFlFmFnFoFpFqFrFsFtFuFvFwFxFyFzFAFBFCFDFEFFFGFHFIFJFKFLFMFNFOFPFQFRFSFTFUFVFWFXFYFZF0F1F2F3F]e4F5F6F7F8F9F!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

166 cdef const char* name 

167 name_err = cydriver.cuGetErrorName(error, &name) 2fikilimiwi9e5bxi1IEjqjrjsjub4b.d

168 if name_err != cydriver.CUresult.CUDA_SUCCESS: 2fikilimiwi9e5bxi1IEjqjrjsjub4b.d

169 raise CUDAError(f"UNEXPECTED ERROR CODE: {error}") 2Ej

170 with gil: 2fikilimiwi9e5bxi1IEjqjrjsjub4b.d

171 # TODO: consider lower this to Cython 

172 expl = DRIVER_CU_RESULT_EXPLANATIONS.get(int(error)) 2fikilimiwi9e5bxi1IEjqjrjsjub4b.d

173 if expl is not None: 2fikilimiwi9e5bxi1IEjqjrjsjub4b.d

174 raise CUDAError(f"{name.decode()}: {expl}") 2fikilimiwi9e5bxi1IEjqjrjsjub4b.d

175 cdef const char* desc 

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

177 if desc_err != cydriver.CUresult.CUDA_SUCCESS: 

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

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

180  

181  

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

183 if error == _RUNTIME_SUCCESS: 26f7f8f9f!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~fagbgcgdgegfggghgigjgkglgmgngogpgqgrgsgtgugWbXbvgwgxgygzgAgYbZb0b1b2b3bPhIP%iQh)e*eRhShThacbcIcJchbibjb6bl m n o p q r s t u v w x y z A kblbmb7bB C D E F G H I J K L M N O P Q nbobpb8bR S T U V W X Y Z 0 1 2 3 4 5 6 ddVc,eed-e.efd:e;e@e[e

184 return 0 26f7f8f9f!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~fagbgcgdgegfggghgigjgkglgmgngogpgqgrgsgtgugWbXbvgwgxgygzgAgYbZb0b1b2b3bPh%iQh)e*eRhShThacbcIcJchbibjb6bl m n o p q r s t u v w x y z A kblbmb7bB C D E F G H I J K L M N O P Q nbobpb8bR S T U V W X Y Z 0 1 2 3 4 5 6 ddVc,eed-e.efd:e;e@e[e

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

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

187 name = error.name 2IP%i

188 expl = RUNTIME_CUDA_ERROR_EXPLANATIONS.get(int(error)) 2IP%i

189 if expl is not None: 2IP%i

190 raise CUDAError(f"{name}: {expl}") 2IP%i

191 desc_err, desc = runtime.cudaGetErrorString(error) 2%i

192 if desc_err != _RUNTIME_SUCCESS: 2%i

193 raise CUDAError(f"{name}") 

194 desc = desc.decode() 2%i

195 raise CUDAError(f"{name}: {desc}") 2%i

196  

197  

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

199 if error == _NVRTC_SUCCESS: 2a }FJPb ccdcecfcgchcicjckclcmcncocpcqcrcsctcucvcwcxcyczcAcBcCcDcEcFcGcHcKc0c1c2c'i3c(i4c)i5cKhLh6c7c/d8c:d;d=dhdidjd?d@d[d]d$c%ckd^dld_d`d{d|dmd}dndodpdqd~daebecedeeefegeherdsdtdc 9cud!c#c7h8h9h!h#h$h*i+i,i-i.i/i:i;i=i?i@i[i]i^i_i`i{i|i}i~iajbjcjdjejfjgjiejekelemeneoepeqere

200 return 0 2a JPb ccdcecfcgchcicjckclcmcncocpcqcrcsctcucvcwcxcyczcAcBcCcDcEcFcGcHcKc0c1c2c'i3c(i4c)i5cKhLh6c7c/d8c:d;d=dhdidjd?d@d[d]d$c%ckd^dld_d`d{d|dmd}dndodpdqd~daebecedeeefegeherdsdtdc 9cud!c#c7h8h9h!h#h$h*i+i,i-i.i/i:i;i=i?i@i[i]i^i_i`i{i|i}i~iajbjcjdjejfjgjiejekelemeneoepeqere

201 err = f"{error}: {nvrtc.nvrtcGetErrorString(error)[1].decode()}" 2}FJP

202 if handle is not None: 2}FJP

203 _, logsize = nvrtc.nvrtcGetProgramLogSize(handle) 2}F

204 log = b" " * logsize 2}F

205 _ = nvrtc.nvrtcGetProgramLog(handle, log) 2}F

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

207 raise NVRTCError(err) 2}FJP

208  

209  

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

211 if isinstance(error, driver.CUresult): 2a ofukvkpfwkqfxkrfyksfzktfAk8 BkvbCkDk9 Ek%dFk! GkHkIkJkKkLkwbMkNkOkxbPkQkRkybSkzbTkUkAbVkBbWkCbXk# Yk$ Zk% 0k' 1k( 2k) 3k* 4k+ 5k, 6k- 7k. 8k/ 9k: !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~kalblcldlelflglufhliljlvfklwfllxfmlyfnlDbolzfplAfqlBfrlsltlEbulFbvlGbwlHbxlIbyl?czl@cAl[cBl]cCl^cDl_cEl`cFl{cGlHlIlJlKlLlMlNlgiOlhiPliiQlRlSlTlUlVlWlXlYlZl0l1l2l3l4l5l6l7l8l9l!l#l$l%l'l(l)l*l+l,l-l.l/l:l;l=l?l@l[l]l^l_l`l{l|l}l~lambmcmdmemfmgmhmimjmkmlmmmnmompmqmrmsmtmumvmwmxmymzmAmBmCmDmEmFmGmHmImJmKmLmMmNmOmPmQmRmSmTmUmVmWmXmYmZm0m1m2m3m4m5m6m7m8m9m!m#m$m%m'm(m)m*m+m,m-m.m/m:m;m=m?m@m[m]m^m_m`m{m|m}m~manbncndnenfngnhninjijnknlnmnnnonpnqnrnsntnunvnwnxnynznAnBn!bCnDnEnCfFnGnHnInJnKnLnMnNnDfOnPnQnRnSnTnUnVnWn'dXn7eYnEfZnFf0nGf1n2nHf3n4n5n6nni7noi8npi9nqi!nri#nsi$nti%nui'n(n)n*n|c+n(d,n-n.n8e/ne :nf ;n=n?nJb@nqb[nKb]n^n_n`n]b{n^b|n_b}n~n`bao{bbo|bco}bdoeofo~bgohoioIfjokoloJfmovinooopoLbqoMbroNbsoObtoPbuovowoxoyod zoAoBo5bCoDo!eEo#eFo$eGoKfHoLfIoMfJog Koh Loi Moj NoabOobbPocbQodbRorbSosbTotbUoNfVoOfWoPfXoQfYoRfZoSf0o1o2oTf3oUf4oeb5oQb6oRb7oVf8oSb9oTb!oUb#o$oVb%ofb'o(oWf)o*oyiziujAiBivjCiDiwjEiFixjGiyj+oNh,oXfzj-o%e.oYfZfAj0f/o1fBj:oOh;oHiCjIiJiDj2f=o3f?o@o[o]o^o_o'e`o6f7f8f)d9f!f#f$f%f'f(f)f*fk +f7 ,f{o|o}o~oapbpcpdpepfpgp-f.f/f:f;f=f?f@f[f]f^f_f`f{f|f}f~fagbgcgdgegfggghgigjgkglgmgng(eogpgqg}crg*dsghpipjp~ctgadugWbkpXblpbdvgcdwgHhxgygIhzgAgYbmpZbnp0bop1bpp2bqp3brpPhsptpupvpwpxpypzpApQh)eBp*eCpDpEpFpGpHpIpJpRhKpLpMpNpOiOpPpQpPiShThRpQiSpTpUpVpWpXpYpZp0p1p2p3p4p5p6p7p8p9pRiSiTi!p#pUiViWi$pXiYiZi%p'p(p)p*p+p,p-p.p/p:p;p=p?p@p[p]p^p_p`p{p|pgbacbc}p~paqbqb cccqdcdqeceqfcfqgcgqhchqiciqjcjqkckqlclqmcmqncnqocoqpcpqqcqqrcrqscsqtctqucuqvcvqwcwqxcxqycyqzczqAcAqBcBqCcCqDcDqEcEqFcFqGcGqHqHcIqJqKqLqUhIcJcVhWhXhYhZh0h1h2h0i1ise'chbLcib(cMcjb6b^e_e`el -bm n o p q r s t u .bv w x y %bz A 'b/bte)ckbNclb*cOcmb7b{e|e}eB :bC D E F G H I J K ;bL M N O (bP Q )b=bue+cnbPcob,cQcpb8b~eafbfR ?bS T U V W X Y Z 0 @b1 2 3 4 *b5 6 +b[b3hMqddNqOqPqQqRqSq+eTqUqVqWqXqYqZqVc0q,e1qed2q-e3q4q5q6q7q8q9q!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~qarbrcrdrerfrgrhrirjrkrlrmrnrorprqrrrsrtrur.evr4h5h/ewrfdxr:e;eyrzrArBrCrDrErFrGrHrIrJrKrLrMrNrOrPrQrRrKcSrTrUr+dVrWrXrYrZr0r1r2r4b3r4r5r6r7r8r9r!r#r$r,d%r'r(r)r*r0c+r1c,r2c-r.r/r:r'i3c;r(i4c=r)i5c?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~zaAbAcAdAeAfAgAhAiAjAkAlAmAnAoApAqArAsAtAuAvAwAxAyAzAAA@eBA[eCADAEAFAGAHAIAJAKALAMANAOAPAQARASATAUAVAWAXAYAZA0A1AKh2A3A4ALh5A6A7A8A9A!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:B;B=B?B@B[B]B^B_B`B{B|B}B~BaCbCcCdCeCfCgChCiCjCkClCmCnCoCpCqCrCsCtCuCvCwCxCyCzCACBCCCDCECFCGCHCICJCKCLCMCNCOCPCQCRCSCTCUCVCWCXCYCZC0C1C6c2C7c3C/d4C8c5C:d6C;d7C=d8Chd9Cid!Cjd#C?d$C@d%C[d'C]d(C$c)C%c*Ckd+C^d,Cld-C_d.C`d/C{d:C|d;Cmd=C}d?Cnd@Cod[Cpd]Cqd^C~d_Cae`Cbe{Cce|Cde}Cee~CfeaDgebDhecDrddDsdeDtdfDc gD9chDiDjDkDlDmDnDoDpDqDrDsDtDuDvDudwDxDyDzDADBDCDDDED!cFDGDHD#cID7h8h9h!h#h$h*i+i,i-i.i/i:i;i=i?i@i[i]i^iJD_i`i{i|i}i~iajbjcjdjKDLDMDNDejfjgjiejekelemeODneoepeqerePDQDRDSDTDUDVDWDXDYDZD0D1D2D3D4D5D6D7D8D9D!D#D$D%D'D(D)D*D+D,D-D.D/D:D;D=D?D@D[D]D^D_D`D{D|D}D~DaEbEcEdEeEfEgEhEiEjEkElEmEnEoEpEqErEsEtEuEvEwExEyEzEAEBECEDEEEFEGEHE,bIEJEKELEMENEOEPEQERESETEUEVEWEXEYEZE0E1E2E3E4E5E6E7E8E9E!E#E$E%E'E(E)E*E+E,E-E.E/E:E;E=E?E@E[E]E^E_E`E$i{E|E}E~EaFbFcFdFeFfFgFhFiFjFkFlFmFnFoFpFqFrFsFtFuFvFwFxFyFzFAFBFCFDFEFFFGFHFIFJFKFLFMFNFOFPFQFRFSFTFUFVFWFXFYFZF0F1F2F3F]e4F5F6F7F8F9F!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

212 return _check_driver_error(error) 2a ofukvkpfwkqfxkrfyksfzktfAk8 BkvbCkDk9 Ek%dFk! GkHkIkJkKkLkwbMkNkOkxbPkQkRkybSkzbTkUkAbVkBbWkCbXk# Yk$ Zk% 0k' 1k( 2k) 3k* 4k+ 5k, 6k- 7k. 8k/ 9k: !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~kalblcldlelflglufhliljlvfklwfllxfmlyfnlDbolzfplAfqlBfrlsltlEbulFbvlGbwlHbxlIbyl?czl@cAl[cBl]cCl^cDl_cEl`cFl{cGlHlIlJlKlLlMlNlgiOlhiPliiQlRlSlTlUlVlWlXlYlZl0l1l2l3l4l5l6l7l8l9l!l#l$l%l'l(l)l*l+l,l-l.l/l:l;l=l?l@l[l]l^l_l`l{l|l}l~lambmcmdmemfmgmhmimjmkmlmmmnmompmqmrmsmtmumvmwmxmymzmAmBmCmDmEmFmGmHmImJmKmLmMmNmOmPmQmRmSmTmUmVmWmXmYmZm0m1m2m3m4m5m6m7m8m9m!m#m$m%m'm(m)m*m+m,m-m.m/m:m;m=m?m@m[m]m^m_m`m{m|m}m~manbncndnenfngnhninjijnknlnmnnnonpnqnrnsntnunvnwnxnynznAnBn!bCnDnEnCfFnGnHnInJnKnLnMnNnDfOnPnQnRnSnTnUnVnWn'dXn7eYnEfZnFf0nGf1n2nHf3n4n5n6nni7noi8npi9nqi!nri#nsi$nti%nui'n(n)n*n|c+n(d,n-n.n8e/ne :nf ;n=n?nJb@nqb[nKb]n^n_n`n]b{n^b|n_b}n~n`bao{bbo|bco}bdoeofo~bgohoioIfjokoloJfmovinooopoLbqoMbroNbsoObtoPbuovowoxoyod zoAoBo5bCoDo!eEo#eFo$eGoKfHoLfIoMfJog Koh Loi Moj NoabOobbPocbQodbRorbSosbTotbUoNfVoOfWoPfXoQfYoRfZoSf0o1o2oTf3oUf4oeb5oQb6oRb7oVf8oSb9oTb!oUb#o$oVb%ofb'o(oWf)o*oyiziujAiBivjCiDiwjEiFixjGiyj+oNh,oXfzj-o%e.oYfZfAj0f/o1fBj:oOh;oHiCjIiJiDj2f=o3f?o@o[o]o^o_o'e`o6f7f8f)d9f!f#f$f%f'f(f)f*fk +f7 ,f{o|o}o~oapbpcpdpepfpgp-f.f/f:f;f=f?f@f[f]f^f_f`f{f|f}f~fagbgcgdgegfggghgigjgkglgmgng(eogpgqg}crg*dsghpipjp~ctgadugWbkpXblpbdvgcdwgHhxgygIhzgAgYbmpZbnp0bop1bpp2bqp3brpsptpupvpwpxpypzpApBpCpDpEpFpGpHpIpJpKpLpMpNpOiOpPpQpPiRpQiSpTpUpVpWpXpYpZp0p1p2p3p4p5p6p7p8p9pRiSiTi!p#pUiViWi$pXiYiZi%p'p(p)p*p+p,p-p.p/p:p;p=p?p@p[p]p^p_p`p{p|pgbacbc}p~paqbqcqdqeqfqgqhqiqjqkqlqmqnqoqpqqqrqsqtquqvqwqxqyqzqAqBqCqDqEqFqGqHqIqJqKqLqUhIcJcVhWhXhYhZh0h1h2h0i1ise'chbLcib(cMcjb6b^e_e`el -bm n o p q r s t u .bv w x y %bz A 'b/bte)ckbNclb*cOcmb7b{e|e}eB :bC D E F G H I J K ;bL M N O (bP Q )b=bue+cnbPcob,cQcpb8b~eafbfR ?bS T U V W X Y Z 0 @b1 2 3 4 *b5 6 +b[b3hMqNqOqPqQqRqSq+eTqUqVqWqXqYqZq0q1q2q3q4q5q6q7q8q9q!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~qarbrcrdrerfrgrhrirjrkrlrmrnrorprqrrrsrtrurvr4h5h/ewrxryrzrArBrCrDrErFrGrHrIrJrKrLrMrNrOrPrQrRrSrTrUr+dVrWrXrYrZr0r1r2r4b3r4r5r6r7r8r9r!r#r$r,d%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~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:B;B=B?B@B[B]B^B_B`B{B|B}B~BaCbCcCdCeCfCgChCiCjCkClCmCnCoCpCqCrCsCtCuCvCwCxCyCzCACBCCCDCECFCGCHCICJCKCLCMCNCOCPCQCRCSCTCUCVCWCXCYCZC0C1C2C3C4C5C6C7C8C9C!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~CaDbDcDdDeDfDgDhDiDjDkDlDmDnDoDpDqDrDsDtDuDvDwDxDyDzDADBDCDDDEDFDGDHDIDJDKDLDMDNDODPDQDRDSDTDUDVDWDXDYDZD0D1D2D3D4D5D6D7D8D9D!D#D$D%D'D(D)D*D+D,D-D.D/D:D;D=D?D@D[D]D^D_D`D{D|D}D~DaEbEcEdEeEfEgEhEiEjEkElEmEnEoEpEqErEsEtEuEvEwExEyEzEAEBECEDEEEFEGEHE,bIEJEKELEMENEOEPEQERESETEUEVEWEXEYEZE0E1E2E3E4E5E6E7E8E9E!E#E$E%E'E(E)E*E+E,E-E.E/E:E;E=E?E@E[E]E^E_E`E$i{E|E}E~EaFbFcFdFeFfFgFhFiFjFkFlFmFnFoFpFqFrFsFtFuFvFwFxFyFzFAFBFCFDFEFFFGFHFIFJFKFLFMFNFOFPFQFRFSFTFUFVFWFXFYFZF0F1F2F3F]e4F5F6F7F8F9F!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

213 elif isinstance(error, runtime.cudaError_t): 2a 6f7f8f9f!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~fagbgcgdgegfggghgigjgkglgmgngogpgqgrgsgtgugWbXbvgwgxgygzgAgYbZb0b1b2b3bPhQh)e*eRhShThacbcb ccdcecfcgchcicjckclcmcncocpcqcrcsctcucvcwcxcyczcAcBcCcDcEcFcGcHcIcJchbibjb6bl m n o p q r s t u v w x y z A kblbmb7bB C D E F G H I J K L M N O P Q nbobpb8bR S T U V W X Y Z 0 1 2 3 4 5 6 ddVc,eed-e.efd:e;eKc0c1c2c'i3c(i4c)i5c@e[eKhLh6c7c/d8c:d;d=dhdidjd?d@d[d]d$c%ckd^dld_d`d{d|dmd}dndodpdqd~daebecedeeefegeherdsdtdc 9cud!c#c7h8h9h!h#h$h*i+i,i-i.i/i:i;i=i?i@i[i]i^i_i`i{i|i}i~iajbjcjdjejfjgjiejekelemeneoepeqere

214 return _check_runtime_error(error) 26f7f8f9f!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~fagbgcgdgegfggghgigjgkglgmgngogpgqgrgsgtgugWbXbvgwgxgygzgAgYbZb0b1b2b3bPhQh)e*eRhShThacbcIcJchbibjb6bl m n o p q r s t u v w x y z A kblbmb7bB C D E F G H I J K L M N O P Q nbobpb8bR S T U V W X Y Z 0 1 2 3 4 5 6 ddVc,eed-e.efd:e;e@e[e

215 elif isinstance(error, nvrtc.nvrtcResult): 2a b ccdcecfcgchcicjckclcmcncocpcqcrcsctcucvcwcxcyczcAcBcCcDcEcFcGcHcKc0c1c2c'i3c(i4c)i5cKhLh6c7c/d8c:d;d=dhdidjd?d@d[d]d$c%ckd^dld_d`d{d|dmd}dndodpdqd~daebecedeeefegeherdsdtdc 9cud!c#c7h8h9h!h#h$h*i+i,i-i.i/i:i;i=i?i@i[i]i^i_i`i{i|i}i~iajbjcjdjejfjgjiejekelemeneoepeqere

216 return _check_nvrtc_error(error, handle=handle) 2a b ccdcecfcgchcicjckclcmcncocpcqcrcsctcucvcwcxcyczcAcBcCcDcEcFcGcHcKc0c1c2c'i3c(i4c)i5cKhLh6c7c/d8c:d;d=dhdidjd?d@d[d]d$c%ckd^dld_d`d{d|dmd}dndodpdqd~daebecedeeefegeherdsdtdc 9cud!c#c7h8h9h!h#h$h*i+i,i-i.i/i:i;i=i?i@i[i]i^i_i`i{i|i}i~iajbjcjdjejfjgjiejekelemeneoepeqere

217 else: 

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

219  

220  

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

222 _check_error(result[0], handle=handle) 2a ofukvkpfwkqfxkrfyksfzktfAk8 BkvbCkDk9 Ek%dFk! GkHkIkJkKkLkwbMkNkOkxbPkQkRkybSkzbTkUkAbVkBbWkCbXk# Yk$ Zk% 0k' 1k( 2k) 3k* 4k+ 5k, 6k- 7k. 8k/ 9k: !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~kalblcldlelflglufhliljlvfklwfllxfmlyfnlDbolzfplAfqlBfrlsltlEbulFbvlGbwlHbxlIbyl?czl@cAl[cBl]cCl^cDl_cEl`cFl{cGlHlIlJlKlLlMlNlgiOlhiPliiQlRlSlTlUlVlWlXlYlZl0l1l2l3l4l5l6l7l8l9l!l#l$l%l'l(l)l*l+l,l-l.l/l:l;l=l?l@l[l]l^l_l`l{l|l}l~lambmcmdmemfmgmhmimjmkmlmmmnmompmqmrmsmtmumvmwmxmymzmAmBmCmDmEmFmGmHmImJmKmLmMmNmOmPmQmRmSmTmUmVmWmXmYmZm0m1m2m3m4m5m6m7m8m9m!m#m$m%m'm(m)m*m+m,m-m.m/m:m;m=m?m@m[m]m^m_m`m{m|m}m~manbncndnenfngnhninjijnknlnmnnnonpnqnrnsntnunvnwnxnynznAnBn!bCnDnEnCfFnGnHnInJnKnLnMnNnDfOnPnQnRnSnTnUnVnWn'dXn7eYnEfZnFf0nGf1n2nHf3n4n5n6nni7noi8npi9nqi!nri#nsi$nti%nui'n(n)n*n|c+n(d,n-n.n8e/ne :nf ;n=n?nJb@nqb[nKb]n^n_n`n]b{n^b|n_b}n~n`bao{bbo|bco}bdoeofo~bgohoioIfjokoloJfmovinooopoLbqoMbroNbsoObtoPbuovowoxoyod zoAoBo5bCoDo!eEo#eFo$eGoKfHoLfIoMfJog Koh Loi Moj NoabOobbPocbQodbRorbSosbTotbUoNfVoOfWoPfXoQfYoRfZoSf0o1o2oTf3oUf4oeb5oQb6oRb7oVf8oSb9oTb!oUb#o$oVb%ofb'o(oWf)o*oyiziujAiBivjCiDiwjEiFixjGiyj+oNh,oXfzj-o%e.oYfZfAj0f/o1fBj:oOh;oHiCjIiJiDj2f=o3f?o@o[o]o^o_o'e`o6f7f8f)d9f!f#f$f%f'f(f)f*fk +f7 ,f{o|o}o~oapbpcpdpepfpgp-f.f/f:f;f=f?f@f[f]f^f_f`f{f|f}f~fagbgcgdgegfggghgigjgkglgmgng(eogpgqg}crg*dsghpipjp~ctgadugWbkpXblpbdvgcdwgHhxgygIhzgAgYbmpZbnp0bop1bpp2bqp3brpPhsptpupvpwpxpypzpApQh)eBp*eCpDpEpFpGpHpIpJpRhKpLpMpNpOiOpPpQpPiShThRpQiSpTpUpVpWpXpYpZp0p1p2p3p4p5p6p7p8p9pRiSiTi!p#pUiViWi$pXiYiZi%p'p(p)p*p+p,p-p.p/p:p;p=p?p@p[p]p^p_p`p{p|pgbacbc}p~paqbqb cccqdcdqeceqfcfqgcgqhchqiciqjcjqkckqlclqmcmqncnqocoqpcpqqcqqrcrqscsqtctqucuqvcvqwcwqxcxqycyqzczqAcAqBcBqCcCqDcDqEcEqFcFqGcGqHqHcIqJqKqLqUhIcJcVhWhXhYhZh0h1h2h0i1ise'chbLcib(cMcjb6b^e_e`el -bm n o p q r s t u .bv w x y %bz A 'b/bte)ckbNclb*cOcmb7b{e|e}eB :bC D E F G H I J K ;bL M N O (bP Q )b=bue+cnbPcob,cQcpb8b~eafbfR ?bS T U V W X Y Z 0 @b1 2 3 4 *b5 6 +b[b3hMqddNqOqPqQqRqSq+eTqUqVqWqXqYqZqVc0q,e1qed2q-e3q4q5q6q7q8q9q!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~qarbrcrdrerfrgrhrirjrkrlrmrnrorprqrrrsrtrur.evr4h5h/ewrfdxr:e;eyrzrArBrCrDrErFrGrHrIrJrKrLrMrNrOrPrQrRrKcSrTrUr+dVrWrXrYrZr0r1r2r4b3r4r5r6r7r8r9r!r#r$r,d%r'r(r)r*r0c+r1c,r2c-r.r/r:r'i3c;r(i4c=r)i5c?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~zaAbAcAdAeAfAgAhAiAjAkAlAmAnAoApAqArAsAtAuAvAwAxAyAzAAA@eBA[eCADAEAFAGAHAIAJAKALAMANAOAPAQARASATAUAVAWAXAYAZA0A1AKh2A3A4ALh5A6A7A8A9A!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:B;B=B?B@B[B]B^B_B`B{B|B}B~BaCbCcCdCeCfCgChCiCjCkClCmCnCoCpCqCrCsCtCuCvCwCxCyCzCACBCCCDCECFCGCHCICJCKCLCMCNCOCPCQCRCSCTCUCVCWCXCYCZC0C1C6c2C7c3C/d4C8c5C:d6C;d7C=d8Chd9Cid!Cjd#C?d$C@d%C[d'C]d(C$c)C%c*Ckd+C^d,Cld-C_d.C`d/C{d:C|d;Cmd=C}d?Cnd@Cod[Cpd]Cqd^C~d_Cae`Cbe{Cce|Cde}Cee~CfeaDgebDhecDrddDsdeDtdfDc gD9chDiDjDkDlDmDnDoDpDqDrDsDtDuDvDudwDxDyDzDADBDCDDDED!cFDGDHD#cID7h8h9h!h#h$h*i+i,i-i.i/i:i;i=i?i@i[i]i^iJD_i`i{i|i}i~iajbjcjdjKDLDMDNDejfjgjiejekelemeODneoepeqerePDQDRDSDTDUDVDWDXDYDZD0D1D2D3D4D5D6D7D8D9D!D#D$D%D'D(D)D*D+D,D-D.D/D:D;D=D?D@D[D]D^D_D`D{D|D}D~DaEbEcEdEeEfEgEhEiEjEkElEmEnEoEpEqErEsEtEuEvEwExEyEzEAEBECEDEEEFEGEHE,bIEJEKELEMENEOEPEQERESETEUEVEWEXEYEZE0E1E2E3E4E5E6E7E8E9E!E#E$E%E'E(E)E*E+E,E-E.E/E:E;E=E?E@E[E]E^E_E`E$i{E|E}E~EaFbFcFdFeFfFgFhFiFjFkFlFmFnFoFpFqFrFsFtFuFvFwFxFyFzFAFBFCFDFEFFFGFHFIFJFKFLFMFNFOFPFQFRFSFTFUFVFWFXFYFZF0F1F2F3F]e4F5F6F7F8F9F!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

223 cdef int out_len = len(result) 2a ofukvkpfwkqfxkrfyksfzktfAk8 BkvbCkDk9 Ek%dFk! GkHkIkJkKkLkwbMkNkOkxbPkQkRkybSkzbTkUkAbVkBbWkCbXk# Yk$ Zk% 0k' 1k( 2k) 3k* 4k+ 5k, 6k- 7k. 8k/ 9k: !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~kalblcldlelflglufhliljlvfklwfllxfmlyfnlDbolzfplAfqlBfrlsltlEbulFbvlGbwlHbxlIbyl?czl@cAl[cBl]cCl^cDl_cEl`cFl{cGlHlIlJlKlLlMlNlgiOlhiPliiQlRlSlTlUlVlWlXlYlZl0l1l2l3l4l5l6l7l8l9l!l#l$l%l'l(l)l*l+l,l-l.l/l:l;l=l?l@l[l]l^l_l`l{l|l}l~lambmcmdmemfmgmhmimjmkmlmmmnmompmqmrmsmtmumvmwmxmymzmAmBmCmDmEmFmGmHmImJmKmLmMmNmOmPmQmRmSmTmUmVmWmXmYmZm0m1m2m3m4m5m6m7m8m9m!m#m$m%m'm(m)m*m+m,m-m.m/m:m;m=m?m@m[m]m^m_m`m{m|m}m~manbncndnenfngnhninjijnknlnmnnnonpnqnrnsntnunvnwnxnynznAnBn!bCnDnEnCfFnGnHnInJnKnLnMnNnDfOnPnQnRnSnTnUnVnWn'dXn7eYnEfZnFf0nGf1n2nHf3n4n5n6nni7noi8npi9nqi!nri#nsi$nti%nui'n(n)n*n|c+n(d,n-n.n8e/ne :nf ;n=n?nJb@nqb[nKb]n^n_n`n]b{n^b|n_b}n~n`bao{bbo|bco}bdoeofo~bgohoioIfjokoloJfmovinooopoLbqoMbroNbsoObtoPbuovowoxoyod zoAoBo5bCoDo!eEo#eFo$eGoKfHoLfIoMfJog Koh Loi Moj NoabOobbPocbQodbRorbSosbTotbUoNfVoOfWoPfXoQfYoRfZoSf0o1o2oTf3oUf4oeb5oQb6oRb7oVf8oSb9oTb!oUb#o$oVb%ofb'o(oWf)o*oyiziujAiBivjCiDiwjEiFixjGiyj+oNh,oXfzj-o%e.oYfZfAj0f/o1fBj:oOh;oHiCjIiJiDj2f=o3f?o@o[o]o^o_o'e`o6f7f8f)d9f!f#f$f%f'f(f)f*fk +f7 ,f{o|o}o~oapbpcpdpepfpgp-f.f/f:f;f=f?f@f[f]f^f_f`f{f|f}f~fagbgcgdgegfggghgigjgkglgmgng(eogpgqg}crg*dsghpipjp~ctgadugWbkpXblpbdvgcdwgHhxgygIhzgAgYbmpZbnp0bop1bpp2bqp3brpPhsptpupvpwpxpypzpApQh)eBp*eCpDpEpFpGpHpIpJpRhKpLpMpNpOiOpPpQpPiShThRpQiSpTpUpVpWpXpYpZp0p1p2p3p4p5p6p7p8p9pRiSiTi!p#pUiViWi$pXiYiZi%p'p(p)p*p+p,p-p.p/p:p;p=p?p@p[p]p^p_p`p{p|pgbacbc}p~paqbqb cccqdcdqeceqfcfqgcgqhchqiciqjcjqkckqlclqmcmqncnqocoqpcpqqcqqrcrqscsqtctqucuqvcvqwcwqxcxqycyqzczqAcAqBcBqCcCqDcDqEcEqFcFqGcGqHqHcIqJqKqLqUhIcJcVhWhXhYhZh0h1h2h0i1ise'chbLcib(cMcjb6b^e_e`el -bm n o p q r s t u .bv w x y %bz A 'b/bte)ckbNclb*cOcmb7b{e|e}eB :bC D E F G H I J K ;bL M N O (bP Q )b=bue+cnbPcob,cQcpb8b~eafbfR ?bS T U V W X Y Z 0 @b1 2 3 4 *b5 6 +b[b3hMqddNqOqPqQqRqSq+eTqUqVqWqXqYqZqVc0q,e1qed2q-e3q4q5q6q7q8q9q!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~qarbrcrdrerfrgrhrirjrkrlrmrnrorprqrrrsrtrur.evr4h5h/ewrfdxr:e;eyrzrArBrCrDrErFrGrHrIrJrKrLrMrNrOrPrQrRrKcSrTrUr+dVrWrXrYrZr0r1r2r4b3r4r5r6r7r8r9r!r#r$r,d%r'r(r)r*r0c+r1c,r2c-r.r/r:r'i3c;r(i4c=r)i5c?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~zaAbAcAdAeAfAgAhAiAjAkAlAmAnAoApAqArAsAtAuAvAwAxAyAzAAA@eBA[eCADAEAFAGAHAIAJAKALAMANAOAPAQARASATAUAVAWAXAYAZA0A1AKh2A3A4ALh5A6A7A8A9A!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:B;B=B?B@B[B]B^B_B`B{B|B}B~BaCbCcCdCeCfCgChCiCjCkClCmCnCoCpCqCrCsCtCuCvCwCxCyCzCACBCCCDCECFCGCHCICJCKCLCMCNCOCPCQCRCSCTCUCVCWCXCYCZC0C1C6c2C7c3C/d4C8c5C:d6C;d7C=d8Chd9Cid!Cjd#C?d$C@d%C[d'C]d(C$c)C%c*Ckd+C^d,Cld-C_d.C`d/C{d:C|d;Cmd=C}d?Cnd@Cod[Cpd]Cqd^C~d_Cae`Cbe{Cce|Cde}Cee~CfeaDgebDhecDrddDsdeDtdfDc gD9chDiDjDkDlDmDnDoDpDqDrDsDtDuDvDudwDxDyDzDADBDCDDDED!cFDGDHD#cID7h8h9h!h#h$h*i+i,i-i.i/i:i;i=i?i@i[i]i^iJD_i`i{i|i}i~iajbjcjdjKDLDMDNDejfjgjiejekelemeODneoepeqerePDQDRDSDTDUDVDWDXDYDZD0D1D2D3D4D5D6D7D8D9D!D#D$D%D'D(D)D*D+D,D-D.D/D:D;D=D?D@D[D]D^D_D`D{D|D}D~DaEbEcEdEeEfEgEhEiEjEkElEmEnEoEpEqErEsEtEuEvEwExEyEzEAEBECEDEEEFEGEHE,bIEJEKELEMENEOEPEQERESETEUEVEWEXEYEZE0E1E2E3E4E5E6E7E8E9E!E#E$E%E'E(E)E*E+E,E-E.E/E:E;E=E?E@E[E]E^E_E`E$i{E|E}E~EaFbFcFdFeFfFgFhFiFjFkFlFmFnFoFpFqFrFsFtFuFvFwFxFyFzFAFBFCFDFEFFFGFHFIFJFKFLFMFNFOFPFQFRFSFTFUFVFWFXFYFZF0F1F2F3F]e4F5F6F7F8F9F!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

224 if out_len == 1: 2a ofukvkpfwkqfxkrfyksfzktfAk8 BkvbCkDk9 Ek%dFk! GkHkIkJkKkLkwbMkNkOkxbPkQkRkybSkzbTkUkAbVkBbWkCbXk# Yk$ Zk% 0k' 1k( 2k) 3k* 4k+ 5k, 6k- 7k. 8k/ 9k: !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~kalblcldlelflglufhliljlvfklwfllxfmlyfnlDbolzfplAfqlBfrlsltlEbulFbvlGbwlHbxlIbyl?czl@cAl[cBl]cCl^cDl_cEl`cFl{cGlHlIlJlKlLlMlNlgiOlhiPliiQlRlSlTlUlVlWlXlYlZl0l1l2l3l4l5l6l7l8l9l!l#l$l%l'l(l)l*l+l,l-l.l/l:l;l=l?l@l[l]l^l_l`l{l|l}l~lambmcmdmemfmgmhmimjmkmlmmmnmompmqmrmsmtmumvmwmxmymzmAmBmCmDmEmFmGmHmImJmKmLmMmNmOmPmQmRmSmTmUmVmWmXmYmZm0m1m2m3m4m5m6m7m8m9m!m#m$m%m'm(m)m*m+m,m-m.m/m:m;m=m?m@m[m]m^m_m`m{m|m}m~manbncndnenfngnhninjijnknlnmnnnonpnqnrnsntnunvnwnxnynznAnBn!bCnDnEnCfFnGnHnInJnKnLnMnNnDfOnPnQnRnSnTnUnVnWn'dXn7eYnEfZnFf0nGf1n2nHf3n4n5n6nni7noi8npi9nqi!nri#nsi$nti%nui'n(n)n*n|c+n(d,n-n.n8e/ne :nf ;n=n?nJb@nqb[nKb]n^n_n`n]b{n^b|n_b}n~n`bao{bbo|bco}bdoeofo~bgohoioIfjokoloJfmovinooopoLbqoMbroNbsoObtoPbuovowoxoyod zoAoBo5bCoDo!eEo#eFo$eGoKfHoLfIoMfJog Koh Loi Moj NoabOobbPocbQodbRorbSosbTotbUoNfVoOfWoPfXoQfYoRfZoSf0o1o2oTf3oUf4oeb5oQb6oRb7oVf8oSb9oTb!oUb#o$oVb%ofb'o(oWf)o*oyiziujAiBivjCiDiwjEiFixjGiyj+oNh,oXfzj-o%e.oYfZfAj0f/o1fBj:oOh;oHiCjIiJiDj2f=o3f?o@o[o]o^o_o'e`o6f7f8f)d9f!f#f$f%f'f(f)f*fk +f7 ,f{o|o}o~oapbpcpdpepfpgp-f.f/f:f;f=f?f@f[f]f^f_f`f{f|f}f~fagbgcgdgegfggghgigjgkglgmgng(eogpgqg}crg*dsghpipjp~ctgadugWbkpXblpbdvgcdwgHhxgygIhzgAgYbmpZbnp0bop1bpp2bqp3brpPhsptpupvpwpxpypzpApQh)eBp*eCpDpEpFpGpHpIpJpRhKpLpMpNpOiOpPpQpPiShThRpQiSpTpUpVpWpXpYpZp0p1p2p3p4p5p6p7p8p9pRiSiTi!p#pUiViWi$pXiYiZi%p'p(p)p*p+p,p-p.p/p:p;p=p?p@p[p]p^p_p`p{p|pgbacbc}p~paqbqb cccqdcdqeceqfcfqgcgqhchqiciqjcjqkckqlclqmcmqncnqocoqpcpqqcqqrcrqscsqtctqucuqvcvqwcwqxcxqycyqzczqAcAqBcBqCcCqDcDqEcEqFcFqGcGqHqHcIqJqKqLqUhIcJcVhWhXhYhZh0h1h2h0i1ise'chbLcib(cMcjb6b^e_e`el -bm n o p q r s t u .bv w x y %bz A 'b/bte)ckbNclb*cOcmb7b{e|e}eB :bC D E F G H I J K ;bL M N O (bP Q )b=bue+cnbPcob,cQcpb8b~eafbfR ?bS T U V W X Y Z 0 @b1 2 3 4 *b5 6 +b[b3hMqddNqOqPqQqRqSq+eTqUqVqWqXqYqZqVc0q,e1qed2q-e3q4q5q6q7q8q9q!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~qarbrcrdrerfrgrhrirjrkrlrmrnrorprqrrrsrtrur.evr4h5h/ewrfdxr:e;eyrzrArBrCrDrErFrGrHrIrJrKrLrMrNrOrPrQrRrKcSrTrUr+dVrWrXrYrZr0r1r2r4b3r4r5r6r7r8r9r!r#r$r,d%r'r(r)r*r0c+r1c,r2c-r.r/r:r'i3c;r(i4c=r)i5c?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~zaAbAcAdAeAfAgAhAiAjAkAlAmAnAoApAqArAsAtAuAvAwAxAyAzAAA@eBA[eCADAEAFAGAHAIAJAKALAMANAOAPAQARASATAUAVAWAXAYAZA0A1AKh2A3A4ALh5A6A7A8A9A!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:B;B=B?B@B[B]B^B_B`B{B|B}B~BaCbCcCdCeCfCgChCiCjCkClCmCnCoCpCqCrCsCtCuCvCwCxCyCzCACBCCCDCECFCGCHCICJCKCLCMCNCOCPCQCRCSCTCUCVCWCXCYCZC0C1C6c2C7c3C/d4C8c5C:d6C;d7C=d8Chd9Cid!Cjd#C?d$C@d%C[d'C]d(C$c)C%c*Ckd+C^d,Cld-C_d.C`d/C{d:C|d;Cmd=C}d?Cnd@Cod[Cpd]Cqd^C~d_Cae`Cbe{Cce|Cde}Cee~CfeaDgebDhecDrddDsdeDtdfDc gD9chDiDjDkDlDmDnDoDpDqDrDsDtDuDvDudwDxDyDzDADBDCDDDED!cFDGDHD#cID7h8h9h!h#h$h*i+i,i-i.i/i:i;i=i?i@i[i]i^iJD_i`i{i|i}i~iajbjcjdjKDLDMDNDejfjgjiejekelemeODneoepeqerePDQDRDSDTDUDVDWDXDYDZD0D1D2D3D4D5D6D7D8D9D!D#D$D%D'D(D)D*D+D,D-D.D/D:D;D=D?D@D[D]D^D_D`D{D|D}D~DaEbEcEdEeEfEgEhEiEjEkElEmEnEoEpEqErEsEtEuEvEwExEyEzEAEBECEDEEEFEGEHE,bIEJEKELEMENEOEPEQERESETEUEVEWEXEYEZE0E1E2E3E4E5E6E7E8E9E!E#E$E%E'E(E)E*E+E,E-E.E/E:E;E=E?E@E[E]E^E_E`E$i{E|E}E~EaFbFcFdFeFfFgFhFiFjFkFlFmFnFoFpFqFrFsFtFuFvFwFxFyFzFAFBFCFDFEFFFGFHFIFJFKFLFMFNFOFPFQFRFSFTFUFVFWFXFYFZF0F1F2F3F]e4F5F6F7F8F9F!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

225 return 28 9 %d! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ ji'd7e|ce f qb]b^b_b`b{b|b}b~b5bg h i j abbbcbdbrbsbtbebfbujvjwjxjyjNhzjAjBjOhCjDj'e6f7f8f)d9f!f#f$f%f'f(f)f*fk +f,f-f.f/f:f;f=f?f@f[f]f^f_f`f{f|f}f~fagbgcgdgegfggghgigjgkglgmgng(eogpgqg}crgsg~ctgadugWbXbbdvgcdwgHhxgygIhzgAgYbZb0b1b2b3bsp)e*egbacbcUhIcJcVhWhXhYhZh0h1h2h0i1ise'chbLcib(cMcjb6b^e_e`el -bm n o p q r s t u .bv w x y %bz A 'b/bte)ckbNclb*cOcmb7b{e|e}eB :bC D E F G H I J K ;bL M N O (bP Q )b=bue+cnbPcob,cQcpb8b~eafbfR ?bS T U V W X Y Z 0 @b1 2 3 4 *b5 6 +b[b3hdd+eVc,eed-e.e4h5h/efd:e;e@e[e

226 elif out_len == 2: 

227 return result[1] 2a ofukvkpfwkqfxkrfyksfzktfAk8 BkvbCkDk9 Ek%dFk! GkHkIkJkKkLkwbMkNkOkxbPkQkRkybSkzbTkUkAbVkBbWkCbXk# Yk$ Zk% 0k' 1k( 2k) 3k* 4k+ 5k, 6k- 7k. 8k/ 9k: !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~kalblcldlelflglufhliljlvfklwfllxfmlyfnlDbolzfplAfqlBfrlsltlEbulFbvlGbwlHbxlIbyl?czl@cAl[cBl]cCl^cDl_cEl`cFl{cGlHlIlJlKlLlMlNlgiOlhiPliiQlRlSlTlUlVlWlXlYlZl0l1l2l3l4l5l6l7l8l9l!l#l$l%l'l(l)l*l+l,l-l.l/l:l;l=l?l@l[l]l^l_l`l{l|l}l~lambmcmdmemfmgmhmimjmkmlmmmnmompmqmrmsmtmumvmwmxmymzmAmBmCmDmEmFmGmHmImJmKmLmMmNmOmPmQmRmSmTmUmVmWmXmYmZm0m1m2m3m4m5m6m7m8m9m!m#m$m%m'm(m)m*m+m,m-m.m/m:m;m=m?m@m[m]m^m_m`m{m|m}m~manbncndnenfngnhninjnknlnmnnnonpnqnrnsntnunvnwnxnynznAnBn!bCnDnEnCfFnGnHnInJnKnLnMnNnDfOnPnQnRnSnTnUnVnWn'dXn7eYnEfZnFf0nGf1n2nHf3n4n5n6nni7noi8npi9nqi!nri#nsi$nti%nui'n(n)n*n|c+n(d,n-n.n8e/ne :nf ;n=n?nJb@nqb[nKb]n^n_n`n]b{n^b|n_b}n~n`bao{bbo|bco}bdoeofo~bgohoioIfjokoloJfmovinooopoLbqoMbroNbsoObtoPbuovowoxoyod zoAoBo5bCoDo!eEo#eFo$eGoKfHoLfIoMfJog Koh Loi Moj NoabOobbPocbQodbRorbSosbTotbUoNfVoOfWoPfXoQfYoRfZoSf0o1o2oTf3oUf4oeb5oQb6oRb7oVf8oSb9oTb!oUb#o$oVb%ofb'o(oWf)o*oyiziujAiBivjCiDiwjEiFixjGiyj+oNh,oXfzj-o%e.oYfZfAj0f/o1fBj:oOh;oHiCjIiJiDj2f=o3f?o@o[o]o^o_o'e`o6f7f8f9f!f#f$f%f'f(f)f*fk +f7 ,f{o|o}o~oapbpcpdpepfpgp-f.f/f:f;f=f?f@f[f]f^f_f`f{f|f}f~fagbgcgdgegfggghgigjgkglgmgng(eogpgqg}crg*dsghpipjp~ctgadugWbkpXblpbdvgcdwgxgygzgAgYbmpZbnp0bop1bpp2bqp3brpPhtpupvpwpxpypzpApQhBpCpDpEpFpGpHpIpJpRhKpLpMpNpOiOpPpQpPiShThRpQiSpTpUpVpWpXpYpZp0p1p2p3p4p5p6p7p8p9pRiSiTi!p#pUiViWi$pXiYiZi%p'p(p)p*p+p,p-p.p/p:p;p=p?p@p[p]p^p_p`p{p|pgbacbc}p~paqbqcqdqeqfqgqhqiqjqkqlqmqnqoqpqqqrqsqtquqvqwqxqyqzqAqBqCqDqEqFqGqHqIqJqKqLqUhIcJcVhWhXhYhZh0h1h2hse'chbLcib(cMcjb6b^e_e`el -bm n o p q r s t u .bv w x y %bz A 'b/bte)ckbNclb*cOcmb7b{e|e}eB :bC D E F G H I J K ;bL M N O (bP Q )b=bue+cnbPcob,cQcpb8b~eafbfR ?bS T U V W X Y Z 0 @b1 2 3 4 *b5 6 +b[b3hMqNqOqPqQqRqSq+eTqUqVqWqXqYqZq0q1q2q3q4q5q6q7q8q9q!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~qarbrcrdrerfrgrhrirjrkrlrmrnrorprqrrrsrtrurvr4h5h/ewrxryrzrArBrCrDrErFrGrHrIrJrKrLrMrNrOrPrQrRrSrTrUr+dVrWrXrYrZr0r1r2r4b3r4r5r6r7r8r9r!r#r$r,d%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~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:B;B=B?B@B[B]B^B_B`B{B|B}B~BaCbCcCdCeCfCgChCiCjCkClCmCnCoCpCqCrCsCtCuCvCwCxCyCzCACBCCCDCECFCGCHCICJCKCLCMCNCOCPCQCRCSCTCUCVCWCXCYCZC0C1C2C3C4C5C6C7C8C9C!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~CaDbDcDdDeDfDgDhDiDjDkDlDmDnDoDpDqDrDsDtDuDvDwDxDyDzDADBDCDDDEDFDGDHDIDJDKDLDMDNDODPDQDRDSDTDUDVDWDXDYDZD0D1D2D3D4D5D6D7D8D9D!D#D$D%D'D(D)D*D+D,D-D.D/D:D;D=D?D@D[D]D^D_D`D{D|D}D~DaEbEcEdEeEfEgEhEiEjEkElEmEnEoEpEqErEsEtEuEvEwExEyEzEAEBECEDEEEFEGEHE,bIEJEKELEMENEOEPEQERESETEUEVEWEXEYEZE0E1E2E3E4E5E6E7E8E9E!E#E$E%E'E(E)E*E+E,E-E.E/E:E;E=E?E@E[E]E^E_E`E$i{E|E}E~EaFbFcFdFeFfFgFhFiFjFkFlFmFnFoFpFqFrFsFtFuFvFwFxFyFzFAFBFCFDFEFFFGFHFIFJFKFLFMFNFOFPFQFRFSFTFUFVFWFXFYFZF0F1F2F3F]e4F5F6F7F8F9F!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

228 else: 

229 return result[1:] 2a 8 9 %d! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ 'debfbb ccdcecfcgchcicjckclcmcncocpcqcrcsctcucvcwcxcyczcAcBcCcDcEcFcGcHcKc0c1c2c'i3c(i4c)i5cKhLh6c7c/d8c:d;d=dhdidjd?d@d[d]d$c%ckd^dld_d`d{d|dmd}dndodpdqd~daebecedeeefegeherdsdtdc 9cud!c#c7h8h9h!h#h$h*i+i,i-i.i/i:i;i=i?i@i[i]i^i_i`i{i|i}i~iajbjcjdjejfjgjiejekelemeneoepeqere

230  

231  

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

233 """ 

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

235 """ 

236 if options is None: 2ofiJpfqfrfsftf8 vbve9 %d! we9PjJxeyewbzekJxbAelJybzb!PAbBbCb# $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ ufvfwfxfyfDbzfAfBfEbFbGbHbIb?c@c[c]c^c_c`c{cBgCgDgmJnJvdEgqJrJwdFguJvJxdGgyJzJydHgCJDJzdIgGJHJAd!bCfKJDfJgKgLgMg'd7eEfFfGfNgHfLJMJNJ|c(dOJPJ8ee f cGOgJbqbKbQJRJSJ]b^b_bTJ`b{b|b}bUJVJ~bPgIfJfLbMbNbObPbd 5b!e#e$eKfLfMfg h i j abbbcbdbrbsbtbNfOfPfQfRfSfTfUfebQbRbVfSbTbUbBeVbfbWJXJWfYJZJ0JXf1J2J%eYfZf0f3J4J1f5J6J2f7J3f8J9J!J#J$J%J'J(J)J*J+J'e,J-JdG.J)d/J:J;J=J?JeG@JfG[J]J^Jk _J7 `JKiLiljmjMiNinjoj7jbJKP{J|J}J~JaKbKcKdKeKfKgKhKiKjKkKlKmKnKoKpKqKrKsKtKuKvKwKxKyKzKAK(eBKCK}cDK*dEK~cFKadGKWbXbbdHKcdIKJKKKLKMKYbZb0b1b2b3b#P$P%P'P(P)P2I)e*e*P+PgGRchG,PNK-POKPKQKRKSKTKScUKVKWKXKiGjGkGCeYKZK0K1K2K3K4K5K6K7K8K9K!K6I7I#K8IBd8j-clG.cmG9I.P!I/P#I$I%I$K'I(I)I*I+I,I-I.I6P7P8P:P;P=P?P@P[P~FgbacbcCdTcDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdUdVdWdXdYdZd0d1d2d3d4dDeEeb cc`Odc{Oec|Ofc}Ogc~OhcaPicbPjccPkcdPlcePmcfPnc~IocaJpcgPqchPrciPscjPtckPuclPvcmPwcnPxcoPycpPzcqPAcrPBcsPCc5fDctPEcuPFcvPGcwPnGxPHcIcJc%K'Kse'chbLcib(cMcjb6b^e_e`el -bm n o p q r s t u .bv w x y %bz A 'b/bte)ckbNclb*cOcmb7b{e|e}eB :bC D E F G H I J K ;bL M N O (bP Q )b=bue+cnbPcob,cQcpb8b~eafbfR ?bS T U V W X Y Z 0 @b1 2 3 4 *b5 6 +b[b(K)K*Kdd+K,K+eoGpG-K.KqGrGsGtGVc,eed-e/K:K;K=K?K@K[K]K^K_K`K{K|K}K~KaLuGbLvGcLwGdLxGeLyGfLzGgLAGhLBGiLjLkLlLmLnLoLpLqLrLsLtLuLvLwLxLyL/I9j:I;I!j=I.e/ezLALBLfd:e;eCLDLELFLCGDGWcXcubYcZcKcEGFG+dGGHGIGJGKGUcLG4bMGNGOGPGQGRGSG9OTGUG,dVGWGXGYG0c1c2cZG0G1G3c4c5c#b$b2G3G4G5G6G7G8G9G!G#G$G%G'G(G)G*G+G,G-G.G/G:G;G=G?G@G[G]G^G_G`G{G|G}G~GaHbHcHdHeHfHgHhHiHjHkHlHmHnHoHpHqHrHsHtHuHvHwHxHyHzHAHBHCHDHJh=e?egd-dGLQgHLILJLEHRgFHLLMLGHNLOLSgPLQLRLHHTgIHTLUgULVLVgXLFeYLZL?I@I0L1L2L3L4L#jGecf5L$j%j6L7L8LHeIeWgXgYgZg0g1g9L!L2g$L3g%L'LJH4gKH)L5g*L+L,LLH6gMH.L7g/L:L8g=L[IJe?L@L[L]L^L_L`L{L|L}LNHKedf~L'jOHaMbMcMdMeMfMgMhMiMjMkMlMmMnM9goMpM!grMsM#gPHLeQHRHSHTHUHVHWHXHYHMeZH2i0H1H2H$gNeOe%g'g(g)g*g5d+g,g/cefff6d:c7d-guMvM8d(j)j3HPe4H5H6H7H8H9H!H#H$H%HQe'H3i(H)H*H.gCMDM/gFM]IReGMHMIMJMKMLMMMNMOMPMQM+HSehfRM*j,HSMTMUMVM:gWMXMYM-H;g.H0M=g1M2M?g4MTe5M6M^I_I7M8M9M!M#M+jUeif$M,j-j%M'M(MVeWe@g[g]g^g_g`g)M*M{g,M|g-M.M/H}g:H:M~g;M=M?M;Hah=H[Mbh]M^Mch`M`IXe{M|M}M~MaNbNcNdNeNfN?HYejfgN.j@HhNiNjNkNlNmNnNoNpNqNrNsNtNuNdhvNwNehyNzNfh[HZe]H^H_H`H{H|H}H~HaI0ebI4icIdIeIgh1e2ehhihjhkhlh9dmhnh;ckflf!d=c#dohBNCN$d/j:jfI3egIhIiIjIkIlImInIoIpI4eqI5irIsItIphJNKNqhMN{I5eNNONPNQNRNSNTNUNVNWNXNuI6enfYN;jvIZN0N1N@e[e=j?jwIxI2Nrh3N4N5NyIshzIAIBICI7N8Nth9N!N#NDIuhEI%N'Nvh(N)N*NFIwhGI,N-Nxh.N/N:NHIyhII=N?Nzh@N[N]NJIAhKI_N`NBh{N|N}NLIChMIaObODhcOdOeONIEhOIPIgOhOFhiOjOkOQIGhRISImO)O*O6c7c|I/d8c:d;d=dhdidjd?d@d[d]d$c%ckd^dld_d`d{d|dmd}dndodpdqd~daebecedeeefegeherdsdtdc 9cTI9bUI.d+O,O-O.O/O:O;O=O?O@O[O]O^O_O@j[jnOVIWIXIudYIoOpOZI}ItjyPzPAPBPCPDPEPFP!cGPHP#ciejekelemeneoe!O#O$O%Ope'Oqere(O]P^P_P`P{P|PqO}P~PaQbQcQdQeQfQgQ0IhQrOsOiQtOjQ]j^j_j`j{j|j}j~jakbkckdk,bek6h6i7i8ifkgkhkik9iuOvOwOxOyOzOjkAOBOCODOEOkklkFOmknkok!ipk#iqkrk4fsktkGOHOIOJOKOLOMONOOOPOQOROSOTOUOVOWOXOYOZO0O1O2O3O4O5O6O7O8O]e

237 if keep_none: 2ofiJpfqfrfsftf8 vbve9 %d! we9PjJxeyewbzekJxbAelJybzb!PAbBbCb# $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ ufvfwfxfyfDbzfAfBfEbFbGbHbIb?c@c[c]c^c_c`c{cmJnJqJrJuJvJyJzJCJDJGJHJCfKJDf'd7eEfFfGfHf|c(d8ee f JbqbKbQJRJSJ]b^b_bTJ`b{b|b}bUJVJ~bIfJfLbMbNbObPbd 5b!e#e$eKfLfMfg h i j abbbcbdbrbsbtbNfOfPfQfRfSfTfUfebQbRbVfSbTbUbBeVbfbWJXJWfYJZJ0JXf1J2J%eYfZf0f3J4J1f5J6J2f7J3f8J9J!J#J$J%J'J(J)J*J+J'ek 7 KiLiljmjMiNinjoj7jbJKP}c*d~cadWbXbbdcdYbZb0b1b2b3b#P$P%P'P(P)P)e*e*P+PgGRchGNKOKPKQKRKSKTKScUKVKWKXKiGjGkGCeYKZK0K1K2K3K4K5K6K7K8K9K!K#KBd-c.c/P$KgbacbcTcIcJc%K'Kse'chbLcib(cMcjb6b^e_e`el -bm n o p q r s t u .bv w x y %bz A 'b/bte)ckbNclb*cOcmb7b{e|e}eB :bC D E F G H I J K ;bL M N O (bP Q )b=bue+cnbPcob,cQcpb8b~eafbfR ?bS T U V W X Y Z 0 @b1 2 3 4 *b5 6 +b[b(K)K*Kdd+K+eoGpG-K.KqGrGsGtGVc,eed-euGbLvGcLwGdLxGeLyGfLzGgLAGhLBGiL9j!j/ezLALBLfdCGDGEGFG+dGGHGIGJGKGUcLG4bMGNGOGPGQGRGSG9OTGUG,dVGWGXGYG3c4c5c#b$b2G3G4G5G6G7G8G9G!G#G$G%G'G(G)G*G+G,G-G.G/G:G;G=G?G@G[G]G^G_G`G{G|G}G~GaHbHcHdHeHfHgHhHiHjHkHlHmHnHoHpHqHrHsHtHuHvHwHxHyHzHAHBHCHDHgdGLHLILJLEHFHLLMLGHNLOLPLQLRLHHIHTLULVLXLFeYLZL?I@I0L1L2L3L4L#jGecf5L$j%j6L7L8LHeIe9L!L$L%L'LJHKH)L*L+L,LLHMH.L/L:L=L[IJe?L@L[L]L^L_L`L{L|L}LNHKedf~L'jOHaMbMcMdMeMfMgMhMiMjMkMlMmMnMoMpMrMsMPHLeQHRHSHTHUHVHWHXHYHMeZH2i0H1H2HNeOeefff6d7duMvM(j)j3HPe4H5H6H7H8H9H!H#H$H%HQe'H3i(H)H*HCMDMFM]IReGMHMIMJMKMLMMMNMOMPMQM+HSehfRM*j,HSMTMUMVMWMXMYM-H.H0M1M2M4MTe5M6M^I_I7M8M9M!M#M+jUeif$M,j-j%M'M(MVeWe)M*M,M-M.M/H:H:M;M=M?M;H=H[M]M^M`M`IXe{M|M}M~MaNbNcNdNeNfN?HYejfgN.j@HhNiNjNkNlNmNnNoNpNqNrNsNtNuNvNwNyNzN[HZe]H^H_H`H{H|H}H~HaI0ebI4icIdIeI1e2ekflf!d#dBNCN/j:jfI3egIhIiIjIkIlImInIoIpI4eqI5irIsItIJNKNMN{I5eNNONPNQNRNSNTNUNVNWNXNuI6enfYN;jvIZN0N1N=j?jwIxI2N3N4N5NyIzIAIBICI7N8N9N!N#NDIEI%N'N(N)N*NFIGI,N-N.N/N:NHIII=N?N@N[N]NJIKI_N`N{N|N}NLIMIaObOcOdOeONIOIPIgOhOiOjOkOQIRISImO7cc 9cTI9bUI.dnOVIWIXIYIoOpOZItj!c_P`P{P|P~PaQbQcQfQ0IhQiQtO,b6h6i7i8i9iuOvOwOxOyOzOAOBOCODOEOFO!i#iGOHOIOJOKOLOMONOOOPOQOROSOTOUOVOWOXOYOZO0O1O2O3O4O5O6O7O8O]e

238 return options 2KJHf8ee f QJRJSJ]b^b_bTJ`b{b|b}bUJVJ~b!e#e$eg h i j abbbcbdbXJYJZJXf2J%eZf3J1f6J7J8J!J$J'J)J+J)e*eBd-c.c%K'K(K)K*K+K+eoGpG-K.KqGrGsGtGVc,eed-euGbLvGcLwGdLxGeLyGfLzGgLAGhLBGiL9j!jzLALBLUc#b$bGLMLOLXLFeYLZL?I@I0L1L2L3L4L#jGecf5L$j%j6L7L8L$L)L[I]IVM4MTe5M6M^I_I7M8M9M!M#M+jUeif$M,j-j%M'M(M,M:M`I{I2N8N'N-N?N`NbOhO,b6huOvOwOxOyOzOAOBOCODOEOFO!i#iGOHOIOJOKOLOMONOOOPOQOROSOTOUOVOWOXOYOZO0O1O2O3O4O5O6O7O8O]e

239 return cls() 2ofiJpfqfrfsftf8 vbve9 %d! we9PjJxeyewbzekJxbAelJybzb!PAbBbCb# $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ ufvfwfxfyfDbzfAfBfEbFbGbHbIb?c@c[c]c^c_c`c{cmJnJqJrJuJvJyJzJCJDJGJHJCfDf'd7eEfFfGf|c(d8ee f JbqbKb]b^b_b`b{b|b}b~bIfJfLbMbNbObPbd 5b!e#e$eKfLfMfg h i j abbbcbdbrbsbtbNfOfPfQfRfSfTfUfebQbRbVfSbTbUbBeVbfbWJWf0J1J%eYf0f4J5J2f3f9J#J%J(J*J'ek 7 KiLiljmjMiNinjoj7jbJKP}c*d~cadWbXbbdcdYbZb0b1b2b3b#P$P%P'P(P)P*P+PgGRchGNKOKPKQKRKSKTKScUKVKWKXKiGjGkGCeYKZK0K1K2K3K4K5K6K7K8K9K!K#KBd-c.c/P$KgbacbcTcIcJcse'chbLcib(cMcjb6b^e_e`el -bm n o p q r s t u .bv w x y %bz A 'b/bte)ckbNclb*cOcmb7b{e|e}eB :bC D E F G H I J K ;bL M N O (bP Q )b=bue+cnbPcob,cQcpb8b~eafbfR ?bS T U V W X Y Z 0 @b1 2 3 4 *b5 6 +b[bddVc/efdCGDGEGFG+dGGHGIGJGKGUcLG4bMGNGOGPGQGRGSG9OTGUG,dVGWGXGYG3c4c5c#b$b2G3G4G5G6G7G8G9G!G#G$G%G'G(G)G*G+G,G-G.G/G:G;G=G?G@G[G]G^G_G`G{G|G}G~GaHbHcHdHeHfHgHhHiHjHkHlHmHnHoHpHqHrHsHtHuHvHwHxHyHzHAHBHCHDHgdHLILJLEHFHLLGHNLPLQLRLHHIHTLULVL?I@I#j%jHeIe9L!L%L'LJHKH*L+L,LLHMH.L/L:L=L[IJe?L@L[L]L^L_L`L{L|L}LNHKedf~L'jOHaMbMcMdMeMfMgMhMiMjMkMlMmMnMoMpMrMsMPHLeQHRHSHTHUHVHWHXHYHMeZH2i0H1H2HNeOeefff6d7duMvM(j)j3HPe4H5H6H7H8H9H!H#H$H%HQe'H3i(H)H*HCMDMFM]IReGMHMIMJMKMLMMMNMOMPMQM+HSehfRM*j,HSMTMUMWMXMYM-H.H0M1M2M^I_I+j-jVeWe)M*M-M.M/H:H;M=M?M;H=H[M]M^M`M`IXe{M|M}M~MaNbNcNdNeNfN?HYejfgN.j@HhNiNjNkNlNmNnNoNpNqNrNsNtNuNvNwNyNzN[HZe]H^H_H`H{H|H}H~HaI0ebI4icIdIeI1e2ekflf!d#dBNCN/j:jfI3egIhIiIjIkIlImInIoIpI4eqI5irIsItIJNKNMN{I5eNNONPNQNRNSNTNUNVNWNXNuI6enfYN;jvIZN0N1N=j?jwIxI3N4N5NyIzIAIBICI7N9N!N#NDIEI%N(N)N*NFIGI,N.N/N:NHIII=N@N[N]NJIKI_N{N|N}NLIMIaOcOdOeONIOIPIgOiOjOkOQIRISImO7cc 9cTI9bUI.dnOVIWIXIYIoOpOZItj!c_P`P{P|P~PaQbQcQfQ0IhQiQtO,b6h6i7i8i9i]e

240 elif isinstance(options, cls): 28 vbve9 ! wexeyewbzexbAeybzbAbBbCb# $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ DbEbFbGbHbIbBgCgDgvdEgwdFgxdGgydHgzdIgAd!bJgKgLgMgNgLJMJNJ|c(dOJPJe f cGOgJbqbKbPgLbMbNbObPbd g h i j abbbcbdbrbsbtbebQbRbSbTbUbBeVbfb,J-JdG.J)d/J:J;J=J?JeG@JfG[J]J^Jk _J7 `JKiLiljmjMiNinjoj7jbJKP{J|J}J~JaKbKcKdKeKfKgKhKiKjKkKlKmKnKoKpKqKrKsKtKuKvKwKxKyKzKAK(eBKCKDKEKFKGKWbXbHKIKJKKKLKMKYbZb0b1b2b3b2IgGRchG,P-PSciGjGkGCe6I7I8IBd8j-clG.cmG9I.P!I#I$I%I'I(I)I*I+I,I-I.I6P7P8P:P;P=P?P@P[P~FgbCdTcDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdUdVdWdXdYdZd0d1d2d3d4dDeEeb cc`Odc{Oec|Ofc}Ogc~OhcaPicbPjccPkcdPlcePmcfPnc~IocaJpcgPqchPrciPscjPtckPuclPvcmPwcnPxcoPycpPzcqPAcrPBcsPCc5fDctPEcuPFcvPGcwPnGxPHcdd,KoGpGqGrGsGtGVced/K:K;K=K?K@K[K]K^K_K`K{K|K}K~KaLuGvGwGxGyGzGAGBGjLkLlLmLnLoLpLqLrLsLtLuLvLwLxLyL/I9j:I;I!j=I.efd:e;eCLDLELFLWcXcubYcZcKcUc0c1c2cZG0G1G#b$bJh=e?egd-dQgRgSgTgUgVgFeGeHeIeWgXgYgZg0g1g2g3g4g5g6g7g8gJeKe9g!g#gLeMe$gNeOe%g'g(g)g*g5d+g,g/c:c-g8dPeQe.g/gReSe:g;g=g?gTeUeVeWe@g[g]g^g_g`g{g|g}g~gahbhchXeYedhehfhZe0egh1e2ehhihjhkhlh9dmhnh;c=coh$d3e4ephqh5e6e@e[e=j?jrhshthuhvhwhxhyhzhAhBhChDhEhFhGh)O*O6c|I/d8c:d;d=dhdidjd?d@d[d]d$c%ckd^dld_d`d{d|dmd}dndodpdqd~daebecedeeefegeherdsdtd+O,O-O.O/O:O;O=O?O@O[O]O^O_O@j[jud}ItjyPzPAPBPCPDPEPFPGPHP#ciejekelemeneoe!O#O$O%Ope'Oqere(O]P^PqO}PdQeQgQ0IrOsOjQ]j^j_j`j{j|j}j~jakbkckdk,bek6h6i7i8ifkgkhkik9ijkkklkmknkok!ipk#iqkrk4fsktk

241 return options 28 vbve9 ! wexeyewbzexbAeybzbAbBbCb# $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ DbEbFbGbHbIbBgCgDgvdEgwdFgxdGgydHgzdIgAd!bJgKgLgMgNgLJMJNJ|c(dOJPJe f cGOgJbqbKbPgLbMbNbObPbd g h i j abbbcbdbrbsbtbebQbRbSbTbUbBeVbfb,J-JdG.J)d/J:J;J=J?JeG@JfG[J]J^Jk _J7 `JljmjnjojKP{J|J}J~JaKbKcKdKeKfKgKhKiKjKkKlKmKnKoKpKqKrKsKtKuKvKwKxKyKzKAK(eBKCKDKEKFKGKWbXbHKIKJKKKLKMKYbZb0b1b2b3bgGRchG,P-PSciGjGkGCe6I7I8IBd8j-clG.cmG9I.P!I#I$I%I'I(I)I*I+I,I-I.I6P7P8P:P;P=P?P@P[P~FgbCdTcDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdUdVdWdXdYdZd0d1d2d3d4dDeEeb cc`Odc{Oec|Ofc}Ogc~OhcaPicbPjccPkcdPlcePmcfPnc~IocaJpcgPqchPrciPscjPtckPuclPvcmPwcnPxcoPycpPzcqPAcrPBcsPCc5fDctPEcuPFcvPGcwPnGxPHcdd,KoGpGqGrGsGtGVced/K:K;K=K?K@K[K]K^K_K`K{K|K}K~KaLuGvGwGxGyGzGAGBGjLkLlLmLnLoLpLqLrLsLtLuLvLwLxLyL.efd:e;eCLDLELFLWcXcubYcZcKcUc0c1c2cZG0G1G#b$bJh=e?egd-dQgRgSgTgUgVgFeGeHeIeWgXgYgZg0g1g2g3g4g5g6g7g8gJeKe9g!g#gLeMe$gNeOe%g'g(g)g*g5d+g,g/c:c-g8dPeQe.g/gReSe:g;g=g?gTeUeVeWe@g[g]g^g_g`g{g|g}g~gahbhchXeYedhehfhZe0egh1e2ehhihjhkhlh9dmhnh;c=coh$d3e4ephqh5e6e@e[erhshthuhvhwhxhyhzhAhBhChDhEhFhGh)O*O6c|I/d8c:d;d=dhdidjd?d@d[d]d$c%ckd^dld_d`d{d|dmd}dndodpdqd~daebecedeeefegeherdsdtd+O,O-O.O/O:O;O=O?O@O[O]O^O_O@j[j}ItjyPzPAPBPCPDPEPFPGPHP#ciejekelemeneoe!O#O$O%Ope'Oqere(O]P^PqO}PdQeQgQ0IrOsOjQ]j^j_j`j{j|j}j~jakbkckdk,bek6h6i7i8ifkgkhkik9ijkkklkmknkok!ipk#iqkrk4fsktk

242 elif isinstance(options, dict): 2KiLiMiNi7jbJ2I~FCd/I9j:I;I!j=I=j?jud4f

243 return cls(**options) 2KiLiMiNi7jbJCd/I9j:I;I!j=I=j?jud

244 else: 

245 raise TypeError( 22I~F4f

246 f"The {options_description} must be provided as an object " 22I~F4f

247 f"of type {cls.__name__} or as a dict with valid {options_description}. " 22I~F4f

248 f"The provided object is '{options}'." 22I~F4f

249 ) 

250  

251  

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

253 """ 

254 Convert a boolean option to a string representation. 

255 """ 

256 return "true" if bool(option) else "false" 2b ccdcecfcgchcicjckclcmcncocpcqcrcsctcucvcwcxcyczcAcBcCcDcEcFcGcHcKc#b$bhdidjdldpdqdrdsdtdv9

257  

258  

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

260 """ 

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

262  

263 Args: 

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

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

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

267  

268 Returns: 

269 Callable: A decorator that creates the wrapping. 

270 """ 

271  

272 def outer(wrapped_function: Callable) -> Callable: 23I

273 """ 

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

275 """ 

276  

277 @functools.wraps(wrapped_function) 23I

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

279 """ 

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

281 """ 

282 checker(*args, **kwargs, what=what) 23I

283 result = wrapped_function(*args, **kwargs) 23I

284  

285 return result 23I

286  

287 return inner 23I

288  

289 return outer 23I

290  

291  

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

293 """ 

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

295 """ 

296 return isinstance(obj, Sequence) 2k 7 RcSc6I7I8I8j-c.c9I!I#I$I%I'I(I)I*I+I,I-I.I6P7P8PgbCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdUdVdWdXdYdZd0d1d2d3d4d~IaJ|I$c%ckdmdndodw9x9y9z9A9@j[jB9C9D9h9i9}I7h8h9h!h#h$hE9F9

297  

298  

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

300 """ 

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

302 """ 

303 return is_sequence(obj) and any(is_sequence(elem) for elem in obj) 2$c%ch9i9

304  

305  

306  

307class Transaction: 

308 """ 

309 A context manager for transactional operations with undo capability. 

310  

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

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

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

314  

315 Usage: 

316 with Transaction() as txn: 

317 txn.append(some_cleanup_function, arg1, arg2) 

318 # ... perform operations ... 

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

320  

321 Methods: 

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

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

324 """ 

325 def __init__(self) -> None: 

326 self._stack = ExitStack() 2WcXcubYcZc

327 self._entered = False 2WcXcubYcZc

328  

329 def __enter__(self): 

330 self._stack.__enter__() 2WcXcubYcZc

331 self._entered = True 2WcXcubYcZc

332 return self 2WcXcubYcZc

333  

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

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

336 self._entered = False 2WcXcubYcZc

337 return self._stack.__exit__(exc_type, exc, tb) 2WcXcubYcZc

338  

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

340 """ 

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

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

343 """ 

344 if not self._entered: 2WcXcubYcZc

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

346 self._stack.callback(partial(fn, *args, **kwargs)) 2WcXcubYcZc

347  

348 def commit(self) -> None: 

349 """ 

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

351 """ 

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

353 self._stack.pop_all() 2WcXcubYcZc

354  

355  

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

357_fork_warning_checked = False 

358  

359  

360def reset_fork_warning() -> None: 

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

362  

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

364 to check the warning behavior. 

365 """ 

366 global _fork_warning_checked 

367 _fork_warning_checked = False 2Jh=e?egdhj-d

368  

369  

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

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

372 cdef unsigned int val 

373 if width == 1: 2VjWjXjYjZj0j1j2j3j4j5j6jbGd 'chbLcib(cMcjbl -bm n o p q r s t u .bv w x y %bz A 'b/b)ckbNclb*cOcmbB :bC D E F G H I J K ;bL M N O (bP Q )b=b+cnbPcob,cQcpbR ?bS T U V W X Y Z 0 @b1 2 3 4 *b5 6 +b[b

374 val = (<uint8_t*>ptr)[0] 2hbm n o v w kbC D E L M nbS T U 1 2

375 elif width == 2: 

376 val = (<uint16_t*>ptr)[0] 2VjXjZj1j3j5jLcibp q r x y %bNclbF G H N O (bPcobV W X 3 4 *b

377 elif width == 4: 

378 val = (<uint32_t*>ptr)[0] 2WjYj0j2j4j6jd Mcjbl s t u z A 'bOcmbB I J K P Q )bQcpbR Y Z 0 5 6 +b

379 else: 

380 raise ValueError(f"value must be 1, 2, or 4 bytes, got {width}") 2bG'c(c-b.b/b)c*c:b;b=b+c,c?b@b[b

381 return (val, <unsigned int>width) 2VjWjXjYjZj0j1j2j3j4j5j6jd hbLcibMcjbl m n o p q r s t u v w x y %bz A 'bkbNclbOcmbB C D E F G H I J K L M N O (bP Q )bnbPcobQcpbR S T U V W X Y Z 0 1 2 3 4 *b5 6 +b

382  

383  

384cpdef tuple _parse_fill_value(value): 

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

386  

387 Parameters 

388 ---------- 

389 value : int or buffer-protocol object 

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

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

392  

393 Returns 

394 ------- 

395 tuple of (int, int) 

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

397  

398 Raises 

399 ------ 

400 OverflowError 

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

402 TypeError 

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

404 ValueError 

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

406 """ 

407 cdef uint8_t byte_val 

408 cdef Py_buffer buf 

409  

410 if isinstance(value, int): 2?c@c[c]c^c_c`c{coJpJVjWjsJtJXjYjwJxJZj0jAJBJ1j2jEJFJ3j4jIJJJ5j6jbGe f ]b^b_b`b{b|b}b~bd 9eg h i j se'chbLcib(cMcjb6b^e_e`el -bm n o p q r s t u .bv w x y %bz A 'b/bte)ckbNclb*cOcmb7b{e|e}eB :bC D E F G H I J K ;bL M N O (bP Q )b=bue+cnbPcob,cQcpb8b~eafbfR ?bS T U V W X Y Z 0 @b1 2 3 4 *b5 6 +b[bKLSLWL$j#L(L-L;L'jqMtM2i5dgf8d(j)jwMxMyMzMAMBM3iEM*jZM3M,j+M/M@M_M.jxNAN4i9dmf$d/j:jDNENFNGNHNIN5iLN;j6N$N+N;N^N~NfOlO

411 byte_val = value 2?c@c[c]c^c_c`c{coJpJsJtJwJxJAJBJEJFJIJJJe f ]b^b_b`b{b|b}b~b9eg h i j 6b^e_e`e7b{e|e}e8b~eafbfKLSLWL$j#L(L-L;L'jqMtM2i5dgf8d(j)jwMxMyMzMAMBM3iEM*jZM3M,j+M/M@M_M.jxNAN4i9dmf$d/j:jDNENFNGNHNIN5iLN;j6N$N+N;N^N~NfOlO

412 return (<unsigned int>byte_val, <unsigned int>1) 2?c@c[c]c^c_c`c{coJpJsJtJwJxJAJBJEJFJIJJJe f ]b^b_b`b{b|b}b~b9eg h i j 6b7b8bKLSLWL$j#L(L-L;L'jqMtM2i5dgf8d(j)jwMxMyMzMAMBM3iEM*jZM3M,j+M/M@M_M.jxNAN4i9dmf$d/j:jDNENFNGNHNIN5iLN;j6N$N+N;N^N~NfOlO

413  

414 if isinstance(value, bytes): 2VjWjXjYjZj0j1j2j3j4j5j6jbGd se'chbLcib(cMcjbl -bm n o p q r s t u .bv w x y %bz A 'b/bte)ckbNclb*cOcmbB :bC D E F G H I J K ;bL M N O (bP Q )b=bue+cnbPcob,cQcpbR ?bS T U V W X Y Z 0 @b1 2 3 4 *b5 6 +b[b

415 return _read_fill_ptr(<const char*><bytes>value, <Py_ssize_t>len(value)) 2VjWjXjYjZj0j1j2j3j4j5j6jbG'chbLcib(cMcjb)ckbNclb*cOcmb+cnbPcob,cQcpb

416  

417 if PyObject_GetBuffer(value, &buf, PyBUF_SIMPLE) != 0: 2d sel -bm n o p q r s t u .bv w x y %bz A 'b/bteB :bC D E F G H I J K ;bL M N O (bP Q )b=bueR ?bS T U V W X Y Z 0 @b1 2 3 4 *b5 6 +b[b

418 raise TypeError( 

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

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

421 ) 

422 try: 2d l -bm n o p q r s t u .bv w x y %bz A 'b/bB :bC D E F G H I J K ;bL M N O (bP Q )b=bR ?bS T U V W X Y Z 0 @b1 2 3 4 *b5 6 +b[b

423 return _read_fill_ptr(<const char*>buf.buf, buf.len) 2d l -bm n o p q r s t u .bv w x y %bz A 'b/bB :bC D E F G H I J K ;bL M N O (bP Q )b=bR ?bS T U V W X Y Z 0 @b1 2 3 4 *b5 6 +b[b

424 finally: 

425 PyBuffer_Release(&buf) 

426  

427  

428def check_multiprocessing_start_method() -> None: 

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

430 global _fork_warning_checked 

431 if _fork_warning_checked: 2LPMPdG)dNPOPPPQPeGfGk 7 KiLiljmjMiNinjojRPSPj9TPk9l9UPm9n9VPo9p9WPq9r9XPs9t9YPu9ZP0P1Ppj2P3P}c*d~cadWbXbbdcdHh4PIh5PYbZb0b1b2b3bJh=e?egdhj-d

432 return 2LPMPdG)dNPOPPPQPeGfGk 7 KiLiljmjMiNinjojRPSPj9TPk9l9UPm9n9VPo9p9WPq9r9XPs9t9YPu9ZP0P1Ppj2P3P}c*d~cadWbXbbdcdHh4PIh5PYbZb0b1b2b3b-d

433 _fork_warning_checked = True 2pjJh=e?egdhj-d

434  

435 # Common warning message parts 

436 common_message = ( 

437 "CUDA does not support. Forked subprocesses exhibit undefined behavior, " 2pjJh=e?egdhj-d

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

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

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

441 ) 

442  

443 try: 2pjJh=e?egdhj-d

444 start_method = multiprocessing.get_start_method() 2pjJh=e?egdhj-d

445 if start_method == "fork": 2pjJh=e?egd-d

446 message = f"multiprocessing start method is 'fork', which {common_message}" 2=e?egd-d

447 warnings.warn(message, UserWarning, stacklevel=3) 2=e?egd-d

448 except RuntimeError: 2hj

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

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

451 if platform.system() == "Linux": 2hj

452 message = ( 

453 f"multiprocessing start method is not set and defaults to 'fork' on Linux, " 2hj

454 f"which {common_message}" 

455 ) 

456 warnings.warn(message, UserWarning, stacklevel=3) 2hj