B
     \7                 @   s8  d dl mZmZmZ d dlmZ d dlmZ d dlm	Z	m
Z
mZmZmZmZ d dlmZ e ZejZejZejZee G dd deZG dd	 d	eZG d
d deZG dd deZG dd deZG dd deZG dd deZG dd deZG dd deZG dd deZ G dd deZ!G dd deZ"G dd deZ#G d d! d!eZ$G d"d# d#eZ%G d$d% d%eZ&G d&d' d'eZ'G d(d) d)eZ(G d*d+ d+eZ)eG d,d- d-e
Z*eG d.d/ d/e
Z+eG d0d1 d1e
Z,eG d2d3 d3e
Z-eG d4d5 d5e
Z.eG d6d7 d7e
Z/eG d8d9 d9e
Z0eG d:d; d;e
Z1eG d<d= d=e
Z2eG d>d? d?e
Z3eG d@dA dAe
Z4eG dBdC dCe
Z5eG dDdE dEe
Z6eG dFdG dGe
Z7eG dHdI dIe
Z8eG dJdK dKe
Z9eG dLdM dMe
Z:eG dNdO dOeZ;eG dPdQ dQeZ<G dRdS dSeZ=eG dTdU dUe=Z>eG dVdW dWe=Z?eG dXdY dYeZ@eG dZd[ d[e	ZAeG d\d] d]e	ZBeG d^d_ d_e	ZCeG d`da dae	ZDeG dbdc dce	ZEeG ddde dee	ZFeG dfdg dge	ZGeG dhdi die	ZHeG djdk dke	ZIeeeJe dlS )m    )print_functiondivisionabsolute_import)types)register_number_classes)AttributeTemplateConcreteTemplateAbstractTemplateMacroTemplate	signatureRegistry)cudac               @   s   e Zd ZejZdS )	Cuda_gridN)__name__
__module____qualname__r   Zgridkey r   r   2lib/python3.7/site-packages/numba/cuda/cudadecl.pyr      s   r   c               @   s   e Zd ZejZdS )Cuda_gridsizeN)r   r   r   r   Zgridsizer   r   r   r   r   r      s   r   c               @   s   e Zd ZejjZdS )Cuda_threadIdx_xN)r   r   r   r   	threadIdxxr   r   r   r   r   r      s   r   c               @   s   e Zd ZejjZdS )Cuda_threadIdx_yN)r   r   r   r   r   yr   r   r   r   r   r      s   r   c               @   s   e Zd ZejjZdS )Cuda_threadIdx_zN)r   r   r   r   r   zr   r   r   r   r   r   "   s   r   c               @   s   e Zd ZejjZdS )Cuda_blockIdx_xN)r   r   r   r   blockIdxr   r   r   r   r   r   r   &   s   r   c               @   s   e Zd ZejjZdS )Cuda_blockIdx_yN)r   r   r   r   r   r   r   r   r   r   r   r   *   s   r   c               @   s   e Zd ZejjZdS )Cuda_blockIdx_zN)r   r   r   r   r   r   r   r   r   r   r   r    .   s   r    c               @   s   e Zd ZejjZdS )Cuda_blockDim_xN)r   r   r   r   blockDimr   r   r   r   r   r   r!   2   s   r!   c               @   s   e Zd ZejjZdS )Cuda_blockDim_yN)r   r   r   r   r"   r   r   r   r   r   r   r#   6   s   r#   c               @   s   e Zd ZejjZdS )Cuda_blockDim_zN)r   r   r   r   r"   r   r   r   r   r   r   r$   :   s   r$   c               @   s   e Zd ZejjZdS )Cuda_gridDim_xN)r   r   r   r   gridDimr   r   r   r   r   r   r%   >   s   r%   c               @   s   e Zd ZejjZdS )Cuda_gridDim_yN)r   r   r   r   r&   r   r   r   r   r   r   r'   B   s   r'   c               @   s   e Zd ZejjZdS )Cuda_gridDim_zN)r   r   r   r   r&   r   r   r   r   r   r   r(   F   s   r(   c               @   s   e Zd ZejZdS )Cuda_warpsizeN)r   r   r   r   Zwarpsizer   r   r   r   r   r)   J   s   r)   c               @   s   e Zd ZejZdS )Cuda_laneidN)r   r   r   r   Zlaneidr   r   r   r   r   r*   N   s   r*   c               @   s   e Zd ZejjZdS )Cuda_shared_arrayN)r   r   r   r   sharedarrayr   r   r   r   r   r+   R   s   r+   c               @   s   e Zd ZejjZdS )Cuda_local_arrayN)r   r   r   r   localr-   r   r   r   r   r   r.   V   s   r.   c               @   s   e Zd ZejjZdS )Cuda_const_arraylikeN)r   r   r   r   constZ
array_liker   r   r   r   r   r0   Z   s   r0   c               @   s   e Zd ZejZeejgZ	dS )Cuda_syncthreadsN)
r   r   r   r   Zsyncthreadsr   r   r   nonecasesr   r   r   r   r2   ^   s   r2   c               @   s"   e Zd ZejZeejejgZ	dS )Cuda_syncthreads_countN)
r   r   r   r   Zsyncthreads_countr   r   r   i4r4   r   r   r   r   r5   d   s   r5   c               @   s"   e Zd ZejZeejejgZ	dS )Cuda_syncthreads_andN)
r   r   r   r   Zsyncthreads_andr   r   r   r6   r4   r   r   r   r   r7   j   s   r7   c               @   s"   e Zd ZejZeejejgZ	dS )Cuda_syncthreads_orN)
r   r   r   r   Zsyncthreads_orr   r   r   r6   r4   r   r   r   r   r8   p   s   r8   c               @   s   e Zd ZejZeejgZ	dS )Cuda_threadfence_deviceN)
r   r   r   r   Zthreadfencer   r   r   r3   r4   r   r   r   r   r9   v   s   r9   c               @   s   e Zd ZejZeejgZ	dS )Cuda_threadfence_blockN)
r   r   r   r   Zthreadfence_blockr   r   r   r3   r4   r   r   r   r   r:   |   s   r:   c               @   s   e Zd ZejZeejgZ	dS )Cuda_threadfence_systemN)
r   r   r   r   Zthreadfence_systemr   r   r   r3   r4   r   r   r   r   r;      s   r;   c               @   s"   e Zd ZejZeejej	gZ
dS )Cuda_syncwarpN)r   r   r   r   Zsyncwarpr   r   r   r3   r6   r4   r   r   r   r   r<      s   r<   c            
   @   s   e Zd ZejZeeej	ej
fej	ej	ej	ej	ej	eeejej
fej	ej	ejej	ej	eeejej
fej	ej	ejej	ej	eeejej
fej	ej	ejej	ej	gZdS )Cuda_shfl_sync_intrinsicN)r   r   r   r   Zshfl_sync_intrinsicr   r   r   Tupler6   b1i8f4f8r4   r   r   r   r   r=      s
   (((r=   c               @   s6   e Zd ZejZeeej	ej
fej	ej	ej
gZdS )Cuda_vote_sync_intrinsicN)r   r   r   r   Zvote_sync_intrinsicr   r   r   r>   r6   r?   r4   r   r   r   r   rC      s   rC   c               @   sV   e Zd ZejZeejejejeejejej	eejejej
eejejejgZdS )Cuda_match_any_syncN)r   r   r   r   Zmatch_any_syncr   r   r   r6   r@   rA   rB   r4   r   r   r   r   rD      s
   rD   c               @   s   e Zd ZejZeeej	ej
fej	ej	eeej	ej
fej	ejeeej	ej
fej	ejeeej	ej
fej	ejgZdS )Cuda_match_all_syncN)r   r   r   r   Zmatch_all_syncr   r   r   r>   r6   r?   r@   rA   rB   r4   r   r   r   r   rE      s
   rE   c            
   @   sz   e Zd ZdZejZeej	ej	eej
ej
eejejeejejeejejeejejeejejeejejgZdS )	Cuda_popcz
    Supported types from `llvm.popc`
    [here](http://docs.nvidia.com/cuda/nvvm-ir-spec/index.html#bit-manipulations-intrinics)
    N)r   r   r   __doc__r   Zpopcr   r   r   int8int16int32int64uint8uint16uint32uint64r4   r   r   r   r   rF      s   rF   c               @   sB   e Zd ZdZejZeej	ej	ej	ej	eej
ej
ej
ej
gZdS )Cuda_fmaz
    Supported types from `llvm.fma`
    [here](https://docs.nvidia.com/cuda/nvvm-ir-spec/index.html#standard-c-library-intrinics)
    N)r   r   r   rG   r   Zfmar   r   r   float32float64r4   r   r   r   r   rP      s   rP   c               @   s.   e Zd ZejZeejejeej	ej	gZ
dS )	Cuda_brevN)r   r   r   r   Zbrevr   r   r   rN   rO   r4   r   r   r   r   rS      s   rS   c            
   @   sz   e Zd ZdZejZeej	ej	eej
ej
eejejeejejeejejeejejeejejeejejgZdS )Cuda_clzz
    Supported types from `llvm.ctlz`
    [here](http://docs.nvidia.com/cuda/nvvm-ir-spec/index.html#bit-manipulations-intrinics)
    N)r   r   r   rG   r   Zclzr   r   r   rH   rI   rJ   rK   rL   rM   rN   rO   r4   r   r   r   r   rT      s   rT   c            
   @   sz   e Zd ZdZejZeej	ej	eej
ej
eejejeejejeejejeejejeejejeejejgZdS )Cuda_ffsz
    Supported types from `llvm.cttz`
    [here](http://docs.nvidia.com/cuda/nvvm-ir-spec/index.html#bit-manipulations-intrinics)
    N)r   r   r   rG   r   Zffsr   r   r   rH   rI   rJ   rK   rL   rM   rN   rO   r4   r   r   r   r   rU      s   rU   c               @   s   e Zd ZejZdd ZdS )	Cuda_selpc             C   sX   |rt |\}}}tjtjtjtjtjtjtjtj	f}||ksF||krJd S t
||||S )N)AssertionErrorr   rR   rQ   rI   rM   rJ   rN   rK   rO   r   )selfargskwsZtestabsupported_typesr   r   r   generic  s    
zCuda_selp.genericN)r   r   r   r   Zselpr   r^   r   r   r   r   rV     s   rV   c               @   s   e Zd ZejjZdd ZdS )Cuda_atomic_addc             C   sP   |rt |\}}}|jdkr0t|j|tj|jS |jdkrLt|j|||jS d S )N   )rW   ndimr   dtyper   intp)rX   rY   rZ   aryidxvalr   r   r   r^     s    


zCuda_atomic_add.genericN)r   r   r   r   atomicaddr   r^   r   r   r   r   r_     s   r_   c               @   s   e Zd Zdd ZdS )Cuda_atomic_maxminc             C   sz   |rt |\}}}tjtjtjtjtjtjf}|j|kr<d S |j	dkrZt
|j|tj|jS |j	dkrvt
|j|||jS d S )Nr`   )rW   r   rR   rQ   rJ   rN   rK   rO   rb   ra   r   rc   )rX   rY   rZ   rd   re   rf   r]   r   r   r   r^   (  s    



zCuda_atomic_maxmin.genericN)r   r   r   r^   r   r   r   r   ri   '  s   ri   c               @   s   e Zd ZejjZdS )Cuda_atomic_maxN)r   r   r   r   rg   maxr   r   r   r   r   rj   :  s   rj   c               @   s   e Zd ZejjZdS )Cuda_atomic_minN)r   r   r   r   rg   minr   r   r   r   r   rl   ?  s   rl   c               @   s   e Zd ZejjZdd ZdS )Cuda_atomic_compare_and_swapc             C   s>   |rt |\}}}|j}|tjkr:|jdkr:t||||S d S )Nr`   )rW   rb   r   rJ   ra   r   )rX   rY   rZ   rd   oldrf   Zdtyr   r   r   r^   H  s
    
