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

1# SPDX-FileCopyrightText: Copyright (c) 2024-2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved. (empty)

2# 

3# SPDX-License-Identifier: Apache-2.0 

4  

5import functools (empty)

6from functools import partial (empty)

7import importlib.metadata (empty)

8import multiprocessing (empty)

9import platform (empty)

10import warnings (empty)

11from collections import namedtuple (empty)

12from collections.abc import Sequence (empty)

13from contextlib import ExitStack (empty)

14from typing import Callable (empty)

15  

16try: (empty)

17 from cuda.bindings import driver, nvrtc, runtime (empty)

18except ImportError: 

19 from cuda import cuda as driver 

20 from cuda import cudart as runtime (empty)

21 from cuda import nvrtc 

22  

23from cuda.bindings.nvvm import nvvmError (empty)

24from cuda.bindings.nvjitlink import nvJitLinkError (empty)

25  

26from cuda.bindings cimport cynvrtc, cynvvm, cynvjitlink 

27  

28from cuda.core._utils.driver_cu_result_explanations import DRIVER_CU_RESULT_EXPLANATIONS (empty)

29from cuda.core._utils.runtime_cuda_error_explanations import RUNTIME_CUDA_ERROR_EXPLANATIONS (empty)

30  

31  

32class CUDAError(Exception): (empty)

33 pass (empty)

34  

35  

36class NVRTCError(CUDAError): (empty)

37 pass (empty)

38  

39  

40  

41ComputeCapability = namedtuple("ComputeCapability", ("major", "minor")) (empty)

42  

43  

44def cast_to_3_tuple(label, cfg): (empty)

45 cfg_orig = cfg 122 ctx2a k l m n o p q c r d e f s g h t u v w x y z A B C D E F G H I J K L M N O P Q R S T U V W i j uchdX Cc5nxe#d6nyevcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c9c!c#c$c%c'c(c)c*c+cwczeAe,c]d^d5d-c.c6dBe/c_d`d7d:c;c8dCeDeEeFeGeHeIeJeleVnWnDl7nme

46 if isinstance(cfg, int): 121 ctx2k l m n o p q c r d e f s g h t u v w x y z A B C D E F G H I J K L M N O P Q R S T U V W i j uchdX Cc5nxe#d6nyevcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c9c!c#c$c%c'c(c)c*c+cwczeAe,c]d^d5d-c.c6dBe/c_d`d7d:c;c8dCeDeEeFeGeHeIeJeleVnWnDl7nme

47 cfg = (cfg,) 89 ctx2k l m n o p q c r d e f s g h t u v w x y z A B C D E F G H I J K L M N O P Q R S T U V W i j uchdX Cc5n#d6nNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c9c!c#c$c%c'c(c)c*c+cwcleVn

48 else: 

49 common = "must be an int, or a tuple with up to 3 ints" 34 ctx2xe#dyevczeAe,c]d^d5d-c.c6dBe/c_d`d7d:c;c8dCeDeEeFeGeHeIeJeleWnDl7nme

50 if not isinstance(cfg, tuple): 34 ctx2xe#dyevczeAe,c]d^d5d-c.c6dBe/c_d`d7d:c;c8dCeDeEeFeGeHeIeJeleWnDl7nme

51 raise ValueError(f"{label} {common} (got {type(cfg)})") 1 ctx27n

52 if len(cfg) > 3: 33 ctx2xe#dyevczeAe,c]d^d5d-c.c6dBe/c_d`d7d:c;c8dCeDeEeFeGeHeIeJeleWnDlme

53 raise ValueError(f"{label} {common} (got tuple with length {len(cfg)})") 1 ctx2Wn

54 if any(not isinstance(val, int) for val in cfg): 32 ctx2xe#dyevczeAe,c]d^d5d-c.c6dBe/c_d`d7d:c;c8dCeDeEeFeGeHeIeJeleDlme

55 raise ValueError(f"{label} {common} (got {cfg})") 1 ctx2Dl

56 if any(val < 1 for val in cfg): 118 ctx2k l m n o p q c r d e f s g h t u v w x y z A B C D E F G H I J K L M N O P Q R S T U V W i j uchdX Cc5nxe#d6nyevcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c9c!c#c$c%c'c(c)c*c+cwczeAe,c]d^d5d-c.c6dBe/c_d`d7d:c;c8dCeDeEeFeGeHeIeJeleVnme

57 plural_s = "" if len(cfg) == 1 else "s" 3 ctx2#dVnme

58 raise ValueError(f"{label} value{plural_s} must be >= 1 (got {cfg_orig})") 3 ctx2#dVnme

59 return cfg + (1,) * (3 - len(cfg)) 116 ctx2k l m n o p q c r d e f s g h t u v w x y z A B C D E F G H I J K L M N O P Q R S T U V W i j uchdX Cc5nxe#d6nyevcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c9c!c#c$c%c'c(c)c*c+cwczeAe,c]d^d5d-c.c6dBe/c_d`d7d:c;c8dCeDeEeFeGeHeIeJele

60  

61  

62def _reduce_3_tuple(t: tuple): (empty)

63 return t[0] * t[1] * t[2] 1 ctx2Cc

64  

65  

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

67 if err != cydriver.CUresult.CUDA_SUCCESS: 1438 ctx2a xqyqzqAqBqCqk Dql Eqm Fq=cGqn Hq?cIqo Jqp Kqq LqDcMqEcNqFcOq@cPq[cQq]cRqc Sqr Tqd Uqe Vqf Wqs Xqg Yqh Zqt 0qu 1qv 2q^c3q_c4q`c5q{c6q|c7q}c8qw 9qx !qy #qz $qA %qB 'qC (qD )qE *qF +qG ,qH -qI .qJ /qK :qL ;qM =qN ?qO @qP [qQ ]qR ^qS _qT `qU {qV |qW }q]p8n^p9nEl!nFl#n_p$n`p%nGl'nHl(ni )nj *n{d~q|darnebroecr}ddr~derpefrqegrKehrLeirIljr{p+n|p,n}p-n~p.naq/nbq:nreidcq;n~c=nad?nkrbd@ncd[nxclrycmrdd]ned^ndq_neq`nfq{ngq|nZ nr0 or1 pr2 qr3 rr4 sr$d}ntr~nuraovrbowrcoxrdoyrzr%djdGcAreoBrCrDrErFrGrHrIrJr'dKrLrMrNrOrPrQrRrSrTrUrVrWrXrYrZr0r1r2r3r4r5r6r7r8r9r!r#r$r%r'r(r)r*r+r,r-r.r/r:r;r=r?r@r[r]r^r_r`r{r|r}r~rasbscsdsesfsgshsisjskslsmsnsospsqsrssstsusvswsxsyszsAsBsCsDsEsFsGsHsIsJsKsLsMsNsOsPsQsRsSsTsUsVsWsXsYsZs0s1s2s3s4s5s6s7s8s9s!s#s$s%s's(s)s*s+s,s-s.s/s:s;s=s?s@s[s]s^s_s`s{s|s}s~satbtctdtetaeftgthtbe(d)dceJlucKlitfojtgoktholtiomtjontkootloptmoqtLlrtMlstNltthdutvtwtxtytztAtBtCtDtseteEtFtnoGtHtItJtX TbUbCcKtLtMtvcNtNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c9c!c#c$c%c'c(c)c*c+c5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | OlOtPt*dVbWb+dQt,d-d.dPlQlRt/d:d;d=ddeeekdStldTtXbUtmdVtYbWtndXtodYtZbZt0b0tpd1tqd2trd3t1b4tsd5t2b6t3b7t4b8t5b9t6b!t7b#t8b$t9b%t!b'ttd(t#b)t$b*t%b+t'b,tud-t(b.t)b/tvd:twd;txd=tyd?t*b@tzd[t+b]tAd^tBd_t,b`t-b{tCd|tDd}tEd~t.bauFdbu/bcu:bdu;beu=bfu?bgu@bhu[biu]bju^bkuGdlu_bmu`bnu{bou|bpuHdqu}bru~bsuIdtuJduuKdvuLdwuacxuMdyubczuNdAuOdBuccCudcDuPdEuQdFuRdGuecHuSdIufcJugcKuhcLuicMujcNukcOulcPumcQuncRuTdSuocTupcUuqcVurcWuUdXuscYutcZuVd0uWd1u?doopoHc2u3uqo4uIc5uXd6uYd7uZd8uro9uso!uto#uuo$uvo%uwo'uxo(uyo)uzo*uAo+uBo,uCo-uDo.uEo/uFo:uGo;uHo=u?u@uIo[u]u^uJo_u`u{uKo|u}u~uLoavbvcvModvevfvNogvhvivOojvkvlvPomvQonvRoovSopvToqvUorvVosvWotvXouvYovvZowv0oxv1oyv2ozv3oAv4oBvXnCvYnDvZnEv0nFv1nGv2nHv5oIv@d[d0d6oJvJcKv7o8ozcAcLvMvY NvBcOvPv} QvRlSlRvTlSvUlwcTvVl~ WlUvXlVvYlfeZlWv0lge1lXv2lYv3lZv4l0v5l1v6l7l8l9l!lbbcb2v#l3v$l4v%l5v'l6v(l7v)l8v*l9v+l!v,l#v-l$v.l%v/l'v:l(v;l)v=l*v?l+v@l,v[l-v]l.v^l/v_l:v`l;v{l=v|l}l~l?vam@vbm[vcm]vdm^vem_vfm`vgm{vhm|vim}vjm~vkmawlmbwmmcwnmdwomewpmfwqmgwrmhwsmiwtmjwumvmkwwmlwxmmwymnwzmowAmpwBmqwCmrwDmswEmtwFmuwGmvwHmwwImxwJmywKmzwLmMm!dAw1dBw2dCwKcDwfdEwNmFwGw9oOmPm!oQmRm#oSmHwIw$oTmUm%ohe,cieVmWm]dXmYmJw'oZm^d0m(oMe)o1m5d2m-c.c3mNe*o+o4m6d5m6mKwLw,o7m8m-oje/cke9m!m_d#m$mMw.o%m`d'm/oOe:o(m7d)m:c;c*mPe;o=o+m8d,m-mNwOw?o.m/m:mLc;m=mMc@o?mPwQw[o@m[m]o]mRwSw^o^m_m_o`mTwUw`o{m|m{o}mVwWw|o~man}obnXwYw~ocndnapenZw0wbpfngnhncpin1w2wdpjnknlnepdb3web4wfpfb5wgb6w7w8w9w!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~waxbxcxdxexfxmngdgxhxixjxkxlxmxnnnxoxpxqxrxsxtxQeReonpnqnabuxrngphpsnvxwxxxhbyxzxAxBxibCxDxExipFxjpGxkpHxlpIxmpJxnpKxopLxppMxqpNxrpOxspPxtpQxupRxvpSxwpTxxpUxtnVxypWxzpXxApYxBpZxCp0xDp1x2x3x4x5x6x7x8x9x!x#x$x%x'x(x)x*x+x,x-x.x/x:x;x=x?xunhqiqjqvnwnkqlqxnynmqnqznoqAnBnpqCnDnEnFnGnHnInJnKnqqLnMnNnrqOnPnQnRnSnTnsqtq@x[x]xUn^x_x`x{x|x}x~xaybycydyeyfygyhyiyjykylymynyoypyqyrysytyuyvywyxyyyzyAyByCyDyEy

