ó
\K]c           @@ s9  d  d l  m Z m Z d  d l Z d  d l Z d  d l Z y d  d l Z Wn e k
 ri d  d l Z n Xd  d l	 Z
 d  d l j Z d  d l m Z m Z m Z m Z m Z m Z m Z d  d l m Z d  d l j j Z d  d l m Z 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
 e j" f d „  ƒ  YZ# d e j" f d „  ƒ  YZ$ d e j" f d „  ƒ  YZ% d „  Z& d „  Z' e' ƒ  Z( e j) e( d ƒ d e j" f d „  ƒ  Yƒ Z* e j) e( d ƒ d e% f d „  ƒ  Yƒ Z+ d e% f d „  ƒ  YZ, d e% f d „  ƒ  YZ- e. d k r5e j/ ƒ  n  d S(   i    (   t   print_functiont   absolute_importN(   t   hsat   Queuet   Programt
   Executablet
   BrigModulet   Contextt   dgpu_present(   R   (   t   float32t   float64t	   vectorize(   t   drvapi(   t   enums(   t	   enums_ext(   t   configt   TestLowLevelApic           B@ s   e  Z d  Z d „  Z RS(   sY   This test checks that all the functions defined in drvapi
    bind properly using ctypes.c         C@ s   g  } x] t  j j ƒ  D]L } y t t | ƒ Wq t k
 ra } | j d j | t | ƒ ƒ ƒ q Xq W|  j	 t
 | ƒ d d d j | ƒ ƒd  S(   Ns
   '{0}': {1}i    t   msgs   
(   R   t   API_PROTOTYPESt   keyst   getattrR   t	   Exceptiont   appendt   formatt   strt   assertEqualt   lent   join(   t   selft   missing_functionst   fnamet   e(    (    sA   lib/python2.7/site-packages/numba/roc/tests/hsadrv/test_driver.pyt   test_functions_available   s    '(   t   __name__t
   __module__t   __doc__R    (    (    (    sA   lib/python2.7/site-packages/numba/roc/tests/hsadrv/test_driver.pyR      s   t
   TestAgentsc           B@ s,   e  Z d  „  Z d „  Z d „  Z d „  Z RS(   c         C@ s   |  j  t t j ƒ d ƒ d  S(   Ni    (   t   assertGreaterR   t   roct   agents(   R   (    (    sA   lib/python2.7/site-packages/numba/roc/tests/hsadrv/test_driver.pyt   test_agents_init+   s    c         C@ sC   x< t  j D]1 } | j r
 | j d ƒ } |  j | t ƒ q
 q
 Wd  S(   Ni   i   i    (   R&   R'   t   is_componentt   create_queue_singlet   assertIsInstanceR   (   R   t   agentt   queue(    (    sA   lib/python2.7/site-packages/numba/roc/tests/hsadrv/test_driver.pyt   test_agents_create_queue_single.   s    	c         C@ sC   x< t  j D]1 } | j r
 | j d ƒ } |  j | t ƒ q
 q
 Wd  S(   Ni   i   i    (   R&   R'   R)   t   create_queue_multiR+   R   (   R   R,   R-   (    (    sA   lib/python2.7/site-packages/numba/roc/tests/hsadrv/test_driver.pyt   test_agents_create_queue_multi4   s    	c         C@ sO   xH t  j D]= } | j r
 | j j ƒ  d k rG |  j | j d ƒ qG q
 q
 Wd  S(   Nt   gfx803t   gfx900i   (   R1   R2   (   R&   R'   R)   t   namet   decodeR   t   wavebits(   R   R,   (    (    sA   lib/python2.7/site-packages/numba/roc/tests/hsadrv/test_driver.pyt   test_agent_wavebits:   s    	(   R!   R"   R(   R.   R0   R6   (    (    (    sA   lib/python2.7/site-packages/numba/roc/tests/hsadrv/test_driver.pyR$   *   s   			t	   _TestBasec           B@ s   e  Z d  „  Z d „  Z RS(   c         C@ sw   g  t  j D] } | j r
 | ^ q
 d |  _ g  t  j D] } | j s6 | ^ q6 d |  _ |  j j |  j j ƒ |  _ d  S(   Ni    (   R&   R'   R)   t   gput   cpuR/   t   queue_max_sizeR-   (   R   t   a(    (    sA   lib/python2.7/site-packages/numba/roc/tests/hsadrv/test_driver.pyt   setUpB   s    ,,c         C@ s   |  `  |  ` |  ` d  S(   N(   R-   R8   R9   (   R   (    (    sA   lib/python2.7/site-packages/numba/roc/tests/hsadrv/test_driver.pyt   tearDownG   s    (   R!   R"   R<   R=   (    (    (    sA   lib/python2.7/site-packages/numba/roc/tests/hsadrv/test_driver.pyR7   A   s   	c          C@ s.   t  j j d ƒ }  t  j j |  ƒ s* t ‚ |  S(   Ns*   /opt/rocm/hsa/sample/vector_copy_full.brig(   t   ost   pathR   t   isfilet   AssertionError(   R?   (    (    sA   lib/python2.7/site-packages/numba/roc/tests/hsadrv/test_driver.pyt   get_brig_fileM   s    c           C@ s$   y t  ƒ  Wn t k
 r t SXt S(   N(   RB   t   BaseExceptiont   Falset   True(    (    (    sA   lib/python2.7/site-packages/numba/roc/tests/hsadrv/test_driver.pyt   _check_example_fileR   s
    s   Brig example not foundt   TestBrigModulec           B@ s   e  Z d  „  Z RS(   c         C@ s2   t  ƒ  } t j | ƒ } |  j t | ƒ d ƒ d  S(   Ni    (   RB   R   t	   from_fileR%   R   (   R   t	   brig_filet   brig_module(    (    sA   lib/python2.7/site-packages/numba/roc/tests/hsadrv/test_driver.pyt   test_from_file^   s    	(   R!   R"   RK   (    (    (    sA   lib/python2.7/site-packages/numba/roc/tests/hsadrv/test_driver.pyRG   \   s   t   TestProgramc           B@ s   e  Z d  „  Z RS(   c         C@ s›   t  ƒ  } d } t j | ƒ } t ƒ  } | j | ƒ | j |  j j ƒ } t ƒ  } | j	 |  j | ƒ | j
 ƒ  | j |  j | ƒ } |  j | j d ƒ d  S(   Ns   &__vector_copy_kerneli    (   RB   R   RH   R   t
   add_modulet   finalizeR8   t   isaR   t   loadt   freezet
   get_symbolR%   t   kernarg_segment_size(   R   RI   t   symbolRJ   t   programt   codet   ext   sym(    (    sA   lib/python2.7/site-packages/numba/roc/tests/hsadrv/test_driver.pyt   test_create_programf   s    			
(   R!   R"   RY   (    (    (    sA   lib/python2.7/site-packages/numba/roc/tests/hsadrv/test_driver.pyRL   d   s   t
   TestMemoryc           B@ sz   e  Z d  „  Z d „  Z d „  Z e j e d ƒ d „  ƒ Z e j e	 d ƒ e j e d ƒ e j
 d ƒ d „  ƒ ƒ ƒ Z RS(   c         C@ sB   |  j  t |  j j j ƒ d ƒ |  j  t |  j j j ƒ d ƒ d  S(   Ni    (   R%   R   R8   t   regionst   globalst   groups(   R   (    (    sA   lib/python2.7/site-packages/numba/roc/tests/hsadrv/test_driver.pyt   test_region_listw   s    c         C@ sT   t  j j d ƒ j t  j ƒ } t j | j j | j ƒ t j	 | j j | j ƒ d  S(   Ni   (
   t   npt   randomt   astypeR	   R&   t   hsa_memory_registert   ctypest   datat   nbytest   hsa_memory_deregister(   R   t   src(    (    sA   lib/python2.7/site-packages/numba/roc/tests/hsadrv/test_driver.pyt   test_register~   s    c   
      C@ sU  |  j  j } |  j t | ƒ d ƒ t ƒ  } x0 | D]( } | j t j ƒ r2 | j | ƒ q2 q2 W|  j t | ƒ d ƒ | d } d } | j	 t
 j t
 j ƒ | ƒ } |  j t
 j | ƒ d d ƒ t j j | ƒ j t j ƒ } t
 j | | j
 j | j ƒ t
 j | j | j ƒ } x/ t | j ƒ D] }	 |  j | |	 | |	 ƒ q"Wt j | ƒ d  S(   Ni    i
   s   pointer must not be NULL(   R8   R[   R%   R   t   listt   supportsR   t   HSA_REGION_GLOBAL_FLAG_KERNARGR   t   allocateRc   t   sizeoft   c_floatt   assertNotEqualt	   addressofR_   R`   Ra   R	   t   memmoveRd   Re   t   from_addresst   valuet   ranget   sizeR   R&   t   hsa_memory_free(
   R   R[   t   kernarg_regionst   rt   kernarg_regiont   nelemt   ptrRg   t   reft   i(    (    sA   lib/python2.7/site-packages/numba/roc/tests/hsadrv/test_driver.pyt   test_allocateƒ   s$    	
s	   dGPU onlyc         C@ sé  |  j  j } t ƒ  } t ƒ  } xI | D]A } | j t j ƒ r% | j rV | j | ƒ qf | j | ƒ q% q% W|  j t	 | ƒ d ƒ |  j t	 | ƒ d ƒ |  j
 j } t ƒ  } x0 | D]( } | j t j ƒ r² | j | ƒ q² q² W|  j t	 | ƒ d ƒ d } | d } | j t j t j ƒ | ƒ }	 |  j t j |	 ƒ d d ƒ | d }
 |
 j t j t j ƒ | ƒ } |  j t j | ƒ d d ƒ | d } | j t j t j ƒ | ƒ } |  j t j | ƒ d d ƒ t j j | ƒ j t j ƒ } t j |	 | j j | j ƒ t j | |	 | j ƒ t j | | | j ƒ t j | j |	 j ƒ } x/ t | j ƒ D] } |  j | | | | ƒ qXWt j | j | j ƒ } x/ t | j ƒ D] } |  j | | | | ƒ q£Wt j | ƒ j t j ƒ } t j |	 | j j | j ƒ t j | |	 | j ƒ x/ t | j ƒ D] } |  j | | | | ƒ q"Wx/ t | j ƒ D] } |  j | | | | ƒ qTWt j | | | j ƒ x/ t | j ƒ D] } |  j | | | | ƒ qœWt j |	 ƒ t j | ƒ t j | ƒ d S(   sA  
        Tests the coarse grained allocation works on a dGPU.
        It performs a data copying round trip via:
        memory
          |
        HSA cpu memory
          |
        HSA dGPU host accessible memory <---|
          |                                 |
        HSA dGPU memory --------------------|
        i    i
   s   pointer must not be NULLN(    R8   R[   Ri   Rj   R   t%   HSA_REGION_GLOBAL_FLAG_COARSE_GRAINEDt   host_accessibleR   R%   R   R9   Rl   Rc   Rm   Rn   Ro   Rp   R_   R`   Ra   R	   R&   t   hsa_memory_copyRd   Re   Rr   Rs   Rt   Ru   R   t   zerosRv   (   R   t   gpu_regionst   gpu_only_coarse_regionst"   gpu_host_accessible_coarse_regionsRx   t   cpu_regionst   cpu_coarse_regionsRz   t
   cpu_regiont   cpu_ptrt   gpu_only_regiont   gpu_only_ptrt   gpu_host_accessible_regiont   gpu_host_accessible_ptrRg   t   cpu_refR}   t
   gpu_ha_reft   z0(    (    sA   lib/python2.7/site-packages/numba/roc/tests/hsadrv/test_driver.pyt   test_coarse_grained_allocate   sn    				



s   Brig example not founds6   Permanently skip? HSA spec violation causes corruptionc   #   	   C@ sQ  d d l  m } m } m } m } t ƒ  } | j | ƒ } |  j t | ƒ d ƒ |  j	 j
 } t ƒ  } t ƒ  }	 xI | D]A }
 |
 j t j ƒ ru |
 j r¦ |	 j |
 ƒ q¶ | j |
 ƒ qu qu W|  j t | ƒ d ƒ |  j t |	 ƒ d ƒ t j d } | ƒ  } | j | ƒ | j | j ƒ } | ƒ  } | j | | ƒ | j ƒ  | j | d ƒ } |  j | j d ƒ |  j | j d ƒ d d l } d d l } d } | j j | ƒ j | j  ƒ } | j! | ƒ } | j" | j# ƒ | } |	 d } | j$ | ƒ } |  j | j% d d ƒ | j$ | ƒ } |  j | j% d d ƒ t j' | | j j( | j) ƒ t j' | | j j( | j) ƒ | d } | j$ | ƒ } |  j | j% d d ƒ | j$ | ƒ } |  j | j% d d ƒ t j' | | | j) ƒ t ƒ  } x0 |	 D]( }
 |
 j t j* ƒ rÝ| j |
 ƒ qÝqÝW|  j t | ƒ d ƒ | d } | j$ d | j" | j+ ƒ ƒ } |  j | d d ƒ d | j, j- | j% ƒ } | j% | d <| j% | d <t j. d ƒ }  | j/ d	 ƒ }! |! j0 | | d
 d d | d d f d d ƒt j' | | | j) ƒ | | j# j- | j% ƒ }" | j1 j2 |" | ƒ t j3 | ƒ t j3 | ƒ t j3 | ƒ t j3 | ƒ d S(   sô   
        This tests the execution of a kernel on a dGPU using coarse memory
        regions for the buffers.
        NOTE: the code violates the HSA spec in that it uses a coarse region
        for kernargs, this is a performance hack.
        i    (   R   R   R   R   s   &__vector_copy_kernelNi   s   pointer must not be NULLi   i   i    t   workgroup_sizei   t	   grid_sizet   signali   (   i   i   i   (4   t   numba.roc.hsadrv.driverR   R   R   R   RB   RH   R%   R   R8   R[   Ri   Rj   R   R   R€   R   R&   t
   componentsRM   RN   RO   RP   RQ   RR   Ro   t   kernel_objectRS   Rc   t   numpyR`   Ra   R	   t
   zeros_likeRm   Rn   Rl   Rs   t   NoneR   Rd   Re   Rk   t   c_void_pt   c_size_tRr   t   create_signalR*   t   dispatcht   testingt   assert_equalRv   (#   R   R   R   R   R   RI   RJ   Rƒ   R„   R…   Rx   R,   t   progRV   RW   RX   Rc   R_   Rz   Rg   R   Re   RŒ   t   host_in_ptrt   host_out_ptrRŠ   t
   gpu_in_ptrt   gpu_out_ptrRw   Ry   t   kernarg_ptrt   argreft   sigR-   R|   (    (    sA   lib/python2.7/site-packages/numba/roc/tests/hsadrv/test_driver.pyt$   test_coarse_grained_kernel_executioný   sŠ    "						


	
(   R!   R"   R^   Rh   R~   t   unittestt
   skipUnlessR   R‘   t   has_brig_examplet   skipR©   (    (    (    sA   lib/python2.7/site-packages/numba/roc/tests/hsadrv/test_driver.pyRZ   v   s   			`t   TestContextc           B@ s×   e  Z d  Z d „  Z e d „ Z d „  Z e j e	 d ƒ d „  ƒ Z
 e j e	 d ƒ d „  ƒ Z d „  Z e j e	 d ƒ d „  ƒ Z e j e	 d ƒ d	 „  ƒ Z e j e	 d ƒ d
 „  ƒ Z e j e	 d ƒ d „  ƒ Z RS(   s-   Tests the Context class behaviour is correct.c         C@ sÇ  d } t  j t  j ƒ | } t rÃ|  j } |  j } t | ƒ } | j | d t ƒ} | j | d t	 ƒ} t | ƒ } | j | d t	 ƒ}	 t
 j j | ƒ j t
 j ƒ }
 t j |	 j |
 j  j |
 j ƒ t j | j |	 j |
 j ƒ t j | j | j |
 j ƒ t
 j |
 ƒ } t j | j | j  j | j ƒ | t  j j | j j ƒ } x( t | ƒ D] } |  j | | d ƒ qWWt j | j | j |
 j ƒ x/ t | ƒ D] } |  j | | |
 | ƒ qžWn  d S(   sÈ   
            Tests Context.memalloc() for a given, in the parlance of HSA,            `component`. Testing includes specialisations for the supported
            components of dGPUs and APUs.
        i
   t   hostAccessiblei    N(   Rc   Rm   t   c_doubleR   R8   R9   R   t   memallocRD   RE   R_   R`   Ra   R
   R&   R   t   device_pointerRd   Re   R™   Rr   Rs   Rt   R   (   R   t   nRe   t
   dGPU_agentt	   CPU_agentt   gpu_ctxt   gpu_only_memt   ha_memt   cpu_ctxt   cpu_memRg   R   R|   t   k(    (    sA   lib/python2.7/site-packages/numba/roc/tests/hsadrv/test_driver.pyt   test_memallocŠ  s.    		c   
      C@ s¾  g  | j  D] } | j r
 | ^ q
 } i g  d 6g  d 6g  d 6g  d 6} x" | D] } | | j j | ƒ qN W| r |  j t | d ƒ d ƒ n |  j t | d ƒ d ƒ |  j t | d ƒ d ƒ |  j t | d ƒ d ƒ |  j t | d ƒ d ƒ |  j t | j  j ƒ t | d ƒ ƒ | d } d  } d  } xD | D]< }	 |	 j t	 j
 ƒ r[|	 } n  |	 j t	 j ƒ r:|	 } q:q:W|  j | ƒ | r|  j | ƒ n |  j | ƒ |  j | | ƒ d  S(   Nt   globalt   readonlyt   privatet   groupi   i   i    (   t   mempoolst   alloc_allowedt   kindR   R   R   R\   Rš   Rj   R   t.   HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_COARSE_GRAINEDt,   HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_FINE_GRAINEDt   assertIsNotNonet   assertIsNonet   assertIsNot(
   R   R,   t   has_fine_graint   mpt   mp_alloc_listt   mpdctt   glbst   coarsegraint	   finegraint   gmp(    (    sA   lib/python2.7/site-packages/numba/roc/tests/hsadrv/test_driver.pyt   check_mempoolsµ  s0    %"&
	c         C@ s   |  j  |  j ƒ d  S(   N(   RÑ   R9   (   R   (    (    sA   lib/python2.7/site-packages/numba/roc/tests/hsadrv/test_driver.pyt   test_cpu_mempool_propertyÙ  s    s	   dGPU onlyc         C@ s   |  j  |  j d t ƒd  S(   NRÉ   (   RÑ   R8   RD   (   R   (    (    sA   lib/python2.7/site-packages/numba/roc/tests/hsadrv/test_driver.pyt   test_gpu_mempool_propertyÜ  s    c         C@ s  d } t  j t  j ƒ | } |  j } |  j } t | ƒ } | j | ƒ } t | ƒ } | j | d | j g ƒ} t j	 j	 | ƒ j
 t j ƒ }	 t j | j |	 j  j |	 j ƒ t j | j | j |	 j ƒ t j |	 ƒ }
 t j | j |
 j  j |
 j ƒ | t  j j | j j ƒ } x( t | ƒ D] } |  j | | d ƒ q Wt j | j | j |	 j ƒ x, t | ƒ D] } |  j | | |	 | ƒ qgWd  S(   Ni
   t   allow_access_toi    (   Rc   Rm   R°   R8   R9   R   t   mempoolallocR,   R_   R`   Ra   R
   R&   R   R²   Rd   Re   R™   Rr   Rs   Rt   R   (   R   R³   Re   R´   Rµ   R¶   R·   R¹   Rº   Rg   R   R|   R»   (    (    sA   lib/python2.7/site-packages/numba/roc/tests/hsadrv/test_driver.pyt   test_mempoolà  s&    		c         C@ sG   |  j  } t | ƒ } |  j } t | ƒ } | j d d | j g ƒd  S(   Ni   RÔ   (   R8   R   R9   RÕ   t   _agent(   R   RÏ   R´   R¶   Rµ   R¹   (    (    sA   lib/python2.7/site-packages/numba/roc/tests/hsadrv/test_driver.pyt   check_mempool_with_flags  s
    		c         C@ s   |  j  d t ƒ d  S(   NRÏ   (   RØ   RE   (   R   (    (    sA   lib/python2.7/site-packages/numba/roc/tests/hsadrv/test_driver.pyt   test_mempool_finegrained  s    c         C@ s   |  j  d t ƒ d  S(   NRÏ   (   RØ   RD   (   R   (    (    sA   lib/python2.7/site-packages/numba/roc/tests/hsadrv/test_driver.pyt   test_mempool_coarsegrained  s    c      
   @ s‹  |  j  } t | ƒ ‰ |  j } t | ƒ ‰ d } | t j t j ƒ ‰ t j d ƒ ‰ t j d ƒ ‰  t d ˆ j	 g d t
 ƒ } ˆ j ˆ |  ‰ ˆ j ˆ |  ‰ d } ˆ j ˆ ƒ ‰ | t j j ˆ j j ƒ } | t j j ˆ j j ƒ } | d t j d | d t j ƒ| (t j | d t j ƒ| (t j ˆ  d ƒ t j ƒ  ‰	 d t j f ‡  ‡	 f d	 †  ƒ  Y} t j ˆ d ƒ t j d
 d ƒ ‰ d t j f ‡  ‡ ‡ ‡ ‡ ‡ ‡ ‡ f d †  ƒ  Y}	 d t j f ‡ ‡ ‡ ‡ ‡ ‡ ‡ f d †  ƒ  Y}
 d } | ƒ  } |	 ƒ  } |
 ƒ  } | j ƒ  | j ƒ  | j ƒ  | j | ƒ | j | ƒ | j | ƒ ˆ	 j ƒ  } |  j | d ƒ t j j | | ƒ d  S(   Ni   i    RÔ   RÏ   ià  t   dtypei   t   validatorThreadc           @ s   e  Z ‡  ‡ f d  †  Z RS(   c         @ s;   t  j ˆ  t j d t j d ƒ t j ƒ } ˆ j | ƒ d  S(   Ni    iÿÿÿÿ(   R&   t   hsa_signal_wait_acquireR   t   HSA_SIGNAL_CONDITION_EQRc   t   c_uint64t   HSA_WAIT_STATE_ACTIVEt   put(   R   t   val(   t   completion_signalt   q(    sA   lib/python2.7/site-packages/numba/roc/tests/hsadrv/test_driver.pyt   run:  s    (   R!   R"   Rå   (    (   Rã   Rä   (    sA   lib/python2.7/site-packages/numba/roc/tests/hsadrv/test_driver.pyRÜ   9  s   Rs   t	   l2hThreadc           @ s,   e  Z ‡  ‡ ‡ ‡ ‡ ‡ ‡ ‡ f d  †  Z RS(   c      	   @ sc   t  j ˆ j ƒ } t j ˆ j j ˆ j j ˆ j j ˆ j j ˆ d t j	 | ƒ ˆ  ƒ ˆ j
 ƒ  d  S(   Ni   (   R   t   hsa_signal_tt   _idR&   t   hsa_amd_memory_async_copyR²   Rs   R×   Rc   t   byreft   release(   R   t
   dep_signal(   Rã   R¹   t   dependent_signalR¶   t	   h2l_startt   host_dstt   kSizet   local_memory(    sA   lib/python2.7/site-packages/numba/roc/tests/hsadrv/test_driver.pyRå   J  s    		(   R!   R"   Rå   (    (   Rã   R¹   Rí   R¶   Rî   Rï   Rð   Rñ   (    sA   lib/python2.7/site-packages/numba/roc/tests/hsadrv/test_driver.pyRæ   I  s   t	   h2lThreadc           @ s)   e  Z ‡  ‡ ‡ ‡ ‡ ‡ ‡ f d  †  Z RS(   c      	   @ sH   ˆ j  ƒ  t j ˆ j j ˆ j j ˆ j j ˆ  j j ˆ d d  ˆ ƒ d  S(   Ni    (   t   acquireR&   Ré   R²   Rs   R×   Rè   Rš   (   R   (   R¹   Rí   R¶   Rî   t   host_srcRð   Rñ   (    sA   lib/python2.7/site-packages/numba/roc/tests/hsadrv/test_driver.pyRå   U  s    
		(   R!   R"   Rå   (    (   R¹   Rí   R¶   Rî   Rô   Rð   Rñ   (    sA   lib/python2.7/site-packages/numba/roc/tests/hsadrv/test_driver.pyRò   T  s   i
   (   R8   R   R9   Rc   Rm   t   c_intR&   R   t   dictR,   RD   RÕ   Rr   R²   Rs   R_   t   aranget   int32R‚   t   hsa_signal_store_relaxedR-   R   t	   threadingt   Threadt	   Semaphoret   startR   t   getR   RŸ   t   assert_allclose(   R   R´   Rµ   t   kNumIntt   flagsR}   t   host_src_viewt   host_dst_viewRÜ   Ræ   Rò   t   timeoutt	   validatort   l2ht   h2lt   wait_res(    (
   Rã   R¹   Rí   R¶   Rî   Rï   Rô   Rð   Rñ   Rä   sA   lib/python2.7/site-packages/numba/roc/tests/hsadrv/test_driver.pyt   test_mempool_amd_example  sJ    		$"41
			


c         C@ s—   d } t  j | ƒ } t  j | ƒ } t d d d ƒd „  ƒ } t j | ƒ t j | ƒ } | | d | ƒ| j ƒ  } t  j j t  j | ƒ | ƒ d S(   s<   
            Tests .to_device() and .copy_to_host()
        i
   s   float64(float64)t   targetR&   c         S@ s   |  d S(   Ni   (    (   t   x(    (    sA   lib/python2.7/site-packages/numba/roc/tests/hsadrv/test_driver.pyt   funcx  s    t   outN(	   R_   R‚   R   t   hsaapit	   to_devicet   copy_to_hostRŸ   R    t   ones(   R   R³   Rd   t   outputR  t
   out_devicet   host_output(    (    sA   lib/python2.7/site-packages/numba/roc/tests/hsadrv/test_driver.pyt   test_to_device_to_hostp  s    (   R!   R"   R#   R¼   RE   RÑ   RÒ   Rª   R«   R   RÓ   RÖ   RØ   RÙ   RÚ   R	  R  (    (    (    sA   lib/python2.7/site-packages/numba/roc/tests/hsadrv/test_driver.pyR®   ‡  s   	+$	"	
\t   __main__(0   t
   __future__R    R   Rc   R>   Rú   R-   t   ImportErrorR   R˜   R_   t   numba.unittest_supportt   unittest_supportRª   R•   R   R   R   R   R   R   R&   t   numba.roc.apit   apiR  t   numbaR	   R
   R   t   numba.roc.hsadrvR   R   R   R   t   TestCaseR   R$   R7   RB   RF   R¬   R«   RG   RL   RZ   R®   R!   t   main(    (    (    sA   lib/python2.7/site-packages/numba/roc/tests/hsadrv/test_driver.pyt   <module>   s@   4			ÿ ü