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

233 statements  

« prev     ^ index     » next       coverage.py v7.15.2, created at 2026-07-19 01:12 +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 vbse9 ! teuevewbwexbxeybzbAbBbCb# $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ DbEbFbGbHbIbngogpgfjtdqgYhudrgvdsgwdtgxdugyd9bvgwgxgygzge f AgJbqbKbBgLbMbNbObPbd g h i j abbbcbdbrbsbtbebQbRbSbTbUbyeVbfb6 7 ScTczezd,c-cgbAd:iOIPIQIgjwhRIhjUcBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdUdVdWdXdYdZd0d1d2dAeBeijSIjjVc!b#bCgDgkjEgFgljGgHgCeDe`eEeFeIgJgKgLgMgNgOgPgQgZhRgSg0hTgUgGeHe{eVgWgXgIeJeYgKeLeZg0g1g2g3g3d4g5g1h.c2h3h|e}e4h5h6h7h8h4d/c9h~e5d!h#h$h6g6dMeNe7g8gOePeaf9g!gmj#g$gQeRebfSeTe%g'g(g)g*g+g,g-g.g%h/g:g'h;g=gUeVecf?g@g[gWeXe]gYeZe^g_g`g{g|g7d}g~g(h:c)h*hdfef+h,h-h.h/h8d;c:hff9d;h=h?hah!d0e1ebhch2e3egfdhehnjfhghojhhihpjjhkhqjlhmhrjnhohsjphqhtjrhshuj;ifIgIKFTI=i

50 if isinstance(cfg, int): 28 vbse9 ! teuevewbwexbxeybzbAbBbCb# $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ DbEbFbGbHbIbngogpgfjtdqgYhudrgvdsgwdtgxdugyd9bvgwgxgygzge f AgJbqbKbBgLbMbNbObPbd g h i j abbbcbdbrbsbtbebQbRbSbTbUbyeVbfb6 7 ScTczezd,c-cgbAd:iOIPIQIgjwhRIhjUcBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdUdVdWdXdYdZd0d1d2dAeBeijSIjjVc!b#bCgDgkjEgFgljGgHgCeDe`eEeFeIgJgKgLgMgNgOgPgQgZhRgSg0hTgUgGeHe{eVgWgXgIeJeYgKeLeZg0g1g2g3g3d4g5g1h.c2h3h|e}e4h5h6h7h8h4d/c9h~e5d!h#h$h6g6dMeNe7g8gOePeaf9g!gmj#g$gQeRebfSeTe%g'g(g)g*g+g,g-g.g%h/g:g'h;g=gUeVecf?g@g[gWeXe]gYeZe^g_g`g{g|g7d}g~g(h:c)h*hdfef+h,h-h.h/h8d;c:hff9d;h=h?hah!d0e1ebhch2e3egfdhehnjfhghojhhihpjjhkhqjlhmhrjnhohsjphqhtjrhshuj;ifIgIKFTI=i

51 cfg = (cfg,) 28 vbse9 ! teuevewbwexbxeybzbAbBbCb# $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ DbEbFbGbHbIbngogpgqgrgsgtgug9bvgwgxgygzge f AgJbqbKbBgLbMbNbObPbd g h i j abbbcbdbrbsbtbebQbRbSbTbUbyeVbfb6 7 ScTczezd,c-cgbAd:iOIPIQIwhRIBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdUdVdWdXdYdZd0d1d2dAeBeSIVc!b#bCgDgEgFgGgHgCeDeEeFeIgJgKgLgMgNgOgPgQgRgSgTgUgGeHeVgWgXgIeJeYgKeLeZg0g1g2g3g3d4g5g.c/c6g6dMeNe7g8gOePe9g!g#g$gQeReSeTe%g'g(g)g*g+g,g-g.g/g:g;g=gUeVe?g@g[gWeXe]gYeZe^g_g`g{g|g7d}g~g:c;cah!d0e1ebhch2e3edhehfhghhhihjhkhlhmhnhohphqhrhsh;ifI

52 else: 

53 common = "must be an int, or a tuple with up to 3 ints" 2fjtdYhudvdwdxdyd:igjwhhjUcijjjkjlj`eZh0h{e1h.c2h3h|e}e4h5h6h7h8h4d/c9h~e5d!h#h$hafmjbf%h'hcf(h:c)h*hdfef+h,h-h.h/h8d;c:hff9d;h=h?hgfnjojpjqjrjsjtjuj;igIKFTI=i

54 if not isinstance(cfg, tuple): 2fjtdYhudvdwdxdyd:igjwhhjUcijjjkjlj`eZh0h{e1h.c2h3h|e}e4h5h6h7h8h4d/c9h~e5d!h#h$hafmjbf%h'hcf(h:c)h*hdfef+h,h-h.h/h8d;c:hff9d;h=h?hgfnjojpjqjrjsjtjuj;igIKFTI=i

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

56 if len(cfg) > 3: 2fjtdYhudvdwdxdyd:igjwhhjUcijjjkjlj`eZh0h{e1h.c2h3h|e}e4h5h6h7h8h4d/c9h~e5d!h#h$hafmjbf%h'hcf(h:c)h*hdfef+h,h-h.h/h8d;c:hff9d;h=h?hgfnjojpjqjrjsjtjuj;igIKF=i

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

58 if any(not isinstance(val, int) for val in cfg): 2fjtdYhudvdwdxdyd:igjwhhjUcijjjkjlj`eZh0h{e1h.c2h3h|e}e4h5h6h7h8h4d/c9h~e5d!h#h$hafmjbf%h'hcf(h:c)h*hdfef+h,h-h.h/h8d;c:hff9d;h=h?hgfnjojpjqjrjsjtjuj;iKF=i

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

60 if any(val < 1 for val in cfg): 28 vbse9 ! teuevewbwexbxeybzbAbBbCb# $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ DbEbFbGbHbIbngogpgfjtdqgYhudrgvdsgwdtgxdugyd9bvgwgxgygzge f AgJbqbKbBgLbMbNbObPbd g h i j abbbcbdbrbsbtbebQbRbSbTbUbyeVbfb6 7 ScTczezd,c-cgbAd:iOIPIQIgjwhRIhjUcBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdUdVdWdXdYdZd0d1d2dAeBeijSIjjVc!b#bCgDgkjEgFgljGgHgCeDe`eEeFeIgJgKgLgMgNgOgPgQgZhRgSg0hTgUgGeHe{eVgWgXgIeJeYgKeLeZg0g1g2g3g3d4g5g1h.c2h3h|e}e4h5h6h7h8h4d/c9h~e5d!h#h$h6g6dMeNe7g8gOePeaf9g!gmj#g$gQeRebfSeTe%g'g(g)g*g+g,g-g.g%h/g:g'h;g=gUeVecf?g@g[gWeXe]gYeZe^g_g`g{g|g7d}g~g(h:c)h*hdfef+h,h-h.h/h8d;c:hff9d;h=h?hah!d0e1ebhch2e3egfdhehnjfhghojhhihpjjhkhqjlhmhrjnhohsjphqhtjrhshuj;ifI=i

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

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

63 return cfg + (1,) * (3 - len(cfg)) 28 vbse9 ! teuevewbwexbxeybzbAbBbCb# $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ DbEbFbGbHbIbngogpgfjtdqgYhudrgvdsgwdtgxdugyd9bvgwgxgygzge f AgJbqbKbBgLbMbNbObPbd g h i j abbbcbdbrbsbtbebQbRbSbTbUbyeVbfb6 7 ScTczezd,c-cgbAd:iOIPIQIgjwhRIhjUcBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdUdVdWdXdYdZd0d1d2dAeBeijSIjjVc!b#bCgDgkjEgFgljGgHgCeDe`eEeFeIgJgKgLgMgNgOgPgQgZhRgSg0hTgUgGeHe{eVgWgXgIeJeYgKeLeZg0g1g2g3g3d4g5g1h.c2h3h|e}e4h5h6h7h8h4d/c9h~e5d!h#h$h6g6dMeNe7g8gOePeaf9g!gmj#g$gQeRebfSeTe%g'g(g)g*g+g,g-g.g%h/g:g'h;g=gUeVecf?g@g[gWeXe]gYeZe^g_g`g{g|g7d}g~g(h:c)h*hdfef+h,h-h.h/h8d;c:hff9d;h=h?hah!d0e1ebhch2e3egfdhehnjfhghojhhihpjjhkhqjlhmhrjnhohsjphqhtjrhshuj;i

64  

65  

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

67 if err != cydriver.CUresult.CUDA_SUCCESS: 2a WPUIXPhfYPifZPjf0Pkf1Plf2P8 3Pvb4Pse5P9 6P#d7P! 8Pte9P!PVI#Pue$Pve%Pwb'Pwe(PWI)Pxb*Pxe+PXI,Pyb-Pzb.P/PAb:PBb;PCb=P# ?P$ @P% [P' ]P( ^P) _P* `P+ {P, |P- }P. ~P/ aQ: bQ; cQ= dQ? eQ@ fQ[ gQ] hQ^ iQ_ jQ` kQ{ lQ| mQ} nQ~ oQpQqQrQsQtQuQvQwQxQ@hyQzQAQBQCQDQEQFQGQHQIQJQKQLQMQNQOQPQQQRQSQTQUQVQWQXQYQmfZQ0Q1Q2Qnf3Qof4Qpf5Qqf6QDb7Qrf8Qsf9Qtf!Q#Q$Q%Q'QEb(QFb)QGb*QHb+QIb,Q=c-Q?c.Q@c/Q[c:Q]c;Q^c=Q_c?Q`c@Q[Q]Q^Q_Q`Q[h{Q]h|Q^h}Q~QaRbRcRdReRngfRoggRhRiRjRpgkRYIZIlRmRnRoRpRtdqR0I1IvjwjrRsRtRuRvRwRxRqgyRzR2IAR3IBRCRDRERFRGRHRIRJRYhudKRLRMR4INR5IORxjPRyjQRRRSRTRURVRWRXRYRZR0R1R2R3R4R5R6R7R8R9R!R#R$R%R'R(Rrg)R6I7I*R+R,R-R.Rvd/R8I9IzjAj:R;R=R?R@R[R]Rsg^R_R`R!I{R#I|R}R~RaSbScSdSeSfSgShSwdiSjSkS$IlS%ImSBjnSCjoSpSqSrSsStSuSvSwStgxSySzS'IAS(IBSCSDSESFSGSHSISJSKSLSxdMSNSOS)IPS*IQSDjRSEjSSTSUSVSWSXSYSZS0Sug1S2S3S+I4S,I5S6S7S8S9S!S#S$S%S'S(Syd)S*S+S-I,S.I-SFj.SGj/S:S;S=S?S@S[S]S^S_S`S{S|S}S~SaTbTcTdTeTfTgThTiTjTkTlTmTnToTpTqTrTsTtTuTvTwTxTyTzTAT_hBTCTDTETFTGTHTITJTKTLTMTNTOTPT`hQTRTSTTTUTLFVT9bWTXTYT{hZTuf0T1T2T3T4T5T6T7T8T9T!T#T$T%T'T(T)T*T+T,T-T.T/Tvg:Twg;Txg=Tyg?Tvf@Twf[Tzg]T/I^T:I_T;I`T{c{T$d|T=I}T?I~Te aUf bUMFcUAgdUJbeUqbfUKbgU@IhU[IiU]IjU[bkU]blU^bmU^InU_boU`bpU{bqU|brU_IsU`ItU}buUBgvUwUxUyUxfzU|hAU{IBUCUDULbEUMbFUNbGUObHUPbIUJUKULUMUNUOUPUQUd RU}hSU4eTU~hUU5eVU6eWU7eXUyfYUzfZUAf0Ug 1Uh 2Ui 3Uj 4Uab5Ubb6Ucb7Udb8Urb9Usb!Utb#UBf$UCf%UDf'UEf(UFf)UGf*U+U,U-Ueb.UQb/URb:USb;UTb=UUb?Uye@UVb[Ufb]U|I}IHf~I^Uaibicidieifigihi_Uii`UaJxh{UbJIfcJdJ8e|UJfKfLfeJfJMfgJhJyh}U~UjikiliNfiJOfjJkJlJmJnJoJpJqJrJsJtJ9eaV-OuJ.OvJNFwJOFxJ/OyJ:OzJ;OAJ=OBJPFCJQFDJEJFJ6 GJ7 HJ~bbVmicV?idV@ieVnifVoigV[ihV]iiVHjjVkVlV?OIJ@OJJKJLJMJ[ONJmVOJPJQJ]ORJnVSJTJUJ^OVJoVWJXJYJ_OZJpV0J1J2J`O3JqV4J5J6J{O7JrV8J9J|O!J}O#J~O$JaP%JbP'JcP(J%d)J|c*JsVtVuV}c+J~c,JWbvVXbwVad-Jbd.JdP/JeP:JfP;JgP=JYbxVZbyV0bzV1bAV2bBV3bCVzhDVEVFVGVHVIVJVKVAh!e#eLVMVNVOVPVQVRVSVTVUVBhVVWVXVYVZV0V1V2V3V4V5V6V7V8V9V!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~WaXbXcXdXeXfXgXhXiXjXkXlXmXnXoXpXpiqXrXsXqiChDhriRFScSFtX?JuX@JvX[JwX]JxX^JyX_JzX`JAXTcBX{JCX|JDX}JEX~JFXTFGXUFHXVFIXzeJXaKbKcKdKeKKXLXMXNXOXPXfKgKhKiKQXjKkKRXlKmKsiSXTXUXhIVXWXiItinKjIuizdIjvi,cWF-cXFXXkIwilIximIyiYXnIziZXoIAioKpIBi0X1X2X3X4X5XqI6XrI7XsI8XtI9XuI!XvI#XwI$X%X'X(X)X*X+X,X-X.X/X:XgbacbcAd;X=X?XUc@XBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdUdVdWdXdYdZd0d1d2dAeBeccdcecfcgchcicjckclcmcncocpcqcrcsctcucvcwcxcyczcAcBcCcDcEcFcGcYFHc[X]X^XEhIcJcFh_XGhHhIhpKqK`XJhKhLhMhCiDi{X|Xhb}X~XibaYbYcYjbdY5beYfYgYhYk iYjYl kYm lYn mYo nYp oYq pYr qYs rYt sYtYu uYv vYw wYx xYyYy zYz AYBYCYDYEYkbFYGYlbHYIYJYmbKY6bLYMYNYOYA PYQYB RYC SYD TYE UYF VYG WYH XYI YYJ ZY0YK 1YL 2YM 3YN 4Y5YO 6YP 7Y8Y9Y!Y#Ynb$Y%Yob'Y(Y)Ypb*Y7b+Y,Y-Y.YQ /Y:YR ;YS =YT ?YU @YV [YW ]YX ^YY _YZ `Y{Y0 |Y1 }Y2 ~Y3 aZbZ4 cZ5 dZeZfZNhrKsKtKgZcdhZiZuKjZvKkZlZmZnZ$eZF0FoZwKpZxKqZ1FrZ2FsZ3FtZ4FuZWcvZ%ewZddxZ'eyZyKzZzKAZAKBZBKCZCKDZDKEZEKFZFKGZGKHZHKIZIKJZJKKZKKLZLKMZMKNZNKOZ5FPZOKQZ6FRZPKSZ7FTZQKUZ8FVZRKWZ9FXZSKYZ!FZZTK0Z#F1ZUK2Z$F3ZVK4ZWK5ZXK6ZYK7ZZK8Z0K9Z1K!Z2K#Z3K$Z4K%Z5K'Z6K(Z7K)Z8K*Z9K+Z!K,Z#K-ZxI.ZJj/ZyI:ZzI;ZKj=ZAI?Z(e@Z[ZOhPh)e$K%K'K]Zed^Z*e+e(K_Z)K`Z*K{Z+K|Z%F}Z'F~ZXcYcuba0Zcb0c0d0e0f0g0h0i0j0k0l0m0n0o0Kcp0(Fq0)F'dr0*F+Fs0,Ft0-Fu0.FVcv0/F4b:Fw0;Fx0=Fy0?Fz0@FA0[FB0]FC0^FD0_FE0`FF0{F|F}F~F0c1c2caGbGcG3c4c5c!bG0#bH0I0dGJ0eGK0fGL0gGM0hGN0iGO0jGP0kGQ0lGR0mGS0nGT0oGU0pGV0qGW0rGX0sGY0tGZ0uG00vG10wG20xG30yG40zG50AGBGCG60DG70EG80FG90GG!0HG#0IG$0JG%0KG'0LG(0MG)0NG*0OG+0PG,0QG-0RG.0SG/0TG:0UG;0VG=0WG?0XGYG@0ZG[00G]01G^02G_03G`04G{05G|06G}07G~08Ga19Gb1!Gc1#Gd1$Ge1%Gf1'G(Gthg1,eh1-eQhfdi1(dj1k1,KCgl1m1n1o1-K.K/Kp1q1r1s1t1)GDgu1:K*Gv1;Kw1x1=K+G?Ky1@KEgz1A1B1C1[K]K^KD1E1F1G1H1,GFgI1_K-GJ1`KK1L1Gg{K|KM1N1O1P1HgQ1}KR1S1~KCeaLbLBICIcLdLeLfLgLLjDe`ehLMjNjiLjLkLEeFeIgJgKgLgMgT1NgU1lLmLV1W1X1Y1OgZ1nL01112131oLPg4151pLqL61718191!1.GQgZh#1rL/G$1%1'1(1sLRg)1*1+1tLuLvL,1-1.1/1:1:GSg0h;1wL;G=1xL?1@1[1TgyLzL]1^1_1`1Ug{1AL|1}1BLDIGeCLDLELFLGLHLILJLKLLL=GHe{eMLOj?GNLOLPLQLRLSLTLULVLWLXLYLZL0LVg1L2L~1a2b2c23Ld2e2f2Wgg2h24L5Li2j2k2l2Xgm26Ln2o2p2q2r2s2t2u2v2w2x2y2@GIe[G]G^G_G`G{G|G}G~GJeaHEibHcHdHYgKeLeZg0g1g2g3g3d4g5g1h.c2h3h|e}e4h5h6h7h8h4d/c9h~e5d!h#h$h6g7L8Lz2A2B2C2D26dPjQj9L!L#L$L%L'LeHMefHgHhHiHjHkHlHmHnHoHNepHFiqHrHsHE27gF2(L)LG2H2I2J28gK2*LL2M2+LEIOe,L-L.L/L:L;L=L?L@L[L]LtHPeaf^LRjuH_L`L{LN2O2|L9gP2Q2R2S2}L~LaMT2U2V2W2X2vH!gY2bMwHZ2cM0212#gdMeM22324252$g62fM7282gMQehMiMFIGIjMkMlMmMnMSjRebfoMTjUjpMqMrMSeTe%g'g(g)g*g92+g!2sMtM#2$2%2'2,g(2uM)2*2+2,2vM-g-2.2wMxM/2:2;2=2?2xH.g%h@2yMyH[2]2^2_2zM/g`2{2|2AMBMCM}2~2a3b3c3zH:g'hd3DMAHe3EMf3g3h3;gFMGMi3j3k3l3=gm3HMn3o3IMHIUeJMKMLMMMNMOMPMQMRMSMBHVecfTMVjCHUMVMWMXMYMZM0M1M2M3M4M5M6M7M?g8M9Mp3q3r3s3!Mt3u3v3@gw3x3#M$My3z3A3B3[gC3%MD3E3F3G3H3I3J3K3L3M3N3O3DHWeEHFHGHHHIHJHKHLHMHXeNHGiOHPHQH]gYeZe^g_g`g{g|g7d}g~g(h:c)h*hdfef+h,h-h.h/h8d;c:hff9d;h=h?hah'M(MP3Q3R3S3T3!dWjXj)M*M+M,M-M.MRH0eSHTHUHVHWHXHYHZH0H1H1e2HHi3H4H5HU3bhV3/M:MW3X3Y3Z3ch03;M1323=MII2e?M@M[M]M^M_M`M{M|M}M~M6H3egfaNYj7HbNcNdN334353.e/eZj0j8H9H63eNdh738393!3fNgNhN#3$3%3'3(3!Heh)3iN#H$Huh%H'Hvh*3jN+3,3-3kNfh.3/3:3;3lNmNnN=3?3@3[3]3(Hgh^3oN)H_3pN`3{3|3qNhh}3~3a4b4rNsNtNc4d4e4f4g4*Hihh4uN+Hi4vNj4k4l4wNjhm4n4o4p4xNyNzNq4r4s4t4u4,Hkhv4AN-Hw4BNx4y4z4CNlhA4B4C4D4DNENFNE4F4G4H4I4.HmhJ4GN/HK4HNL4M4N4INnhO4P4Q4R4JNKNLNS4T4U4V4W4:HohX4MN;HY4NNZ40414ONph24344454PNQNRN64748494!4=Hqh#4SN?H@H$4TN%4'4(4UNrh)4*4+4,4VNWNXN-4.4/4:4;4[Hsh=4YN]H^H?4ZN@4[4]46c^4JI7c_48c`4{4|4}4~4a5b5c5d5e5f5g5h5i5j5k5l5m5n5o5p5q5r5s5t5u5v5w5x5y5z5A5B5C5D5E5F5G5c H5I5_H8b`H)dJ5K5L5M5N5O5P5Q5R5S5T5U5V5W51j2j0N{H|H}HLcX5~H1N2NaIKI^iY5Z505152535455565759c!c8595!5#5$5%5'5(5)5*5+5,5-53N.5/5:5;5=5?5@5[5]5^5bI_5`54N{55N|5}56N~5a6b6Iic6d6e6f6g6h6i6j6k6l6m6n6o6p6q6r6s6t6u6v6w6x6y6z6A6B6C6D6E6F6G6H6I6J6K6L6M6N6O6P6Q6R63jS64jT65jU66jV67jW68jX69jY6!jZ6#j06$j16%j26'j36+b46(j56Rh66Ji76Ki8696!6#6$6%6'6Li(6)j)6*j*6+j+6,j,6Mi-6.6/6:6;67N=68N?69N@6!N[6#N]6$N^6-j_6%N`6'N{6(N|6)N}6*N~6.ja7/jb7+Nc7:jd7;je7=jf7Nig7?jh7Oii7@jj7[jk7l7m7Pfn7o7]jp7q7^jr7Pis7t7u7v7w7x7y7z7A7B7C7D7E7F7G7H7I7J7K7L7M7N7O7P7Q7R7S7T7U7,N-NV7.N/N:N;N=N?NW7X7Y7@NZ7[N07]N^N_N17`N{N|N}N~NaO27bOcO3747dO5767eO7787fOgO97!7#7hO$7%7'7(7)7*7Qf+7,7-7.7/7:7;7=7?7@7[7]7^7_7`7{7|7}7~7a8b8c8d8e8f8g8h8i8j8k8l8m8n8o8p8q8r8

