J/PhVdZddlZddlZddlZddlZddlZddlZddlZddlZddl Z ddl Z ddl m Z ddl mZmZddlmZmZmZmZmZmZmZmZmZddlZddlZddlZddlmZmZddl m!Z!ddl"m#Z#m$Z$m%Z%d d l&m'Z'm(Z(d d l)m*Z*d d l)m+Z+m,Z,m-Z-dd l.m/Z/m)Z)m0Z0m1Z1e%j2Z3e3rddl4m4Z5dZ6dZ7ej89dZ:ej;j<Z=ej;j>Z?ej@ge=_Aej@ge?_AdZBGddeCZDGddeCZEGdde(ZFdZGdZHdZIdZJdZKdZLdZMdZNd ZOeNZPGd!d"eQZRGd#d$eQZSeRZTd%ZUeUZVGd&d'eQZWd(ZXGd)d*eQe+ZYGd,d-eYZZGd.d/Z[Gd0d1e[eZZ\d Z]da^d2Z_d3Z`Gd4d5eaZbebZbGd6d7eQZced8d9Zd Gd:d;eQZed<Zfd=Zgd>Zhd?Zid@ZjdAZkdBZldCZmdDZnGdEdFeQZoGdGdHeQZpGdIdJeQZqGdKdLeQZrGdMdNerZsGdOdPesZtGdQdRe!juZvGdSdTesZwGdUdVeQZxGdWdXexe!juZyGdYdZexe!juZzGd[d\eQZ{Gd]d^eQZ|d_Z}Gd`dae+Z~Gdbdce~ZGdddee~ZedfgdgZGdhdie+ZGdjdkeZGdldmeZ ddoZe3r/e5jZejejejejejejdpZn'e/je/je/je/je/je/jdpZGdqdre+ZdsZGdtdueZGdvdweZGdxdyeZdzZd{Zd|Zd}Zd~ZddZdZdZdZdZdZdZdZdZddZddZddZddZdZdZejdZdZdS)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)linuxcptjt}|st t j}tt|d}t|ts tj }| |t jr^tj tj}d}|tj|||n&|tj|S)Nz;== CUDA [%(relativeCreated)d] %(levelname)5s -- %(message)s)fmt)logging getLogger__name__ hasHandlersstrrCUDA_LOG_LEVELuppergetattr isinstanceintCRITICALsetLevel StreamHandlersysstderr setFormatter Formatter addHandler NullHandler)loggerlvlhandlerr&s Y/var/www/html/test/jupyter/venv/lib/python3.11/site-packages/numba/cuda/cudadrv/driver.py make_loggerr>;s  x ( (F     5&'((..00gsD))#s## #"C   5+CJ77GOC  !2s!;!;!; < < <   g & & & &   g133 4 4 4 MceZdZdS)DeadMemoryErrorNr) __module__ __qualname__r?r=rArASDr?rAceZdZdS) LinkerErrorNrBrEr?r=rHrHWrFr?rHc$eZdZfdZdZxZS) CudaAPIErrorcv||_||_tt|||dSN)codemsgsuperrJ__init__)selfrMrN __class__s r=rPzCudaAPIError.__init__\s7  lD!!**455555r?c&d|jd|jS)N[z] )rMrNrQs r=__str__zCudaAPIError.__str__as IIItxx00r?)r)rCrDrPrV __classcell__rRs@r=rJrJ[sG66666 1111111r?rJc(tj}|dkrttjdkrt j}dg}dg}n7tjdkrt j}dg}dg}nt j}dd g}d d g}|rv tj |}n #t$rtd |zwxYwtj |std |z|g}n|dt||Dz}||fS)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 drivercTg|]%\}}tj||&SrE)ospathjoin).0xys r= z,locate_driver_and_loader..sB E E E$(Aq!# Q 2 2 E E Er?)r CUDA_DRIVER_raise_driver_not_foundr4platformctypesWinDLLCDLLr^r_abspath ValueErrorisfiler)envpathdlloaderdldirdlnames candidatess r=locate_driver_and_loaderrsesp G#~~!!! |w=&'.  ! !;&'"#;\*0E &goog..GG & & &G$%&& & &w~~g&& D9;BCDD DY  E E,3E7,C,C E E EE  Z s >BB;cg}g}|D]r} ||}||fcS#t$rR}|tj| ||Yd}~kd}~wwxYwt |rt dSdd|D}t|dS)N c34K|]}t|VdSrLr+)raes r= zload_driver..s(==a3q66======r?) OSErrorappendr^r_rmallrfr`_raise_driver_error)rorrpath_not_existdriver_load_errorr_dllrxerrmsgs r= load_driverrsN (4..C 9     ( ( (  ! !bgnnT&:&:": ; ; ;  $ $Q ' ' ' ' ' ' ' ' ( >$!!!!!==+<=====F#####s  A7AA22A7cNt\}}t||\}}|SrL)rsr)rorrrr_s r= find_driverrs*355HjHj11IC 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*ttrL)rDRIVER_NOT_FOUND_MSGrEr?r=rfrfs / 0 00r?c0tt|zrL)rDRIVER_LOAD_ERROR_MSG)rxs r=r}r}s 014 5 55r?cd}tj}ttD]1}||rt t|}|||<2|S)N CUDA_ERROR)r UniqueDictdirr startswithr.)prefixmapnamerMs r=_build_reverse_error_maprs\ F    CE  ??6 " " 5$''DCI Jr?c(tjSrL)r^getpidrEr?r=_getpidrs 9;;r?ceZdZdZdZdZdZdZdZe dZ dZ dd Z d Z d Zd Zd ZdZddZdZdZdZdZdZdZdS)Driverz0 Driver API functions are lazily bound. Nc^|j}||St|}||_|SrL) _singletonobject__new__)clsobjs r=rzDriver.__new__s0n ?J..%%C CN r?ctj|_d|_d|_d|_ t jrd}t|t|_ dS#t$r}d|_|j |_Yd}~dSd}~wwxYw)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_CUDArrlibrN)rQrNrxs r=rPzDriver.__init__s')) #$(! ." ,)'s+++"}}DHHH . . ."&D ()D % % % % % % % .s0A!! B +BB cl|jrdStad|_ td|dt |_n=#t$r0}|jd|j d}||_ td|d}~wwxYw| dS)NTinitrz ()zError at driver init: ) rr>_loggerinfocuInitrrrJrNrMrr_initialize_extras)rQrx descriptions r=ensure_initializedzDriver.ensure_initializeds    F--" ! LL KKNNN yyDHH  K K KU//af///K(3D %"#IK#I#IJJ J K !!!!!s/A## B-+BBctrdStjdt}|tj}||dtjttjtj tjtj tj }|tj }d|_|d|}||_dS)NcuIpcOpenMemHandlecall_cuIpcOpenMemHandle)USE_NV_BINDINGrh CFUNCTYPEr rset_cuIpcOpenMemHandle _find_apirPOINTERr cu_device_ptrcu_ipc_mem_handlerrr)_ctypes_wrap_fnr)rQ set_protor call_protor safe_calls r=rzDriver._initialize_extrass   F$T844 !*7+I!J!Jt~~.BCCDDD%e&,nV5I&J&J&,nV5M&N&N&,m55 #-*W-L"M"M+D((()B)@BB #,r?c<||jduSrL)rrrUs r= is_availablezDriver.is_availables" !!!(D00r?c||jtd|jztr||S||S)NzError at driver init: %s:)rrrr_cuda_python_wrap_fnrrQfnames r= __getattr__zDriver.__getattr__"so !!!  $ 0"#@#'#<$=>> >  /,,U33 3''.. .r?cla t}n#t$rtwxYw|d}|dd}|_|_fd}fd}t jr|}n|}tj |} t| | S)Nrrcdd|D}tdj||}|dS)N, c,g|]}t|SrErwraargs r=rdzIDriver._ctypes_wrap_fn..verbose_cuda_api_call..?999SC999r?call driver api: %s(%s))r`rdebugr)_check_ctypes_error)argsargstrretcoderlibfnrQs r=verbose_cuda_api_callz5Driver._ctypes_wrap_fn..verbose_cuda_api_call>scYY99D999::F MM3U^V L L LeTlG  $ $UG 4 4 4 4 4r?c~tdj|}|dSNzcall driver api: %s)rrr)r)rrrrrQs r=safe_cuda_api_callz2Driver._ctypes_wrap_fn..safe_cuda_api_callDsA MM/ @ @ @eTlG  $ $UG 4 4 4 4 4r?) rKeyErrorAttributeErrorrrestypeargtypesrCUDA_LOG_API_ARGS functoolswrapssetattr) rQrrprotorrrrwrapperrs ``` r=rzDriver._ctypes_wrap_fn/s  = ,&u- , , ,$U+++ ,AhGQRRyHNN5))E#EM%EN 5 5 5 5 5 5 5  5 5 5 5 5 5 5  # )+GG(G*IOE**733 eY'''s /cttfd}fd}tjr|}n|}t j|}t ||S)Ncdd|D}tdj||S)Nrc,g|]}t|SrErwrs r=rdzNDriver._cuda_python_wrap_fn..verbose_cuda_api_call..Vrr?r)r`rrr)_check_cuda_python_error)rrrrrQs r=rz:Driver._cuda_python_wrap_fn..verbose_cuda_api_callUsXYY99D999::F MM3U^V L L L00t EE Er?cvtdj|Sr)rrr)r)rrrrQs r=rz7Driver._cuda_python_wrap_fn..safe_cuda_api_callZs6 MM/ @ @ @00t EE Er?)r.bindingrrrrr)rQrrrrrrs`` @r=rzDriver._cuda_python_wrap_fnRs'' F F F F F F F  F F F F F F F  # )+GG(G*IOE**733 eY'''r?ctjr tsd}nd}|D]-} t|j|cS#t $rY*wxYwfd}t |||S)N)_v2_ptds_v2_ptsz_ptds_ptsz_v2)rrc(td)NzDriver missing function: )r)rkwsrs r=absent_functionz)Driver._find_api..absent_functionxs!"Ee"E"EFF Fr?)rCUDA_PER_THREAD_DEFAULT_STREAMrr.rrr)rQrvariantsvariantrs ` r=rzDriver._find_apigs  0 # #LHH"H  G txE)<7)<)<=====!      G G G G G e_---s: AAc|jUt|jkr@d}t|t|jt ddSdS)Nz0pid %s forked from pid %s after CUDA driver initzCUDA initialized before forking)rrrcriticalr)rQrNs r= _detect_forkzDriver._detect_fork~sZ 8 GII$9$9DC   S'))TX 6 6 6!"CDD D $9$9r?c|tjkrqt|d}d|d|}t||tjkr|t||dS)NUNKNOWN_CUDA_ERRORCall to results in ) r CUDA_SUCCESS ERROR_MAPgetrerrorCUDA_ERROR_NOT_INITIALIZEDrrJ)rQrrerrnamerNs r=rzDriver._check_ctypes_errors e( ( (mmG-ABBGG05ww?C MM#   %:::!!###w,, , ) (r?cJ|d}|dd}t|dkr|d}|tjjkr`d|d|j}t ||tjjkr|t|||S)Nrrrr) lenrCUresultrrrrrrrJ)rQrreturnedrretvalrNs r=rzDriver._check_cuda_python_errors1+!"" v;;!  AYF g&3 3 3 305w||DC MM#   '*EEE!!###w,, , r?rc|j|}|t|}||j|<tj|SrL)rrDeviceweakrefproxy)rQdevnumdevs r= get_devicezDriver.get_devicesClv&& ;..C#&DL }S!!!r?ctr|St}|t||jSrL)rcuDeviceGetCountrrvalue)rQcounts r=get_device_countzDriver.get_device_countsG  +((** * eEll+++{r?cNt|jS)z)Returns a list of active devices )listrvaluesrUs r= list_deviceszDriver.list_devicess DL''))***r?cf|jD]}|dS)zReset all devices N)rrreset)rQrs r=rz Driver.resets:<&&((  C IIKKKK  r?cX|5}|jttr%tcdddSt j}tt||cdddS ddddS#1swxYwYdS)znPop the active CUDA context and return the handle. If no CUDA context is active, return None. N)get_active_contextrrdrivercuCtxPopCurrentr cu_contextr)rQacpoppeds r=pop_active_contextzDriver.pop_active_contexts $ $ & & ""y$!"!1133 " " " " " " " " $.00F**5==999! " " " " " " " "$ " " " " " " " " " " " " " " " " " "s'B ;BB#&B#ctS)z3Returns an instance of ``_ActiveContext``. )_ActiveContextrUs r=rzDriver.get_active_contextsr?ctrt}nGtjd}ttj||j}|dz}||dzz dz}||fS)zM Returns the CUDA Runtime version as a tuple (major, minor). ri )rrcuDriverGetVersionrhrrr )rQversiondvmajorminors r= get_versionzDriver.get_versionsy  //11GGaB  % %fl2&6&6 7 7 7hG4EDL)b0u~r?rLr)r)rCrD__doc__rrrPrrpropertyrrrrrrrrr rrrrrr&rEr?r=rrs\J... """*,,,*11X1 / / /!!!!F*.EEE ---   """"+++  " " "   r?rcFeZdZdZejZdZdZdZ e Z dS)rzAn 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. cd}t|jdr|jj\}}ntr/t}t |dkrd}nFtjd}tt||j r|nd}|d}ntr't t }nAtj }t t||j }||f|j_d}||_ ||_||_|S)NF ctx_devnumrT)hasattr _tls_cacher,rrcuCtxGetCurrentr0rrrr cuCtxGetDevice cu_device_is_topcontext_handler)rQis_tophctxrhdevices r= __enter__z_ActiveContext.__enter__s! 4?L 1 1 ?5LD&& 4--//t99>>D(++&&uT{{333#z3ttt|!+ !6!6!8!899FF$.00G))%..999$]F.2F^* "  r?cB|jrt|jddSdS)Nr,)r2delattrr.)rQexc_typeexc_valexc_tbs r=__exit__z_ActiveContext.__exit__s. < 3 DO\ 2 2 2 2 2 3 3r?c|jduS)zAReturns True is there's a valid and active CUDA context. N)r3rUs r=__bool__z_ActiveContext.__bool__s"$..r?N) r)rCrDr( threadinglocalr.r7r=r? __nonzero__rEr?r=rrs^ !""J   D333/// KKKr?rcd}tj}ttD]D}||r-t t|||t |d<E|S)NCU_DEVICE_ATTRIBUTE_)rrrrrr.r)rrrs r=_build_reverse_device_attrsrEsg #F    CE ;; ??6 " " ;&-eT&:&:CS[[\\" # Jr?czeZdZdZedZdZdZdZdZ dZ dZ d Z d Z d Zd Zed ZdS)rz The device object owns the CUDA contexts. This is owned by the driver object. User should not construct devices directly. cttD]8}t|}||kr|cS9d|}t |)zgCreate 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.)rangerrr get_device_identityformat RuntimeError)rQidentitydeviddrs r= from_identityzDevice.from_identity$s 6224455 ' 'E!!%((A$$&&(223DfX v&& &r?c tr1t|}||_t |}nDt }tt |||j}||_d|d|}||krt|i|_ |j |j f|_ d}trIt ||j}|dd}n9t!|z}t |||j|j}||_tr4t|j}t'|j} nWt+}tt ||jt't)|} d} | dz} | dz} | d z} d | d | d | d | d | }|| z|_d|_dS) NzDriver returned device z instead of utf-8z%02xzGPU--)rr cuDeviceGetidr0rrr rK attributesCOMPUTE_CAPABILITY_MAJORCOMPUTE_CAPABILITY_MINORcompute_capabilitycuDeviceGetNamedecoderstripr rcuDeviceGetUuidtuplebytesruuidprimary_context)rQrresult got_devnumrNbufszbufrrd uuid_valsbb2b4b6r&s r=rPzDevice.__init__4s  !''//FDGVJJWWF   uV}}f 5 5 5J DGH HHHH Z  s## #$(#@#'#@#B  ((88C::g&&--d33DDE>$$C  " "3tw 7 7 79D   +))$'22Ddj))II99D  " "5;; 8 8 8eDkk**I  U U U-R--"--r--B----)O #r?c,|j|j|jdS)N) pci_domain_id pci_bus_id pci_device_id) PCI_DOMAIN_ID PCI_BUS_ID PCI_DEVICE_IDrUs r=rIzDevice.get_device_identityhs"!//!/   r?c$d|j|jfzS)Nz)rYrrUs r=__repr__zDevice.__repr__os&$'49)===r?ctr>ttjd|}t||j}no t|}n#t$rt|wxYwt}tt|||j|j }t||||S)zRead attributes lazily rD)rr.rCUdevice_attributercuDeviceGetAttributerYDEVICE_ATTRIBUTESrrrrr r)rQattrrMr rfs r=rzDevice.__getattr__rs  !758$88::D//dg>>EE +(. + + +$T*** +WWF  ' 'f tTW E E ELEdE""" s  AA/c*t|jSrL)hashrYrUs r=__hash__zDevice.__hash__sDG}}r?cPt|tr|j|jkSdSNF)r/rrYrQothers r=__eq__z Device.__eq__s( eV $ $ '7eh& &ur?c||k SrLrErs r=__ne__z Device.__ne__sEM""r?c`|j|jSt|tr t|j}n@t j}tt||jttj ||}||_|S)zo Returns the primary context for the device. Note: it is not pushed to the CPU thread. ) remet_requirement_for_devicerrcuDevicePrimaryCtxRetainrYrrrContextrr)rQr5ctxs r=get_primary_contextzDevice.get_primary_contexts   +' '"4(((  B2247;;DD$&&D  + +E$KK A A AgmD))400" r?cd|jr(t|jd|_dSdS)zO Release reference to primary context if it has been retained. N)rercuDevicePrimaryCtxReleaserYrUs r=release_primary_contextzDevice.release_primary_contexts=   (  , ,TW 5 5 5#'D  ( (r?c |j|j|t|jdS#t|jwxYwrL)rerrrcuDevicePrimaryCtxResetrYrUs r=rz Device.resetsq 4#/$**,,,  ( ( * * *  * *47 3 3 3 3 3F * *47 3 3 3 3s 4A!A8c|jdkS)N)r#r")r]rUs r=supports_float16zDevice.supports_float16s&&00r?N)r)rCrDr( classmethodrOrPrIrwrrrrrrrr)rrEr?r=rrs ' '[ '2$2$2$h   >>>( ###((((44411X111r?rcX|jtkrt|dtdS)Nz has compute capability < )r]MIN_REQUIRED_CCr)devices r=rrs:  ?22 & 9:: :32r?ceZdZdZdZedZedZedZedZ edZ edZ ed Z ed Z eed Zd S) BaseCUDAMemoryManagerzAAbstract base class for External Memory Management (EMM) Plugins.c`d|vrtd|d|_dS)Ncontextz!Memory manager requires a context)rKpopr)rQrkwargss r=rPzBaseCUDAMemoryManager.__init__s3 F " "BCC Czz),, r?cdS)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` NrE)rQsizes r=memalloczBaseCUDAMemoryManager.memallocr?cdS)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` NrE)rQrmappedportablewcs r= memhostallocz"BaseCUDAMemoryManager.memhostallocrr?cdS)aZ 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` NrErQownerpointerrrs r=mempinzBaseCUDAMemoryManager.mempinrr?cdS)z Perform any initialization required for the EMM plugin instance to be ready to use. :return: None NrErUs r= initializez BaseCUDAMemoryManager.initializerr?cdS)a 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` NrErQmemorys r=get_ipc_handlez$BaseCUDAMemoryManager.get_ipc_handlerr?cdS)a 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` NrErUs r=get_memory_infoz%BaseCUDAMemoryManager.get_memory_inforr?cdS)zX Clears up all memory allocated in this context. :return: None NrErUs r=rzBaseCUDAMemoryManager.resetrr?cdS)z Returns a context manager that ensures the implementation of deferred cleanup whilst it is active. :return: Context manager NrErUs r= defer_cleanupz#BaseCUDAMemoryManager.defer_cleanup"rr?cdS)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. NrErUs r=interface_versionz'BaseCUDAMemoryManager.interface_version+rr?N)r)rCrDr(rPrrrrrrrrrr)rrEr?r=rrs&KK---   ^   ^ (  ^ "  ^   ^   ^   ^   ^   ^X   r?r) metaclasscheZdZdZfdZdZ d dZd dZdZdZ e j d Z xZ S) HostOnlyCUDAMemoryManagera]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. ctj|i|tj|_t |_dSrL)rOrPrr allocations_PendingDeallocs deallocationsrQrrrRs r=rPz"HostOnlyCUDAMemoryManager.__init__GsD$)&))) +---//r?c |S#t$r^}trtjj}n t j}|j|kr(|j|cYd}~Sd}~wwxYw)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) rJrrrCUDA_ERROR_OUT_OF_MEMORYrrMrclear)rQ allocatorrxoom_codes r=_attempt_allocationz-HostOnlyCUDAMemoryManager._attempt_allocationLs  9;;     :"+D 9v!!"((*** y{{"""""" s! A4AA/(A4.A//A4Fc6 d |r tjz |r tjz |r tjz tr+ fd}|r|| n | }n>t fd}|r||n | j}t| ||}tj |j }|r1t| |} | |j |<| St| |S)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`. rc:tSrL)rcuMemHostAlloc)flagsrsr=rz9HostOnlyCUDAMemoryManager.memhostalloc..allocatorss,,T5999r?cZttdSrL)rrrrrrsr=rz9HostOnlyCUDAMemoryManager.memhostalloc..allocators'%%eGnndEBBBBBr? finalizer)rCU_MEMHOSTALLOC_DEVICEMAPCU_MEMHOSTALLOC_PORTABLECU_MEMHOSTALLOC_WRITECOMBINEDrrr r _hostalloc_finalizerrrr MappedMemoryrown PinnedMemory) rQrrrrr alloc_keyrrmemrrs ` @@r=rz&HostOnlyCUDAMemoryManager.memhostalloccs  5 U4 4E  4 U3 3E  9 U8 8E  & : : : : : : &229==#)++IIjjG C C C C C C C ((3333  I(w 4PP mDL))  IsGTYGGGC*-D Y '7799 WdiHHH Hr?c ttrtsttr}nj}d |r t jz fd}|r||n |t|||}tj |j }|r2t|||} | |j |<| St|||S)zImplements the pinning of host memory. It is recommended that this method is not overridden by EMM Plugin implementations - instead, use the :class:`BaseCUDAMemoryManager`. rc@tdSrL)rcuMemHostRegisterrsr=rz3HostOnlyCUDAMemoryManager.mempin..allocators!  $ $WdE : : : : :r?)rr)r/r0rr r rCU_MEMHOSTREGISTER_DEVICEMAPr_pin_finalizerrrrrrrr) rQrrrrrrrrrrs `` @r=rz HostOnlyCUDAMemoryManager.mempins6 gs # # (N (w''G  &II I   8 U7 7E ; ; ; ; ; ; ;    $ $Y / / / / IKKK"4)VDD mDL))  5sGT)2444C*-D Y '7799 Wd%*3555 5r?cvtrfd}||}n6tjfd}||j}t ||}t j|j}t||}||j |<| S)Nctj}r |jj}n |jj}t |SrL)rCUmemAttach_flagsCU_MEM_ATTACH_GLOBALr CU_MEM_ATTACH_HOSTrcuMemAllocManaged)ma_flagsr attach_globalrs r=rz.allocatorsA"4 >$9?EE$7=E//e<<.allocatorsJ 5!6EE!4E((sT5AAAAAr?r) rrrrr _alloc_finalizerrrr ManagedMemoryrr) rQrrrrrrrrs `` @r=memallocmanagedz)HostOnlyCUDAMemoryManager.memallocmanageds  " = = = = = =**955CII&((C B B B B B B B  $ $Y / / / I$T3 4@@ mDL))Cdi@@@&)#wwyyr?cj|j|jdS)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)rrrrUs r=rzHostOnlyCUDAMemoryManager.resets4       """""r?c#vK|j5dVddddS#1swxYwYdS)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)rdisablerUs r=rz'HostOnlyCUDAMemoryManager.defer_cleanups  ' ' ) )   EEE                  s .22FFFF)r)rCrDr(rPrrrrr contextlibcontextmanagerrrWrXs@r=rr5s"00000 .9>.I.I.I.I`(5(5(5(5T$$$L###r?rceZdZdZdZdS)GetIpcHandleMixinzLA class that provides a default implementation of ``get_ipc_handle()``. ct|\}}tr?t|}t |jt |z }nJt j}tt|||jj |z }|j j }t|||j||S)a[Open 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_extentsrrcuIpcGetMemHandler0handlerrrr rrrI IpcHandler)rQrbaseend ipchandler source_infos r=rz GetIpcHandleMixin.get_ipc_handles#6** c  00066I''#d))3FF022I  $ $U9%5%5t < < <](4/Fl)==?? FK &((( (r?N)r)rCrDr(rrEr?r=rrs-(((((r?rc:eZdZdZdZdZdZedZdS)NumbaCUDAMemoryManagerzInternal on-device memory management for Numba. This is implemented using the EMM Plugin interface, but is not part of the public API.cz|jjtkr%|j|j_dSdSrL)rmemory_capacity _SizeNotSetrtotalrUs r=rz!NumbaCUDAMemoryManager.initializes=   - < <151E1E1G1G1MD  . . . = .allocators((...r?cXttdSrL)rrr)rrsr=rz2NumbaCUDAMemoryManager.memalloc..allocator%s%!!%**d33333r?r) rrrrr rrrrAutoFreePointerrr)rQrrrrrrrs ` @r=rzNumbaCUDAMemoryManager.memallocs  " / / / / /**955CII&((C 4 4 4 4 4 4  $ $Y / / / I$T3 4@@ mDL))c3 BBB&)#wwyyr?c*trt\}}n_t}t}tt |t ||j}|j}t ||S)N)freer)rr cuMemGetInforrr  MemoryInfo)rQr rs r=rz&NumbaCUDAMemoryManager.get_memory_info1sv   --//KD%%::DJJE   d U5\\ : : ::DKEt51111r?ctSrL) _SUPPORTED_EMM_INTERFACE_VERSIONrUs r=rz(NumbaCUDAMemoryManager.interface_version=s//r?N) r)rCrDr(rrrr)rrEr?r=rrskDDNNN , 2 2 200X000r?rctrdStjdkr tadS t jtj}t |jdS#t$rtdtjzwxYw)Ndefaultz$Failed to use memory manager from %s) _memory_managerrCUDA_MEMORY_MANAGERr importlib import_moduleset_memory_manager_numba_memory_manager ExceptionrK) mgr_modules r=_ensure_memory_managerrGs !Y..07,V-GHH :;<<<<< 777A!5677 77s 2A'A?c||d}|j}|tkrd|tfz}t||adS)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)rr rKr) mm_plugindummyiverrs r=rrYsT Id # # #E  B ---I5673OOOr?c(eZdZdZfdZdZxZS)rzC Dummy object for _PendingDeallocs when *size* is not set. cHt|dSNr)rOr)rrrrRs r=rz_SizeNotSet.__new__sswwsA&&&r?cdS)N?rErUs r=rVz_SizeNotSet.__str__vssr?)r)rCrDr(rrVrWrXs@r=rrnsQ'''''r?rc~eZdZdZefdZedZefdZdZ e j dZ edZ dZd S) rz 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. cVt|_d|_d|_||_dSr!)r_cons_disable_count_sizer)rQcapacitys r=rPz_PendingDeallocs.__init__s*WW  'r?cDt|jtjzSrL)r0rrCUDA_DEALLOCS_RATIOrUs r=_max_pending_bytesz#_PendingDeallocs._max_pending_bytess4'&*DDEEEr?cLtd|j||j|||f|xjt |z c_t|jtj ks|j|j kr| dSdS)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)r&r{r(r0rrCUDA_DEALLOCS_COUNTr,rrQdtorrrs r=add_itemz_PendingDeallocs.add_items  7MMM 4./// c$ii  OOf8 8 8 T444 JJLLLLL54r?c|js`|jrP|j\}}}td|j||||jPd|_dSdS)zh Flush any pending deallocations unless it is disabled. Do nothing if disabled. zdealloc: %s %s bytesrN) is_disabledr&popleftrrr)r(r/s r=rz_PendingDeallocs.clears|  * '+z'9'9';';$vt 3T]DIIIV * DJJJ   r?c#K|xjdz c_ dV|xjdzc_|jdksJdS#|xjdzc_|jdksJwxYw)zs Context manager to temporarily disable flushing pending deallocation. This can be nested. rNrr'rUs r=rz_PendingDeallocs.disables q  , EEE   1 $  &!++++++   1 $  &!++++++++s 7Ac|jdkSr!r6rUs r=r3z_PendingDeallocs.is_disableds"Q&&r?c*t|jS)z: Returns number of pending deallocations. )rr&rUs r=__len__z_PendingDeallocs.__len__s4:r?N)r)rCrDr(rrPr)r,r1rrrrr3r9rEr?r=rr}s !,(((( FFXF+6    , , ,''X'r?rr z free,totalc eZdZdZdZdZdZ d(dZdZdZ d(d Z d Z d Z d Z d ZdZdZd)dZd*dZd+dZdZdZd,dZdZdZdZdZdZdZdZd Zd!Zd)d"Z d#Z!e"j#d$Z$d%Z%d&Z&d'Z'dS)-rzs This object wraps a CUDA Context resource. Contexts should not be constructed directly by user code. c||_||_tj|_t |_tt||_ tj|_ i|_ dS)Nr) rrrrrrrrrmemory_managermodulesextras)rQrrs r=rPzContext.__init__si   +---//   -d;;;'))  r?ctd|jj|j|j|jdS)z? Clean up all owned resources in this context. zreset context of device %sN) rrrrYr<rr=rrrUs r=rz Context.resetsa  14;>BBB !!###    """""r?c4|jS)z>Returns (free, total) memory in bytes in the context. )r<rrUs r=rzContext.get_memory_infos"22444r?NcD||||f}tr |j|S|j|S)aReturn 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)rQfunc blocksizememsizerrs r=$get_active_blocks_per_multiprocessorz,Context.get_active_blocks_per_multiprocessors;i%0  HE4EtL L@4@$G Gr?c||j||g}|stj|S||tj|SrL)rr+cuOccupancyMaxActiveBlocksPerMultiprocessorr{4cuOccupancyMaxActiveBlocksPerMultiprocessorWithFlags)rQrDrErFrpss r=rBz5Context._cuda_python_active_blocks_per_multiprocessorsGk9g . KErJ J %JBOOr?ct}t||j||f}|stj|ntj||jSrL)rrrrrIrJr )rQrDrErFrrrs r=rCz0Context._ctypes_active_blocks_per_multiprocessorsQf t{Iw? O  > E E E  G N N|r?cF|||||f}tr |j|S|j|S)ajSuggest 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)rQrDb2d_funcrFblocksizelimitrrs r=get_max_potential_block_sizez$Context.get_max_potential_block_sizes=h?  @=4=tD D848$? ?r?c*t}t}t|}t|t||j|||g} |st j| n#| |t j| |j|jfSrL) rrrrr cuOccupancyMaxPotentialBlockSizer{)cuOccupancyMaxPotentialBlockSizeWithFlagsr ) rQrDrPrFrQrgridsizerEb2d_cbrs r=rOz(Context._ctypes_max_potential_block_size&s77GG &x00hy!1!14;) D  3T : : : KK     ;FHHOO$ggO  ( ()?)?)- G G GO$$$r?ct|tr|d}tr|}nt |}||SNutf8)r/r+encoderr create_module_image)rQptximages r=create_module_ptxzContext.create_module_ptxsS c3   %**V$$C  "EESMME''...r?ct||}tr|j}n |jj}||j|<t j|SrL)load_module_imagerrr r=rr)rQrmodulekeys r=rzContext.create_module_imagesH"4//  &-CC-%C" S}V$$$r?cLtr|j}n |jj}|j|=dSrL)rrr r=)rQrrs r= unload_modulezContext.unload_modules-  &-CC-%C L   r?ctrtjt}nt jtj}t tj||dSrL) rrCUstreamCU_STREAM_DEFAULTr cu_streamStreamrrrQrs r=get_default_streamzContext.get_default_streamsM  @%&788FF%f&>??FgmD))64888r?ctrtjtj}nt jtj}t tj||dSrL) rrrCU_STREAM_LEGACYrrrrrrs r=get_legacy_default_streamz!Context.get_legacy_default_streamsN  ?%g&>??FF%f&=>>FgmD))64888r?ctrtjtj}nt jtj}t tj||dSrL) rrrCU_STREAM_PER_THREADrrrrrrs r=get_per_thread_default_streamz%Context.get_per_thread_default_streamsO  C%g&BCCFF%f&ABBFgmD))64888r?cTtr1tjjj}t |}n;tj}t t|dttj ||t|j|Sr!)rrCUstream_flagsrr rcuStreamCreaterrrrrr_stream_finalizerr)rQrrs r= create_streamzContext.create_streams  4 *<BE**511FF%''F  ! !%-- 3 3 3gmD))6'(:FCCEE Er?ct|tstdtrt j|}nt j|}ttj ||ddS)Nz&ptr for external stream must be an intT)external) r/r0 TypeErrorrrrrrrrr)rQrrs r=create_external_streamzContext.create_external_streamsx#s## FDEE E  +%c**FF%c**FgmD))64#%%% %r?cPd}|s|tjz}trt|}n;t j}tt||ttj ||t|j |SNrr) rCU_EVENT_DISABLE_TIMINGrr cuEventCreatercu_eventrEventrr_event_finalizerr)rQtimingrrs r= create_eventzContext.create_events 3 U2 2E  7))%00FF_&&F  v 6 6 6W]4((&/0BFKKMMM Mr?c8tdSrL)rcuCtxSynchronizerUs r= synchronizezContext.synchronizes!!!!!r?c#K|j5|j5dVdddn #1swxYwYddddS#1swxYwYdSrL)r<rrrrUs r=rzContext.defer_cleanups  . . 0 0  #++--                                   s3AA AA AA AA#&A#c.d|j|jjfzS)Nz)rrrYrUs r=rwzContext.__repr__s/4; 2OOOr?cZt|tr|j|jkStSrL)r/rrNotImplementedrs r=rzContext.__eq__s) eW % % ";%,. .! !r?c.|| SrL)rrs r=rzContext.__ne__s;;u%%%%r?rL)Trrr')(r)rCrDr(rPrrrGrBrCrRrOrNr`rcrrrrrrrprtrzrrrrrrrrrrrrrrwrrrEr?r=rrs~    # # #555 48 H H H HPPP   <@@@@@$111 K K K)))  5 5 5666LLLLPPPPHHHH ::: ( ( (:::: % % %///%%%999999999 E E E%%% M M M M""" PPP""" &&&&&r?rcPtrt||St||S)! image must be a pointer )rload_module_image_cuda_pythonload_module_image_ctypes)rrs r=rrs,8,We<<<'777r?c tj}t|z}t|z}tjt |tjt|tjt |tj t|tj ttj i}tj t|z|}tt|z|}tj} t$t)||t|||nD#t*$r7} d|jdz} t+| j| d} ~ wwxYw|j} t3t5j||| t9||S)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_modulercuModuleLoadDataExrrJr r_rM CtypesModulerr_module_finalizer) rrlogszjitinfo jiterrorsoptions option_keys option_valsrrxrNinfo_logs r=rrs  E~  G%""I $i&8&8 /% %y';'; 0(5//  (6+F"G"G G'#g,,6HKc'll*W^^-=-=>K    F(!!%--G "-{ < < < < (((- 0F0Fv0N0NN163'''(}H  g..)'6:: < <.6s---1---r?cg|]}|SrErE)ravs r=rdz1load_module_image_cuda_python..7s///1///r?rRrN)rr bytearrayr CUjit_optionrrrrrrrrrrrrJr_rMCudaPythonModulerrr)rrrrr jit_optionrrrrrx err_stringrNrs r=rr$sX  EG%  I%J)74e*I5u%v'B G.-gllnn---K//gnn..///K(**5#g,, +688 (((%%g.. - :163'''( ~~g&&H GM'22FH-gv>> @ @@s#*C D /DD c>|j|jfd}|S)NcTr=tjdSrL)r1r cuMemFree)rrrrrsr=corez_alloc_finalizer..coreKs4  'I&v/d;;;;;r?)rr)r<rrrrrrs ``` @@r=rrGsK ,K"0M<<<<<<<<< Kr?cT|j|jstfd}|S)a[ 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. cXrr=tjdSrL)r1r cuMemFreeHost)rrrrrrsr=rz"_hostalloc_finalizer..corebs;  'k 'I&v3S$?????r?)rrr)r<rrrrrrrs ```` @@r=rrSsj!,K"0M @@@@@@@@@@ Kr?c,|jfd}|S)aB 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. cJrr=tdSrL)rcuMemHostUnregister)rrrrsr=rz_pin_finalizer..coreys4  'k 'I&""3'''''r?)r)r<rrrrrs ``` @r=rrjs?!,K(((((((( Kr?cfd}|S)NcHtjdSrL)r1rcuEventDestroydeallocsrsr=rz_event_finalizer..cores"&/88888r?rErrrs`` r=rrs)999999 Kr?cfd}|S)NcHtjdSrL)r1rcuStreamDestroyrsr=rz_stream_finalizer..cores"&0&99999r?rErs`` r=rrs):::::: Kr?c\|j|jtrnjfd}|S)Nc\tjfd}|dS)Nc\svsJt|dSrL)rcuModuleUnload)rrr= shutting_downs r= module_unloadz6_module_finalizer..core..module_unloads;!=?? 8c&8&8&8&8  ! !& ) ) ) ) )r?)rrr1)rrdeallocrrr=s @r=rz_module_finalizer..coresM+  * * * * * * * /////r?)rr=rr )rrrrrr=s ` @@@r=rrsZ#GoGl 0 0 0 0 0 0 0 0 Kr?c$eZdZdZdZdZdZdS) _CudaIpcImplzjImplementation of GPU IPC using CUDA driver API. This requires the devices to be peer accessible. ct|j|_|j|_|j|_|j|_d|_dSrL)rrrr _opened_mem)rQparents r=rPz_CudaIpcImpl.__init__s4K m K m r?c|jtd|jtd||j|j|jz}||_||jS)T Import the IPC memory and returns a raw CUDA memory pointer object Nz'opening IpcHandle from original processIpcHandle is already opened) rrlrrprrrrview)rQrrs r=openz_CudaIpcImpl.opensy 9 FGG G   ':;; ;%%dk4;3JKKwwyy~~dk***r?c|jtdt|jjd|_dSNzIpcHandle not opened)rrlrcuIpcCloseMemHandlerrUs r=closez_CudaIpcImpl.closesB   #344 4""4#3#:;;;r?Nr)rCrDr(rPrrrEr?r=rrsK   +++"     r?rc$eZdZdZdZdZdZdS)_StagedIpcImplzImplementation of GPU IPC using custom staging logic to workaround CUDA IPC limitation on peer accessibility between devices. cj||_|j|_|j|_|j|_||_dSrL)rrrrr)rQrrs r=rPz_StagedIpcImpl.__init__s2 K m K &r?c4ddlm}t|j}t rt |j}n|j}t|j }|j |5| |j }dddn #1swxYwY||j}t!|||j|j |5|dddn #1swxYwY|S)Nrr )r)numbar!rrOrrr0rYrrgpusrr get_contextrrdevice_to_devicer)rQrr!srcdev srcdev_idimpl source_ptrnewmems r=rz_StagedIpcImpl.opens%%d&677  "FIII I4;/// Yy ! ? ?4<#;#;#=#=>>J ? ? ? ? ? ? ? ? ? ? ? ? ? ? ?!!$),, TY777Yy !   JJLLL                s$+-B$$B(+B(,D  DDcdSrLrErUs r=rz_StagedIpcImpl.closes r?NrrEr?r=rrsK'''4     r?rcbeZdZdZddZdZdZdZdZd Z dd Z d Z d Z e d ZdS)rao 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 NrcZ||_||_||_||_d|_||_dSrL)rrrr_implr)rQrrrrrs r=rPzIpcHandle.__init__s1   &  r?c2|jtddS)Nz#IPC handle doesn't have source info)rrKrUs r=_sentry_source_infozIpcHandle._sentry_source_infos"   #DEE E $ #r?c||j|jkrdSt|j}||jS)zdReturns a bool indicating whether the active context can peer access the IPC handle T)rrrrIrrOrzrY)rQr source_devices r=rzzIpcHandle.can_access_peersd   """  w~AACC C C4,,T-=>> &&}'7888r?c||jtdt||j|_|j|S)zCOpen the IPC by allowing staging on the host memory first. Nr)rr rlrrrrQrs r= open_stagedzIpcHandle.open_stagedsU   """ : !:;; ;#D$*:;; zw'''r?c|jtdt||_|j|S)rNr)r rlrrrs r= open_directzIpcHandle.open_direct's> : !:;; ;!$'' zw'''r?cn|j||r|j}n|j}||S)aOpen 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(). )rrzrr)rQrfns r=rzIpcHandle.open1s@   #t';';G'D'D #!BB!Br'{{r?c|ddlm}||j}||}|||||S)zC Similar to `.open()` but returns an device array. r) devicearrayN)shapestridesdtypegpu_data)rritemsizer DeviceNDArray)rQrrrrrros r= open_arrayzIpcHandle.open_array@s^ "!!!!! ?nGyy!!((ug/4t)EE Er?cr|jtd|jd|_dSr)r rlrrUs r=rzIpcHandle.closeNs8 : 344 4  r?ctr |jj}nt|j}|j||j|j|jf}tj |fSrL) rrreservedrbrRrrrr_rebuild_reduction)rQpreprocessed_handlers r= __reduce__zIpcHandle.__reduce__TsW  5"&+"6  "' "4"4  N  I   K  ,d33r?ctrtj}||_nt j|}|d||||S)N)rrrrr)rrCUipcMemHandler$rr)r handle_aryrrrrs r=_rebuildzIpcHandle._rebuildcsR  ;+--F(FOO-z:FsV$*6;;; ;r?r!rL)r)rCrDr(rPrrzrrrr!rr'rr+rEr?r=rrs FFF999 ( ( ((((    E E E E 4 4 4;;[;;;r?rc|eZdZdZdZd dZedZdZdZ dd Z dd Z ed Z ed Z dS)rnaA 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 TNc||_||_||_||_|du|_d|_|j|_||_|tj |||_ dSdSNr) rdevice_pointerr_cuda_memsize_ is_managedrefctr_ownerrfinalize _finalizerrQrrrrrs r=rPzMemoryPointer.__init__sj % "#4/ )   %.tY??DOOO ! r?c"|j|n|jSrL)r3rUs r=rzMemoryPointer.owners{*tt ;r?cDttj|SrL) OwnedPointerrrrUs r=rzMemoryPointer.ownsGM$//000r?c|jr=|jjstd||jjrJdSdS)z8 Forces the device memory to the trash. zFreeing dead memoryN)r1r5aliverKrUs r=r zMemoryPointer.frees] ? -?( :"#8999 OO   , , , ,  - - - ,r?rc||jn|}|r)t|j|||jdSt|j||dSrL)rrcuMemsetD8Asyncr/r cuMemsetD8)rQbyter streams r=memsetzMemoryPointer.memsetsm"]   @  " "4#6e#)= 2 2 2 2 2   d14 ? ? ? ? ?r?c>| |j|z }n||z }|js|dkrtd|}n|j|z}|dkrtdtrLt j}t j| }||_ nt j|}t|j |||j }t|j ttfr'tt!j|j |S|S)Nrz non-empty slice into empty slicezsize cannot be negative)r)rdevice_pointer_valuerKrr CUdeviceptrrr from_addressgetPtrr rnrrr/r9rr)rQstartstoprrrr ctypes_ptrs r=rzMemoryPointer.views <9u$DD%>w~~?O?OPP #'    .t44 wDJOOOD dj=,"? @ @  dj 9 94@@ @Kr?c|jSrL)r/rUs r=device_ctypes_pointerz#MemoryPointer.device_ctypes_pointers ""r?cTtrt|jpdS|jjSrL)rr0r/r rUs r=rCz"MemoryPointer.device_pointer_values,  -t*++3t 3&, ,r?NNr!rL)r)rCrDr(__cuda_memory__rPr)rrr rArrKrCrEr?r=rnrnns<O @ @ @ @<<X<111---@@@@>##X#--X---r?rnc"eZdZdZfdZxZS)raModifies 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`. chtt|j|i||xjdzc_dSr.)rOrrPr2rs r=rPzAutoFreePointer.__init__s:-ot$$-t>v>>> a r?)r)rCrDr(rPrWrXs@r=rrsBr?rc.eZdZdZdZdfd ZdZxZS)ra=A 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 TNc||_||_tr(t|d}|j|_nMt j}tt||d|jj |_||_ tt| |||||j|_|j|_dSr)owned host_pointerrrcuMemHostGetDevicePointer_bufptr_rrrr r/rOrrPrr_buflen_rQrrrrrdevptrrRs r=rPzMappedMemory.__init__s #  455gqAAF -DMM)++F  , ,U6]]GQ G G G -3DM$ lD!!**7FD5> + @ @ @'   r?cDttj|SrL)MappedOwnedPointerrrrUs r=rzMappedMemory.owns!'-"5"5666r?rMr)rCrDr(rNrPrrWrXs@r=rrs]&O""""""(7777777r?rc eZdZdZddZdZdS)ra*A 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||_||_||_||_|du|_|j|_|j|_tr |j|_n|jj |_|tj ||dSdSrL) rrSrrTr1rrWrrVr rr4r6s r=rPzPinnedMemory.__init__0s   ##4/'    4 -DMM -3DM   T9 - - - - - ! r?c|SrLrErUs r=rzPinnedMemory.ownBs r?rM)r)rCrDr(rPrrEr?r=rrsA$....$r?rc.eZdZdZdZdfd ZdZxZS)raIA 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 TNc||_|}t|||||j|_t r|j|_dS|jj|_dS)Nr) rSrOrPrrWrr/rVr rXs r=rPzManagedMemory.__init__\sb  &$)DDD   6 /DMMM /5DMMMr?cDttj|SrL)ManagedOwnedPointerrrrUs r=rzManagedMemory.ownhs"7=#6#6777r?rMr\rXs@r=rrFs]&O 6 6 6 6 6 68888888r?rceZdZddZdZdS)r9Nc||_| |j|_n|jrJ||_|jfd}|jxjdz c_t j||dS)Nc xjdzc_jdksJjdkrdSdS#t$rYdSwxYw)Nrr)r2r ReferenceError)rsr=derefz$OwnedPointer.__init__..derefxsk  Q yA~~~~9>>HHJJJJJ">!    s      T9 - - - - - ! r?chtrt|jS|jjp tjSrL)rr0rr rrrUs r=__int__zStream.__int__s0  At{## #;$@(@ @r?cbtr8tdtjdtjdi}t |jpd}n>tjdtjdtjdi}|jjp tj}||vr|||j zS|j r d||j fzSd||j fzS)Nzz"z&rzz) rrrrrr0rrr rr)rQdefault_streamsrs r=rwzStream.__repr__s  @!#@(8,< Odk""'aCC(*G')M+< O +#?v'?C / ! !"3'$,6 6 ] B4T\7JJ J+sDL.AA Ar?cDt|jdS)zy Wait for all commands in this stream to execute. This will commit any pending memory transfers. N)rcuStreamSynchronizerrUs r=rzStream.synchronizes ""4;/////r?c#:K|V|dS)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)rrUs r=auto_synchronizezStream.auto_synchronizes)   r?Nc|||f}t|trEt|jd}t j|}t|}n|j}t |j ||ddS)ai 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. rYrZrN) _py_increfrr0r\_stream_callbackrCUstreamCallbackrYrcuStreamAddCallbackr)rQcallbackrdatarstream_callbacks r= add_callbackzStream.add_callbacks(h$4  4..!6(.KKC%6s;;Od88DD"3O""4;qIIIIIr?c |\}}}||||n.#t$r!}tjd|Yd}~nd}~wwxYwt|dS#t|wxYw)NzException in stream callback: )rwarningswarn _py_decref)rstatusrr@r~rrxs r=r{zStream._stream_callbacks $( !FHc HVVS ) ) ) ) @ @ @ M>1>> ? ? ? ? ? ? ? ? @ t     Jt    s(A A<AAAA&returnctj}fdfd}|||S)z Return an awaitable that resolves once all preceding stream operations are complete. The result of the awaitable is the current stream. c|rdS|dkr|dS|td|dS)Nrz Stream error )done set_result set_exceptionr)futurerrQs r=resolverz#Stream.async_done..resolversg{{}} J1!!$'''''$$Y/Gv/G/G%H%HIIIIIr?c6||dSrL)call_soon_threadsafe)r@rrlooprs r=r~z#Stream.async_done..callback s!  % %h ? ? ? ? ?r?)asyncioget_running_loop create_futurer)rQrr~rrs` @@r= async_donezStream.async_dones '))##%% J J J J J @ @ @ @ @ @ (F+++ r?rrL)r)rCrDrPrrrwrrrrxr staticmethodrr{rfuturesFuturerrEr?r=rrs....AAABBB4000JJJJ>\GO2r?rc8eZdZd dZdZd dZdZd dZdZdS) rNcT||_||_|tj||dSdSrL)rrrr4)rQrrrs r=rPzEvent.__init__ s7    T9 - - - - - ! r?c t|jdS#t$r!}|jt jkrYd}~dSd}~wwxYw)zy Returns True if all work before the most recent record has completed; otherwise, returns False. TNF)r cuEventQueryrrJrMrCUDA_ERROR_NOT_READY)rQrxs r=queryz Event.query se      , , ,4    v333uuuuu  s# AA A  Arctr|r|jntjd}n |r|jnd}t|j|dS)a 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)rrrrr cuEventRecord)rQr@hstreams r=recordz Event.record# sW  5'-Ffmm73CA3F3FGG'-4fmm1GT['22222r?cDt|jdS)zN Synchronize the host thread for the completion of the event. N)rcuEventSynchronizerrUs r=rzEvent.synchronize2 s  !!$+.....r?ctr|r|jntjd}n |r|jnd}d}t||j|dS)zZ All future works submitted to stream will wait util the event completes. rN)rrrrrcuStreamWaitEvent)rQr@rrs r=waitz Event.wait8 s^  5'-Ffmm73CA3F3FGG'-4fmm1G  $+u=====r?c"t||SrL)event_elapsed_time)rQevtends r= elapsed_timezEvent.elapsed_timeC s!$///r?rLr') r)rCrDrPrrrrrrEr?r=rr s}....    3 3 3 3/// > > > >00000r?rctr%t|j|jSt }tt ||j|j|jS)zF Compute the elapsed time between two events in milliseconds. )rrcuEventElapsedTimerr rr )evtstartrmsecs r=rrG sY((&-HHHyy!!%++x NNNzr?cLeZdZdZddZdZedZedZdS)ModulezAbstract base class for modulesNcl||_||_||_|tj|||_dSdSrL)rrrrr4r5)rQrrrrs r=rPzModule.__init__V s=     %.tY??DOOO ! r?c:|j|dS)z#Unload this module from the contextN)rrrUs r=unloadz Module.unload] s ""4(((((r?cdS)z:Returns a Function object encapsulating the named functionNrErQrs r= get_functionzModule.get_functiona rr?cdS)z4Return a MemoryPointer referring to the named symbolNrErs r=get_global_symbolzModule.get_global_symbole rr?rL) r)rCrDr(rPrrrrrEr?r=rrS s))@@@@)))II^ICC^CCCr?rceZdZdZdZdS)rctj}tt ||j|dttj |||Sr|) r cu_functionrcuModuleGetFunctionrrr~CtypesFunctionrrrQrrs r=rzCtypesModule.get_functionl s_#%%""5==$+#';;v#6#6 8 8 8gmD1164@@@r?c&tj}tj}tt |t ||j|dt|j |||j fSr|) rrrrcuModuleGetGlobalrrr~rnrr rQrrrs r=rzCtypesModule.get_global_symbolr ss"$$    sU4[[$+!%V!4!4 6 6 6T\355tzAAr?Nr)rCrDrrrEr?r=rrj s:AAA BBBBBr?rceZdZdZdZdS)rct|j|d}t t j|||Sr|)rrrr~CudaPythonFunctionrrrs r=rzCudaPythonModule.get_function| sA++DKV9L9LMM!'-"5"5vtDDDr?ct|j|d\}}t |j|||fSr|)rrrr~rnrrs r=rz"CudaPythonModule.get_global_symbol sB,,T[$++f:M:MNN TT\355t;;r?NrrEr?r=rrz s5EEE<<<<)rrUs r=rwzFunction.__repr__ s#di//r?c$|jjjSrL)rrrrUs r=rzFunction.device s{"))r?FcdS)z.Set the cache configuration for this function.NrE)rQ prefer_equal prefer_cache prefer_shareds r= cache_configzFunction.cache_config rr?cdS)z0Return the value of the attribute with given ID.NrErQattrids r=read_func_attrzFunction.read_func_attr rr?cdS)zPReturn a FuncAttr object with the values of various function attributes.NrErUs r=rzFunction.read_func_attr_all rr?Nr)r)rCrDgriddimblockdimr@ sharedmemrPrwr)rrrrrrEr?r=rr sGH FI/// 000**X*>> + + + + +r?rFc  d| D} tt| z| } trt| } d}n| } d}| r%t|||||||||| dSt|||||||||| | dS)Nc,g|]}t|SrE)r rs r=rdz!launch_kernel.. s111S)C..111r?r)r rrr rcuLaunchCooperativeKernelcuLaunchKernel) cufunc_handlegxgygzbxbybzrrr cooperative param_ptrsparamsparams_for_launchextras r= launch_kernelr  s21D111JZ(: 6F%f--"%(()+R)+R)2)0):  < < < < < m "b "b'%/#  % % % % %r?)orarcubinfatbinceZdZdZeddZedZeedZ eedZ ed Z d Z ed Z d Zd ZedZdS)LinkerzAbstract base class for linkersrFNctjrt|||Strt |||St |||SrL)r'CUDA_ENABLE_MINOR_VERSION_COMPATIBILITY MVCLinkerrCudaPythonLinker CtypesLinker)r max_registerslineinfoccs r=newz Linker.new$ sK  9 =]Hb99 9  =#M8R@@ @ x<< >^X>55^5---$//^/000 & & &  ^   r?rzYMinor version compatibility requires ptxcompiler and cubinlinker packages to be availablecdeZdZdZd fd ZedZedZd dZd Z d Z xZ S) rz_ Linker supporting Minor Version Compatibility, backed by the cubinlinker package. NFc ddlm}n'#t$r}tt|d}~wwxYw|t dt |||d|ddz|dz}d|dg}|rd |}|||r|d t||_ |d ||_ dS) Nr) CubinLinkerzEMVCLinker requires Compute Capability to be specified, but cc is Nonesm_r rz --gpu-namez-cz--maxrregcount=z--generate-line-infoz--arch=) cubinlinkerrA ImportError_MVC_ERROR_MESSAGErKrOrPr{rbptx_compile_options_linker) rQrrrrArarchptx_compile_optsrrRs r=rPzMVCLinker.__init__~ s) ; / / / / / / / ; ; ;011s : ; : ;<< < "555)RURZ"Q%'))($5  )3M33C  # #C ( ( (  <  # #$: ; ; ;#()9#:#: "{#3T#3#344 s .).c|jjSrL)rGrrUs r=rzMVCLinker.info_log s |$$r?c|jjSrL)rGr!rUs r=r!zMVCLinker.error_log s |%%r? c" ddlm}ddlm}n'#t$r}t t |d}~wwxYw|||j} |j |j |dS#|$r }t|d}~wwxYw)Nr) compile_ptxCubinLinkerError) ptxcompilerrNrCrPrDrEr_rFrG add_cubincompiled_programrH)rQrrrNrPrcompile_resultrxs r=r#zMVCLinker.add_ptx s ; / / / / / / 4 4 4 4 4 4 4 ; ; ;011s : ;$SZZ\\43KLL % L " ">#BD I I I I I % % %1 $ %s(  3.3 A<<BB  Bc ddlm}n'#t$r}tt|d}~wwxYw t |d5}|}dddn #1swxYwYn #t $rt|dwxYwtj |j }|tdkr |j j }nz|tdkr |j j}n\|tdkrtd||td kr|||Std| |||dS#|$r } t| d} ~ wwxYw) NrrOr3 not foundrrrzDon't know how to link r)rCrPrDrErr4FileNotFoundErrorrHpathlibPathrr9rGrR add_fatbinr#) rQr_r0rPrr6rrrrxs r=r1zMVCLinker.add_file s ; 4 4 4 4 4 4 4 ; ; ;011s : ; 3dD!! Qvvxx                 3 3 311122 2 3|D!!& %g. . .'BB '1 1 1(BB ', , ,>>>?? ? '. . .<<d++ +>>>?? ? % BtTNNNNN % % %1 $ %sW -(-A.A" A."A&&A.)A&*A..B  EE EE c ddlm}n'#t$r}tt|d}~wwxYw |jS#|$r }t |d}~wwxYw)NrrO)rCrPrDrErGr=rH)rQrPrrxs r=r=zMVCLinker.complete s ; 4 4 4 4 4 4 4 ; ; ;011s : ; %<((** * % % %1 $ %s' -(-A AAA)NFNrL r)rCrDr(rPr)rr!r#r1r=rWrXs@r=rry s555555.%%X%&&X& % % % %%%%: % % % % % % %r?rcdeZdZdZd fd ZedZedZd d Zd Z d Z xZ S)r1 Links for current device if no CC given rFNc t|||tj}t |z}t |z}t jt|t jt|t j t|t j t|t j tdi}|rt||t j <|rtd|t j<|d|t j<n0|ddz|dz}t||t j<t#|} t#|} t)jt-| z| } tt-| z| } t)jx|_} t2t-| | | t7|jt9j|t2j| ||_||_ ||| | g|_!dS)Nrrr )"rOrPrrr rrr rr rrrCU_JIT_MAX_REGISTERSCU_JIT_GENERATE_LINE_INFOCU_JIT_TARGET_FROM_CUCONTEXT CU_JIT_TARGETrrrrrr cu_link_staterr cuLinkCreaterrr4 cuLinkDestroylinker_info_buflinker_errors_buf _keep_alive)rQrrrr linkerinfo linkererrorsrcc_valraw_keys raw_valuesrrrrRs r=rPzCtypesLinker.__init__ s "555$un'' ))   ()J*?*?  3Xe__  )9\+B+B  4huoo  $hqkk    J2:=2I2IGE. /  C7?{{GE3 4 ::;GE6 7 7URZ"Q%'F+3F+;+;GE' ( '''..**++ +c(mm;hG #j//1J? %3555 fCMM; !$+.. 0 0 0 v3V<<<)!-& k;Or?c@|jjdSr|)rhr r_rUs r=rzCtypesLinker.info_log s#)00888r?c@|jjdSr|)rir r_rUs r=r!zCtypesLinker.error_log s%+226:::r?rLc ^t|}t|d}|xj||gz c_ t|jt j|t||ddddS#t$r}t|d|j d}~wwxYwNr}rru) r r~rjr cuLinkAddDatarrCU_JIT_INPUT_PTXrrJrHr!)rQrrptxbufnamebufrxs r=r#zCtypesLinker.add_ptx s#4;;v..// VW-- >  e.D!'S7AtT K K K K K > > >!!!T^^<== = >s=B B,B''B,cXt|d}|j| t|j||ddddS#t$r;}|jtj kr|d}n |d|j }t|d}~wwxYwNr}rrVru) r r~rjr{r cuLinkAddFilerrJrMrCUDA_ERROR_FILE_NOT_FOUNDr!rHrQr_r0pathbufrxrNs r=r1zCtypesLinker.add_file s4;;v..// ((( #  dGQd K K K K K # # #v888)))"#!!T^^4c"" "  #s$A$$ B).6B$$B)ctd}td} t|jt |t |n+#t $r}t|d|jd}~wwxYw|j }|dks Jd|j dd=tj |tj tj}tt j||fS)Nrru"linker returned a zero sized cubinr)r rrcuLinkCompleterrrJrHr!r rjrhcastrr rcnp ctypeslibas_arrayrQ cubin_bufrrx cubin_ptrs r=r=zCtypesLinker.complete sQKK {{ >  ! !$+uY/?/?t M M M M > > >!!!T^^<== = >zaxxx=xxx  QQQ K 6>&-+H+HII R\**9TG*DDEEEs;A B&A??Br>r\r]rXs@r=rr s)P)P)P)P)P)PV99X9;;X;>>>> # # #FFFFFFFr?rcdeZdZdZd fd ZedZedZd d Zd Z d Z xZ S)rr_rFNc  t|||tj}t |}t |}t j}|j||j||j ||j ||j di}|r |||j <|r d||j <| d||j<n;|ddz|dz} tt jd| } | ||j<t%|} t%|} t*t/| | | |_t3j|t*j|j||_||_||| | g|_dS)Nrrr CU_TARGET_COMPUTE_)rOrPrrrrrrrrrrrarbrcr. CUjit_targetrdrrrrrfrrrr4rgrhrirj)rQrrrrrkrlrrrmcc_enumrnrorRs r=rPzCudaPythonLinker.__init__6 s "555$u%%  '' )   -z  8%  .  95  )1    E7DGJ3 4  ><=GJ8 9 :?@GJ; < <URZ"Q%'Fg2;6;;==G07GJ, - '''..**++ ))#h--:NN v3T[AAA)!-& h Kr?c6|jdSr|)rhr_rUs r=rzCudaPythonLinker.info_log` s#**6222r?c6|jdSr|)rir_rUs r=r!zCudaPythonLinker.error_logd s%,,V444r?rLc 4|d}|xj||gz c_ tjj}t |j||t||dggdS#t$r}t|d|j d}~wwxYwrs) r~rjrCUjitInputTyperurrtrrrJrHr!)rQrrrw input_ptxrxs r=r#zCudaPythonLinker.add_ptxh s++f%% S'N* >.?I  ic#hh!(!R 5 5 5 5 5 > > >!!!T^^<== = >sAA// B9BBcH|d}|j| t|j||dggdS#t $r@}|jtj j kr|d}n |d|j }t|d}~wwxYwry) r~rjr{rrzrrJrMrrr{r!rHr|s r=r1zCudaPythonLinker.add_filer s++f%% ((( #  dGQB G G G G G # # #v)CCC)))"#!!T^^4c"" "  #s$A B!!;BB!c t|j\}}n+#t$r}t |d|jd}~wwxYw|dks Jd|jdd=tj|tj tj }ttj ||fS)Nrurrr)rrrrJrHr!rjrhrrr rcrrrrs r=r=zCudaPythonLinker.complete s >$33DK@@OItt > > >!!!T^^<== = >axxx=xxx  QQQ K 6>&-+H+HII R\**9TG*DDEEEs"% A AA r>r\r]rXs@r=rr2 s(L(L(L(L(L(LT33X355X5>>>> # # # F F F F F F Fr?rc|dkrtrBtj}|j}tj|}t ||Stj}tj}t t||||StrtjStjS)zZQuery the device pointer usable in the current context from an arbitrary pointer. r) rrCUpointer_attribute#CU_POINTER_ATTRIBUTE_DEVICE_POINTERrDrcuPointerGetAttributerrrr)r ptr_attrsr|ptrobjrYs r=get_devptr_for_active_ctxr s axx  3I@D(--F//f== =)++Ft|dd}|duo|jdvS)z?Returns True if the obj.dtype is datetime64 or timedelta64 rNMm)r.char)rrs r=_is_datetime_dtyper s, C$ ' 'E   3t!33r?cbt|r|tj}|S)z^Workaround for numpy#4983: buffer protocol doesn't support datetime64 or timedelta64. )rrrint64rs r=_workaround_for_datetimer s,#!hhrx   Jr?ct|tr|Sd}|s)t|tjpt |}t |}t j|||S)aGet 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/r0rvoidrrrmemoryview_get_buffer)rreadonly forcewritables r=rTrT sj#s M L"300K4Fs4K4K "3 ' 'C  )#}h G GGr?cHt|}tj|S)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?ct|t|ks Jdt|}tj||||\}}||z S)z_Get the byte size of a contiguous memory buffer given the shape, strides and itemsize. z# dim mismatch)rrmemoryview_get_extents_info)rrrndimrrxs r=memory_size_from_infor sX u::W % % %'7 % % % u::D  /wh O ODAq q5Lr?cPt|\}}||ks Jd||z S)zGet the size of the memoryzmemory extend of negative size)r)rrrxs r=host_memory_sizer s0 s # #DAq 6663666 q5Lr?cFtr|jSt|jS)z$Get the device pointer as an integer)rrKr rs r=r/r/ s$0(($S))//r?cP|tdSt||jS)z,Get the ctypes object for the device pointerNr)r require_device_memoryrKrs r=rKrK s* {{{#  $$r?c$t|ddS)a_All 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. rNF)r.rs r=is_device_memoryr s 3)5 1 11r?cBt|stddS)z9A sentry for methods that accept CUDA memory object. zNot a CUDA memory object.N)rrrs r=rr s- C 5344455r?cRt|dg}||dS)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)robjsdepsets r=device_memory_dependsr& s- V[" - -F MM$r?cg}|r>t|tsJtj}||jn tj}|t|t|d|g|RdS) 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/rrcuMemcpyHtoDAsyncr{r cuMemcpyHtoDr/rTdstsrcrr@varargsrs r=host_to_devicer0 s G !&&)))))  %v}%%%%  B~cLt<<t|tsJtj}||jn tj}|t|t||g|RdSrN) r/rrcuMemcpyDtoHAsyncr{r cuMemcpyDtoHrTr/rs r=device_to_hostrB s G !&&)))))  %v}%%%%  B|C.--t>g>>>>>>r?cg}|r>t|tsJtj}||jn tj}|t|t||g|RdSr)r/rrcuMemcpyDtoDAsyncr{r cuMemcpyDtoDr/rs r=rrT s G !&&)))))  %v}%%%%  B~cN3//@@@@@@@r?cg}|r>t|tsJtj}||jn tj}|t|||g|RdS)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/rrr=r{rr>r/)rvalrr@rrs r= device_memsetrf sxG &&)))))  #v}%%%%  B~cC0000000r?c8tdS)z; Enable profile collection in the current context. N)rcuProfilerStartrEr?r= profile_startr{ s r?c8tdS)z< Disable profile collection in the current context. N)rcuProfilerStoprEr?r= profile_stopr s r?c#JKtdVtdS)z] Context manager that enables profiling on entry and disables profiling on exit. N)rrrEr?r= profilingr s' OOO EEENNNNNr?c4tS)z@ Return the driver version as a tuple of (major, minor) )rr&rEr?r=r&r& s     r?rr')r(r4r^rhrrrr'r@rrX itertoolsrabcrrrrrr r r r r rrrnumpyr collectionsrrrr numba.corerrrrrrrrrrrnumba.cuda.cudadrvrrrCUDA_USE_NVIDIA_BINDINGrr!rrrrgrrl pythonapi Py_DecRefr Py_IncRefrz py_objectrr>rKrArHrJrsrrrrrfr}rrrrrrrrEr{rrrrrrr rrrr0rrr rrrrrrrrrrrrrrnrrMemAllocrrr9r[rcrrrrrrrrrrr rjittyCU_JIT_INPUT_OBJECTruCU_JIT_INPUT_LIBRARYCU_JIT_INPUT_CUBINCU_JIT_INPUT_FATBINARYr9rrErrrrrrrrrTrrrr/rKrrrrrrrrrrrr&rEr?r=rs    ''''''''//////////////////////))))))))//////////44444444""""""LLLLLLLLLL<<<<<<<<<<<</$$$$$$|&&w//   '   ' '( '( 0     l        ,   11111?111( ( ( V$$$0  111666 % $ & & BBBBBVBBBJ44444V444n 0/11Y1Y1Y1Y1Y1VY1Y1Y1x::: q q q q q Fgq q q q h@@@@@ 5@@@F((((((((2.0.0.0.0.0.0I.0.0.0b$% 777$   *     #   kmm BBBBBvBBBJZ l 3 3  i&i&i&i&i&fi&i&i&X 888<<<> @ @ @F   ..0! ! ! ! ! 6! ! ! H' ' ' ' ' V' ' ' Tv;v;v;v;v;v;v;v;ro-o-o-o-o-Fo-o-o-d     m   +7+7+7+7+7?+7+7+7\&&&&&8$&&&R#8#8#8#8#8O#8#8#8L*****6***>     x'8        ,(9   wwwwwVwwwt7070707070F707070t   CCCCCwCCCC. B B B B B6 B B B <<<<++++++++J$ % % % %F  "E  &%  ')).  &%  ')). O O O O O wO O O O f V%V%V%V%V%V%V%V%r]F]F]F]F]F6]F]F]F@WFWFWFWFWFvWFWFWFz***,&"444HHHH(000 000%%%222555NNNN$????$AAAA$1111*      r?