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
« 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
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 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
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
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
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
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: 2'HoOpO_H8b`H)dqOrOsOtOuOvOwOxOyOzOAOBOCODO1j2j
109 return 0 2'HoOpO_H8b`H)dqOrOsOtOuOvOwOxOyOzOAOBOCODO1j2j
110 with gil: 28b
111 _raise_nvvm_error(prog, err) 28b
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
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
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
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 _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()}")
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
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
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}")
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
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 )
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
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: 2eI
273 """
274 A decorator that actually wraps the function for checking preconditions.
275 """
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
285 return result 2eI
287 return inner 2eI
289 return outer 2eI
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
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
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() 2XcYcubreZc
327 self._entered = False 2XcYcubreZc
329 def __enter__(self):
330 self._stack.__enter__() 2XcYcubreZc
331 self._entered = True 2XcYcubreZc
332 return self 2XcYcubreZc
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
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
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
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 2th,e-efd/i(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: 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
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`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
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
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)
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
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 )
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