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 ucgdX BcYnxe!dZnyevcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c9c!c#c$c%c'c(c)c*cY zeAe+c[d]d4d,c-c5dBe.c^d_d6d/c:c7dCeDeEeFeGeHeIeJeneOnPnAl0noe

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 ucgdX BcYnxe!dZnyevcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c9c!c#c$c%c'c(c)c*cY zeAe+c[d]d4d,c-c5dBe.c^d_d6d/c:c7dCeDeEeFeGeHeIeJeneOnPnAl0noe

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 ucgdX BcYn!dZnMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c9c!c#c$c%c'c(c)c*cY neOn

48 else: 

49 common = "must be an int, or a tuple with up to 3 ints" 34 ctx2xe!dyevczeAe+c[d]d4d,c-c5dBe.c^d_d6d/c:c7dCeDeEeFeGeHeIeJenePnAl0noe

50 if not isinstance(cfg, tuple): 34 ctx2xe!dyevczeAe+c[d]d4d,c-c5dBe.c^d_d6d/c:c7dCeDeEeFeGeHeIeJenePnAl0noe

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

52 if len(cfg) > 3: 33 ctx2xe!dyevczeAe+c[d]d4d,c-c5dBe.c^d_d6d/c:c7dCeDeEeFeGeHeIeJenePnAloe

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

54 if any(not isinstance(val, int) for val in cfg): 32 ctx2xe!dyevczeAe+c[d]d4d,c-c5dBe.c^d_d6d/c:c7dCeDeEeFeGeHeIeJeneAloe

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

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 ucgdX BcYnxe!dZnyevcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c9c!c#c$c%c'c(c)c*cY zeAe+c[d]d4d,c-c5dBe.c^d_d6d/c:c7dCeDeEeFeGeHeIeJeneOnoe

57 plural_s = "" if len(cfg) == 1 else "s" 3 ctx2!dOnoe

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

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 ucgdX BcYnxe!dZnyevcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c9c!c#c$c%c'c(c)c*cY zeAe+c[d]d4d,c-c5dBe.c^d_d6d/c:c7dCeDeEeFeGeHeIeJene

60  

61  

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

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

64  

65  

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