68 return _check_driver_error(err) 4 ctx2sete~ gd

69 return 0 1436 ctx2a xqyqzqAqBqCqk Dql Eqm Fq=cGqn Hq?cIqo Jqp Kqq LqDcMqEcNqFcOq@cPq[cQq]cRqc Sqr Tqd Uqe Vqf Wqs Xqg Yqh Zqt 0qu 1qv 2q^c3q_c4q`c5q{c6q|c7q}c8qw 9qx !qy #qz $qA %qB 'qC (qD )qE *qF +qG ,qH -qI .qJ /qK :qL ;qM =qN ?qO @qP [qQ ]qR ^qS _qT `qU {qV |qW }q]p8n^p9nEl!nFl#n_p$n`p%nGl'nHl(ni )nj *n{d~q|darnebroecr}ddr~derpefrqegrKehrLeirIljr{p+n|p,n}p-n~p.naq/nbq:nreidcq;n~c=nad?nkrbd@ncd[nxclrycmrdd]ned^ndq_neq`nfq{ngq|nZ nr0 or1 pr2 qr3 rr4 sr$d}ntr~nuraovrbowrcoxrdoyrzr%djdGcAreoBrCrDrErFrGrHrIrJr'dKrLrMrNrOrPrQrRrSrTrUrVrWrXrYrZr0r1r2r3r4r5r6r7r8r9r!r#r$r%r'r(r)r*r+r,r-r.r/r:r;r=r?r@r[r]r^r_r`r{r|r}r~rasbscsdsesfsgshsisjskslsmsnsospsqsrssstsusvswsxsyszsAsBsCsDsEsFsGsHsIsJsKsLsMsNsOsPsQsRsSsTsUsVsWsXsYsZs0s1s2s3s4s5s6s7s8s9s!s#s$s%s's(s)s*s+s,s-s.s/s:s;s=s?s@s[s]s^s_s`s{s|s}s~satbtctdtetaeftgthtbe(d)dceJlucKlitfojtgoktholtiomtjontkootloptmoqtLlrtMlstNltthdutvtwtxtytztAtBtCtDtEtFtnoGtHtItJtX TbUbCcKtLtMtvcNtNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c9c!c#c$c%c'c(c)c*c+c5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | OlOtPt*dVbWb+dQt,d-d.dPlQlRt/d:d;d=ddeeekdStldTtXbUtmdVtYbWtndXtodYtZbZt0b0tpd1tqd2trd3t1b4tsd5t2b6t3b7t4b8t5b9t6b!t7b#t8b$t9b%t!b'ttd(t#b)t$b*t%b+t'b,tud-t(b.t)b/tvd:twd;txd=tyd?t*b@tzd[t+b]tAd^tBd_t,b`t-b{tCd|tDd}tEd~t.bauFdbu/bcu:bdu;beu=bfu?bgu@bhu[biu]bju^bkuGdlu_bmu`bnu{bou|bpuHdqu}bru~bsuIdtuJduuKdvuLdwuacxuMdyubczuNdAuOdBuccCudcDuPdEuQdFuRdGuecHuSdIufcJugcKuhcLuicMujcNukcOulcPumcQuncRuTdSuocTupcUuqcVurcWuUdXuscYutcZuVd0uWd1u?doopoHc2u3uqo4uIc5uXd6uYd7uZd8uro9uso!uto#uuo$uvo%uwo'uxo(uyo)uzo*uAo+uBo,uCo-uDo.uEo/uFo:uGo;uHo=u?u@uIo[u]u^uJo_u`u{uKo|u}u~uLoavbvcvModvevfvNogvhvivOojvkvlvPomvQonvRoovSopvToqvUorvVosvWotvXouvYovvZowv0oxv1oyv2ozv3oAv4oBvXnCvYnDvZnEv0nFv1nGv2nHv5oIv@d[d0d6oJvJcKv7o8ozcAcLvMvY NvBcOvPv} QvRlSlRvTlSvUlwcTvVl~ WlUvXlVvYlfeZlWv0lge1lXv2lYv3lZv4l0v5l1v6l7l8l9l!lbbcb2v#l3v$l4v%l5v'l6v(l7v)l8v*l9v+l!v,l#v-l$v.l%v/l'v:l(v;l)v=l*v?l+v@l,v[l-v]l.v^l/v_l:v`l;v{l=v|l}l~l?vam@vbm[vcm]vdm^vem_vfm`vgm{vhm|vim}vjm~vkmawlmbwmmcwnmdwomewpmfwqmgwrmhwsmiwtmjwumvmkwwmlwxmmwymnwzmowAmpwBmqwCmrwDmswEmtwFmuwGmvwHmwwImxwJmywKmzwLmMm!dAw1dBw2dCwKcDwfdEwNmFwGw9oOmPm!oQmRm#oSmHwIw$oTmUm%ohe,cieVmWm]dXmYmJw'oZm^d0m(oMe)o1m5d2m-c.c3mNe*o+o4m6d5m6mKwLw,o7m8m-oje/cke9m!m_d#m$mMw.o%m`d'm/oOe:o(m7d)m:c;c*mPe;o=o+m8d,m-mNwOw?o.m/m:mLc;m=mMc@o?mPwQw[o@m[m]o]mRwSw^o^m_m_o`mTwUw`o{m|m{o}mVwWw|o~man}obnXwYw~ocndnapenZw0wbpfngnhncpin1w2wdpjnknlnepdb3web4wfpfb5wgb6w7w8w9w!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~waxbxcxdxexfxmngdgxhxixjxkxlxmxnnnxoxpxqxrxsxtxQeReonpnqnabuxrngphpsnvxwxxxhbyxzxAxBxibCxDxExipFxjpGxkpHxlpIxmpJxnpKxopLxppMxqpNxrpOxspPxtpQxupRxvpSxwpTxxpUxtnVxypWxzpXxApYxBpZxCp0xDp1x2x3x4x5x6x7x8x9x!x#x$x%x'x(x)x*x+x,x-x.x/x:x;x=x?xunhqiqjqvnwnkqlqxnynmqnqznoqAnBnpqCnDnEnFnGnHnInJnKnqqLnMnNnrqOnPnQnRnSnTnsqtq@x[x]xUn^x_x`x{x|x}x~xaybycydyeyfygyhyiyjykylymynyoypyqyrysytyuyvywyxyyyzyAyByCyDyEy

70  

71  

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

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

74 if err == cynvrtc.nvrtcResult.NVRTC_SUCCESS: 307 ctx2k l m n o p q c r d e f s g h t u v w x y z A B C D E F G H I J K L M N O P Q R S T U V W i j uchdX CcvcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c9c!c#c$c%c'c(c)c*c+cb 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | Ol} RlSlTlUlwcVl~ WlXlYlZl0l1l2l3l4l5l6l7l8l9l!lbbcb#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~lambmcmdmemfmgmhmimjmkmlmmmnmompmqmrmsmtmumvmwmxmymzmAmBmCmDmEmFmGmHmImJmKmLmMmOmPmRmTmUmheieWmXmZm0m1m2m-c.c3m4m5m7m8mjeke!m#m%m'm(m)m:c;c*m+m,m.m/m:m;m@m[m^m_m{m|m~mancndnfngnhnjnknlndbebfbgbybzblbmbnbAbBbCbDbjbkbobEbpbFbGbHbIbqbJbrbsbtbubKbLbMbNbObPbQbRbSbvbwbxbonpnqnabrnsnhbib