z$Cuda_atomic_compare_and_swap.genericN)r   r   r   r   rg   Zcompare_and_swapr   r^   r   r   r   r   rn   D  s   rn   c               @   s0   e Zd ZeejZdd Zdd Z	dd Z
dS )Cuda_threadIdxc             C   s
   t tS )N)r   Macror   )rX   modr   r   r   	resolve_xU  s    zCuda_threadIdx.resolve_xc             C   s
   t tS )N)r   rq   r   )rX   rr   r   r   r   	resolve_yX  s    zCuda_threadIdx.resolve_yc             C   s
   t tS )N)r   rq   r   )rX   rr   r   r   r   	resolve_z[  s    zCuda_threadIdx.resolve_zN)r   r   r   r   Moduler   r   r   rs   rt   ru   r   r   r   r   rp   Q  s   rp   c               @   s0   e Zd ZeejZdd Zdd Z	dd Z
dS )Cuda_blockIdxc             C   s
   t tS )N)r   rq   r   )rX   rr   r   r   r   rs   c  s    zCuda_blockIdx.resolve_xc             C   s
   t tS )N)r   rq   r   )rX   rr   r   r   r   rt   f  s    zCuda_blockIdx.resolve_yc             C   s
   t tS )N)r   rq   r    )rX   rr   r   r   r   ru   i  s    zCuda_blockIdx.resolve_zN)r   r   r   r   rv   r   r   r   rs   rt   ru   r   r   r   r   rw   _  s   rw   c               @   s0   e Zd ZeejZdd Zdd Z	dd Z
