ó
\K]c           @` sW  d  d l  m Z m Z m Z d  d l Z d  d l 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 „  Z d „  Z d „  Z d	 „  Z d
 „  Z d „  Z d „  Z d „  Z d „  Z d „  Z d „  Z d „  Z d „  Z d „  Z d „  Z d „  Z  d „  Z! d „  Z" d e e j# f d „  ƒ  YZ$ e% d k rSe j& ƒ  n  d S(   i    (   t   print_functiont   absolute_importt   divisionN(   t   cudat   int32t   float32(   t   unittestt   SerialMixint   skip_on_cudasim(   t   IS_PY3c         C` s   t  j j } | |  d <d  S(   Ni    (   R   t	   threadIdxt   x(   t   aryt   i(    (    sF   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_intrinsics.pyt   simple_threadidx
   s    c         C` s   t  j j } | |  | <d  S(   N(   R   R
   R   (   R   R   (    (    sF   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_intrinsics.pyt   fill_threadidx   s    c         C` sO   t  j j } t  j j } t  j j } | d | d | d |  | | | f <d  S(   Ni   (   R   R
   R   t   yt   z(   R   R   t   jt   k(    (    sF   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_intrinsics.pyt   fill3d_threadidx   s    c         C` s   t  j d ƒ } | |  | <d  S(   Ni   (   R   t   grid(   R   R   (    (    sF   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_intrinsics.pyt   simple_grid1d   s    c         C` s-   t  j d ƒ \ } } | | |  | | f <d  S(   Ni   (   R   R   (   R   R   R   (    (    sF   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_intrinsics.pyt   simple_grid2d!   s    c         C` s;   t  j d ƒ } t  j d ƒ } | d k r7 | |  d <n  d  S(   Ni   i    (   R   R   t   gridsize(   R   R   R   (    (    sF   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_intrinsics.pyt   simple_gridsize1d&   s    c         C` s]   t  j d ƒ \ } } t  j d ƒ \ } } | d k rY | d k rY | |  d <| |  d <n  d  S(   Ni   i    i   (   R   R   R   (   R   R   R   R   R   (    (    sF   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_intrinsics.pyt   simple_gridsize2d-   s
    
c   	      C` s¢   t  j d ƒ \ } } t  j j t  j j } t  j j t  j j } |  j \ } } xK t | | | ƒ D]7 } x. t | | | ƒ D] } | | |  | | f <q| Wqc Wd  S(   Ni   (   R   R   t   gridDimR   t   blockDimR   t   shapet   range(	   t   ct   startXt   startYt   gridXt   gridYt   heightt   widthR   R   (    (    sF   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_intrinsics.pyt   intrinsic_forloop_step5   s    c         C` s   t  j | ƒ |  d <d  S(   Ni    (   R   t   popc(   R   R   (    (    sF   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_intrinsics.pyt   simple_popc@   s    c         C` s   t  j | | | ƒ |  d <d  S(   Ni    (   R   t   fma(   R   t   at   bR   (    (    sF   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_intrinsics.pyt
   simple_fmaD   s    c         C` s   t  j | ƒ |  d <d  S(   Ni    (   R   t   brev(   R   R   (    (    sF   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_intrinsics.pyt   simple_brevH   s    c         C` s   t  j | ƒ |  d <d  S(   Ni    (   R   t   clz(   R   R   (    (    sF   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_intrinsics.pyt
   simple_clzL   s    c         C` s   t  j | ƒ |  d <d  S(   Ni    (   R   t   ffs(   R   R   (    (    sF   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_intrinsics.pyt
   simple_ffsP   s    c         C` s   t  | ƒ |  d <d  S(   Ni    (   t   round(   R   R   (    (    sF   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_intrinsics.pyt   simple_roundT   s    c         C` s[   t  j d ƒ } |  | d k rM | d d k r@ | | |  | <qW d |  | <n
 d |  | <d  S(   Ni   i   i   i    i   i   (   R   R   (   R*   R+   R   R   (    (    sF   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_intrinsics.pyt   branching_with_ifsX   s    c         C` sY   t  j d ƒ } t  j | d d k | | d ƒ } t  j |  | d k | d ƒ |  | <d  S(   Ni   i   i    i   i   i   (   R   R   t   selp(   R*   R+   R   R   t   inner(    (    sF   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_intrinsics.pyt   branching_with_selpsd   s    #c         C` s    t  j d ƒ } t  j |  | <d  S(   Ni   (   R   R   t   laneid(   R   R   (    (    sF   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_intrinsics.pyt   simple_laneidk   s    c         C` s   t  j |  d <d  S(   Ni    (   R   t   warpsize(   R   (    (    sF   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_intrinsics.pyt   simple_warpsizep   s    t   TestCudaIntrinsicc           B` ss  e  Z d  „  Z d „  Z d „  Z d „  Z d „  Z d „  Z e d ƒ d „  ƒ Z	 d „  Z
 d	 „  Z d
 „  Z d „  Z d „  Z d „  Z d „  Z d „  Z d „  Z e d ƒ d „  ƒ Z d „  Z d „  Z d „  Z d „  Z e d ƒ d „  ƒ Z d „  Z d „  Z d „  Z d „  Z e d ƒ d „  ƒ Z d „  Z d „  Z e  j! e" d ƒ d  „  ƒ Z# e  j! e" d ƒ d! „  ƒ Z$ RS("   c         C` sR   t  j d ƒ t ƒ } t j d d t j ƒ} | | ƒ |  j | d d k ƒ d  S(   Ns   void(int32[:])i   t   dtypei    (   R   t   jitR   t   npt   onesR   t
   assertTrue(   t   selft   compiledR   (    (    sF   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_intrinsics.pyt   test_simple_threadidxu   s    
c         C` s   t  j d ƒ t ƒ } d } t j | d t j ƒ} t j | d t j ƒ} | d | f | ƒ |  j t j | | k ƒ ƒ d  S(   Ns   void(int32[:])i
   R>   i   (	   R   R?   R   R@   RA   R   t   arangeRB   t   all(   RC   RD   t   NR   t   exp(    (    sF   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_intrinsics.pyt   test_fill_threadidx{   s    c         ` sk   d \ ‰  ‰ ‰ ‡  ‡ ‡ f d †  } ‡  ‡ ‡ f d †  } | ƒ  } | ƒ  } |  j  t j | | k ƒ ƒ d  S(   Ni   i   i   c          ` sW   t  j d ƒ t ƒ }  t j ˆ  ˆ ˆ f d t j ƒ} |  d ˆ  ˆ ˆ f f | ƒ | S(   Ns   void(int32[:,:,::1])R>   i   (   R   R?   R   R@   t   zerosR   (   RD   R   (   t   Xt   Yt   Z(    sF   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_intrinsics.pyt   c_contigous†   s    !c          ` s`   t  j d ƒ t ƒ }  t j t j ˆ  ˆ ˆ f d t j ƒƒ } |  d ˆ  ˆ ˆ f f | ƒ | S(   Ns   void(int32[::1,:,:])R>   i   (   R   R?   R   R@   t   asfortranarrayRK   R   (   RD   R   (   RL   RM   RN   (    sF   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_intrinsics.pyt   f_contigousŒ   s    *(   i   i   i   (   RB   R@   RG   (   RC   RO   RQ   t   c_rest   f_res(    (   RL   RM   RN   sF   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_intrinsics.pyt   test_fill3d_threadidxƒ   s    		c         C` s€   t  j d ƒ t ƒ } d \ } } | | } t j | d t j ƒ} | | | f | ƒ |  j t j | t j | ƒ k ƒ ƒ d  S(   Ns   void(int32[::1])i   i   R>   (   i   i   (	   R   R?   R   R@   t   emptyR   RB   RG   RF   (   RC   RD   t   ntidt   nctaidt   nelemR   (    (    sF   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_intrinsics.pyt   test_simple_grid1d–   s    
c   	      C` sí   t  j d ƒ t ƒ } d	 } d
 } | d | d | d | d f } t j | d t j ƒ} | j ƒ  } | | | f | ƒ xM t | j d ƒ D]8 } x/ t | j d ƒ D] } | | | | | f <q« Wq‘ W|  j	 t j
 | | k ƒ ƒ d  S(   Ns   void(int32[:,::1])i   i   i   i   i    i   R>   (   i   i   (   i   i   (   R   R?   R   R@   RU   R   t   copyR   R   RB   RG   (	   RC   RD   RV   RW   R   R   RI   R   R   (    (    sF   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_intrinsics.pyt   test_simple_grid2dž   s    $c         C` si   t  j d ƒ t ƒ } d \ } } t j d d t j ƒ} | | | f | ƒ |  j | d | | ƒ d  S(   Ns   void(int32[::1])i   i   i   R>   i    (   i   i   (   R   R?   R   R@   RK   R   t   assertEqual(   RC   RD   RV   RW   R   (    (    sF   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_intrinsics.pyt   test_simple_gridsize1d­   s
    s   Tests PTX emissionc   	      C` sc  t  j d ƒ t ƒ } t  j d ƒ t ƒ } d } d } t j d d d d d t j ƒ } | j ƒ  } d | d	 *t j | d t j ƒ} | | d
 f | | | ƒ | j	 ƒ  } |  j
 d t t j d | ƒ ƒ ƒ t j j | | d d ƒt j | d t j ƒ} | | d
 f | | | ƒ | j	 ƒ  } |  j
 d t t j d | ƒ ƒ ƒ t j j | | d d ƒd  S(   Ns   void(i8[:], i8, i8[:])i    i   R   t
   fill_valuei   R>   i   i   i   i   s	   \s+bra\s+t   err_msgt	   branchingi    R6   (   R   R?   R5   R8   R@   t   fullt   int64RZ   RF   t   inspect_asmR\   t   lent   ret   findallt   testingt   assert_array_equal(	   RC   t   cu_branching_with_ifst   cu_branching_with_selpst   nR+   R   t   expectedR*   t   ptx(    (    sF   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_intrinsics.pyt	   test_selp´   s"    !
""c         C` s‘   t  j d ƒ t ƒ } d
 } d } t j d d t j ƒ} | | | f | ƒ |  j | d | d | d ƒ |  j | d	 | d	 | d	 ƒ d  S(   Ns   void(int32[::1])i   i   i   i   i   R>   i    i   (   i   i   (   i   i   (   R   R?   R   R@   RK   R   R\   (   RC   RD   RV   RW   R   (    (    sF   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_intrinsics.pyt   test_simple_gridsize2dÌ   s     c         C` sK  t  j d ƒ t ƒ } d	 } d
 } | d | d | d | d f } t j | d t j ƒ} | | | f | ƒ | \ } } | j \ } }	 x¸ t t | d ƒ t | d ƒ ƒ D]“ \ }
 } | |
 | | } } xo t | |	 | ƒ D][ } xR t | | | ƒ D]> } |  j	 | | | f | | k | | | f | | f ƒ qý Wqä Wq° Wd  S(   Ns   void(float32[:,::1])i   i   i   i   i    i   R>   (   i   i   (   i   i   (
   R   R?   R&   R@   RU   R   R   t   zipR   RB   (   RC   RD   RV   RW   R   R   R"   R#   R$   R%   R   R   R    R!   R   R   (    (    sF   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_intrinsics.pyt   test_intrinsic_forloop_stepÖ   s    $0c         C` sd   t  j d „  ƒ } t j d d t j ƒj d d d ƒ } | d d f | ƒ t j j | d ƒ d  S(	   Nc         S` sO   t  j d ƒ \ } } } t  j d ƒ \ } } } | | | |  | | | f <d  S(   Ni   (   R   R   R   (   t   outR   R   R   R*   R+   R   (    (    sF   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_intrinsics.pyt   fooè   s    i	   i   R>   iÙ  (   i   i   i   (   i   i   i   iÙ  (   R   R?   R@   RK   R   t   reshapeRg   t   assert_equal(   RC   Rs   t   arr(    (    sF   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_intrinsics.pyt   test_3dgridç   s    'c         C` sƒ   t  j d „  ƒ } d d d } } } t j | | | d t j ƒj | | | ƒ } | d	 d
 f | ƒ |  j t j | ƒ ƒ d  S(   Nc   	      S` s  t  j d ƒ \ } } } t  j d ƒ \ } } } | t  j j t  j j t  j j k oŸ | t  j j t  j j t  j j k oŸ | t  j j t  j j t  j j k } | t  j j t  j	 j k oó | t  j j t  j	 j k oó | t  j j t  j	 j k } | oÿ | |  | | | f <d  S(   Ni   (
   R   R   R   R
   R   t   blockIdxR   R   R   R   (	   Rr   R   R   R   R*   R+   R   t   grid_is_rightt   gridsize_is_right(    (    sF   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_intrinsics.pyRs   ô   s    &&&i   i   i   R>   i   i   i   (   i   i   i   (   i   i   i   (   R   R?   R@   RK   t   boolRt   RB   RG   (   RC   Rs   R   R   R   Rv   (    (    sF   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_intrinsics.pyt   test_3dgrid_2ó   s
    /c         C` sR   t  j d ƒ t ƒ } t j d d t j ƒ} | | d ƒ |  j | d d ƒ d  S(   Ns   void(int32[:], uint32)i   R>   ið   i    i   (   R   R?   R(   R@   RK   R   t   assertEquals(   RC   RD   R   (    (    sF   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_intrinsics.pyt   test_popc_u4  s    c         C` sR   t  j d ƒ t ƒ } t j d d t j ƒ} | | d ƒ |  j | d d ƒ d  S(   Ns   void(int32[:], uint64)i   R>   I     ð  i    i   (   R   R?   R(   R@   RK   R   R}   (   RC   RD   R   (    (    sF   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_intrinsics.pyt   test_popc_u8  s    c         C` s[   t  j d ƒ t ƒ } t j d d t j ƒ} | | d d d ƒ t j j | d d ƒ d  S(   Ns   void(f4[:], f4, f4, f4)i   R>   g       @g      @g      @i    i   i   i   i   i
   (   R   R?   R,   R@   RK   R   Rg   t   assert_allclose(   RC   RD   R   (    (    sF   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_intrinsics.pyt   test_fma_f4  s    c         C` s[   t  j d ƒ t ƒ } t j d d t j ƒ} | | d d d ƒ t j j | d d ƒ d  S(   Ns   void(f8[:], f8, f8, f8)i   R>   g       @g      @g      @i    i   i   i   i   i
   (   R   R?   R,   R@   RK   t   float64Rg   R€   (   RC   RD   R   (    (    sF   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_intrinsics.pyt   test_fma_f8  s    c         C` sR   t  j d ƒ t ƒ } t j d d t j ƒ} | | d ƒ |  j | d d ƒ d  S(   Ns   void(uint32[:], uint32)i   R>   ið0  i    i  (   R   R?   R.   R@   RK   t   uint32R}   (   RC   RD   R   (    (    sF   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_intrinsics.pyt   test_brev_u4   s    s.   only get given a Python "int", assumes 32 bitsc         C` sR   t  j d ƒ t ƒ } t j d d t j ƒ} | | d ƒ |  j | d d ƒ d  S(   Ns   void(uint64[:], uint64)i   R>   Ið0  ð0  i    I    (   R   R?   R.   R@   RK   t   uint64R}   (   RC   RD   R   (    (    sF   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_intrinsics.pyt   test_brev_u8&  s    c         C` sR   t  j d ƒ t ƒ } t j d d t j ƒ} | | d ƒ |  j | d d ƒ d  S(   Ns   void(int32[:], int32)i   R>   i   i    i   (   R   R?   R0   R@   RK   R   R}   (   RC   RD   R   (    (    sF   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_intrinsics.pyt   test_clz_i4-  s    c         C` sR   t  j d ƒ t ƒ } t j d d t j ƒ} | | d ƒ |  j | d d ƒ d S(   s²  
        Although the CUDA Math API (http://docs.nvidia.com/cuda/cuda-math-api/group__CUDA__MATH__INTRINSIC__INT.html)
        only says int32 & int64 arguments are supported in C code, the LLVM
        IR input supports i8, i16, i32 & i64 (LLVM doesn't have a concept of
        unsigned integers, just unsigned operations on integers).
        http://docs.nvidia.com/cuda/nvvm-ir-spec/index.html#bit-manipulations-intrinics
        s   void(int32[:], uint32)i   R>   i   i    i   N(   R   R?   R0   R@   RK   R„   R}   (   RC   RD   R   (    (    sF   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_intrinsics.pyt   test_clz_u43  s    c         C` sR   t  j d ƒ t ƒ } t j d d t j ƒ} | | d ƒ |  j | d d ƒ d  S(   Ns   void(int32[:], int32)i   R>   Iÿÿÿÿ    i    (   R   R?   R0   R@   RK   R   R}   (   RC   RD   R   (    (    sF   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_intrinsics.pyt   test_clz_i4_1s@  s    c         C` sU   t  j d ƒ t ƒ } t j d d t j ƒ} | | d ƒ |  j | d d d ƒ d  S(   Ns   void(int32[:], int32)i   R>   i    i    s   CUDA semantics(   R   R?   R0   R@   RK   R   R}   (   RC   RD   R   (    (    sF   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_intrinsics.pyt   test_clz_i4_0sF  s    c         C` sR   t  j d ƒ t ƒ } t j d d t j ƒ} | | d ƒ |  j | d d ƒ d  S(   Ns   void(int32[:], int64)i   R>   i   i    i/   (   R   R?   R0   R@   RK   R   R}   (   RC   RD   R   (    (    sF   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_intrinsics.pyt   test_clz_i8L  s    c         C` sR   t  j d ƒ t ƒ } t j d d t j ƒ} | | d ƒ |  j | d d ƒ d  S(   Ns   void(int32[:], int32)i   R>   i   i    i   (   R   R?   R2   R@   RK   R   R}   (   RC   RD   R   (    (    sF   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_intrinsics.pyt   test_ffs_i4S  s    c         C` sR   t  j d ƒ t ƒ } t j d d t j ƒ} | | d ƒ |  j | d d ƒ d  S(   Ns   void(int32[:], uint32)i   R>   i   i    i   (   R   R?   R2   R@   RK   R„   R}   (   RC   RD   R   (    (    sF   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_intrinsics.pyt   test_ffs_u4Y  s    c         C` sR   t  j d ƒ t ƒ } t j d d t j ƒ} | | d ƒ |  j | d d ƒ d  S(   Ns   void(int32[:], int32)i   R>   Iÿÿÿÿ    i    (   R   R?   R2   R@   RK   R   R}   (   RC   RD   R   (    (    sF   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_intrinsics.pyt   test_ffs_i4_1s_  s    c         C` sU   t  j d ƒ t ƒ } t j d d t j ƒ} | | d ƒ |  j | d d d ƒ d  S(   Ns   void(int32[:], int32)i   R>   i    i    s   CUDA semantics(   R   R?   R2   R@   RK   R   R}   (   RC   RD   R   (    (    sF   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_intrinsics.pyt   test_ffs_i4_0se  s    c         C` sR   t  j d ƒ t ƒ } t j d d t j ƒ} | | d ƒ |  j | d d ƒ d  S(   Ns   void(int32[:], int64)i   R>   i   i    i   (   R   R?   R2   R@   RK   R   R}   (   RC   RD   R   (    (    sF   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_intrinsics.pyt   test_ffs_i8k  s    c         C` s“   t  j d ƒ t ƒ } d } t j | d d t j ƒ} t j t j d d t j ƒ| ƒ } | d | d f | ƒ |  j t j	 | | k ƒ ƒ d  S(   Ns   void(int32[:])i   i    R>   i   (
   R   R?   R:   R@   RK   R   t   tileRF   RB   RG   (   RC   RD   t   countR   RI   (    (    sF   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_intrinsics.pyt   test_simple_laneidr  s    $c         C` sR   t  j d ƒ t ƒ } t j d d t j ƒ} | | ƒ |  j | d d d ƒ d  S(   Ns   void(int32[:])i   R>   i    i    s   CUDA semantics(   R   R?   R<   R@   RK   R   R}   (   RC   RD   R   (    (    sF   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_intrinsics.pyt   test_simple_warpsizez  s    
s   round() returns float on Py2c         C` s   t  j d ƒ t ƒ } t j d d t j ƒ} xM d d d d d d	 d
 d g D]- } | | | ƒ |  j | d t | ƒ ƒ qL Wd  S(   Ns   void(int64[:], float32)i   R>   g      Àg      Àg      Àg      ø¿g      ø?g      @g      @g      @i    (   R   R?   R4   R@   RK   R   R}   R3   (   RC   RD   R   R   (    (    sF   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_intrinsics.pyt   test_round_f4€  s
    %c         C` s   t  j d ƒ t ƒ } t j d d t j ƒ} xM d d d d d d	 d
 d g D]- } | | | ƒ |  j | d t | ƒ ƒ qL Wd  S(   Ns   void(int64[:], float64)i   R>   g      Àg      Àg      Àg      ø¿g      ø?g      @g      @g      @i    (   R   R?   R4   R@   RK   R   R}   R3   (   RC   RD   R   R   (    (    sF   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_intrinsics.pyt   test_round_f8‰  s
    %(%   t   __name__t
   __module__RE   RJ   RT   RY   R[   R]   R   Rn   Ro   Rq   Rw   R|   R~   R   R   Rƒ   R…   R‡   Rˆ   R‰   RŠ   R‹   RŒ   R   RŽ   R   R   R‘   R”   R•   R   t
   skipUnlessR	   R–   R—   (    (    (    sF   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_intrinsics.pyR=   t   s>   							
																			t   __main__('   t
   __future__R    R   R   t   numpyR@   Re   t   numbaR   R   R   t   numba.cuda.testingR   R   R   t   numba.utilsR	   R   R   R   R   R   R   R   R&   R(   R,   R.   R0   R2   R4   R5   R8   R:   R<   t   TestCaseR=   R˜   t   main(    (    (    sF   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_intrinsics.pyt   <module>   s6   																		ÿ  