75 return 0 307 ctx2k l m n o p q c r d e f s g h t u v w x y z A B C D E F G H I J K L M N O P Q R S T U V W i j uchdX CcvcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c9c!c#c$c%c'c(c)c*c+cb 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | Ol} RlSlTlUlwcVl~ WlXlYlZl0l1l2l3l4l5l6l7l8l9l!lbbcb#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~lambmcmdmemfmgmhmimjmkmlmmmnmompmqmrmsmtmumvmwmxmymzmAmBmCmDmEmFmGmHmImJmKmLmMmOmPmRmTmUmheieWmXmZm0m1m2m-c.c3m4m5m7m8mjeke!m#m%m'm(m)m:c;c*m+m,m.m/m:m;m@m[m^m_m{m|m~mancndnfngnhnjnknlndbebfbgbybzblbmbnbAbBbCbDbjbkbobEbpbFbGbHbIbqbJbrbsbtbubKbLbMbNbObPbQbRbSbvbwbxbonpnqnabrnsnhbib

76 with gil: 

77 _raise_nvrtc_error(prog, err) 

78  

79  

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

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

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

83 cdef size_t logsize = 0 

84 if prog != NULL: 

85 cynvrtc.nvrtcGetProgramLogSize(prog, &logsize) 

86 cdef bytes log_bytes 

87 cdef str log_str = "" 

88 if logsize > 1 and prog != NULL: 

89 log_bytes = b" " * logsize 

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

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

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

93 if log_str: 

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

95 raise NVRTCError(err_msg) 

96  

97  

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

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

100 if err == cynvvm.nvvmResult.NVVM_SUCCESS: 21 ctx2=mEpFpmngdGpHpIpJpKpLpMpnnNpOpPpQpRpSpQeRe

101 return 0 21 ctx2=mEpFpmngdGpHpIpJpKpLpMpnnNpOpPpQpRpSpQeRe

102 with gil: 

103 _raise_nvvm_error(prog, err) 

104  

105  

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

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

108 cdef size_t logsize = 0 

109 if prog != NULL: 

110 cynvvm.nvvmGetProgramLogSize(prog, &logsize) 

111 cdef bytes log_bytes 

112 cdef str log_str = "" 

113 if logsize > 1 and prog != NULL: 

114 log_bytes = b" " * logsize 

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

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

117 cdef object exc = nvvmError(err) 

118 if log_str: 

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

120 raise exc 

121  

122  

123cdef int HANDLE_RETURN_NVJITLINK( (empty)

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

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

126 if err == cynvjitlink.nvJitLinkResult.NVJITLINK_SUCCESS: 40 ctx2b TpUpVpWpXpYpZp0p1p2p3n4n3p4p5p6p7p8p9p!p#p$p%p'p(p9d)p*p+p,p-p.p/p:p;p=p?p@p[p

127 return 0 39 ctx2b TpUpVpWpXpYpZp0p1p2p3n4n3p4p5p6p7p8p9p!p#p$p%p'p(p)p*p+p,p-p.p/p:p;p=p?p@p[p

128 with gil: 2 ctx2b 9d

129 _raise_nvjitlink_error(handle, err) 2 ctx2b 9d

130  

131  

132cdef int _raise_nvjitlink_error( (empty)

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

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

135 cdef size_t logsize = 0 2 ctx2b 9d

136 if handle != NULL: 2 ctx2b 9d

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

138 cdef bytes log_bytes 

139 cdef str log_str = "" 2 ctx2b 9d

140 if logsize > 1 and handle != NULL: 2 ctx2b 9d

141 log_bytes = b" " * logsize 1 ctx1b

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

143 cynvjitlink.nvJitLinkResult.NVJITLINK_SUCCESS: 1 ctx1b

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

145 cdef object exc = nvJitLinkError(err) 2 ctx2b 9d

146 if log_str: 2 ctx2b 9d

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

148 raise exc 2 ctx2b 9d

149  

150  

151cdef object _RUNTIME_SUCCESS = runtime.cudaError_t.cudaSuccess (empty)

152cdef object _NVRTC_SUCCESS = nvrtc.nvrtcResult.NVRTC_SUCCESS (empty)

153  

154  

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

156 if error == cydriver.CUresult.CUDA_SUCCESS: 1001 ctx2a SeTeUeVeWeXeYeZe0e1e2e3e4e5e6e7e8e9e!e#e$e%e'e(e)e*ek +el ,em -e=c.en /e?c:eo ;ep =eq ?eDcEcFc@c[c]cc r d e f s g h t u v ^c_c`c{c|c}cw @ex [ey ]ez ^eA _eB `eC {eD |eE }eF ~eG afH bfI cfJ dfK efL ffM gfN hfO ifP jfQ kfR lfS mfT nfU ofV pfW qfi j id~cadbdcdxcycddedZ 0 1 2 3 4 rfsftfufvfwfxfueyfzfGcAfBfCfDfEfFfGfHfIfJfKfLfaeMfNfOfbePfceucQfRfSfTfUfVfWfXfYfZf0f1f2fseteX TbUb3f4f5f6fNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c9c!c#c$c%c'c(c)c*c+cb 5 7f6 8f7 9f8 !f9 #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*dVbWb+d,d-d.dPlQl/d:d;d=ddeeekdldXbmdYbndodZb0bpdqdrd1bsd2b3b4b5b6b7b8b9b!btd#b$b%b'bud(b)bvdwdxdyd*bzd+bAdBd,b-bCdDdEd.bFd/b:b;b=b?b@b[b]b^bGd_b`b{b|bHd}b~bIdJdKdLdacMdbcNdOdccdcPdQdRdecSdfcgchcicjckclcmcncTdocpcqcrcUdsctcVdWd?ddgegfggghgigjgkglgmgngogpgqgrgsgtgugvgwgxgygzgAgBgCgDgEgFgGgHgIgJgKgLgMgNgOgPgQgRgSgTgUgVgWgXgYgZg0g1g2g3g4g5g6g7g8g9g!g#g@d[d0d$g%gzcAcY Bc} 'g(g)g*g+gwc,g-g~ .g/gfe:g;gge=g?g@g[g]g^g_g`g{g|g}g3dbb~g4dcbahbhchdhehfhghhhihjhkhlhmhnhohphqhrhshthuhvhwhxhyhzhAhBhChDhEhFhGhHhIhJhKhLhMhNhOhPhQhRhShThUhVhWhXhYhZh0h1h2h3h4h5h6h7h8h9h!h#h$h%hNm'h(h)h*h+h,h-hQm.h/h:hSm;h=h?h@h[h]h^hhe_h,c`hie{hVm|h}h~haiYmbicidieifigihiMeiijikiliminioipiNeqirisitiuivi6mwixiyiziAiBiCijeDi/cEikeFi9mGiHiIiJi$mKiLiMiNiOiPiQiOeRiSiTiUiViWiXiYiPeZi0i1i2i3i4i-m5i6i7i8i9i!i#iLc$i%i'iMc(i)i?m*i+i,i-i.i/i:i]m;i=i?i@i[i]i^i`m_i`i{i|i}i~iaj}mbjcjdjejfjgjhjbnijjjkjljmjnjojenpjqjrjsjtjujvjwjinxjyjzjAjBjCjDjEjdbFjebGjfbHjgbIjybJjzbKjlbLjmbMjnbNjAbOjBbPjCbQjDbRjjbSjkbTjobUjEbVjpbWjFbXjGbYjHbZjIb0jqb1jJb2jrb3jsb4jtb5jub6jKb7jLb8jMb9jNb!jOb#jPb$jQb%jRb'jSb(jvb)jwb*jxb+jgd,j-j.j/j:j;j=j?j@j[j]j^j_j`jab{j|j}j~jhbakbkckdkekibfkgkhkikjkkklkmknkokpkqkrksktkukvkwkxkykzkAkBkCkDkEkFkGkHkIkJkKkLkMkNkOkPkQkRkSkTkUkVkWkXkYkZk0k1k2k3k4k5k6k7k8k9k!k#k$k%k'k(k)k*k+k,k-k.k/k:k;k=k?k@kunhqiqjqvnwnkqlqxnynmqnqznoqAnBnpqCnDnEnFnGnHnInJnKnqqLnMnNnrqOnPnQnRnSnTnsqtq[k]k^k_k`kUn{k|k}k~kalblcldlelflglhliljlklllmlnlolplqlrlsltlulvlwlxlylzlAlBl

157 return 0 998 ctx2a SeTeUeVeWeXeYeZe0e1e2e3e4e5e6e7e8e9e!e#e$e%e'e(e)e*ek +el ,em -e=c.en /e?c:eo ;ep =eq ?eDcEcFc@c[c]cc r d e f s g h t u v ^c_c`c{c|c}cw @ex [ey ]ez ^eA _eB `eC {eD |eE }eF ~eG afH bfI cfJ dfK efL ffM gfN hfO ifP jfQ kfR lfS mfT nfU ofV pfW qfi j id~cadbdcdxcycddedZ 0 1 2 3 4 rfsftfufvfwfxfueyfzfGcAfBfCfDfEfFfGfHfIfJfKfLfaeMfNfOfbePfceucQfRfSfTfUfVfWfXfYfZf0f1f2fX TbUb3f4f5f6fNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c9c!c#c$c%c'c(c)c*c+cb 5 7f6 8f7 9f8 !f9 #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*dVbWb+d,d-d.dPlQl/d:d;d=ddeeekdldXbmdYbndodZb0bpdqdrd1bsd2b3b4b5b6b7b8b9b!btd#b$b%b'bud(b)bvdwdxdyd*bzd+bAdBd,b-bCdDdEd.bFd/b:b;b=b?b@b[b]b^bGd_b`b{b|bHd}b~bIdJdKdLdacMdbcNdOdccdcPdQdRdecSdfcgchcicjckclcmcncTdocpcqcrcUdsctcVdWd?ddgegfggghgigjgkglgmgngogpgqgrgsgtgugvgwgxgygzgAgBgCgDgEgFgGgHgIgJgKgLgMgNgOgPgQgRgSgTgUgVgWgXgYgZg0g1g2g3g4g5g6g7g8g9g!g#g@d[d0d$g%gzcAcY Bc} 'g(g)g*g+gwc,g-g~ .g/gfe:g;gge=g?g@g[g]g^g_g`g{g|g}g3dbb~g4dcbahbhchdhehfhghhhihjhkhlhmhnhohphqhrhshthuhvhwhxhyhzhAhBhChDhEhFhGhHhIhJhKhLhMhNhOhPhQhRhShThUhVhWhXhYhZh0h1h2h3h4h5h6h7h8h9h!h#h$h%hNm'h(h)h*h+h,h-hQm.h/h:hSm;h=h?h@h[h]h^hhe_h,c`hie{hVm|h}h~haiYmbicidieifigihiMeiijikiliminioipiNeqirisitiuivi6mwixiyiziAiBiCijeDi/cEikeFi9mGiHiIiJi$mKiLiMiNiOiPiQiOeRiSiTiUiViWiXiYiPeZi0i1i2i3i4i-m5i6i7i8i9i!i#iLc$i%i'iMc(i)i?m*i+i,i-i.i/i:i]m;i=i?i@i[i]i^i`m_i`i{i|i}i~iaj}mbjcjdjejfjgjhjbnijjjkjljmjnjojenpjqjrjsjtjujvjwjinxjyjzjAjBjCjDjEjdbFjebGjfbHjgbIjybJjzbKjlbLjmbMjnbNjAbOjBbPjCbQjDbRjjbSjkbTjobUjEbVjpbWjFbXjGbYjHbZjIb0jqb1jJb2jrb3jsb4jtb5jub6jKb7jLb8jMb9jNb!jOb#jPb$jQb%jRb'jSb(jvb)jwb*jxb+j,j-j.j/j:j;j=j?j@j[j]j^j_j`jab{j|j}j~jhbakbkckdkekibfkgkhkikjkkklkmknkokpkqkrksktkukvkwkxkykzkAkBkCkDkEkFkGkHkIkJkKkLkMkNkOkPkQkRkSkTkUkVkWkXkYkZk0k1k2k3k4k5k6k7k8k9k!k#k$k%k'k(k)k*k+k,k-k.k/k:k;k=k?k@kunhqiqjqvnwnkqlqxnynmqnqznoqAnBnpqCnDnEnFnGnHnInJnKnqqLnMnNnrqOnPnQnRnSnTnsqtq[k]k^k_k`kUn{k|k}k~kalblcldlelflglhliljlklllmlnlolplqlrlsltlulvlwlxlylzlAlBl

158 cdef const char* name 

159 name_err = cydriver.cuGetErrorName(error, &name) 6 ctx2ueseteY ~ gd

160 if name_err != cydriver.CUresult.CUDA_SUCCESS: 6 ctx2ueseteY ~ gd

161 raise CUDAError(f"UNEXPECTED ERROR CODE: {error}") 1 ctx2ue

162 with gil: 6 ctx2ueseteY ~ gd

163 # TODO: consider lower this to Cython 

164 expl = DRIVER_CU_RESULT_EXPLANATIONS.get(int(error)) 6 ctx2ueseteY ~ gd

165 if expl is not None: 6 ctx2ueseteY ~ gd

166 raise CUDAError(f"{name.decode()}: {expl}") 6 ctx2ueseteY ~ gd

167 cdef const char* desc 

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

169 if desc_err != cydriver.CUresult.CUDA_SUCCESS: 

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

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

172  

173  

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

175 if error == _RUNTIME_SUCCESS: 86 ctx2Z 0 1 2 3 4 $dveweCl%djdGc'd(d)dTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!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~bacbcccdcecfcgchcicjckclcmcncocpcqcrcsctcHcIcXdYdZdJc

176 return 0 86 ctx2Z 0 1 2 3 4 $dveweCl%djdGc'd(d)dTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!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~bacbcccdcecfcgchcicjckclcmcncocpcqcrcsctcHcIcXdYdZdJc

177 name_err, name = runtime.cudaGetErrorName(error) 1 ctx2Cl

178 if name_err != _RUNTIME_SUCCESS: 1 ctx2Cl

179 raise CUDAError(f"UNEXPECTED ERROR CODE: {error}") 

180 name = name.decode() 1 ctx2Cl

181 expl = RUNTIME_CUDA_ERROR_EXPLANATIONS.get(int(error)) 1 ctx2Cl

182 if expl is not None: 1 ctx2Cl

183 raise CUDAError(f"{name}: {expl}") 1 ctx2Cl

184 desc_err, desc = runtime.cudaGetErrorString(error) 

185 if desc_err != _RUNTIME_SUCCESS: 

186 raise CUDAError(f"{name}") 

187 desc = desc.decode() 

188 raise CUDAError(f"{name}: {desc}") 

189  

190  

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

192 if error == _NVRTC_SUCCESS: 82 ctx2a b 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } 3dbb4dcbLcMcdbebfbgbybzblbmbnbAbBbCbDbjbkbobEbpbFbGbHbIbqbJbrbsbtbubKbLbMbNbObPbQbRbSbvbwbxbabhbib

193 return 0 82 ctx2a b 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } 3dbb4dcbLcMcdbebfbgbybzblbmbnbAbBbCbDbjbkbobEbpbFbGbHbIbqbJbrbsbtbubKbLbMbNbObPbQbRbSbvbwbxbabhbib

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