dS )Cuda_blockDimc             C   s
   t tS )N)r   rq   r!   )rX   rr   r   r   r   rs   q  s    zCuda_blockDim.resolve_xc             C   s
   t tS )N)r   rq   r#   )rX   rr   r   r   r   rt   t  s    zCuda_blockDim.resolve_yc             C   s
   t tS )N)r   rq   r$   )rX   rr   r   r   r   ru   w  s    zCuda_blockDim.resolve_zN)r   r   r   r   rv   r   r"   r   rs   rt   ru   r   r   r   r   rx   m  s   rx   c               @   s0   e Zd ZeejZdd Zdd Z	dd Z
dS )Cuda_gridDimc             C   s
   t tS )N)r   rq   r%   )rX   rr   r   r   r   rs     s    zCuda_gridDim.resolve_xc             C   s
   t tS )N)r   rq   r'   )rX   rr   r   r   r   rt     s    zCuda_gridDim.resolve_yc             C   s
   t tS )N)r   rq   r(   )rX   rr   r   r   r   ru     s    zCuda_gridDim.resolve_zN)r   r   r   r   rv   r   r&   r   rs   rt   ru   r   r   r   r   ry   {  s   ry   c               @   s    e Zd ZeejZdd ZdS )CudaSharedModuleTemplatec             C   s
   t tS )N)r   rq   r+   )rX   rr   r   r   r   resolve_array  s    z&CudaSharedModuleTemplate.resolve_arrayN)	r   r   r   r   rv   r   r,   r   r{   r   r   r   r   rz     s   rz   c               @   s    e Zd ZeejZdd ZdS )CudaConstModuleTemplatec             C   s
   t tS )N)r   rq   r0   )rX   rr   r   r   r   resolve_array_like  s    z*CudaConstModuleTemplate.resolve_array_likeN)	r   r   r   r   rv   r   r1   r   r}   r   r   r   r   r|     s   r|   c               @   s    e Zd ZeejZdd ZdS )CudaLocalModuleTemplatec             C   s
   t tS )N)r   rq   r.   )rX   rr   r   r   r   r{     s    z%CudaLocalModuleTemplate.resolve_arrayN)	r   r   r   r   rv   r   r/   r   r{   r   r   r   r   r~     s   r~   c               @   s8   e Zd ZeejZdd Zdd Z	dd Z
dd Zd	S )
CudaAtomicTemplatec             C   s
   t tS )N)r   Functionr_   )rX   rr   r   r   r   resolve_add  s    zCudaAtomicTemplate.resolve_addc             C   s
   t tS )N)r   r   rj   )rX   rr   r   r   r   resolve_max  s    zCudaAtomicTemplate.resolve_maxc             C   s
   t tS )N)r   r   rl   )rX   rr   r   r   r   resolve_min  s    zCudaAtomicTemplate.resolve_minc             C   s
   t tS )N)r   r   rn   )rX   rr   r   r   r   resolve_compare_and_swap  s    z+CudaAtomicTemplate.resolve_compare_and_swapN)r   r   r   r   rv   r   rg   r   r   r   r   r   r   r   r   r   r     s
   r   c               @   s  e Zd ZeeZdd Zdd Zdd Z	dd Z
