B
     \                 @   sd   d dl mZmZ d dlmZ d dlmZ d dlmZ dZ	G dd dej
ZdZG d	d
 d
ejZdS )    )print_functionabsolute_import)roc)deviceufunc)dispatcha  
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__

c               @   s@   e Zd Zdd Zdd Zdd Zdd Zd	d
 Zedd Z	dS )HsaVectorizec             C   s"   t j|dd| j}||jjjfS )NT)device)r   jitpyfuncZcresZ	signatureZreturn_type)selfsigZhsadevfn r   4lib/python3.7/site-packages/numba/roc/vectorizers.py_compile_coreL   s    zHsaVectorize._compile_corec             C   s   | j j}|t|d |S )N)__hsa____core__)r
   __globals__updater   )r   corefnZglblr   r   r   _get_globalsP   s    
zHsaVectorize._get_globalsc             C   s   t ||S )N)r   r	   )r   fnobjr   r   r   r   _compile_kernelV   s    zHsaVectorize._compile_kernelc          	      sf   dd t t|jD   fdd}t|d |dd|dd|d	d|d
dd}|jf |}|S )Nc             S   s   g | ]}d | qS )za%dr   ).0ir   r   r   
<listcomp>Z   s    z3HsaVectorize._get_kernel_source.<locals>.<listcomp>c                s6   d  fddD }tdk r.d|S |S d S )Nz, c             3   s   | ]}d | f V  qdS )z%s[__tid%d__]Nr   )r   r   )nr   r   	<genexpr>]   s    zIHsaVectorize._get_kernel_source.<locals>.make_argitems.<locals>.<genexpr>   z({0},))joinlenformat)r   out)args)r   r   make_argitems\   s    
z6HsaVectorize._get_kernel_source.<locals>.make_argitemsz, r   )r      r      )namer"   Z
argitems_0Z
argitems_1Z
argitems_2Z
argitems_3)ranger   r"   dictr   r    )r   templater   funcnamer#   Zfmtssrcr   )r"   r   _get_kernel_sourceY   s    
zHsaVectorize._get_kernel_sourcec             C   s   t | jS )N)r   ZHsaUFuncDispatcher	kernelmap)r   r   r   r   build_ufunco   s    zHsaVectorize.build_ufuncc             C   s   t S )N)vectorizer_stager_source)r   r   r   r   _kernel_templater   s    zHsaVectorize._kernel_templateN)
__name__
__module____qualname__r   r   r   r,   r.   propertyr0   r   r   r   r   r   K   s   r   z
def __gufunc_{name}({args}):
    __tid__ = __hsa__.get_global_id(0)
    if __tid__ < {checkedarg}:
        __core__({argitems})
c               @   s0   e Zd Zdd Zdd Zedd Zdd Zd	S )
HsaGUFuncVectorizec             C   s    t | j| j}tj| j|dS )N)r-   engine)r   ZGUFuncEngineZinputsigZ	outputsigr   ZHSAGenerializedUFuncr-   )r   r6   r   r   r   r.      s    zHsaGUFuncVectorize.build_ufuncc             C   s   t ||S )N)r   r	   )r   r   r   r   r   r   r      s    z"HsaGUFuncVectorize._compile_kernelc             C   s   t S )N)_gufunc_stager_source)r   r   r   r   r0      s    z#HsaGUFuncVectorize._kernel_templatec             C   s4   t j|dd| j}| jj }|t |d |S )NT)r   )r   r   )r   r	   r
   Zpy_funcr   copyr   )r   r   r   Zglblsr   r   r   r      s
    
zHsaGUFuncVectorize._get_globalsN)r1   r2   r3   r.   r   r4   r0   r   r   r   r   r   r5      s   r5   N)Z
__future__r   r   Znumbar   Znumba.npyufuncr   Z	numba.rocr   r/   ZDeviceVectorizer   r7   ZDeviceGUFuncVectorizer5   r   r   r   r   <module>   s   B4