195 if handle is not None: 

196 _, logsize = nvrtc.nvrtcGetProgramLogSize(handle) 

197 log = b" " * logsize 

198 _ = nvrtc.nvrtcGetProgramLog(handle, log) 

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

200 raise NVRTCError(err) 

201  

202  

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

204 if isinstance(error, driver.CUresult): 905 ctx2a SeTeUeVeWeXeYeZe0e1e2e3e4e5e6e7e8e9e!e#e$e%e'e(e)e*ek +el ,em -e=c.en /e?c:eo ;ep =eq ?eDcEcFc@c[c]cc r d e f s g h t u v ^c_c`c{c|c}cw @ex [ey ]ez ^eA _eB `eC {eD |eE }eF ~eG afH bfI cfJ dfK efL ffM gfN hfO ifP jfQ kfR lfS mfT nfU ofV pfW qfi j id~cadbdcdxcycddedZ 0 1 2 3 4 $dverfwesftfufvfwfxfyf%djdzfGcAfBfCfDfEfFfGfHf'dIfJfKfLfaeMfNfOfbe(d)dPfceQfRfSfTfUfVfWfXfYfZf0f1f2fX TbUb3f4f5f6fb 5 7f6 8f7 9f8 !f9 #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*dVbWb+d,d-d.d/d:d;d=ddeeekdldXbmdYbndodZb0bpdqdrd1bsd2b3b4b5b6b7b8b9b!btd#b$b%b'bud(b)bvdwdxdyd*bzd+bAdBd,b-bCdDdEd.bFd/b:b;b=b?b@b[b]b^bGd_b`b{b|bHd}b~bIdJdKdLdacMdbcNdOdccdcPdQdRdecSdfcgchcicjckclcmcncTdocpcqcrcUdsctcVdWd?dHcdgegfgIcggXdhgYdigZdjgkglgmgngogpgqgrgsgtgugvgwgxgygzgAgBgCgDgEgFgGgHgIgJgKgLgMgNgOgPgQgRgSgTgUgVgWgXgYgZg0g1g2g3g4g5g6g7g8g9g!g#g@d[d0d$gJc%g} 'g(g)g*g+g,g-g~ .g/gfe:g;gge=g?g@g[g]g^g_g`g{g|g}g3dbb~g4dcbahbhchdhehfhghhhihjhkhlhmhnhohphqhrhshthuhvhwhxhyhzhAhBhChDhEhFhGhHhIhJhKhLhMhNhOhPhQhRhShThUhVhWhXhYhZh0h1h2h3h4h5h6h7h8h9h!h#h$h%h'h(h)h*h+h,h-h.h/h:h;h=h?h@h[h]h^h_h`h{h|h}h~haibicidieifigihiiijikiliminioipiqirisitiuiviwixiyiziAiBiCiDiEiFiGiHiIiJiKiLiMiNiOiPiQiRiSiTiUiViWiXiYiZi0i1i2i3i4i5i6i7i8i9i!i#iLc$i%i'iMc(i)i*i+i,i-i.i/i:i;i=i?i@i[i]i^i_i`i{i|i}i~iajbjcjdjejfjgjhjijjjkjljmjnjojpjqjrjsjtjujvjwjxjyjzjAjBjCjDjEjdbFjebGjfbHjgbIjybJjzbKjlbLjmbMjnbNjAbOjBbPjCbQjDbRjjbSjkbTjobUjEbVjpbWjFbXjGbYjHbZjIb0jqb1jJb2jrb3jsb4jtb5jub6jKb7jLb8jMb9jNb!jOb#jPb$jQb%jRb'jSb(jvb)jwb*jxb+j,j-j.j/j:j;j=j?j@j[j]j^j_j`jab{j|j}j~jhbakbkckdkekibfkgkhkikjkkklkmknkokpkqkrksktkukvkwkxkykzkAkBkCkDkEkFkGkHkIkJkKkLkMkNkOkPkQkRkSkTkUkVkWkXkYkZk0k1k2k3k4k5k6k7k8k9k!k#k$k%k'k(k)k*k+k,k-k.k/k:k;k=k?k@k[k]k^k_k`k{k|k}k~kalblcldlelflglhliljlklllmlnlolplqlrlsltlulvlwlxlylzlAlBl