d	d
 Zdd Zdd Zdd Zdd Zdd Zdd Zdd Zdd Zdd Zdd Zdd  Zd!d" Zd#d$ Zd%d& Zd'd( Zd)d* Zd+d, Zd-d. Zd/d0 Zd1d2 Zd3d4 Z d5d6 Z!d7d8 Z"d9d: Z#d;d< Z$d=S )>CudaModuleTemplatec             C   s
   t tS )N)r   rq   r   )rX   rr   r   r   r   resolve_grid  s    zCudaModuleTemplate.resolve_gridc             C   s
   t tS )N)r   rq   r   )rX   rr   r   r   r   resolve_gridsize  s    z#CudaModuleTemplate.resolve_gridsizec             C   s   t tjS )N)r   rv   r   r   )rX   rr   r   r   r   resolve_threadIdx  s    z$CudaModuleTemplate.resolve_threadIdxc             C   s   t tjS )N)r   rv   r   r   )rX   rr   r   r   r   resolve_blockIdx  s    z#CudaModuleTemplate.resolve_blockIdxc             C   s   t tjS )N)r   rv   r   r"   )rX   rr   r   r   r   resolve_blockDim  s    z#CudaModuleTemplate.resolve_blockDimc             C   s   t tjS )N)r   rv   r   r&   )rX   rr   r   r   r   resolve_gridDim  s    z"CudaModuleTemplate.resolve_gridDimc             C   s
   t tS )N)r   rq   r)   )rX   rr   r   r   r   resolve_warpsize  s    z#CudaModuleTemplate.resolve_warpsizec             C   s
   t tS )N)r   rq   r*   )rX   rr   r   r   r   resolve_laneid  s    z!CudaModuleTemplate.resolve_laneidc             C   s   t tjS )N)r   rv   r   r,   )rX   rr   r   r   r   resolve_shared  s    z!CudaModuleTemplate.resolve_sharedc             C   s
   t tS )N)r   r   rF   )rX   rr   r   r   r   resolve_popc  s    zCudaModuleTemplate.resolve_popcc             C   s
   t tS )N)r   r   rS   )rX   rr   r   r   r   resolve_brev  s    zCudaModuleTemplate.resolve_brevc             C   s
   t tS )N)r   r   rT   )rX   rr   r   r   r   resolve_clz  s    zCudaModuleTemplate.resolve_clzc             C   s
   t tS )N)r   r   rU   )rX   rr   r   r   r   resolve_ffs  s    zCudaModuleTemplate.resolve_ffsc             C   s
   t tS )N)r   r   rP   )rX   rr   r   r   r   resolve_fma  s    zCudaModuleTemplate.resolve_fmac             C   s
   t tS )N)r   r   r2   )rX   rr   r   r   r   resolve_syncthreads  s    z&CudaModuleTemplate.resolve_syncthreadsc             C   s
   t tS )N)r   r   r5   )rX   rr   r   r   r   resolve_syncthreads_count  s    z,CudaModuleTemplate.resolve_syncthreads_countc             C   s
   t tS )N)r   r   r7   )rX   rr   r   r   r   resolve_syncthreads_and  s    z*CudaModuleTemplate.resolve_syncthreads_andc             C   s
   t tS )N)r   r   r8   )rX   rr   r   r   r   resolve_syncthreads_or  s    z)CudaModuleTemplate.resolve_syncthreads_orc             C   s
   t tS )N)r   r   r9   )rX   rr   r   r   r   resolve_threadfence  s    z&CudaModuleTemplate.resolve_threadfencec             C   s
   t tS )N)r   r   r:   )rX   rr   r   r   r   resolve_threadfence_block  s    z,CudaModuleTemplate.resolve_threadfence_blockc             C   s
   t tS )N)r   r   r;   )rX   rr   r   r   r   resolve_threadfence_system  s    z-CudaModuleTemplate.resolve_threadfence_systemc             C   s
   t tS )N)r   r   r<   )rX   rr   r   r   r   resolve_syncwarp  s    z#CudaModuleTemplate.resolve_syncwarpc             C   s
   t tS )N)r   r   r=   )rX   rr   r   r   r   resolve_shfl_sync_intrinsic  s    z.CudaModuleTemplate.resolve_shfl_sync_intrinsicc             C   s
   t tS )N)r   r   rC   )rX   rr   r   r   r   resolve_vote_sync_intrinsic  s    z.CudaModuleTemplate.resolve_vote_sync_intrinsicc             C   s
   t tS )N)r   r   rD   )rX   rr   r   r   r   resolve_match_any_sync  s    z)CudaModuleTemplate.resolve_match_any_syncc             C   s
   t tS )N)r   r   rE   )rX   rr   r   r   r   resolve_match_all_sync  s    z)CudaModuleTemplate.resolve_match_all_syncc             C   s
   t tS )N)r   r   rV   )rX   rr   r   r   r   resolve_selp  s    zCudaModuleTemplate.resolve_selpc             C   s   t tjS )N)r   rv   r   rg   )rX   rr   r   r   r   resolve_atomic  s    z!CudaModuleTemplate.resolve_atomicc             C   s   t tjS )N)r   rv   r   r1   )rX   rr   r   r   r   resolve_const
  s    z CudaModuleTemplate.resolve_constc             C   s   t tjS )N)r   rv   r   r/   )rX   rr   r   r   r   resolve_local  s    z CudaModuleTemplate.resolve_localN)%r   r   r   r   rv   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r     s>   
r   N)KZ
__future__r   r   r   Znumbar   Znumba.typing.npydeclr   Znumba.typing.templatesr   r   r	   r
   r   r   r   registryregisterZ	intrinsicZregister_attrZintrinsic_attrZregister_globalZintrinsic_globalr   r   r   r   r   r   r   r    r!   r#   r$   r%   r'   r(   r)   r*   r+   r.   r0   r2   r5   r7   r8   r9   r:   r;   r<   r=   rC   rD   rE   rF   rP   rS   rT   rU   rV   r_   ri   rj   rl   rn   rp   rw   rx   ry   rz   r|   r~   r   r   rv   r   r   r   r   <module>   s    


^