σ
Ϋ³Y]c           @@ s   d  d l  m Z m Z d  d l m Z d  d l m Z d  d l m Z d Z	 d e j
 f d     YZ d Z d	 e j f d
     YZ d S(   i    (   t   print_functiont   absolute_import(   t   roc(   t   deviceufunc(   t   dispatchsέ  
def __vectorized_{name}({args}, __out__):

    __tid__ = __hsa__.get_local_id(0)
    __blksz__ = __hsa__.get_local_size(0)
    __blkid__ = __hsa__.get_group_id(0)

    __tid0__ = __tid__ + __blksz__ * (4 * __blkid__)
    __tid1__ = __tid__ + __blksz__ * (4 * __blkid__ + 1)
    __tid2__ = __tid__ + __blksz__ * (4 * __blkid__ + 2)
    __tid3__ = __tid__ + __blksz__ * (4 * __blkid__ + 3)

    __ilp0__ = __tid0__ < __out__.shape[0]
    if not __ilp0__:
        # Early escape
        return
    __ilp1__ = __tid1__ < __out__.shape[0]
    __ilp2__ = __tid2__ < __out__.shape[0]
    __ilp3__ = __tid3__ < __out__.shape[0]

    if __ilp3__:
        __args0__ = {argitems_0}
        __args1__ = {argitems_1}
        __args2__ = {argitems_2}
        __args3__ = {argitems_3}

        __r0__ = __core__(*__args0__)
        __r1__ = __core__(*__args1__)
        __r2__ = __core__(*__args2__)
        __r3__ = __core__(*__args3__)

        __out__[__tid0__] = __r0__
        __out__[__tid1__] = __r1__
        __out__[__tid2__] = __r2__
        __out__[__tid3__] = __r3__

    elif __ilp2__:
        __args0__ = {argitems_0}
        __args1__ = {argitems_1}
        __args2__ = {argitems_2}

        __r0__ = __core__(*__args0__)
        __r1__ = __core__(*__args1__)
        __r2__ = __core__(*__args2__)

        __out__[__tid0__] = __r0__
        __out__[__tid1__] = __r1__
        __out__[__tid2__] = __r2__

    elif __ilp1__:
        __args0__ = {argitems_0}
        __args1__ = {argitems_1}

        __r0__ = __core__(*__args0__)
        __r1__ = __core__(*__args1__)

        __out__[__tid0__] = __r0__
        __out__[__tid1__] = __r1__

    else:
        __args0__ = {argitems_0}
        __r0__ = __core__(*__args0__)
        __out__[__tid0__] = __r0__

t   HsaVectorizec           B@ sD   e  Z d    Z d   Z d   Z d   Z d   Z e d    Z RS(   c         C@ s1   t  j | d t |  j  } | | j j j f S(   Nt   device(   R   t   jitt   Truet   pyfunct   crest	   signaturet   return_type(   t   selft   sigt   hsadevfn(    (    s4   lib/python2.7/site-packages/numba/roc/vectorizers.pyt   _compile_coreL   s    c         C@ s+   |  j  j } | j i t d 6| d 6 | S(   Nt   __hsa__t   __core__(   R	   t   __globals__t   updateR   (   R   t   corefnt   glbl(    (    s4   lib/python2.7/site-packages/numba/roc/vectorizers.pyt   _get_globalsP   s    c         C@ s   t  j |  |  S(   N(   R   R   (   R   t   fnobjR   (    (    s4   lib/python2.7/site-packages/numba/roc/vectorizers.pyt   _compile_kernelV   s    c         @ s¨   g  t  t | j   D] } d | ^ q     f d   } t d | d d j    d | d d  d	 | d d
  d | d d  d | d d   } | j |   } | S(   Ns   a%dc         @ sF   d j    f d    D  } t   d k  r> d j |  S| Sd  S(   Ns   , c         3@ s   |  ] } d  |   f Vq d S(   s   %s[__tid%d__]N(    (   t   .0t   i(   t   n(    s4   lib/python2.7/site-packages/numba/roc/vectorizers.pys	   <genexpr>]   s    i   s   ({0},)(   t   joint   lent   format(   R   t   out(   t   args(   R   s4   lib/python2.7/site-packages/numba/roc/vectorizers.pyt   make_argitems\   s    t   nameR!   s   , t
   argitems_0R   i    t
   argitems_1i   t
   argitems_2i   t
   argitems_3i   (   t   rangeR   R!   t   dictR   R   (   R   t   templateR   t   funcnameR   R"   t   fmtst   src(    (   R!   s4   lib/python2.7/site-packages/numba/roc/vectorizers.pyt   _get_kernel_sourceY   s    ,
c         C@ s   t  j |  j  S(   N(   R   t   HsaUFuncDispatchert	   kernelmap(   R   (    (    s4   lib/python2.7/site-packages/numba/roc/vectorizers.pyt   build_ufunco   s    c         C@ s   t  S(   N(   t   vectorizer_stager_source(   R   (    (    s4   lib/python2.7/site-packages/numba/roc/vectorizers.pyt   _kernel_templater   s    (	   t   __name__t
   __module__R   R   R   R.   R1   t   propertyR3   (    (    (    s4   lib/python2.7/site-packages/numba/roc/vectorizers.pyR   K   s   					s   
def __gufunc_{name}({args}):
    __tid__ = __hsa__.get_global_id(0)
    if __tid__ < {checkedarg}:
        __core__({argitems})
t   HsaGUFuncVectorizec           B@ s2   e  Z d    Z d   Z e d    Z d   Z RS(   c         C@ s1   t  j |  j |  j  } t j d |  j d |  S(   NR0   t   engine(   R   t   GUFuncEnginet   inputsigt	   outputsigR   t   HSAGenerializedUFuncR0   (   R   R8   (    (    s4   lib/python2.7/site-packages/numba/roc/vectorizers.pyR1      s    c         C@ s   t  j |  |  S(   N(   R   R   (   R   R   R   (    (    s4   lib/python2.7/site-packages/numba/roc/vectorizers.pyR      s    c         C@ s   t  S(   N(   t   _gufunc_stager_source(   R   (    (    s4   lib/python2.7/site-packages/numba/roc/vectorizers.pyR3      s    c         C@ sO   t  j | d t |  j  } |  j j j   } | j i t  d 6| d 6 | S(   NR   R   R   (   R   R   R   R	   t   py_funcR   t   copyR   (   R   R   R   t   glbls(    (    s4   lib/python2.7/site-packages/numba/roc/vectorizers.pyR      s
    (   R4   R5   R1   R   R6   R3   R   (    (    (    s4   lib/python2.7/site-packages/numba/roc/vectorizers.pyR7      s   		N(   t
   __future__R    R   t   numbaR   t   numba.npyufuncR   t	   numba.rocR   R2   t   DeviceVectorizeR   R=   t   DeviceGUFuncVectorizeR7   (    (    (    s4   lib/python2.7/site-packages/numba/roc/vectorizers.pyt   <module>   s   B4