205 return _check_driver_error(error) 890 ctx2a SeTeUeVeWeXeYeZe0e1e2e3e4e5e6e7e8e9e!e#e$e%e'e(e)e*ek +el ,em -e=c.en /e?c:eo ;ep =eq ?eDcEcFc@c[c]cc r d e f s g h t u v ^c_c`c{c|c}cw @ex [ey ]ez ^eA _eB `eC {eD |eE }eF ~eG afH bfI cfJ dfK efL ffM gfN hfO ifP jfQ kfR lfS mfT nfU ofV pfW qfi j id~cadbdcdxcycddedZ 0 1 2 3 4 rfsftfufvfwfxfyfzfAfBfCfDfEfFfGfHfIfJfKfLfaeMfNfOfbePfceQfRfSfTfUfVfWfXfYfZf0f1f2fX TbUb3f4f5f6fb 5 7f6 8f7 9f8 !f9 #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*dVbWb+d,d-d.d/d:d;d=ddeeekdldXbmdYbndodZb0bpdqdrd1bsd2b3b4b5b6b7b8b9b!btd#b$b%b'bud(b)bvdwdxdyd*bzd+bAdBd,b-bCdDdEd.bFd/b:b;b=b?b@b[b]b^bGd_b`b{b|bHd}b~bIdJdKdLdacMdbcNdOdccdcPdQdRdecSdfcgchcicjckclcmcncTdocpcqcrcUdsctcVdWd?ddgegfggghgigjgkglgmgngogpgqgrgsgtgugvgwgxgygzgAgBgCgDgEgFgGgHgIgJgKgLgMgNgOgPgQgRgSgTgUgVgWgXgYgZg0g1g2g3g4g5g6g7g8g9g!g#g@d[d0d$g%g} 'g(g)g*g+g,g-g~ .g/gfe:g;gge=g?g@g[g]g^g_g`g{g|g}g3dbb~g4dcbahbhchdhehfhghhhihjhkhlhmhnhohphqhrhshthuhvhwhxhyhzhAhBhChDhEhFhGhHhIhJhKhLhMhNhOhPhQhRhShThUhVhWhXhYhZh0h1h2h3h4h5h6h7h8h9h!h#h$h%h'h(h)h*h+h,h-h.h/h:h;h=h?h@h[h]h^h_h`h{h|h}h~haibicidieifigihiiijikiliminioipiqirisitiuiviwixiyiziAiBiCiDiEiFiGiHiIiJiKiLiMiNiOiPiQiRiSiTiUiViWiXiYiZi0i1i2i3i4i5i6i7i8i9i!i#iLc$i%i'iMc(i)i*i+i,i-i.i/i:i;i=i?i@i[i]i^i_i`i{i|i}i~iajbjcjdjejfjgjhjijjjkjljmjnjojpjqjrjsjtjujvjwjxjyjzjAjBjCjDjEjdbFjebGjfbHjgbIjybJjzbKjlbLjmbMjnbNjAbOjBbPjCbQjDbRjjbSjkbTjobUjEbVjpbWjFbXjGbYjHbZjIb0jqb1jJb2jrb3jsb4jtb5jub6jKb7jLb8jMb9jNb!jOb#jPb$jQb%jRb'jSb(jvb)jwb*jxb+j,j-j.j/j:j;j=j?j@j[j]j^j_j`jab{j|j}j~jhbakbkckdkekibfkgkhkikjkkklkmknkokpkqkrksktkukvkwkxkykzkAkBkCkDkEkFkGkHkIkJkKkLkMkNkOkPkQkRkSkTkUkVkWkXkYkZk0k1k2k3k4k5k6k7k8k9k!k#k$k%k'k(k)k*k+k,k-k.k/k:k;k=k?k@k[k]k^k_k`k{k|k}k~kalblcldlelflglhliljlklllmlnlolplqlrlsltlulvlwlxlylzlAlBl

206 elif isinstance(error, runtime.cudaError_t): 167 ctx2a Z 0 1 2 3 4 $dvewe%djdGc'd(d)dTbUbb 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | VbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!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~bacbcccdcecfcgchcicjckclcmcncocpcqcrcsctcHcIcXdYdZdJc} 3dbb4dcbLcMcdbebfbgbybzblbmbnbAbBbCbDbjbkbobEbpbFbGbHbIbqbJbrbsbtbubKbLbMbNbObPbQbRbSbvbwbxbabhbib

207 return _check_runtime_error(error) 85 ctx2Z 0 1 2 3 4 $dvewe%djdGc'd(d)dTbUbVbWbXbYbZb0b1b2b3b4b5b6b7b8b9b!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~bacbcccdcecfcgchcicjckclcmcncocpcqcrcsctcHcIcXdYdZdJc

208 elif isinstance(error, nvrtc.nvrtcResult): 82 ctx2a b 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } 3dbb4dcbLcMcdbebfbgbybzblbmbnbAbBbCbDbjbkbobEbpbFbGbHbIbqbJbrbsbtbubKbLbMbNbObPbQbRbSbvbwbxbabhbib

209 return _check_nvrtc_error(error, handle=handle) 82 ctx2a b 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } 3dbb4dcbLcMcdbebfbgbybzblbmbnbAbBbCbDbjbkbobEbpbFbGbHbIbqbJbrbsbtbubKbLbMbNbObPbQbRbSbvbwbxbabhbib

210 else: 

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

212  

213  

214def handle_return(tuple result, handle=None): (empty)

215 _check_error(result[0], handle=handle) 905 ctx2a SeTeUeVeWeXeYeZe0e1e2e3e4e5e6e7e8e9e!e#e$e%e'e(e)e*ek +el ,em -e=c.en /e?c:eo ;ep =eq ?eDcEcFc@c[c]cc r d e f s g h t u v ^c_c`c{c|c}cw @ex [ey ]ez ^eA _eB `eC {eD |eE }eF ~eG afH bfI cfJ dfK efL ffM gfN hfO ifP jfQ kfR lfS mfT nfU ofV pfW qfi j id~cadbdcdxcycddedZ 0 1 2 3 4 $dverfwesftfufvfwfxfyf%djdzfGcAfBfCfDfEfFfGfHf'dIfJfKfLfaeMfNfOfbe(d)dPfceQfRfSfTfUfVfWfXfYfZf0f1f2fX TbUb3f4f5f6fb 5 7f6 8f7 9f8 !f9 #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*dVbWb+d,d-d.d/d:d;d=ddeeekdldXbmdYbndodZb0bpdqdrd1bsd2b3b4b5b6b7b8b9b!btd#b$b%b'bud(b)bvdwdxdyd*bzd+bAdBd,b-bCdDdEd.bFd/b:b;b=b?b@b[b]b^bGd_b`b{b|bHd}b~bIdJdKdLdacMdbcNdOdccdcPdQdRdecSdfcgchcicjckclcmcncTdocpcqcrcUdsctcVdWd?dHcdgegfgIcggXdhgYdigZdjgkglgmgngogpgqgrgsgtgugvgwgxgygzgAgBgCgDgEgFgGgHgIgJgKgLgMgNgOgPgQgRgSgTgUgVgWgXgYgZg0g1g2g3g4g5g6g7g8g9g!g#g@d[d0d$gJc%g} 'g(g)g*g+g,g-g~ .g/gfe:g;gge=g?g@g[g]g^g_g`g{g|g}g3dbb~g4dcbahbhchdhehfhghhhihjhkhlhmhnhohphqhrhshthuhvhwhxhyhzhAhBhChDhEhFhGhHhIhJhKhLhMhNhOhPhQhRhShThUhVhWhXhYhZh0h1h2h3h4h5h6h7h8h9h!h#h$h%h'h(h)h*h+h,h-h.h/h:h;h=h?h@h[h]h^h_h`h{h|h}h~haibicidieifigihiiijikiliminioipiqirisitiuiviwixiyiziAiBiCiDiEiFiGiHiIiJiKiLiMiNiOiPiQiRiSiTiUiViWiXiYiZi0i1i2i3i4i5i6i7i8i9i!i#iLc$i%i'iMc(i)i*i+i,i-i.i/i:i;i=i?i@i[i]i^i_i`i{i|i}i~iajbjcjdjejfjgjhjijjjkjljmjnjojpjqjrjsjtjujvjwjxjyjzjAjBjCjDjEjdbFjebGjfbHjgbIjybJjzbKjlbLjmbMjnbNjAbOjBbPjCbQjDbRjjbSjkbTjobUjEbVjpbWjFbXjGbYjHbZjIb0jqb1jJb2jrb3jsb4jtb5jub6jKb7jLb8jMb9jNb!jOb#jPb$jQb%jRb'jSb(jvb)jwb*jxb+j,j-j.j/j:j;j=j?j@j[j]j^j_j`jab{j|j}j~jhbakbkckdkekibfkgkhkikjkkklkmknkokpkqkrksktkukvkwkxkykzkAkBkCkDkEkFkGkHkIkJkKkLkMkNkOkPkQkRkSkTkUkVkWkXkYkZk0k1k2k3k4k5k6k7k8k9k!k#k$k%k'k(k)k*k+k,k-k.k/k:k;k=k?k@k[k]k^k_k`k{k|k}k~kalblcldlelflglhliljlklllmlnlolplqlrlsltlulvlwlxlylzlAlBl