68 return _check_driver_error(err) 2@h_h`h{h}h4e~h4b)d

69 return 0 2a WPUIXPhfYPifZPjf0Pkf1Plf2P8 3Pvb4Pse5P9 6P#d7P! 8Pte9P!PVI#Pue$Pve%Pwb'Pwe(PWI)Pxb*Pxe+PXI,Pyb-Pzb.P/PAb:PBb;PCb=P# ?P$ @P% [P' ]P( ^P) _P* `P+ {P, |P- }P. ~P/ aQ: bQ; cQ= dQ? eQ@ fQ[ gQ] hQ^ iQ_ jQ` kQ{ lQ| mQ} nQ~ oQpQqQrQsQtQuQvQwQxQ@hyQzQAQBQCQDQEQFQGQHQIQJQKQLQMQNQOQPQQQRQSQTQUQVQWQXQYQmfZQ0Q1Q2Qnf3Qof4Qpf5Qqf6QDb7Qrf8Qsf9Qtf!Q#Q$Q%Q'QEb(QFb)QGb*QHb+QIb,Q=c-Q?c.Q@c/Q[c:Q]c;Q^c=Q_c?Q`c@Q[Q]Q^Q_Q`Q[h{Q]h|Q^h}Q~QaRbRcRdReRngfRoggRhRiRjRpgkRYIZIlRmRnRoRpRtdqR0I1IvjwjrRsRtRuRvRwRxRqgyRzR2IAR3IBRCRDRERFRGRHRIRJRYhudKRLRMR4INR5IORxjPRyjQRRRSRTRURVRWRXRYRZR0R1R2R3R4R5R6R7R8R9R!R#R$R%R'R(Rrg)R6I7I*R+R,R-R.Rvd/R8I9IzjAj:R;R=R?R@R[R]Rsg^R_R`R!I{R#I|R}R~RaSbScSdSeSfSgShSwdiSjSkS$IlS%ImSBjnSCjoSpSqSrSsStSuSvSwStgxSySzS'IAS(IBSCSDSESFSGSHSISJSKSLSxdMSNSOS)IPS*IQSDjRSEjSSTSUSVSWSXSYSZS0Sug1S2S3S+I4S,I5S6S7S8S9S!S#S$S%S'S(Syd)S*S+S-I,S.I-SFj.SGj/S:S;S=S?S@S[S]S^S_S`S{S|S}S~SaTbTcTdTeTfTgThTiTjTkTlTmTnToTpTqTrTsTtTuTvTwTxTyTzTAT_hBTCTDTETFTGTHTITJTKTLTMTNTOTPT`hQTRTSTTTUTLFVT9bWTXTYT{hZTuf0T1T2T3T4T5T6T7T8T9T!T#T$T%T'T(T)T*T+T,T-T.T/Tvg:Twg;Txg=Tyg?Tvf@Twf[Tzg]T/I^T:I_T;I`T{c{T$d|T=I}T?I~Te aUf bUMFcUAgdUJbeUqbfUKbgU@IhU[IiU]IjU[bkU]blU^bmU^InU_boU`bpU{bqU|brU_IsU`ItU}buUBgvUwUxUyUxfzU|hAU{IBUCUDULbEUMbFUNbGUObHUPbIUJUKULUMUNUOUPUQUd RU}hSU4eTU~hUU5eVU6eWU7eXUyfYUzfZUAf0Ug 1Uh 2Ui 3Uj 4Uab5Ubb6Ucb7Udb8Urb9Usb!Utb#UBf$UCf%UDf'UEf(UFf)UGf*U+U,U-Ueb.UQb/URb:USb;UTb=UUb?Uye@UVb[Ufb]U|I}IHf~I^Uaibicidieifigihi_Uii`UaJxh{UbJIfcJdJ8e|UJfKfLfeJfJMfgJhJyh}U~UjikiliNfiJOfjJkJlJmJnJoJpJqJrJsJtJ9eaV-OuJ.OvJNFwJOFxJ/OyJ:OzJ;OAJ=OBJPFCJQFDJEJFJ6 GJ7 HJ~bbVmicV?idV@ieVnifVoigV[ihV]iiVHjjVkVlV?OIJ@OJJKJLJMJ[ONJmVOJPJQJ]ORJnVSJTJUJ^OVJoVWJXJYJ_OZJpV0J1J2J`O3JqV4J5J6J{O7JrV8J9J|O!J}O#J~O$JaP%JbP'JcP(J%d)J|c*JsVtVuV}c+J~c,JWbvVXbwVad-Jbd.JdP/JeP:JfP;JgP=JYbxVZbyV0bzV1bAV2bBV3bCVzhDVEVFVGVHVIVJVKVAh!e#eLVMVNVOVPVQVRVSVTVUVBhVVWVXVYVZV0V1V2V3V4V5V6V7V8V9V!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~WaXbXcXdXeXfXgXhXiXjXkXlXmXnXoXpXpiqXrXsXqiChDhriRFScSFtX?JuX@JvX[JwX]JxX^JyX_JzX`JAXTcBX{JCX|JDX}JEX~JFXTFGXUFHXVFIXzeJXaKbKcKdKeKKXLXMXNXOXPXfKgKhKiKQXjKkKRXlKmKsiSXTXUXhIVXWXiItinKjIuizdIjvi,cWF-cXFXXkIwilIximIyiYXnIziZXoIAioKpIBi0X1X2X3X4X5XqI6XrI7XsI8XtI9XuI!XvI#XwI$X%X'X(X)X*X+X,X-X.X/X:XgbacbcAd;X=X?XUc@XBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdUdVdWdXdYdZd0d1d2dAeBeccdcecfcgchcicjckclcmcncocpcqcrcsctcucvcwcxcyczcAcBcCcDcEcFcGcYFHc[X]X^XEhIcJcFh_XGhHhIhpKqK`XJhKhLhMhCiDi{X|Xhb}X~XibaYbYcYjbdY5beYfYgYhYk iYjYl kYm lYn mYo nYp oYq pYr qYs rYt sYtYu uYv vYw wYx xYyYy zYz AYBYCYDYEYkbFYGYlbHYIYJYmbKY6bLYMYNYOYA PYQYB RYC SYD TYE UYF VYG WYH XYI YYJ ZY0YK 1YL 2YM 3YN 4Y5YO 6YP 7Y8Y9Y!Y#Ynb$Y%Yob'Y(Y)Ypb*Y7b+Y,Y-Y.YQ /Y:YR ;YS =YT ?YU @YV [YW ]YX ^YY _YZ `Y{Y0 |Y1 }Y2 ~Y3 aZbZ4 cZ5 dZeZfZNhrKsKtKgZcdhZiZuKjZvKkZlZmZnZ$eZF0FoZwKpZxKqZ1FrZ2FsZ3FtZ4FuZWcvZ%ewZddxZ'eyZyKzZzKAZAKBZBKCZCKDZDKEZEKFZFKGZGKHZHKIZIKJZJKKZKKLZLKMZMKNZNKOZ5FPZOKQZ6FRZPKSZ7FTZQKUZ8FVZRKWZ9FXZSKYZ!FZZTK0Z#F1ZUK2Z$F3ZVK4ZWK5ZXK6ZYK7ZZK8Z0K9Z1K!Z2K#Z3K$Z4K%Z5K'Z6K(Z7K)Z8K*Z9K+Z!K,Z#K-ZxI.ZJj/ZyI:ZzI;ZKj=ZAI?Z(e@Z[ZOhPh)e$K%K'K]Zed^Z*e+e(K_Z)K`Z*K{Z+K|Z%F}Z'F~ZXcYcuba0Zcb0c0d0e0f0g0h0i0j0k0l0m0n0o0Kcp0(Fq0)F'dr0*F+Fs0,Ft0-Fu0.FVcv0/F4b:Fw0;Fx0=Fy0?Fz0@FA0[FB0]FC0^FD0_FE0`FF0{F|F}F~F0c1c2caGbGcG3c4c5c!bG0#bH0I0dGJ0eGK0fGL0gGM0hGN0iGO0jGP0kGQ0lGR0mGS0nGT0oGU0pGV0qGW0rGX0sGY0tGZ0uG00vG10wG20xG30yG40zG50AGBGCG60DG70EG80FG90GG!0HG#0IG$0JG%0KG'0LG(0MG)0NG*0OG+0PG,0QG-0RG.0SG/0TG:0UG;0VG=0WG?0XGYG@0ZG[00G]01G^02G_03G`04G{05G|06G}07G~08Ga19Gb1!Gc1#Gd1$Ge1%Gf1'G(Gthg1,eh1-eQhfdi1(dj1k1,KCgl1m1n1o1-K.K/Kp1q1r1s1t1)GDgu1:K*Gv1;Kw1x1=K+G?Ky1@KEgz1A1B1C1[K]K^KD1E1F1G1H1,GFgI1_K-GJ1`KK1L1Gg{K|KM1N1O1P1HgQ1}KR1S1~KCeaLbLBICIcLdLeLfLgLLjDe`ehLMjNjiLjLkLEeFeIgJgKgLgMgT1NgU1lLmLV1W1X1Y1OgZ1nL01112131oLPg4151pLqL61718191!1.GQgZh#1rL/G$1%1'1(1sLRg)1*1+1tLuLvL,1-1.1/1:1:GSg0h;1wL;G=1xL?1@1[1TgyLzL]1^1_1`1Ug{1AL|1}1BLDIGeCLDLELFLGLHLILJLKLLL=GHe{eMLOj?GNLOLPLQLRLSLTLULVLWLXLYLZL0LVg1L2L~1a2b2c23Ld2e2f2Wgg2h24L5Li2j2k2l2Xgm26Ln2o2p2q2r2s2t2u2v2w2x2y2@GIe[G]G^G_G`G{G|G}G~GJeaHEibHcHdHYgKeLeZg0g1g2g3g3d4g5g1h.c2h3h|e}e4h5h6h7h8h4d/c9h~e5d!h#h$h6g7L8Lz2A2B2C2D26dPjQj9L!L#L$L%L'LeHMefHgHhHiHjHkHlHmHnHoHNepHFiqHrHsHE27gF2(L)LG2H2I2J28gK2*LL2M2+LEIOe,L-L.L/L:L;L=L?L@L[L]LtHPeaf^LRjuH_L`L{LN2O2|L9gP2Q2R2S2}L~LaMT2U2V2W2X2vH!gY2bMwHZ2cM0212#gdMeM22324252$g62fM7282gMQehMiMFIGIjMkMlMmMnMSjRebfoMTjUjpMqMrMSeTe%g'g(g)g*g92+g!2sMtM#2$2%2'2,g(2uM)2*2+2,2vM-g-2.2wMxM/2:2;2=2?2xH.g%h@2yMyH[2]2^2_2zM/g`2{2|2AMBMCM}2~2a3b3c3zH:g'hd3DMAHe3EMf3g3h3;gFMGMi3j3k3l3=gm3HMn3o3IMHIUeJMKMLMMMNMOMPMQMRMSMBHVecfTMVjCHUMVMWMXMYMZM0M1M2M3M4M5M6M7M?g8M9Mp3q3r3s3!Mt3u3v3@gw3x3#M$My3z3A3B3[gC3%MD3E3F3G3H3I3J3K3L3M3N3O3DHWeEHFHGHHHIHJHKHLHMHXeNHGiOHPHQH]gYeZe^g_g`g{g|g7d}g~g(h:c)h*hdfef+h,h-h.h/h8d;c:hff9d;h=h?hah'M(MP3Q3R3S3T3!dWjXj)M*M+M,M-M.MRH0eSHTHUHVHWHXHYHZH0H1H1e2HHi3H4H5HU3bhV3/M:MW3X3Y3Z3ch03;M1323=MII2e?M@M[M]M^M_M`M{M|M}M~M6H3egfaNYj7HbNcNdN334353.e/eZj0j8H9H63eNdh738393!3fNgNhN#3$3%3'3(3!Heh)3iN#H$Huh%H'Hvh*3jN+3,3-3kNfh.3/3:3;3lNmNnN=3?3@3[3]3(Hgh^3oN)H_3pN`3{3|3qNhh}3~3a4b4rNsNtNc4d4e4f4g4*Hihh4uN+Hi4vNj4k4l4wNjhm4n4o4p4xNyNzNq4r4s4t4u4,Hkhv4AN-Hw4BNx4y4z4CNlhA4B4C4D4DNENFNE4F4G4H4I4.HmhJ4GN/HK4HNL4M4N4INnhO4P4Q4R4JNKNLNS4T4U4V4W4:HohX4MN;HY4NNZ40414ONph24344454PNQNRN64748494!4=Hqh#4SN?H@H$4TN%4'4(4UNrh)4*4+4,4VNWNXN-4.4/4:4;4[Hsh=4YN]H^H?4ZN@4[4]46c^4JI7c_48c`4{4|4}4~4a5b5c5d5e5f5g5h5i5j5k5l5m5n5o5p5q5r5s5t5u5v5w5x5y5z5A5B5C5D5E5F5G5c H5I5_H8b`H)dJ5K5L5M5N5O5P5Q5R5S5T5U5V5W51j2j0N{H|H}HLcX5~H1N2NaIKI^iY5Z505152535455565759c!c8595!5#5$5%5'5(5)5*5+5,5-53N.5/5:5;5=5?5@5[5]5^5bI_5`54N{55N|5}56N~5a6b6Iic6d6e6f6g6h6i6j6k6l6m6n6o6p6q6r6s6t6u6v6w6x6y6z6A6B6C6D6E6F6G6H6I6J6K6L6M6N6O6P6Q6R63jS64jT65jU66jV67jW68jX69jY6!jZ6#j06$j16%j26'j36+b46(j56Rh66Ji76Ki8696!6#6$6%6'6Li(6)j)6*j*6+j+6,j,6Mi-6.6/6:6;67N=68N?69N@6!N[6#N]6$N^6-j_6%N`6'N{6(N|6)N}6*N~6.ja7/jb7+Nc7:jd7;je7=jf7Nig7?jh7Oii7@jj7[jk7l7m7Pfn7o7]jp7q7^jr7Pis7t7u7v7w7x7y7z7A7B7C7D7E7F7G7H7I7J7K7L7M7N7O7P7Q7R7S7T7U7,N-NV7.N/N:N;N=N?NW7X7Y7@NZ7[N07]N^N_N17`N{N|N}N~NaO27bOcO3747dO5767eO7787fOgO97!7#7hO$7%7'7(7)7*7Qf+7,7-7.7/7:7;7=7?7@7[7]7^7_7`7{7|7}7~7a8b8c8d8e8f8g8h8i8j8k8l8m8n8o8p8q8r8

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: 23j4j5j6j7j8j9j!j#j$j%j'j+b(jRhJiKiLi)j*j+j,jMi-j.j/j:j;j=jNi?jOi@j[jPf]j^j

76 HANDLE_RETURN(cydriver.cuCtxGetDevice(&dev)) 23j4j5j6j7j8j9j!j#j$j%j'j+b(jRhJiKiLi)j*j+j,jMi-j.j/j:j;j=jNi?jOi@j[jPf]j^j

77 return <int>dev 23j4j5j6j7j8j9j!j#j$j%j'j+b(jRhJiKiLi)j*j+j,jMi-j.j/j:j;j=jNi?jOi@j[jPf]j^j

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 vbse9 ! teuevewbwexbxeybzbAbBbCb# $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ DbEbFbGbHbIbngogpgtdqgudrgvdsgwdtgxdugyd9bvgwgxgygzge f MFAgJbqbKbBgLbMbNbObPbd g h i j abbbcbdbrbsbtbebQbRbSbTbUbyeVbfb6 7 ScTczeIjWFXFgbAdUcBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdUdVdWdXdYdZd0d1d2dAeBeb ccdcecfcgchcicjckclcmcncocpcqcrcsctcucvcwcxcyczcAcBcCcDcEcFcGcYFHcKc(F)F'd*F+F,F-F.FVc/F4b:F;F=F?F@F[F]F^F_F`F{F|F}F~F0c1c2caGbGcG3c4c5c!b#bdGeGfGgGhGiGjGkGlGmGnGoGpGqGrGsGtGuGvGwGxGyGzGAGBGCGDGEGFGGGHGIGJGKGLGMGNGOGPGQGRGSGTGUGVGWGXGYGZG0G1G2G3G4G5G6G7G8G9G!G#G$G%G'G(GCg)GDg*G+GEg,GFg-GGgHgCeLjDeNjEeFeIgJgKgLgMgNgOgPg.GQg/GRg:GSg;GTgUgGe=GHe?GVgWgXg@GIe[G]G^G_G`G{G|G}G~GJeaHEibHcHdHYgKeLeZg0g1g2g3g3d4g5g.c4d/c5d6g6deHMefHgHhHiHjHkHlHmHnHoHNepHFiqHrHsH7g8gOetHPeuH9gvH!gwH#g$gQeSjReUjSeTe%g'g(g)g*g+g,g-gxH.gyH/gzH:gAH;g=gUeBHVeCH?g@g[gDHWeEHFHGHHHIHJHKHLHMHXeNHGiOHPHQH]gYeZe^g_g`g{g|g7d}g~g:c8d;c9dah!dRH0eSHTHUHVHWHXHYHZH0H1H1e2HHi3H4H5Hbhch2e6H3e7H8H9Hdh!Heh#H$H%Hfh(Hgh)Hhh*Hih+Hjh,Hkh-Hlh.Hmh/Hnh:Hoh;Hph=Hqh?H@Hrh[Hsh]H^H*d6c7c8c+d,d-dgdhdid.d/d:d;d#c$cjd=dkd?d@d[d]dld^dmdndodpd_d`d{d|d}d~daebeceqdrdsdc de{H|H}HLc~HaI9c!ceefegeheiejekeiOjOkOlOlemOmenenO

83 return 0 28 vbse9 ! teuevewbwexbxeybzbAbBbCb# $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ DbEbFbGbHbIbngogpgtdqgudrgvdsgwdtgxdugyd9bvgwgxgygzge f MFAgJbqbKbBgLbMbNbObPbd g h i j abbbcbdbrbsbtbebQbRbSbTbUbyeVbfb6 7 ScTczeIjWFXFgbAdUcBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdUdVdWdXdYdZd0d1d2dAeBeb ccdcecfcgchcicjckclcmcncocpcqcrcsctcucvcwcxcyczcAcBcCcDcEcFcGcYFHcKc(F)F'd*F+F,F-F.FVc/F4b:F;F=F?F@F[F]F^F_F`F{F|F}F~F0c1c2caGbGcG3c4c5c!b#bdGeGfGgGhGiGjGkGlGmGnGoGpGqGrGsGtGuGvGwGxGyGzGAGBGCGDGEGFGGGHGIGJGKGLGMGNGOGPGQGRGSGTGUGVGWGXGYGZG0G1G2G3G4G5G6G7G8G9G!G#G$G%G'G(GCg)GDg*G+GEg,GFg-GGgHgCeLjDeNjEeFeIgJgKgLgMgNgOgPg.GQg/GRg:GSg;GTgUgGe=GHe?GVgWgXg@GIe[G]G^G_G`G{G|G}G~GJeaHEibHcHdHYgKeLeZg0g1g2g3g3d4g5g.c4d/c5d6g6deHMefHgHhHiHjHkHlHmHnHoHNepHFiqHrHsH7g8gOetHPeuH9gvH!gwH#g$gQeSjReUjSeTe%g'g(g)g*g+g,g-gxH.gyH/gzH:gAH;g=gUeBHVeCH?g@g[gDHWeEHFHGHHHIHJHKHLHMHXeNHGiOHPHQH]gYeZe^g_g`g{g|g7d}g~g:c8d;c9dah!dRH0eSHTHUHVHWHXHYHZH0H1H1e2HHi3H4H5Hbhch2e6H3e7H8H9Hdh!Heh#H$H%Hfh(Hgh)Hhh*Hih+Hjh,Hkh-Hlh.Hmh/Hnh:Hoh;Hph=Hqh?H@Hrh[Hsh]H^H*d6c7c8c+d,d-dgdhdid.d/d:d;d#c$cjd=dkd?d@d[d]dld^dmdndodpd_d`d{d|d}d~daebeceqdrdsdc de{H|H}HLc~HaI9c!ceefegeheiejekeiOjOkOlOlemOmenenO

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: 2'HoOpO_H8b`H)dqOrOsOtOuOvOwOxOyOzOAOBOCODO1j2j

109 return 0 2'HoOpO_H8b`H)dqOrOsOtOuOvOwOxOyOzOAOBOCODO1j2j

110 with gil: 28b

111 _raise_nvvm_error(prog, err) 28b

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 28b

117 if prog != NULL: 28b

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

119 cdef bytes log_bytes 

120 cdef str log_str = "" 28b

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

122 log_bytes = b" " * logsize 28b

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

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

125 cdef object exc = nvvmError(err) 28b

126 if log_str: 28b

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

128 raise exc 28b

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 EOFOGOHOIOJOKOLOMONOOOLIMIPOQOROSOTOUOVOWOXOYOZO0O1ORf2O3O4O5O6O!b#b^i7O8O9O!O#O$O%O'O(O)O

135 return 0 2b EOFOGOHOIOJOKOLOMONOOOLIMIPOQOROSOTOUOVOWOXOYOZO0O1O2O3O4O5O6O!b#b^i7O8O9O!O#O$O%O'O(O)O

136 with gil: 2b Rf

137 _raise_nvjitlink_error(handle, err) 2b Rf

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 Rf

144 if handle != NULL: 2b Rf

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

146 cdef bytes log_bytes 

147 cdef str log_str = "" 2b Rf

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

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 Rf

154 if log_str: 2b Rf

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

156 raise exc 2b Rf

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 _jhf`jif{jjf|jkf}jlf~j8 akvbbkck9 dk#dek! fkgkhkikjkkkwblkmknkxbokpkqkybrkzbsktkAbukBbvkCbwk# xk$ yk% zk' Ak( Bk) Ck* Dk+ Ek, Fk- Gk. Hk/ Ik: Jk; Kk= Lk? Mk@ Nk[ Ok] Pk^ Qk_ Rk` Sk{ Tk| Uk} Vk~ WkXkYkZk0k1k@h2k3k4k5k6k7k8k9k!k#k$k%k'k(k)kmf*k+k,knf-kof.kpf/kqf:kDb;krf=ksf?ktf@k[k]kEb^kFb_kGb`kHb{kIb|k=c}k?c~k@cal[cbl]ccl^cdl_cel`cflglhliljlklllml[hnl]hol^hplqlrlsltlulvlwlxlylzlAlBlClDlElFlGlHlIlJlKlLlMlNlOlPlQlRlSlTlUlVlWlXlYlZl0l1l2l3l4l5l6l7l8l9l!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_h?m@m[m]m^m_m`m{m`h|m}m~man9bbncn{hdnufenfngnhninjnknlnmnnnonpnqnrnsntnvfunwfvnwnxnynzn{cAn$dBnCnDne Enf FnGnHnJbInqbJnKbKnLnMnNn[bOn]bPn^bQnRn_bSn`bTn{bUn|bVnWnXn}bYnZn0n1nxf2n|h3n4n5nLb6nMb7nNb8nOb9nPb!n#n$n%n'nd (n}h)n4e*n~h+n5e,n6e-n7e.nyf/nzf:nAf;ng =nh ?ni @nj [nab]nbb^ncb_ndb`nrb{nsb|ntb}nBf~nCfaoDfboEfcoFfdoGfeofogoebhoQbioRbjoSbkoTbloUbmonoVboofbpoqoHfrosoaibi_icidi`ieifi{igihi|iii}itoxhuoIf~ivo8ewoJfKfajLfxoMfbjyoyhzojicjkilidjNfAoOfBoCoDoEoFoGo9eHoSfTfUfVfWfXfYfZf0f1f2f3f6 4f7 5f~bIoJoKoLoMoNoOoPoQoRoSo6f7f8f9f!f#f$f%f'f(f)f*f+f,f-f.f/f:f;f=f?f@f[f]fQi^f_f`f{f|f}f~fagbgcg%ddg|cegToUoVo}cfg~cggWbWoXbXoadhgbdigjgkglgmgYbYoZbZo0b0o1b1o2b2o3b3o4o5o6o7o8o9o!ocIej#o$o%o'o(o)o*o+o,o-o.o/o:o;opi=o?o@oqi[oriSc]o^o_o`o{o|o}o~oTcapbpcpdpepfpgphpipsijptiuizdvi,ckp-clpwixiyimpziAiBinpopppqprpsptpupvpwpxpypzpApBpCpDpEpFpGpHpIpgbacbcJpKpLpMpBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdUdVdWdXdYdZd0d1d2dAeBeNpOpPpQpRpSpTpUpVpWpXpYpZp0p1p2p3p4p5p6p7p8p9p!p#p$p%p'p(p)p*p+p,p-p.p/pEhIcJcFhGhHhIhJhKhLhMhCiDioe%chbMcib'cNcjb5b:e;e=ek ,bl m n o p q r s t -bu v w x $by z %b.bpe(ckbOclb)cPcmb6b?e@e[eA /bB C D E F G H I J :bK L M N 'bO P (b;bqe*cnbQcob+cRcpb7b]e^e_eQ =bR S T U V W X Y Z ?b0 1 2 3 )b4 5 *b@bNh:p;p=p?p@p[p]p$e^p_p`p{p|p}p~paqbqcqdqeqfqgqhqiqjqkqlqmqnqoqpqqqrqsqtquqvqwqxqyqzqAqBqCqDqEqFqGqHqIqJqKqLqMqNqOqPqQqRqSqTqUqVqWqXqYqZq0q1q2q3q4q5q6qOhPh)e7q8q9q!q#q$q%F%q'F'qXcYcubre(qZc)q*q+q,q-q.q/q:q;q=q?q@q[q]q^q_q'd`q{q|q}q~qarVcbrcr4bdrerfrgrhrirjrkrlrmrnrorprqrrrsrtrurvrwrxryrzr!bAr#bBrCrDrErFrGrHrIrJrKrLrMrNrOrPrQrRrSrTrUrVrWrXrYrZr0r1r2r3r4r5r6r7r8r9r!r#r$r%r'r(r)r*r+r,r-r.r/r:r;r=r?r@r[r]r^r_r`r{r|r}r~rasbscsdsesfsQhgshsisjskslsmsnsospsqsrssstsusvswsxsyszsAsBsCsDsEsFsGsHsIsJsKsLsMsNsOsPsQsRsSsTsUsVsWsXsYsZs0s1s2s3s4s5s6s7s8s9s!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~BaCbCcCdCeCfCgChCiCjCkClCmCnCoCpCqCrCsCtCuCvCwCxCyCzCACBCCCDCECFCGCHCICJCKCLCMCNCOCPCQCRC)dSCTCUCVCWCXCYCZC0C1C2C3C4C5C6C7C8C9C!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~CaDbDcDdDeDfDgDhDiDjDkDlDmDnDoDIipDqDrDsDtDuDvDwDxDyDzDADBDCDDDEDFDGDHDIDJDKDLDMDNDODPDQDRDSDTDUDVDWDXDYDZD0D1D2D3D4D5D6D7D8D9D!D#D$D%D'D(D)D*D+b+D,D-D.D/D:D;D=D?D@D[D]D^D_D`D{D|D}D~DaEbEcEdEeEfEgEhEiEjEkElEmEnEoEpEqErEsEtEuEvEwExEyEzEAEBECEDEEEFEPiGEHEIEJEKELEMENEOEPEQERESETEUEVEWEXEYEZE0E1E2E3E4E5E6E7E8E9E!E#E$E%E'E(E)E*E+E,E-E.E/E:E;E=E?E@E[E]E^E_E`E{E|E}E~EaFbFcFQfdFeFfFgFhFiFjFkFlFmFnFoFpFqFrFsFtFuFvFwFxFyFzFAFBFCFDFEFFFGFHF

165 return 0 2a _jhf`jif{jjf|jkf}jlf~j8 akvbbkck9 dk#dek! fkgkhkikjkkkwblkmknkxbokpkqkybrkzbsktkAbukBbvkCbwk# xk$ yk% zk' Ak( Bk) Ck* Dk+ Ek, Fk- Gk. Hk/ Ik: Jk; Kk= Lk? Mk@ Nk[ Ok] Pk^ Qk_ Rk` Sk{ Tk| Uk} Vk~ WkXkYkZk0k1k2k3k4k5k6k7k8k9k!k#k$k%k'k(k)kmf*k+k,knf-kof.kpf/kqf:kDb;krf=ksf?ktf@k[k]kEb^kFb_kGb`kHb{kIb|k=c}k?c~k@cal[cbl]ccl^cdl_cel`cflglhliljlklllml[hnl]hol^hplqlrlsltlulvlwlxlylzlAlBlClDlElFlGlHlIlJlKlLlMlNlOlPlQlRlSlTlUlVlWlXlYlZl0l1l2l3l4l5l6l7l8l9l!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~man9bbncndnufenfngnhninjnknlnmnnnonpnqnrnsntnvfunwfvnwnxnynzn{cAn$dBnCnDne Enf FnGnHnJbInqbJnKbKnLnMnNn[bOn]bPn^bQnRn_bSn`bTn{bUn|bVnWnXn}bYnZn0n1nxf2n|h3n4n5nLb6nMb7nNb8nOb9nPb!n#n$n%n'nd (n)n4e*n+n5e,n6e-n7e.nyf/nzf:nAf;ng =nh ?ni @nj [nab]nbb^ncb_ndb`nrb{nsb|ntb}nBf~nCfaoDfboEfcoFfdoGfeofogoebhoQbioRbjoSbkoTbloUbmonoVboofbpoqoHfrosoaibi_icidi`ieifi{igihi|iii}itoxhuoIf~ivo8ewoJfKfajLfxoMfbjyoyhzojicjkilidjNfAoOfBoCoDoEoFoGo9eHoSfTfUfVfWfXfYfZf0f1f2f3f6 4f7 5f~bIoJoKoLoMoNoOoPoQoRoSo6f7f8f9f!f#f$f%f'f(f)f*f+f,f-f.f/f:f;f=f?f@f[f]fQi^f_f`f{f|f}f~fagbgcg%ddg|cegToUoVo}cfg~cggWbWoXbXoadhgbdigjgkglgmgYbYoZbZo0b0o1b1o2b2o3b3o4o5o6o7o8o9o!oej#o$o%o'o(o)o*o+o,o-o.o/o:o;opi=o?o@oqi[oriSc]o^o_o`o{o|o}o~oTcapbpcpdpepfpgphpipsijptiuizdvi,ckp-clpwixiyimpziAiBinpopppqprpsptpupvpwpxpypzpApBpCpDpEpFpGpHpIpgbacbcJpKpLpMpBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdUdVdWdXdYdZd0d1d2dAeBeNpOpPpQpRpSpTpUpVpWpXpYpZp0p1p2p3p4p5p6p7p8p9p!p#p$p%p'p(p)p*p+p,p-p.p/pEhIcJcFhGhHhIhJhKhLhMhCiDioe%chbMcib'cNcjb5b:e;e=ek ,bl m n o p q r s t -bu v w x $by z %b.bpe(ckbOclb)cPcmb6b?e@e[eA /bB C D E F G H I J :bK L M N 'bO P (b;bqe*cnbQcob+cRcpb7b]e^e_eQ =bR S T U V W X Y Z ?b0 1 2 3 )b4 5 *b@bNh:p;p=p?p@p[p]p$e^p_p`p{p|p}p~paqbqcqdqeqfqgqhqiqjqkqlqmqnqoqpqqqrqsqtquqvqwqxqyqzqAqBqCqDqEqFqGqHqIqJqKqLqMqNqOqPqQqRqSqTqUqVqWqXqYqZq0q1q2q3q4q5q6qOhPh)e7q8q9q!q#q$q%F%q'F'qXcYcubre(qZc)q*q+q,q-q.q/q:q;q=q?q@q[q]q^q_q'd`q{q|q}q~qarVcbrcr4bdrerfrgrhrirjrkrlrmrnrorprqrrrsrtrurvrwrxryrzr!bAr#bBrCrDrErFrGrHrIrJrKrLrMrNrOrPrQrRrSrTrUrVrWrXrYrZr0r1r2r3r4r5r6r7r8r9r!r#r$r%r'r(r)r*r+r,r-r.r/r:r;r=r?r@r[r]r^r_r`r{r|r}r~rasbscsdsesfsQhgshsisjskslsmsnsospsqsrssstsusvswsxsyszsAsBsCsDsEsFsGsHsIsJsKsLsMsNsOsPsQsRsSsTsUsVsWsXsYsZs0s1s2s3s4s5s6s7s8s9s!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~CaDbDcDdDeDfDgDhDiDjDkDlDmDnDoDIipDqDrDsDtDuDvDwDxDyDzDADBDCDDDEDFDGDHDIDJDKDLDMDNDODPDQDRDSDTDUDVDWDXDYDZD0D1D2D3D4D5D6D7D8D9D!D#D$D%D'D(D)D*D+b+D,D-D.D/D:D;D=D?D@D[D]D^D_D`D{D|D}D~DaEbEcEdEeEfEgEhEiEjEkElEmEnEoEpEqErEsEtEuEvEwExEyEzEAEBECEDEEEFEPiGEHEIEJEKELEMENEOEPEQERESETEUEVEWEXEYEZE0E1E2E3E4E5E6E7E8E9E!E#E$E%E'E(E)E*E+E,E-E.E/E:E;E=E?E@E[E]E^E_E`E{E|E}E~EaFbFcFQfdFeFfFgFhFiFjFkFlFmFnFoFpFqFrFsFtFuFvFwFxFyFzFAFBFCFDFEFFFGFHF

166 cdef const char* name 

167 name_err = cydriver.cuGetErrorName(error, &name) 2@h_h`h{h}h4e~hcIejub4b)d

168 if name_err != cydriver.CUresult.CUDA_SUCCESS: 2@h_h`h{h}h4e~hcIejub4b)d

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