67 if err != cydriver.CUresult.CUDA_SUCCESS: 1436 ctx2a uqvqwqxqyqzqk `dl Aqm Bq;cCqn Dq=cEqo Fqp Gqq HqCcIqDcJqEcKq?cLq@cMq[cNqc Oqr Pqd Qqe Rqf Sqs Tqg Uqh Vqt Wqu Xqv Yq]cZq^c0q_c1q`c2q{c3q|c4qw 5qx 6qy 7qz 8qA 9qB !qC #qD $qE %qF 'qG (qH )qI *qJ +qK ,qL -qM .qN /qO :qP ;qQ =qR ?qS @qT [qU ]qV ^qW _q:p1n;p2nBl3nCl4n=p5n?p6nDl7nEl8ni 9nj !n{d`q|d{qpe|qqe}q}d~q~darrebrsecrKedrLeerFlfr@p#n[p$n]p%n^p'n_p(n`p)n{p*n|p+n}c,n~c-nad.nbd/nwcgrxcaecd:ndd;n}p=n~p?naq@ntehd0 hr1 ir2 jr3 kr4 lr5 mr#d[nnr]nor^npr_nqr`nrr{nsrtr$didFcur|nvrwrxryrzrArBrCrDr%dErFrGrHrIrJrKrLrMrNrOrPrQrRrSrTrUrVrWrXrYrZr0r1r2r3r4r5r6r7r8r9r!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}sbe~satbtce'd(ddeGlucHlct}ndt~netaoftbogtcohtdoiteojtfoktIlltJlmtKlntgdotptqtrtstttutvtwtxtytztAtBtgoCteeDtEtX TbUbBcFtGtHtvcItMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c9c!c#c$c%c'c(c)c*c6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } LlJtKt)dVbWb*dLt+d,d-dMlNlMt.d/d:d;dfegejdNtkdOtXbPtldQtYbRtmdStndTtZbUt0bVtodWtpdXtqdYt1bZtrd0t2b1t3b2t4b3t5b4t6b5t7b6t8b7t9b8t!b9tsd!t#b#t$b$t%b%t'b'ttd(t(b)t)b*tud+tvd,twd-txd.t*b/tyd:t+b;tzd=tAd?t,b@t-b[tBd]tCd^tDd_t.b`tEd{t/b|t:b}t;b~t=bau?bbu@bcu[bdu]beu^bfuFdgu_bhu`biu{bju|bkuGdlu}bmu~bnuHdouIdpuJdquKdruacsuLdtubcuuMdvuNdwuccxudcyuOdzuPdAuQdBuecCuRdDufcEugcFuhcGuicHujcIukcJulcKumcLuncMuSdNuocOupcPuqcQurcRuTdSuscTutcUuUdVuVdWu=dhoioGcXuYujoZuHc0uWd1uXd2uYd3uko4ulo5umo6uno7uoo8upo9uqo!uro#uso$uto%uuo'uvo(uwo)uxo*uyo+uzo,uAo-u.u/uBo:u;u=uCo?u@u[uDo]u^u_uEo`u{u|uFo}u~uavGobvcvdvHoevfvgvIohvJoivKojvLokvMolvNomvOonvPoovQopvRoqvSorvTosvUotvVouvWovvXowvQnxvRnyvSnzvTnAvUnBvVnCvYoDv?d@dZdZoEvIcFv0o1oyczcGvHvZ IvAcJvKv~ LvOlPlMvQlNvRlY OvSlabTlPvUlQvVlRvWlSvXlTvYlheZlUv0lVv1lWv2lie3l4l5l6l7lcbdbXv8lYv9lZv!l0v#l1v$l2v%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`l{l.v|l/v}l:v~l;vam=vbm?vcm@vdm[vem]vfm^vgm_vhm`vim{vjm|vkm}vlm~vmmawnmbwomcwpmdwqmewrmsmfwtmgwumhwvmiwwmjwxmkwymlwzmmwAmnwBmowCmpwDmqwEmrwFmswGmtwHmuwImJm9dvw0dww1dxwJcywedzwKmAwBw2oLmMm3oNmOm4oPmCwDw5oQmRm6oje+ckeSmTm[dUmVmEw7oWm]dXm8oMe9oYm4dZm,c-c0mNe!o#o1m5d2m3mFwGw$o4m5m%ole.cme6m7m^d8m9mHw'o!m_d#m(oOe)o$m6d%m/c:c'mPe*o+o(m7d)m*mIwJw,o+m,m-mKc.m/mLc-o:mKwLw.o;m=m/o?mMwNw:o@m[m;o]mOwPw=o^m_m?o`mQwRw@o{m|m[o}mSwTw]o~man^obnUwVw_ocndnen`ofnWwXw{ognhnin|oebYwfbZw}ogb0whb1w2w3w4w5w6w7w8w9w!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~waxjnfdbxcxdxexfxgxhxixjxkxlxmxnxoxQeReknlnmnbbpxnn~oaponqxrxibsxtxuxvxwxxxyxzxAxbpBxcpCxdpDxepExfpFxgpGxhpHxipIxjpJxkpKxlpLxmpMxnpNxopOxppPxqpQxpnRxrpSxspTxtpUxupVxvpWxwpXxYxZx0x1x2x3x4x5x6x7x8x9x!x#x$x%x'x(x)x*x+x,x-x.x/xqnrnbqsntnuncqdqvnwnxneqynznfqAnBnCngqhqiqDnjqkqEnlqmqFnnqGnoqHnInJnKnpqLnMnqq:x;x=xNn?x@x[x]x^x_x`x{x|x}x~xaybycydyeyfygyhyiyjykylymynyoypyqyrysytyuyvywyxyyyzyAy

68 return _check_driver_error(err) 2 ctx2abfd

69 return 0 1436 ctx2a uqvqwqxqyqzqk `dl Aqm Bq;cCqn Dq=cEqo Fqp Gqq HqCcIqDcJqEcKq?cLq@cMq[cNqc Oqr Pqd Qqe Rqf Sqs Tqg Uqh Vqt Wqu Xqv Yq]cZq^c0q_c1q`c2q{c3q|c4qw 5qx 6qy 7qz 8qA 9qB !qC #qD $qE %qF 'qG (qH )qI *qJ +qK ,qL -qM .qN /qO :qP ;qQ =qR ?qS @qT [qU ]qV ^qW _q:p1n;p2nBl3nCl4n=p5n?p6nDl7nEl8ni 9nj !n{d`q|d{qpe|qqe}q}d~q~darrebrsecrKedrLeerFlfr@p#n[p$n]p%n^p'n_p(n`p)n{p*n|p+n}c,n~c-nad.nbd/nwcgrxcaecd:ndd;n}p=n~p?naq@ntehd0 hr1 ir2 jr3 kr4 lr5 mr#d[nnr]nor^npr_nqr`nrr{nsrtr$didFcur|nvrwrxryrzrArBrCrDr%dErFrGrHrIrJrKrLrMrNrOrPrQrRrSrTrUrVrWrXrYrZr0r1r2r3r4r5r6r7r8r9r!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}sbe~satbtce'd(ddeGlucHlct}ndt~netaoftbogtcohtdoiteojtfoktIlltJlmtKlntgdotptqtrtstttutvtwtxtytztAtBtgoCteeDtEtX TbUbBcFtGtHtvcItMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c9c!c#c$c%c'c(c)c*c6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } LlJtKt)dVbWb*dLt+d,d-dMlNlMt.d/d:d;dfegejdNtkdOtXbPtldQtYbRtmdStndTtZbUt0bVtodWtpdXtqdYt1bZtrd0t2b1t3b2t4b3t5b4t6b5t7b6t8b7t9b8t!b9tsd!t#b#t$b$t%b%t'b'ttd(t(b)t)b*tud+tvd,twd-txd.t*b/tyd:t+b;tzd=tAd?t,b@t-b[tBd]tCd^tDd_t.b`tEd{t/b|t:b}t;b~t=bau?bbu@bcu[bdu]beu^bfuFdgu_bhu`biu{bju|bkuGdlu}bmu~bnuHdouIdpuJdquKdruacsuLdtubcuuMdvuNdwuccxudcyuOdzuPdAuQdBuecCuRdDufcEugcFuhcGuicHujcIukcJulcKumcLuncMuSdNuocOupcPuqcQurcRuTdSuscTutcUuUdVuVdWu=dhoioGcXuYujoZuHc0uWd1uXd2uYd3uko4ulo5umo6uno7uoo8upo9uqo!uro#uso$uto%uuo'uvo(uwo)uxo*uyo+uzo,uAo-u.u/uBo:u;u=uCo?u@u[uDo]u^u_uEo`u{u|uFo}u~uavGobvcvdvHoevfvgvIohvJoivKojvLokvMolvNomvOonvPoovQopvRoqvSorvTosvUotvVouvWovvXowvQnxvRnyvSnzvTnAvUnBvVnCvYoDv?d@dZdZoEvIcFv0o1oyczcGvHvZ IvAcJvKv~ LvOlPlMvQlNvRlY OvSlabTlPvUlQvVlRvWlSvXlTvYlheZlUv0lVv1lWv2lie3l4l5l6l7lcbdbXv8lYv9lZv!l0v#l1v$l2v%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`l{l.v|l/v}l:v~l;vam=vbm?vcm@vdm[vem]vfm^vgm_vhm`vim{vjm|vkm}vlm~vmmawnmbwomcwpmdwqmewrmsmfwtmgwumhwvmiwwmjwxmkwymlwzmmwAmnwBmowCmpwDmqwEmrwFmswGmtwHmuwImJm9dvw0dww1dxwJcywedzwKmAwBw2oLmMm3oNmOm4oPmCwDw5oQmRm6oje+ckeSmTm[dUmVmEw7oWm]dXm8oMe9oYm4dZm,c-c0mNe!o#o1m5d2m3mFwGw$o4m5m%ole.cme6m7m^d8m9mHw'o!m_d#m(oOe)o$m6d%m/c:c'mPe*o+o(m7d)m*mIwJw,o+m,m-mKc.m/mLc-o:mKwLw.o;m=m/o?mMwNw:o@m[m;o]mOwPw=o^m_m?o`mQwRw@o{m|m[o}mSwTw]o~man^obnUwVw_ocndnen`ofnWwXw{ognhnin|oebYwfbZw}ogb0whb1w2w3w4w5w6w7w8w9w!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~waxjnfdbxcxdxexfxgxhxixjxkxlxmxnxoxQeReknlnmnbbpxnn~oaponqxrxibsxtxuxvxwxxxyxzxAxbpBxcpCxdpDxepExfpFxgpGxhpHxipIxjpJxkpKxlpLxmpMxnpNxopOxppPxqpQxpnRxrpSxspTxtpUxupVxvpWxwpXxYxZx0x1x2x3x4x5x6x7x8x9x!x#x$x%x'x(x)x*x+x,x-x.x/xqnrnbqsntnuncqdqvnwnxneqynznfqAnBnCngqhqiqDnjqkqEnlqmqFnnqGnoqHnInJnKnpqLnMnqq:x;x=xNn?x@x[x]x^x_x`x{x|x}x~xaybycydyeyfygyhyiyjykylymynyoypyqyrysytyuyvywyxyyyzyAy

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: 306 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 ucgdX BcvcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c9c!c#c$c%c'c(c)c*cb 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } Ll~ OlPlQlRlY SlabTlUlVlWlXlYlZl0l1l2l3l4l5l6l7lcbdb8l9l!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~lambmcmdmemfmgmhmimjmkmlmmmnmompmqmrmsmtmumvmwmxmymzmAmBmCmDmEmFmGmHmImJmLmMmOmQmRmjekeTmUmWmXmYmZm,c-c0m1m2m4m5mleme7m8m!m#m$m%m/c:c'm(m)m+m,m-m.m;m=m@m[m^m_m{m|m~mancndnengnhninebfbgbhbybzblbmbnbAbBbCbDbjbkbobEbpbFbGbHbIbqbJbrbsbtbubKbLbMbNbObPbQbRbSbvbwbxbknlnmnbbnnonib

75 return 0 306 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 ucgdX BcvcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c9c!c#c$c%c'c(c)c*cb 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } Ll~ OlPlQlRlY SlabTlUlVlWlXlYlZl0l1l2l3l4l5l6l7lcbdb8l9l!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~lambmcmdmemfmgmhmimjmkmlmmmnmompmqmrmsmtmumvmwmxmymzmAmBmCmDmEmFmGmHmImJmLmMmOmQmRmjekeTmUmWmXmYmZm,c-c0m1m2m4m5mleme7m8m!m#m$m%m/c:c'm(m)m+m,m-m.m;m=m@m[m^m_m{m|m~mancndnengnhninebfbgbhbybzblbmbnbAbBbCbDbjbkbobEbpbFbGbHbIbqbJbrbsbtbubKbLbMbNbObPbQbRbSbvbwbxbknlnmnbbnnonib

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/mxpypjnfdzpApBpCpDpEpFpGpHpIpJpKpLpMpQeRe

101 return 0 21 ctx2/mxpypjnfdzpApBpCpDpEpFpGpHpIpJpKpLpMpQeRe

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 NpOpPpQpRpSpTpUpVpWpWnXnXpYpZp0p1p2p3p4p5p6p7p8p9p8d!p#p$p%p'p(p)p*p+p,p-p.p/p

127 return 0 39 ctx2b NpOpPpQpRpSpTpUpVpWpWnXnXpYpZp0p1p2p3p4p5p6p7p8p9p!p#p$p%p'p(p)p*p+p,p-p.p/p

128 with gil: 2 ctx2b 8d

129 _raise_nvjitlink_error(handle, err) 2 ctx2b 8d

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 8d

136 if handle != NULL: 2 ctx2b 8d

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

138 cdef bytes log_bytes 

139 cdef str log_str = "" 2 ctx2b 8d

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

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 8d

146 if log_str: 2 ctx2b 8d

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

148 raise exc 2 ctx2b 8d

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: 998 ctx2a SeTeUeVeWeXeYeZe0e1e2e3e4e5e6e7e8e9e!e#e$e%e'e(ek `d)el *em +e;c,en -e=c.eo /ep :eq ;eCcDcEc?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 }eH ~eI afJ bfK cfL dfM efN ffO gfP hfQ ifR jfS kfT lfU mfV nfW ofi j }c~cadbdwcxcaecdddhd0 1 2 3 4 5 pfqfrfsftfufvfuewfFcxfyfzfAfBfCfDfEfFfGfHfIfbeJfKfLfceMfdeucNfOfPfQfRfSfTfUfVfWfXfYfZfeeX TbUb0f1f2f3fMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c9c!c#c$c%c'c(c)c*cb 6 4f7 5f8 6f9 7f! 8f# 9f$ !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~f)dVbWb*d+d,d-dMlNl.d/d:d;dfegejdkdXbldYbmdndZb0bodpdqd1brd2b3b4b5b6b7b8b9b!bsd#b$b%b'btd(b)budvdwdxd*byd+bzdAd,b-bBdCdDd.bEd/b:b;b=b?b@b[b]b^bFd_b`b{b|bGd}b~bHdIdJdKdacLdbcMdNdccdcOdPdQdecRdfcgchcicjckclcmcncSdocpcqcrcTdsctcUdVd=dagbgcgdgegfggghgigjgkglgmgngogpgqgrgsgtgugvgwgxgygzgAgBgCgDgEgFgGgHgIgJgKgLgMgNgOgPgQgRgSgTgUgVgWgXgYgZg0g1g2g3g4g5g6g7g8g?d@dZd9g!gyczcZ Ac~ #g$g%g'g(gY )g*gab+g,g-g.g/ghe:g;g=g?gie@g[g]g^g_g`g2dcb{g3ddb|g}g~gahbhchdhehfhghhhihjhkhlhmhnhohphqhrhshthuhvhwhxhyhzhAhBhChDhEhFhGhHhIhJhKhLhMhNhOhPhQhRhShThUhVhWhXhYhZh0h1h2h3h4h5h6h7h8h9h!hKm#h$h%h'h(h)h*hNm+h,h-hPm.h/h:h;h=h?h@hje[h+c]hke^hSm_h`h{h|hVm}h~haibicidieiMefigihiiijikilimiNenioipiqirisi3mtiuiviwixiyizileAi.cBimeCi6mDiEiFiGi9mHiIiJiKiLiMiNiOeOiPiQiRiSiTiUiViPeWiXiYiZi0i1i*m2i3i4i5i6i7i8iKc9i!i#iLc$i%i:m'i(i)i*i+i,i-i?m.i/i:i;i=i?i@i]m[i]i^i_i`i{i|i`m}i~iajbjcjdjej}mfjgjhjijjjkjljbnmjnjojpjqjrjsjtjfnujvjwjxjyjzjAjBjebCjfbDjgbEjhbFjybGjzbHjlbIjmbJjnbKjAbLjBbMjCbNjDbOjjbPjkbQjobRjEbSjpbTjFbUjGbVjHbWjIbXjqbYjJbZjrb0jsb1jtb2jub3jKb4jLb5jMb6jNb7jOb8jPb9jQb!jRb#jSb$jvb%jwb'jxb(jfd)j*j+j,j-j.j/j:j;j=j?j@j[j]jbb^j_j`jib{j|j}j~jakbkckdkekfkgkhkikjkkklkmknkokpkqkrksktkukvkwkxkykzkAkBkCkDkEkFkGkHkIkJkKkLkMkNkOkPkQkRkSkTkUkVkWkXkYkZk0k1k2k3k4k5k6k7k8k9k!k#k$k%k'k(k)k*k+k,k-k.k/k:k;kqnrnbqsntnuncqdqvnwnxneqynznfqAnBnCngqhqiqDnjqkqEnlqmqFnnqGnoqHnInJnKnpqLnMnqq=k?k@k[k]kNn^k_k`k{k|k}k~kalblcldlelflglhliljlklllmlnlolplqlrlsltlulvlwlxlyl

157 return 0 997 ctx2a SeTeUeVeWeXeYeZe0e1e2e3e4e5e6e7e8e9e!e#e$e%e'e(ek `d)el *em +e;c,en -e=c.eo /ep :eq ;eCcDcEc?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 }eH ~eI afJ bfK cfL dfM efN ffO gfP hfQ ifR jfS kfT lfU mfV nfW ofi j }c~cadbdwcxcaecdddhd0 1 2 3 4 5 pfqfrfsftfufvfuewfFcxfyfzfAfBfCfDfEfFfGfHfIfbeJfKfLfceMfdeucNfOfPfQfRfSfTfUfVfWfXfYfZfeeX TbUb0f1f2f3fMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c9c!c#c$c%c'c(c)c*cb 6 4f7 5f8 6f9 7f! 8f# 9f$ !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~f)dVbWb*d+d,d-dMlNl.d/d:d;dfegejdkdXbldYbmdndZb0bodpdqd1brd2b3b4b5b6b7b8b9b!bsd#b$b%b'btd(b)budvdwdxd*byd+bzdAd,b-bBdCdDd.bEd/b:b;b=b?b@b[b]b^bFd_b`b{b|bGd}b~bHdIdJdKdacLdbcMdNdccdcOdPdQdecRdfcgchcicjckclcmcncSdocpcqcrcTdsctcUdVd=dagbgcgdgegfggghgigjgkglgmgngogpgqgrgsgtgugvgwgxgygzgAgBgCgDgEgFgGgHgIgJgKgLgMgNgOgPgQgRgSgTgUgVgWgXgYgZg0g1g2g3g4g5g6g7g8g?d@dZd9g!gyczcZ Ac~ #g$g%g'g(gY )g*gab+g,g-g.g/ghe:g;g=g?gie@g[g]g^g_g`g2dcb{g3ddb|g}g~gahbhchdhehfhghhhihjhkhlhmhnhohphqhrhshthuhvhwhxhyhzhAhBhChDhEhFhGhHhIhJhKhLhMhNhOhPhQhRhShThUhVhWhXhYhZh0h1h2h3h4h5h6h7h8h9h!hKm#h$h%h'h(h)h*hNm+h,h-hPm.h/h:h;h=h?h@hje[h+c]hke^hSm_h`h{h|hVm}h~haibicidieiMefigihiiijikilimiNenioipiqirisi3mtiuiviwixiyizileAi.cBimeCi6mDiEiFiGi9mHiIiJiKiLiMiNiOeOiPiQiRiSiTiUiViPeWiXiYiZi0i1i*m2i3i4i5i6i7i8iKc9i!i#iLc$i%i:m'i(i)i*i+i,i-i?m.i/i:i;i=i?i@i]m[i]i^i_i`i{i|i`m}i~iajbjcjdjej}mfjgjhjijjjkjljbnmjnjojpjqjrjsjtjfnujvjwjxjyjzjAjBjebCjfbDjgbEjhbFjybGjzbHjlbIjmbJjnbKjAbLjBbMjCbNjDbOjjbPjkbQjobRjEbSjpbTjFbUjGbVjHbWjIbXjqbYjJbZjrb0jsb1jtb2jub3jKb4jLb5jMb6jNb7jOb8jPb9jQb!jRb#jSb$jvb%jwb'jxb(j)j*j+j,j-j.j/j:j;j=j?j@j[j]jbb^j_j`jib{j|j}j~jakbkckdkekfkgkhkikjkkklkmknkokpkqkrksktkukvkwkxkykzkAkBkCkDkEkFkGkHkIkJkKkLkMkNkOkPkQkRkSkTkUkVkWkXkYkZk0k1k2k3k4k5k6k7k8k9k!k#k$k%k'k(k)k*k+k,k-k.k/k:k;kqnrnbqsntnuncqdqvnwnxneqynznfqAnBnCngqhqiqDnjqkqEnlqmqFnnqGnoqHnInJnKnpqLnMnqq=k?k@k[k]kNn^k_k`k{k|k}k~kalblcldlelflglhliljlklllmlnlolplqlrlsltlulvlwlxlyl

158 cdef const char* name 

159 name_err = cydriver.cuGetErrorName(error, &name) 4 ctx2ueZ abfd

160 if name_err != cydriver.CUresult.CUDA_SUCCESS: 4 ctx2ueZ abfd

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

162 with gil: 4 ctx2ueZ abfd

163 # TODO: consider lower this to Cython 

164 expl = DRIVER_CU_RESULT_EXPLANATIONS.get(int(error)) 4 ctx2ueZ abfd

165 if expl is not None: 4 ctx2ueZ abfd

166 raise CUDAError(f"{name.decode()}: {expl}") 4 ctx2ueZ abfd

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 ctx20 1 2 3 4 5 #dvewezl$didFc%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~bacbcccdcecfcgchcicjckclcmcncocpcqcrcsctcGcHcWdXdYdIc

176 return 0 86 ctx20 1 2 3 4 5 #dvewezl$didFc%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~bacbcccdcecfcgchcicjckclcmcncocpcqcrcsctcGcHcWdXdYdIc

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

178 if name_err != _RUNTIME_SUCCESS: 1 ctx2zl

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

180 name = name.decode() 1 ctx2zl

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

182 if expl is not None: 1 ctx2zl

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

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: 81 ctx2a b 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ 2dcb3ddbKcLcebfbgbhbybzblbmbnbAbBbCbDbjbkbobEbpbFbGbHbIbqbJbrbsbtbubKbLbMbNbObPbQbRbSbvbwbxbbbib

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

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(ek `d)el *em +e;c,en -e=c.eo /ep :eq ;eCcDcEc?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 }eH ~eI afJ bfK cfL dfM efN ffO gfP hfQ ifR jfS kfT lfU mfV nfW ofi j }c~cadbdwcxcaecdddhd0 1 2 3 4 5 #dvepfweqfrfsftfufvf$didwfFcxfyfzfAfBfCfDfEf%dFfGfHfIfbeJfKfLfce'd(dMfdeNfOfPfQfRfSfTfUfVfWfXfYfZfeeX TbUb0f1f2f3fb 6 4f7 5f8 6f9 7f! 8f# 9f$ !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~f)dVbWb*d+d,d-d.d/d:d;dfegejdkdXbldYbmdndZb0bodpdqd1brd2b3b4b5b6b7b8b9b!bsd#b$b%b'btd(b)budvdwdxd*byd+bzdAd,b-bBdCdDd.bEd/b:b;b=b?b@b[b]b^bFd_b`b{b|bGd}b~bHdIdJdKdacLdbcMdNdccdcOdPdQdecRdfcgchcicjckclcmcncSdocpcqcrcTdsctcUdVd=dGcagbgcgHcdgWdegXdfgYdgghgigjgkglgmgngogpgqgrgsgtgugvgwgxgygzgAgBgCgDgEgFgGgHgIgJgKgLgMgNgOgPgQgRgSgTgUgVgWgXgYgZg0g1g2g3g4g5g6g7g8g?d@dZd9gIc!g~ #g$g%g'g(gY )g*gab+g,g-g.g/ghe:g;g=g?gie@g[g]g^g_g`g2dcb{g3ddb|g}g~gahbhchdhehfhghhhihjhkhlhmhnhohphqhrhshthuhvhwhxhyhzhAhBhChDhEhFhGhHhIhJhKhLhMhNhOhPhQhRhShThUhVhWhXhYhZh0h1h2h3h4h5h6h7h8h9h!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~haibicidieifigihiiijikiliminioipiqirisitiuiviwixiyiziAiBiCiDiEiFiGiHiIiJiKiLiMiNiOiPiQiRiSiTiUiViWiXiYiZi0i1i2i3i4i5i6i7i8iKc9i!i#iLc$i%i'i(i)i*i+i,i-i.i/i:i;i=i?i@i[i]i^i_i`i{i|i}i~iajbjcjdjejfjgjhjijjjkjljmjnjojpjqjrjsjtjujvjwjxjyjzjAjBjebCjfbDjgbEjhbFjybGjzbHjlbIjmbJjnbKjAbLjBbMjCbNjDbOjjbPjkbQjobRjEbSjpbTjFbUjGbVjHbWjIbXjqbYjJbZjrb0jsb1jtb2jub3jKb4jLb5jMb6jNb7jOb8jPb9jQb!jRb#jSb$jvb%jwb'jxb(j)j*j+j,j-j.j/j:j;j=j?j@j[j]jbb^j_j`jib{j|j}j~jakbkckdkekfkgkhkikjkkklkmknkokpkqkrksktkukvkwkxkykzkAkBkCkDkEkFkGkHkIkJkKkLkMkNkOkPkQkRkSkTkUkVkWkXkYkZk0k1k2k3k4k5k6k7k8k9k!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~kalblcldlelflglhliljlklllmlnlolplqlrlsltlulvlwlxlyl

205 return _check_driver_error(error) 890 ctx2a SeTeUeVeWeXeYeZe0e1e2e3e4e5e6e7e8e9e!e#e$e%e'e(ek `d)el *em +e;c,en -e=c.eo /ep :eq ;eCcDcEc?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 }eH ~eI afJ bfK cfL dfM efN ffO gfP hfQ ifR jfS kfT lfU mfV nfW ofi j }c~cadbdwcxcaecdddhd0 1 2 3 4 5 pfqfrfsftfufvfwfxfyfzfAfBfCfDfEfFfGfHfIfbeJfKfLfceMfdeNfOfPfQfRfSfTfUfVfWfXfYfZfeeX TbUb0f1f2f3fb 6 4f7 5f8 6f9 7f! 8f# 9f$ !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~f)dVbWb*d+d,d-d.d/d:d;dfegejdkdXbldYbmdndZb0bodpdqd1brd2b3b4b5b6b7b8b9b!bsd#b$b%b'btd(b)budvdwdxd*byd+bzdAd,b-bBdCdDd.bEd/b:b;b=b?b@b[b]b^bFd_b`b{b|bGd}b~bHdIdJdKdacLdbcMdNdccdcOdPdQdecRdfcgchcicjckclcmcncSdocpcqcrcTdsctcUdVd=dagbgcgdgegfggghgigjgkglgmgngogpgqgrgsgtgugvgwgxgygzgAgBgCgDgEgFgGgHgIgJgKgLgMgNgOgPgQgRgSgTgUgVgWgXgYgZg0g1g2g3g4g5g6g7g8g?d@dZd9g!g~ #g$g%g'g(gY )g*gab+g,g-g.g/ghe:g;g=g?gie@g[g]g^g_g`g2dcb{g3ddb|g}g~gahbhchdhehfhghhhihjhkhlhmhnhohphqhrhshthuhvhwhxhyhzhAhBhChDhEhFhGhHhIhJhKhLhMhNhOhPhQhRhShThUhVhWhXhYhZh0h1h2h3h4h5h6h7h8h9h!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~haibicidieifigihiiijikiliminioipiqirisitiuiviwixiyiziAiBiCiDiEiFiGiHiIiJiKiLiMiNiOiPiQiRiSiTiUiViWiXiYiZi0i1i2i3i4i5i6i7i8iKc9i!i#iLc$i%i'i(i)i*i+i,i-i.i/i:i;i=i?i@i[i]i^i_i`i{i|i}i~iajbjcjdjejfjgjhjijjjkjljmjnjojpjqjrjsjtjujvjwjxjyjzjAjBjebCjfbDjgbEjhbFjybGjzbHjlbIjmbJjnbKjAbLjBbMjCbNjDbOjjbPjkbQjobRjEbSjpbTjFbUjGbVjHbWjIbXjqbYjJbZjrb0jsb1jtb2jub3jKb4jLb5jMb6jNb7jOb8jPb9jQb!jRb#jSb$jvb%jwb'jxb(j)j*j+j,j-j.j/j:j;j=j?j@j[j]jbb^j_j`jib{j|j}j~jakbkckdkekfkgkhkikjkkklkmknkokpkqkrksktkukvkwkxkykzkAkBkCkDkEkFkGkHkIkJkKkLkMkNkOkPkQkRkSkTkUkVkWkXkYkZk0k1k2k3k4k5k6k7k8k9k!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~kalblcldlelflglhliljlklllmlnlolplqlrlsltlulvlwlxlyl

206 elif isinstance(error, runtime.cudaError_t): 166 ctx2a 0 1 2 3 4 5 #dvewe$didFc%d'd(dTbUbb 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~bacbcccdcecfcgchcicjckclcmcncocpcqcrcsctcGcHcWdXdYdIc~ 2dcb3ddbKcLcebfbgbhbybzblbmbnbAbBbCbDbjbkbobEbpbFbGbHbIbqbJbrbsbtbubKbLbMbNbObPbQbRbSbvbwbxbbbib

207 return _check_runtime_error(error) 85 ctx20 1 2 3 4 5 #dvewe$didFc%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~bacbcccdcecfcgchcicjckclcmcncocpcqcrcsctcGcHcWdXdYdIc

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

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

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(ek `d)el *em +e;c,en -e=c.eo /ep :eq ;eCcDcEc?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 }eH ~eI afJ bfK cfL dfM efN ffO gfP hfQ ifR jfS kfT lfU mfV nfW ofi j }c~cadbdwcxcaecdddhd0 1 2 3 4 5 #dvepfweqfrfsftfufvf$didwfFcxfyfzfAfBfCfDfEf%dFfGfHfIfbeJfKfLfce'd(dMfdeNfOfPfQfRfSfTfUfVfWfXfYfZfeeX TbUb0f1f2f3fb 6 4f7 5f8 6f9 7f! 8f# 9f$ !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~f)dVbWb*d+d,d-d.d/d:d;dfegejdkdXbldYbmdndZb0bodpdqd1brd2b3b4b5b6b7b8b9b!bsd#b$b%b'btd(b)budvdwdxd*byd+bzdAd,b-bBdCdDd.bEd/b:b;b=b?b@b[b]b^bFd_b`b{b|bGd}b~bHdIdJdKdacLdbcMdNdccdcOdPdQdecRdfcgchcicjckclcmcncSdocpcqcrcTdsctcUdVd=dGcagbgcgHcdgWdegXdfgYdgghgigjgkglgmgngogpgqgrgsgtgugvgwgxgygzgAgBgCgDgEgFgGgHgIgJgKgLgMgNgOgPgQgRgSgTgUgVgWgXgYgZg0g1g2g3g4g5g6g7g8g?d@dZd9gIc!g~ #g$g%g'g(gY )g*gab+g,g-g.g/ghe:g;g=g?gie@g[g]g^g_g`g2dcb{g3ddb|g}g~gahbhchdhehfhghhhihjhkhlhmhnhohphqhrhshthuhvhwhxhyhzhAhBhChDhEhFhGhHhIhJhKhLhMhNhOhPhQhRhShThUhVhWhXhYhZh0h1h2h3h4h5h6h7h8h9h!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~haibicidieifigihiiijikiliminioipiqirisitiuiviwixiyiziAiBiCiDiEiFiGiHiIiJiKiLiMiNiOiPiQiRiSiTiUiViWiXiYiZi0i1i2i3i4i5i6i7i8iKc9i!i#iLc$i%i'i(i)i*i+i,i-i.i/i:i;i=i?i@i[i]i^i_i`i{i|i}i~iajbjcjdjejfjgjhjijjjkjljmjnjojpjqjrjsjtjujvjwjxjyjzjAjBjebCjfbDjgbEjhbFjybGjzbHjlbIjmbJjnbKjAbLjBbMjCbNjDbOjjbPjkbQjobRjEbSjpbTjFbUjGbVjHbWjIbXjqbYjJbZjrb0jsb1jtb2jub3jKb4jLb5jMb6jNb7jOb8jPb9jQb!jRb#jSb$jvb%jwb'jxb(j)j*j+j,j-j.j/j:j;j=j?j@j[j]jbb^j_j`jib{j|j}j~jakbkckdkekfkgkhkikjkkklkmknkokpkqkrksktkukvkwkxkykzkAkBkCkDkEkFkGkHkIkJkKkLkMkNkOkPkQkRkSkTkUkVkWkXkYkZk0k1k2k3k4k5k6k7k8k9k!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~kalblcldlelflglhliljlklllmlnlolplqlrlsltlulvlwlxlyl

216 cdef int out_len = len(result) 905 ctx2a SeTeUeVeWeXeYeZe0e1e2e3e4e5e6e7e8e9e!e#e$e%e'e(ek `d)el *em +e;c,en -e=c.eo /ep :eq ;eCcDcEc?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 }eH ~eI afJ bfK cfL dfM efN ffO gfP hfQ ifR jfS kfT lfU mfV nfW ofi j }c~cadbdwcxcaecdddhd0 1 2 3 4 5 #dvepfweqfrfsftfufvf$didwfFcxfyfzfAfBfCfDfEf%dFfGfHfIfbeJfKfLfce'd(dMfdeNfOfPfQfRfSfTfUfVfWfXfYfZfeeX TbUb0f1f2f3fb 6 4f7 5f8 6f9 7f! 8f# 9f$ !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~f)dVbWb*d+d,d-d.d/d:d;dfegejdkdXbldYbmdndZb0bodpdqd1brd2b3b4b5b6b7b8b9b!bsd#b$b%b'btd(b)budvdwdxd*byd+bzdAd,b-bBdCdDd.bEd/b:b;b=b?b@b[b]b^bFd_b`b{b|bGd}b~bHdIdJdKdacLdbcMdNdccdcOdPdQdecRdfcgchcicjckclcmcncSdocpcqcrcTdsctcUdVd=dGcagbgcgHcdgWdegXdfgYdgghgigjgkglgmgngogpgqgrgsgtgugvgwgxgygzgAgBgCgDgEgFgGgHgIgJgKgLgMgNgOgPgQgRgSgTgUgVgWgXgYgZg0g1g2g3g4g5g6g7g8g?d@dZd9gIc!g~ #g$g%g'g(gY )g*gab+g,g-g.g/ghe:g;g=g?gie@g[g]g^g_g`g2dcb{g3ddb|g}g~gahbhchdhehfhghhhihjhkhlhmhnhohphqhrhshthuhvhwhxhyhzhAhBhChDhEhFhGhHhIhJhKhLhMhNhOhPhQhRhShThUhVhWhXhYhZh0h1h2h3h4h5h6h7h8h9h!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~haibicidieifigihiiijikiliminioipiqirisitiuiviwixiyiziAiBiCiDiEiFiGiHiIiJiKiLiMiNiOiPiQiRiSiTiUiViWiXiYiZi0i1i2i3i4i5i6i7i8iKc9i!i#iLc$i%i'i(i)i*i+i,i-i.i/i:i;i=i?i@i[i]i^i_i`i{i|i}i~iajbjcjdjejfjgjhjijjjkjljmjnjojpjqjrjsjtjujvjwjxjyjzjAjBjebCjfbDjgbEjhbFjybGjzbHjlbIjmbJjnbKjAbLjBbMjCbNjDbOjjbPjkbQjobRjEbSjpbTjFbUjGbVjHbWjIbXjqbYjJbZjrb0jsb1jtb2jub3jKb4jLb5jMb6jNb7jOb8jPb9jQb!jRb#jSb$jvb%jwb'jxb(j)j*j+j,j-j.j/j:j;j=j?j@j[j]jbb^j_j`jib{j|j}j~jakbkckdkekfkgkhkikjkkklkmknkokpkqkrksktkukvkwkxkykzkAkBkCkDkEkFkGkHkIkJkKkLkMkNkOkPkQkRkSkTkUkVkWkXkYkZk0k1k2k3k4k5k6k7k8k9k!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~kalblcldlelflglhliljlklllmlnlolplqlrlsltlulvlwlxlyl

217 if out_len == 1: 905 ctx2a SeTeUeVeWeXeYeZe0e1e2e3e4e5e6e7e8e9e!e#e$e%e'e(ek `d)el *em +e;c,en -e=c.eo /ep :eq ;eCcDcEc?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 }eH ~eI afJ bfK cfL dfM efN ffO gfP hfQ ifR jfS kfT lfU mfV nfW ofi j }c~cadbdwcxcaecdddhd0 1 2 3 4 5 #dvepfweqfrfsftfufvf$didwfFcxfyfzfAfBfCfDfEf%dFfGfHfIfbeJfKfLfce'd(dMfdeNfOfPfQfRfSfTfUfVfWfXfYfZfeeX TbUb0f1f2f3fb 6 4f7 5f8 6f9 7f! 8f# 9f$ !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~f)dVbWb*d+d,d-d.d/d:d;dfegejdkdXbldYbmdndZb0bodpdqd1brd2b3b4b5b6b7b8b9b!bsd#b$b%b'btd(b)budvdwdxd*byd+bzdAd,b-bBdCdDd.bEd/b:b;b=b?b@b[b]b^bFd_b`b{b|bGd}b~bHdIdJdKdacLdbcMdNdccdcOdPdQdecRdfcgchcicjckclcmcncSdocpcqcrcTdsctcUdVd=dGcagbgcgHcdgWdegXdfgYdgghgigjgkglgmgngogpgqgrgsgtgugvgwgxgygzgAgBgCgDgEgFgGgHgIgJgKgLgMgNgOgPgQgRgSgTgUgVgWgXgYgZg0g1g2g3g4g5g6g7g8g?d@dZd9gIc!g~ #g$g%g'g(gY )g*gab+g,g-g.g/ghe:g;g=g?gie@g[g]g^g_g`g2dcb{g3ddb|g}g~gahbhchdhehfhghhhihjhkhlhmhnhohphqhrhshthuhvhwhxhyhzhAhBhChDhEhFhGhHhIhJhKhLhMhNhOhPhQhRhShThUhVhWhXhYhZh0h1h2h3h4h5h6h7h8h9h!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~haibicidieifigihiiijikiliminioipiqirisitiuiviwixiyiziAiBiCiDiEiFiGiHiIiJiKiLiMiNiOiPiQiRiSiTiUiViWiXiYiZi0i1i2i3i4i5i6i7i8iKc9i!i#iLc$i%i'i(i)i*i+i,i-i.i/i:i;i=i?i@i[i]i^i_i`i{i|i}i~iajbjcjdjejfjgjhjijjjkjljmjnjojpjqjrjsjtjujvjwjxjyjzjAjBjebCjfbDjgbEjhbFjybGjzbHjlbIjmbJjnbKjAbLjBbMjCbNjDbOjjbPjkbQjobRjEbSjpbTjFbUjGbVjHbWjIbXjqbYjJbZjrb0jsb1jtb2jub3jKb4jLb5jMb6jNb7jOb8jPb9jQb!jRb#jSb$jvb%jwb'jxb(j)j*j+j,j-j.j/j:j;j=j?j@j[j]jbb^j_j`jib{j|j}j~jakbkckdkekfkgkhkikjkkklkmknkokpkqkrksktkukvkwkxkykzkAkBkCkDkEkFkGkHkIkJkKkLkMkNkOkPkQkRkSkTkUkVkWkXkYkZk0k1k2k3k4k5k6k7k8k9k!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~kalblcldlelflglhliljlklllmlnlolplqlrlsltlulvlwlxlyl

218 return 197 ctx2k `dl m ;cn =co p q CcDcEc?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 aehd0 1 2 3 4 5 idFceeX TbUb)dVbWb*d+d,d-d.d/d:d;dfegejdkdXbldYbmdndZb0bodpdqd1brd2b3b4b5b6b7b8b9b!bsd#b$b%b'btd(b)budvdwdxd*byd+bzdAd,b-bBdCdDd.bEd/b:b;b=b?b@b[b]b^bFd_b`b{b|bGd}b~bHdIdJdKdacLdbcMdNdccdcOdPdQdecRdfcgchcicjckclcmcncSdocpcqcrcTdsctcUdVd=dGcHcWdXdYd?d@dZdIcY

219 elif out_len == 2: 

220 return result[1] 891 ctx2a SeTeUeVeWeXeYeZe0e1e2e3e4e5e6e7e8e9e!e#e$e%e'e(ek )el *em +e;c,en -e=c.eo /ep :eq ;eCcDcEc?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 }eH ~eI afJ bfK cfL dfM efN ffO gfP hfQ ifR jfS kfT lfU mfV nfW ofi j }c~cadbdwcxccdddhd0 1 2 3 4 5 #dvepfweqfrfsftfufvf$dwfxfyfzfAfBfCfDfEf%dFfGfHfIfbeJfKfLfce'd(dMfdeNfOfPfQfRfSfTfUfVfWfXfYfZfX TbUb0f1f2f3fb 6 4f7 5f8 6f9 7f! 8f# 9f$ !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~f)dVbWb*d+d,d-d.d/d:d;djdkdXbldYbmdndZb0bodpdqd1brd2b3b4b5b6b7b8b9b!bsd#b$b%b'btd(b)budvdwdxd*byd+bzdAd,b-bBdCdDd.bEd/b:b;b=b?b@b[b]b^bFd_b`b{b|bGd}b~bHdIdJdKdacLdbcMdNdccdcOdPdQdecRdfcgchcicjckclcmcncSdocpcqcrcTdsctcUdVd=dagbgcgdgegfggghgigjgkglgmgngogpgqgrgsgtgugvgwgxgygzgAgBgCgDgEgFgGgHgIgJgKgLgMgNgOgPgQgRgSgTgUgVgWgXgYgZg0g1g2g3g4g5g6g7g8g?d@dZd9g!g~ #g$g%g'g(g)g*gab+g,g-g.g/ghe:g;g=g?gie@g[g]g^g_g`g2dcb{g3ddb|g}g~gahbhchdhehfhghhhihjhkhlhmhnhohphqhrhshthuhvhwhxhyhzhAhBhChDhEhFhGhHhIhJhKhLhMhNhOhPhQhRhShThUhVhWhXhYhZh0h1h2h3h4h5h6h7h8h9h!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~haibicidieifigihiiijikiliminioipiqirisitiuiviwixiyiziAiBiCiDiEiFiGiHiIiJiKiLiMiNiOiPiQiRiSiTiUiViWiXiYiZi0i1i2i3i4i5i6i7i8iKc9i!i#iLc$i%i'i(i)i*i+i,i-i.i/i:i;i=i?i@i[i]i^i_i`i{i|i}i~iajbjcjdjejfjgjhjijjjkjljmjnjojpjqjrjsjtjujvjwjxjyjzjAjBjebCjfbDjgbEjhbFjybGjzbHjlbIjmbJjnbKjAbLjBbMjCbNjDbOjjbPjkbQjobRjEbSjpbTjFbUjGbVjHbWjIbXjqbYjJbZjrb0jsb1jtb2jub3jKb4jLb5jMb6jNb7jOb8jPb9jQb!jRb#jSb$jvb%jwb'jxb(j)j*j+j,j-j.j/j:j;j=j?j@j[j]jbb^j_j`jib{j|j}j~jakbkckdkekfkgkhkikjkkklkmknkokpkqkrksktkukvkwkxkykzkAkBkCkDkEkFkGkHkIkJkKkLkMkNkOkPkQkRkSkTkUkVkWkXkYkZk0k1k2k3k4k5k6k7k8k9k!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~kalblcldlelflglhliljlklllmlnlolplqlrlsltlulvlwlxlyl

221 else: 

222 return result[1:] 138 ctx2a k l m ;cn =co p q CcDcEc?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 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ 2dcb3ddbKcLcebfbgbhbybzblbmbnbAbBbCbDbjbkbobEbpbFbGbHbIbqbJbrbsbtbubKbLbMbNbObPbQbRbSbvbwbxbbbib

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: 746 ctx2k l m ;cn =co p q CcDcEc?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 1n2nBl3nCl4n5n6nDl7nEl8ni 9nj !n{d|dpeqe}d~dreseKeLeFl#n$n%n'n(n)n*n+n}c,n~c-nad.nbd/nwcxccd:ndd;n=n?n@nhd0 1 2 3 4 5 [n]n^n_n`n{nidFcrq|nGlucHlsq}ntq~naobocodoeofoIlJlKlgdgoX TbUbBcvcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c9c!c#c$c%c'c(c)c*cb 6 Np7 Op8 Pp9 Qp! Rp# Sp$ Tp% Up' Vp( Wp) Wn* Xn+ Xp, Yp- Zp. 0p/ 1p: 2p; 3p= 4p? 5p@ 6p[ 7p] 8p^ 9p_ 8d` !p{ #p| $p} %pLlVbWbMlNljdkdXbldYbmdndZb0bodpdqd1brd2b3b4b5b6b7b8b9b!bsd#b$b%b'btd(b)budvdwdxd*byd+bzdAd,b-bBdCdDd.bEd/b:b;b=b?b@b[b]b^bFd_b`b{b|bGd}b~bHdIdJdKdacLdbcMdNdccdcOdPdQdecRdfcgchcicjckclcmcncSdocpcqcrcTdsctcUdVdhoioGcjoHcWdXdYdkolomonooopoqorosotouovowoxoyozoAoBoCoDoEoFoGoHoIoJoKoLoMoNoOoPoQoRoSoToUoVoWoXoQnRnSnTnUnVnYoZdZoIc0o1oyczcZ Ac~ OlPlQlRlY SlabTlUlVlWlXlYlZl0l1l2l3l4l5l6l7lcbdb8l9l!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~lambmcmdmemfmgmhmimjmkmlmmmnmompmqmrmsmtmumvmwmxmymzmAmBmCmDmEmFmGmHmImJm9d0d1dJcedKm2oLmMm3oNmOm4oPm5oQmRm6oje+ckeSmTmUmVm7oWmXm8oMe9oYm4dZm,c-c0mNe!o#o1m5d2m3m$o4m5m%ole.cme6m7m8m9m'o!m#m(oOe)o$m6d%m/c:c'mPe*o+o(m7d)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{m|m[o}m]o~man^obn_ocndnen`ofn{ognhnin|oxpypebfb}ogbhbybzblbmbnbAbBbCbDbjbkbobEbpbFbGbHbIbqbJbrbsbtbubKbLbMbNbObPbQbRbSbvbwbxbjnfdzpApBpCpDpEpFpGpHpIpJpKpLpMpQeReknlnmnbbnn~oapon'p(p)pib*p+p,p-p.p/pbpcpdpepfpgphpipjpkplpmpnpopppqppnrpsptpupvpwpqnrnsntnunvnwnxnynznAnBnCnDnEnFnGnHnInJnKnLnMnNn

230 if keep_none: 478 ctx2k l m ;cn =co p q CcDcEc?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|dpeqe}d~dreseKeLeFl}c~cadbdwcxccddd0 1 2 3 4 5 [n]n^n_n`n{nidFcrq|nGlucHl}n~naobocodoeofoIlJlKlgdgoX TbUbvcVbWbMlNljdkdXbldYbmdndZb0bodpdqd1brd2b3b4b5b6b7b8b9b!bsd#b$b%b'btd(b)budvdwdxd*byd+bzdAd,b-bBdCdDd.bEd/b:b;b=b?b@b[b]b^bFd_b`b{b|bGd}b~bHdIdJdKdacLdbcMdNdccdcOdPdQdecRdfcgchcicjckclcmcncSdocpcqcrcTdsctcUdVdhoioGcjoHcWdYdZdZoIcOlPlQlRlY SlabTlUlVlWlXlYlZl0l1l2l3l4l5lcbdb8l9l!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~lambmcmdmemfmgmhmimjmkmlmmmnmompmqmrmsmtmumvmwmxmymzmAmBmCmDmEmFmGmHmImJmJcKm2oLmMm3oNmOm4oPm5oQmRm6oje+ckeSmTmUmVm7oWmXm8oMe9oYm4dZm,c-c0mNe!o#o1m5d2m3m$o4m5m%ole.cme6m7m8m9m'o!m#m(oOe)o$m6d%m/c:c'mPe*o+o(m7d)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{m|m[o}m]o~man^obn_ocndnen`ofn{ognhnin|ofbjnfdknlnmnnn~oaponibdpepfpgpjpkplpmppppnrpupvpqnrnsntnunvnwnxnynznAnBnCnDnEnFnGnHnInJnKnLnMnNn

231 return options 72 ctx2CcDcEcc r d e f s g h idFcMlNlhoiojoWdYdZoY KmNmPmje+ckeSmVmMeNe3mle.cme6m9mOePe*m:m?m]m`m}mbnfnqnrnsntnunvnwnxnynznAnBnCnDnEnFnGnHnInJnKnLnMnNn

232 return cls() 426 ctx2k l m ;cn =co p q CcDcEc?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|dpeqe}d~dreseKeLeFl}c~cadbdwcxccddd0 1 2 3 4 5 [n]n^n_n`n{nrq|nGlucHl}n~naobocodoeofoIlJlKlgdgoX TbUbvcVbWbjdkdXbldYbmdndZb0bodpdqd1brd2b3b4b5b6b7b8b9b!bsd#b$b%b'btd(b)budvdwdxd*byd+bzdAd,b-bBdCdDd.bEd/b:b;b=b?b@b[b]b^bFd_b`b{b|bGd}b~bHdIdJdKdacLdbcMdNdccdcOdPdQdecRdfcgchcicjckclcmcncSdocpcqcrcTdsctcUdVdGcHcZdIcOlPlQlRlY SlabTlUlVlWlXlYlZl0l1l2l3l4l5lcbdb8l9l!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~lambmcmdmemfmgmhmimjmkmlmmmnmompmqmrmsmtmumvmwmxmymzmAmBmCmDmEmFmGmHmImJmJc2oLmMm3oOm4o5oQmRm6ojekeTmUm7oWmXm8oMe9oYm4dZm,c-c0mNe!o#o1m5d2m$o4m5m%oleme7m8m'o!m#m(oOe)o$m6d%m/c:c'mPe*o+o(m7d)m,o+m,m-m.m/m-o.o;m=m/o:o@m[m;o=o^m_m?o@o{m|m[o]o~man^o_ocndnen`o{ognhnin|ofbjnfdknlnmnnn~oaponibdpepfpgpjpkplpmppppnrpupvp

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 1n2nBl3nCl4n5n6nDl7nEl8ni 9nj !n{d|dpeqe}d~dreseKeLeFl#n$n%n'n(n)n*n+n,n-n.n/nwcxc:n;n=n?n@nhd0 1 2 3 4 5 GlucHlsqtqIlJlKlgdX BcvcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c9c!c#c$c%c'c(c)c*cb 6 Np7 Op8 Pp9 Qp! Rp# Sp$ Tp% Up' Vp( Wp) Wn* Xn+ Xp, Yp- Zp. 0p/ 1p: 2p; 3p= 4p? 5p@ 6p[ 7p] 8p^ 9p_ 8d` !p{ #p| $p} %pLlGcHcXdkolomonooopoqorosotouovowoxoyozoAoBoCoDoEoFoGoHoIoJoKoLoMoNoOoPoQoRoSoToUoVoWoXoQnRnSnTnUnVnYoIc0o1oyczcZ Ac~ Y 6l7l9d0d1dJcedxpypeb}ogbhbybzblbmbnbAbBbCbDbjbkbobEbpbFbGbHbIbqbJbrbsbtbubKbLbMbNbObPbQbRbSbvbwbxbzpApBpCpDpEpFpGpHpIpJpKpLpMpQeRebb'p(p)p*p+p,p-p.p/pbpcphpipnpopqppnsptpwp

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 1n2nBl3nCl4n5n6nDl7nEl8ni 9nj !npeqereseFl#n$n%n'n(n)n*n+n,n-n.n/nwcxc:n;n=n?n@nhd0 1 2 3 4 5 GlucHlsqtqIlJlKlgdX BcvcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c9c!c#c$c%c'c(c)c*cb 6 Np7 Op8 Pp9 Qp! Rp# Sp$ Tp% Up' Vp( Wp) Wn* Xn+ Xp, Yp- Zp. 0p/ 1p: 2p; 3p= 4p? 5p@ 6p[ 7p] 8p^ 9p_ 8d` !p{ #p| $p} %pLlGcHcXdkolomonooopoqorosotouovowoxoyozoAoBoCoDoEoFoGoHoIoJoKoLoMoNoOoPoQoRoSoToUoVoWoXoYoIc0o1oyczcZ Ac~ Y 6l7l9d0d1dJcedxpypeb}ogbhbybzblbmbnbAbBbCbDbjbkbobEbpbFbGbHbIbqbJbrbsbtbubKbLbMbNbObPbQbRbSbvbwbxbzpApBpCpDpEpFpGpHpIpJpKpLpMpQeRe'p(p)p*p+p,p-p.p/pbpcphpipnpopqppnsptpwp

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

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

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 6 7 8 9 ! # $ % ' ( ) * + , - . / : ; = ? @ [ ] ^ _ ` { | } ~ lbmbnbpbtbubvbwbxbBy

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 BcMcNcOcPcQcRcScTcUcVcWcXcYcZc0c1c2c3c4c5c6c7c8c9c!c#c$c%c'c(c)c*cWnXnjbkbobqbrbsbQeRe

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 ctx2yczcZ Ac

328 self._entered = False 4 ctx2yczcZ Ac

329  

330 def __enter__(self): (empty)

331 self._stack.__enter__() 4 ctx2yczcZ Ac

332 self._entered = True 4 ctx2yczcZ Ac

333 return self 4 ctx2yczcZ Ac

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 ctx2yczcZ Ac

338 return self._stack.__exit__(exc_type, exc, tb) 4 ctx2yczcZ Ac

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 ctx2yczcZ Ac

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

347 self._stack.callback(partial(fn, *args, **kwargs)) 4 ctx2yczcZ Ac

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 ctx2yczcZ Ac

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 ctx29d0d1dJced

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;pBlCl=p?pDlEli j {d|dpeqe}d~drese@p[p]p^p_p`p{p|p}c~cadbdwcxccddd}p~paqte0 1 2 3 4 5 9d0d1dJced

375 return 45 ctx2:p;pBlCl=p?pDlEli j {d|dpeqe}d~drese@p[p]p^p_p`p{p|p}c~cadbdwcxccddd}p~paqte0 1 2 3 4 5 ed

376 _fork_warning_checked = True 6 ctx2te9d0d1dJced

377  

378 # Common warning message parts 

379 common_message = ( 

380 "CUDA does not support. Forked subprocesses exhibit undefined behavior, " 6 ctx2te9d0d1dJced

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 ctx2te9d0d1dJced

387 start_method = multiprocessing.get_start_method() 6 ctx2te9d0d1dJced

388 if start_method == "fork": 6 ctx2te9d0d1dJced

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

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

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)