σ
\K]c           @` sΝ   d  d l  m Z m Z m Z d  d l Z d  d l m Z m Z m	 Z	 d  d l
 m Z m Z e j rr d \ Z Z n d \ Z Z e e Z e e f Z d	 e e j f d
     YZ e d k rΙ e j   n  d S(   i    (   t   print_functiont   divisiont   absolute_importN(   t   cudat   configt   float32(   t   unittestt   SerialMixini   i   i2   i    t   TestCudaMatMulc           B` s   e  Z d    Z RS(   c   
   	   C` s  t  j d t d  d   d  d  d  f t d  d   d  d  d  f t d  d   d  d  d  f g  d    } t j j d  t j t j j t t f  d t j } t j t j j t t f  d t j } t j |  } t  j	   } | j
   t t  j | |  } t  j | |  } t  j | |  } | t t f t t f | f | | |  | j | |  Wd  QXt j | |  }	 t j j | |	 d d d  S(   Nt   argtypesi   c         S` sΉ  t  j j d t d t  } t  j j d t t f d t  } t  j j } t  j j } t  j	 j } t  j	 j } t  j
 j }	 t  j
 j }
 | | |	 } | | |
 } t d  } xΫ t t  D]Ν } | t k  r| t k  r|  | | | t f | | | f <| | | t | f | | | f <n  t  j   | t k  r|| t k  r|x9 t t  D]( } | | | | f | | | f 7} qMWn  t  j   qΉ W| t k  r΅| t k  r΅| | | | f <n  d  S(   Nt   shapet   dtypei    (   R   t   sharedt   arrayt   SM_SIZER   t   tpbt	   threadIdxt   xt   yt   blockIdxt   blockDimt   ranget   bpgt   nt   syncthreads(   t   At   Bt   Ct   sAt   sBt   txt   tyt   bxt   byt   bwt   bhR   R   t   acct   it   j(    (    sB   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_matmul.pyt   cu_square_matrix_mul   s,    !"%
)i*   R   t   rtolgρhγ΅ψδ>(   R   t   jitR   t   npt   randomt   seedR   R   t
   empty_liket   streamt   auto_synchronizet	   to_deviceR   R   t   copy_to_hostt   dott   testingt   assert_allclose(
   t   selfR'   R   R   R   R.   t   dAt   dBt   dCt   Cans(    (    sB   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_matmul.pyt	   test_func   s    o **)(   t   __name__t
   __module__R:   (    (    (    sB   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_matmul.pyR      s   t   __main__(   i   i   (   i2   i    (   t
   __future__R    R   R   t   numpyR*   t   numbaR   R   R   t   numba.cuda.testingR   R   t   ENABLE_CUDASIMR   R   R   R   t   TestCaseR   R;   t   main(    (    (    sB   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_matmul.pyt   <module>   s   	
8