170 with gil: 2@h_h`h{h}h4e~hcIejub4b)d

171 # TODO: consider lower this to Cython 

172 expl = DRIVER_CU_RESULT_EXPLANATIONS.get(int(error)) 2@h_h`h{h}h4e~hcIejub4b)d

173 if expl is not None: 2@h_h`h{h}h4e~hcIejub4b)d

174 raise CUDAError(f"{name.decode()}: {expl}") 2@h_h`h{h}h4e~hcIejub4b)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: 2SfTfUfVfWfXfYfZf0f1f2f3f4f5f6f7f8f9f!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~fagbgcgdgegfgggWbXbhgigjgkglgmgYbZb0b1b2b3bzh*ORiAh!e#eBhChDhacbcIcJchbibjb5bk l m n o p q r s t u v w x y z kblbmb6bA B C D E F G H I J K L M N O P nbobpb7bQ R S T U V W X Y Z 0 1 2 3 4 5 cdWc%edd'e(eed*e+e.e/e

184 return 0 2SfTfUfVfWfXfYfZf0f1f2f3f4f5f6f7f8f9f!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~fagbgcgdgegfgggWbXbhgigjgkglgmgYbZb0b1b2b3bzhRiAh!e#eBhChDhacbcIcJchbibjb5bk l m n o p q r s t u v w x y z kblbmb6bA B C D E F G H I J K L M N O P nbobpb7bQ R S T U V W X Y Z 0 1 2 3 4 5 cdWc%edd'e(eed*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 2*ORi

