ó
\K]c           @@ s‡  d  d l  m Z m Z d  d l Z d  d l m Z m Z m Z m Z m	 Z	 d  d l j
 j Z d  d l j Z d  d l m Z m Z 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  d	 l m Z d  d
 l m  Z  d d l! m" Z" d d l# m$ Z$ m% Z% d d l& m' Z' d e j f d „  ƒ  YZ( e j) d e j* ƒ Z+ d e f d „  ƒ  YZ, d e f d „  ƒ  YZ- d S(   i    (   t   print_functiont   absolute_importN(   t   Typet   Buildert   LINKAGE_INTERNALt   Constantt   ICMP_EQ(   t   typingt   typest   cgutilst	   debuginfot
   dispatcher(   t   cached_property(   t   BaseContext(   t   MinimalCallConv(   t	   cmathimpl(   t	   cmathdecl(   t   itanium_mangleri   (   t   nvvm(   t   codegent	   nvvmutils(   t	   jitdevicet   CUDATypingContextc           B@ s   e  Z d  „  Z d „  Z RS(   c         C@ sJ   d d l  m } m } |  j | j ƒ |  j | j ƒ |  j t j ƒ d  S(   Ni   (   t   cudadeclt   cudamath(   t    R   R   t   install_registryt   registryR   (   t   selfR   R   (    (    s0   lib/python2.7/site-packages/numba/cuda/target.pyt   load_additional_registries   s    c         C@ s‘   t  | t j ƒ r{ y | j } Wq{ t k
 rw | j sG t d ƒ ‚ n  t | d | j j	 d ƒ ƒ} | | _ | } q{ Xn  t
 t |  ƒ j | ƒ S(   Ns<   using cpu function on device but its compilation is disabledt   debug(   t
   isinstanceR   t
   Dispatchert!   _CUDATypingContext__cudajitdevicet   AttributeErrort   _can_compilet
   ValueErrorR   t   targetoptionst   gett   superR   t   resolve_value_type(   R   t   valt   jd(    (    s0   lib/python2.7/site-packages/numba/cuda/target.pyR(   !   s    		(   t   __name__t
   __module__R   R(   (    (    (    s0   lib/python2.7/site-packages/numba/cuda/target.pyR      s   	s	   [^a-z0-9]t   CUDATargetContextc           B@ s§   e  Z e Z e Z e j Z d  „  Z d „  Z	 d „  Z
 d „  Z e d „  ƒ Z e d „  ƒ Z d „  Z d „  Z d „  Z d	 „  Z d
 „  Z d „  Z d „  Z d „  Z RS(   c         C@ s   |  j  j | ƒ S(   N(   t   _internal_codegent   _create_empty_module(   R   t   name(    (    s0   lib/python2.7/site-packages/numba/cuda/target.pyt   create_module@   s    c         C@ s+   t  j d ƒ |  _ t j t j ƒ |  _ d  S(   Ns   numba.cuda.jit(   R   t   JITCUDACodegenR.   t   llt   create_target_dataR   t   default_data_layoutt   _target_data(   R   (    (    s0   lib/python2.7/site-packages/numba/cuda/target.pyt   initC   s    c         C@ s`   d d l  m } m } m } |  j | j ƒ |  j | j ƒ |  j | j ƒ |  j t j ƒ d  S(   Ni   (   t   cudaimplt	   printimplt	   libdevice(   R   R8   R9   R:   R   R   R   (   R   R8   R9   R:   (    (    s0   lib/python2.7/site-packages/numba/cuda/target.pyR   G   s
    c         C@ s   |  j  S(   N(   R.   (   R   (    (    s0   lib/python2.7/site-packages/numba/cuda/target.pyR   N   s    c         C@ s   |  j  S(   N(   R6   (   R   (    (    s0   lib/python2.7/site-packages/numba/cuda/target.pyt   target_dataQ   s    c         C@ s
   t  |  ƒ S(   N(   t   CUDACallConv(   R   (    (    s0   lib/python2.7/site-packages/numba/cuda/target.pyt	   call_convU   s    c         C@ s   t  j | | ƒ S(   N(   R   t   mangle(   R   R0   t   argtypes(    (    s0   lib/python2.7/site-packages/numba/cuda/target.pyt   manglerY   s    c         C@ sW   |  j  ƒ  j d ƒ } | j | ƒ |  j | | | d | ƒ} t j | j ƒ | | f S(   sY  
        Adapt a code library ``codelib`` with the numba compiled CUDA kernel
        with name ``fname`` and arguments ``argtypes`` for NVVM.
        A new library is created with a wrapper function that can be used as
        the kernel entry point for the given kernel.

        Returns the new code library and the wrapper function.
        R   R   (   R   t   create_libraryt   add_linking_libraryt   generate_kernel_wrapperR   t   fix_data_layoutt   _final_module(   R   t   codelibt   fnameR?   R   t   libraryt   wrapper(    (    s0   lib/python2.7/site-packages/numba/cuda/target.pyt   prepare_cuda_kernel\   s    		c         @ s.  |  j  | ƒ } t | j ƒ } t j t j ƒ  | ƒ } |  j d ƒ ‰ t j t j ƒ  |  j j	 t
 j ƒ g | ƒ } ˆ j | d | ƒ}	 t j |	 j d d ƒ}
 ˆ j | d |
 ƒ‰  t ˆ  j d ƒ ƒ } ‡  ‡ f d †  } | d ƒ } g  } g  } x< d D]4 } | j | d	 | ƒ ƒ | j | d
 | ƒ ƒ qû W| j | ˆ  j ƒ } |  j j | |	 t
 j | | ƒ \ } } | rêt j | | j ƒ  | j ƒ  Wd QX| j | j | j ƒ ƒ /t j | j j  ƒ } t! j j | j | j | j | j g ƒ } ˆ j | d d ƒ} | j" | | | | j# g ƒ } | j$ t% | | ƒ } t& j' | ƒ } | j | ƒ ƒ x< t( d | ƒ D]+ \ } } | j) | ƒ } | j* | | ƒ qmWx< t( d | ƒ D]+ \ } } | j+ | ƒ } | j* | | ƒ q¬WWd QXWd QXn  | j ƒ  t, j- ˆ  ƒ | j. ˆ ƒ | j/ ƒ  | j0 ˆ  j ƒ ‰  ˆ  S(   sÎ   
        Generate the kernel wrapper in the given ``library``.
        The function being wrapped have the name ``fname`` and argument types
        ``argtypes``.  The wrapper function is returned.
        s   cuda.kernel.wrapperR0   t   nst   cudapyR   c         @ s>   ˆ j  t j ƒ  d ˆ  j |  ƒ} t j | j j ƒ | _ | S(   NR0   (	   t   add_global_variableR   t   intR0   R   t   nullt   typet   pointeet   initializer(   t   postfixt   gv(   t   wrapfnt   wrapper_module(    s0   lib/python2.7/site-packages/numba/cuda/target.pyt   define_error_gv   s    t   __errcode__t   xyzs	   __tid%s__s   __ctaid%s__Nt   ___numba_cas_hack(1   t   get_arg_packert   listt   argument_typesR   t   functiont   voidR1   RN   R=   t   get_return_typeR   t   pyobjectt   add_functionR   t   prepend_namespaceR0   R   t   append_basic_blockt   appendt   from_argumentst   argst   call_functionR	   t	   if_likelyt   is_okt   ret_voidt   if_thent   not_t   is_python_excR   RO   RP   RQ   t   lct   callt   codet   icmpR   R   t   SRegBuildert   zipt   tidt   storet   ctaidR   t   set_cuda_kernelt   add_ir_modulet   finalizet   get_function(   R   RH   RG   R?   R   t   arginfot   argtyst   wrapfntyt   fntyt   funct   prefixedt   builderRW   t   gv_exct   gv_tidt   gv_ctaidt   it   callargst   statust   _t   oldt   casfntyt   casfnt   xchgt   changedt   sregt   dimt   ptrR)   (    (   RU   RV   s0   lib/python2.7/site-packages/numba/cuda/target.pyRC   l   sZ    		#

c         C@ s"   |  j  | ƒ |  | ƒ } | j ƒ  S(   sn   
        Return dummy value.

        XXX: We should be able to move cuda.const.array_like into here.
        (   t
   make_arrayt	   _getvalue(   R   R‚   t   typt   aryt   a(    (    s0   lib/python2.7/site-packages/numba/cuda/target.pyt   make_constant_array¸   s    c         C@ sµ   t  j | ƒ } d j d t j | ƒ g ƒ } | j j | ƒ } | d k rŠ | j | j	 d | d t
 j ƒ} t | _ t | _ | | _ n  | j	 j j } t  j | | j t
 j ƒ ƒ S(   si   
        Unlike the parent version.  This returns a a pointer in the constant
        addrspace.
        t   $t   __conststring__R0   t	   addrspaceN(   R   t   stringzt   joinR   t   mangle_identifiert   globalsR&   t   NoneRM   RP   R   t   ADDRSPACE_CONSTANTR   t   linkaget   Truet   global_constantRR   RQ   t   elementt   bitcastt
   as_pointer(   R   t   modt   stringt   textR0   RT   t   charty(    (    s0   lib/python2.7/site-packages/numba/cuda/target.pyt   insert_const_stringÂ   s    				c         C@ s1   | j  } |  j | | ƒ } |  j | | t j ƒ S(   s«   
        Insert a constant string in the constant addresspace and return a
        generic i8 pointer to the data.

        This function attempts to deduplicate.
        (   t   moduleR«   t   insert_addrspace_convR   R    (   R   R‚   R¨   t   lmodRT   (    (    s0   lib/python2.7/site-packages/numba/cuda/target.pyt   insert_string_const_addrspaceÙ   s    	c         C@ s=   | j  } | j j } t j | | | ƒ } | j | | g ƒ S(   sI   
        Perform addrspace conversion according to the NVVM spec
        (   R¬   RP   RQ   R   R­   Rp   (   R   R‚   R‘   Rš   R®   t	   base_typet   conv(    (    s0   lib/python2.7/site-packages/numba/cuda/target.pyR­   å   s    	c         C@ s   d S(   s   Run O1 function passes
        N(    (   R   R€   (    (    s0   lib/python2.7/site-packages/numba/cuda/target.pyt   optimize_functionî   s    (   R+   R,   R¢   t   implement_powi_as_math_callt   strict_alignmentR
   t   NvvmDIBuildert	   DIBuilderR1   R7   R   R   t   propertyR;   R   R=   R@   RJ   RC   R—   R«   R¯   R­   R²   (    (    (    s0   lib/python2.7/site-packages/numba/cuda/target.pyR-   :   s"   								L	
				R<   c           B@ s   e  Z RS(    (   R+   R,   (    (    (    s0   lib/python2.7/site-packages/numba/cuda/target.pyR<   ü   s   (.   t
   __future__R    R   t   ret   llvmlite.llvmpy.coreR   R   R   R   R   t   llvmpyt   coreRo   t   llvmlite.bindingt   bindingR3   t   numbaR   R   R	   R
   R   t   numba.utilsR   t   numba.targets.baseR   t   numba.targets.callconvR   t   numba.targetsR   t   numba.typingR   R   t   cudadrvR   R   R   R   t
   decoratorsR   R   t   compilet   It   VALID_CHARSR-   R<   (    (    (    s0   lib/python2.7/site-packages/numba/cuda/target.pyt   <module>   s$   ((Â