216 cdef int out_len = len(result) 905 ctx2a SeTeUeVeWeXeYeZe0e1e2e3e4e5e6e7e8e9e!e#e$e%e'e(e)e*ek +el ,em -e=c.en /e?c:eo ;ep =eq ?eDcEcFc@c[c]cc r d e f s g h t u v ^c_c`c{c|c}cw @ex [ey ]ez ^eA _eB `eC {eD |eE }eF ~eG afH bfI cfJ dfK efL ffM gfN hfO ifP jfQ kfR lfS mfT nfU ofV pfW qfi j id~cadbdcdxcycddedZ 0 1 2 3 4 $dverfwesftfufvfwfxfyf%djdzfGcAfBfCfDfEfFfGfHf'dIfJfKfLfaeMfNfOfbe(d)dPfceQfRfSfTfUfVfWfXfYfZf0f1f2fX TbUb3f4f5f6fb 5 7f6 8f7 9f8 !f9 #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*dVbWb+d,d-d.d/d:d;d=ddeeekdldXbmdYbndodZb0bpdqdrd1bsd2b3b4b5b6b7b8b9b!btd#b$b%b'bud(b)bvdwdxdyd*bzd+bAdBd,b-bCdDdEd.bFd/b:b;b=b?b@b[b]b^bGd_b`b{b|bHd}b~bIdJdKdLdacMdbcNdOdccdcPdQdRdecSdfcgchcicjckclcmcncTdocpcqcrcUdsctcVdWd?dHcdgegfgIcggXdhgYdigZdjgkglgmgngogpgqgrgsgtgugvgwgxgygzgAgBgCgDgEgFgGgHgIgJgKgLgMgNgOgPgQgRgSgTgUgVgWgXgYgZg0g1g2g3g4g5g6g7g8g9g!g#g@d[d0d$gJc%g} 'g(g)g*g+g,g-g~ .g/gfe:g;gge=g?g@g[g]g^g_g`g{g|g}g3dbb~g4dcbahbhchdhehfhghhhihjhkhlhmhnhohphqhrhshthuhvhwhxhyhzhAhBhChDhEhFhGhHhIhJhKhLhMhNhOhPhQhRhShThUhVhWhXhYhZh0h1h2h3h4h5h6h7h8h9h!h#h$h%h'h(h)h*h+h,h-h.h/h:h;h=h?h@h[h]h^h_h`h{h|h}h~haibicidieifigihiiijikiliminioipiqirisitiuiviwixiyiziAiBiCiDiEiFiGiHiIiJiKiLiMiNiOiPiQiRiSiTiUiViWiXiYiZi0i1i2i3i4i5i6i7i8i9i!i#iLc$i%i'iMc(i)i*i+i,i-i.i/i:i;i=i?i@i[i]i^i_i`i{i|i}i~iajbjcjdjejfjgjhjijjjkjljmjnjojpjqjrjsjtjujvjwjxjyjzjAjBjCjDjEjdbFjebGjfbHjgbIjybJjzbKjlbLjmbMjnbNjAbOjBbPjCbQjDbRjjbSjkbTjobUjEbVjpbWjFbXjGbYjHbZjIb0jqb1jJb2jrb3jsb4jtb5jub6jKb7jLb8jMb9jNb!jOb#jPb$jQb%jRb'jSb(jvb)jwb*jxb+j,j-j.j/j:j;j=j?j@j[j]j^j_j`jab{j|j}j~jhbakbkckdkekibfkgkhkikjkkklkmknkokpkqkrksktkukvkwkxkykzkAkBkCkDkEkFkGkHkIkJkKkLkMkNkOkPkQkRkSkTkUkVkWkXkYkZk0k1k2k3k4k5k6k7k8k9k!k#k$k%k'k(k)k*k+k,k-k.k/k:k;k=k?k@k[k]k^k_k`k{k|k}k~kalblcldlelflglhliljlklllmlnlolplqlrlsltlulvlwlxlylzlAlBl

217 if out_len == 1: 905 ctx2a SeTeUeVeWeXeYeZe0e1e2e3e4e5e6e7e8e9e!e#e$e%e'e(e)e*ek +el ,em -e=c.en /e?c:eo ;ep =eq ?eDcEcFc@c[c]cc r d e f s g h t u v ^c_c`c{c|c}cw @ex [ey ]ez ^eA _eB `eC {eD |eE }eF ~eG afH bfI cfJ dfK efL ffM gfN hfO ifP jfQ kfR lfS mfT nfU ofV pfW qfi j id~cadbdcdxcycddedZ 0 1 2 3 4 $dverfwesftfufvfwfxfyf%djdzfGcAfBfCfDfEfFfGfHf'dIfJfKfLfaeMfNfOfbe(d)dPfceQfRfSfTfUfVfWfXfYfZf0f1f2fX TbUb3f4f5f6fb 5 7f6 8f7 9f8 !f9 #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*dVbWb+d,d-d.d/d:d;d=ddeeekdldXbmdYbndodZb0bpdqdrd1bsd2b3b4b5b6b7b8b9b!btd#b$b%b'bud(b)bvdwdxdyd*bzd+bAdBd,b-bCdDdEd.bFd/b:b;b=b?b@b[b]b^bGd_b`b{b|bHd}b~bIdJdKdLdacMdbcNdOdccdcPdQdRdecSdfcgchcicjckclcmcncTdocpcqcrcUdsctcVdWd?dHcdgegfgIcggXdhgYdigZdjgkglgmgngogpgqgrgsgtgugvgwgxgygzgAgBgCgDgEgFgGgHgIgJgKgLgMgNgOgPgQgRgSgTgUgVgWgXgYgZg0g1g2g3g4g5g6g7g8g9g!g#g@d[d0d$gJc%g} 'g(g)g*g+g,g-g~ .g/gfe:g;gge=g?g@g[g]g^g_g`g{g|g}g3dbb~g4dcbahbhchdhehfhghhhihjhkhlhmhnhohphqhrhshthuhvhwhxhyhzhAhBhChDhEhFhGhHhIhJhKhLhMhNhOhPhQhRhShThUhVhWhXhYhZh0h1h2h3h4h5h6h7h8h9h!h#h$h%h'h(h)h*h+h,h-h.h/h:h;h=h?h@h[h]h^h_h`h{h|h}h~haibicidieifigihiiijikiliminioipiqirisitiuiviwixiyiziAiBiCiDiEiFiGiHiIiJiKiLiMiNiOiPiQiRiSiTiUiViWiXiYiZi0i1i2i3i4i5i6i7i8i9i!i#iLc$i%i'iMc(i)i*i+i,i-i.i/i:i;i=i?i@i[i]i^i_i`i{i|i}i~iajbjcjdjejfjgjhjijjjkjljmjnjojpjqjrjsjtjujvjwjxjyjzjAjBjCjDjEjdbFjebGjfbHjgbIjybJjzbKjlbLjmbMjnbNjAbOjBbPjCbQjDbRjjbSjkbTjobUjEbVjpbWjFbXjGbYjHbZjIb0jqb1jJb2jrb3jsb4jtb5jub6jKb7jLb8jMb9jNb!jOb#jPb$jQb%jRb'jSb(jvb)jwb*jxb+j,j-j.j/j:j;j=j?j@j[j]j^j_j`jab{j|j}j~jhbakbkckdkekibfkgkhkikjkkklkmknkokpkqkrksktkukvkwkxkykzkAkBkCkDkEkFkGkHkIkJkKkLkMkNkOkPkQkRkSkTkUkVkWkXkYkZk0k1k2k3k4k5k6k7k8k9k!k#k$k%k'k(k)k*k+k,k-k.k/k:k;k=k?k@k[k]k^k_k`k{k|k}k~kalblcldlelflglhliljlklllmlnlolplqlrlsltlulvlwlxlylzlAlBl