188 expl = RUNTIME_CUDA_ERROR_EXPLANATIONS.get(int(error)) 2*ORi

189 if expl is not None: 2*ORi

190 raise CUDAError(f"{name}: {expl}") 2*ORi

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

192 if desc_err != _RUNTIME_SUCCESS: 2Ri

193 raise CUDAError(f"{name}") 

194 desc = desc.decode() 2Ri

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

196  

197  

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

199 if error == _NVRTC_SUCCESS: 2a IF+Ob ccdcecfcgchcicjckclcmcncocpcqcrcsctcucvcwcxcyczcAcBcCcDcEcFcGcHcKc0c1c2cSi3cTi4cUi5cuhvh*d6c7c8c+d,d-dgdhdid.d/d:d;d#c$cjd=dkd?d@d[d]dld^dmdndodpd_d`d{d|d}d~daebeceqdrdsdc deLc9c!cShThUhVhWhXhViWiXiYiZi0i1i2i3i4i5i6i7i8i9i!i#i$i%i'i(i)i*i+i,i-i.ieefegeheiejekelemene

200 return 0 2a +Ob ccdcecfcgchcicjckclcmcncocpcqcrcsctcucvcwcxcyczcAcBcCcDcEcFcGcHcKc0c1c2cSi3cTi4cUi5cuhvh*d6c7c8c+d,d-dgdhdid.d/d:d;d#c$cjd=dkd?d@d[d]dld^dmdndodpd_d`d{d|d}d~daebeceqdrdsdc deLc9c!cShThUhVhWhXhViWiXiYiZi0i1i2i3i4i5i6i7i8i9i!i#i$i%i'i(i)i*i+i,i-i.ieefegeheiejekelemene

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

202 if handle is not None: 2IF+O

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

204 log = b" " * logsize 2IF

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

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

207 raise NVRTCError(err) 2IF+O

208  

209  

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

211 if isinstance(error, driver.CUresult): 2a _jhf`jif{jjf|jkf}jlf~j8 akvbbkck9 dk#dek! fkgkhkikjkkkwblkmknkxbokpkqkybrkzbsktkAbukBbvkCbwk# xk$ yk% zk' Ak( Bk) Ck* Dk+ Ek, Fk- Gk. Hk/ Ik: Jk; Kk= Lk? Mk@ Nk[ Ok] Pk^ Qk_ Rk` Sk{ Tk| Uk} Vk~ WkXkYkZk0k1k2k3k4k5k6k7k8k9k!k#k$k%k'k(k)kmf*k+k,knf-kof.kpf/kqf:kDb;krf=ksf?ktf@k[k]kEb^kFb_kGb`kHb{kIb|k=c}k?c~k@cal[cbl]ccl^cdl_cel`cflglhliljlklllml[hnl]hol^hplqlrlsltlulvlwlxlylzlAlBlClDlElFlGlHlIlJlKlLlMlNlOlPlQlRlSlTlUlVlWlXlYlZl0l1l2l3l4l5l6l7l8l9l!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~man9bbncndnufenfngnhninjnknlnmnnnonpnqnrnsntnvfunwfvnwnxnynzn{cAn$dBnCnDne Enf FnGnHnJbInqbJnKbKnLnMnNn[bOn]bPn^bQnRn_bSn`bTn{bUn|bVnWnXn}bYnZn0n1nxf2n|h3n4n5nLb6nMb7nNb8nOb9nPb!n#n$n%n'nd (n)n*n+n5e,n6e-n7e.nyf/nzf:nAf;ng =nh ?ni @nj [nab]nbb^ncb_ndb`nrb{nsb|ntb}nBf~nCfaoDfboEfcoFfdoGfeofogoebhoQbioRbjoSbkoTbloUbmonoVboofbpoqoHfrosoaibi_icidi`ieifi{igihi|iii}itoxhuoIf~ivo8ewoJfKfajLfxoMfbjyoyhzojicjkilidjNfAoOfBoCoDoEoFoGo9eHoSfTfUfVfWfXfYfZf0f1f2f3f6 4f7 5f~bIoJoKoLoMoNoOoPoQoRoSo6f7f8f9f!f#f$f%f'f(f)f*f+f,f-f.f/f:f;f=f?f@f[f]fQi^f_f`f{f|f}f~fagbgcg%ddg|cegToUoVo}cfg~cggWbWoXbXoadhgbdigjgkglgmgYbYoZbZo0b0o1b1o2b2o3b3ozh4o5o6o7o8o9o!o#oAh!e$o#e%o'o(o)o*o+o,o-oBh.o/o:o;opi=o?o@oqiChDh[ori]o^o_o`o{o|o}o~oapbpcpdpepfpgphpipsijptiuivikplpwixiyimpziAiBinpopppqprpsptpupvpwpxpypzpApBpCpDpEpFpGpHpIpgbacbcJpKpLpMpb ccNpdcOpecPpfcQpgcRphcSpicTpjcUpkcVplcWpmcXpncYpocZppc0pqc1prc2psc3ptc4puc5pvc6pwc7pxc8pyc9pzc!pAc#pBc$pCc%pDc'pEc(pFc)pGc*p+pHc,p-p.p/pEhIcJcFhGhHhIhJhKhLhMhCiDioe%chbMcib'cNcjb5b:e;e=ek ,bl m n o p q r s t -bu v w x $by z %b.bpe(ckbOclb)cPcmb6b?e@e[eA /bB C D E F G H I J :bK L M N 'bO P (b;bqe*cnbQcob+cRcpb7b]e^e_eQ =bR S T U V W X Y Z ?b0 1 2 3 )b4 5 *b@bNh:pcd;p=p?p@p[p]p$e^p_p`p{p|p}p~pWcaq%ebqddcq'edqeqfqgqhqiqjqkqlqmqnqoqpqqqrqsqtquqvqwqxqyqzqAqBqCqDqEqFqGqHqIqJqKqLqMqNqOqPqQqRqSqTqUqVqWqXqYqZq0q1q2q3q4q5q(e6qOhPh)e7qed8q*e+e9q!q#q$q%q'q(q)q*q+q,q-q.q/q:q;q=q?q@q[qKc]q^q_q'd`q{q|q}q~qarbrcr4bdrerfrgrhrirjrkrlrmrnrorprqr0crr1csr2ctrurvrwrSi3cxrTi4cyrUi5czrArBrCrDrErFrGrHrIrJrKrLrMrNrOrPrQrRrSrTrUrVrWrXrYrZr0r1r2r3r4r5r6r7r8r9r!r#r$r%r'r(r)r*r+r,r-r.r/r:r;r=r?r@r[r]r^r_r`r{r|r}r~rasbscsdsesfsQhgshsisjskslsmsnsospsqsrssstsusvswsxsyszsAsBsCsDsEsFsGsHsIsJsKsLsMsNsOsPsQsRsSsTsUsVsWsXsYsZs0s1s2s3s4s5s6s7s8s9s!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.e#z/e$z%z'z(z)z*z+z,z-z.z/z:z;z=z?z@z[z]z^z_z`z{z|z}z~zaAuhbAcAdAvheAfAgAhAiAjAkAlAmAnAoApAqArAsAtAuAvAwAxAyAzAAABACADAEAFAGAHAIAJAKALAMANAOAPAQARASATAUAVAWAXAYAZA0A1A2A3A4A5A6A7A8A9A!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~BaC*dbC6ccC7cdC8ceC+dfC,dgC-dhCgdiChdjCidkC.dlC/dmC:dnC;doC#cpC$cqCjdrC=dsCkdtC?duC@dvC[dwC]dxCldyC^dzCmdACndBCodCCpdDC_dEC`dFC{dGC|dHC}dIC~dJCaeKCbeLCceMCqdNCrdOCsdPCc QCdeRCSCTCUCVCWCXCYCZC0C1C2C3C4C5CLc6C7C8C9C!C#C$C%C'C(C)C9c*C!c+CShThUhVhWhXhViWiXiYiZi0i1i2i3i4i5i6i7i8i,C9i!i#i$i%i'i(i)i*i+i-C.C/C:C,i-i.ieefegeheie;Cjekelemene=C?C@C[C]C^C_C`C{C|C}C~CaDbDcDdDeDfDgDhDiDjDkDlDmDnDoDIipDqDrDsDtDuDvDwDxDyDzDADBDCDDDEDFDGDHDIDJDKDLDMDNDODPDQDRDSDTDUDVDWDXDYDZD0D1D2D3D4D5D6D7D8D9D!D#D$D%D'D(D)D*D+b+D,D-D.D/D:D;D=D?D@D[D]D^D_D`D{D|D}D~DaEbEcEdEeEfEgEhEiEjEkElEmEnEoEpEqErEsEtEuEvEwExEyEzEAEBECEDEEEFEPiGEHEIEJEKELEMENEOEPEQERESETEUEVEWEXEYEZE0E1E2E3E4E5E6E7E8E9E!E#E$E%E'E(E)E*E+E,E-E.E/E:E;E=E?E@E[E]E^E_E`E{E|E}E~EaFbFcFQfdFeFfFgFhFiFjFkFlFmFnFoFpFqFrFsFtFuFvFwFxFyFzFAFBFCFDFEFFFGFHF

