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
« 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
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
15from cuda.bindings import driver as driver, nvrtc as nvrtc, runtime as runtime
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'
25from cuda.bindings.nvvm import nvvmError
26from cuda.bindings.nvjitlink import nvJitLinkError
28from cpython.buffer cimport PyObject_GetBuffer, PyBuffer_Release, Py_buffer, PyBUF_SIMPLE
30from cuda.bindings cimport cynvrtc, cynvvm, cynvjitlink
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
36class CUDAError(Exception):
37 pass
40class NVRTCError(CUDAError):
41 pass
45ComputeCapability = namedtuple("ComputeCapability", ("major", "minor"))
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
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
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
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
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
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
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
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
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
159cdef object _RUNTIME_SUCCESS = runtime.cudaError_t.cudaSuccess
160cdef object _NVRTC_SUCCESS = nvrtc.nvrtcResult.NVRTC_SUCCESS
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()}")
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
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
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}")
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
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 )
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
259def precondition(checker: Callable[..., None], str what="") -> Callable[..., Any]:
260 """
261 A decorator that adds checks to ensure any preconditions are met.
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.
268 Returns:
269 Callable: A decorator that creates the wrapping.
270 """
272 def outer(wrapped_function: Callable) -> Callable: 23I
273 """
274 A decorator that actually wraps the function for checking preconditions.
275 """
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
285 return result 23I
287 return inner 23I
289 return outer 23I
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
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
307class Transaction:
308 """
309 A context manager for transactional operations with undo capability.
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.
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
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
329 def __enter__(self):
330 self._stack.__enter__() 2WcXcubYcZc
331 self._entered = True 2WcXcubYcZc
332 return self 2WcXcubYcZc
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
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
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
356# Track whether we've already warned about fork method
357_fork_warning_checked = False
360def reset_fork_warning() -> None:
361 """Reset the fork warning check flag for testing purposes.
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
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
384cpdef tuple _parse_fill_value(value):
385 """Parse a fill/memset value into (raw_value, element_size).
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.
393 Returns
394 -------
395 tuple of (int, int)
396 (raw_value, element_size) where element_size is 1, 2, or 4.
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
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
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
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)
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
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 )
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