shVSrSSKrSSKrSSKrSSKrSSKrSSKrSSKrSSKrSSK r SSK r SSK J r SSK JrJr SSKJrJrJrJrJrJrJrJrJr SSKrSSKrSSKrSSKJrJr SSK J!r! SSK"J#r#J$r$J%r% S S K&J'r'J(r( S S K)J*r* S S K)J+r+J,r,J-r- SS K.J/r/J)r)J0r0J1r1 \%Rdr3\3(aSSK4J4r5 Sr6Sr7\RpRsS5r:\RvRxr=\RvR|r?\R/\=lA\R/\?lASrB"SS\C5rD"SS\C5rE"SS\(5rFSrGSrHSrISrJSrKSrLSrMSrNS rO\N"5rP"S!S"\Q5rR"S#S$\Q5rS\R"5rTS%rU\U"5rV"S&S'\Q5rWS(rX"S)S*\Q\S+9rY"S,S-\Y5rZ"S.S/5r["S0S1\[\Z5r\S r]Sq^S2r_S3r`"S4S5\a5rb\b"5rb"S6S7\Q5rc\"S8S95rd"S:S;\Q5reS<rfS=rgS>rhS?riS@rjSArkSBrlSCrmSDrn"SESF\Q5ro"SGSH\Q5rp"SISJ\Q5rq"SKSL\Q5rr"SMSN\r5rs"SOSP\s5rt"SQSR\!R5rv"SSST\s5rw"SUSV\Q5rx"SWSX\x\!R5ry"SYSZ\x\!R5rz"S[S\\Q5r{"S]S^\Q5r|S_r}"S`Sa\S+9r~"SbSc\~5r"SdSe\~5r\"Sf/SgQ5r"ShSi\S+9r"SjSk\5r"SlSm\5rSSnjr\3(aY\5GR r\GR\GR\GR\GR\GR\GRSo.rOK\/GR\/GR\/GR\/GR\/GR\/GRSo.r"SpSq\S+9rSrr"SsSt\5r"SuSv\5r"SwSx\5rSyrSzrS{rS|rS}rSS~jrSrSrSrSrSrSrSrSrSSjrSSjrSSjrSSjrSrSr\GRNS5rSrg)a CUDA driver bridge implementation NOTE: The new driver implementation uses a *_PendingDeallocs* that help prevents a crashing the system (particularly OSX) when the CUDA context is corrupted at resource deallocation. The old approach ties resource management directly into the object destructor; thus, at corruption of the CUDA context, subsequent deallocation could further corrupt the CUDA context and causes the system to freeze in some cases. N)product)ABCMetaabstractmethod) c_intbyrefc_size_tc_charc_char_p addressofc_void_pc_floatc_uint) namedtupledeque)mviewbuf)utils serializeconfig)CudaSupportErrorCudaDriverError)API_PROTOTYPES)cu_occupancy_b2d_sizecu_stream_callback_pyobjcu_uuid)enumsdrvapinvrtc_extrascuda)linuxcl[R"[5nUR5(Gd[ [ R 5R5n[[US5n[U[5(d[RnURU5 [ R (a\[R"[R5nSnUR![R""US95 UR%U5 U$UR%[R&"55 U$)Nz;== CUDA [%(relativeCreated)d] %(levelname)5s -- %(message)s)fmt)logging getLogger__name__ hasHandlersstrrCUDA_LOG_LEVELuppergetattr isinstanceintCRITICALsetLevel StreamHandlersysstderr setFormatter Formatter addHandler NullHandler)loggerlvlhandlerr&s m/Users/tiagomarins/Projetos/claudeai/copy_bank/venv/lib/python3.13/site-packages/numba/cuda/cudadrv/driver.py make_loggerr>;s   x (F     &''(..0gsD)#s##""C  ++CJJ7GOC  !2!2s!; <   g & M   g113 4 Mc\rSrSrSrg)DeadMemoryErrorSNr) __module__ __qualname____firstlineno____static_attributes__rCr?r=rArASr?rAc\rSrSrSrg) LinkerErrorWrCNrDrCr?r=rKrKWrIr?rKc.^\rSrSrU4SjrSrSrU=r$) CudaAPIError[cD>XlX l[[U]X5 gN)codemsgsuperrN__init__)selfrRrS __class__s r=rUCudaAPIError.__init__\s  lD*45r?c>SUR<SUR<3$)N[z] rRrSrVs r=__str__CudaAPIError.__str__as IItxx00r?r[)r)rErFrGrUr]rH __classcell__rWs@r=rNrN[s6 11r?rNc [RnUS:Xa [5 [RS:Xa[ R nS/nS/nOC[RS:Xa[ RnS/nS/nO[ RnSS /nS S /nU(aX[RRU5n[RRU5(d[S U-5eU/nX4$U[X#5VVs/sH$upV[RRXV5PM& snn-nX4$![a [S U-5ef=fs snnf)N0win32z\windows\system32z nvcuda.dlldarwinz/usr/local/cuda/libz libcuda.dylibz/usr/libz /usr/lib64z libcuda.soz libcuda.so.1z(NUMBA_CUDA_DRIVER %s is not a valid pathzoNUMBA_CUDA_DRIVER %s is not a valid file path. Note it must be a filepath of the .so/.dll/.dylib or the driver)r CUDA_DRIVER_raise_driver_not_foundr4platformctypesWinDLLCDLLospathabspath ValueErrorisfilerjoin)envpathdlloaderdldirdlnames candidatesxys r=locate_driver_and_loaderrxes[  G#~! ||w==&'.  !;;&'"#;;\*0 &ggoog.Gww~~g&&9;BCD DY   ,3E,C E,CDA!# Q 2,C EE    &G$%& & & EsD1=+E 1E cj/n/nUHnU"U5nXT4s $ [ U5(a [ 5 gSRSU55n[U5 g![aOnUR[RR U5(+5 URU5 SnAMSnAff=f)N c38# UHn[U5v M g7frQ)r+).0es r= load_driver..s=+!=+<==F# (  ! !bggnnT&:": ;  $ $Q ' ' (sA B2#AB--B2c8[5up[X5up#U$rQ)rxr)rrrurrls r= find_driverrs35HH1IC Jr?z CUDA driver library cannot be found. If you are sure that a CUDA driver is installed, try setting environment variable NUMBA_CUDA_DRIVER with the file path of the CUDA driver shared library. zM Possible CUDA driver libraries are found but error occurred during load: %s c [[5erQ)rDRIVER_NOT_FOUND_MSGrCr?r=rfrfs / 00r?c&[[U-5erQ)rDRIVER_LOAD_ERROR_MSG)r}s r=rrs 014 55r?cSn[R"5n[[5H/nUR U5(dM[ [U5nX!U'M1 U$)N CUDA_ERROR)r UniqueDictdirr startswithr.)prefixmapnamerRs r=_build_reverse_error_maprsL F    CE  ??6 " "5$'DI Jr?c,[R"5$rQ)rkgetpidrCr?r=_getpidrs 99;r?c\rSrSrSrSrSrSrSrSr \ S5r S r SS jr S rS rS rSrSrSSjrSrSrSrSrSrSrSrg)Driverz( Driver API functions are lazily bound. Nc^URnUbU$[RU5nXlU$rQ) _singletonobject__new__)clsobjs r=rDriver.__new__s-nn ?J..%C N r?c$[R"5UlSUlSUlSUl[ R(a Sn[U5e[5Ul g![a"nSUlURUlSnAgSnAff=f)NFzzCUDA is disabled due to setting NUMBA_DISABLE_CUDA=1 in the environment, or because CUDA is unsupported on 32-bit systems.T) rrdevicesis_initializedinitialization_errorpidr DISABLE_CUDArrlibrS)rVrSr}s r=rUDriver.__init__s{'') #$(! ."")'s++"}DH ."&D ()D % % .s1A## B-B  BcZUR(ag[5qSUl[RS5 UR S5 [ 5UlUR5 g![a5nURSURS3nX l [SU35eSnAff=f)NTinitrz ()zError at driver init: ) rr>_loggerinfocuInitrrrNrSrRrr_initialize_extras)rVr} descriptions r=ensure_initializedDriver.ensure_initializeds    -" ! LL KKN yDH ! KUUG2affXQ/K(3 %"%;K=#IJ J Ks&A++ B*50B%%B*c[(ag[R"S[5nU"[R 5nU"UR S55 [R"[[R"[R5[R"[R5[R5nU"[R5nSUlURSU5nXPlg)NcuIpcOpenMemHandlecall_cuIpcOpenMemHandle)USE_NV_BINDINGrh CFUNCTYPEr rset_cuIpcOpenMemHandle _find_apirPOINTERr cu_device_ptrcu_ipc_mem_handlerrr)_ctypes_wrap_fnr)rV set_protor call_protor safe_calls r=rDriver._initialize_extrass > $$T84 !*7+I+I!Jt~~.BCD%%e&,nnV5I5I&J&,nnV5M5M&N&,mm5 #-W-L-L"M+D((()B)@B #,r?c>UR5 URSL$rQ)rrr\s r= is_availableDriver.is_availables !((D00r?cUR5 URb[SUR-5e[(aUR U5$UR U5$)NzError at driver init: %s:)rrrr_cuda_python_wrap_fnrrVfnames r= __getattr__Driver.__getattr__"sa !  $ $ 0"#@#'#<#<$=> > >,,U3 3''. .r?cP^^^Tc3[TnUSnUSSnTRT5mUTlUTlUUU4SjnUUU4Sjn[ R(aUnOUn[R"T5"U5n [TTU 5 U $![a [T5ef=f)Nrrc>SRUVs/sHn[U5PM sn5n[RSTRU5 T"U6nTR TU5 gs snfNz, zcall driver api: %s(%s))rpr+rdebugr)_check_ctypes_error)argsargargstrretcoderlibfnrVs r=verbose_cuda_api_call5Driver._ctypes_wrap_fn..verbose_cuda_api_call>sWYYD9DSCD9:F MM3U^^V LTlG  $ $UG 4 :sA&ct>[RSTR5 T"U6nTRTU5 gNzcall driver api: %s)rrr)r)rrrrrVs r=safe_cuda_api_call2Driver._ctypes_wrap_fn..safe_cuda_api_callDs/ MM/ @TlG  $ $UG 4r?) rKeyErrorAttributeErrorrrestypeargtypesrCUDA_LOG_API_ARGS functoolswrapssetattr) rVrrprotorrrrwrapperrs ``` r=rDriver._ctypes_wrap_fn/s = ,&u-AhGQRyHNN5)E#EM%EN 5  5  # #+G(GOOE*73 eY'9 ,$U++ ,s BB%c^^^[[T5mUUU4SjnUUU4Sjn[R(aUnOUn[R "T5"U5n[ TTU5 U$)Nc>SRUVs/sHn[U5PM sn5n[RSTRU5 TR TT"U65$s snfr)rpr+rrr)_check_cuda_python_error)rrrrrrVs r=r:Driver._cuda_python_wrap_fn..verbose_cuda_api_callUsUYYD9DSCD9:F MM3U^^V L00t E E :sA#cn>[RSTR5 TRTT"U65$r)rrr)r)rrrrVs r=r7Driver._cuda_python_wrap_fn..safe_cuda_api_callZs- MM/ @00t E Er?)r.bindingrrrrr)rVrrrrrrs`` @r=rDriver._cuda_python_wrap_fnRsU' F  F  # #+G(GOOE*73 eY'r?c^[R(a[(dSnOSnUHn[URTU35s $ U4Sjn[ UTU5 U$![ a MDf=f)N)_v2_ptds_v2_ptsz_ptds_ptsz_v2)rrc >[ST35e)NzDriver missing function: )r)rkwsrs r=absent_function)Driver._find_api..absent_functionxs!$=eW"EF Fr?)rCUDA_PER_THREAD_DEFAULT_STREAMrr.rrr)rVrvariantsvariantrs ` r=rDriver._find_apigss  0 0LH"HG txxE77))<==  G e_-"  sA  A.-A.cURbO[5UR:wa6Sn[RU[5UR5 [ S5egg)Nz0pid %s forked from pid %s after CUDA driver initzCUDA initialized before forking)rrrcriticalr)rVrSs r= _detect_forkDriver._detect_fork~sJ 88 GI$9DC   S')TXX 6!"CD D%: r?cU[R:wad[RUS5nSU<SU<3n[R U5 U[R :XaUR5 [X$5eg)NUNKNOWN_CUDA_ERRORCall to results in ) r CUDA_SUCCESS ERROR_MAPgetrerrorCUDA_ERROR_NOT_INITIALIZEDrrN)rVrrerrnamerSs r=rDriver._check_ctypes_errorsa e(( (mmG-ABG05w?C MM# %:::!!#w, , )r?cBUSnUSSn[U5S:XaUSnU[RR:wabSU<SUR<3n[ R U5 U[RR:XaUR5 [X55eU$)Nrrrr) lenrCUresultr rrr r rrN)rVrreturnedrretvalrSs r=rDriver._check_cuda_python_errors1+!" v;! AYF g&&33 305w||DC MM# '**EEE!!#w, , r?cURRU5nUc[U5nX RU'[R"U5$rQ)rr Deviceweakrefproxy)rVdevnumdevs r= get_deviceDriver.get_devices?llv& ;.C#&LL }}S!!r?c[(aUR5$[5nUR[U55 UR$rQ)rcuDeviceGetCountrrvalue)rVcounts r=get_device_countDriver.get_device_counts9 >((* * eEl+{{r?cH[URR55$)z)Returns a list of active devices )listrvaluesr\s r= list_devicesDriver.list_devicessDLL'')**r?cfURR5HnUR5 M g)zReset all devices N)rr&reset)rVrs r=r* Driver.resets$<<&&(C IIK)r?c@UR5nURbf[(a[R 5sSSS5 $[ R "5n[R [U55 UsSSS5 $SSS5 g!,(df  g=f)z^Pop the active CUDA context and return the handle. If no CUDA context is active, return None. N)get_active_contextrrdrivercuCtxPopCurrentr cu_contextr)rVacpoppeds r=pop_active_contextDriver.pop_active_contextsq $ $ &"yy$!>!113' & $..0F**5=9!' &$' & &s,B4B Bc[5$)z3Returns an instance of ``_ActiveContext``. )_ActiveContextr\s r=r-Driver.get_active_contexts r?c[(a[R5nOK[R"S5n[R[R "U55 UR nUS-nXS-- S-nX44$)z= Returns the CUDA Runtime version as a tuple (major, minor). ri )rr.cuDriverGetVersionrhrrr )rVversiondvmajorminors r= get_versionDriver.get_versionsh >//1GaB  % %fll2&6 7hhG4DL)b0~r?)rrrrrrrQr)r)rErFrG__doc__rrrUrrpropertyrrrrrrrrrr"r'r*r3r-r?rHrCr?r=rrs~J. "*,*11 /!F*.E - "+  " r?rcP\rSrSrSr\R "5rSrSr Sr \ r Sr g)r6izAn contextmanager object to cache active context to reduce dependency on querying the CUDA driver API. Once entering the context, it is assumed that the active CUDA context is not changed until the context is exited. cSn[URS5(aURRup#O[(a&[R 5n[ U5S:XaSnOI[R"S5n[R [U55 UR(aUOSnUcSnO|[(a[ [R55nO?[R"5n[R[U55 URnX#4URlSnXl X lX0lU$)NF ctx_devnumrT)hasattr _tls_cacherFrr.cuCtxGetCurrentr0rr0rr cuCtxGetDevice cu_device_is_topcontext_handler)rVis_tophctxrhdevices r= __enter___ActiveContext.__enter__s 4??L 1 1??55LD&~--/t9>D((+&&uT{3#zztt|!> !6!6!89F$..0G))%.9$]]F.2^* "  r?cTUR(a[URS5 gg)NrF)rLdelattrrH)rVexc_typeexc_valexc_tbs r=__exit___ActiveContext.__exit__s << DOO\ 2 r?cURSL$)zAReturns True is there's a valid and active CUDA context. N)rMr\s r=__bool___ActiveContext.__bool__s""$..r?)rLrMrN) r)rErFrGrB threadinglocalrHrQrXr[ __nonzero__rHrCr?r=r6r6s, "J D3/ Kr?r6cSn[R"5n[[5H9nUR U5(dM[ [U5X[ U5S'M; U$)NCU_DEVICE_ATTRIBUTE_)rrrrrr.r)rrrs r=_build_reverse_device_attrsrbsS #F    CE  ??6 " "&-eT&:CS[\" # Jr?ct\rSrSrSr\S5rSrSrSr Sr Sr S r S r S rS rS r\S5rSrg)riz} The device object owns the CUDA contexts. This is owned by the driver object. User should not construct devices directly. c[[R55H0n[RU5nUR 5U:XdM.Us $ SR U5n[ U5e)zWCreate Device object from device identity created by ``Device.get_device_identity()``. zKNo device of {} is found. Target device may not be visible in this process.)ranger.r"rget_device_identityformat RuntimeError)rVidentitydeviddrs r= from_identityDevice.from_identity$sd 62245E!!%(A$$&(26 DfX v& &r?c [(a'[RU5nX l[ U5nO;[ 5n[R[ U5U5 URnX0lSUSU3nX:wa [U5e0Ul URUR4Ul Sn[(a@[RXPR5nURS5RS5nO:[ U-"5n[RXeUR5 URnXpl[(a5[R%UR5n['UR(5n OG[+5n[R%[ U5UR5 ['[)U55n Sn U S-n U S-n U S -n S U S U S U S U S U 3 nX-UlSUlg) NzDriver returned device z instead of utf-8z%02xzGPU--)rr. cuDeviceGetidr0rrr rh attributesCOMPUTE_CAPABILITY_MAJORCOMPUTE_CAPABILITY_MINORcompute_capabilitycuDeviceGetNamedecoderstripr rcuDeviceGetUuidtuplebytesruuidprimary_context)rVrresult got_devnumrSbufszbufrr uuid_valsbb2b4b6r&s r=rUDevice.__init__4s >''/FGVJWF   uV}f 5J G' |<xH  s# #$(#@#@#'#@#@#B >((8C::g&--d3DE>$C  " "3tww 799D  >))$''2Ddjj)I9D  " "5; 8eDk*I  U U URD"Qrd!B4q-O #r?cJURURURS.$)N) pci_domain_id pci_bus_id pci_device_id) PCI_DOMAIN_ID PCI_BUS_ID PCI_DEVICE_IDr\s r=rfDevice.get_device_identityhs%!////!//  r?c8SURUR4-$)Nz)rwrr\s r=__repr__Device.__repr__os&$''499)===r?cr[(a=[[RSU35n[R X R 5nOI[Un[5n[R [U5X R 5 URn[XU5 U$![a [U5ef=f)zRead attributes lazily ra)rr.rCUdevice_attributer.cuDeviceGetAttributerwDEVICE_ATTRIBUTESrrrrr r)rVattrrRr rs r=rDevice.__getattr__rs >7551$8:D//gg>E +(.WF  ' 'f tWW ELLEE"  +$T** +s B B6c,[UR5$rQ)hashrwr\s r=__hash__Device.__hash__sDGG}r?c`[U[5(aURUR:H$gNF)r/rrwrVothers r=__eq__ Device.__eq__s% eV $ $77ehh& &r?c X:X+$rQrCrs r=__ne__ Device.__ne__s ""r?clURb UR$[U5 [(a [R UR 5nO>[ R"5n[R [U5UR 5 [[R"U5U5nX lU$)zW Returns the primary context for the device. Note: it is not pushed to the CPU thread. ) rmet_requirement_for_devicerr.cuDevicePrimaryCtxRetainrwrr0rContextrr)rVrOctxs r=get_primary_contextDevice.get_primary_contexts    +'' '"4( >22477;D$$&D  + +E$K AgmmD)40" r?ctUR(a'[RUR5 SUlgg)z? Release reference to primary context if it has been retained. N)rr.cuDevicePrimaryCtxReleaserwr\s r=release_primary_contextDevice.release_primary_contexts,     , ,TWW 5#'D  r?cURbURR5 UR5 [R UR 5 g![R UR 5 f=frQ)rr*rr.cuDevicePrimaryCtxResetrwr\s r=r* Device.resetsX 4##/$$**,  ( ( *  * *477 3F * *477 3s 7A!A:c URS:$)N)r#r")r{r\s r=supports_float16Device.supports_float16s&&&00r?)rxr{rwrrrN)r)rErFrGrB classmethodrlrUrfrrrrrrrr*rCrrHrCr?r=rrsb ' '2$h >( #((411r?rcXUR[:a[U<S[<35eg)Nz has compute capability < )r{MIN_REQUIRED_CCr)devices r=rrs,   ?2 & 9: :3r?c\rSrSrSrSr\S5r\S5r\S5r \S5r \S5r \S 5r \S 5r \S 5r\\S 55rS rg)BaseCUDAMemoryManagerizAAbstract base class for External Memory Management (EMM) Plugins.cRSU;a [S5eURS5Ulg)Ncontextz!Memory manager requires a context)rhpopr)rVrkwargss r=rUBaseCUDAMemoryManager.__init__s& F "BC Czz), r?cg)z Allocate on-device memory in the current context. :param size: Size of allocation in bytes :type size: int :return: A memory pointer instance that owns the allocated memory :rtype: :class:`MemoryPointer` NrC)rVsizes r=memallocBaseCUDAMemoryManager.memallocr?cg)a Allocate pinned host memory. :param size: Size of the allocation in bytes :type size: int :param mapped: Whether the allocated memory should be mapped into the CUDA address space. :type mapped: bool :param portable: Whether the memory will be considered pinned by all contexts, and not just the calling context. :type portable: bool :param wc: Whether to allocate the memory as write-combined. :type wc: bool :return: A memory pointer instance that owns the allocated memory. The return type depends on whether the region was mapped into device memory. :rtype: :class:`MappedMemory` or :class:`PinnedMemory` NrC)rVrmappedportablewcs r= memhostalloc"BaseCUDAMemoryManager.memhostallocrr?cg)a Pin a region of host memory that is already allocated. :param owner: The object that owns the memory. :param pointer: The pointer to the beginning of the region to pin. :type pointer: int :param size: The size of the region in bytes. :type size: int :param mapped: Whether the region should also be mapped into device memory. :type mapped: bool :return: A memory pointer instance that refers to the allocated memory. :rtype: :class:`MappedMemory` or :class:`PinnedMemory` NrCrVownerpointerrrs r=mempinBaseCUDAMemoryManager.mempinrr?cg)zd Perform any initialization required for the EMM plugin instance to be ready to use. :return: None NrCr\s r= initialize BaseCUDAMemoryManager.initializerr?cg)z Return an IPC handle from a GPU allocation. :param memory: Memory for which the IPC handle should be created. :type memory: :class:`MemoryPointer` :return: IPC handle for the allocation :rtype: :class:`IpcHandle` NrCrVmemorys r=get_ipc_handle$BaseCUDAMemoryManager.get_ipc_handlerr?cg)z Returns ``(free, total)`` memory in bytes in the context. May raise :class:`NotImplementedError`, if returning such information is not practical (e.g. for a pool allocator). :return: Memory info :rtype: :class:`MemoryInfo` NrCr\s r=get_memory_info%BaseCUDAMemoryManager.get_memory_inforr?cg)z@ Clears up all memory allocated in this context. :return: None NrCr\s r=r*BaseCUDAMemoryManager.resetrr?cg)z~ Returns a context manager that ensures the implementation of deferred cleanup whilst it is active. :return: Context manager NrCr\s r= defer_cleanup#BaseCUDAMemoryManager.defer_cleanup"rr?cg)z Returns an integer specifying the version of the EMM Plugin interface supported by the plugin implementation. Should always return 1 for implementations of this version of the specification. NrCr\s r=interface_version'BaseCUDAMemoryManager.interface_version+rr?rN)r)rErFrGrBrUrrrrrrrr*rrCrrHrCr?r=rrsK-     (  "            r?r) metaclasscz^\rSrSrSrU4SjrSrS SjrS SjrSr Sr \ RS 5r S rU=r$) HostOnlyCUDAMemoryManageri5a)Base class for External Memory Management (EMM) Plugins that only implement on-device allocation. A subclass need not implement the ``memhostalloc`` and ``mempin`` methods. This class also implements ``reset`` and ``defer_cleanup`` (see :class:`numba.cuda.BaseCUDAMemoryManager`) for its own internal state management. If an EMM Plugin based on this class also implements these methods, then its implementations of these must also call the method from ``super()`` to give ``HostOnlyCUDAMemoryManager`` an opportunity to do the necessary work for the host allocations it is managing. This class does not implement ``interface_version``, as it will always be consistent with the version of Numba in which it is implemented. An EMM Plugin subclassing this class should implement ``interface_version`` instead. cx>[TU]"U0UD6 [R"5Ul[ 5UlgrQ)rTrUrr allocations_PendingDeallocs deallocationsrVrrrWs r=rU"HostOnlyCUDAMemoryManager.__init__Gs1 $)&) ++--/r?cU"5$![arn[(a[RRnO[ RnUR U:Xa&URR5 U"5sSnA$eSnAff=f)z Attempt allocation by calling *allocator*. If an out-of-memory error is raised, the pending deallocations are flushed and the allocation is retried. If it fails in the second attempt, the error is reraised. N) rNrrrCUDA_ERROR_OUT_OF_MEMORYrrRrclear)rV allocatorr}oom_codes r=_attempt_allocation-HostOnlyCUDAMemoryManager._attempt_allocationLsi  ;  ~"++DD 99vv!""((* {" s! BA&B9B?BBcf^^ ^ Sm U(aT [R-m U(aT [R-m U(aT [R-m [(a*U U4SjnU(aUR U5m OU"5m T nO>[ 5m U U U4SjnU(aUR U5 OU"5 T Rn[UT UTU5n[R"UR5nU(a*[UT TUS9n XRU'U R5$[UT TUS9$)zImplements the allocation of pinned host memory. It is recommended that this method is not overridden by EMM Plugin implementations - instead, use the :class:`BaseCUDAMemoryManager`. rc0>[RTT5$rQ)r.cuMemHostAlloc)flagsrsr=r9HostOnlyCUDAMemoryManager.memhostalloc..allocatorss,,T599r?cF>[R[T5TT5 grQ)r.rrrrrsr=rrs%%eGndEBr? finalizer)rCU_MEMHOSTALLOC_DEVICEMAPCU_MEMHOSTALLOC_PORTABLECU_MEMHOSTALLOC_WRITECOMBINEDrrr r _hostalloc_finalizerrrr MappedMemoryrown PinnedMemory) rVrrrrr alloc_keyrrmemrrs ` @@r=r&HostOnlyCUDAMemoryManager.memhostalloccs  U44 4E  U33 3E U88 8E > :229=#+IjG C((3  I(w 4P mmDLL) sGTYGC*-  Y '779 WdiH Hr?c^^^ [T[5(a[(d [T5m[(aTnO TRnSm U(aT [ R -m U UU4SjnU(aURU5 OU"5 [UTXT5n[R"UR5nU(a+[UTTUUS9n XRU'U R5$[UTTUUS9$)zImplements the pinning of host memory. It is recommended that this method is not overridden by EMM Plugin implementations - instead, use the :class:`BaseCUDAMemoryManager`. rc4>[RTTT5 grQ)r.cuMemHostRegisterrsr=r3HostOnlyCUDAMemoryManager.mempin..allocators  $ $WdE :r?)rr)r/r0rr r rCU_MEMHOSTREGISTER_DEVICEMAPr_pin_finalizerrrrr rr r ) rVrrrrr rrrrrs `` @r=r HostOnlyCUDAMemoryManager.mempins gs # #NNw'G >I I   U77 7E ;   $ $Y / K"4)D mmDLL) sGT)24C*-  Y '779 Wd%*35 5r?cx^^^[(aUU4SjnURU5mTnO:[R"5mUUU4SjnURU5 TRn[ UTUT5n[ R"UR5n[UTTUS9nXpRU'UR5$)Nc>[RnT(aURRnOURRn[ R TU5$rQ)rCUmemAttach_flagsCU_MEM_ATTACH_GLOBALr CU_MEM_ATTACH_HOSTr.cuMemAllocManaged)ma_flagsr attach_globalrs r=r.allocatorsF"44 $99??E$77==E//e<[5nT(a[RnO[Rn[R [ T5TU5 grQ)rrrrr.rr)rrptrrs r=rrs9 !66E!44E((sT5Ar?r) rrrrr _alloc_finalizerrrr ManagedMemoryrr ) rVrrrr rrrr!s `` @r=memallocmanaged)HostOnlyCUDAMemoryManager.memallocmanageds > =**95CI&&(C B  $ $Y / I$T3 4@ mmDLL)Cdi@&)#wwyr?clURR5 URR5 g)zClears up all host memory (mapped and/or pinned) in the current context. EMM Plugins that override this method must call ``super().reset()`` to ensure that host allocations are also cleaned up.N)rrrr\s r=r*HostOnlyCUDAMemoryManager.resets(     "r?c#|# URR5 Sv SSS5 g!,(df  g=f7f)a Returns a context manager that disables cleanup of mapped or pinned host memory in the current context whilst it is active. EMM Plugins that override this method must obtain the context manager from this method before yielding to ensure that cleanup of host allocations is also deferred.N)rdisabler\s r=r'HostOnlyCUDAMemoryManager.defer_cleanups'   ' ' ) * ) )s<+ < 9<rrFFFF)r)rErFrGrBrUrrrr$r* contextlibcontextmanagerrrHr_r`s@r=rr5sL"0 .9>.I`(5T$L#r?rc\rSrSrSrSrSrg)GetIpcHandleMixinizLA class that provides a default implementation of ``get_ipc_handle()``. c[U5up#[(a7[RU5n[ UR 5[ U5- nOM[ R"5n[R[U5U5 UR RU- nURRR5n[XURUUS9$)a3Open an IPC memory handle by using ``cuMemGetAddressRange`` to determine the base pointer of the allocation. An IPC handle of type ``cu_ipc_mem_handle`` is constructed and initialized with ``cuIpcGetMemHandle``. A :class:`numba.cuda.IpcHandle` is returned, populated with the underlying ``ipc_mem_handle``. )offset)device_extentsrr.cuIpcGetMemHandler0handlerrrr rrrf IpcHandler)rVrbaseend ipchandler3 source_infos r=r GetIpcHandleMixin.get_ipc_handles#6*  >006I'#d)3F002I  $ $U9%5t <]]((4/Fll))==? FKK &( (r?rCN)r)rErFrGrBrrHrCr?r=r1r1s (r?r1c:\rSrSrSrSrSrSr\S5r Sr g) NumbaCUDAMemoryManagerizInternal on-device memory management for Numba. This is implemented using the EMM Plugin interface, but is not part of the public API.cURR[:Xa*UR5RURlggrQ)rmemory_capacity _SizeNotSetrtotalr\s r=r!NumbaCUDAMemoryManager.initializes:    - - <151E1E1G1M1MD   . =r?cr^^[(aU4SjnURU5mTnO9[R"5mUU4SjnURU5 TRn[ UTUT5n[ R"UR5n[UTTUS9nX`RU'UR5$)Nc.>[RT5$rQ)r. cuMemAlloc)rsr=r2NumbaCUDAMemoryManager.memalloc..allocators((..r?cD>[R[T5T5 grQ)r.rFr)r!rsr=rrG%s!!%*d3r?r) rrrrr r"rrrAutoFreePointerrr )rVrrr rrrr!s ` @r=rNumbaCUDAMemoryManager.memallocs > /**95CI&&(C 4  $ $Y / I$T3 4@ mmDLL)c3 B&)#wwyr?c[(a[R5upOT[5n[5n[R[ U5[ U55 UR nUR n[ XS9$)N)freerB)rr. cuMemGetInforrr  MemoryInfo)rVrLrBs r=r&NumbaCUDAMemoryManager.get_memory_info1s\ > --/KD%:DJE   d U5\ :::DKKEt11r?c[$rQ) _SUPPORTED_EMM_INTERFACE_VERSIONr\s r=r(NumbaCUDAMemoryManager.interface_version=s//r?rCN) r)rErFrGrBrrrrCrrHrCr?r=r>r>s,DN , 200r?r>c[(ag[RS:Xa[qg[R "[R5n[ UR5 g![a [S[R-5ef=f)Ndefaultz$Failed to use memory manager from %s) _memory_managerrCUDA_MEMORY_MANAGERr> importlib import_moduleset_memory_manager_numba_memory_manager Exceptionrh) mgr_modules r=_ensure_memory_managerr]Gsx !!Y.07,,V-G-GH :;;< 7A!5567 77s 9A##'B clU"SS9nURnU[:waSU[4-n[U5eUqg)a"Configure Numba to use an External Memory Management (EMM) Plugin. If the EMM Plugin version does not match one supported by this version of Numba, a RuntimeError will be raised. :param mm_plugin: The class implementing the EMM Plugin. :type mm_plugin: BaseCUDAMemoryManager :return: None Nrz9EMM Plugin interface has version %d - version %d required)rrQrhrU) mm_plugindummyiverrs r=rYrYYsE d #E  B --I5673Or?c2^\rSrSrSrU4SjrSrSrU=r$)rAinz; Dummy object for _PendingDeallocs when *size* is not set. c$>[TU]US5$Nr)rTr)rrrrWs r=r_SizeNotSet.__new__sswsA&&r?cg)N?rCr\s r=r]_SizeNotSet.__str__vsr?rC) r)rErFrGrBrr]rHr_r`s@r=rArAns'r?rAc\rSrSrSr\4Sjr\S5r\4Sjr Sr \ RS5r \S5rS rS rg ) ri}z Pending deallocations of a context (or device since we are using the primary context). The capacity defaults to being unset (_SizeNotSet) but can be modified later once the driver is initialized and the total memory capacity known. cJ[5UlSUlSUlXlgre)r_cons_disable_count_sizer@)rVcapacitys r=rU_PendingDeallocs.__init__s W  'r?cN[UR[R-5$rQ)r0r@rCUDA_DEALLOCS_RATIOr\s r=_max_pending_bytes#_PendingDeallocs._max_pending_bytess4''&*D*DDEEr?c`[RSURU5 URR XU45 U=R [ U5- sl[UR5[R:dUR UR:aUR5 gg)a/ Add a pending deallocation. The *dtor* arg is the destructor function that takes an argument, *handle*. It is used as ``dtor(handle)``. The *size* arg is the byte size of the resource added. It is an optional argument. Some resources (e.g. CUModule) has an unknown memory footprint on the device. z add pending dealloc: %s %s bytesN) rrr)rlrrnr0rrCUDA_DEALLOCS_COUNTrsrrVdtorr6rs r=add_item_PendingDeallocs.add_itemsx  7M 4./ c$i  Of88 8 T444 JJL5r?c UR(drUR(aYURR5upn[R SUR U5 U"U5 UR(aMYSUlgg)zP Flush any pending deallocations unless it is disabled. Do nothing if disabled. zdealloc: %s %s bytesrN) is_disabledrlpopleftrrr)rnrws r=r_PendingDeallocs.clears` **'+zz'9'9';$t 3T]]DIV ***DJ  r?c## U=RS- slSv U=RS-slURS:deg!U=RS-slURS:def=f7f)z[ Context manager to temporarily disable flushing pending deallocation. This can be nested. rNrrmr\s r=r)_PendingDeallocs.disablesm q  ,    1 $ &&!+ ++   1 $ &&!+ ++sA1A(A1)A..A1c URS:$rerr\s r=r|_PendingDeallocs.is_disableds""Q&&r?c,[UR5$)z* Returns number of pending deallocations. )rrlr\s r=__len___PendingDeallocs.__len__s4::r?)rlrmrnr@N)r)rErFrGrBrArUrCrsryrr.r/r)r|rrHrCr?r=rr}si !,( FF+6  , ,''r?rrNz free,totalc"\rSrSrSrSrSrSrS'SjrSr S r S'S jr S r S r S rSrSrSrS(SjrS)SjrS*SjrSrSrS+SjrSrSrSrSrSrSrSrSrSr S(S jr!S!r"\#RHS"5r%S#r&S$r'S%r(S&r)g),rizg This object wraps a CUDA Context resource. Contexts should not be constructed directly by user code. cXlX l[R"5Ul[ 5Ul[5 [US9Ul [R"5Ul 0Ul g)Nr) rr6rrrrrr]rUmemory_managermodulesextras)rVrr6s r=rUContext.__init__sR   ++--/ -d;'')  r?c[RSURR5 URR 5 UR R5 URR5 g)z/ Clean up all owned resources in this context. zreset context of device %sN) rrrrwrr*rrrr\s r=r* Context.resetsO  14;;>>B !!#    "r?c6URR5$)z>Returns (free, total) memory in bytes in the context. )rrr\s r=rContext.get_memory_infos""2244r?Nc\XX44n[(aUR"U6$UR"U6$)zReturn occupancy of a function. :param func: kernel for which occupancy is calculated :param blocksize: block size the kernel is intended to be launched with :param memsize: per-block dynamic shared memory usage intended, in bytes )r-_cuda_python_active_blocks_per_multiprocessor(_ctypes_active_blocks_per_multiprocessor)rVfunc blocksizememsizerrs r=$get_active_blocks_per_multiprocessor,Context.get_active_blocks_per_multiprocessors50 >EEtL L@@$G Gr?cURX#/nU(d[R"U6$URU5 [R"U6$rQ)r6r.+cuOccupancyMaxActiveBlocksPerMultiprocessorr4cuOccupancyMaxActiveBlocksPerMultiprocessorWithFlags)rVrrrrpss r=r5Context._cuda_python_active_blocks_per_multiprocessorsAkk9 .EErJ J %JJBOOr?c[5n[U5URX#4nU(d[R"U6 UR $[R "U6 UR $rQ)rrr6r.rrr )rVrrrrrrs r=r0Context._ctypes_active_blocks_per_multiprocessorsTf t{{I?  > > E||  G G N||r?c^XX4U4n[(aUR"U6$UR"U6$)a"Suggest a launch configuration with reasonable occupancy. :param func: kernel for which occupancy is calculated :param b2d_func: function that calculates how much per-block dynamic shared memory 'func' uses based on the block size. Can also be the address of a C function. Use `0` to pass `NULL` to the underlying CUDA API. :param memsize: per-block dynamic shared memory usage intended, in bytes :param blocksizelimit: maximum block size the kernel is designed to handle )r%_cuda_python_max_potential_block_size _ctypes_max_potential_block_size)rVrb2d_funcrblocksizelimitrrs r=get_max_potential_block_size$Context.get_max_potential_block_sizes7? >==tD D88$? ?r?c4[5n[5n[U5n[U5[U5URUX4/n U(d[R "U 6 O$U R U5 [R"U 6 URUR4$rQ) rrrr6r. cuOccupancyMaxPotentialBlockSizer)cuOccupancyMaxPotentialBlockSizeWithFlagsr ) rVrrrrrgridsizerb2d_cbrs r=r(Context._ctypes_max_potential_block_size&sz7G &x0hy!14;;)  3 3T : KK   < v;#dkk"22 22<<4;;#4#44 44r?c8URRU5$rQ)rr)rVbytesizes r=rContext.memalloc[s""++H55r?c8URRX5$rQ)rr$)rVrrs r=r$Context.memallocmanaged^s""228KKr?c:URRXX45$rQ)rr)rVrrrrs r=rContext.memhostallocas""//(OOr?cU(a3URR(d[SUR-5eURR XX45$)Nz%s cannot map host memory)rCAN_MAP_HOST_MEMORYrrrrs r=rContext.mempinds? $++99!"= "KL L""))%$GGr?cd[(d [S5eURRU5$)z/ Returns an *IpcHandle* from a GPU allocation. zOS does not support CUDA IPC) SUPPORTS_IPCrrrrs r=rContext.get_ipc_handleis,|89 9""11&99r?cSn[(a[RX5nO4[R"5n[R[ U5X5 [ [R"U5UUS9$)Nr)rrr) rr.rrrr MemoryPointerrr)rVr6rrdptrs r=open_ipc_handleContext.open_ipc_handleqs] >,,V;D'')D  % %eDk6 AW]]4%8$"&( (r?cHUS:XdS5e[RX5 g)zLEnable peer access between the current context and the peer context rz$*flags* is reserved and MUST be zeroN)r.cuCtxEnablePeerAccess)rV peer_contextrs r=enable_peer_accessContext.enable_peer_access~s$zAAAz$$\9r?c,[(aA[R"U5n[R UR R U5nO>[5n[R [U5UR R U5 [U5$)zcReturns a bool indicating whether the peer access between the current and peer device is possible. ) rrCUdevicer.cuDeviceCanAccessPeerrrwrrbool)rV peer_devicecan_access_peers r=rContext.can_access_peersp >!**;7K$::4;;>>;FHO$gO  ( ()?)- GO$$r?c[U[5(aURS5n[(aUnO [ U5nUR U5$Nutf8)r/r+encoderr create_module_image)rVptximages r=create_module_ptxContext.create_module_ptxsA c3  **V$C >ESME''..r?c[X5n[(a URnOURRnX RU'[ R "U5$rQ)load_module_imagerr6r rrr)rVrmodulekeys r=rContext.create_module_imagesD"4/ >--C--%%C" S}}V$$r?cz[(a URnOURRnURU grQ)rr6r r)rVrrs r= unload_moduleContext.unload_modules+ >--C--%%C LL r?c[(a[R"[5nO$[R "[R5n[ [R"U5US5$rQ) rrCUstreamCU_STREAM_DEFAULTr cu_streamStreamrrrVr6s r=get_default_streamContext.get_default_streamsF >%%&78F%%f&>&>?FgmmD)6488r?c[(a%[R"[R5nO$[R "[R5n[ [R"U5US5$rQ) rrrCU_STREAM_LEGACYrrrrrrs r=get_legacy_default_stream!Context.get_legacy_default_streamsK >%%g&>&>?F%%f&=&=>FgmmD)6488r?c[(a%[R"[R5nO$[R "[R5n[ [R"U5US5$rQ) rrrCU_STREAM_PER_THREADrrrrrrs r=get_per_thread_default_stream%Context.get_per_thread_default_streamsK >%%g&B&BCF%%f&A&ABFgmmD)6488r?c^[(a:[RRRn[ R U5nO4[R"5n[ R [U5S5 [[R"U5U[URU55$re)rrCUstream_flagsrr r.cuStreamCreaterrrrrr_stream_finalizerr)rVrr6s r= create_streamContext.create_streams~ > **<<BBE**51F%%'F  ! !%- 3gmmD)6'(:(:FCE Er?c[U[5(d [S5e[(a[R "U5nO[ R"U5n[[R"U5USSS9$)Nz&ptr for external stream must be an intT)external) r/r0 TypeErrorrrrrrrrr)rVr!r6s r=create_external_streamContext.create_external_streams_#s##DE E >%%c*F%%c*FgmmD)64#% %r?cJSnU(dU[R-n[(a[R U5nO4[ R "5n[R [U5U5 [[R"U5U[URU5S9$Nrr) rCU_EVENT_DISABLE_TIMINGrr. cuEventCreatercu_eventrEventrr_event_finalizerr)rVtimingrr6s r= create_eventContext.create_events{ U22 2E >))%0F__&F  v 6W]]4(&/0B0BFKM Mr?c,[R5 grQ)r.cuCtxSynchronizer\s r= synchronizeContext.synchronizes!r?c## URR5 URR5 Sv SSS5 SSS5 g!,(df  N=f!,(df  g=f7frQ)rrrr)r\s r=rContext.defer_cleanupsH  . . 0##++-.1 0--1 0s2A0AAA A0 A A A-)A0cLSURURR4-$)Nz)r6rrwr\s r=rContext.__repr__s/4;; 2OOOr?cj[U[5(aURUR:H$[$rQ)r/rr6NotImplementedrs r=rContext.__eq__s( eW % %;;%,,. .! !r?c.URU5(+$rQ)rrs r=rContext.__ne__s;;u%%%r?)rrrrr6rrrQ)Tr,r-rA)*r)rErFrGrBrUr*rrrrrrrrrrrr$rrrrrrrrrrrrrrrrr.r/rrrrrHrCr?r=rrs  #5 48 HP <@@$1 K)  56LPH : (: %/%999 E% M" P" &r?rcD[(a [X5$[X5$) image must be a pointer )rload_module_image_cuda_pythonload_module_image_ctypes)rrs r=rrs~,W<<'77r?c [Rn[U-"5n[U-"5n[R[ U5[R [U5[R[ U5[R[U5[R[[R50n[R[U5-"UR56n[[U5-"UR!56n[R""5n[$R'[)U5U[U5Xg5 UR,n [3[4R6"U5X[9X55$![*a9n SUR,R/S5-n [+U R0U 5eSn A ff=f)NcuModuleLoadDataEx error: %sr)r CUDA_LOG_SIZEr rCU_JIT_INFO_LOG_BUFFERr !CU_JIT_INFO_LOG_BUFFER_SIZE_BYTESr CU_JIT_ERROR_LOG_BUFFER"CU_JIT_ERROR_LOG_BUFFER_SIZE_BYTESCU_JIT_LOG_VERBOSECUDA_VERBOSE_JIT_LOGr cu_jit_optionrkeysr& cu_moduler.cuModuleLoadDataExrrNr r}rR CtypesModulerr_module_finalizer) rrlogszjitinfo jiterrorsoptions option_keys option_valsr6r}rSinfo_logs r=r"r"sX  E~ G%"I $$i&8 //% %%y'; 00(5/   (6+F+F"G G''#g,6HKc'l*W^^-=>K    F(!!%-G "- < }}H  g.)': << (- 0F0Fv0NN1663''(s *F G 4F>>Gc [Rn[U5n[U5n[RnUR UUR UURUURUUR[R0nUR5Vs/sHowPM nnUR5V s/sHoPM n n [RU[U5UU 5n UR#S5n['[(R*"U5X[-X 55$s snfs sn f![ a1n UR#S5n SU -n[!U R$U5eSn A ff=f)r rpr$N)rr% bytearrayr CUjit_optionr&r'r(r)r*r+r-r&r.r/rrNr}rRCudaPythonModulerrr1)rrr2r3r4 jit_optionr5kr6vr7r6r} err_stringrSr8s r=r!r!$sF  EG% I%%J))744e**I55u%%v'B'B G&lln-n1nK-%nn./.1.K/(**5#g, +68~~g&H GMM'2F-g> @@./ (%%g. - :1663''(s$ D4 D#!D(( E#2,EE#cT^^^^^URmURmUUUUU4SjnU$)Nc\>T(aTT TR[RTT5 grQ)ryr. cuMemFree)r rrr!rsr=core_alloc_finalizer..coreKs& I&v//d;r?r+)rr!r rrDrrs ``` @@r=r"r"Gs* ,,K"00M<< Kr?cr^^^^^^URmURmT(d[mUUUUUU4SjnU$)aC Finalize page-locked host memory allocated by `context.memhostalloc`. This memory is managed by CUDA, and finalization entails deallocation. The issues noted in `_pin_finalizer` are not relevant in this case, and the finalization is placed in the `context.deallocations` queue along with finalization of device objects. cj>T(a T(aTT TR[RTT5 grQ)ryr. cuMemFreeHost)r rrrr!rsr=rD"_hostalloc_finalizer..corebs( kI&v33S$?r?)rrrA)rr!r rrrDrrs ```` @@r=r r Ss6!,,K"00M @@ Kr?c8^^^^URmUUUU4SjnU$)a Finalize temporary page-locking of host memory by `context.mempin`. This applies to memory not otherwise managed by CUDA. Page-locking can be requested multiple times on the same memory, and must therefore be lifted as soon as finalization is requested, otherwise subsequent calls to `mempin` may fail with `CUDA_ERROR_HOST_MEMORY_ALREADY_REGISTERED`, leading to unexpected behavior for the context managers `cuda.{pinned,mapped}`. This function therefore carries out finalization immediately, bypassing the `context.deallocations` queue. cR>T(a T(aTT [RT5 grQ)r.cuMemHostUnregister)r rrr!sr=rD_pin_finalizer..coreys kI&""3'r?)r)rr!r rrDrs ``` @r=rrjs!!,,K(( Kr?c^^UU4SjnU$)NcF>TR[RT5 grQ)ryr.cuEventDestroydeallocsr6sr=rD_event_finalizer..cores&//8r?rCrRr6rDs`` r=rrs9 Kr?c^^UU4SjnU$)NcF>TR[RT5 grQ)ryr.cuStreamDestroyrQsr=rD_stream_finalizer..cores&00&9r?rCrTs`` r=rrs: Kr?c^^^^URmURm[(aTmO TRmUUUU4SjnU$)Nc\>^[RmUUU4SjnTRUT5 g)NcX>T"5(dTT;de[RU5 grQ)r.cuModuleUnload)r6rr shutting_downs r= module_unload6_module_finalizer..core..module_unloads'!??c&8 88  ! !& )r?)rr]ry)r^r]deallocr6rrs @r=rD_module_finalizer..cores%++  * /r?)rrrr )rr6rDr`rrs ` @@@r=r1r1s;##GooG~ll 0 0 Kr?c*\rSrSrSrSrSrSrSrg) _CudaIpcImplizbImplementation of GPU IPC using CUDA driver API. This requires the devices to be peer accessible. cURUlURUlURUlURUlSUlgrQ)r8r6rr3 _opened_mem)rVparents r=rU_CudaIpcImpl.__init__s6KK mm KK mm r?c&URb [S5eURb [S5eURURUR UR -5nX lUR5RUR 5$)D Import the IPC memory and returns a raw CUDA memory pointer object z'opening IpcHandle from original processIpcHandle is already opened) r8rnrerr6r3rr view)rVrrs r=open_CudaIpcImpl.opensv 99 FG G    ':; ;%%dkk4;;3JKwwy~~dkk**r?cURc [S5e[RURR5 SUlgNzIpcHandle not opened)rernr.cuIpcCloseMemHandler6r\s r=close_CudaIpcImpl.closes<    #34 4""4#3#3#:#:;r?)rer8r6r3rN r)rErFrGrBrUrlrqrHrCr?r=rcrcs +" r?rcc*\rSrSrSrSrSrSrSrg)_StagedIpcImpliz~Implementation of GPU IPC using custom staging logic to workaround CUDA IPC limitation on peer accessibility between devices. cXlURUlURUlURUlX lgrQ)rfr8r6rr;)rVrfr;s r=rU_StagedIpcImpl.__init__s. KK mm KK &r?cLSSKJn [RUR5n[ (a[ UR5nO URn[URS9nURU URURR55nSSS5 URUR5n[!UWUR5 URU UR#5 SSS5 U$!,(df  Nj=f!,(df  U$=f)Nrr )rf)numbar!rrlr;rr0rwrcrfgpusrlr get_contextrrdevice_to_devicerq)rVrr!srcdev srcdev_idimpl source_ptrnewmems r=rl_StagedIpcImpl.opens%%d&6&67 >FIII I4;;/ YYy !4<<#;#;#=>J"!!$)), TYY7YYy ! JJL" " !" ! s5*D(D D D#cgrQrCr\s r=rq_StagedIpcImpl.closes r?)r8r6rfrr;NrsrCr?r=rurus'4 r?rucf\rSrSrSrSSjrSrSrSrSr S r SS jr S r S r \S 5rSrg)r7ia7 CUDA IPC handle. Serialization of the CUDA IPC handle object is implemented here. :param base: A reference to the original allocation to keep it alive :type base: MemoryPointer :param handle: The CUDA IPC handle, as a ctypes array of bytes. :param size: Size of the original allocation :type size: int :param source_info: The identity of the device on which the IPC handle was opened. :type source_info: dict :param offset: The offset into the underlying allocation of the memory referred to by this IPC handle. :type offset: int NcNXlX lX0lX@lSUlXPlgrQ)r8r6rr;_implr3)rVr8r6rr;r3s r=rUIpcHandle.__init__s#   &  r?c4URc [S5eg)Nz#IPC handle doesn't have source info)r;rhr\s r=_sentry_source_infoIpcHandle._sentry_source_infos    #DE E $r?cUR5 URURR5:Xag[R UR5nUR UR5$)zTReturns a bool indicating whether the active context can peer access the IPC handle T)rr;rrfrrlrrw)rVr source_devices r=rIpcHandle.can_access_peers[   "   w~~AAC C,,T-=-=> &&}'7'788r?cUR5 URb [S5e[XR5UlURR U5$)zCOpen the IPC by allowing staging on the host memory first. rj)rrrnrur;rlrVrs r= open_stagedIpcHandle.open_stagedsK   " :: !:; ;#D*:*:; zzw''r?cURb [S5e[U5UlURRU5$)rirj)rrnrcrlrs r= open_directIpcHandle.open_direct's9 :: !:; ;!$' zzw''r?cURbURU5(a URnO URnU"U5$)ayOpen the IPC handle and import the memory for usage in the given context. Returns a raw CUDA memory pointer object. This is enhanced over CUDA IPC that it will work regardless of whether the source device is peer-accessible by the destination device. If the devices are peer-accessible, it uses .open_direct(). If the devices are not peer-accessible, it uses .open_staged(). )r;rrr)rVrfns r=rlIpcHandle.open1s@    #t';';G'D'D!!B!!B'{r?cnSSKJn Uc URnURU5nUR X$X6S9$)z3 Similar to `.open()` but returns an device array. r) devicearray)shapestridesdtypegpu_data)rritemsizerl DeviceNDArray)rVrrrrrrs r= open_arrayIpcHandle.open_array@sD " ?nnGyy!((u/4)E Er?cvURc [S5eURR5 SUlgro)rrnrqr\s r=rqIpcHandle.closeNs/ :: 34 4  r?c[(aURRnO[UR5nURUUR UR UR4n[RU4$rQ) rr6reservedrrWrr;r3r_rebuild_reduction)rVpreprocessed_handlers r= __reduce__IpcHandle.__reduce__Tsa >"&++"6"6 "' "4  NN  II    KK  ,,d33r?c[(a[R"5nXlO[R "U6nU"SXRX4S9$)N)r8r6rr;r3)rrCUipcMemHandlerrr)r handle_aryrr;r3r6s r=_rebuildIpcHandle._rebuildcs> >++-F(O--z:FV*; ;r?)rr8r6r3rr;rerQ)r)rErFrGrBrUrrrrrlrrqrrrrHrCr?r=r7r7sJ F9 ((  E 4;;r?r7cv\rSrSrSrSrSSjr\S5rSr Sr SS jr SS jr \S 5r \S 5rS rg)rina?A memory pointer that owns a buffer, with an optional finalizer. Memory pointers provide reference counting, and instances are initialized with a reference count of 1. The base ``MemoryPointer`` class does not use the reference count for managing the buffer lifetime. Instead, the buffer lifetime is tied to the memory pointer instance's lifetime: - When the instance is deleted, the finalizer will be called. - When the reference count drops to 0, no action is taken. Subclasses of ``MemoryPointer`` may modify these semantics, for example to tie the buffer lifetime to the reference count, so that the buffer is freed when there are no more references. :param context: The context in which the pointer was allocated. :type context: Context :param pointer: The address of the buffer. :type pointer: ctypes.c_void_p :param size: The size of the allocation in bytes. :type size: int :param owner: The owner is sometimes set by the internals of this class, or used for Numba's internal memory management. It should not be provided by an external user of the ``MemoryPointer`` class (e.g. from within an EMM Plugin); the default of `None` should always suffice. :type owner: NoneType :param finalizer: A function that is called when the buffer is to be freed. :type finalizer: function TNcXlX lX0lX0lUSLUlSUlURUlX@lUb[R"X5Ul ggNr) rdevice_pointerr_cuda_memsize_ is_managedrefctr6_ownerrfinalize _finalizerrVrrrrrs r=rUMemoryPointer.__init__sZ % "#4/ ))   %..t?DO !r?c8URcU$UR$rQ)rr\s r=rMemoryPointer.owners{{*t; ;r?c@[[R"U55$rQ) OwnedPointerrrr\s r=r MemoryPointer.ownsGMM$/00r?cUR(aTURR(d [S5eUR5 URR(aegg)z( Forces the device memory to the trash. zFreeing dead memoryN)rraliverhr\s r=rLMemoryPointer.freesI ????(("#899 OO ,, ,,, r?cUc UROUnU(a,[RURXUR5 g[R URX5 grQ)rr.cuMemsetD8Asyncrr6 cuMemsetD8)rVbyter!streams r=memsetMemoryPointer.memsetsK"]    " "4#6#6#)== 2   d114 ?r?cnUcURU- nOX!- nUR(dUS:wa [S5eUnOURU-nUS:a [S5e[(aI[R "5n[ RRUR55nXWl O[ R"U5n[URXcURS9n[UR[[45(a*[[ R""UR5U5$U$)Nrz non-empty slice into empty slicezsize cannot be negative)r)rdevice_pointer_valuerhrr CUdeviceptrrr from_addressgetPtrr rrrr/rrr)rVstartstoprrkr8r ctypes_ptrs r=rkMemoryPointer.views <99u$D>w~~?OP #'  ..t4 wDJJOD djj=,"? @ @ djj 94@ @Kr?cUR$rQ)rr\s r=device_ctypes_pointer#MemoryPointer.device_ctypes_pointers"""r?c[(a[UR5=(d S$URR$rQ)rr0rr r\s r=r"MemoryPointer.device_pointer_values/ >t**+3t 3&&,, ,r?) rrrrrr6rrrNNrerQ)r)rErFrGrB__cuda_memory__rUrCrr rLrrkrrrHrCr?r=rrnse<O @<<1-@>##--r?rc,^\rSrSrSrU4SjrSrU=r$)rIiaModifies the ownership semantic of the MemoryPointer so that the instance lifetime is directly tied to the number of references. When the reference count reaches zero, the finalizer is invoked. Constructor arguments are the same as for :class:`MemoryPointer`. cX>[[U] "U0UD6 U=RS-slgr)rTrIrUrrs r=rUAutoFreePointer.__init__s& ot-t>v> a r?rC)r)rErFrGrBrUrHr_r`s@r=rIrIsr?rIc:^\rSrSrSrSrSU4SjjrSrSrU=r $)r iaA memory pointer that refers to a buffer on the host that is mapped into device memory. :param context: The context in which the pointer was mapped. :type context: Context :param pointer: The address of the buffer. :type pointer: ctypes.c_void_p :param size: The size of the buffer in bytes. :type size: int :param owner: The owner is sometimes set by the internals of this class, or used for Numba's internal memory management. It should not be provided by an external user of the ``MappedMemory`` class (e.g. from within an EMM Plugin); the default of `None` should always suffice. :type owner: NoneType :param finalizer: A function that is called when the buffer is to be freed. :type finalizer: function Tc>X@lX l[(a([R US5nURUlOP[ R"5n[R [U5US5 URRUlX`l [[U]7XUUS9 URUlURUlgr )owned host_pointerrr.cuMemHostGetDevicePointer_bufptr_rrrr rrTr rUr6r_buflen_rVrrrrrdevptrrWs r=rUMappedMemory.__init__s # >55gqAF --DM))+F  , ,U6]GQ G --33DM$ lD*7D5> + @''   r?c@[[R"U55$rQ)MappedOwnedPointerrrr\s r=r MappedMemory.owns!'--"566r?)rrrr6rrr r)rErFrGrBrrUr rHr_r`s@r=r r s&O"(77r?r c(\rSrSrSrSSjrSrSrg)r iaA pointer to a pinned buffer on the host. :param context: The context in which the pointer was mapped. :type context: Context :param owner: The object owning the memory. For EMM plugin implementation, this ca :param pointer: The address of the buffer. :type pointer: ctypes.c_void_p :param size: The size of the buffer in bytes. :type size: int :param owner: An object owning the buffer that has been pinned. For EMM plugin implementation, the default of ``None`` suffices for memory allocated in ``memhostalloc`` - for ``mempin``, it should be the owner passed in to the ``mempin`` method. :param finalizer: A function that is called when the buffer is to be freed. :type finalizer: function Nc.XlX@lX0lX lUSLUlURUlURUl[(aURUlOURRUlUb[R"X5 ggrQ) rrrrrr6rrrr rrrs r=rUPinnedMemory.__init__0s{   ##4/''   > --DM --33DM    T - !r?cU$rQrCr\s r=r PinnedMemory.ownBs r?)rrrr6rrrrr)r)rErFrGrBrUr rHrCr?r=r r s$.$r?r c:^\rSrSrSrSrSU4SjjrSrSrU=r $)r#iFa A memory pointer that refers to a managed memory buffer (can be accessed on both host and device). :param context: The context in which the pointer was mapped. :type context: Context :param pointer: The address of the buffer. :type pointer: ctypes.c_void_p :param size: The size of the buffer in bytes. :type size: int :param owner: The owner is sometimes set by the internals of this class, or used for Numba's internal memory management. It should not be provided by an external user of the ``ManagedMemory`` class (e.g. from within an EMM Plugin); the default of `None` should always suffice. :type owner: NoneType :param finalizer: A function that is called when the buffer is to be freed. :type finalizer: function Tc>X@lUn[TU] XX5S9 URUl[ (aUR UlgUR RUlg)Nr) rrTrUrrrrrr rs r=rUManagedMemory.__init__\sQ  $D  > //DM //55DMr?c@[[R"U55$rQ)ManagedOwnedPointerrrr\s r=r ManagedMemory.ownhs"7==#677r?)rrrrrr`s@r=r#r#Fs&O 688r?r#c$\rSrSrSSjrSrSrg)rilNc^XlUcURUlOUR(aeX lURmU4SjnUR=RS- sl[R "X5 g)Nc>T=RS-slTRS:deTRS:XaTR5 gg![a gf=f)Nrr)rrLReferenceError)rsr=deref$OwnedPointer.__init__..derefxsP  Q yyA~%~99>HHJ"!  sAA AAr)_mem_viewrrrr)rVmemptrrkrrs @r=rUOwnedPointer.__init__msZ <DJ &&Jii  1%r?c.[URU5$)z$Proxy MemoryPointer methods )r.rrs r=rOwnedPointer.__getattr__stzz5))r?)rrrQ)r)rErFrGrUrrHrCr?r=rrls &0*r?rc\rSrSrSrg)rirCNrDrCr?r=rrrIr?rc\rSrSrSrg)rirCNrDrCr?r=rrrIr?rc\rSrSrS SjrSrSrSr\RS5r SSjr \ \ S 55rS \R R"4S jrS rg)ric\XlX lX@lUb[R"X5 ggrQ)rr6rrr)rVrr6rrs r=rUStream.__init__s*      T - !r?c[(a[UR5$URR=(d [R $rQ)rr0r6r rrr\s r=__int__Stream.__int__s3 >t{{# #;;$$@(@(@ @r?c[(aG[S[RS[RS0n[ UR 5=(d SnO_[RS[RS[RS0nUR R=(d [RnX!;aXUR-$UR(aSX R4-$SX R4-$)Nzz"z&rzz) rrrrrr0r6rr rr)rVdefault_streamsr!s r=rStream.__repr__s >!#@((8,,< Odkk"'aC((*G'')M++< O ++##?v'?'?C  !"'$,,6 6 ]]4\\7JJ J+sLL.AA Ar?cB[RUR5 g)za Wait for all commands in this stream to execute. This will commit any pending memory transfers. N)r.cuStreamSynchronizer6r\s r=rStream.synchronizes ""4;;/r?c#4# Uv UR5 g7f)z A context manager that waits for all commands in this stream to execute and commits any pending memory transfers upon exiting the context. N)rr\s r=auto_synchronizeStream.auto_synchronizes   sNcXU4n[U5 [(a@[RURSS9n[ R "U5n[U5nO URn[RURXSS5 g)a Add a callback to a compute stream. The user provided function is called from a driver thread once all preceding stream operations are complete. Callback functions are called from a CUDA driver thread, not from the thread that invoked `add_callback`. No CUDA API functions may be called from within the callback function. The duration of a callback function should be kept short, as the callback will block later work in the stream and may block other callbacks from being executed. Note: The driver function underlying this method is marked for eventual deprecation and may be replaced in a future CUDA release. :param callback: Callback function with arguments (stream, status, arg). :param arg: Optional user data to be passed to the callback function. rrrN) _py_increfrr0r_stream_callbackrCUstreamCallbackrwr.cuStreamAddCallbackr6)rVcallbackrdatar!stream_callbacks r= add_callbackStream.add_callbacksm($4 >..!6!6(.KC%66s;Od8D"33O""4;;qIr?cUup4nU"X1U5 [U5 g![a#n[R"SU35 SnAN4SnAff=f![U5 f=f)NzException in stream callback: )r[warningswarn _py_decref)r6statusrrrrr}s r=rStream._stream_callbacks[ $( !Fc VS ) t  @ MM:1#> ? ? @ t s% A AA A  A Areturnc^^^[R"5mTR5nU4SjmUU4SjnTRX!5 U$)z Return an awaitable that resolves once all preceding stream operations are complete. The result of the awaitable is the current stream. c>UR5(agUS:XaURT5 gUR[SU355 g)Nrz Stream error )done set_result set_exceptionr[)futurer"rVs r=resolver#Stream.async_done..resolvers@{{}}1!!$'$$Yvh/G%HIr?c*>TRTX!5 grQ)call_soon_threadsafe)rr"r*loopr+s r=r#Stream.async_done..callback s  % %h ?r?)asyncioget_running_loop create_futurer)rVr*rr/r+s` @@r= async_doneStream.async_donesB '')##% J @ (+ r?)rrr6r-rQ)r)rErFrGrUr rrr.r/rr staticmethodrrr1futuresFuturer4rHrCr?r=rrsj.AB40J>GOO22r?rcD\rSrSrS SjrSrS SjrSrS SjrSr S r g) r i NcPXlX lUb[R"X5 ggrQ)rr6rr)rVrr6rs r=rUEvent.__init__ s%     T - !r?c[RUR5 g![a)nUR[ R :XaSnAgeSnAff=f)za Returns True if all work before the most recent record has completed; otherwise, returns False. TNF)r. cuEventQueryr6rNrRrCUDA_ERROR_NOT_READY)rVr}s r=query Event.query sH      ,  vv333  s" AAAAc[(a*U(a URO[R"S5nOU(a UROSn[R URU5 g)z Set the record point of the event to the current point in the given stream. The event will be considered to have occurred when all work that was queued in the stream at the time of the call to ``record()`` has been completed. rN)rr6rrr. cuEventRecord)rVrhstreams r=record Event.record# sB >'-fmm73C3CA3FG'-fmm1GT[['2r?cB[RUR5 g)z> Synchronize the host thread for the completion of the event. N)r.cuEventSynchronizer6r\s r=rEvent.synchronize2 s !!$++.r?c[(a*U(a URO[R"S5nOU(a UROSnSn[R X RU5 g)zJ All future works submitted to stream will wait util the event completes. rN)rr6rrr.cuStreamWaitEvent)rVrrCrs r=wait Event.wait8 sG >'-fmm73C3CA3FG'-fmm1G  ++u=r?c[X5$rQ)event_elapsed_time)rVevtends r= elapsed_timeEvent.elapsed_timeC s !$//r?)rr6rQrA) r)rErFrGrUr?rDrrKrPrHrCr?r=r r  s .  3/ >0r?r c[(a*[RURUR5$[ 5n[R[ U5URUR5 UR $)z> Compute the elapsed time between two events in milliseconds. )rr.cuEventElapsedTimer6r rr )evtstartrOmsecs r=rNrNG sS~((&--HHy!!%+x Nzzr?cH\rSrSrSrS SjrSr\S5r\S5r Sr g) ModuleiS zAbstract base class for modulesNcfXlX lX0lUb[R"X5UlggrQ)rr6r8rrr)rVrr6r8rs r=rUModule.__init__V s.    %..t?DO !r?c:URRU5 g)z#Unload this module from the contextN)rrr\s r=unload Module.unload] s ""4(r?cg)z:Returns a Function object encapsulating the named functionNrCrVrs r= get_functionModule.get_functiona rr?cg)z4Return a MemoryPointer referring to the named symbolNrCr^s r=get_global_symbolModule.get_global_symbole rr?)rrr6r8rQ) r)rErFrGrBrUr[rr_rbrHrCr?r=rWrWS s<)@)IICCr?rWc \rSrSrSrSrSrg)r0ij c[R"5n[R[ U5UR UR S55 [[R"U5X!5$r) r cu_functionr.cuModuleGetFunctionrr6rCtypesFunctionrrrVrr6s r=r_CtypesModule.get_functionl sL##%""5=$++#';;v#6 8gmmD16@@r?c [R"5n[R"5n[R [ U5[ U5UR URS55 [URX#5UR4$r) rrrr.cuModuleGetGlobalrr6rrrr rVrr!rs r=rbCtypesModule.get_global_symbolr sb""$   sU4[$++!%V!4 6T\\35tzzAAr?rCNr)rErFrGr_rbrHrCr?r=r0r0j sA Br?r0c \rSrSrSrSrSrg)r<iz c[RURURS55n[ [ R "U5X!5$r)r.rgr6rCudaPythonFunctionrrris r=r_CudaPythonModule.get_function| s7++DKKV9LM!'--"5vDDr?c[RURURS55up#[ UR X#5U4$r)r.rlr6rrrrms r=rb"CudaPythonModule.get_global_symbol s9,,T[[$++f:MN T\\35t;;r?rCNrorCr?r=r<r<z s E)rr\s r=rFunction.__repr__ s#dii//r?cBURRR$rQ)rrrr\s r=rFunction.device s{{"")))r?cg)z.Set the cache configuration for this function.NrC)rV prefer_equal prefer_cache prefer_shareds r= cache_configFunction.cache_config rr?cg)z0Return the value of the attribute with given ID.NrCrVattrids r=read_func_attrFunction.read_func_attr rr?cg)zHReturn a FuncAttr object with the values of various function attributes.NrCr\s r=r~Function.read_func_attr_all rr?)rr6rrNr,)r)rErFrGgriddimblockdimr sharedmemrUrrCrrrrr~rHrCr?r=r|r| svGH FI/ 0** +r?rrc *U V s/sHn [U 5PM n n [[U 5-"U 6n [(a[U 5nSnOU nSnU (a[R UXUXEUUUU5 g[R UXUXEUUUUU5 gs sn fre)r r rrr.cuLaunchCooperativeKernelcuLaunchKernel) cufunc_handlegxgygzbxbybzrrCr cooperativer param_ptrsparamsparams_for_launchextras r= launch_kernelr s-11DS)C.DJ1Z(: 6F~%f-"(()+)+)2)0):  < m b b'%/#  %%2sB)orarcubinfatbinc\rSrSrSr\SSj5r\S5r\ \S55r \ \S55r \S5r S r \S 5rS rS r\S 5rSrg)Linkeri! zAbstract base class for linkersNc[R(a [XU5$[(a [ XU5$[ XU5$rQ)r'CUDA_ENABLE_MINOR_VERSION_COMPATIBILITY MVCLinkerrCudaPythonLinker CtypesLinker)r max_registerslineinfoccs r=new Linker.new$ s7  9 9]b9 9 ^#MR@ @ < =f)Nrb)rlreadrrkrlbasename)rVrlfrs r= add_cu_fileLinker.add_cu_fileW sD $ B B((./ s A A$c[RRU5SSSnUS:Xa [S5eUS:XaUR U5 g[ R US5nUc[SU35eURX5 g)z=Add a file to the link, guessing its type from its extension.rNrz-Don't know how to link file with no extensionrz,Don't know how to link file with extension .)rkrlrrhrFILE_EXTENSION_MAPr r)rVrlextrs r=add_file_guess_extLinker.add_file_guess_ext\ sggt$Q'+ "9NO O D[   T "%))#t4D|"$''*e$-.. MM$ %r?cg)zComplete the link. Returns (cubin, size) cubin is a pointer to a internal buffer of cubin owned by the linker; thus, it should be loaded before the linker is destroyed. NrCr\s r=completeLinker.completej rr?rrFN)r)rErFrGrBrrrrUrCr8rrrrrrrrHrCr?r=rr! s)== ==>>55-$//0 &  r?rzYMinor version compatibility requires ptxcompiler and cubinlinker packages to be availablecf^\rSrSrSrS U4Sjjr\S5r\S5rS Sjr Sr Sr S r U=r $) riy zS Linker supporting Minor Version Compatibility, backed by the cubinlinker package. cj>SSKJn Uc [ S5e[ T U]XU5 SUSS-US-3nSUS/nU(aS U3nURU5 U(aURS 5 [U5Ul U"S U35Ul g![an[[5UeSnAff=f) Nr) CubinLinkerzEMVCLinker requires Compute Capability to be specified, but cc is Nonesm_r9rz --gpu-namez-cz--maxrregcount=z--generate-line-infoz--arch=) cubinlinkerr ImportError_MVC_ERROR_MESSAGErhrTrUrrptx_compile_options_linker) rVrrrrrbarchptx_compile_optsrrWs r=rUMVCLinker.__init__~ s ; / : ;< < "5RURZ"Q%'()($5 #M?3C  # #C (   # #$: ;#()9#: "WTF#34 % ;01s : ;sB B2B--B2c.URR$rQ)r r8r\s r=r8MVCLinker.info_log s||$$$r?c.URR$rQ)r rr\s r=rMVCLinker.error_log s||%%%r?cSSKJn SSKJn U"UR 5UR5nURRURU5 g![an[ [ 5UeSnAff=f!Ua n[UeSnAff=f)Nr) compile_ptxCubinLinkerError) ptxcompilerrrrrr r}r r  add_cubincompiled_programrK)rVrrrrrbcompile_resultr}s r=rMVCLinker.add_ptx s ; / 4%SZZ\43K3KL % LL " ">#B#BD I  ;01s : ;   %1 $ %s. A&A9 A6!A11A69B ?BB cSSKJn [ US5nUR 5nSSS5 [R"U5RnU[S:XaURRnOlU[S:XaURRnOHU[S:Xa[SU35eU[S :XaURWU5$[SU35eU"WU5 g![an[[5UeSnAff=f!,(df  N=f![ a [US35ef=f!Ua n [U eSn A ff=f) Nrrr not foundrrrzDon't know how to link r)rrrr rlrFileNotFoundErrorrKpathlibPathrrr r add_fatbinr) rVrlrrrbrrrrr}s r=rMVCLinker.add_file s> ; 4 3dD!Qvvx" ||D!&& %g. .''B '1 1((B ', , 7v>? ? '. .<<d+ + 7v>? ? % tTN- ;01s : ;"!  3j12 2 3"  %1 $ %sVC) DD D D8) D3DD DDDD58E >EE cSSKJn URR 5$![an[[5UeSnAff=f!Ua n[ UeSnAff=f)Nrr)rrrr r rrK)rVrrbr}s r=rMVCLinker.complete s[ ; 4 %<<((* *  ;01s : ;   %1 $ %s)#A A=AA AA)r r )NFNz r)rErFrGrBrUrCr8rrrrrHr_r`s@r=rry sI5.%%&& %%: % %r?rcf^\rSrSrSrS U4Sjjr\S5r\S5rS Sjr Sr Sr S r U=r $) ri ) Links for current device if no CC given c j>[TU]XU5 [Rn[U-"5n[U-"5n[ R [U5[ R[U5[ R[U5[ R[U5[ R[S50nU(a[U5U[ R'U(a[S5U[ R'UcSU[ R'O*USS-US-n[U5U[ R '[#UR%55n [#UR'55n [(R*[-U 5-"U 6n [[-U 5-"U 6n [(R."5=Uln [2R5[-U 5X[7UR055 [8R:"U[2R<U 5 XPlX`l XVX/Ul!g)Nrrr9)"rTrUrr%r rr&r r'r r(r)r*CU_JIT_MAX_REGISTERSCU_JIT_GENERATE_LINE_INFOCU_JIT_TARGET_FROM_CUCONTEXT CU_JIT_TARGETr%r-r&rr,r cu_link_stater6r. cuLinkCreaterrr cuLinkDestroylinker_info_buflinker_errors_buf _keep_alive)rVrrrr2 linkerinfo linkererrorsr5cc_valraw_keys raw_valuesr6r7r6rWs r=rUCtypesLinker.__init__ s "5$$un' )   ( ()J*?  3 3Xe_  ) )9\+B  4 4huo  $ $hqk   2:=2IGE.. / 7?{GE33 4 ::;GE66 7URZ"Q%'F+3F+;GE'' ( ''..*+ ++c(m;hG #j/1J? %3355 fCM;!$++. 0 v33V<)!-&kOr?cLURRRS5$r)r1r r}r\s r=r8CtypesLinker.info_log s##))0088r?cLURRRS5$r)r2r r}r\s r=rCtypesLinker.error_log s%%++226::r?c R[U5n[URS55nU=RX4/- sl[R UR [ RU[U5USSS5 g![a!n[U<SUR<35eSnAff=fNrrrz) r rr3r. cuLinkAddDatar6rCU_JIT_INPUT_PTXrrNrKr)rVrrptxbufnamebufr}s r=rCtypesLinker.add_ptx s#4;;v./ V-- >  e.D.D!'S7AtT K >!T^^<= = >s=A;; B&B!!B&c^[URS55nURRU5 [R UR X#SSS5 g![aGnUR[R:XaUS3nOU<SUR<3n[U5eSnAff=fNrrrrz) r rr3rr. cuLinkAddFiler6rNrRrCUDA_ERROR_FILE_NOT_FOUNDrrKrVrlrpathbufr}rSs r=rCtypesLinker.add_file s4;;v./ ( #  dQd K #vv888j)"#T^^4c" "  #s#A B,%AB''B,c[S5n[S5n[RUR[ U5[ U55 URnUS:dS5eURSS2 [R"U[R"[R55n[[ R"R%XB4S95$![ a!n[U<SUR<35eSnAff=f)Nrrz"linker returned a zero sized cubinr)r rr.cuLinkCompleter6rrNrKrr r3rhcastrr rnp ctypeslibas_arrayrV cubin_bufrr} cubin_ptrs r=rCtypesLinker.complete sQK { >  ! !$++uY/?t Mzzax===x   Q KK 6>>&--+HI R\\**9G*DEE >!T^^<= = >s3C C>C99C>r3r6r2r1rr%r&r`s@r=rr sM)PV99;;> #FFr?rcf^\rSrSrSrS U4Sjjr\S5r\S5rS Sjr Sr Sr S r U=r $) ri2 r(c .>[T U]XU5 [Rn[ U5n[ U5n[ R nURUURUURUURUURS0nU(aXUR'U(aSXR'UcSXR'O9USS-US-n [[ R SU 35n XUR"'[%UR'55n [%UR)55n [*R-[/U 5X5Ul[2R4"U[*R6UR05 XPlX`lXVX/Ulg)Nrrr9CU_TARGET_COMPUTE_)rTrUrr%r:rr;r&r'r(r)r*r*r+r,r. CUjit_targetr-r%r-r&r.r/rr6rrr0r1r2r3)rVrrrr2r4r5r=r5r6cc_enumr7r8rWs r=rUCudaPythonLinker.__init__6 sa "5$$u%  ' ))   - -z  8 8%  . .  9 95  ) )1   7DJ33 4 <=G88 9 :?@G;; <URZ"Q%'Fg22 26(;=G07J,, - ''..*+ ))#h-N v33T[[A)!-&hKr?c8URRS5$r)r1r}r\s r=r8CudaPythonLinker.info_log` s##**622r?c8URRS5$r)r2r}r\s r=rCudaPythonLinker.error_logd s%%,,V44r?c @URS5nU=RX/- sl[RRn[ R URXA[U5US//5 g![a!n[U<SUR<35eSnAff=fr?) rr3rCUjitInputTyperAr.r@r6rrNrKr)rVrrrC input_ptxr}s r=rCudaPythonLinker.add_ptxh s++f% SN* >..??I  ic#h!(!R 5 >!T^^<= = >sAA22 B<BBc`URS5nURRU5 [R UR X#S//5 g![ aQnUR[RR:XaUS3nOU<SUR<3n[U5eSnAff=frF) rr3rr.rGr6rNrRrrrHrrKrIs r=rCudaPythonLinker.add_filer s++f% ( #  dQB G #vv))CCCj)"#T^^4c" "  #s#A B-A B((B-c[RUR5upUS:dS5eUR SS2 [R"U[R"[R55n[[RRXB4S95$![a!n[ U<SUR <35eSnAff=f)NrzrrMrN)r.rOr6rNrKrr3rhrPrr rrQrRrSrTs r=rCudaPythonLinker.complete s >$33DKK@OIax===x   Q KK 6>>&--+HI R\\**9G*DEE >!T^^<= = >s!B C )CC rXrr%r&r`s@r=rr2 sM(LT3355> # F Fr?rcUS:wa[(aG[RnURn[R"U5n[ R X#5$[R"5n[Rn[ R [U5X 5 U$[(a[R"5$[R"5$)zRQuery the device pointer usable in the current context from an arbitrary pointer. r) rrCUpointer_attribute#CU_POINTER_ATTRIBUTE_DEVICE_POINTERrr.cuPointerGetAttributerrrr)r! ptr_attrsrptrobjrs r=get_devptr_for_active_ctxrq s ax >33I@@D((-F//= =))+F<&&( ('') )r?cn[U5n[(a;[RU5up#U[R "[ U5U-54$[R"5n[5n[R[U5[U5U5 URURp2X"U-4$)zFind the extents (half open begin and end pointer) of the underlying device memory allocation. NOTE: it always returns the extents of the allocation but the extents of the device memory view that can be a subsection of the entire allocation. ) rrr.cuMemGetAddressRangerrr0rrrrr )devmemrsns r=r4r4 s#6 *F~**62'%%c!fqj111  " J##E!HeAh?ww1a%xr?c[USS5nUc:[U5up#[(a[U5[U5- nOX2- nXlUS:dSR U55eU$)zCheck the memory size of the device memory. The result is cached in the device memory object. It may query the driver for the memory size of the device memory allocation. rNrz{} length array)r.r4rr0rrg)rtszrur}s r=device_memory_sizery si )4 0B zf% >Q#a&BB " 70%,,R007 Ir?cP[USS5nUSL=(a URS;$)z?Returns True if the obj.dtype is datetime64 or timedelta64 rNMm)r.char)rrs r=_is_datetime_dtyper} s+ C$ 'E   3t!33r?cd[U5(aUR[R5nU$)zVWorkaround for numpy#4983: buffer protocol doesn't support datetime64 or timedelta64. )r}rkrQint64rs r=_workaround_for_datetimer s&#hhrxx  Jr?c[U[5(aU$SnU(d,[U[R5=(d [ U5n[ U5n[ R"XU5$)zGet host pointer from an obj. If `readonly` is False, the buffer must be writable. NOTE: The underlying data pointer from the host data buffer is used and it should not be changed until the operation which can be asynchronous completes. F)r/r0rQvoidr}rrmemoryview_get_buffer)rreadonly forcewritables r=rr sV#s M "30K4Fs4K "3 'C  ) )#h GGr?cD[U5n[R"U5$)zHReturns (start, end) the start and end pointer of the array (half open).)rrmemoryview_get_extentsrs r=host_memory_extentsr s "3 'C  * *3 //r?c[U5[U5:XdS5e[U5n[R"XX25upEXT- $)zWGet the byte size of a contiguous memory buffer given the shape, strides and itemsize. z# dim mismatch)rrmemoryview_get_extents_info)rrrndimrur}s r=memory_size_from_infor sD u:W %7'77 % u:D  / / ODA 5Lr?c<[U5upX!:dS5eX!- $)zGet the size of the memoryzmemory extend of negative size)r)rrur}s r=host_memory_sizer s% s #DA 63336 5Lr?cZ[(a UR$[U5R$)z$Get the device pointer as an integer)rrr rs r=rr s$~((($S)///r?cLUc [S5$[U5 UR$)z,Get the ctypes object for the device pointerr)r require_device_memoryrrs r=rr s% {{#  $ $$r?c[USS5$)aKAll CUDA memory object is recognized as an instance with the attribute "__cuda_memory__" defined and its value evaluated to True. All CUDA memory object should also define an attribute named "device_pointer" which value is an int object carrying the pointer value of the device memory address. This is not tested in this method. rF)r.rs r=is_device_memoryr s 3)5 11r?c:[U5(d [S5eg)z9A sentry for methods that accept CUDA memory object. zNot a CUDA memory object.N)rr[rs r=rr s  C 344 !r?c@[US/5nURU5 g)zAdd dependencies to the device memory. Mainly used for creating structures that points to other device memory, so that the referees are not GC and released. _depends_N)r.extend)rtobjsdepsets r=device_memory_dependsr& s V[" -F MM$r?c/nU(aC[U[5(de[RnUR UR 5 O[R nU"[U5[USS9U/UQ76 g) NOTE: The underlying data pointer from the host data buffer is used and it should not be changed until the operation which can be asynchronous completes. T)rN) r/rr.cuMemcpyHtoDAsyncrr6 cuMemcpyHtoDrrdstsrcrrvarargsrs r=host_to_devicer0 sb G &&))))  % %v}}%  ~cLtg>r?c/nU(aC[U[5(de[RnUR UR 5 O[R nU"[U5[U5U/UQ76 gr)r/rr.cuMemcpyDtoDAsyncrr6 cuMemcpyDtoDrrs r=r|r|T s` G &&))))  % %v}}%  ~cN3/@@r?c/nU(aC[U[5(de[RnUR UR 5 O[R nU"[U5X/UQ76 g)zMemset on the device. If stream is not zero, asynchronous mode is used. dst: device memory val: byte value to be written size: number of byte to be written stream: a CUDA stream N)r/rr.rrr6rr)rvalrrrrs r= device_memsetrf sYG &&))))  # #v}}%   ~cC00r?c,[R5 g)z3 Enable profile collection in the current context. N)r.cuProfilerStartrCr?r= profile_startr{ s r?c,[R5 g)z4 Disable profile collection in the current context. N)r.cuProfilerStoprCr?r= profile_stopr s r?c#<# [5 Sv [5 g7f)zQ Context manager that enables profiling on entry and disables profiling on exit. N)rrrCr?r= profilingr s O Nsc*[R5$)z8 Return the driver version as a tuple of (major, minor) )r.r?rCr?r=r?r? s    r?r-rA)rBr4rkrhrrrr'r]r1r itertoolsrabcrrrrrr r r r r rr.rWnumpyrQ collectionsrrryr numba.corerrrr rrrrrrrnumba.cuda.cudadrvrrrCUDA_USE_NVIDIA_BINDINGrr!rrrrgrr pythonapi Py_DecRefr! Py_IncRefr py_objectrr>rhrArKrNrxrrrrrfrrrr rrr6r.rbrrrrrr1r>rQrUr]rYr0rArrNrrr"r!r"r rrrr1rcrur7rrIr MemAllocr r#rrrrr rNrWr0r<rvr|rhrrrrdjittyCU_JIT_INPUT_OBJECTrACU_JIT_INPUT_LIBRARYCU_JIT_INPUT_CUBINCU_JIT_INPUT_FATBINARYrrr rrrrqr4ryr}rrrrrrrrrrrrr|rrrr/rr?rCr?r=rs  '///)//4"LL<<//$||&&w/    ' '    ' ' ''( ''( 0 l  , 1?1( V$0  16 % & BVBJ4V4n 01Y1VY1x: q Fgq h@ 5@F((2.0.0I.0b$% 7$ * # m BvBJ l 3  i&fi&X 8<> @F ..0! 6! H' V' Tv;v;ro-Fo-d m +7?+7\&8$$&R#8O#8L*6*> x'8'8  ,(9(9 wVwt70F70t CwC. B6 B ++J$ %F  " "E  & &%%  ' '))))..  & &%%  ' ')))).. O wO f V%V%r]F6]F@WFvWFz*,&"4H(0 0%25N$?$A$1*  r?