212 return _check_driver_error(error) 2a _jhf`jif{jjf|jkf}jlf~j8 akvbbkck9 dk#dek! fkgkhkikjkkkwblkmknkxbokpkqkybrkzbsktkAbukBbvkCbwk# xk$ yk% zk' Ak( Bk) Ck* Dk+ Ek, Fk- Gk. Hk/ Ik: Jk; Kk= Lk? Mk@ Nk[ Ok] Pk^ Qk_ Rk` Sk{ Tk| Uk} Vk~ WkXkYkZk0k1k2k3k4k5k6k7k8k9k!k#k$k%k'k(k)kmf*k+k,knf-kof.kpf/kqf:kDb;krf=ksf?ktf@k[k]kEb^kFb_kGb`kHb{kIb|k=c}k?c~k@cal[cbl]ccl^cdl_cel`cflglhliljlklllml[hnl]hol^hplqlrlsltlulvlwlxlylzlAlBlClDlElFlGlHlIlJlKlLlMlNlOlPlQlRlSlTlUlVlWlXlYlZl0l1l2l3l4l5l6l7l8l9l!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~man9bbncndnufenfngnhninjnknlnmnnnonpnqnrnsntnvfunwfvnwnxnynzn{cAn$dBnCnDne Enf FnGnHnJbInqbJnKbKnLnMnNn[bOn]bPn^bQnRn_bSn`bTn{bUn|bVnWnXn}bYnZn0n1nxf2n|h3n4n5nLb6nMb7nNb8nOb9nPb!n#n$n%n'nd (n)n*n+n5e,n6e-n7e.nyf/nzf:nAf;ng =nh ?ni @nj [nab]nbb^ncb_ndb`nrb{nsb|ntb}nBf~nCfaoDfboEfcoFfdoGfeofogoebhoQbioRbjoSbkoTbloUbmonoVboofbpoqoHfrosoaibi_icidi`ieifi{igihi|iii}itoxhuoIf~ivo8ewoJfKfajLfxoMfbjyoyhzojicjkilidjNfAoOfBoCoDoEoFoGo9eHoSfTfUfVfWfXfYfZf0f1f2f3f6 4f7 5f~bIoJoKoLoMoNoOoPoQoRoSo6f7f8f9f!f#f$f%f'f(f)f*f+f,f-f.f/f:f;f=f?f@f[f]fQi^f_f`f{f|f}f~fagbgcg%ddg|cegToUoVo}cfg~cggWbWoXbXoadhgbdigjgkglgmgYbYoZbZo0b0o1b1o2b2o3b3o4o5o6o7o8o9o!o#o$o%o'o(o)o*o+o,o-o.o/o:o;opi=o?o@oqi[ori]o^o_o`o{o|o}o~oapbpcpdpepfpgphpipsijptiuivikplpwixiyimpziAiBinpopppqprpsptpupvpwpxpypzpApBpCpDpEpFpGpHpIpgbacbcJpKpLpMpNpOpPpQpRpSpTpUpVpWpXpYpZp0p1p2p3p4p5p6p7p8p9p!p#p$p%p'p(p)p*p+p,p-p.p/pEhIcJcFhGhHhIhJhKhLhMhCiDioe%chbMcib'cNcjb5b:e;e=ek ,bl m n o p q r s t -bu v w x $by z %b.bpe(ckbOclb)cPcmb6b?e@e[eA /bB C D E F G H I J :bK L M N 'bO P (b;bqe*cnbQcob+cRcpb7b]e^e_eQ =bR S T U V W X Y Z ?b0 1 2 3 )b4 5 *b@bNh:p;p=p?p@p[p]p$e^p_p`p{p|p}p~paqbqcqdqeqfqgqhqiqjqkqlqmqnqoqpqqqrqsqtquqvqwqxqyqzqAqBqCqDqEqFqGqHqIqJqKqLqMqNqOqPqQqRqSqTqUqVqWqXqYqZq0q1q2q3q4q5q6qOhPh)e7q8q9q!q#q$q%q'q(q)q*q+q,q-q.q/q:q;q=q?q@q[q]q^q_q'd`q{q|q}q~qarbrcr4bdrerfrgrhrirjrkrlrmrnrorprqrrrsrtrurvrwrxryrzrArBrCrDrErFrGrHrIrJrKrLrMrNrOrPrQrRrSrTrUrVrWrXrYrZr0r1r2r3r4r5r6r7r8r9r!r#r$r%r'r(r)r*r+r,r-r.r/r:r;r=r?r@r[r]r^r_r`r{r|r}r~rasbscsdsesfsQhgshsisjskslsmsnsospsqsrssstsusvswsxsyszsAsBsCsDsEsFsGsHsIsJsKsLsMsNsOsPsQsRsSsTsUsVsWsXsYsZs0s1s2s3s4s5s6s7s8s9s!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~CaDbDcDdDeDfDgDhDiDjDkDlDmDnDoDIipDqDrDsDtDuDvDwDxDyDzDADBDCDDDEDFDGDHDIDJDKDLDMDNDODPDQDRDSDTDUDVDWDXDYDZD0D1D2D3D4D5D6D7D8D9D!D#D$D%D'D(D)D*D+b+D,D-D.D/D:D;D=D?D@D[D]D^D_D`D{D|D}D~DaEbEcEdEeEfEgEhEiEjEkElEmEnEoEpEqErEsEtEuEvEwExEyEzEAEBECEDEEEFEPiGEHEIEJEKELEMENEOEPEQERESETEUEVEWEXEYEZE0E1E2E3E4E5E6E7E8E9E!E#E$E%E'E(E)E*E+E,E-E.E/E:E;E=E?E@E[E]E^E_E`E{E|E}E~EaFbFcFQfdFeFfFgFhFiFjFkFlFmFnFoFpFqFrFsFtFuFvFwFxFyFzFAFBFCFDFEFFFGFHF

213 elif isinstance(error, runtime.cudaError_t): 2a SfTfUfVfWfXfYfZf0f1f2f3f4f5f6f7f8f9f!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~fagbgcgdgegfgggWbXbhgigjgkglgmgYbZb0b1b2b3bzhAh!e#eBhChDhacbcb ccdcecfcgchcicjckclcmcncocpcqcrcsctcucvcwcxcyczcAcBcCcDcEcFcGcHcIcJchbibjb5bk l m n o p q r s t u v w x y z kblbmb6bA B C D E F G H I J K L M N O P nbobpb7bQ R S T U V W X Y Z 0 1 2 3 4 5 cdWc%edd'e(eed*e+eKc0c1c2cSi3cTi4cUi5c.e/euhvh*d6c7c8c+d,d-dgdhdid.d/d:d;d#c$cjd=dkd?d@d[d]dld^dmdndodpd_d`d{d|d}d~daebeceqdrdsdc deLc9c!cShThUhVhWhXhViWiXiYiZi0i1i2i3i4i5i6i7i8i9i!i#i$i%i'i(i)i*i+i,i-i.ieefegeheiejekelemene

214 return _check_runtime_error(error) 2SfTfUfVfWfXfYfZf0f1f2f3f4f5f6f7f8f9f!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~fagbgcgdgegfgggWbXbhgigjgkglgmgYbZb0b1b2b3bzhAh!e#eBhChDhacbcIcJchbibjb5bk l m n o p q r s t u v w x y z kblbmb6bA B C D E F G H I J K L M N O P nbobpb7bQ R S T U V W X Y Z 0 1 2 3 4 5 cdWc%edd'e(eed*e+e.e/e

215 elif isinstance(error, nvrtc.nvrtcResult): 2a b ccdcecfcgchcicjckclcmcncocpcqcrcsctcucvcwcxcyczcAcBcCcDcEcFcGcHcKc0c1c2cSi3cTi4cUi5cuhvh*d6c7c8c+d,d-dgdhdid.d/d:d;d#c$cjd=dkd?d@d[d]dld^dmdndodpd_d`d{d|d}d~daebeceqdrdsdc deLc9c!cShThUhVhWhXhViWiXiYiZi0i1i2i3i4i5i6i7i8i9i!i#i$i%i'i(i)i*i+i,i-i.ieefegeheiejekelemene

216 return _check_nvrtc_error(error, handle=handle) 2a b ccdcecfcgchcicjckclcmcncocpcqcrcsctcucvcwcxcyczcAcBcCcDcEcFcGcHcKc0c1c2cSi3cTi4cUi5cuhvh*d6c7c8c+d,d-dgdhdid.d/d:d;d#c$cjd=dkd?d@d[d]dld^dmdndodpd_d`d{d|d}d~daebeceqdrdsdc deLc9c!cShThUhVhWhXhViWiXiYiZi0i1i2i3i4i5i6i7i8i9i!i#i$i%i'i(i)i*i+i,i-i.ieefegeheiejekelemene

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 _jhf`jif{jjf|jkf}jlf~j8 akvbbkck9 dk#dek! fkgkhkikjkkkwblkmknkxbokpkqkybrkzbsktkAbukBbvkCbwk# xk$ yk% zk' Ak( Bk) Ck* Dk+ Ek, Fk- Gk. Hk/ Ik: Jk; Kk= Lk? Mk@ Nk[ Ok] Pk^ Qk_ Rk` Sk{ Tk| Uk} Vk~ WkXkYkZk0k1k2k3k4k5k6k7k8k9k!k#k$k%k'k(k)kmf*k+k,knf-kof.kpf/kqf:kDb;krf=ksf?ktf@k[k]kEb^kFb_kGb`kHb{kIb|k=c}k?c~k@cal[cbl]ccl^cdl_cel`cflglhliljlklllml[hnl]hol^hplqlrlsltlulvlwlxlylzlAlBlClDlElFlGlHlIlJlKlLlMlNlOlPlQlRlSlTlUlVlWlXlYlZl0l1l2l3l4l5l6l7l8l9l!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~man9bbncndnufenfngnhninjnknlnmnnnonpnqnrnsntnvfunwfvnwnxnynzn{cAn$dBnCnDne Enf FnGnHnJbInqbJnKbKnLnMnNn[bOn]bPn^bQnRn_bSn`bTn{bUn|bVnWnXn}bYnZn0n1nxf2n|h3n4n5nLb6nMb7nNb8nOb9nPb!n#n$n%n'nd (n)n*n+n5e,n6e-n7e.nyf/nzf:nAf;ng =nh ?ni @nj [nab]nbb^ncb_ndb`nrb{nsb|ntb}nBf~nCfaoDfboEfcoFfdoGfeofogoebhoQbioRbjoSbkoTbloUbmonoVboofbpoqoHfrosoaibi_icidi`ieifi{igihi|iii}itoxhuoIf~ivo8ewoJfKfajLfxoMfbjyoyhzojicjkilidjNfAoOfBoCoDoEoFoGo9eHoSfTfUfVfWfXfYfZf0f1f2f3f6 4f7 5f~bIoJoKoLoMoNoOoPoQoRoSo6f7f8f9f!f#f$f%f'f(f)f*f+f,f-f.f/f:f;f=f?f@f[f]fQi^f_f`f{f|f}f~fagbgcg%ddg|cegToUoVo}cfg~cggWbWoXbXoadhgbdigjgkglgmgYbYoZbZo0b0o1b1o2b2o3b3ozh4o5o6o7o8o9o!o#oAh!e$o#e%o'o(o)o*o+o,o-oBh.o/o:o;opi=o?o@oqiChDh[ori]o^o_o`o{o|o}o~oapbpcpdpepfpgphpipsijptiuivikplpwixiyimpziAiBinpopppqprpsptpupvpwpxpypzpApBpCpDpEpFpGpHpIpgbacbcJpKpLpMpb ccNpdcOpecPpfcQpgcRphcSpicTpjcUpkcVplcWpmcXpncYpocZppc0pqc1prc2psc3ptc4puc5pvc6pwc7pxc8pyc9pzc!pAc#pBc$pCc%pDc'pEc(pFc)pGc*p+pHc,p-p.p/pEhIcJcFhGhHhIhJhKhLhMhCiDioe%chbMcib'cNcjb5b:e;e=ek ,bl m n o p q r s t -bu v w x $by z %b.bpe(ckbOclb)cPcmb6b?e@e[eA /bB C D E F G H I J :bK L M N 'bO P (b;bqe*cnbQcob+cRcpb7b]e^e_eQ =bR S T U V W X Y Z ?b0 1 2 3 )b4 5 *b@bNh:pcd;p=p?p@p[p]p$e^p_p`p{p|p}p~pWcaq%ebqddcq'edqeqfqgqhqiqjqkqlqmqnqoqpqqqrqsqtquqvqwqxqyqzqAqBqCqDqEqFqGqHqIqJqKqLqMqNqOqPqQqRqSqTqUqVqWqXqYqZq0q1q2q3q4q5q(e6qOhPh)e7qed8q*e+e9q!q#q$q%q'q(q)q*q+q,q-q.q/q:q;q=q?q@q[qKc]q^q_q'd`q{q|q}q~qarbrcr4bdrerfrgrhrirjrkrlrmrnrorprqr0crr1csr2ctrurvrwrSi3cxrTi4cyrUi5czrArBrCrDrErFrGrHrIrJrKrLrMrNrOrPrQrRrSrTrUrVrWrXrYrZr0r1r2r3r4r5r6r7r8r9r!r#r$r%r'r(r)r*r+r,r-r.r/r:r;r=r?r@r[r]r^r_r`r{r|r}r~rasbscsdsesfsQhgshsisjskslsmsnsospsqsrssstsusvswsxsyszsAsBsCsDsEsFsGsHsIsJsKsLsMsNsOsPsQsRsSsTsUsVsWsXsYsZs0s1s2s3s4s5s6s7s8s9s!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.e#z/e$z%z'z(z)z*z+z,z-z.z/z:z;z=z?z@z[z]z^z_z`z{z|z}z~zaAuhbAcAdAvheAfAgAhAiAjAkAlAmAnAoApAqArAsAtAuAvAwAxAyAzAAABACADAEAFAGAHAIAJAKALAMANAOAPAQARASATAUAVAWAXAYAZA0A1A2A3A4A5A6A7A8A9A!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~BaC*dbC6ccC7cdC8ceC+dfC,dgC-dhCgdiChdjCidkC.dlC/dmC:dnC;doC#cpC$cqCjdrC=dsCkdtC?duC@dvC[dwC]dxCldyC^dzCmdACndBCodCCpdDC_dEC`dFC{dGC|dHC}dIC~dJCaeKCbeLCceMCqdNCrdOCsdPCc QCdeRCSCTCUCVCWCXCYCZC0C1C2C3C4C5CLc6C7C8C9C!C#C$C%C'C(C)C9c*C!c+CShThUhVhWhXhViWiXiYiZi0i1i2i3i4i5i6i7i8i,C9i!i#i$i%i'i(i)i*i+i-C.C/C:C,i-i.ieefegeheie;Cjekelemene=C?C@C[C]C^C_C`C{C|C}C~CaDbDcDdDeDfDgDhDiDjDkDlDmDnDoDIipDqDrDsDtDuDvDwDxDyDzDADBDCDDDEDFDGDHDIDJDKDLDMDNDODPDQDRDSDTDUDVDWDXDYDZD0D1D2D3D4D5D6D7D8D9D!D#D$D%D'D(D)D*D+b+D,D-D.D/D:D;D=D?D@D[D]D^D_D`D{D|D}D~DaEbEcEdEeEfEgEhEiEjEkElEmEnEoEpEqErEsEtEuEvEwExEyEzEAEBECEDEEEFEPiGEHEIEJEKELEMENEOEPEQERESETEUEVEWEXEYEZE0E1E2E3E4E5E6E7E8E9E!E#E$E%E'E(E)E*E+E,E-E.E/E:E;E=E?E@E[E]E^E_E`E{E|E}E~EaFbFcFQfdFeFfFgFhFiFjFkFlFmFnFoFpFqFrFsFtFuFvFwFxFyFzFAFBFCFDFEFFFGFHF