218 return 196 ctx2Ue!ek l m =cn ?co p q DcEcFc@c[c]cc r d e f s g h t u v ^c_c`c{c|c}cw x y z A B C D E F G H I J K L M N O P Q R S T U V W idZ 0 1 2 3 4 yfjdGcX TbUb*dVbWb+d,d-d.d/d:d;d=ddeeekdldXbmdYbndodZb0bpdqdrd1bsd2b3b4b5b6b7b8b9b!btd#b$b%b'bud(b)bvdwdxdyd*bzd+bAdBd,b-bCdDdEd.bFd/b:b;b=b?b@b[b]b^bGd_b`b{b|bHd}b~bIdJdKdLdacMdbcNdOdccdcPdQdRdecSdfcgchcicjckclcmcncTdocpcqcrcUdsctcVdWd?dHcIcXdYdZd@d[d0dJc

219 elif out_len == 2: 

220 return result[1] 892 ctx2a SeTeVeWeXeYeZe0e1e2e3e4e5e6e7e8e9e#e$e%e'e(e)e*ek +el ,em -e=c.en /e?c:eo ;ep =eq ?eDcEcFc@c[c]cc r d e f s g h t u v ^c_c`c{c|c}cw @ex [ey ]ez ^eA _eB `eC {eD |eE }eF ~eG afH bfI cfJ dfK efL ffM gfN hfO ifP jfQ kfR lfS mfT nfU ofV pfW qfi j id~cadbdcdxcycddedZ 0 1 2 3 4 $dverfwesftfufvfwfxf%dzfAfBfCfDfEfFfGfHf'dIfJfKfLfaeMfNfOfbe(d)dPfceQfRfSfTfUfVfWfXfYfZf0f1f2fX TbUb3f4f5f6fb 5 7f6 8f7 9f8 !f9 #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*dVbWb+d,d-d.d/d:d;d=dkdldXbmdYbndodZb0bpdqdrd1bsd2b3b4b5b6b7b8b9b!btd#b$b%b'bud(b)bvdwdxdyd*bzd+bAdBd,b-bCdDdEd.bFd/b:b;b=b?b@b[b]b^bGd_b`b{b|bHd}b~bIdJdKdLdacMdbcNdOdccdcPdQdRdecSdfcgchcicjckclcmcncTdocpcqcrcUdsctcVdWd?ddgegfggghgigjgkglgmgngogpgqgrgsgtgugvgwgxgygzgAgBgCgDgEgFgGgHgIgJgKgLgMgNgOgPgQgRgSgTgUgVgWgXgYgZg0g1g2g3g4g5g6g7g8g9g!g#g@d[d0d$g%g} 'g(g)g*g+g,g-g~ .g/gfe:g;gge=g?g@g[g]g^g_g`g{g|g}g3dbb~g4dcbahbhchdhehfhghhhihjhkhlhmhnhohphqhrhshthuhvhwhxhyhzhAhBhChDhEhFhGhHhIhJhKhLhMhNhOhPhQhRhShThUhVhWhXhYhZh0h1h2h3h4h5h6h7h8h9h!h#h$h%h'h(h)h*h+h,h-h.h/h:h;h=h?h@h[h]h^h_h`h{h|h}h~haibicidieifigihiiijikiliminioipiqirisitiuiviwixiyiziAiBiCiDiEiFiGiHiIiJiKiLiMiNiOiPiQiRiSiTiUiViWiXiYiZi0i1i2i3i4i5i6i7i8i9i!i#iLc$i%i'iMc(i)i*i+i,i-i.i/i:i;i=i?i@i[i]i^i_i`i{i|i}i~iajbjcjdjejfjgjhjijjjkjljmjnjojpjqjrjsjtjujvjwjxjyjzjAjBjCjDjEjdbFjebGjfbHjgbIjybJjzbKjlbLjmbMjnbNjAbOjBbPjCbQjDbRjjbSjkbTjobUjEbVjpbWjFbXjGbYjHbZjIb0jqb1jJb2jrb3jsb4jtb5jub6jKb7jLb8jMb9jNb!jOb#jPb$jQb%jRb'jSb(jvb)jwb*jxb+j,j-j.j/j:j;j=j?j@j[j]j^j_j`jab{j|j}j~jhbakbkckdkekibfkgkhkikjkkklkmknkokpkqkrksktkukvkwkxkykzkAkBkCkDkEkFkGkHkIkJkKkLkMkNkOkPkQkRkSkTkUkVkWkXkYkZk0k1k2k3k4k5k6k7k8k9k!k#k$k%k'k(k)k*k+k,k-k.k/k:k;k=k?k@k[k]k^k_k`k{k|k}k~kalblcldlelflglhliljlklllmlnlolplqlrlsltlulvlwlxlylzlAlBl

221 else: 

222 return result[1:] 139 ctx2a k l m =cn ?co p q DcEcFc@c[c]cc d e f g h t u v ^c_c`c{c|c}cw x y z A B C D E F G H I J K L M N O P Q R S T U V W b 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } 3dbb4dcbLcMcdbebfbgbybzblbmbnbAbBbCbDbjbkbobEbpbFbGbHbIbqbJbrbsbtbubKbLbMbNbObPbQbRbSbvbwbxbabhbib

223  

224  

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

226 """ 

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

228 """ 

229 if options is None: 750 ctx2k l m =cn ?co p q DcEcFc@c[c]cc r d e f s g h t u v ^c_c`c{c|c}cw x y z A B C D E F G H I J K L M N O P Q R S T U V W 8n9nEl!nFl#n$n%nGl'nHl(ni )nj *n{d|dneoe}d~dpeqeKeLeIl+n,n-n.n/n:nid;n~c=nad?nbd@ncd[nxcycdd]ned^n_n`n{n|nZ 0 1 2 3 4 }n~naobocodojdGcuqeoJlucKlvqfowqgohoiojokolomoLlMlNlhdnoX TbUbCcvcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c9c!c#c$c%c'c(c)c*c+cb 5 Tp6 Up7 Vp8 Wp9 Xp! Yp# Zp$ 0p% 1p' 2p( 3n) 4n* 3p+ 4p, 5p- 6p. 7p/ 8p: 9p; !p= #p? $p@ %p[ 'p] (p^ 9d_ )p` *p{ +p| ,pOlVbWbPlQlkdldXbmdYbndodZb0bpdqdrd1bsd2b3b4b5b6b7b8b9b!btd#b$b%b'bud(b)bvdwdxdyd*bzd+bAdBd,b-bCdDdEd.bFd/b:b;b=b?b@b[b]b^bGd_b`b{b|bHd}b~bIdJdKdLdacMdbcNdOdccdcPdQdRdecSdfcgchcicjckclcmcncTdocpcqcrcUdsctcVdWdoopoHcqoIcXdYdZdrosotouovowoxoyozoAoBoCoDoEoFoGoHoIoJoKoLoMoNoOoPoQoRoSoToUoVoWoXoYoZo0o1o2o3o4oXnYnZn0n1n2n5o0d6oJc7o8ozcAcY Bc} RlSlTlUlwcVl~ WlXlYlZl0l1l2l3l4l5l6l7l8l9l!lbbcb#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~lambmcmdmemfmgmhmimjmkmlmmmnmompmqmrmsmtmumvmwmxmymzmAmBmCmDmEmFmGmHmImJmKmLmMm!d1d2dKcfdNm9oOmPm!oQmRm#oSm$oTmUm%ohe,cieVmWmXmYm'oZm0m(oMe)o1m5d2m-c.c3mNe*o+o4m6d5m6m,o7m8m-oje/cke9m!m#m$m.o%m'm/oOe:o(m7d)m:c;c*mPe;o=o+m8d,m-m?o.m/m:m;m=m@o?m[o@m[m]o]m^o^m_m_o`m`o{m|m{o}m|o~man}obn~ocndnapenbpfngnhncpindpjnknlnepEpFpdbebfpfbgbybzblbmbnbAbBbCbDbjbkbobEbpbFbGbHbIbqbJbrbsbtbubKbLbMbNbObPbQbRbSbvbwbxbmngdGpHpIpJpKpLpMpnnNpOpPpQpRpSpQeReonpnqnabrngphpsn-p.p/p:phb;p=p?p@p[pibipjpkplpmpnpopppqprpsptpupvpwpxptnypzpApBpCpDpunvnwnxnynznAnBnCnDnEnFnGnHnInJnKnLnMnNnOnPnQnRnSnTnUn

230 if keep_none: 482 ctx2k l m =cn ?co p q DcEcFc@c[c]cc r d e f s g h t u v ^c_c`c{c|c}cw x y z A B C D E F G H I J K L M N O P Q R S T U V W i j {d|dneoe}d~dpeqeKeLeIl~cadbdcdxcycddedZ 0 1 2 3 4 }n~naobocodojdGcuqeoJlucKlfogohoiojokolomoLlMlNlhdnoX TbUbvcVbWbPlQlkdldXbmdYbndodZb0bpdqdrd1bsd2b3b4b5b6b7b8b9b!btd#b$b%b'bud(b)bvdwdxdyd*bzd+bAdBd,b-bCdDdEd.bFd/b:b;b=b?b@b[b]b^bGd_b`b{b|bHd}b~bIdJdKdLdacMdbcNdOdccdcPdQdRdecSdfcgchcicjckclcmcncTdocpcqcrcUdsctcVdWdoopoHcqoIcXdZd0d6oJcRlSlTlUlwcVl~ WlXlYlZl0l1l2l3l4l5l6l7l8lbbcb#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~lambmcmdmemfmgmhmimjmkmlmmmnmompmqmrmsmtmumvmwmxmymzmAmBmCmDmEmFmGmHmImJmKmLmMmKcNm9oOmPm!oQmRm#oSm$oTmUm%ohe,cieVmWmXmYm'oZm0m(oMe)o1m5d2m-c.c3mNe*o+o4m6d5m6m,o7m8m-oje/cke9m!m#m$m.o%m'm/oOe:o(m7d)m:c;c*mPe;o=o+m8d,m-m?o.m/m:m;m=m@o?m[o@m[m]o]m^o^m_m_o`m`o{m|m{o}m|o~man}obn~ocndnapenbpfngnhncpindpjnknlnepebmngdonpnqnrngphpsnhbibkplpmpnpqprpsptpwptnypBpCpunvnwnxnynznAnBnCnDnEnFnGnHnInJnKnLnMnNnOnPnQnRnSnTnUn

231 return options 75 ctx2DcEcFcc r d e f s g h jdGcPlQloopoqoXdZd6owcNmQmSmhe,cieVmYmMeNe6mje/cke9m$mOePe-m?m]m`m}mbneninunvnwnxnynznAnBnCnDnEnFnGnHnInJnKnLnMnNnOnPnQnRnSnTnUn

232 return cls() 427 ctx2k l m =cn ?co p q DcEcFc@c[c]cc r d e f s g h t u v ^c_c`c{c|c}cw x y z A B C D E F G H I J K L M N O P Q R S T U V W i j {d|dneoe}d~dpeqeKeLeIl~cadbdcdxcycddedZ 0 1 2 3 4 }n~naobocodouqeoJlucKlfogohoiojokolomoLlMlNlhdnoX TbUbvcVbWbkdldXbmdYbndodZb0bpdqdrd1bsd2b3b4b5b6b7b8b9b!btd#b$b%b'bud(b)bvdwdxdyd*bzd+bAdBd,b-bCdDdEd.bFd/b:b;b=b?b@b[b]b^bGd_b`b{b|bHd}b~bIdJdKdLdacMdbcNdOdccdcPdQdRdecSdfcgchcicjckclcmcncTdocpcqcrcUdsctcVdWdHcIc0dJcRlSlTlUlwcVl~ WlXlYlZl0l1l2l3l4l5l6l7l8lbbcb#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~lambmcmdmemfmgmhmimjmkmlmmmnmompmqmrmsmtmumvmwmxmymzmAmBmCmDmEmFmGmHmImJmKmLmMmKc9oOmPm!oRm#o$oTmUm%oheieWmXm'oZm0m(oMe)o1m5d2m-c.c3mNe*o+o4m6d5m,o7m8m-ojeke!m#m.o%m'm/oOe:o(m7d)m:c;c*mPe;o=o+m8d,m?o.m/m:m;m=m@o[o@m[m]o^o^m_m_o`o{m|m{o|o~man}o~ocndnapbpfngnhncpdpjnknlnepebmngdonpnqnrngphpsnhbibkplpmpnpqprpsptpwptnypBpCp

