
\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 m Z d d l m Z m Z d d l m Z d d l 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 l" m# Z# d d l$ Z% e j& e'  Z( d e f d     YZ) d   Z* d Z+ d   Z, d Z- d   Z. d Z/ d   Z0 d Z1 d   Z2 d Z3 d e4 f d     YZ5 d e4 f d     YZ6 e6   Z7 d e4 f d      YZ8 e d! e8 f d"     Y Z9 d# e# f d$     YZ: d% e8 f d&     YZ; d' e8 f d(     YZ< d) e4 f d*     YZ= d+ e4 f d,     YZ> d- e4 f d.     YZ? d/ e4 f d0     YZ@ d1 e4 f d2     YZA d3 e4 f d4     YZB d5 e4 f d6     YZC d7 e8 f d8     YZD d9 e4 f d:     YZE d; e jF f d<     YZG d= e4 f d>     YZH d? e4 f d@     YZI dA e4 f dB     YZJ dC   ZK dD   ZL dE   ZM dF   ZN dG   ZO dH   ZP dI   ZQ dJ   ZR dK   ZS dL   ZT dM   ZU dN   ZV dO   ZW dP   ZX eX   d k ZY d S(Q   s"   
HSA driver bridge implementation
i    (   t   absolute_importt   print_functiont   divisionN(   t   contextmanager(   t   defaultdictt   deque(   t   total_ordering(   t   mviewbuf(   t   utils(   t   configi   (   t   HsaSupportErrort   HsaDriverErrort   HsaApiError(   t   enumst	   enums_extt   drvapi(   t   longint(   t   Sequencet   HsaKernelTimedOutc           B` s   e  Z RS(    (   t   __name__t
   __module__(    (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyR       s   c         C` s.   y d d d g |  SWn t  k
 r) d SXd  S(   Nt   CPUt   GPUt   DSPt   Unknown(   t
   IndexError(   t   device(    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyt   _device_type_to_string$   s    s!   /opt/rocm/lib/libhsa-runtime64.soc          C` s  t  j j d t  }  |  d k r+ t   n  t j d  d k s^ t j d k s^ t j d k rh t	   n t
 j } d d g } d	 } |  d  k	 r y t  j j |   }  Wn! t k
 r t d
 |    n Xt  j j |   s t d |    n  |  g } n/ | g g  | D] } t  j j | |  ^ q} g  } g  } x_ | D]W } y | |  }	 Wn: t k
 r}
 | j t  j j |   | j |
  q<X|	 Sq<Wt |  rt   n# d j d   | D  } t |  d  S(   Nt   NUMBA_HSA_DRIVERt   0t   Pi   t   win32t   darwins   /usr/libs
   /usr/lib64s   libhsa-runtime64.sos'   NUMBA_HSA_DRIVER %s is not a valid pathsn   NUMBA_HSA_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(   t   str(   t   .0t   e(    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pys	   <genexpr>a   s    (   t   ost   environt   gett   DEFAULT_HSA_DRIVERt   _raise_driver_not_foundt   structt   calcsizet   syst   platformt   _raise_platform_not_supportedt   ctypest   CDLLt   Nonet   patht   abspatht
   ValueErrorR
   t   isfilet   joint   OSErrort   appendt   allt   _raise_driver_error(   t   envpatht   dlloadert   dldirt   dlnamet
   candidatest   xt   path_not_existt   driver_load_errorR1   t   dllR#   t   errmsg(    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyt   _find_driver.   sD    

	/
s8   
HSA is not currently supported on this platform ({0}).
c           C` s   t  t j t j    d  S(   N(   R
   t   PLATFORM_NOT_SUPPORTED_ERRORt   formatR+   R,   (    (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyR-   j   s    s   
The HSA runtime library cannot be found.

If you are sure that the HSA is installed, try setting environment
variable NUMBA_HSA_DRIVER with the file path of the HSA runtime shared
library.
c           C` s   t  t   d  S(   N(   R
   t   DRIVER_NOT_FOUND_MSG(    (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyR(   w   s    sD   
A HSA runtime library was found, but failed to load with error:
%s
c         C` s   t  t |    d  S(   N(   R
   t   DRIVER_LOAD_ERROR_MSG(   R#   (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyR9      s    s   driver missing function: %s.
t   Recyclerc           B` s5   e  Z d    Z d   Z d   Z d   Z d   Z RS(   c         C` s   g  |  _  t |  _ d  S(   N(   t   _garbaget   Truet   enabled(   t   self(    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyt   __init__   s    	c         C` s   |  j  j |  |  j   d  S(   N(   RJ   R7   t   service(   RM   t   obj(    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyt   free   s    c         C` s,   x |  j  D] } | j |  q
 W|  j  2d  S(   N(   RJ   t
   _finalizer(   RM   RP   (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyt   _cleanup   s    c         C` s2   |  j  r. t |  j  d k r. |  j   q. n  d  S(   Ni
   (   RL   t   lenRJ   RS   (   RM   (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyRO      s    	c         C` s   |  j    t |  _ d  S(   N(   RS   t   FalseRL   (   RM   (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyt   drain   s    
(   R   R   RN   RQ   RS   RO   RV   (    (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyRI      s
   				t   Driverc           B` s'  e  Z d  Z d Z d Z e j Z i e	 j
 e j f d 6e	 j e j f d 6e	 j e j f d 6e	 j e j f d 6e	 j e j f d 6Z d   Z d   Z d   Z d	   Z e d
    Z e d    Z e	 j e	 j e	 j d d  Z d d  Z d   Z d   Z  e d    Z! d   Z" d   Z# RS(   s0   
    Driver API functions are lazily bound.
    t   version_majort   version_minort	   timestampt   timestamp_frequencyt   signal_max_waitc         C` s5   |  j  } | d  k	 r | St j |   } | |  _  | S(   N(   t
   _singletonR0   t   objectt   __new__(   t   clsRP   (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyR_      s    		c         C` s   y: t  j r t d   n  t   |  _ t |  _ d  |  _ Wn% t k
 ra } t	 |  _ | |  _ n Xd  |  _
 i  |  _ t   |  _ t j   |  _ d  S(   Ns   HSA disabled by user(   R	   t   DISABLE_HSAR
   RD   t   libRU   t   is_initializedR0   t   initialization_errorRK   t
   _agent_mapt	   _programsRI   t	   _recyclert   weakreft   WeakSett   _active_streams(   RM   R#   (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyRN      s    					c         ` so     j  r d  St   _  y   j   Wn, t k
 rR } |   _ t d |   n Xt j   f d    } d  S(   Ns   Error at driver init: 
%s:c          ` sG   y" x   j  D] }  |  j   q WWn t k
 r5 n X  j j   d  S(   N(   t   agentst   releaset   AttributeErrorRg   RV   (   t   agent(   RM   (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyt   shutdown   s    (   Rc   RK   t   hsa_initR   Rd   R   t   atexitt   register(   RM   R#   Ro   (    (   RM   s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyt   _initialize_api   s    			c         ` st   |  j  d  k	 r d  S|  j   g      f d   } t j |  } |  j | d   t d     D  } | |  _  d  S(   Nc         ` s     j  |   t j S(   N(   R7   R   t   HSA_STATUS_SUCCESS(   t   agent_idt   ctxt(   t	   agent_ids(    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyt   on_agent   s    c         s` s!   |  ] } | t  |  f Vq d  S(   N(   t   Agent(   R"   Ru   (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pys	   <genexpr>   s    (   Re   R0   Rs   R   t   HSA_ITER_AGENT_CALLBACK_FUNCt   hsa_iterate_agentst   dict(   RM   Rx   t   callbackt	   agent_map(    (   Rw   s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyt   _initialize_agents   s    
c         C` s   |  j    |  j d  k S(   N(   Rs   Rd   R0   (   RM   (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyt   is_available   s    
c         C` s   |  j    |  j j   S(   N(   R   Re   t   values(   RM   (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyRk      s    
c         C` sJ   t  j   } | d  k s t  |  j | | | | t j |   t |  S(   N(   R   t   hsa_ext_program_tR0   t   AssertionErrort   hsa_ext_program_createR.   t   byreft   Program(   RM   t   modelt   profilet   rounding_modet   optionst   program(    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyt   create_program   s
    c         C` s   | d  k r t |  j  } n  t |  } t j | } | g  | D] } | j ^ qA   } t j   } |  j | | | t	 j
 |   t | j  S(   N(   R0   t   tupleRk   RT   R   t   hsa_agent_tt   _idt   hsa_signal_tt   hsa_signal_createR.   R   t   Signalt   value(   RM   t   initial_valuet	   consumerst   consumers_lent   consumers_typet   ct   result(    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyt   create_signal
  s    "c         ` s  |  j    y@ |  j   \ } } |   } |  j | t j |   | j SWn t k
 r] n Xy |  j   } Wn t k
 r t     n X|  j	 d  k	 r t d |  j	   n  |  j    } x* | j   D] \ } } t | | |  q W  f d   }	 |	 |  }
 t |    |
  |
 S(   Ns   Error at driver init: 
%s:c         ` s      f d   } | S(   Nc          ` s   t  j d     |  |   S(   Ns   call driver api: %s(   t   _loggert   debug(   t   argst   kwargs(   t   fnt   fname(    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyt   wrapped5  s    (    (   R   R   (   R   (   R   s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyt   driver_wrapper4  s    (   Rs   t   _hsa_propertiest   hsa_system_get_infoR.   R   R   t   KeyErrort   _api_prototypesRm   Rd   R0   R
   t	   _find_apit   itemst   setattr(   RM   R   t   enumt   typR   t   protot   libfnt   keyt   valR   t   retval(    (   R   s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyt   __getattr__  s,    
	c         ` sK   y t  |  j    SWn t k
 r' n X  f d   } t |    |  | S(   Nc          ` s   t  t     d  S(   N(   R   t   MISSING_FUNCTION_ERRMSG(   R   t   kws(   R   (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyt   absent_functionG  s    (   t   getattrRb   Rm   R   (   RM   R   R   (    (   R   s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyR   >  s    c         C` s%   t  t d   t t |  j     S(   s^   Returns a ordered list of components

        The first device should be picked first
        c         S` s   |  j  S(   N(   t   is_component(   t   a(    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyt   <lambda>S  t    (   t   listt   filtert   reversedt   sortedRk   (   RM   (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyt
   componentsM  s    c         C` s   t    } |  j j |  | S(   N(   t   StreamRj   t   add(   RM   t   st(    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyt   create_streamV  s    	c         C` s/   t  j d  x |  j D] } | j   q Wd S(   sc   
        Implicit synchronization for all asynchronous streams
        across all devices.
        s   implicit syncN(   R   t   infoRj   t   synchronize(   RM   R   (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyt   implicit_sync[  s    N($   R   R   t   __doc__R0   R]   Re   R   t   API_PROTOTYPESR   R   t   HSA_SYSTEM_INFO_VERSION_MAJORR.   t   c_uint16t   HSA_SYSTEM_INFO_VERSION_MINORt   HSA_SYSTEM_INFO_TIMESTAMPt   c_uint64t#   HSA_SYSTEM_INFO_TIMESTAMP_FREQUENCYt   HSA_SYSTEM_INFO_SIGNAL_MAX_WAITR   R_   RN   Rs   R   t   propertyR   Rk   t   HSA_MACHINE_MODEL_LARGEt   HSA_PROFILE_FULLt'   HSA_DEFAULT_FLOAT_ROUNDING_MODE_DEFAULTR   R   R   R   R   R   R   (    (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyRW      s4   							'			t
   HsaWrapperc           B` s   e  Z d    Z d   Z RS(   c         C` s   y |  j  | \ } } Wn* t k
 rC t d |  j | f   n Xt t |  j  } |   } t | d  } | rz | n t j	 |  } | |  j
 | |  | s | j t j k r | j St |  Sd  S(   Ns   %r object has no attribute %rt   _length_(   R   R   Rm   t	   __class__R   t   hsat   _hsa_info_functiont   hasattrR.   R   R   t   _type_t   c_charR   R   (   RM   R   R   R   t   funcR   t   is_array_typet   result_buff(    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyR   h  s    	c         C` s6   t  t t t |    |  j j   |  j j     S(   N(   R   t   sett   dirt   typet   __dict__t   keysR   (   RM   (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyt   __dir__{  s    (   R   R   R   R   (    (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyR   g  s   	Ry   c           B` s  e  Z d  Z d Z i e j e j d f d 6e j e j d f d 6e j	 e
 j f d 6e j e j f d 6e j e j d f d 6e j e
 j f d	 6e j e j f d
 6e j e j f d 6e j e j f d 6e j e j f d 6e j e
 j f d 6e j e j f d 6e j e
 j f d 6e j e j d f d 6e j e
 j f d 6Z d   Z e  d    Z! e  d    Z" e  d    Z# e  d    Z$ e  d    Z% d   Z& d   Z' d' d' d' d' d' d  Z) d   Z* d   Z+ d   Z, d    Z- d!   Z. d"   Z/ d#   Z0 d$   Z1 d%   Z2 d&   Z3 RS((   sr   Abstracts a HSA compute agent.

    This will wrap and provide an OO interface for hsa_agent_t C-API elements
    t   hsa_agent_get_infoi@   t   namet   vendor_namet   featuret   wavefront_sizei   t   workgroup_max_dimt   grid_max_dimt   grid_max_sizet   fbarrier_max_sizet
   queues_maxt   queue_max_sizet
   queue_typet   nodet   _devicei   t
   cache_sizet   isac         C` s9   | |  _  t j |  _ t   |  _ |  j   |  j   d  S(   N(   R   R   Rg   R   t   _queuest   _initialize_regionst   _initialize_mempools(   RM   Ru   (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyRN     s
    	
c         C` s   t  |  j  S(   N(   R   R   (   RM   (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyR     s    c         C` s   |  j  t j @d k S(   Ni    (   R   R   t!   HSA_AGENT_FEATURE_KERNEL_DISPATCH(   RM   (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyR     s    c         C` s   |  j  S(   N(   t   _regions(   RM   (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyt   regions  s    c         C` s   |  j  S(   N(   t	   _mempools(   RM   (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyt   mempools  s    c         C` s#   t  |  j  d d d  j d  S(   s&   
        log2(wavefront_size)
        Nit   1(   t   binR   t   index(   RM   (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyt   wavebits  s    c         ` sl   g      f d   } t  j |  } t j |  j | d   t g    D] } t j |  |  ^ qD  |  _	 d  S(   Nc         ` s     j  |   t j S(   N(   R7   R   Rt   (   t	   region_idRv   (   t
   region_ids(    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyt	   on_region  s    (
   R   t'   HSA_AGENT_ITERATE_REGIONS_CALLBACK_FUNCR   t   hsa_agent_iterate_regionsR   R0   t   _RegionListt	   MemRegiont   instance_forR   (   RM   R  R}   R   (    (   R  s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyR     s    c         ` so   g    d    f d  } t j |  } t j |  j | d   t g    D] } t j |  |  ^ qG  |  _	 d  S(   Nc         ` s     j  |   t j S(   N(   R7   R   Rt   (   R   Rv   (   t   mempool_ids(    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyR    s    (
   R0   R   t+   HSA_AMD_AGENT_ITERATE_MEMORY_POOLS_CALLBACKR   t"   hsa_amd_agent_iterate_memory_poolsR   R  t   MemPoolR  R   (   RM   R  R}   t
   mempool_id(    (   R  s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyR     s    c      
   C` s  | d  k	 s t  | |  j k s' t  t j } | d  k rN t j d  |  n	 | |  } t j t j    }	 | d  k r t j	 d  n | } | d  k r t j	 d  n | } t
 j |  j | | | | | | t j |	   t |  |	  }
 |  j j |
  t j |
  S(   Ni(   R0   R   R   R   t   HSA_QUEUE_CALLBACK_FUNCR.   t   castt   POINTERt   hsa_queue_tt   c_uint32R   t   hsa_queue_createR   R   t   QueueR   R   Rh   t   proxy(   RM   t   sizeR}   t   datat   private_segment_sizet   group_segment_sizeR   t   cb_typt   cbR   t   q(    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyt   _create_queue  s    	*c         O` s   t  j | d <|  j | |   S(   NR   (   R   t   HSA_QUEUE_TYPE_SINGLER  (   RM   R   R   (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyt   create_queue_single  s    c         O` s   t  j | d <|  j | |   S(   NR   (   R   t   HSA_QUEUE_TYPE_MULTIR  (   RM   R   R   (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyt   create_queue_multi  s    c         C` s(   x! t  |  j  D] } | j   q Wd S(   sJ   
        Release all resources

        Called at system teardown
        N(   R   R   Rl   (   RM   R  (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyRl     s    c         C` s$   |  j  j |  |  j j |  d  S(   N(   R   t   removeRg   RQ   (   RM   t   queue(    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyt   release_queue
  s    c         C` s4   d j  |  j |  j |  j |  j |  j r- d n d  S(   Ns#   <HSA agent ({0}): {1} {2} '{3}'{4}>s    (component)R   (   RF   R   R   R   R   R   (   RM   (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyt   __repr__  s
    c         C` s   |  j  |  j |  j f S(   N(   R   R   R   (   RM   (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyt   _rank  s    c         C` s-   t  |  t  r% |  j   | j   k  St Sd  S(   N(   t
   isinstanceRy   R%  t   NotImplemented(   RM   t   other(    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyt   __lt__  s    c         C` s-   t  |  t  r% |  j   | j   k St Sd  S(   N(   R&  Ry   R%  R'  (   RM   R(  (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyt   __eq__  s    c         C` s   t  |  j    S(   N(   t   hashR%  (   RM   (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyt   __hash__$  s    c         C` s
   t  |   S(   N(   t   Context(   RM   (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyt   create_context'  s    N(4   R   R   R   R   R   t   HSA_AGENT_INFO_NAMER.   R   t   HSA_AGENT_INFO_VENDOR_NAMEt   HSA_AGENT_INFO_FEATURER   t   hsa_agent_feature_tt   HSA_AGENT_INFO_WAVEFRONT_SIZER  t    HSA_AGENT_INFO_WORKGROUP_MAX_DIMR   t   HSA_AGENT_INFO_GRID_MAX_DIMt
   hsa_dim3_tt   HSA_AGENT_INFO_GRID_MAX_SIZEt    HSA_AGENT_INFO_FBARRIER_MAX_SIZEt   HSA_AGENT_INFO_QUEUES_MAXt   HSA_AGENT_INFO_QUEUE_MAX_SIZEt   HSA_AGENT_INFO_QUEUE_TYPEt   hsa_queue_type_tt   HSA_AGENT_INFO_NODEt   HSA_AGENT_INFO_DEVICEt   hsa_device_type_tt   HSA_AGENT_INFO_CACHE_SIZEt   HSA_AGENT_INFO_ISAt	   hsa_isa_tR   RN   R   R   R   R   R   R   R   R   R0   R  R  R   Rl   R#  R$  R%  R)  R*  R,  R.  (    (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyRy     sN   	
												R  c           B` s;   e  Z d
 Z d   Z d   Z d   Z d   Z d	   Z RS(   t   _allt   globalst	   readonlyst   privatest   groupsc         C` sw   t  |  |  _ t  d   | D  |  _ t  d   | D  |  _ t  d   | D  |  _ t  d   | D  |  _ d  S(   Nc         s` s$   |  ] } | j  d  k r | Vq d S(   t   globalN(   t   kind(   R"   R?   (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pys	   <genexpr>0  s    c         s` s$   |  ] } | j  d  k r | Vq d S(   t   readonlyN(   RI  (   R"   R?   (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pys	   <genexpr>1  s    c         s` s$   |  ] } | j  d  k r | Vq d S(   t   privateN(   RI  (   R"   R?   (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pys	   <genexpr>2  s    c         s` s$   |  ] } | j  d  k r | Vq d S(   t   groupN(   RI  (   R"   R?   (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pys	   <genexpr>3  s    (   R   RC  RD  RE  RF  RG  (   RM   t   lst(    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyRN   .  s
    c         C` s   t  |  j  S(   N(   RT   RC  (   RM   (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyt   __len__5  s    c         C` s   | |  j  k S(   N(   RC  (   RM   t   item(    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyt   __contains__8  s    c         C` s   t  |  j  S(   N(   R   RC  (   RM   (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyt   __reversed__;  s    c         C` s   |  j  | S(   N(   RC  (   RM   t   idx(    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyt   __getitem__>  s    (   RC  RD  RE  RF  RG  (   R   R   t	   __slots__RN   RN  RP  RQ  RS  (    (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyR  +  s   				R  c           B` s  e  Z d  Z d Z i e j e j f d 6e j e	 j
 f d 6e j e	 j f d 6e j e	 j f d 6e j e	 j f d 6e j e	 j f d 6e j e	 j f d 6Z i d	 e j 6d
 e j 6d e j 6d e j 6Z d   Z e d    Z e d    Z d   Z d   Z i  Z e d    Z  RS(   s{   Abstracts a HSA mem pool.

    This will wrap and provide an OO interface for hsa_amd_memory_pool_t
    C-API elements
    t   hsa_amd_memory_pool_get_infot   segmentt   _flagsR  t   alloc_allowedt   alloc_granulet   alloc_alignmentt   accessible_by_allRH  RJ  RK  RL  c         C` s"   | |  _  | |  _ |  j  |  _ d S(   s{   Do not instantiate MemPool objects directly, use the factory class
        method 'instance_for' to ensure MemPool identityN(   R   t   _owner_agentt   _as_parameter_(   RM   Rn   t   pool(    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyRN   f  s    		c         C` s   |  j  |  j S(   N(   t   _segment_name_mapRV  (   RM   (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyRI  m  s    c         C` s   |  j  S(   N(   R\  (   RM   (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyRn   q  s    c         C` s"   |  j  d k r |  j | @St Sd S(   s  
            Determines if a given feature is supported by this MemRegion.
            Feature flags are found in "./enums_exp.py" under:
                * hsa_amd_memory_pool_global_flag_t
                Params:
                check_flag: Feature flag to test
        RH  N(   RI  RW  RU   (   RM   t
   check_flag(    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyt   supportsu  s    c         C` s   |  j  s t  | d k s! t  t j   } t j d  } t j |  j | | t j |   | j	 d  k r t d j |     n  | S(   Ni    s   Failed to allocate from {}(   RX  R   R.   t   c_void_pR  R   t   hsa_amd_memory_pool_allocateR   R   R   R0   R   RF   (   RM   t   nbytest   bufft   flags(    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyt   allocate  s    "c         C` sD   y |  j  | SWn. t k
 r? |  | |  } | |  j  | <| SXd  S(   N(   t   _instance_dictR   (   R`   t   ownerR   t   new_instance(    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyR    s    (!   R   R   R   R   R   t    HSA_AMD_MEMORY_POOL_INFO_SEGMENTR   t   hsa_amd_segment_tt%   HSA_AMD_MEMORY_POOL_INFO_GLOBAL_FLAGSR.   R  t   HSA_AMD_MEMORY_POOL_INFO_SIZEt   c_size_tt.   HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_ALLOWEDt   c_boolt.   HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_GRANULEt0   HSA_AMD_MEMORY_POOL_INFO_RUNTIME_ALLOC_ALIGNMENTt*   HSA_AMD_MEMORY_POOL_INFO_ACCESSIBLE_BY_ALLR   t   HSA_AMD_SEGMENT_GLOBALt   HSA_AMD_SEGMENT_READONLYt   HSA_AMD_SEGMENT_PRIVATEt   HSA_AMD_SEGMENT_GROUPR_  RN   R   RI  Rn   Ra  Rg  Rh  t   classmethodR  (    (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyR  B  s:   


			
R  c           B` s7  e  Z d  Z d Z i e j e j f d 6e j e j	 f d 6e
 j e j f d 6e j e j f d 6e j e j f d 6e j e j f d 6e j e j f d 6e j e j f d	 6Z i d
 e j 6d e j 6d e j 6d e j 6Z d   Z e d    Z e d    Z d   Z d   Z d   Z  i  Z! e" d    Z# RS(   ss   Abstracts a HSA memory region.

    This will wrap and provide an OO interface for hsa_region_t C-API elements
    t   hsa_region_get_infoRV  RW  t   host_accessibleR  t   alloc_max_sizeRZ  RY  RX  RH  RJ  RK  RL  c         C` s"   | |  _  | |  _ |  j  |  _ d S(   s   Do not instantiate MemRegion objects directly, use the factory class
        method 'instance_for' to ensure MemRegion identityN(   R   R\  R]  (   RM   Rn   R   (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyRN     s    		c         C` s   |  j  |  j S(   N(   R_  RV  (   RM   (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyRI    s    c         C` s   |  j  S(   N(   R\  (   RM   (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyRn     s    c         C` s"   |  j  d k r |  j | @St Sd S(   s  
            Determines if a given feature is supported by this MemRegion.
            Feature flags are found in "./enums.py" under:
                * hsa_region_global_flag_t
                Params:
                check_flag: Feature flag to test
        RH  N(   RI  RW  RU   (   RM   R`  (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyRa    s    c         C` se   |  j  s t  | |  j k s$ t  | d k s6 t  t j   } t j |  j | t j |   | S(   Ni    (	   RX  R   R|  R.   Rb  R   t   hsa_memory_allocateR   R   (   RM   Rd  Re  (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyRg    s    c         C` s   t  j |  d  S(   N(   R   t   hsa_memory_free(   RM   t   ptr(    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyRQ     s    c         C` sD   y |  j  | SWn. t k
 r? |  | |  } | |  j  | <| SXd  S(   N(   Rh  R   (   R`   Ri  R   Rj  (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyR    s    ($   R   R   R   R   R   t   HSA_REGION_INFO_SEGMENTR   t   hsa_region_segment_tt   HSA_REGION_INFO_GLOBAL_FLAGSt   hsa_region_global_flag_tR   t#   HSA_AMD_REGION_INFO_HOST_ACCESSIBLER.   Rq  t   HSA_REGION_INFO_SIZERo  t   HSA_REGION_INFO_ALLOC_MAX_SIZEt'   HSA_REGION_INFO_RUNTIME_ALLOC_ALIGNMENTt%   HSA_REGION_INFO_RUNTIME_ALLOC_GRANULEt%   HSA_REGION_INFO_RUNTIME_ALLOC_ALLOWEDR   t   HSA_REGION_SEGMENT_GLOBALt   HSA_REGION_SEGMENT_READONLYt   HSA_REGION_SEGMENT_PRIVATEt   HSA_REGION_SEGMENT_GROUPR_  RN   R   RI  Rn   Ra  Rg  RQ   Rh  Ry  R  (    (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyR    s@   


				R  c           B` s_   e  Z d    Z d   Z d   Z e d    Z d   Z d d d d  Z	 d   Z
 d   Z RS(	   c         C` s7   t  j |  |  _ | |  _ |  j |  _ t j |  _ d S(   s   The id in a queue is a pointer to the queue object returned by hsa_queue_create.
        The Queue object has ownership on that queue objectN(   Rh   R  t   _agentR   R]  R   t   hsa_queue_destroyRR   (   RM   Rn   t	   queue_ptr(    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyRN     s    	c         C` s   |  j  j |   d  S(   N(   R  R#  (   RM   (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyRl     s    c         C` s   t  |  j j |  S(   N(   R   R   t   contents(   RM   R   (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyR     s    c   
      c` s  |  j  j } | j d } t j |  t j t j  k s@ t  | | j } t j	 |  j  d  } xC t
 r t j |  j   } | | k o | | j k  n re Pqe qe W| | @} | j | j  } | | }	 t j t j |	  d t j |   |	 Vt j |  j  j j |  d  S(   Ni   i    (   R   R  R  R.   t   sizeofR   t   hsa_kernel_dispatch_packet_tR   R   t!   hsa_queue_add_write_index_acq_relRK   t!   hsa_queue_load_read_index_acquiret   from_addresst   base_addresst   memsett	   addressoft   hsa_signal_store_releaset   doorbell_signal(
   RM   t   packet_typet   queue_structt
   queue_maskt   packet_array_tR   t   read_offsett   queue_offsetR"  t   packet(    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyt   _get_packet  s     	#

%c         C` s   |  j  t j  o } | j | _ d } | t j t j >O} | t j t j >O} | t j	 t j
 >O} | d t j >O} | | _ Wd  QXd  S(   Ni    i   (   R  R   t   hsa_barrier_and_packet_tR   t   dep_signal0R   t   HSA_FENCE_SCOPE_SYSTEMt%   HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPEt%   HSA_PACKET_HEADER_RELEASE_FENCE_SCOPEt   HSA_PACKET_TYPE_BARRIER_ANDt   HSA_PACKET_HEADER_TYPEt   HSA_PACKET_HEADER_BARRIERt   header(   RM   t
   dep_signalR  R  (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyt   insert_barrier  s    c         C` s  t  j d | j  t |  } | t |  k s7 t  d | k  oN d k n sY t  | | k sk t  | t |  j j  |  k r d } t | j	 | t |  j j  |     n  | d  k	 r | n t j d  } |  j t j  L}	 |	 j | t j >O_ | d |	 _ | d k r'| d n d |	 _ | d k rF| d n d |	 _ | d |	 _ | d k rr| d n d |	 _ | d k r| d n d |	 _ | j |	 _ | j |	 _ | d  k rd n | j |	 _ | j |	 _ | j |	 _ d }
 |
 t j  t j! >O}
 |
 t j  t j" >O}
 |
 t j# t j$ >O}
 |
 |	 _% Wd  QX| d  k rt  j d  d } | j& d	 |  sd
 } t' | j	 d	 |    qn  d  S(   Ns   dispatch %si    i   s"   workgroupsize is too big {0} > {1}i   i   s&   wait for sychronous kernel to completei
   t   timeouts'   Kernel timed out after {timeout} second((   R   R   R   RT   R   R   R  R   R   RF   R0   R   R   R  R   R  t   setupR   t+   HSA_KERNEL_DISPATCH_PACKET_SETUP_DIMENSIONSt   workgroup_size_xt   workgroup_size_yt   workgroup_size_zt   grid_size_xt   grid_size_yt   grid_size_zR   t   completion_signalt   kernel_objectR   t   kernarg_addressR  R  R  R  R  t   HSA_PACKET_TYPE_KERNEL_DISPATCHR  R  t   wait_until_ne_oneR   (   RM   t   symbolt   kernargst   workgroup_sizet	   grid_sizet   signalt   dimst   msgt   sR  R  R  (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyt   dispatch+  sF    "!c         C` s)   t  t t |  j j  |  j j     S(   N(   R   R   R   R   R  R   R   (   RM   (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyR   g  s    c         C` s
   t  |   S(   N(   t   ManagedQueueProxy(   RM   (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyt   ownedk  s    N(   R   R   RN   Rl   R   R   R  R  R0   R  R   R  (    (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyR    s   				9	R  c           B` s   e  Z d    Z d   Z RS(   c         C` s   t  j |  |  _ d  S(   N(   Rh   t   reft   _queue(   RM   R"  (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyRN   p  s    c         C` s   t  |  j   |  S(   N(   R   R  (   RM   RO  (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyR   s  s    (   R   R   RN   R   (    (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyR  o  s   	R   c           B` s5   e  Z d  Z d   Z d   Z d   Z d d  Z RS(   s   The id for the signal is going to be the hsa_signal_t returned by create_signal.
    Lifetime of the underlying signal will be tied with this object".
    Note that it is likely signals will have lifetime issues.c         C` s2   | |  _  |  j  |  _ t j |  t j |  j   d  S(   N(   R   R]  R   t   finalizeR   t   hsa_signal_destroy(   RM   t	   signal_id(    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyRN   |  s    	c         C` s   t  j |  j  S(   N(   R   t   hsa_signal_load_relaxedR   (   RM   (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyt   load_relaxed  s    c         C` s   t  j |  j  S(   N(   R   t   hsa_signal_load_acquireR   (   RM   (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyt   load_acquire  s    c         C` sd   d } d } | d k r! d } n | t j | } t j |  j t j | | t j  |  j   | k S(   sL   
        Returns a boolean to indicate whether the wait has timeout
        i   i
   i   ii@B N(	   R0   R   R[   t   hsa_signal_wait_acquireR   R   t   HSA_SIGNAL_CONDITION_NEt   HSA_WAIT_STATE_ACTIVER  (   RM   R  t   onet   mhzt   expire(    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyR    s    	
N(   R   R   R   RN   R  R  R0   R  (    (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyR   w  s
   			t
   BrigModulec           B` s2   e  Z d    Z e d    Z d   Z d   Z RS(   c         C` s=   t  j |  } | |  _ t  j t  j |  t j  |  _ d S(   s5   
        Take a byte buffer of a Brig module
        N(   R.   t   create_string_buffert   _bufferR  R  R   t   hsa_ext_module_tR   (   RM   t   brig_buffert   buf(    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyRN     s    	c         C` s.   t  | d   } | j   } Wd  QXt |  S(   Nt   rb(   t   opent   readR  (   R`   t	   file_namet   finR  (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyt	   from_file  s    c         C` s   t  |  j  S(   N(   RT   R  (   RM   (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyRN    s    c         C` s"   d j  t t |    t |    S(   Ns!   <BrigModule id={0} size={1}bytes>(   RF   t   hext   idRT   (   RM   (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyR$    s    (   R   R   RN   Ry  R  RN  R$  (    (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyR    s   			R   c           B` sD   e  Z e j e j e j d d  d d  Z d   Z d d d  Z	 RS(   i   i    c   
      C` s  t  j   |  _ | d  k s! t  d   } t j d  } t j t	 j
 | | t j |   | j sz t d | | f   t  j   |  _ t j t	 j
 | | t j |  j   |  j j | | | | t j |  j   }	 | |	  |  j |  _ t j |  |  j j |  j  d  S(   Nc         S` s_   |  t  j k	 r[ t j   } t j |  t j |   t j | j	 j
 d   t |   n  d  S(   Ns   utf-8(   R   Rt   R.   t   c_char_pR   t   hsa_status_stringR   R   R   R   t   decodet   exit(   t
   hsa_statusR  (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyt   check_fptr_return  s
    i    s(   HSA system extension %s.%s not supported(   R   R   R   R0   R   R.   Rq  R   t   hsa_system_extension_supportedR   t   HSA_EXTENSION_FINALIZERR   R   t   hsa_ext_finalizer_1_00_pfn_tt   _ftablt   hsa_system_get_extension_tableR   R]  R   R  t   hsa_ext_program_destroy(
   RM   R   R   R   R   RX   RY   R  t   supportt   ret(    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyRN     s,    	
c         C` s   |  j  j |  j | j  d  S(   N(   R  t   hsa_ext_program_add_moduleR   (   RM   t   module(    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyt
   add_module  s    c      	   C` sx   t  j   } t  j   } t j t j |  d t j |   |  j j |  j	 | | | | t
 j t j |   t |  S(   sN   
        The program object is safe to be deleted after ``finalize``.
        i    (   R   t   hsa_code_object_tt   hsa_ext_control_directives_tR.   R  R   R  R  t   hsa_ext_program_finalizeR   R   t   HSA_CODE_OBJECT_TYPE_PROGRAMt
   CodeObject(   RM   R   t   callconvR   t   code_objectt   control_directives(    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyR    s    N(
   R   R   R   R   R   R   R0   RN   R  R  (    (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyR     s   '	R  c           B` s   e  Z d    Z RS(   c         C` s2   | |  _  |  j  |  _ t j |  t j |  j   d  S(   N(   R   R]  R   R  R   t   hsa_code_object_destroy(   RM   R  (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyRN     s    	(   R   R   RN   (    (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyR    s   t
   Executablec           B` s,   e  Z d    Z d   Z d   Z d   Z RS(   c         C` sc   t  j   } t j t j t j d  t j	 |   | |  _
 |  j
 |  _ t j |  t j |  j
  d  S(   N(   R   t   hsa_executable_tR   t   hsa_executable_createR   R   t   HSA_EXECUTABLE_STATE_UNFROZENR0   R.   R   R   R]  R   R  t   hsa_executable_destroy(   RM   t   ex(    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyRN     s    	c         C` s#   t  j |  j | j | j d   d  S(   N(   R   t   hsa_executable_load_code_objectR   R0   (   RM   Rn   R  (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyt   load  s    c         C` s   t  j |  j d  d S(   s0   Freeze executable before we can query for symbolN(   R   t   hsa_executable_freezeR   R0   (   RM   (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyt   freeze
  s    c         C` sV   t  j   } t j |  j d  t j | j d   | j d t j	 |   t
 | |  S(   Nt   asciii    (   R   t   hsa_executable_symbol_tR   t   hsa_executable_get_symbolR   R0   R.   R  t   encodeR   t   Symbol(   RM   Rn   R   R  (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyt
   get_symbol  s    	(   R   R   RN   R  R  R  (    (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyR    s   	
		R  c           B` si   e  Z d  Z i e j e j f d 6e j e j f d 6e j	 e j f d 6e j
 e j f d 6Z d   Z RS(   t   hsa_executable_symbol_get_infoR  t   kernarg_segment_sizeR  R  c         C` s   | |  _  | |  _ d  S(   N(   R   R   (   RM   R   t	   symbol_id(    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyRN   -  s    	(   R   R   R   R   t(   HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_OBJECTR.   R   t6   HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_KERNARG_SEGMENT_SIZER  t4   HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_GROUP_SEGMENT_SIZEt6   HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_PRIVATE_SEGMENT_SIZER   RN   (    (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyR    s   t   MemoryPointerc           B` sV   e  Z e Z d d   Z d   Z d   Z d   Z d   Z	 e
 d    Z d   Z RS(   c         C` sg   t  | t  s t  | |  _ | |  _ | |  _ | |  _ | |  _ | d  k	 |  _	 t
 |  _ d |  _ d  S(   Ni    (   R&  R-  R   t   contextt   device_pointerR  t   _hsa_memsize_t	   finalizerR0   t
   is_managedRK   t   is_alivet   refct(   RM   R  t   pointerR  R  (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyRN   5  s    						c         C` s;   y# |  j  r" |  j r" |  j   n  Wn t j   n Xd  S(   N(   R  R  R  t	   tracebackt	   print_exc(   RM   (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyt   __del__@  s
    c         C` s   t  t j |    S(   N(   t   OwnedPointerRh   R  (   RM   (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyt   ownG  s    c         C` s;   |  j  r7 |  j s! t d   n  |  j   t |  _ n  d S(   s8   
        Forces the device memory to the trash.
        s   Freeing dead memoryN(   R  R  t   RuntimeErrorR  RU   (   RM   (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyRQ   J  s
    		
c         C` s:   |  j  j } t |  j | |  j  } t t j |   |  S(   N(   R  R   R  R  R  R"  Rh   R  (   RM   R  t   view(    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyR%  T  s    c         C` s   |  j  S(   N(   R  (   RM   (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyt   device_ctypes_pointerY  s    c         G` sb   t  |  } | d k r d S| t j g  | D] } | j ^ q-   } t j | | d |  j  d S(   s   
        Grant access to given *agents*.
        Upon return, only the listed-agents and the owner agent have direct
        access to this pointer.
        i    N(   RT   R   R   R   R   t   hsa_amd_agents_allow_accessR0   R  (   RM   Rk   t   ctR   t   agent_array(    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyt   allow_access_to]  s    )N(   R   R   RK   t   __hsa_memory__R0   RN   R!  R#  RQ   R%  R   R&  R*  (    (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyR  2  s   			
	t
   HostMemoryc           B` s   e  Z d    Z d   Z RS(   c         C` sO   | |  _  | |  _ | |  _ | |  _ |  j |  _ |  j |  _ |  j j |  _ d  S(   N(   R  R  R  t   host_pointert   handlet   _buflen_R   t   _bufptr_(   RM   R  Ri  R  R  (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyRN   l  s    				c         C` s   |  S(   N(    (   RM   (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyR#  w  s    (   R   R   RN   R#  (    (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyR,  k  s   	R"  c           B` s&   e  Z d d   Z d   Z d   Z RS(   c         C` sS   | |  _  |  j  j d 7_ | d  k r6 |  j  |  _ n | j sF t  | |  _ d  S(   Ni   (   t   _memR  R0   t   _viewR  R   (   RM   t   memptrR%  (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyRN   |  s    	c         C` sx   yP |  j  j d 8_ |  j  j d k s- t  |  j  j d k rO |  j  j   n  Wn! t k
 rc n t j   n Xd  S(   Ni   i    (   R1  R  R   RQ   t   ReferenceErrorR  R   (   RM   (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyR!    s    c         C` s   t  |  j |  S(   s$   Proxy MemoryPointer methods
        (   R   R2  (   RM   R   (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyR     s    N(   R   R   R0   RN   R!  R   (    (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyR"  {  s   		R-  c           B` s   e  Z d  Z e j   Z d   Z d   Z e d    Z	 e d    Z
 e d    Z e d    Z e d    Z d e d  Z d e d	  Z d
   Z RS(   s2   
    A context is associated with a component
    c         C` s   t  j |  |  _ |  j j rT | j } |  j j | d |  j } | j   |  _ n  t	 j
   |  _ t j } t j } g  | j j D] } | j r | ^ q } d  |  _ d  |  _ xD | D]< } | j |  r | |  _ n  | j |  r | |  _ q q Wd  S(   NR}   (   Rh   R  R  R   R   R   t	   _callbackR  t   _defaultqueueR   t
   UniqueDictt   allocationsR   t.   HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_COARSE_GRAINEDt,   HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_FINE_GRAINEDR   RD  RX  R0   t   _coarsegrain_mempoolt   _finegrain_mempoolRa  (   RM   Rn   t   qst   defqt   coarse_flagt	   fine_flagt   mpt	   alloc_mps(    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyRN     s     			(		c         C` s!   t  j | |  t j d  d  S(   Ni   (   R   t   _check_errorR+   R  (   RM   t   statusR"  (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyR5    s    c         C` s   |  S(   N(    (   RM   (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyt   unproxy  s    c         C` s   |  j  S(   N(   R6  (   RM   (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyt   default_queue  s    c         C` s   |  j  S(   N(   R  (   RM   (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyRn     s    c         C` s7   |  j  d  k r0 d j |  j  } t |   n  |  j  S(   Ns*   coarsegrain mempool is not available in {}(   R;  R0   RF   R  R3   (   RM   R  (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyt   coarsegrain_mempool  s    c         C` s7   |  j  d  k r0 d j |  j  } t |   n  |  j  S(   Ns(   finegrain mempool is not available in {}(   R<  R0   RF   R  R3   (   RM   R  (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyt   finegrain_mempool  s    c         C` sF  |  j  j } |  j  j } t   } t   } | d k sB | d k r1| d
 k	 r xl | D][ } d }	 x* | D]" }
 | j |
  rh |	 d 7}	 qh qh W|	 t |  k rU | j |  qU qU Wn | } x | D]f } | d k r | j t j	  r q n  | r| j
 r*| j |  q*q | j
 s | j |  q q Wn t d |   t |  d k s_t d   d
 } xE | D]= } y" t j |  j  |  j |  } Wn t k
 rqlXPqlW| d
 k rt d | | | f   n  t t j  } t t j |   | | d | |  |  } | j d
 k r,t d	   n  | |  j | j <| j   S(   s  
        Allocates memory.
        Parameters:
        nbytes the number of bytes to allocate.
        memTypeFlags the flags for which the memory region must have support,                     due to the inherent rawness of the underlying call, the                     validity of the flag is not checked, cf. C language.
        hostAccessible boolean as to whether the region in which the                       allocation takes place should be host accessible
        R   R   i    i   s   Unknown device type string "%s"s!   No suitable memory regions found.s   Memory allocation failed. No agent/region               combination could meet allocation restraints               (hardware = %s, size = %s, flags = %s).R  s   MemoryPointer has no valueN(   R  R   R   R   R0   Ra  RT   R7   R   t%   HSA_REGION_GLOBAL_FLAG_COARSE_GRAINEDR{  R$  R   R  R  Rg  R   t   _make_mem_finalizerR   R~  R  Rh   R  R   R8  R#  (   RM   Rd  t   memTypeFlagst   hostAccessiblet   hwt   all_regt	   flag_ok_rR   t   rt   countRf  t   memR   R  R  (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyt   memalloc  sX    				c         C` s   | r |  j  n |  j } | j |  } t t j  } t t j |   | | d | |  |  } | j	 |   | |  j
 | j <| j   S(   s   
        Allocates memory in a memory pool.
        Parameters:
        *nbytes* the number of bytes to allocate.
        *allow_acces_to*
        *finegrain*
        R  (   RH  RG  Rg  RJ  R   t   hsa_amd_memory_pool_freeR  Rh   R  R*  R8  R   R#  (   RM   Rd  R*  t	   finegraint   mempoolRe  R  RA  (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyt   mempoolalloc%  s    		c         C` sF   |  j  | d | d | } t t j |   d | d | j d | j S(   NR*  RU  Ri  R  R  (   RW  R,  Rh   R  R  R  (   RM   R  RU  R*  RR  (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyt   memhostalloc9  s    	N(    (   R   R   R   Rh   Ri   Rj   RN   R5  R   RE  RF  Rn   RG  RH  R0   RK   RS  RU   RW  RX  (    (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyR-    s   		KR   c           B` sS   e  Z d  Z d   Z d   Z d   Z d   Z d   Z d   Z e	 d    Z
 RS(   s.   
    An asynchronous stream for async API
    c         C` s   t    |  _ t t  |  _ d  S(   N(   R   t   _signalsR   R   t
   _callbacks(   RM   (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyRN   D  s    c         C` s9   t  |  j  d k r% |  j d  n  |  j j |  d S(   sA   
        Add a signal that corresponds to an async task.
        id   i2   N(   RT   RY  t   _syncR7   (   RM   R  (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyt   _add_signalH  s    c         C` s0   t  |  s t  |  j |  j   j |  d  S(   N(   t   callableR   RZ  t   _get_last_signalR7   (   RM   R}   (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyt   _add_callbackQ  s    c         C` s   |  j  r |  j  d Sd S(   s&   
        Get the last signal.
        iN(   RY  R0   (   RM   (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyR^  U  s    c         C` s   |  j  t |  j   d S(   s)   
        Synchronize the stream.
        N(   R[  RT   RY  (   RM   (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyR   [  s    c         C` s   d } x~ |  j  r | | k r" Pn  |  j  j   } | j   d k rP | j   n  x |  j | D] } |   q^ W|  j | =| d 7} q	 Wd  S(   Ni    i   (   RY  t   popleftR  R  RZ  (   RM   t   limitR(  t   sigR  (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyR[  a  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   (   RM   (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyt   auto_synchronizen  s    (   R   R   R   RN   R\  R_  R^  R   R[  R   Rc  (    (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyR   @  s   							c         ` s     f d   } | S(   s   
    finalises memory
    Parameters:
    dtor a function that will delete/free held memory from a reference

    Returns:
    Finalising function
    c         ` s.   |  j    t j       f d   } | S(   Nc           ` sL   t  j d      r7 t  j d  j     j =n        d  S(   Ns   Current allocations: %ss   Attempting delete on %s(   R   R   R   (    (   R8  t   dtorR.  t   sync(    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyt   core  s    (   R8  R   R   (   R  R.  Rf  (   Rd  (   R8  R.  Re  s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyt   mem_finalize  s    		(    (   Rd  Rg  (    (   Rd  s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyRJ  x  s    	c         C` s   t  |   j S(   s$   Get the device pointer as an integer(   R&  R   (   RP   (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyR    s    c         C` s'   |  d k r t d  St |   |  j S(   s,   Get the ctypes object for the device pointeri    N(   R0   Rb  t   require_device_memoryR&  (   RP   (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyR&    s    

c         C` s   t  |  d t  S(   sj  All HSA dGPU memory object is recognized as an instance with the
    attribute "__hsa_memory__" defined and its value evaluated to True.

    All HSA 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   RU   (   RP   (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyt   is_device_memory  s    c         C` s   t  |   s t d   n  d S(   s8   A sentry for methods that accept HSA memory object.
    s   Not a HSA memory object.N(   Ri  t	   Exception(   RP   (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyRh    s    c         C` s;   t  |  t t f  r |  St  |  t j  } t j |  |  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&  t   intt   longt   npt   voidR   t   memoryview_get_buffer(   RP   t   forcewritable(    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyR-    s    c         C` sO   t  j d  | d k  r, t d |   n  t j t |  t |  |  d S(   s  
    Copy data from a host memory region to a dGPU.
    Parameters:
    context the dGPU context
    dst a pointer to the destination location in dGPU memory
    src a pointer to the source location in host memory
    size the size (in bytes) of data to transfer
    s	   CPU->dGPUi    s   Invalid size given: %sN(   R   R   R3   R   t   hsa_memory_copyR  R-  (   R  t   dstt   srcR  (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyt   host_to_dGPU  s    	c         C` sO   t  j d  | d k  r, t d |   n  t j t |  t |  |  d S(   s  
    Copy data from a host memory region to a dGPU.
    Parameters:
    context the dGPU context
    dst a pointer to the destination location in dGPU memory
    src a pointer to the source location in host memory
    size the size (in bytes) of data to transfer
    s	   dGPU->CPUi    s   Invalid size given: %sN(   R   R   R3   R   Rq  R-  R  (   R  Rr  Rs  R  (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyt   dGPU_to_host  s    	c         C` sO   t  j d  | d k  r, t d |   n  t j t |  t |  |  d  S(   Ns
   dGPU->dGPUi    s   Invalid size given: %s(   R   R   R3   R   Rq  R  (   R  Rr  Rs  R  (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyt   dGPU_to_dGPU  s    c         C` sH   t  j d  t d |  d | d t |  d t |  d | d |  d  S(   Ns   Async CPU->dGPUt   dst_ctxt   src_ctxRs  Rr  R  t   stream(   R   R   t   async_copy_dgpuR-  R  (   Rw  Rx  Rr  Rs  R  Ry  (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyt   async_host_to_dGPU  s    c         C` sH   t  j d  t d |  d | d t |  d t |  d | d |  d  S(   Ns   Async dGPU->CPURw  Rx  Rr  Rs  R  Ry  (   R   R   Rz  R-  R  (   Rw  Rx  Rr  Rs  R  Ry  (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyt   async_dGPU_to_host  s    c         C` sH   t  j d  t d |  d | d t |  d t |  d | d |  d  S(   Ns   Async dGPU->dGPURw  Rx  Rr  Rs  R  Ry  (   R   R   Rz  R  (   Rw  Rx  Rr  Rs  R  Ry  (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyt   async_dGPU_to_dGPU  s    c   
      C` s   | d k  r t  d |   n  t j d  } | j   } | d  k	 rs t j | j  } d t j	 |  | f }	 n d d  | f }	 t j
 | |  j j | | j j | |	  | j |  d  S(   Ni    s   Invalid size given: %si   (   R3   R   R   R^  R0   R   R   R   R.   R   t   hsa_amd_memory_async_copyR  R\  (
   Rw  Rx  Rr  Rs  R  Ry  R  t   dependent_signalt   dsignalt   signals(    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyRz    s    
c          C` sQ   d }  y= x6 t  j D]+ } | j r | j d k r |  d 7}  q q WWn n X|  S(   sM   
    Returns the number of discrete GPUs present on the current machine.
    i    R   i   (   R   Rk   R   R   (   t   ngpusR   (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyt
   dgpu_count  s    (Z   R   t
   __future__R    R   R   R+   Rq   R$   R.   R)   R  Rh   t   loggingt
   contextlibR   t   collectionsR   R   t   numba.utilsR   t   numbaR   R   R	   t   errorR
   R   R   R   R   R   R   R   Rl  t	   numba.sixR   t   numpyRm  t	   getLoggerR   R   R   R   R'   RD   RE   R-   RG   R(   RH   R9   R   R^   RI   RW   R   R   Ry   R  R  R  R  R  R   R  R   R  R  R  R  t   MemAllocR,  R"  R-  R   RJ  R  R&  Ri  Rh  R-  Rt  Ru  Rv  R{  R|  R}  Rz  R  t   dgpu_present(    (    (    s6   lib/python2.7/site-packages/numba/roc/hsadrv/driver.pyt   <module>   s   		9	
			VW$@98														