223 cdef int out_len = len(result) 2a _jhf`jif{jjf|jkf}jlf~j8 akvbbkck9 dk#dek! fkgkhkikjkkkwblkmknkxbokpkqkybrkzbsktkAbukBbvkCbwk# xk$ yk% zk' Ak( Bk) Ck* Dk+ Ek, Fk- Gk. Hk/ Ik: Jk; Kk= Lk? Mk@ Nk[ Ok] Pk^ Qk_ Rk` Sk{ Tk| Uk} Vk~ WkXkYkZk0k1k2k3k4k5k6k7k8k9k!k#k$k%k'k(k)kmf*k+k,knf-kof.kpf/kqf:kDb;krf=ksf?ktf@k[k]kEb^kFb_kGb`kHb{kIb|k=c}k?c~k@cal[cbl]ccl^cdl_cel`cflglhliljlklllml[hnl]hol^hplqlrlsltlulvlwlxlylzlAlBlClDlElFlGlHlIlJlKlLlMlNlOlPlQlRlSlTlUlVlWlXlYlZl0l1l2l3l4l5l6l7l8l9l!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~man9bbncndnufenfngnhninjnknlnmnnnonpnqnrnsntnvfunwfvnwnxnynzn{cAn$dBnCnDne Enf FnGnHnJbInqbJnKbKnLnMnNn[bOn]bPn^bQnRn_bSn`bTn{bUn|bVnWnXn}bYnZn0n1nxf2n|h3n4n5nLb6nMb7nNb8nOb9nPb!n#n$n%n'nd (n)n*n+n5e,n6e-n7e.nyf/nzf:nAf;ng =nh ?ni @nj [nab]nbb^ncb_ndb`nrb{nsb|ntb}nBf~nCfaoDfboEfcoFfdoGfeofogoebhoQbioRbjoSbkoTbloUbmonoVboofbpoqoHfrosoaibi_icidi`ieifi{igihi|iii}itoxhuoIf~ivo8ewoJfKfajLfxoMfbjyoyhzojicjkilidjNfAoOfBoCoDoEoFoGo9eHoSfTfUfVfWfXfYfZf0f1f2f3f6 4f7 5f~bIoJoKoLoMoNoOoPoQoRoSo6f7f8f9f!f#f$f%f'f(f)f*f+f,f-f.f/f:f;f=f?f@f[f]fQi^f_f`f{f|f}f~fagbgcg%ddg|cegToUoVo}cfg~cggWbWoXbXoadhgbdigjgkglgmgYbYoZbZo0b0o1b1o2b2o3b3ozh4o5o6o7o8o9o!o#oAh!e$o#e%o'o(o)o*o+o,o-oBh.o/o:o;opi=o?o@oqiChDh[ori]o^o_o`o{o|o}o~oapbpcpdpepfpgphpipsijptiuivikplpwixiyimpziAiBinpopppqprpsptpupvpwpxpypzpApBpCpDpEpFpGpHpIpgbacbcJpKpLpMpb ccNpdcOpecPpfcQpgcRphcSpicTpjcUpkcVplcWpmcXpncYpocZppc0pqc1prc2psc3ptc4puc5pvc6pwc7pxc8pyc9pzc!pAc#pBc$pCc%pDc'pEc(pFc)pGc*p+pHc,p-p.p/pEhIcJcFhGhHhIhJhKhLhMhCiDioe%chbMcib'cNcjb5b:e;e=ek ,bl m n o p q r s t -bu v w x $by z %b.bpe(ckbOclb)cPcmb6b?e@e[eA /bB C D E F G H I J :bK L M N 'bO P (b;bqe*cnbQcob+cRcpb7b]e^e_eQ =bR S T U V W X Y Z ?b0 1 2 3 )b4 5 *b@bNh:pcd;p=p?p@p[p]p$e^p_p`p{p|p}p~pWcaq%ebqddcq'edqeqfqgqhqiqjqkqlqmqnqoqpqqqrqsqtquqvqwqxqyqzqAqBqCqDqEqFqGqHqIqJqKqLqMqNqOqPqQqRqSqTqUqVqWqXqYqZq0q1q2q3q4q5q(e6qOhPh)e7qed8q*e+e9q!q#q$q%q'q(q)q*q+q,q-q.q/q:q;q=q?q@q[qKc]q^q_q'd`q{q|q}q~qarbrcr4bdrerfrgrhrirjrkrlrmrnrorprqr0crr1csr2ctrurvrwrSi3cxrTi4cyrUi5czrArBrCrDrErFrGrHrIrJrKrLrMrNrOrPrQrRrSrTrUrVrWrXrYrZr0r1r2r3r4r5r6r7r8r9r!r#r$r%r'r(r)r*r+r,r-r.r/r:r;r=r?r@r[r]r^r_r`r{r|r}r~rasbscsdsesfsQhgshsisjskslsmsnsospsqsrssstsusvswsxsyszsAsBsCsDsEsFsGsHsIsJsKsLsMsNsOsPsQsRsSsTsUsVsWsXsYsZs0s1s2s3s4s5s6s7s8s9s!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.e#z/e$z%z'z(z)z*z+z,z-z.z/z:z;z=z?z@z[z]z^z_z`z{z|z}z~zaAuhbAcAdAvheAfAgAhAiAjAkAlAmAnAoApAqArAsAtAuAvAwAxAyAzAAABACADAEAFAGAHAIAJAKALAMANAOAPAQARASATAUAVAWAXAYAZA0A1A2A3A4A5A6A7A8A9A!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~BaC*dbC6ccC7cdC8ceC+dfC,dgC-dhCgdiChdjCidkC.dlC/dmC:dnC;doC#cpC$cqCjdrC=dsCkdtC?duC@dvC[dwC]dxCldyC^dzCmdACndBCodCCpdDC_dEC`dFC{dGC|dHC}dIC~dJCaeKCbeLCceMCqdNCrdOCsdPCc QCdeRCSCTCUCVCWCXCYCZC0C1C2C3C4C5CLc6C7C8C9C!C#C$C%C'C(C)C9c*C!c+CShThUhVhWhXhViWiXiYiZi0i1i2i3i4i5i6i7i8i,C9i!i#i$i%i'i(i)i*i+i-C.C/C:C,i-i.ieefegeheie;Cjekelemene=C?C@C[C]C^C_C`C{C|C}C~CaDbDcDdDeDfDgDhDiDjDkDlDmDnDoDIipDqDrDsDtDuDvDwDxDyDzDADBDCDDDEDFDGDHDIDJDKDLDMDNDODPDQDRDSDTDUDVDWDXDYDZD0D1D2D3D4D5D6D7D8D9D!D#D$D%D'D(D)D*D+b+D,D-D.D/D:D;D=D?D@D[D]D^D_D`D{D|D}D~DaEbEcEdEeEfEgEhEiEjEkElEmEnEoEpEqErEsEtEuEvEwExEyEzEAEBECEDEEEFEPiGEHEIEJEKELEMENEOEPEQERESETEUEVEWEXEYEZE0E1E2E3E4E5E6E7E8E9E!E#E$E%E'E(E)E*E+E,E-E.E/E:E;E=E?E@E[E]E^E_E`E{E|E}E~EaFbFcFQfdFeFfFgFhFiFjFkFlFmFnFoFpFqFrFsFtFuFvFwFxFyFzFAFBFCFDFEFFFGFHF

224 if out_len == 1: 2a _jhf`jif{jjf|jkf}jlf~j8 akvbbkck9 dk#dek! fkgkhkikjkkkwblkmknkxbokpkqkybrkzbsktkAbukBbvkCbwk# xk$ yk% zk' Ak( Bk) Ck* Dk+ Ek, Fk- Gk. Hk/ Ik: Jk; Kk= Lk? Mk@ Nk[ Ok] Pk^ Qk_ Rk` Sk{ Tk| Uk} Vk~ WkXkYkZk0k1k2k3k4k5k6k7k8k9k!k#k$k%k'k(k)kmf*k+k,knf-kof.kpf/kqf:kDb;krf=ksf?ktf@k[k]kEb^kFb_kGb`kHb{kIb|k=c}k?c~k@cal[cbl]ccl^cdl_cel`cflglhliljlklllml[hnl]hol^hplqlrlsltlulvlwlxlylzlAlBlClDlElFlGlHlIlJlKlLlMlNlOlPlQlRlSlTlUlVlWlXlYlZl0l1l2l3l4l5l6l7l8l9l!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~man9bbncndnufenfngnhninjnknlnmnnnonpnqnrnsntnvfunwfvnwnxnynzn{cAn$dBnCnDne Enf FnGnHnJbInqbJnKbKnLnMnNn[bOn]bPn^bQnRn_bSn`bTn{bUn|bVnWnXn}bYnZn0n1nxf2n|h3n4n5nLb6nMb7nNb8nOb9nPb!n#n$n%n'nd (n)n*n+n5e,n6e-n7e.nyf/nzf:nAf;ng =nh ?ni @nj [nab]nbb^ncb_ndb`nrb{nsb|ntb}nBf~nCfaoDfboEfcoFfdoGfeofogoebhoQbioRbjoSbkoTbloUbmonoVboofbpoqoHfrosoaibi_icidi`ieifi{igihi|iii}itoxhuoIf~ivo8ewoJfKfajLfxoMfbjyoyhzojicjkilidjNfAoOfBoCoDoEoFoGo9eHoSfTfUfVfWfXfYfZf0f1f2f3f6 4f7 5f~bIoJoKoLoMoNoOoPoQoRoSo6f7f8f9f!f#f$f%f'f(f)f*f+f,f-f.f/f:f;f=f?f@f[f]fQi^f_f`f{f|f}f~fagbgcg%ddg|cegToUoVo}cfg~cggWbWoXbXoadhgbdigjgkglgmgYbYoZbZo0b0o1b1o2b2o3b3ozh4o5o6o7o8o9o!o#oAh!e$o#e%o'o(o)o*o+o,o-oBh.o/o:o;opi=o?o@oqiChDh[ori]o^o_o`o{o|o}o~oapbpcpdpepfpgphpipsijptiuivikplpwixiyimpziAiBinpopppqprpsptpupvpwpxpypzpApBpCpDpEpFpGpHpIpgbacbcJpKpLpMpb ccNpdcOpecPpfcQpgcRphcSpicTpjcUpkcVplcWpmcXpncYpocZppc0pqc1prc2psc3ptc4puc5pvc6pwc7pxc8pyc9pzc!pAc#pBc$pCc%pDc'pEc(pFc)pGc*p+pHc,p-p.p/pEhIcJcFhGhHhIhJhKhLhMhCiDioe%chbMcib'cNcjb5b:e;e=ek ,bl m n o p q r s t -bu v w x $by z %b.bpe(ckbOclb)cPcmb6b?e@e[eA /bB C D E F G H I J :bK L M N 'bO P (b;bqe*cnbQcob+cRcpb7b]e^e_eQ =bR S T U V W X Y Z ?b0 1 2 3 )b4 5 *b@bNh:pcd;p=p?p@p[p]p$e^p_p`p{p|p}p~pWcaq%ebqddcq'edqeqfqgqhqiqjqkqlqmqnqoqpqqqrqsqtquqvqwqxqyqzqAqBqCqDqEqFqGqHqIqJqKqLqMqNqOqPqQqRqSqTqUqVqWqXqYqZq0q1q2q3q4q5q(e6qOhPh)e7qed8q*e+e9q!q#q$q%q'q(q)q*q+q,q-q.q/q:q;q=q?q@q[qKc]q^q_q'd`q{q|q}q~qarbrcr4bdrerfrgrhrirjrkrlrmrnrorprqr0crr1csr2ctrurvrwrSi3cxrTi4cyrUi5czrArBrCrDrErFrGrHrIrJrKrLrMrNrOrPrQrRrSrTrUrVrWrXrYrZr0r1r2r3r4r5r6r7r8r9r!r#r$r%r'r(r)r*r+r,r-r.r/r:r;r=r?r@r[r]r^r_r`r{r|r}r~rasbscsdsesfsQhgshsisjskslsmsnsospsqsrssstsusvswsxsyszsAsBsCsDsEsFsGsHsIsJsKsLsMsNsOsPsQsRsSsTsUsVsWsXsYsZs0s1s2s3s4s5s6s7s8s9s!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.e#z/e$z%z'z(z)z*z+z,z-z.z/z:z;z=z?z@z[z]z^z_z`z{z|z}z~zaAuhbAcAdAvheAfAgAhAiAjAkAlAmAnAoApAqArAsAtAuAvAwAxAyAzAAABACADAEAFAGAHAIAJAKALAMANAOAPAQARASATAUAVAWAXAYAZA0A1A2A3A4A5A6A7A8A9A!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~BaC*dbC6ccC7cdC8ceC+dfC,dgC-dhCgdiChdjCidkC.dlC/dmC:dnC;doC#cpC$cqCjdrC=dsCkdtC?duC@dvC[dwC]dxCldyC^dzCmdACndBCodCCpdDC_dEC`dFC{dGC|dHC}dIC~dJCaeKCbeLCceMCqdNCrdOCsdPCc QCdeRCSCTCUCVCWCXCYCZC0C1C2C3C4C5CLc6C7C8C9C!C#C$C%C'C(C)C9c*C!c+CShThUhVhWhXhViWiXiYiZi0i1i2i3i4i5i6i7i8i,C9i!i#i$i%i'i(i)i*i+i-C.C/C:C,i-i.ieefegeheie;Cjekelemene=C?C@C[C]C^C_C`C{C|C}C~CaDbDcDdDeDfDgDhDiDjDkDlDmDnDoDIipDqDrDsDtDuDvDwDxDyDzDADBDCDDDEDFDGDHDIDJDKDLDMDNDODPDQDRDSDTDUDVDWDXDYDZD0D1D2D3D4D5D6D7D8D9D!D#D$D%D'D(D)D*D+b+D,D-D.D/D:D;D=D?D@D[D]D^D_D`D{D|D}D~DaEbEcEdEeEfEgEhEiEjEkElEmEnEoEpEqErEsEtEuEvEwExEyEzEAEBECEDEEEFEPiGEHEIEJEKELEMENEOEPEQERESETEUEVEWEXEYEZE0E1E2E3E4E5E6E7E8E9E!E#E$E%E'E(E)E*E+E,E-E.E/E:E;E=E?E@E[E]E^E_E`E{E|E}E~EaFbFcFQfdFeFfFgFhFiFjFkFlFmFnFoFpFqFrFsFtFuFvFwFxFyFzFAFBFCFDFEFFFGFHF

225 return 28 9 #d! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ {ce f qb[b]b^b_b`b{b|b}bg h i j abbbcbdbrbsbtbebfb_i`i{i|i}ixh~iajbjyhcjdj9eSfTfUfVfWfXfYfZf0f1f2f3f4f5f~b6f7f8f9f!f#f$f%f'f(f)f*f+f,f-f.f/f:f;f=f?f@f[f]fQi^f_f`f{f|f}f~fagbgcgdg|ceg}cfg~cggWbXbadhgbdigjgkglgmgYbZb0b1b2b3b!e#esigbacbcEhIcJcFhGhHhIhJhKhLhMhCiDioe%chbMcib'cNcjb5b:e;e=ek ,bl m n o p q r s t -bu v w x $by z %b.bpe(ckbOclb)cPcmb6b?e@e[eA /bB C D E F G H I J :bK L M N 'bO P (b;bqe*cnbQcob+cRcpb7b]e^e_eQ =bR S T U V W X Y Z ?b0 1 2 3 )b4 5 *b@bNhcd$eWc%edd'e(eOhPh)eed*e+eQh.e/eIi

226 elif out_len == 2: 

