ó
\K]c           @  s/  d  d l  m Z d  d l Z d  d l m Z d  d l m Z m Z d  d l	 m
 Z
 e j g  ƒ Z e j d d e j ƒd Z e j e j d	 d e j ƒj d d ƒ ƒ Z e j d) d e j ƒj d
 d
 d
 ƒ d d Z e j d d e j ƒZ e j g  d d e f d e f g ƒZ e j d* d+ g d d e f d e f g ƒZ e j d, d- g d e j d d e j f d e j f d e j f d e j f d e j f g d e ƒ ƒ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 r+e j+ ƒ  n  d S(.   iÿÿÿÿ(   t   print_functionN(   t   cuda(   t   unittestt   SerialMixin(   t   ENABLE_CUDASIMi
   t   dtypeg       @id   i   y              ð?y               @i   t   xt   yg      ð?i   g      @i   i   Iï¾­Þ    i   i   I­Þï¾    t   at   bt   zt   alignc         C  s5   t  j j t ƒ } t  j d ƒ } t | ƒ |  | <d  S(   Ni   (   R   t   constt
   array_liket   CONST_EMPTYt   gridt   len(   t   At   Ct   i(    (    sD   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_constmem.pyt   cuconstEmpty$   s    c         C  s7   t  j j t ƒ } t  j d ƒ } | | d |  | <d  S(   Ni   g      ð?(   R   R   R   t   CONST1DR   (   R   R   R   (    (    sD   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_constmem.pyt   cuconst*   s    c         C  sE   t  j j t ƒ } t  j d ƒ \ } } | | | f |  | | f <d  S(   Ni   (   R   R   R   t   CONST2DR   (   R   R   R   t   j(    (    sD   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_constmem.pyt	   cuconst2d2   s    c         C  sZ   t  j j t ƒ } t  j j } t  j j } t  j j } | | | | f |  | | | f <d  S(   N(   R   R   R   t   CONST3Dt	   threadIdxR   R   R
   (   R   R   R   R   t   k(    (    sD   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_constmem.pyt	   cuconst3d8   s
    c         C  s5   t  j j t ƒ } t  j d ƒ } t | ƒ |  | <d  S(   Ni   (   R   R   R   t   CONST_RECORD_EMPTYR   R   (   R   R   R   (    (    sD   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_constmem.pyt   cuconstRecEmpty@   s    c         C  sI   t  j j t ƒ } t  j d ƒ } | | d |  | <| | d | | <d  S(   Ni   R   R   (   R   R   R   t   CONST_RECORDR   (   R   t   BR   R   (    (    sD   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_constmem.pyt
   cuconstRecF   s    c         C  s   t  j j t ƒ } t  j d ƒ } | | d |  | <| | d | | <| | d | | <| | d | | <| | d | | <d  S(   Ni   R   R	   R   R   R
   (   R   R   R   t   CONST_RECORD_ALIGNR   (   R   R!   R   t   Dt   Et   ZR   (    (    sD   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_constmem.pyt   cuconstRecAlignM   s    c         C  sM   t  j j t ƒ } t  j j t ƒ } t  j d ƒ } | | | | |  | <d  S(   Ni   (   R   R   R   t   CONST3BYTESR   R   (   R
   R   R	   R   (    (    sD   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_constmem.pyt   cuconstAlignW   s    t   TestCudaConstantMemoryc           B  sP   e  Z d  „  Z d „  Z d „  Z d „  Z d „  Z d „  Z d „  Z d „  Z	 RS(   c         C  su   t  j d ƒ t ƒ } t j t ƒ } | d | ƒ |  j t j | t d k ƒ ƒ t sq |  j	 d | j
 d ƒ n  d  S(   Ns   void(float64[:])i   i   i   s   ld.const.f64s'   as we're adding to it, load as a double(   i   i   (   R   t   jitR   t   npt
   zeros_likeR   t
   assertTruet   allR   t   assertInt   ptx(   t   selft   jcuconstR   (    (    sD   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_constmem.pyt   test_const_array_   s     c         C  s^   t  j d ƒ t ƒ } t j d d d d t ƒ} | d | ƒ |  j t j | d k ƒ ƒ d  S(   Ns   void(float64[:])i   t
   fill_valueiÿÿÿÿR   i    (   i   i   (   R   R+   R   R,   t   fullt   intR.   R/   (   R2   t   jcuconstEmptyR   (    (    sD   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_constmem.pyt   test_const_emptyk   s    c         C  si   t  j d ƒ t ƒ } t j d d t j d t ƒ} | d | ƒ |  j t j | t	 t
 d  k ƒ ƒ d  S(   Ns   void(float64[:])i   R5   R   i   (   i   i   (   R   R+   R)   R,   R6   t   nant   floatR.   R/   R(   R   (   R2   t   jcuconstAlignR   (    (    sD   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_constmem.pyt   test_const_alignq   s    c         C  s}   t  j d ƒ t ƒ } t j t d d ƒ} | d d	 f | ƒ |  j t j | t k ƒ ƒ t sy |  j	 d | j
 d ƒ n  d  S(
   Ns   void(int32[:,:])t   orderR   i   i   s   ld.const.u32s   load the ints as ints(   i   i   (   i   i   (   R   R+   R   R,   R-   R   R.   R/   R   R0   R1   (   R2   t
   jcuconst2dR   (    (    sD   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_constmem.pyt   test_const_array_2dw   s    c         C  s}   t  j d ƒ t ƒ } t j t d d ƒ} | d d f | ƒ |  j t j | t k ƒ ƒ t sy |  j	 d | j
 d ƒ n  d  S(	   Ns   void(complex64[:,:,:])R>   t   Fi   i   s   ld.const.v2.u32s*   load the two halves of the complex as u32s(   i   i   i   (   R   R+   R   R,   R-   R   R.   R/   R   R0   R1   (   R2   t
   jcuconst3dR   (    (    sD   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_constmem.pyt   test_const_array_3dƒ   s    c         C  s^   t  j d ƒ t ƒ } t j d d d d t ƒ} | d | ƒ |  j t j | d k ƒ ƒ d  S(   Ns   void(float64[:])i   R5   iÿÿÿÿR   i    (   i   i   (   R   R+   R   R,   R6   R7   R.   R/   (   R2   t   jcuconstRecEmptyR   (    (    sD   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_constmem.pyt   test_const_record_empty   s    c           sÂ   t  j d d t ƒ} t  j d d t ƒ} t j t ƒ j | | ƒ ‰  t s t	 ‡  f d †  d d g Dƒ ƒ s t
 d ƒ ‚ q n  ˆ  d
 | | ƒ t  j j | t d ƒ t  j j | t d	 ƒ d  S(   Ni   R   c         3  s   |  ] } | ˆ  j  k Vq d  S(   N(   R1   (   t   .0t   c(   R3   (    sD   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_constmem.pys	   <genexpr>›   s    s   ld.const.v2.u64s   ld.const.u32sL   the compiler should realise it doesn't need to interpret the bytes as float!i   R   R   (   i   i   (   R,   t   zerosR;   R7   R   R+   R"   t
   specializeR   t   anyt   AssertionErrort   testingt   assert_allcloseR    (   R2   R   R!   (    (   R3   sD   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_constmem.pyt   test_const_record•   s    c         C  sx  t  j d d t  j ƒ} t  j d d t  j ƒ} t  j d d t  j ƒ} t  j d d t  j ƒ} t  j d d t  j ƒ} t j t ƒ j | | | | | ƒ } t sç |  j d | j	 d ƒ |  j d | j	 d ƒ |  j d | j	 d ƒ n  | d | | | | | ƒ t  j
 j | t d
 ƒ t  j
 j | t d ƒ t  j
 j | t d ƒ t  j
 j | t d ƒ t  j
 j | t d ƒ d  S(   Ni   R   s   ld.const.v4.u8s&   load the first three bytes as a vectors   ld.const.u32s   load the uint32 nativelys   ld.const.u8s   load the last byte by itselfi   R   R	   R   R   R
   (   i   i   (   R,   RH   t   float64R   R+   R'   RI   R   R0   R1   RL   RM   R#   (   R2   R   R!   R   R$   R%   R3   (    (    sD   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_constmem.pyt   test_const_record_align¯   s2    $
(
   t   __name__t
   __module__R4   R9   R=   R@   RC   RE   RN   RP   (    (    (    sD   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_constmem.pyR*   ^   s   							t   __main__i   i}   (   g      ð?i   (   g      @i   (   i   i   i   Iï¾­Þ    i   (   i   i   i   I­Þï¾    i
   (,   t
   __future__R    t   numpyR,   t   numbaR   t   numba.cuda.testingR   R   t   numba.configR   t   arrayR   t   arangeRO   R   t   asfortranarrayt   int32t   reshapeR   t	   complex64R   t   uint8R(   R;   R7   R   R    R   t   uint32t   TrueR#   R   R   R   R   R   R"   R'   R)   t   TestCaseR*   RQ   t   main(    (    (    sD   lib/python2.7/site-packages/numba/cuda/tests/cudapy/test_constmem.pyt   <module>   sJ   '(								
	q