233 elif isinstance(options, cls): 349 ctx2k l m n o p q c r d e f s g h t u v w x y z A B C D E F G H I J K L M N O P Q R S T U V W 8n9nEl!nFl#n$n%nGl'nHl(ni )nj *n{d|dneoe}d~dpeqeKeLeIl+n,n-n.n/n:nid;n=n?n@n[nxcyc]n^n_n`n{n|nZ 0 1 2 3 4 JlucKlvqwqLlMlNlhdX CcvcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c9c!c#c$c%c'c(c)c*c+cb 5 Tp6 Up7 Vp8 Wp9 Xp! Yp# Zp$ 0p% 1p' 2p( 3n) 4n* 3p+ 4p, 5p- 6p. 7p/ 8p: 9p; !p= #p? $p@ %p[ 'p] (p^ 9d_ )p` *p{ +p| ,pOlHcIcYdrosotouovowoxoyozoAoBoCoDoEoFoGoHoIoJoKoLoMoNoOoPoQoRoSoToUoVoWoXoYoZo0o1o2o3o4oXnYnZn0n1n2n5oJc7o8ozcAcY Bc} wc9l!l!d1d2dKcfdEpFpdbfpfbgbybzblbmbnbAbBbCbDbjbkbobEbpbFbGbHbIbqbJbrbsbtbubKbLbMbNbObPbQbRbSbvbwbxbGpHpIpJpKpLpMpnnNpOpPpQpRpSpQeReab-p.p/p:p;p=p?p@p[pipjpopppupvpxptnzpApDp

234 return options 336 ctx2k l m n o p q c r d e f s g h t u v w x y z A B C D E F G H I J K L M N O P Q R S T U V W 8n9nEl!nFl#n$n%nGl'nHl(ni )nj *nneoepeqeIl+n,n-n.n/n:nid;n=n?n@n[nxcyc]n^n_n`n{n|nZ 0 1 2 3 4 JlucKlvqwqLlMlNlhdX CcvcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c9c!c#c$c%c'c(c)c*c+cb 5 Tp6 Up7 Vp8 Wp9 Xp! Yp# Zp$ 0p% 1p' 2p( 3n) 4n* 3p+ 4p, 5p- 6p. 7p/ 8p: 9p; !p= #p? $p@ %p[ 'p] (p^ 9d_ )p` *p{ +p| ,pOlHcIcYdrosotouovowoxoyozoAoBoCoDoEoFoGoHoIoJoKoLoMoNoOoPoQoRoSoToUoVoWoXoYoZo0o1o2o3o4o5oJc7o8ozcAcY Bc} wc9l!l!d1d2dKcfdEpFpdbfpfbgbybzblbmbnbAbBbCbDbjbkbobEbpbFbGbHbIbqbJbrbsbtbubKbLbMbNbObPbQbRbSbvbwbxbGpHpIpJpKpLpMpnnNpOpPpQpRpSpQeRe-p.p/p:p;p=p?p@p[pipjpopppupvpxptnzpApDp

235 elif isinstance(options, dict): 14 ctx2{d|d}d~dKeLeCcXnYnZn0n1n2nab

236 return cls(**options) 14 ctx2{d|d}d~dKeLeCcXnYnZn0n1n2nab

237 else: 

238 raise TypeError( 

239 f"The {options_description} must be provided as an object " 

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

241 f"The provided object is '{options}'." 

242 ) 

243  

244  

245def _handle_boolean_option(option: bool) -> str: (empty)

246 """ 

247 Convert a boolean option to a string representation. 

248 """ 

249 return "true" if bool(option) else "false" 42 ctx2b 5 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } lbmbnbpbtbubvbwbxbFy

250  

251  

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

253 """ 

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

255  

256 Args: 

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

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

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

260  

261 Returns: 

262 Callable: A decorator that creates the wrapping. 

263 """ 

264  

265 def outer(wrapped_function): 

266 """ 

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

268 """ 

269  

270 @functools.wraps(wrapped_function) 

271 def inner(*args, **kwargs): 

272 """ 

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

274 """ 

275 checker(*args, **kwargs, what=what) 

276 result = wrapped_function(*args, **kwargs) 

277  

278 return result 

279  

280 return inner 

281  

282 return outer 

283  

284  

285def is_sequence(obj): (empty)

286 """ 

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

288 """ 

289 return isinstance(obj, Sequence) 47 ctx2i j ucX CcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c9c!c#c$c%c'c(c)c*c+c3n4njbkbobqbrbsbQeRe

290  

291  

292def is_nested_sequence(obj): (empty)

293 """ 

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

295 """ 

296 return is_sequence(obj) and any(is_sequence(elem) for elem in obj) 2 ctx2jbkb

297  

298  

299@functools.lru_cache (empty)

300def get_binding_version(): 

301 try: (empty)

302 major_minor = importlib.metadata.version("cuda-bindings").split(".")[:2] (empty)

303 except importlib.metadata.PackageNotFoundError: 

304 major_minor = importlib.metadata.version("cuda-python").split(".")[:2] 

305 return tuple(int(v) for v in major_minor) (empty)

306  

307  

308class Transaction: (empty)

309 """ 

310 A context manager for transactional operations with undo capability. 

311  

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

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

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

315  

316 Usage: 

317 with Transaction() as txn: 

318 txn.append(some_cleanup_function, arg1, arg2) 

319 # ... perform operations ... 

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

321  

322 Methods: 

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

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

325 """ 

326 def __init__(self): (empty)

327 self._stack = ExitStack() 4 ctx2zcAcY Bc

328 self._entered = False 4 ctx2zcAcY Bc

329  

330 def __enter__(self): (empty)

331 self._stack.__enter__() 4 ctx2zcAcY Bc

332 self._entered = True 4 ctx2zcAcY Bc

333 return self 4 ctx2zcAcY Bc

334  

335 def __exit__(self, exc_type, exc, tb): (empty)

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

337 self._entered = False 4 ctx2zcAcY Bc

338 return self._stack.__exit__(exc_type, exc, tb) 4 ctx2zcAcY Bc

339  

340 def append(self, fn, /, *args, **kwargs): (empty)

341 """ 

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

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

344 """ 

345 if not self._entered: 4 ctx2zcAcY Bc

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

347 self._stack.callback(partial(fn, *args, **kwargs)) 4 ctx2zcAcY Bc

348  

349 def commit(self): (empty)

350 """ 

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

352 """ 

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

354 self._stack.pop_all() 4 ctx2zcAcY Bc

355  

356  

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

358_fork_warning_checked = False (empty)

359  

360  

361def reset_fork_warning(): (empty)

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

363  

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

365 to check the warning behavior. 

366 """ 

367 global _fork_warning_checked 

368 _fork_warning_checked = False 5 ctx2!d1d2dKcfd

369  

370  

371def check_multiprocessing_start_method(): (empty)

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

373 global _fork_warning_checked 

374 if _fork_warning_checked: 49 ctx2]p^pElFl_p`pGlHli j {d|dneoe}d~dpeqe{p|p}p~paqbqrecq~cadbdcdxcycddeddqeqfqgqZ 0 1 2 3 4 !d1d2dKcfd

375 return 45 ctx2]p^pElFl_p`pGlHli j {d|dneoe}d~dpeqe{p|p}p~paqbqrecq~cadbdcdxcycddeddqeqfqgqZ 0 1 2 3 4 fd

376 _fork_warning_checked = True 6 ctx2re!d1d2dKcfd

377  

378 # Common warning message parts 

379 common_message = ( 

380 "CUDA does not support. Forked subprocesses exhibit undefined behavior, " 6 ctx2re!d1d2dKcfd

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

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

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

384 ) 

385  

386 try: 6 ctx2re!d1d2dKcfd

387 start_method = multiprocessing.get_start_method() 6 ctx2re!d1d2dKcfd

388 if start_method == "fork": 6 ctx2re!d1d2dKcfd

389 message = f"multiprocessing start method is 'fork', which {common_message}" 4 ctx21d2dKcfd

390 warnings.warn(message, UserWarning, stacklevel=3) 4 ctx21d2dKcfd

391 except RuntimeError: 

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

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

394 if platform.system() == "Linux": 

395 message = ( 

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

397 f"which {common_message}" 

398 ) 

399 warnings.warn(message, UserWarning, stacklevel=3)