227 return result[1] 2a _jhf`jif{jjf|jkf}jlf~j8 akvbbkck9 dk#dek! fkgkhkikjkkkwblkmknkxbokpkqkybrkzbsktkAbukBbvkCbwk# xk$ yk% zk' Ak( Bk) Ck* Dk+ Ek, Fk- Gk. Hk/ Ik: Jk; Kk= Lk? Mk@ Nk[ Ok] Pk^ Qk_ Rk` Sk{ Tk| Uk} Vk~ WkXkYkZk0k1k2k3k4k5k6k7k8k9k!k#k$k%k'k(k)kmf*k+k,knf-kof.kpf/kqf:kDb;krf=ksf?ktf@k[k]kEb^kFb_kGb`kHb{kIb|k=c}k?c~k@cal[cbl]ccl^cdl_cel`cflglhliljlklllml[hnl]hol^hplqlrlsltlulvlwlxlylzlAlBlClDlElFlGlHlIlJlKlLlMlNlOlPlQlRlSlTlUlVlWlXlYlZl0l1l2l3l4l5l6l7l8l9l!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~man9bbncndnufenfngnhninjnknlnmnnnonpnqnrnsntnvfunwfvnwnxnynzn{cAn$dBnCnDne Enf FnGnHnJbInqbJnKbKnLnMnNn[bOn]bPn^bQnRn_bSn`bTn{bUn|bVnWnXn}bYnZn0n1nxf2n|h3n4n5nLb6nMb7nNb8nOb9nPb!n#n$n%n'nd (n)n*n+n5e,n6e-n7e.nyf/nzf:nAf;ng =nh ?ni @nj [nab]nbb^ncb_ndb`nrb{nsb|ntb}nBf~nCfaoDfboEfcoFfdoGfeofogoebhoQbioRbjoSbkoTbloUbmonoVboofbpoqoHfrosoaibi_icidi`ieifi{igihi|iii}itoxhuoIf~ivo8ewoJfKfajLfxoMfbjyoyhzojicjkilidjNfAoOfBoCoDoEoFoGo9eHoSfTfUfVfWfXfYfZf0f1f2f3f6 4f7 5fIoJoKoLoMoNoOoPoQoRoSo6f7f8f9f!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~fagbgcg%ddg|cegToUoVo}cfg~cggWbWoXbXoadhgbdigjgkglgmgYbYoZbZo0b0o1b1o2b2o3b3ozh4o5o6o7o8o9o!o#oAh$o%o'o(o)o*o+o,o-oBh.o/o:o;opi=o?o@oqiChDh[ori]o^o_o`o{o|o}o~oapbpcpdpepfpgphpipjptiuivikplpwixiyimpziAiBinpopppqprpsptpupvpwpxpypzpApBpCpDpEpFpGpHpIpgbacbcJpKpLpMpNpOpPpQpRpSpTpUpVpWpXpYpZp0p1p2p3p4p5p6p7p8p9p!p#p$p%p'p(p)p*p+p,p-p.p/pEhIcJcFhGhHhIhJhKhLhMhoe%chbMcib'cNcjb5b:e;e=ek ,bl m n o p q r s t -bu v w x $by z %b.bpe(ckbOclb)cPcmb6b?e@e[eA /bB C D E F G H I J :bK L M N 'bO P (b;bqe*cnbQcob+cRcpb7b]e^e_eQ =bR S T U V W X Y Z ?b0 1 2 3 )b4 5 *b@bNh:p;p=p?p@p[p]p$e^p_p`p{p|p}p~paqbqcqdqeqfqgqhqiqjqkqlqmqnqoqpqqqrqsqtquqvqwqxqyqzqAqBqCqDqEqFqGqHqIqJqKqLqMqNqOqPqQqRqSqTqUqVqWqXqYqZq0q1q2q3q4q5q6qOhPh)e7q8q9q!q#q$q%q'q(q)q*q+q,q-q.q/q:q;q=q?q@q[q]q^q_q'd`q{q|q}q~qarbrcr4bdrerfrgrhrirjrkrlrmrnrorprqrrrsrtrurvrwrxryrzrArBrCrDrErFrGrHrIrJrKrLrMrNrOrPrQrRrSrTrUrVrWrXrYrZr0r1r2r3r4r5r6r7r8r9r!r#r$r%r'r(r)r*r+r,r-r.r/r:r;r=r?r@r[r]r^r_r`r{r|r}r~rasbscsdsesfsQhgshsisjskslsmsnsospsqsrssstsusvswsxsyszsAsBsCsDsEsFsGsHsIsJsKsLsMsNsOsPsQsRsSsTsUsVsWsXsYsZs0s1s2s3s4s5s6s7s8s9s!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+b+D,D-D.D/D:D;D=D?D@D[D]D^D_D`D{D|D}D~DaEbEcEdEeEfEgEhEiEjEkElEmEnEoEpEqErEsEtEuEvEwExEyEzEAEBECEDEEEFEPiGEHEIEJEKELEMENEOEPEQERESETEUEVEWEXEYEZE0E1E2E3E4E5E6E7E8E9E!E#E$E%E'E(E)E*E+E,E-E.E/E:E;E=E?E@E[E]E^E_E`E{E|E}E~EaFbFcFQfdFeFfFgFhFiFjFkFlFmFnFoFpFqFrFsFtFuFvFwFxFyFzFAFBFCFDFEFFFGFHF

228 else: 

229 return result[1:] 2a 8 9 #d! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ ebfbb ccdcecfcgchcicjckclcmcncocpcqcrcsctcucvcwcxcyczcAcBcCcDcEcFcGcHcKc0c1c2cSi3cTi4cUi5cuhvh*d6c7c8c+d,d-dgdhdid.d/d:d;d#c$cjd=dkd?d@d[d]dld^dmdndodpd_d`d{d|d}d~daebeceqdrdsdc deLc9c!cShThUhVhWhXhViWiXiYiZi0i1i2i3i4i5i6i7i8i9i!i#i$i%i'i(i)i*i+i,i-i.ieefegeheiejekelemene

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: 2UIhfifjfkflf8 vbse9 #d! tekPVIuevewbweWIxbxeXIybzblPAbBbCb# $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ mfnfofpfqfDbrfsftfEbFbGbHbIb=c?c@c[c]c^c_c`cngogpgYIZItdqg2I3Iudrg6I7Ivdsg!I#Iwdtg'I(Ixdug+I,Iyd9bufvgwgxgygvfwfzg/I:I;I{c$d=I?Ie f MFAgJbqbKb@I[I]I[b]b^b^I_b`b{b|b_I`I}bBgxf{ILbMbNbObPbd 5e6e7eyfzfAfg h i j abbbcbdbrbsbtbBfCfDfEfFfGfebQbRbSbTbUbyeVbfb|I}IHf~IaJbJIfcJdJ8eJfKfLfeJfJMfgJhJNfiJOfjJkJlJmJnJoJpJqJrJsJtJ9euJvJNFwJOFxJyJzJAJBJPFCJQFDJEJFJ6 GJ7 HJ~bmi?i@inioi[i]iHjNI,OIJJJKJLJMJNJOJPJQJRJSJTJUJVJWJXJYJZJ0J1J2J3J4J5J6J7J8J9J!J#J$J%J'J(J%d)J|c*J}c+J~c,JWbXbad-Jbd.J/J:J;J=JYbZb0b1b2b3bmPnPoPpPqPrPdI!e#esPtPRFScSFuP?JvP@J[J]J^J_J`JTc{J|J}J~JTFUFVFzeaKbKcKdKeKfKgKhKiKjKkKlKmKhIiInKjIzdIj,cWF-cXFkIwPlIxPmInIoIoKpIqIrIsItIuIvIwIhPiPjPyPzPAPBPCPDPJFgbacbcAdUcBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdUdVdWdXdYdZd0d1d2dAeBeb ccEOdcFOecGOfcHOgcIOhcJOicKOjcLOkcMOlcNOmcOOncLIocMIpcPOqcQOrcROscSOtcTOucUOvcVOwcWOxcXOycYOzcZOAc0OBc1OCcRfDc2OEc3OFc4OGc5OYF6OHcIcJcpKqKoe%chbMcib'cNcjb5b:e;e=ek ,bl m n o p q r s t -bu v w x $by z %b.bpe(ckbOclb)cPcmb6b?e@e[eA /bB C D E F G H I J :bK L M N 'bO P (b;bqe*cnbQcob+cRcpb7b]e^e_eQ =bR S T U V W X Y Z ?b0 1 2 3 )b4 5 *b@brKsKtKcduKvK$eZF0FwKxK1F2F3F4FWc%edd'eyKzKAKBKCKDKEKFKGKHKIKJKKKLKMKNK5FOK6FPK7FQK8FRK9FSK!FTK#FUK$FVKWKXKYKZK0K1K2K3K4K5K6K7K8K9K!K#KxIJjyIzIKjAI(e)e$K%K'Ked*e+e(K)K*K+K%F'FXcYcubreZcKc(F)F'd*F+F,F-F.FVc/F4b:F;F=F?F@F[F]F^F_F`F{F|F}F~F0c1c2caGbGcG3c4c5c!b#bdGeGfGgGhGiGjGkGlGmGnGoGpGqGrGsGtGuGvGwGxGyGzGAGBGCGDGEGFGGGHGIGJGKGLGMGNGOGPGQGRGSGTGUGVGWGXGYGZG0G1G2G3G4G5G6G7G8G9G!G#G$G%G'G(Gth,e-efd(d,KCg-K.K/K)GDg*G;K=K+G?K@KEg[K]K^K,GFg-G`KGg{K|KHg~KCeaLbLBICIcLdLeLfLgLLjDe`ehLMjNjiLjLkLEeFeIgJgKgLgMgNglLmLOgoLPgpLqL.GQg/GsLRgtLuLvL:GSg;GxLTgyLzLUgBLDIGeCLDLELFLGLHLILJLKLLL=GHe{eMLOj?GNLOLPLQLRLSLTLULVLWLXLYLZL0LVg1L2LWg4L5LXg@GIe[G]G^G_G`G{G|G}G~GJeaHEibHcHdHYgKeLeZg0g1g2g3g3d4g5g.c|e}e4d/c5d6g7L8L6dPjQjeHMefHgHhHiHjHkHlHmHnHoHNepHFiqHrHsH7g(L)L8g+LEIOe,L-L.L/L:L;L=L?L@L[L]LtHPeaf^LRjuH_L`L{L|L9g}L~LaMvH!gwHcM#gdMeM$ggMQehMiMFIGIjMkMlMmMnMSjRebfoMTjUjpMqMrMSeTe%g'g(g)g*g+gsMtM,gvM-gwMxMxH.gyHzM/gAMBMCMzH:gAHEM;gFMGM=gIMHIUeJMKMLMMMNMOMPMQMRMSMBHVecfTMVjCHUMVMWMXMYMZM0M1M2M3M4M5M6M7M?g8M9M@g#M$M[gDHWeEHFHGHHHIHJHKHLHMHXeNHGiOHPHQH]gYeZe^g_g`g{g|g7d}g~g:cdfef8d;c9dah'M(M!dWjXjRH0eSHTHUHVHWHXHYHZH0H1H1e2HHi3H4H5Hbh/M:Mch=MII2e?M@M[M]M^M_M`M{M|M}M~M6H3egfaNYj7HbNcNdN.e/eZj0j8H9HeNdhfNgNhN!Heh#H$H%H'HjNkNfhlNmNnN(Hgh)HpNqNhhrNsNtN*Hih+HvNwNjhxNyNzN,Hkh-HBNCNlhDNENFN.Hmh/HHNINnhJNKNLN:Hoh;HNNONphPNQNRN=Hqh?H@HTNUNrhVNWNXN[Hsh]H^HZNoOpO*d6cJI7c8c+d,d-dgdhdid.d/d:d;d#c$cjd=dkd?d@d[d]dld^dmdndodpd_d`d{d|d}d~daebeceqdrdsdc de_H8b`H)dqOrOsOtOuOvOwOxOyOzOAOBOCODO1j2j0N{H|H}HLc~H1N2NaIKI^i7O8O9O!O#O$O%O'O(O)O9c!ceefegeheiejekeiOjOkOlOlemOmenenOEPFPGPHPIPJP3NKPLPMPNPOPPPQPRPSPbITP4N5NUP6NVP3j4j5j6j7j8j9j!j#j$j%j'j+b(jRhJiKiLi)j*j+j,jMi7N8N9N!N#N$N-j%N'N(N)N*N.j/j+N:j;j=jNi?jOi@j[jPf]j^j,N-N.N/N:N;N=N?N@N[N]N^N_N`N{N|N}N~NaObOcOdOeOfOgOhOQf

237 if keep_none: 2UIhfifjfkflf8 vbse9 #d! tekPVIuevewbweWIxbxeXIybzblPAbBbCb# $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ mfnfofpfqfDbrfsftfEbFbGbHbIb=c?c@c[c]c^c_c`cYIZI2I3I6I7I!I#I'I(I+I,Iufvfwf{c$de f JbqbKb@I[I]I[b]b^b^I_b`b{b|b_I`I}bxf{ILbMbNbObPbd 5e6e7eyfzfAfg h i j abbbcbdbrbsbtbBfCfDfEfFfGfebQbRbSbTbUbyeVbfb|I}IHf~IaJbJIfcJdJ8eJfKfLfeJfJMfgJhJNfiJOfjJkJlJmJnJoJpJqJrJsJtJ9e6 7 ~bmi?i@inioi[i]iHjNI,O%d|c}c~cWbXbadbdYbZb0b1b2b3bmPnPoPpPqPrP!e#esPtPRFScSF?J@J[J]J^J_J`JTc{J|J}J~JTFUFVFzeaKbKcKdKeKfKgKhKiKjKkKlKmKnKzd,c-cxPoKgbacbcUcIcJcpKqKoe%chbMcib'cNcjb5b:e;e=ek ,bl m n o p q r s t -bu v w x $by z %b.bpe(ckbOclb)cPcmb6b?e@e[eA /bB C D E F G H I J :bK L M N 'bO P (b;bqe*cnbQcob+cRcpb7b]e^e_eQ =bR S T U V W X Y Z ?b0 1 2 3 )b4 5 *b@brKsKtKcduK$eZF0FwKxK1F2F3F4FWc%edd'e5FOK6FPK7FQK8FRK9FSK!FTK#FUK$FVKJjKj)e$K%K'Ked%F'F(F)F'd*F+F,F-F.FVc/F4b:F;F=F?F@F[F]F^F_F`F{F|F}F~F3c4c5c!b#bdGeGfGgGhGiGjGkGlGmGnGoGpGqGrGsGtGuGvGwGxGyGzGAGBGCGDGEGFGGGHGIGJGKGLGMGNGOGPGQGRGSGTGUGVGWGXGYGZG0G1G2G3G4G5G6G7G8G9G!G#G$G%G'G(Gfd,K-K.K/K)G*G;K=K+G?K@K[K]K^K,G-G`K{K|K~KCeaLbLBICIcLdLeLfLgLLjDe`ehLMjNjiLjLkLEeFelLmLoLpLqL.G/GsLtLuLvL:G;GxLyLzLBLDIGeCLDLELFLGLHLILJLKLLL=GHe{eMLOj?GNLOLPLQLRLSLTLULVLWLXLYLZL0L1L2L4L5L@GIe[G]G^G_G`G{G|G}G~GJeaHEibHcHdHKeLe|e}e4d5d7L8LPjQjeHMefHgHhHiHjHkHlHmHnHoHNepHFiqHrHsH(L)L+LEIOe,L-L.L/L:L;L=L?L@L[L]LtHPeaf^LRjuH_L`L{L|L}L~LaMvHwHcMdMeMgMQehMiMFIGIjMkMlMmMnMSjRebfoMTjUjpMqMrMSeTesMtMvMwMxMxHyHzMAMBMCMzHAHEMFMGMIMHIUeJMKMLMMMNMOMPMQMRMSMBHVecfTMVjCHUMVMWMXMYMZM0M1M2M3M4M5M6M7M8M9M#M$MDHWeEHFHGHHHIHJHKHLHMHXeNHGiOHPHQHYeZedfef8d9d'M(MWjXjRH0eSHTHUHVHWHXHYHZH0H1H1e2HHi3H4H5H/M:M=MII2e?M@M[M]M^M_M`M{M|M}M~M6H3egfaNYj7HbNcNdNZj0j8H9HeNfNgNhN!H#H$H%H'HjNkNlNmNnN(H)HpNqNrNsNtN*H+HvNwNxNyNzN,H-HBNCNDNENFN.H/HHNINJNKNLN:H;HNNONPNQNRN=H?H@HTNUNVNWNXN[H]H^HZN6cc de_H8b`H)d0N{H|H}H~H1N2NaI^i9cGPHPIPJPLPMPNPOPRPbITPUP6N+bRhJiKiLiMi7N8N9N!N#N$N%N'N(N)N*N+NNiOi,N-N.N/N:N;N=N?N@N[N]N^N_N`N{N|N}N~NaObOcOdOeOfOgOhOQf

238 return options 2e f @I[I]I[b]b^b^I_b`b{b|b_I`I}b{I5e6e7eg h i j abbbcbdb}I~IaJIfdJ8eKfeJMfhJiJjJlJnJpJrJtJ!e#ezd,c-cpKqKrKsKtKuK$eZF0FwKxK1F2F3F4FWc%edd'e5FOK6FPK7FQK8FRK9FSK!FTK#FUK$FVKJjKj$K%K'KVc!b#b,K=K@K~KCeaLbLBICIcLdLeLfLgLLjDe`ehLMjNjiLjLkLoLsLDIEI|LgMQehMiMFIGIjMkMlMmMnMSjRebfoMTjUjpMqMrMvMzMHIIIeNkNqNwNCNINONUN+bRh7N8N9N!N#N$N%N'N(N)N*N+NNiOi,N-N.N/N:N;N=N?N@N[N]N^N_N`N{N|N}N~NaObOcOdOeOfOgOhO

