
\K]c           @` s  d  Z  d d l m Z m Z m Z d d l Z d d l Z d d l Z d d l Z d d l	 Z	 d d l
 Z
 d d l Z d d l Z d d l Z d d l m Z d d l m Z m Z m Z m Z m Z m Z m Z m Z d d l Z d d l Z d d l m Z m Z d d l m Z m  Z  d d	 l! m" Z" m# Z# d d
 l$ m% Z% d d l$ m& Z& d d l' m( Z( m$ Z$ m) Z) d d l m* Z* m+ Z+ m, Z, d d l- m. Z/ d d l0 m1 Z1 m2 Z2 e3 e1 d d   Z4 d d f Z5 e j6 j7 d  Z8 d   Z9 d e: f d     YZ; d e: f d     YZ< d e# f d     YZ= d   Z> d Z? d Z@ d   ZA d   ZB d   ZC d    ZD eC   ZE d! ZF d" eG f d#     YZH d$ eG f d%     YZI eH   ZJ d&   ZK eK   ZL d' eG f d(     YZM d)   ZN d* eG f d+     YZO eO   ZO d, eG f d-     YZP e d. d/  ZQ d0 eG f d1     YZR d2   ZS d3   ZT d4   ZU d5   ZV d6   ZW d7   ZX d8   ZY d9 eG f d:     YZZ d; eG f d<     YZ[ d= eG f d>     YZ\ d? eG f d@     YZ] dA e] f dB     YZ^ dC e^ f dD     YZ_ dE e  j` f dF     YZa dG eG f dH     YZb dI eb e  j` f dJ     YZc dK eG f dL     YZd dM eG f dN     YZe dO   Zf dP eG f dQ     YZg e dR dS dT dU dV dW g  Zh dX eG f dY     YZi dZ   Zj i e( jk d[ 6e( jl d\ 6e( jm d] 6e( jn d^ 6e( jo d_ 6Zp d` eG f da     YZq db   Zr dc   Zs dd   Zt de   Zu df   Zv dg   Zw dh   Zx ey di  Zz dj   Z{ dk   Z| dl   Z} dm   Z~ dn   Z do   Z dp   Z dq   Z d dr  Z d ds  Z d dt  Z d du  Z dv   Z dw   Z e j dx    Z d S(y   s  
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.

i    (   t   absolute_importt   print_functiont   divisionN(   t   product(   t   c_intt   byreft   c_size_tt   c_chart   c_char_pt	   addressoft   c_void_pt   c_float(   t
   namedtuplet   deque(   t   utilst   mviewbufi   (   t   CudaSupportErrort   CudaDriverError(   t   API_PROTOTYPES(   t   cu_occupancy_b2d_size(   t   enumst   drvapit   _extras(   t   configt	   serializet   errors(   t   longint(   t   get_numba_envvart   get_numbapro_envvart   VERBOSE_CU_JIT_LOGi   t   linuxc          C` s   t  j t  }  t j |   s t t j  j   } t	 t  | d   } t | t  s` t  j } n  |  j |  t j r t  j t j  } d } | j t  j d |   |  j |  q |  j t  j    n  |  S(   Ns;   == CUDA [%(relativeCreated)d] %(levelname)5s -- %(message)st   fmt(   t   loggingt	   getLoggert   __name__R   t   logger_hasHandlerst   strR   t   CUDA_LOG_LEVELt   uppert   getattrt   Nonet
   isinstancet   intt   CRITICALt   setLevelt   StreamHandlert   syst   stderrt   setFormattert	   Formattert
   addHandlert   NullHandler(   t   loggert   lvlt   handlerR   (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   _make_logger.   s    	t   DeadMemoryErrorc           B` s   e  Z RS(    (   R"   t
   __module__(    (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyR8   F   s   t   LinkerErrorc           B` s   e  Z RS(    (   R"   R9   (    (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyR:   J   s   t   CudaAPIErrorc           B` s   e  Z d    Z d   Z RS(   c         C` s/   | |  _  | |  _ t t |   j | |  d  S(   N(   t   codet   msgt   superR;   t   __init__(   t   selfR<   R=   (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyR?   O   s    		c         C` s   d |  j  |  j f S(   Ns   [%s] %s(   R<   R=   (   R@   (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   __str__T   s    (   R"   R9   R?   RA   (    (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyR;   N   s   	c          C` s  t  d  }  |  d k r" t   n  t j d k rO t j } d g } d g } nN t j d k r| t j } d g } d g } n! t j } d	 d
 g } d d g } |  d  k	 ry t j	 j
 |   }  Wn! t k
 r t d |    n Xt j	 j |   st d |    n  |  g } n; | g  t | |  D]! \ } } t j	 j | |  ^ q'} g  } g  } x_ | D]W }	 y | |	  }
 Wn: t k
 r} | j t j	 j |	   | j |  qbX|
 SqbWt |  rt   n# d j d   | D  } t |  d  S(   Nt   CUDA_DRIVERt   0t   win32s   \windows\system32s
   nvcuda.dllt   darwins   /usr/local/cuda/libs   libcuda.dylibs   /usr/libs
   /usr/lib64s
   libcuda.sos   libcuda.so.1s(   NUMBA_CUDA_DRIVER %s is not a valid pathso   NUMBA_CUDA_DRIVER %s is not a valid file path.  Note it must be a filepath of the .so/.dll/.dylib or the drivers   
c         s` s   |  ] } t  |  Vq d  S(   N(   R$   (   t   .0t   e(    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pys	   <genexpr>   s    (   R   t   _raise_driver_not_foundR.   t   platformt   ctypest   WinDLLt   CDLLR(   t   ost   patht   abspatht
   ValueErrort   isfileR   t   joint   OSErrort   appendt   allt   _raise_driver_error(   t   envpatht   dlloadert   dldirt   dlnamest
   candidatest   xt   yt   path_not_existt   driver_load_errorRN   t   dllRG   t   errmsg(    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   find_driverX   sN    
					5
s   
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.
sM   
Possible CUDA driver libraries are found but error occurred during load:
%s
c           C` s   t  t   d  S(   N(   R   t   DRIVER_NOT_FOUND_MSG(    (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyRH      s    c         C` s   t  t |    d  S(   N(   R   t   DRIVER_LOAD_ERROR_MSG(   RG   (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyRV      s    c          C` sX   d }  t  j   } x? t t  D]1 } | j |   r t t |  } | | | <q q W| S(   Nt
   CUDA_ERROR(   R   t
   UniqueDictt   dirR   t
   startswithR'   (   t   prefixt   mapt   nameR<   (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   _build_reverse_error_map   s    c           C` s
   t  j   S(   N(   RM   t   getpid(    (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   _getpid   s    s9   driver missing function: %s.
Requires CUDA 8.0 or above.
t   Driverc           B` s   e  Z d  Z d Z d   Z d   Z d   Z d   Z e	 d    Z
 d   Z d   Z d   Z d	   Z d
 d  Z d   Z d   Z d   Z d   Z d   Z RS(   s0   
    Driver API functions are lazily bound.
    c         C` s5   |  j  } | d  k	 r | St j |   } | |  _  | S(   N(   t
   _singletonR(   t   objectt   __new__(   t   clst   obj(    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyRr      s    		c         C` s   t  j   |  _ t |  _ d  |  _ d  |  _ y. t j	 rK d } t
 |   n  t   |  _ Wn% t
 k
 r } t |  _ | |  _ n Xd  S(   Nsz   CUDA is disabled due to setting NUMBA_DISABLE_CUDA=1 in the environment, or because CUDA is unsupported on 32-bit systems.(   R   Rf   t   devicest   Falset   is_initializedR(   t   initialization_errort   pidR   t   DISABLE_CUDAR   Rb   t   libt   True(   R@   R=   RG   (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyR?      s    					c         C` sy   t    a t |  _ y t j d  |  j d  Wn, t k
 r^ } | |  _ t d |   n Xt	   |  _
 |  j   d  S(   Nt   initi    s   Error at driver init: 
%s:(   R7   t   _loggerR|   Rw   t   infot   cuInitR;   Rx   R   Rn   Ry   t   _initialize_extras(   R@   RG   (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt
   initialize   s    			c         C` s   t  j d  t  } | t j  } | |  j d   t  j t t  j t	 j
  t  j t	 j  t  j  } | t j  } d | _ |  j d |  } | |  _ d  S(   Nt   cuIpcOpenMemHandlet   call_cuIpcOpenMemHandle(   RJ   t	   CFUNCTYPER(   R
   R   t   set_cuIpcOpenMemHandlet	   _find_apiR   t   POINTERR   t   cu_device_ptrt   cu_ipc_mem_handlet   c_uintR   R"   t   _wrap_api_callR   (   R@   t	   set_protoR   t
   call_protoR   t	   safe_call(    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyR      s    				c         C` s#   |  j  s |  j   n  |  j d  k S(   N(   Rw   R   Rx   R(   (   R@   (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   is_available  s    	c         C` s   y t  | } Wn t k
 r- t |   n X| d } | d } |  j sX |  j   n  |  j d  k	 r} t d |  j   n  |  j |  } | | _	 | | _
 |  j | |  } t |  | |  | S(   Ni    i   s   Error at driver init: 
%s:(   R   t   KeyErrort   AttributeErrorRw   R   Rx   R(   R   R   t   restypet   argtypesR   t   setattr(   R@   t   fnamet   protoR   R   t   libfnR   (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   __getattr__  s"    

			c         ` s(   t  j       f d    } | S(   Nc          ` s3   t  j d  j   |    }  j   |  d  S(   Ns   call driver api: %s(   R~   t   debugR"   t   _check_error(   t   argst   retcode(   R   R   R@   (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   safe_cuda_api_call"  s    (   t	   functoolst   wraps(   R@   R   R   R   (    (   R   R   R@   s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyR   !  s    $c         ` sw   y t  |  j   d  SWn t k
 r+ n Xy t  |  j    SWn t k
 rS n X  f d   } t |    |  | S(   Nt   _v2c          ` s   t  t     d  S(   N(   R   t   MISSING_FUNCTION_ERRMSG(   R   t   kws(   R   (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   absent_function8  s    (   R'   R{   R   R   (   R@   R   R   (    (   R   s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyR   )  s    c         C` s   | t  j k r t j | d  } d | | f } t j |  | t  j k r |  j d  k	 r t	   |  j k r d } t j
 | t	   |  j  t d   q n  t | |   n  d  S(   Nt   UNKNOWN_CUDA_ERRORs   Call to %s results in %ss0   pid %s forked from pid %s after CUDA driver inits   CUDA initialized before forking(   R   t   CUDA_SUCCESSt	   ERROR_MAPt   getR~   t   errort   CUDA_ERROR_NOT_INITIALIZEDRy   R(   Rn   t   criticalR   R;   (   R@   R   R   t   errnameR=   (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyR   >  s    !i    c         C` sG   |  j  j |  } | d  k r: t |  } | |  j  | <n  t j |  S(   N(   Ru   R   R(   t   Devicet   weakreft   proxy(   R@   t   devnumt   dev(    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt
   get_deviceK  s
    c         C` s#   t    } |  j t |   | j S(   N(   R   t   cuDeviceGetCountR   t   value(   R@   t   count(    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   get_device_countR  s    	c         C` s   t  |  j j    S(   s)   Returns a list of active devices
        (   t   listRu   t   values(   R@   (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   list_devicesW  s    c         C` s(   x! |  j  j   D] } | j   q Wd S(   s   Reset all devices
        N(   Ru   R   t   reset(   R@   R   (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyR   \  s    c         C` sK   |  j    9 } | j d k	 rA t j   } t j t |   | SWd QXd S(   sn   Pop the active CUDA context and return the handle.
        If no CUDA context is active, return None.
        N(   t   get_active_contextR   R(   R   t
   cu_contextt   drivert   cuCtxPopCurrentR   (   R@   t   act   popped(    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   pop_active_contextb  s
    c         C` s   t    S(   s3   Returns an instance of ``_ActiveContext``.
        (   t   _ActiveContext(   R@   (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyR   l  s    N(   R"   R9   t   __doc__R(   Rp   Rr   R?   R   R   t   propertyR   R   R   R   R   R   R   R   R   R   R   (    (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyRo      s"   													
R   c           B` s;   e  Z d  Z e j   Z d   Z d   Z d   Z e Z	 RS(   s   An 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.
    c         C` s   t  } t |  j d  r- |  j j \ } } n t j d  } t j t |   | j	 r^ | n d  } | d  k ry d  } n@ t j   } t j t |   | j	 } | | f |  j _ t } | |  _ | |  _ | |  _ |  S(   Nt
   ctx_devnumi    (   Rv   t   hasattrt
   _tls_cacheR   R   R   R   t   cuCtxGetCurrentR   R   R(   t	   cu_devicet   cuCtxGetDeviceR|   t   _is_topt   context_handleR   (   R@   t   is_topt   hctxR   t   hdevice(    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt	   __enter__{  s"    					c         C` s    |  j  r t |  j d  n  d  S(   NR   (   R   t   delattrR   (   R@   t   exc_typet   exc_valt   exc_tb(    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   __exit__  s    	c         C` s   |  j  d k	 S(   sA   Returns True is there's a valid and active CUDA context.
        N(   R   R(   (   R@   (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   __bool__  s    (
   R"   R9   R   t	   threadingt   localR   R   R   R   t   __nonzero__(    (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyR   r  s   			c          C` s\   d }  t  j   } xC t t  D]5 } | j |   r t t |  | | t |   <q q W| S(   Nt   CU_DEVICE_ATTRIBUTE_(   R   Rf   Rg   R   Rh   R'   t   len(   Ri   Rj   Rk   (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   _build_reverse_device_attrs  s    $R   c           B` s   e  Z d  Z e d    Z d   Z d   Z e d    Z d   Z	 d   Z
 d   Z d   Z d	   Z d
   Z d   Z d   Z RS(   s   
    The device object owns the CUDA contexts.  This is owned by the driver
    object.  User should not construct devices directly.
    c         C` sa   xZ t  t j    D]+ } t j |  } | j   | k r | Sq Wd j |  } t |   d S(   sg   Create Device object from device identity created by
        ``Device.get_device_identity()``.
        sK   No device of {} is found. Target device may not be visible in this process.N(   t   rangeR   R   R   t   get_device_identityt   formatt   RuntimeError(   R@   t   identityt   devidt   dRa   (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   from_identity  s    	c         C` s   t    } t j t |  |  | | j k s: t d   | j |  _ i  |  _ t    } t    } t j t |  t |  |  j  | j | j f |  _	 d } t
 |   } t j | | |  j  | j |  _ d  |  _ d  S(   Ns   Driver returned another devicei   (   R   R   t   cuDeviceGetR   R   t   AssertionErrort   idt
   attributest   cuDeviceComputeCapabilityt   compute_capabilityR   t   cuDeviceGetNameRk   R(   t   primary_context(   R@   R   t
   got_devnumt   cc_majort   cc_minort   bufszt   buf(    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyR?     s    				
c         C` s"   i |  j  d 6|  j d 6|  j d 6S(   Nt   pci_domain_idt
   pci_bus_idt   pci_device_id(   t   PCI_DOMAIN_IDt
   PCI_BUS_IDt   PCI_DEVICE_ID(   R@   (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyR     s    

c         C` s   t  j d t  |  j S(   s,   
        For backward compatibility
        sA   Deprecated attribute 'COMPUTE_CAPABILITY'; use lower case version(   t   warningst   warnt   DeprecationWarningR   (   R@   (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   COMPUTE_CAPABILITY  s    	c         C` s   d |  j  |  j f S(   Ns   <CUDA device %d '%s'>(   R   Rk   (   R@   (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   __repr__  s    c         C` sm   y t  | } Wn t k
 r- t |   n Xt   } t j t |  | |  j  t |  | | j	  | j	 S(   s   Read attributes lazily
        (
   t   DEVICE_ATTRIBUTESR   R   R   R   t   cuDeviceGetAttributeR   R   R   R   (   R@   t   attrR<   R   (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyR     s    	c         C` s   t  |  j  S(   N(   t   hashR   (   R@   (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   __hash__  s    c         C` s#   t  | t  r |  j | j k St S(   N(   R)   R   R   Rv   (   R@   t   other(    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   __eq__  s    c         C` s   |  | k S(   N(    (   R@   R  (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   __ne__  s    c         C` sj   |  j  d k	 r |  j  St |   t j   } t j t |  |  j  t	 t
 j |   |  } | |  _  | S(   so   
        Returns the primary context for the device.
        Note: it is not pushed to the CPU thread.
        N(   R   R(   t   met_requirement_for_deviceR   R   R   t   cuDevicePrimaryCtxRetainR   R   t   ContextR   R   (   R@   R   t   ctx(    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   get_primary_context  s    
	c         C` s   t  j |  j  d |  _ d S(   s6   
        Release reference to primary context
        N(   R   t   cuDevicePrimaryCtxReleaseR   R(   R   (   R@   (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   release_primary_context  s    c         C` sE   z- |  j  d  k	 r" |  j  j   n  |  j   Wd  t j |  j  Xd  S(   N(   R   R(   R   R  R   t   cuDevicePrimaryCtxResetR   (   R@   (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyR     s
    (   R"   R9   R   t   classmethodR   R?   R   R   R   R   R   R  R  R  R	  R  R   (    (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyR     s   										c         C` s,   |  j  t k  r( t d |  t f   n  d  S(   Ns   %s has compute capability < %s(   R   t   MIN_REQUIRED_CCR   (   t   device(    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyR  '  s    t   _SizeNotSetc           B` s    e  Z d  Z d   Z d   Z RS(   sC   
    Dummy object for _PendingDeallocs when *size* is not set.
    c         C` s   d S(   Nt   ?(    (   R@   (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyRA   1  s    c         C` s   d S(   Ni    (    (   R@   (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   __int__4  s    (   R"   R9   R   RA   R  (    (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyR  -  s   	t   _PendingDeallocsc           B` se   e  Z d  Z d   Z e d    Z e d  Z d   Z e	 j
 d    Z e d    Z d   Z RS(   sd   
    Pending deallocations of a context (or device since we are using the primary
    context).
    c         C` s+   t    |  _ d |  _ d |  _ | |  _ d  S(   Ni    (   R   t   _const   _disable_countt   _sizet   _memory_capacity(   R@   t   capacity(    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyR?   ?  s    		c         C` s   t  |  j t j  S(   N(   R*   R  R   t   CUDA_DEALLOCS_RATIO(   R@   (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   _max_pending_bytesE  s    c         C` s   t  j d | j |  |  j j | | | f  |  j t |  7_ t |  j  t j	 k sn |  j |  j
 k r{ |  j   n  d S(   s_  
        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.
        s    add pending dealloc: %s %s bytesN(   R~   R   R"   R  RT   R  R*   R   R   t   CUDA_DEALLOCS_COUNTR  t   clear(   R@   t   dtort   handlet   size(    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   add_itemI  s    	c         C` sa   |  j  s] xE |  j rP |  j j   \ } } } t j d | j |  | |  q Wd |  _ n  d S(   sh   
        Flush any pending deallocations unless it is disabled.
        Do nothing if disabled.
        s   dealloc: %s %s bytesi    N(   t   is_disabledR  t   popleftR~   R   R"   R  (   R@   R  R  R  (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyR  Y  s    	c         c` sD   |  j  d 7_  z	 d VWd |  j  d 8_  |  j  d k s? t  Xd S(   ss   
        Context manager to temporarily disable flushing pending deallocation.
        This can be nested.
        i   Ni    (   R  R   (   R@   (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   disablee  s
    	c         C` s   |  j  d k S(   Ni    (   R  (   R@   (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyR!  r  s    c         C` s   t  |  j  S(   s:   
        Returns number of pending deallocations.
        (   R   R  (   R@   (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   __len__v  s    (   R"   R9   R   R?   R   R  R  R   R  t
   contextlibt   contextmanagerR#  R!  R$  (    (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyR  :  s   		t   _MemoryInfos
   free,totalR  c           B` s  e  Z d  Z d   Z d   Z d   Z d d  Z d d  Z d   Z	 d   Z
 d   Z d	   Z d
   Z e e e d  Z e d  Z d   Z d   Z d   Z d d  Z d   Z d   Z d   Z d   Z d   Z e d  Z d   Z d   Z d   Z d   Z RS(   ss   
    This object wraps a CUDA Context resource.

    Contexts should not be constructed directly by user code.
    c         C` sF   | |  _  | |  _ t j   |  _ d  |  _ t j   |  _ i  |  _ d  S(   N(	   R  R  R   Rf   t   allocationsR(   t   deallocationst   modulest   extras(   R@   R  R  (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyR?     s    			c         C` sM   t  j d |  j j  |  j j   |  j j   |  j rI |  j j   n  d S(   s?   
        Clean up all owned resources in this context.
        s   reset context of device %sN(   R~   R   R  R   R(  R  R*  R)  (   R@   (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyR     s
    	c         C` sG   t    } t    } t j t |  t |   t d | j d | j  S(   s>   Returns (free, total) memory in bytes in the context.
        t   freet   total(   R   R   t   cuMemGetInfoR   R'  R   (   R@   R,  R-  (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   get_memory_info  s    		c         C` sZ   t    } | s1 t j t |  | j | |  n" t j t |  | j | | |  | j S(   s   Return 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   R   t+   cuOccupancyMaxActiveBlocksPerMultiprocessorR   R  t4   cuOccupancyMaxActiveBlocksPerMultiprocessorWithFlagsR   (   R@   t   funct	   blocksizet   memsizet   flagst   retval(    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt$   get_active_blocks_per_multiprocessor  s
    	""c   	      C` s   t    } t    } t |  } | sR t j t |  t |  | j | | |  n. t j t |  t |  | j | | | |  | j | j f S(   s  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.
        :param memsize: per-block dynamic shared memory usage intended, in bytes
        :param blocksizelimit: maximum block size the kernel is designed to handle(   R   R   R   t    cuOccupancyMaxPotentialBlockSizeR   R  t)   cuOccupancyMaxPotentialBlockSizeWithFlagsR   (	   R@   R2  t   b2d_funcR4  t   blocksizelimitR5  t   gridsizeR3  t   b2d_cb(    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   get_max_potential_block_size  s    			c         C` s.   |  j  d k r* t |  j   j  |  _  n  d S(   sW   Initialize the context for use.
        It's safe to be called multiple times.
        N(   R)  R(   R  R/  R-  (   R@   (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   prepare_for_use  s    c         C` s   t  j |  j  |  j   d S(   s@   
        Pushes this context on the current CPU Thread.
        N(   R   t   cuCtxPushCurrentR  R?  (   R@   (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   push  s    c         C` s+   t  j   } | j |  j j k s' t  d S(   s   
        Pops this context off the current CPU thread. Note that this context must
        be at the top of the context stack, otherwise an error will occur.
        N(   R   R   R   R  R   (   R@   R   (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   pop  s    c         C` sQ   y |   Wn? t  k
 rL } | j t j k rF |  j j   |   qM   n Xd S(   s   
        Attempt allocation by calling *allocator*.  If a 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(   R;   R<   R   t   CUDA_ERROR_OUT_OF_MEMORYR)  R  (   R@   t	   allocatorRG   (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   _attempt_allocation  s    
c         ` su   t  j       f d   } |  j |  t |      } t t j |      |  } | |  j  j <| j	   S(   Nc           ` s   t  j t      d  S(   N(   R   t
   cuMemAllocR   (    (   t   bytesizet   ptr(    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyRD    s    (
   R   R   RE  t   _alloc_finalizert   AutoFreePointerR   R   R(  R   t   own(   R@   RG  RD  t	   finalizert   mem(    (   RG  RH  s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   memalloc  s    c   	      ` s  t     d  | r%  t j O n  | r;  t j O n  | rQ  t j O n      f d   } | r| |  j |  n |   d  } t |     |  } | r t t	 j
 |   |    d | } | |  j | j j <| j   St t	 j
 |   |    d | } | Sd  S(   Ni    c           ` s   t  j t       d  S(   N(   R   t   cuMemHostAllocR   (    (   RG  R5  t   pointer(    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyRD    s    RL  (   R
   R   t   CU_MEMHOSTALLOC_DEVICEMAPt   CU_MEMHOSTALLOC_PORTABLEt   CU_MEMHOSTALLOC_WRITECOMBINEDRE  R(   t   _hostalloc_finalizert   MappedMemoryR   R   R(  R  R   RK  t   PinnedMemory(	   R@   RG  t   mappedt   portablet   wcRD  t   ownerRL  RM  (    (   RG  R5  RP  s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   memhostalloc  s,    		
	c         ` s   t   t t f  r$ t    n  | rM |  j j rM t d |  j   n  d   | ri   t j O  n      f d   } | r |  j	 |  n |   t
 |   |  } | r t t j |   |   d | } | |  j | j j <| j   St t j |   |   d | } | Sd  S(   Ns   %s cannot map host memoryi    c           ` s   t  j      d  S(   N(   R   t   cuMemHostRegister(    (   R5  RP  R  (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyRD  ,  s    RL  (   R)   R*   t   longR
   R  t   CAN_MAP_HOST_MEMORYR   R   t   CU_MEMHOSTREGISTER_DEVICEMAPRE  t   _pin_finalizerRU  R   R   R(  R  R   RK  RV  (   R@   RZ  RP  R  RW  RD  RL  RM  (    (   R5  RP  R  s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   mempin  s(    	
	c         C` s
   t   d  S(   N(   t   NotImplementedError(   R@   RP  (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   memunpin@  s    c         C` s   t  s t d   n  t j   } t j t j |  | j j	  |  j
 j   } | j	 j | j j	 j } t | | | j | d | S(   s>   
        Returns a *IpcHandle* from a GPU allocation.
        s   OS does not support CUDA IPCt   offset(   t   SUPPORTS_IPCRS   R   R   R   t   cuIpcGetMemHandleRJ   R   RZ  R  R  R   R   t	   IpcHandleR  (   R@   t   memoryt	   ipchandlet   source_infoRd  (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   get_ipc_handleC  s    c         C` sM   t  j   } d } t j t |  | |  t d t j |   d | d |  S(   Ni   t   contextRP  R  (   R   R   R   R   R   t   MemoryPointerR   R   (   R@   R  R  t   dptrR5  (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   open_ipc_handleS  s
    i    c         C` s,   | d k s t  d   t j | |  d S(   sL   Enable peer access between the current context and the peer context
        i    s$   *flags* is reserved and MUST be zeroN(   R   R   t   cuCtxEnablePeerAccess(   R@   t   peer_contextR5  (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   enable_peer_access\  s    c         C` s2   t    } t j t |  |  j j |  t |  S(   ss   Returns a bool indicating whether the peer access between the
        current and peer device is possible.
        (   R   R   t   cuDeviceCanAccessPeerR   R  R   t   bool(   R@   t   peer_devicet   can_access_peer(    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyRv  b  s    			c         C` s:   t  | t  r! | j d  } n  t |  } |  j |  S(   Nt   utf8(   R)   R$   t   encodeR   t   create_module_image(   R@   t   ptxt   image(    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   create_module_ptxn  s    c         C` s/   t  |  |  } | |  j | j j <t j |  S(   N(   t   load_module_imageR*  R  R   R   R   (   R@   R{  t   module(    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyRy  t  s    c         C` s   |  j  | j j =d  S(   N(   R*  R  R   (   R@   R~  (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   unload_moduley  s    c         C` sG   t  j   } t j t |  d  t t j |   | t |  j	 |   S(   Ni    (
   R   t	   cu_streamR   t   cuStreamCreateR   t   StreamR   R   t   _stream_finalizerR)  (   R@   R  (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   create_stream|  s    c         C` sf   t  j   } d } | s( | t j O} n  t j t |  |  t t j	 |   | d t
 |  j |  S(   Ni    RL  (   R   t   cu_eventR   t   CU_EVENT_DISABLE_TIMINGR   t   cuEventCreateR   t   EventR   R   t   _event_finalizerR)  (   R@   t   timingR  R5  (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   create_event  s    c         C` s   t  j   d  S(   N(   R   t   cuCtxSynchronize(   R@   (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   synchronize  s    c         C` s   d |  j  |  j j f S(   Ns   <CUDA context %s of device %d>(   R  R  R   (   R@   (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyR     s    c         C` s'   t  | t  r |  j | j k St Sd  S(   N(   R)   R  R  t   NotImplemented(   R@   R  (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyR    s    c         C` s   |  j  |  S(   N(   R  (   R@   R  (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyR    s    N(    R"   R9   R   R?   R   R/  R(   R7  R>  R?  RA  RB  RE  RN  Rv   R[  Ra  Rc  Rk  Ro  Rr  Rv  R|  Ry  R  R  R|   R  R  R   R  R  (    (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyR    s6   	
							 #													c         C` sa  t  t d d   } t |   } t |   } i t |  t j 6t |  t j 6t |  t j 6t |  t j	 6t t
  t j 6} t j t |  | j     } t t |  | j     } t j   } y) t j t |  | t |  | |  Wn; t k
 r2}	 d | j j d  }
 t |	 j |
   n X| j } t t j |   | | t |  |   S(   s!   
    image must be a pointer
    t   CUDA_LOG_SIZEi   s   cuModuleLoadDataEx error:
%sRw  (   R*   R   R   R	   R   t   CU_JIT_INFO_LOG_BUFFERR
   t!   CU_JIT_INFO_LOG_BUFFER_SIZE_BYTESt   CU_JIT_ERROR_LOG_BUFFERt"   CU_JIT_ERROR_LOG_BUFFER_SIZE_BYTESt   VERBOSE_JIT_LOGt   CU_JIT_LOG_VERBOSER   t   cu_jit_optionR   t   keysR   t	   cu_moduleR   t   cuModuleLoadDataExR   R;   R   t   decodeR<   t   ModuleR   R   t   _module_finalizer(   Rl  R{  t   logszt   jitinfot	   jiterrorst   optionst   option_keyst   option_valsR  RG   R=   t   info_log(    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyR}    s*    	c         ` s.   |  j    |  j       f d   } | S(   Nc           ` s-     r    j  =n   j t j    d  S(   N(   R   R   R   t	   cuMemFree(    (   R(  R)  R  R  (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   core  s    (   R(  R)  (   Rl  R  R  R  (    (   R(  R)  R  R  s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyRI    s    		c         ` s@   |  j    |  j   s! t  n        f d   } | S(   s[  
    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.

    c           ` s3    r   r    j  =n   j t j    d  S(   N(   R   R   R   t   cuMemFreeHost(    (   R(  R)  R  RW  R  (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyR    s    (   R(  R)  R  (   Rl  R  R  RW  R  (    (   R(  R)  R  RW  R  s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyRT    s    
			c         ` s"   |  j        f d   } | S(   sB  
    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.

    c           ` s*    r   r    j  =n  t j   d  S(   N(   R   R   t   cuMemHostUnregister(    (   R(  R  RW  (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyR    s    (   R(  (   Rl  R  RW  R  (    (   R(  R  RW  s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyR`    s    	c         ` s      f d   } | S(   Nc           ` s     j  t j   d  S(   N(   R   R   t   cuEventDestroy(    (   t   deallocsR  (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyR    s    (    (   R  R  R  (    (   R  R  s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyR    s    c         ` s      f d   } | S(   Nc           ` s     j  t j   d  S(   N(   R   R   t   cuStreamDestroy(    (   R  R  (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyR    s    (    (   R  R  R  (    (   R  R  s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyR    s    c         ` s+   |  j    |  j      f d   } | S(   Nc          ` s/   t  j      f d   }   j |    d  S(   Nc         ` s/      s |  j    k s t  t j |   d  S(   N(   R   R   R   t   cuModuleUnload(   R  (   R*  t   shutting_down(    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   module_unload  s    (   R   R  R   (   R  (   t   deallocR  R*  (   R  s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyR  	  s    	(   R)  R*  (   Rl  R  R  (    (   R  R  R*  s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyR    s    		t   _CudaIpcImplc           B` s)   e  Z d  Z d   Z d   Z d   Z RS(   sj   Implementation of GPU IPC using CUDA driver API.
    This requires the devices to be peer accessible.
    c         C` s=   | j  |  _  | j |  _ | j |  _ | j |  _ d  |  _ d  S(   N(   t   baseR  R  Rd  R(   t   _opened_mem(   R@   t   parent(    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyR?     s
    c         C` sz   |  j  d k	 r t d   n  |  j d k	 r< t d   n  | j |  j |  j |  j  } | |  _ | j   j	 |  j  S(   sT   
        Import the IPC memory and returns a raw CUDA memory pointer object
        s'   opening IpcHandle from original processs   IpcHandle is already openedN(
   R  R(   RP   R  Ro  R  Rd  R  RK  t   view(   R@   Rl  RM  (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   open$  s    	c         C` s>   |  j  d  k r t d   n  t j |  j  j  d  |  _  d  S(   Ns   IpcHandle not opened(   R  R(   RP   R   t   cuIpcCloseMemHandleR  (   R@   (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   close5  s    (   R"   R9   R   R?   R  R  (    (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyR    s   		t   _StagedIpcImplc           B` s)   e  Z d  Z d   Z d   Z d   Z RS(   s   Implementation of GPU IPC using custom staging logic to workaround
    CUDA IPC limitation on peer accessibility between devices.
    c         C` s:   | |  _  | j |  _ | j |  _ | j |  _ | |  _ d  S(   N(   R  R  R  R  Rj  (   R@   R  Rj  (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyR?   @  s
    	c         C` s   d d l  m } t j |  j  } t d |  j  } | j | j  | j	 | j
 j    } Wd  QX| j |  j  } t | | |  j  | j | j  | j   Wd  QX| j   S(   Ni    (   t   cudaR  (   t   numbaR  R   R   Rj  R  R  t   gpusR   R  Ru   t   get_contextRN  R  t   device_to_deviceR  RK  (   R@   Rl  R  t   srcdevt   implt
   source_ptrt   newmem(    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyR  G  s    c         C` s   d  S(   N(    (   R@   (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyR  ]  s    (   R"   R9   R   R?   R  R  (    (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyR  <  s   		Rg  c           B` sw   e  Z d  Z d d d  Z d   Z d   Z d   Z d   Z d   Z	 d d  Z
 d	   Z d
   Z e d    Z RS(   s  
    Internal IPC handle.

    Serialization of the CUDA IPC handle object is implemented here.

    The *base* attribute is a reference to the original allocation to keep it
    alive.  The *handle* is a ctypes object of the CUDA IPC handle. The *size*
    is the allocation size.
    i    c         C` s:   | |  _  | |  _ | |  _ | |  _ d  |  _ | |  _ d  S(   N(   R  R  R  Rj  R(   t   _implRd  (   R@   R  R  R  Rj  Rd  (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyR?   l  s    					c         C` s"   |  j  d  k r t d   n  d  S(   Ns#   IPC handle doesn't have source info(   Rj  R(   R   (   R@   (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   _sentry_source_infot  s    c         C` sH   |  j    |  j | j j   k r& t St j |  j  } | j | j  S(   sd   Returns a bool indicating whether the active context can peer
        access the IPC handle
        (	   R  Rj  R  R   R|   R   R   Rv  R   (   R@   Rl  t   source_device(    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyRv  x  s
    
c         C` sM   |  j    |  j d k	 r( t d   n  t |  |  j  |  _ |  j j |  S(   sC   Open the IPC by allowing staging on the host memory first.
        s   IpcHandle is already openedN(   R  R  R(   RP   R  Rj  R  (   R@   Rl  (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   open_staged  s
    
c         C` s=   |  j  d k	 r t d   n  t |   |  _  |  j  j |  S(   sT   
        Import the IPC memory and returns a raw CUDA memory pointer object
        s   IpcHandle is already openedN(   R  R(   RP   R  R  (   R@   Rl  (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   open_direct  s    c         C` s=   |  j  d k s |  j |  r* |  j } n	 |  j } | |  S(   s  Open 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().
        N(   Rj  R(   Rv  R  R  (   R@   Rl  t   fn(    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyR    s    		c      	   C` sY   d d l  m } | d k r( | j } n  |  j |  } | j d | d | d | d |  S(   sC   
        Simliar to `.open()` but returns an device array.
        i   (   t   devicearrayt   shapet   stridest   dtypet   gpu_dataN(   t    R  R(   t   itemsizeR  t   DeviceNDArray(   R@   Rl  R  R  R  R  Rn  (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt
   open_array  s    c         C` s8   |  j  d  k r t d   n  |  j  j   d  |  _  d  S(   Ns   IpcHandle not opened(   R  R(   RP   R  (   R@   (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyR    s    c         C` s=   t  |  j  } |  j | |  j |  j |  j f } t j | f S(   N(   t   tupleR  t	   __class__R  Rj  Rd  R   t   _rebuild_reduction(   R@   t   preprocessed_handleR   (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt
   __reduce__  s    c         C` s4   t  j |   } |  d d  d | d | d | d |  S(   NR  R  R  Rj  Rd  (   R   R   R(   (   Rs   t
   handle_aryR  Rj  Rd  R  (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   _rebuild  s    N(   R"   R9   R   R(   R?   R  Rv  R  R  R  R  R  R  R  R  (    (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyRg  b  s   			
		
			Rm  c           B` sn   e  Z d  Z e Z d	 d	 d  Z e d    Z d   Z	 d   Z
 d	 d d  Z d	 d  Z e d    Z RS(
   so  A memory pointer that owns the buffer with an optional finalizer.

    When an instance is deleted, the finalizer will be called regardless
    of the `.refct`.

    An instance is created with `.refct=1`.  The buffer lifetime
    is tied to the MemoryPointer instance's lifetime.  The finalizer is invoked
    only if the MemoryPointer instance's lifetime ends.
    c         C` sy   | |  _  | |  _ | |  _ | |  _ | d  k	 |  _ d |  _ |  j |  _ | |  _ | d  k	 ru t	 j
 |  |  |  _ n  d  S(   Ni   (   Rl  t   device_pointerR  t   _cuda_memsize_R(   t
   is_managedt   refctR  t   _ownerR   t   finalizet
   _finalizer(   R@   Rl  RP  R  RL  RZ  (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyR?     s    						c         C` s   |  j  d  k r |  S|  j  S(   N(   R  R(   (   R@   (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyRZ    s    c         C` s   t  t j |    S(   N(   t   OwnedPointerR   R   (   R@   (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyRK    s    c         C` sH   |  j  rD |  j j s$ t d   n  |  j   |  j j sD t  n  d S(   s8   
        Forces the device memory to the trash.
        s   Freeing dead memoryN(   R  R  t   aliveR   R   (   R@   (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyR,    s
    	
i    c         C` sZ   | d  k r |  j n | } | r@ t j |  j | | | j  n t j |  j | |  d  S(   N(   R(   R  R   t   cuMemsetD8AsyncR  R  t
   cuMemsetD8(   R@   t   byteR   t   stream(    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   memset  s
    c         C` s   | d  k r |  j | } n
 | | } |  j j d  k r\ | d k rS t d   n  |  } nX |  j j | } | d k  r t d   n  t j |  } t |  j | | d |  j	 } t
 |  j	 t t f  r t t j |  j	  |  S| Sd  S(   Ni    s    non-empty slice into empty slices   size cannot be negativeRZ  (   R(   R  R  R   R   R   R   Rm  Rl  RZ  R)   R  R   R   (   R@   t   startt   stopR  R  R  RP  (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyR    s    
	c         C` s   |  j  S(   N(   R  (   R@   (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   device_ctypes_pointer  s    N(   R"   R9   R   R|   t   __cuda_memory__R(   R?   R   RZ  RK  R,  R  R  R  (    (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyRm    s   			
RJ  c           B` s   e  Z d  Z d   Z RS(   s   Modifies the ownership semantic of the MemoryPointer so that the
    instance lifetime is directly tied to the number of references.

    When `.refct` reaches zero, the finalizer is invoked.
    c         O` s,   t  t |   j | |   |  j d 8_ d  S(   Ni   (   R>   RJ  R?   R  (   R@   R   t   kwargs(    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyR?   $  s    (   R"   R9   R   R?   (    (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyRJ    s   RU  c           B` s#   e  Z e Z d d   Z d   Z RS(   c         C` s   | |  _  | |  _ t j   } t j t |  | d  | |  _ t t	 |   j
 | | | d | |  j |  _ |  j |  _ |  j j |  _ d  S(   Ni    RL  (   t   ownedt   host_pointerR   R   R   t   cuMemHostGetDevicePointerR   R  R>   RU  R?   R  R  t   _buflen_R   t   _bufptr_(   R@   Rl  RZ  t   hostpointerR  RL  t   devptr(    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyR?   .  s    			c         C` s   t  t j |    S(   N(   t   MappedOwnedPointerR   R   (   R@   (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyRK  =  s    N(   R"   R9   R|   R  R(   R?   RK  (    (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyRU  +  s   RV  c           B` s   e  Z d d   Z d   Z RS(   c         C` s}   | |  _  | |  _ | |  _ | |  _ | d  k	 |  _ |  j |  _ |  j |  _ |  j j |  _	 | d  k	 ry t
 j |  |  n  d  S(   N(   Rl  R  R  R  R(   R  R  R  R   R  R   R  (   R@   Rl  RZ  RP  R  RL  (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyR?   B  s    				c         C` s   |  S(   N(    (   R@   (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyRK  Q  s    N(   R"   R9   R(   R?   RK  (    (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyRV  A  s   R  c           B` s   e  Z d d   Z d   Z RS(   c         ` s{   | |  _  | d  k r$ |  j  |  _ n | j s4 t  | |  _ |  j      f d   } |  j  j d 7_ t j |  |  d  S(   Nc           ` s\   yD   j  d 8_    j  d k s' t    j  d k rC   j   n  Wn t k
 rW n Xd  S(   Ni   i    (   R  R   R,  t   ReferenceError(    (   RM  (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   derefa  s    i   (   t   _memR(   t   _viewR  R   R  R   R  (   R@   t   memptrR  R  (    (   RM  s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyR?   V  s    			
c         C` s   t  |  j |  S(   s$   Proxy MemoryPointer methods
        (   R'   R  (   R@   R   (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyR   n  s    N(   R"   R9   R(   R?   R   (    (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyR  U  s   R  c           B` s   e  Z RS(    (   R"   R9   (    (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyR  t  s   R  c           B` s>   e  Z d    Z d   Z d   Z d   Z e j d    Z RS(   c         C` s5   | |  _  | |  _ | d  k	 r1 t j |  |  n  d  S(   N(   Rl  R  R(   R   R  (   R@   Rl  R  RL  (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyR?   y  s    		c         C` s
   |  j  j S(   N(   R  R   (   R@   (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyR    s    c         C` s   d |  j  j |  j f S(   Ns   <CUDA stream %d on %s>(   R  R   Rl  (   R@   (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyR     s    c         C` s   t  j |  j  d S(   sy   
        Wait for all commands in this stream to execute. This will commit any
        pending memory transfers.
        N(   R   t   cuStreamSynchronizeR  (   R@   (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyR    s    c         c` s   |  V|  j    d S(   s   
        A context manager that waits for all commands in this stream to execute
        and commits any pending memory transfers upon exiting the context.
        N(   R  (   R@   (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   auto_synchronize  s    (	   R"   R9   R?   R  R   R  R%  R&  R  (    (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyR  x  s
   				R  c           B` sG   e  Z d d   Z d   Z d d  Z d   Z d d  Z d   Z RS(   c         C` s5   | |  _  | |  _ | d  k	 r1 t j |  |  n  d  S(   N(   Rl  R  R(   R   R  (   R@   Rl  R  RL  (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyR?     s    		c         C` sK   y t  j |  j  Wn, t k
 rB } | j t j k r< t S  n Xt Sd S(   sy   
        Returns True if all work before the most recent record has completed;
        otherwise, returns False.
        N(	   R   t   cuEventQueryR  R;   R<   R   t   CUDA_ERROR_NOT_READYRv   R|   (   R@   RG   (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   query  s    i    c         C` s,   | r | j  n d } t j |  j  |  d S(   s  
        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.
        i    N(   R  R   t   cuEventRecord(   R@   R  t   hstream(    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   record  s    	c         C` s   t  j |  j  d S(   sN   
        Synchronize the host thread for the completion of the event.
        N(   R   t   cuEventSynchronizeR  (   R@   (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyR    s    c         C` s5   | r | j  n d } d } t j | |  j  |  d S(   sZ   
        All future works submitted to stream will wait util the event completes.
        i    N(   R  R   t   cuStreamWaitEvent(   R@   R  R  R5  (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   wait  s    c         C` s   t  |  |  S(   N(   t   event_elapsed_time(   R@   t   evtend(    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   elapsed_time  s    N(	   R"   R9   R(   R?   R  R  R  R  R  (    (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyR    s   		c         C` s/   t    } t j t |  |  j | j  | j S(   sF   
    Compute the elapsed time between two events in milliseconds.
    (   R   R   t   cuEventElapsedTimeR   R  R   (   t   evtstartR  t   msec(    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyR    s    	R  c           B` s/   e  Z d d   Z d   Z d   Z d   Z RS(   c         C` sC   | |  _  | |  _ | |  _ | d  k	 r? t j |  |  |  _ n  d  S(   N(   Rl  R  R  R(   R   R  R  (   R@   Rl  R  R  RL  (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyR?     s
    			c         C` s   |  j  j |   d  S(   N(   Rl  R  (   R@   (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   unload  s    c         C` sJ   t  j   } t j t |  |  j | j d   t t j	 |   | |  S(   NRw  (
   R   t   cu_functionR   t   cuModuleGetFunctionR   R  Rx  t   FunctionR   R   (   R@   Rk   R  (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   get_function  s    c         C` sb   t  j   } t  j   } t j t |  t |  |  j | j d   t |  j	 | |  | j
 f S(   NRw  (   R   R   R   R   t   cuModuleGetGlobalR   R  Rx  Rm  Rl  R   (   R@   Rk   RH  R  (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   get_global_symbol  s
    N(   R"   R9   R(   R?   R  R  R  (    (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyR    s   		t   FuncAttrt   regst   sharedR   t   constt
   maxthreadsR  c           B` s}   e  Z d
 Z d Z d Z d Z d   Z d   Z e e e d  Z	 d d d  Z
 d   Z e d    Z d   Z d	   Z RS(   i   i    c         C` s.   | |  _  | |  _ | |  _ |  j   |  _ d  S(   N(   R~  R  Rk   t   _read_func_attr_allt   attrs(   R@   R~  R  Rk   (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyR?     s    			c         C` s   d |  j  S(   Ns   <CUDA function %s>(   Rk   (   R@   (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyR     s    c         C` sh   | p | o | } | r$ t  j } n- | r6 t  j } n | rH t  j } n	 t  j } t j |  j |  d  S(   N(   R   t   CU_FUNC_CACHE_PREFER_EQUALt   CU_FUNC_CACHE_PREFER_L1t   CU_FUNC_CACHE_PREFER_SHAREDt   CU_FUNC_CACHE_PREFER_NONER   t   cuFuncSetCacheConfigR  (   R@   t   prefer_equalt   prefer_cachet   prefer_sharedt   flag(    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   cache_config  s    	c         C` s   x  t  |  d k  r" | d 7} q Wx  t  |  d k  rE | d 7} q& Wt j |   } | | _ | | _ | | _ | r | | _ n	 d | _ | S(   Ni   i   i    (   i   (   i   (   R   t   copyt   griddimt   blockdimt	   sharedmemR  (   R@   R!  R"  R#  R  t   inst(    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt	   configure  s    				c         G` sG   |  j  r |  j  j } n d } t |  j |  j |  j |  j | |  d S(   sS   
        *args -- Must be either ctype objects of DevicePointer instances.
        N(   R  R  R(   t   launch_kernelR!  R"  R#  (   R@   R   t   streamhandle(    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   __call__  s
    	c         C` s   |  j  j j S(   N(   R~  Rl  R  (   R@   (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyR  )  s    c         C` s,   t    } t j t |  | |  j  | j S(   s,   
        Read CUfunction attributes
        (   R   R   t   cuFuncGetAttributeR   R  R   (   R@   t   attridR6  (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   _read_func_attr-  s    	c         C` s   |  j  t j  } |  j  t j  } |  j  t j  } |  j  t j  } |  j  t j  } t d | d | d | d | d |  S(   NR  R  R   R  R  (   R+  R   t   CU_FUNC_ATTRIBUTE_NUM_REGSt"   CU_FUNC_ATTRIBUTE_CONST_SIZE_BYTESt"   CU_FUNC_ATTRIBUTE_LOCAL_SIZE_BYTESt#   CU_FUNC_ATTRIBUTE_SHARED_SIZE_BYTESt'   CU_FUNC_ATTRIBUTE_MAX_THREADS_PER_BLOCKR  (   R@   t   nregst   cmemt   lmemt   smemt   maxtpb(    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyR  5  s    (   i   i   i   (   i   i   i   (   R"   R9   R!  R"  R  R#  R?   R   Rv   R  R%  R(  R   R  R+  R  (    (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyR    s   				c         C` s   | \ } } } | \ }	 }
 } g  } xI | D]A } t  |  rY | j t t |    q+ | j t |   q+ Wt t |  |   } t j |  | | | |	 |
 | | | | d   d  S(   N(	   t   is_device_memoryRT   R	   R  R
   R   R   t   cuLaunchKernelR(   (   t   cufunc_handleR!  R"  R#  R  R   t   gxt   gyt   gzt   bxt   byt   bzt
   param_valst   argt   params(    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyR&  @  s    			t   oRz  t   at   cubint   fatbint   Linkerc           B` sY   e  Z d  d  Z e d    Z e d    Z d d  Z d   Z d   Z d   Z	 RS(	   i    c         C` sz  t  t d d   } t |   } t |   } i t |  t j 6t |  t j 6t |  t j 6t |  t j	 6t d  t j
 6} | r t |  | t j <n  t | j    t j g } t | j    } ~ t j t |  |   } t t |  |   }	 t j   |  _ }
 t j t |  | |	 t |  j   t j |  t j |
  | |  _ | |  _ | | | |	 g |  _ d  S(   NR  i   i   (   R*   R   R   R	   R   R  R
   R  R  R  R  t   CU_JIT_MAX_REGISTERSR   R  t   CU_JIT_TARGET_FROM_CUCONTEXTR   R   R  R   t   cu_link_stateR  R   t   cuLinkCreateR   R   R  t   cuLinkDestroyt   linker_info_buft   linker_errors_buft   _keep_alive(   R@   t   max_registersR  t
   linkerinfot   linkererrorsR  t   raw_keyst
   raw_valuesR  R  R  (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyR?   `  s.    		c         C` s   |  j  j j d  S(   NRw  (   RL  R   R  (   R@   (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyR    s    c         C` s   |  j  j j d  S(   NRw  (   RM  R   R  (   R@   (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt	   error_log  s    s   <cudapy-ptx>c      	   C` s   t  |  } t  | j d   } |  j | | g 7_ y2 t j |  j t j | t |  | d d  d   Wn, t
 k
 r } t d | |  j f   n Xd  S(   NRw  i    s   %s
%s(   R   Rx  RN  R   t   cuLinkAddDataR  R   t   CU_JIT_INPUT_PTXR   R(   R;   R:   RT  (   R@   Rz  Rk   t   ptxbuft   namebufRG   (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   add_ptx  s     c         C` s{   t  | j d   } |  j j |  y# t j |  j | | d d  d   Wn, t k
 rv } t	 d | |  j
 f   n Xd  S(   NRw  i    s   %s
%s(   R   Rx  RN  RT   R   t   cuLinkAddFileR  R(   R;   R:   RT  (   R@   RN   t   kindt   pathbufRG   (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   add_file  s    #c         C` s4   | j  d d  d } t | } |  j | |  d  S(   Nt   .i   (   t   rsplitt   FILE_EXTENSION_MAPR]  (   R@   RN   t   extR[  (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   add_file_guess_ext  s    
c         C` s   t  d  } t d  } y& t j |  j t |  t |   Wn, t k
 rl } t d | |  j f   n X| j	 } | d k s t
 d   |  j 2| | f S(   s   
        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.
        i    s   %s
%ss"   linker returned a zero sized cubin(   R
   R   R   t   cuLinkCompleteR  R   R;   R:   RT  R   R   RN  (   R@   RD  R  RG   (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   complete  s    &	(
   R"   R9   R?   R   R  RT  RY  R]  Rb  Rd  (    (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyRF  _  s   !
			c         C` s5   t  j t |  | t |    } t  j | d  d S(   s*   Query attribute on the device pointer
    s!   Failed to query pointer attributeN(   R   t   cuPointerGetAttributeR   R  t   check_error(   t   devmemR   t   odataR   (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   _device_pointer_attr  s    c         C` sX   t  d  } t |  t j |  i d t j 6d t j 6d t j 6d t j 6} | | j S(   sA   Query the device pointer type: host, device, array, unified?
    i    t   hostR  t   arrayt   unified(	   R   Ri  R   t    CU_POINTER_ATTRIBUTE_MEMORY_TYPEt   CU_MEMORYTYPE_HOSTt   CU_MEMORYTYPE_DEVICEt   CU_MEMORYTYPE_ARRAYt   CU_MEMORYTYPE_UNIFIEDR   (   Rg  t   ptrtypeRj   (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   device_pointer_type  s    


c         C` s2   t  d  } t j } t j t |  | |   | S(   sZ   Query the device pointer usable in the current context from an arbitrary
    pointer.
    i    (   R
   R   t#   CU_POINTER_ATTRIBUTE_DEVICE_POINTERR   Re  R   (   RH  R  R   (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   get_devptr_for_active_ctx  s    	c         C` sa   t  j   } t   } t |   } t j t |  t |  |  | j | j } } | | | f S(   s  Find 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.
    (   R   R   R   R  R   t   cuMemGetAddressRangeR   R   (   Rg  t   st   nR  (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   device_extents  s    	c         C` sk   t  |  d d  } | d k rF t |   \ } } | | } | |  _ n  | d k sg t d j |    | S(   s   Check 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.
    R  i    s   {} length arrayN(   R'   R(   Ry  R  R   R   (   Rg  t   szRw  RG   (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   device_memory_size  s    
!c         C` s+   t  |  d d  } | d k	 o* | j d k S(   s?   Returns True if the obj.dtype is datetime64 or timedelta64
    R  t   MmN(   R'   R(   t   char(   Rt   R  (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   _is_datetime_dtype  s    c         C` s%   t  |   r! |  j t j  }  n  |  S(   s^   Workaround for numpy#4983: buffer protocol doesn't support
    datetime64 or timedelta64.
    (   R~  R  t   npt   int64(   Rt   (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   _workaround_for_datetime  s    c         C` se   t  |  t t f  r |  St } | sF t  |  t j  p@ t |   } n  t |   }  t j	 |  | |  S(   s  Get 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.
    (
   R)   R*   R]  Rv   R  t   voidR~  R  R   t   memoryview_get_buffer(   Rt   t   readonlyt   forcewritable(    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyR    s    	!c         C` s   t  |   }  t j |   S(   sH   Returns (start, end) the start and end pointer of the array (half open).(   R  R   t   memoryview_get_extents(   Rt   (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   host_memory_extents  s    c         C` sV   t  |   t  |  k s$ t d   t  |   } t j |  | | |  \ } } | | S(   s_   Get the byte size of a contiguous memory buffer given the shape, strides
    and itemsize.
    s   # dim mismatch(   R   R   R   t   memoryview_get_extents_info(   R  R  R  t   ndimRw  RG   (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   memory_size_from_info  s    $c         C` s2   t  |   \ } } | | k s* t d   | | S(   s   Get the size of the memorys   memory extend of negative size(   R  R   (   Rt   Rw  RG   (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   host_memory_size)  s    c         C` s   t  |   j S(   s$   Get the device pointer as an integer(   R  R   (   Rt   (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyR  0  s    c         C` s'   |  d k r t d  St |   |  j S(   s,   Get the ctypes object for the device pointeri    N(   R(   R
   t   require_device_memoryR  (   Rt   (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyR  5  s    

c         C` s   t  |  d t  S(   sh  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(or long) object carrying the pointer
    value of the device memory address.  This is not tested in this method.
    R  (   R'   Rv   (   Rt   (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyR6  =  s    c         C` s   t  |   s t d   n  d S(   s9   A sentry for methods that accept CUDA memory object.
    s   Not a CUDA memory object.N(   R6  t	   Exception(   Rt   (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyR  H  s    c         G` s#   t  |  d g   } | j |  d S(   s   Add 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.
    t	   _depends_N(   R'   t   extend(   Rg  t   objst   depset(    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   device_memory_dependsO  s    c         C` so   g  } | r= t  | t  s! t  t j } | j | j  n	 t j } | t |   t	 | d t
 | |  d S(   s   
    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.
    R  N(   R)   R  R   R   t   cuMemcpyHtoDAsyncRT   R  t   cuMemcpyHtoDR  R  R|   (   t   dstt   srcR  R  t   varargsR  (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   host_to_deviceY  s    		c         C` si   g  } | r= t  | t  s! t  t j } | j | j  n	 t j } | t |   t	 |  | |  d S(   s   
    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.
    N(
   R)   R  R   R   t   cuMemcpyDtoHAsyncRT   R  t   cuMemcpyDtoHR  R  (   R  R  R  R  R  R  (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   device_to_hostk  s    		c         C` si   g  } | r= t  | t  s! t  t j } | j | j  n	 t j } | t |   t |  | |  d S(   s   
    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.
    N(	   R)   R  R   R   t   cuMemcpyDtoDAsyncRT   R  t   cuMemcpyDtoDR  (   R  R  R  R  R  R  (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyR  }  s    		c         C` sc   g  } | r= t  | t  s! t  t j } | j | j  n	 t j } | t |   | | |  d S(   s   Memset 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)   R  R   R   R  RT   R  R  R  (   R  t   valR  R  R  R  (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   device_memset  s    			c           C` s   t  j   d S(   s;   
    Enable profile collection in the current context.
    N(   R   t   cuProfilerStart(    (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   profile_start  s    c           C` s   t  j   d S(   s<   
    Disable profile collection in the current context.
    N(   R   t   cuProfilerStop(    (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   profile_stop  s    c           c` s   t    d Vt   d S(   s]   
    Context manager that enables profiling on entry and disables profiling on
    exit.
    N(   R  R  (    (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt	   profiling  s    (   R   t
   __future__R    R   R   R.   RM   RJ   R   R   R   R   R    R   t	   itertoolsR   R   R   R   R   R   R	   R
   R   R%  t   numpyR  t   collectionsR   R   R  R   R   R   R   R   R   R   R   R  R   R   R   R   R   t   numba.utilsR   R]  t   numba.cuda.envvarsR   R   R*   R  R  RI   Rh   Re  R7   R   R8   R:   R;   Rb   Rc   Rd   RH   RV   Rl   Rn   R   R   Rq   Ro   R   R   R   R   R   R  R  R  R'  R  R}  RI  RT  R`  R  R  R  R  R  Rg  Rm  RJ  RU  t   MemAllocRV  R  R  R  R  R  R  R  R  R&  t   CU_JIT_INPUT_OBJECTRV  t   CU_JIT_INPUT_LIBRARYt   CU_JIT_INPUT_CUBINt   CU_JIT_INPUT_FATBINARR`  RF  Ri  Rs  Ru  Ry  R{  R~  R  Rv   R  R  R  R  R  R  R6  R  R  R  R  R  R  R  R  R&  R  (    (    (    s8   lib/python2.7/site-packages/numba/cuda/cudadrv/driver.pyt   <module>   s   :	
	C			
		/				w	
	C 	"						$&kQ4		P	



Z			
						
						
		