239 return cls() 2UIhfifjfkflf8 vbse9 #d! tekPVIuevewbweWIxbxeXIybzblPAbBbCb# $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ mfnfofpfqfDbrfsftfEbFbGbHbIb=c?c@c[c]c^c_c`cYIZI2I3I6I7I!I#I'I(I+I,Iufvfwf{c$de f JbqbKb[b]b^b_b`b{b|b}bxfLbMbNbObPbd 5e6e7eyfzfAfg h i j abbbcbdbrbsbtbBfCfDfEfFfGfebQbRbSbTbUbyeVbfb|IHfbJcJ8eJfLffJgJNfOfkJmJoJqJsJ9e6 7 ~bmi?i@inioi[i]iHjNI,O%d|c}c~cWbXbadbdYbZb0b1b2b3bmPnPoPpPqPrPsPtPRFScSF?J@J[J]J^J_J`JTc{J|J}J~JTFUFVFzeaKbKcKdKeKfKgKhKiKjKkKlKmKnKzd,c-cxPoKgbacbcUcIcJcoe%chbMcib'cNcjb5b:e;e=ek ,bl m n o p q r s t -bu v w x $by z %b.bpe(ckbOclb)cPcmb6b?e@e[eA /bB C D E F G H I J :bK L M N 'bO P (b;bqe*cnbQcob+cRcpb7b]e^e_eQ =bR S T U V W X Y Z ?b0 1 2 3 )b4 5 *b@bcdWc)eed%F'F(F)F'd*F+F,F-F.FVc/F4b:F;F=F?F@F[F]F^F_F`F{F|F}F~F3c4c5c!b#bdGeGfGgGhGiGjGkGlGmGnGoGpGqGrGsGtGuGvGwGxGyGzGAGBGCGDGEGFGGGHGIGJGKGLGMGNGOGPGQGRGSGTGUGVGWGXGYGZG0G1G2G3G4G5G6G7G8G9G!G#G$G%G'G(Gfd-K.K/K)G*G;K+G?K[K]K^K,G-G`K{K|KBICILjNjEeFelLmLpLqL.G/GtLuLvL:G;GxLyLzLBLDIGeCLDLELFLGLHLILJLKLLL=GHe{eMLOj?GNLOLPLQLRLSLTLULVLWLXLYLZL0L1L2L4L5L@GIe[G]G^G_G`G{G|G}G~GJeaHEibHcHdHKeLe|e}e4d5d7L8LPjQjeHMefHgHhHiHjHkHlHmHnHoHNepHFiqHrHsH(L)L+LEIOe,L-L.L/L:L;L=L?L@L[L]LtHPeaf^LRjuH_L`L{L}L~LaMvHwHcMdMeMFIGISjUjSeTesMtMwMxMxHyHAMBMCMzHAHEMFMGMIMHIUeJMKMLMMMNMOMPMQMRMSMBHVecfTMVjCHUMVMWMXMYMZM0M1M2M3M4M5M6M7M8M9M#M$MDHWeEHFHGHHHIHJHKHLHMHXeNHGiOHPHQHYeZedfef8d9d'M(MWjXjRH0eSHTHUHVHWHXHYHZH0H1H1e2HHi3H4H5H/M:M=MII2e?M@M[M]M^M_M`M{M|M}M~M6H3egfaNYj7HbNcNdNZj0j8H9HfNgNhN!H#H$H%H'HjNlNmNnN(H)HpNrNsNtN*H+HvNxNyNzN,H-HBNDNENFN.H/HHNJNKNLN:H;HNNPNQNRN=H?H@HTNVNWNXN[H]H^HZN6cc de_H8b`H)d0N{H|H}H~H1N2NaI^i9cGPHPIPJPLPMPNPOPRPbITPUP6N+bRhJiKiLiMiQf

240 elif isinstance(options, cls): 28 vbse9 ! teuevewbwexbxeybzbAbBbCb# $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ DbEbFbGbHbIbngogpgtdqgudrgvdsgwdtgxdugyd9bvgwgxgygzg/I:I;I{c$d=I?Ie f MFAgJbqbKbBgLbMbNbObPbd g h i j abbbcbdbrbsbtbebQbRbSbTbUbyeVbfbuJvJNFwJOFxJyJzJAJBJPFCJQFDJEJFJ6 GJ7 HJ~bmi?i@inioi[i]iHjNI,OIJJJKJLJMJNJOJPJQJRJSJTJUJVJWJXJYJZJ0J1J2J3J4J5J6J7J8J9J!J#J$J%J'J(J)J*J+J,JWbXb-J.J/J:J;J=JYbZb0b1b2b3bdIRFScSFuPvPTcTFUFVFzehIiIjIzdIj,cWF-cXFkIwPlImInIoIpIqIrIsItIuIvIwIhPiPjPyPzPAPBPCPDPJFgbAdUcBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdUdVdWdXdYdZd0d1d2dAeBeb ccEOdcFOecGOfcHOgcIOhcJOicKOjcLOkcMOlcNOmcOOncLIocMIpcPOqcQOrcROscSOtcTOucUOvcVOwcWOxcXOycYOzcZOAc0OBc1OCcRfDc2OEc3OFc4OGc5OYF6OHccdvKZF0F1F2F3F4FWcddyKzKAKBKCKDKEKFKGKHKIKJKKKLKMKNK5F6F7F8F9F!F#F$FWKXKYKZK0K1K2K3K4K5K6K7K8K9K!K#KxIJjyIzIKjAI(eed*e+e(K)K*K+KXcYcubreZcKcVc0c1c2caGbGcG!b#bth,e-efd(dCgDgEgFgGgHgCeDeEeFeIgJgKgLgMgNgOgPgQgRgSgTgUgGeHeVgWgXgIeJeYgKeLeZg0g1g2g3g3d4g5g.c/c6g6dMeNe7g8gOePe9g!g#g$gQeReSeTe%g'g(g)g*g+g,g-g.g/g:g;g=gUeVe?g@g[gWeXe]gYeZe^g_g`g{g|g7d}g~g:c;cah!d0e1ebhch2e3e.e/eZj0jdhehfhghhhihjhkhlhmhnhohphqhrhshoOpO*dJI7c8c+d,d-dgdhdid.d/d:d;d#c$cjd=dkd?d@d[d]dld^dmdndodpd_d`d{d|d}d~daebeceqdrdsdqOrOsOtOuOvOwOxOyOzOAOBOCODO1j2jLcKI^i7O8O9O!O#O$O%O'O(O)O!ceefegeheiejekeiOjOkOlOlemOmenenOEPFP3NKPPPQPSPbI4N5NVP3j4j5j6j7j8j9j!j#j$j%j'j+b(jRhJiKiLi)j*j+j,jMi-j.j/j:j;j=jNi?jOi@j[jPf]j^j

241 return options 28 vbse9 ! teuevewbwexbxeybzbAbBbCb# $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ DbEbFbGbHbIbngogpgtdqgudrgvdsgwdtgxdugyd9bvgwgxgygzg/I:I;I{c$d=I?Ie f MFAgJbqbKbBgLbMbNbObPbd g h i j abbbcbdbrbsbtbebQbRbSbTbUbyeVbfbuJvJNFwJOFxJyJzJAJBJPFCJQFDJEJFJ6 GJ7 HJ?i@i[i]i,OIJJJKJLJMJNJOJPJQJRJSJTJUJVJWJXJYJZJ0J1J2J3J4J5J6J7J8J9J!J#J$J%J'J(J)J*J+J,JWbXb-J.J/J:J;J=JYbZb0b1b2b3bRFScSFuPvPTcTFUFVFzehIiIjIzdIj,cWF-cXFkIwPlImInIoIpIqIrIsItIuIvIwIhPiPjPyPzPAPBPCPDPJFgbAdUcBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdUdVdWdXdYdZd0d1d2dAeBeb ccEOdcFOecGOfcHOgcIOhcJOicKOjcLOkcMOlcNOmcOOncLIocMIpcPOqcQOrcROscSOtcTOucUOvcVOwcWOxcXOycYOzcZOAc0OBc1OCcRfDc2OEc3OFc4OGc5OYF6OHccdvKZF0F1F2F3F4FWcddyKzKAKBKCKDKEKFKGKHKIKJKKKLKMKNK5F6F7F8F9F!F#F$FWKXKYKZK0K1K2K3K4K5K6K7K8K9K!K#K(eed*e+e(K)K*K+KXcYcubreZcKcVc0c1c2caGbGcG!b#bth,e-efd(dCgDgEgFgGgHgCeDeEeFeIgJgKgLgMgNgOgPgQgRgSgTgUgGeHeVgWgXgIeJeYgKeLeZg0g1g2g3g3d4g5g.c/c6g6dMeNe7g8gOePe9g!g#g$gQeReSeTe%g'g(g)g*g+g,g-g.g/g:g;g=gUeVe?g@g[gWeXe]gYeZe^g_g`g{g|g7d}g~g:c;cah!d0e1ebhch2e3e.e/edhehfhghhhihjhkhlhmhnhohphqhrhshoOpO*dJI7c8c+d,d-dgdhdid.d/d:d;d#c$cjd=dkd?d@d[d]dld^dmdndodpd_d`d{d|d}d~daebeceqdrdsdqOrOsOtOuOvOwOxOyOzOAOBOCODO1j2jKI^i7O8O9O!O#O$O%O'O(O)O!ceefegeheiejekeiOjOkOlOlemOmenenOEPFP3NKPPPQPSPbI4N5NVP3j4j5j6j7j8j9j!j#j$j%j'j+b(jRhJiKiLi)j*j+j,jMi-j.j/j:j;j=jNi?jOi@j[jPf]j^j

242 elif isinstance(options, dict): 2~bminioiHjNIdIJFAdxIJjyIzIKjAIZj0jLcPf

243 return cls(**options) 2~bminioiHjNIAdxIJjyIzIKjAIZj0jLc

244 else: 

245 raise TypeError( 2dIJFPf

246 f"The {options_description} must be provided as an object " 2dIJFPf

247 f"of type {cls.__name__} or as a dict with valid {options_description}. " 2dIJFPf

248 f"The provided object is '{options}'." 2dIJFPf

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#bgdhdidkdodpdqdrdsdF8

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: 2eI

273 """ 

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

275 """ 

276  

277 @functools.wraps(wrapped_function) 2eI

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) 2eI

283 result = wrapped_function(*args, **kwargs) 2eI

284  

285 return result 2eI

286  

287 return inner 2eI

288  

289 return outer 2eI

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) 26 7 ScTchIiIjIIj,c-ckIlImInIoIpIqIrIsItIuIvIwIhPiPjPgbAdBdCdDdEdFdGdHdIdJdKdLdMdNdOdPdQdRdSdTdUdVdWdXdYdZd0d1d2dLIMIJI#c$cjdldmdndG8H8I8J8K81j2jL8M8N8s8t8KIShThUhVhWhXhO8P8

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$cs8t8

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() 2XcYcubreZc

327 self._entered = False 2XcYcubreZc

328  

329 def __enter__(self): 

330 self._stack.__enter__() 2XcYcubreZc

331 self._entered = True 2XcYcubreZc

332 return self 2XcYcubreZc

333  

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

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

336 self._entered = False 2XcYcubreZc

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

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: 2XcYcubreZc

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

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

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() 2XcYcubreZc

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 2th,e-efd/i(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: 2vjwjxjyjzjAjBjCjDjEjFjGjLFd %chbMcib'cNcjbk ,bl m n o p q r s t -bu v w x $by z %b.b(ckbOclb)cPcmbA /bB C D E F G H I J :bK L M N 'bO P (b;b*cnbQcob+cRcpbQ =bR S T U V W X Y Z ?b0 1 2 3 )b4 5 *b@b

374 val = (<uint8_t*>ptr)[0] 2hbl m n u v kbB C D K L nbR S T 0 1

375 elif width == 2: 

376 val = (<uint16_t*>ptr)[0] 2vjxjzjBjDjFjMcibo p q w x $bOclbE F G M N 'bQcobU V W 2 3 )b

377 elif width == 4: 

378 val = (<uint32_t*>ptr)[0] 2wjyjAjCjEjGjd Ncjbk r s t y z %bPcmbA H I J O P (bRcpbQ X Y Z 4 5 *b

379 else: 

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

381 return (val, <unsigned int>width) 2vjwjxjyjzjAjBjCjDjEjFjGjd hbMcibNcjbk l m n o p q r s t u v w x $by z %bkbOclbPcmbA B C D E F G H I J K L M N 'bO P (bnbQcobRcpbQ R S T U V W X Y Z 0 1 2 3 )b4 5 *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`c0I1Ivjwj4I5Ixjyj8I9IzjAj$I%IBjCj)I*IDjEj-I.IFjGjLFe f [b]b^b_b`b{b|b}bd 4eg h i j oe%chbMcib'cNcjb5b:e;e=ek ,bl m n o p q r s t -bu v w x $by z %b.bpe(ckbOclb)cPcmb6b?e@e[eA /bB C D E F G H I J :bK L M N 'bO P (b;bqe*cnbQcob+cRcpb7b]e^e_eQ =bR S T U V W X Y Z ?b0 1 2 3 )b4 5 *b@b:K_K}KMjnLrLwLALOj3L6LEi3d~e6dPjQj9L!L#L$L%L'LFi*LRjbMfMTjuMyMDMHMVj!M%MGi7dff!dWjXj)M*M+M,M-M.MHi;MYjiNoNuNANGNMNSNYN

411 byte_val = value 2=c?c@c[c]c^c_c`c0I1I4I5I8I9I$I%I)I*I-I.Ie f [b]b^b_b`b{b|b}b4eg h i j 5b:e;e=e6b?e@e[e7b]e^e_e:K_K}KMjnLrLwLALOj3L6LEi3d~e6dPjQj9L!L#L$L%L'LFi*LRjbMfMTjuMyMDMHMVj!M%MGi7dff!dWjXj)M*M+M,M-M.MHi;MYjiNoNuNANGNMNSNYN

412 return (<unsigned int>byte_val, <unsigned int>1) 2=c?c@c[c]c^c_c`c0I1I4I5I8I9I$I%I)I*I-I.Ie f [b]b^b_b`b{b|b}b4eg h i j 5b6b7b:K_K}KMjnLrLwLALOj3L6LEi3d~e6dPjQj9L!L#L$L%L'LFi*LRjbMfMTjuMyMDMHMVj!M%MGi7dff!dWjXj)M*M+M,M-M.MHi;MYjiNoNuNANGNMNSNYN

413  

414 if isinstance(value, bytes): 2vjwjxjyjzjAjBjCjDjEjFjGjLFd oe%chbMcib'cNcjbk ,bl m n o p q r s t -bu v w x $by z %b.bpe(ckbOclb)cPcmbA /bB C D E F G H I J :bK L M N 'bO P (b;bqe*cnbQcob+cRcpbQ =bR S T U V W X Y Z ?b0 1 2 3 )b4 5 *b@b

415 return _read_fill_ptr(<const char*><bytes>value, <Py_ssize_t>len(value)) 2vjwjxjyjzjAjBjCjDjEjFjGjLF%chbMcib'cNcjb(ckbOclb)cPcmb*cnbQcob+cRcpb

416  

417 if PyObject_GetBuffer(value, &buf, PyBUF_SIMPLE) != 0: 2d oek ,bl m n o p q r s t -bu v w x $by z %b.bpeA /bB C D E F G H I J :bK L M N 'bO P (b;bqeQ =bR S T U V W X Y Z ?b0 1 2 3 )b4 5 *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 k ,bl m n o p q r s t -bu v w x $by z %b.bA /bB C D E F G H I J :bK L M N 'bO P (b;bQ =bR S T U V W X Y Z ?b0 1 2 3 )b4 5 *b@b

423 return _read_fill_ptr(<const char*>buf.buf, buf.len) 2d k ,bl m n o p q r s t -bu v w x $by z %b.bA /bB C D E F G H I J :bK L M N 'bO P (b;bQ =bR S T U V W X Y Z ?b0 1 2 3 )b4 5 *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: 2-O.ONFOF/O:O;O=OPFQF6 7 ~bmi?i@inioi[i]i?O@Ou8[Ov8w8]Ox8y8^Oz8A8_OB8C8`OD8Qi{OE8|O}O~OaPbPcP%d|c}c~cWbXbadbddPePfPgPYbZb0b1b2b3bth,e-efd/i(d

432 return 2-O.ONFOF/O:O;O=OPFQF6 7 ~bmi?i@inioi[i]i?O@Ou8[Ov8w8]Ox8y8^Oz8A8_OB8C8`OD8Qi{OE8|O}O~OaPbPcP%d|c}c~cWbXbadbddPePfPgPYbZb0b1b2b3b(d

433 _fork_warning_checked = True 2th,e-efd/i(d

434  

435 # Common warning message parts 

436 common_message = ( 

437 "CUDA does not support. Forked subprocesses exhibit undefined behavior, " 2th,e-efd/i(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: 2th,e-efd/i(d

444 start_method = multiprocessing.get_start_method() 2th,e-efd/i(d

445 if start_method == "fork": 2th,e-efd(d

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

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

448 except RuntimeError: 2/i

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": 2/i

452 message = ( 

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

454 f"which {common_message}" 

455 ) 

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