U
    hdV[                     @   s  d dl Z d dlmZ d dlmZmZmZmZmZ d dl	m
Z
mZmZmZmZmZ d dlmZmZ d dlmZ d dl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eG dd deZeG dd deZeG dd deZ eG dd deZ!eG dd deZ"eG dd deZ#eG dd deZ$eG dd deZ%e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 d#d$ d$eZ*eG d%d& d&eZ+eG d'd( d(eZ,eG d)d* d*eZ-eG d+d, d,eZ.eG d-d. d.eZ/eG d/d0 d0eZ0eG d1d2 d2eZ1eG d3d4 d4eZ2eG d5d6 d6eZ3eG d7d8 d8eZ4eG d9d: d:eZ5eG d;d< d<eZ6d=d> Z7d?d@ Z8dAdB Z9ee:G dCdD dDeZ;dEdF Z<dGdH Z=dIdJ Z>dKdL Z?e9ej@jAZBe?e jCZDe?e jEZFe9ej@jGZHe?e jIZJe?e jKZLe9ej@jMZNe?e jOZPe?e jQZRe9ej@jSZTe9ej@jUZVe7ej@jWZXe8e jYZZe7ej@j[Z\e8e]Z^e<ej@j_Z`e>e ja e<ej@jbZce>e jd e<ej@jeZfe>e jg e<ej@jhZie>e jj e<ej@jkZle>e jm e<ej@jnZoe>e jp e?e jq e?e jr dMdN ZsdOdP ZtesdQZuesdRZvesdSZwesdTZxesdUZyesdVZzesdWZ{esdXZ|esdYZ}esdZZ~esd[Zesd\Zesd]Zesd^Zesd_Zetd`Zdadb ZejejejejejejfZejejejejfZejejfZeejjCeZeejjIeZeejjeZeejjeZeejjeZeejjeZeejjeZeejjeZeejjeZeejjeZeejjeZeejjeZeG dcdd ddeZeG dedf dfeZeG dgdh dheZeG didj dje
ZeG dkdl dle
ZeG dmdn dne
ZeG dodp dpe
ZeG dqdr dre
ZeG dsdt dte
ZeG dudv dve
Zeeee eD ]Zeee qdS )w    N)types)parse_dtypeparse_shaperegister_number_classesregister_numpy_ufunctrigonometric_functions)AttributeTemplateConcreteTemplateAbstractTemplateCallableTemplate	signatureRegistry)dim3
grid_group)
Conversion)cuda) declare_device_function_templatec                   @   s   e Zd Zdd ZdS )Cuda_array_declc                 C   s   dd }|S )Nc                 S   s   t | tjrt | tjsLd S n.t | tjtjfrHtdd | D rLd S nd S t| }t|}|d k	r||d k	r|tj	||ddS d S )Nc                 S   s   g | ]}t |tj qS  )
isinstancer   IntegerLiteral).0sr   r   L/home/sam/Atlas/atlas_env/lib/python3.8/site-packages/numba/cuda/cudadecl.py
<listcomp>!   s   z:Cuda_array_decl.generic.<locals>.typer.<locals>.<listcomp>C)dtypendimZlayout)
r   r   Integerr   TupleZUniTupleanyr   r   ZArray)shaper   r   Znb_dtyper   r   r   typer   s    z&Cuda_array_decl.generic.<locals>.typerr   selfr"   r   r   r   generic   s    zCuda_array_decl.genericN__name__
__module____qualname__r%   r   r   r   r   r      s   r   c                   @   s   e Zd ZejjZdS )Cuda_shared_arrayN)r'   r(   r)   r   sharedarraykeyr   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.   4   s   r.   c                   @   s   e Zd ZejjZdd ZdS )Cuda_const_array_likec                 C   s   dd }|S )Nc                 S   s   | S Nr   )Zndarrayr   r   r   r"   >   s    z,Cuda_const_array_like.generic.<locals>.typerr   r#   r   r   r   r%   =   s    zCuda_const_array_like.genericN)r'   r(   r)   r   constZ
array_liker-   r%   r   r   r   r   r0   9   s   r0   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   nonecasesr   r   r   r   r3   C   s   r3   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   r4   r5   r   r   r   r   r6   I   s   r6   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   r4   r5   r   r   r   r   r7   O   s   r7   c                   @   s*   e Zd ZejZeejeejej	gZ
dS )Cuda_syncwarpN)r'   r(   r)   r   Zsyncwarpr-   r   r   r4   i4r5   r   r   r   r   r8   U   s   r8   c                   @   s   e Zd ZejjZeegZ	dS )Cuda_cg_this_gridN)
r'   r(   r)   r   cgZ	this_gridr-   r   r   r5   r   r   r   r   r:   [   s   r:   c                   @   s    e Zd ZeejZdd ZdS )CudaCgModuleTemplatec                 C   s
   t tS r1   )r   Functionr:   r$   modr   r   r   resolve_this_gride   s    z&CudaCgModuleTemplate.resolve_this_gridN)	r'   r(   r)   r   Moduler   r;   r-   r@   r   r   r   r   r<   a   s   r<   c                   @   s   e Zd ZdZdd ZdS )Cuda_grid_group_synczGridGroup.syncc                 C   s   t tj| jdS )N)Zrecvr)r   r   int32thisr$   argskwsr   r   r   r%   l   s    zCuda_grid_group_sync.genericNr'   r(   r)   r-   r%   r   r   r   r   rB   i   s   rB   c                   @   s   e Zd ZeZdd ZdS )GridGroup_attrsc                 C   s   t ttS r1   )r   ZBoundFunctionrB   r   r>   r   r   r   resolve_synct   s    zGridGroup_attrs.resolve_syncN)r'   r(   r)   r   r-   rJ   r   r   r   r   rI   p   s   rI   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   r   r9   b1i8f4f8r5   r   r   r   r   rK   x   s<                   rK   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   r9   rL   r5   r   r   r   r   rP      s     rP   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   r9   rM   rN   rO   r5   r   r   r   r   rQ      s   rQ   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   r9   rL   rM   rN   rO   r5   r   r   r   r   rR      s   rR   c                   @   s   e Zd ZejZeejgZ	dS )Cuda_activemaskN)
r'   r(   r)   r   Z
activemaskr-   r   r   uint32r5   r   r   r   r   rS      s   rS   c                   @   s   e Zd ZejZeejgZ	dS )Cuda_lanemask_ltN)
r'   r(   r)   r   Zlanemask_ltr-   r   r   rT   r5   r   r   r   r   rU      s   rU   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int16rC   int64uint8uint16rT   uint64r5   r   r   r   r   rV      s   rV   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)   rW   r   fmar-   r   r   float32float64r5   r   r   r   r   r^      s
   r^   c                   @   s,   e Zd ZejjZeej	ej	ej	ej	gZ
dS )	Cuda_hfmaN)r'   r(   r)   r   fp16Zhfmar-   r   r   float16r5   r   r   r   r   rb      s   rb   c                   @   s.   e Zd ZejZeejejeej	ej	gZ
dS )	Cuda_cbrtN)r'   r(   r)   r   Zcbrtr-   r   r   r`   ra   r5   r   r   r   r   re      s   re   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   rT   r]   r5   r   r   r   r   rf      s   rf   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)   rW   r   Zclzr-   r   r   rX   rY   rC   rZ   r[   r\   rT   r]   r5   r   r   r   r   rg      s   rg   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)   rW   r   Zffsr-   r   r   rT   rX   rY   rC   rZ   r[   r\   r]   r5   r   r   r   r   rh      s   rh   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 r1   )AssertionErrorr   ra   r`   rY   r\   rC   rT   rZ   r]   r   )r$   rF   rG   testabsupported_typesr   r   r   r%     s    
   zCuda_selp.genericN)r'   r(   r)   r   Zselpr-   r%   r   r   r   r   ri     s   ri   c                    s   t G  fdddt}|S )Nc                       s    e Zd Z ZeejejgZdS )z'_genfp16_unary.<locals>.Cuda_fp16_unaryNr'   r(   r)   r-   r   r   rd   r5   r   l_keyr   r   Cuda_fp16_unary'  s   rr   registerr	   rq   rr   r   rp   r   _genfp16_unary&  s    rv   c                    s    t  G  fdddt}|S )Nc                       s   e Zd Z Zdd ZdS )z0_genfp16_unary_operator.<locals>.Cuda_fp16_unaryc                 S   s4   |rt t|dkr0|d tjkr0ttjtjS d S )N   r   )rj   lenr   rd   r   rE   r   r   r   r%   4  s    z8_genfp16_unary_operator.<locals>.Cuda_fp16_unary.genericNrH   r   rp   r   r   rr   0  s   rr   register_globalr
   ru   r   rp   r   _genfp16_unary_operator/  s    r{   c                    s   t G  fdddt}|S )Nc                       s$   e Zd Z ZeejejejgZdS )z)_genfp16_binary.<locals>.Cuda_fp16_binaryNro   r   rp   r   r   Cuda_fp16_binary=  s   r|   rs   )rq   r|   r   rp   r   _genfp16_binary<  s    r}   c                   @   s   e Zd Zdd ZdS )Floatc                 C   s&   |rt |\}|tjkr"t||S d S r1   )rj   r   rd   r   )r$   rF   rG   argr   r   r   r%   H  s    
zFloat.genericNr&   r   r   r   r   r~   E  s   r~   c                    s   t G  fdddt}|S )Nc                       s$   e Zd Z ZeejejejgZdS )z1_genfp16_binary_comparison.<locals>.Cuda_fp16_cmpN)	r'   r(   r)   r-   r   r   rL   rd   r5   r   rp   r   r   Cuda_fp16_cmpR  s   r   rs   )rq   r   r   rp   r   _genfp16_binary_comparisonQ  s    r   c                    s"   t  G  fdddt}|S )Nc                       s   e Zd Z ZfddZdS )z1_fp16_binary_operator.<locals>.Cuda_fp16_operatorc                    s   |rt t|dkr|d tjks0|d tjkr|d tjkrV| j|d |d }n| j|d |d }|tjks|tjks|tj	krt
 tjtjS d S )N   r   rw   )rj   rx   r   rd   contextZcan_convertr   exactZpromotesafer   )r$   rF   rG   Zconvertible)rettyr   r   r%   n  s    

z9_fp16_binary_operator.<locals>.Cuda_fp16_operator.genericNrH   r   rq   r   r   r   Cuda_fp16_operatorj  s   r   ry   )rq   r   r   r   r   r   _fp16_binary_operatori  s    r   c                 C   s   t | tjS r1   )r   r   rL   opr   r   r   _genfp16_comparison_operator  s    r   c                 C   s   t | tjS r1   )r   r   rd   r   r   r   r   _genfp16_binary_operator  s    r   c                 C   s"   t d|  tjtjf}t|S NZ__numba_wrapper_r   r   rd   r=   fnamedeclr   r   r   _resolve_wrapped_unary  s
    
r   c                 C   s&   t d|  tjtjtjf}t|S r   r   r   r   r   r   _resolve_wrapped_binary  s
    

r   ZhsinZhcosZhlogZhlog10Zhlog2ZhexpZhexp10Zhexp2ZhsqrtZhrsqrtZhfloorZhceilZhrcpZhrintZhtruncZhdivc                    s   t G  fdddt}|S )Nc                       s   e Zd Z ZfddZdS )z_gen.<locals>.Cuda_atomicc                    s^   |rt |\}}}|j kr d S |jdkr>t|j|tj|jS |jdkrZt|j|||jS d S Nrw   )rj   r   r   r   r   intp)r$   rF   rG   aryidxval)rn   r   r   r%     s    



z!_gen.<locals>.Cuda_atomic.genericNrH   r   rq   rn   r   r   Cuda_atomic  s   r   )rt   r
   )rq   rn   r   r   r   r   _gen  s    r   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kr8|jdkr8t||||S d S r   )rj   r   integer_numba_typesr   r   )r$   rF   rG   r   oldr   dtyr   r   r   r%     s
    
z$Cuda_atomic_compare_and_swap.genericN)r'   r(   r)   r   atomicZcompare_and_swapr-   r%   r   r   r   r   r     s   r   c                   @   s   e Zd ZejjZdd ZdS )Cuda_atomic_casc                 C   s`   |rt |\}}}}|j}|tkr&d S |jdkrBt||tj||S |jdkr\t|||||S d S r   )rj   r   r   r   r   r   r   )r$   rF   rG   r   r   r   r   r   r   r   r   r%     s    

zCuda_atomic_cas.genericN)r'   r(   r)   r   r   Zcasr-   r%   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_nanosleepN)r'   r(   r)   r   Z	nanosleepr-   r   r   voidrT   r5   r   r   r   r   r     s   r   c                   @   s(   e Zd ZeZdd Zdd Zdd ZdS )
Dim3_attrsc                 C   s   t jS r1   r   rC   r>   r   r   r   	resolve_x$  s    zDim3_attrs.resolve_xc                 C   s   t jS r1   r   r>   r   r   r   	resolve_y'  s    zDim3_attrs.resolve_yc                 C   s   t jS r1   r   r>   r   r   r   	resolve_z*  s    zDim3_attrs.resolve_zN)r'   r(   r)   r   r-   r   r   r   r   r   r   r   r      s   r   c                   @   s    e Zd ZeejZdd ZdS )CudaSharedModuleTemplatec                 C   s
   t tS r1   )r   r=   r*   r>   r   r   r   resolve_array2  s    z&CudaSharedModuleTemplate.resolve_arrayN)	r'   r(   r)   r   rA   r   r+   r-   r   r   r   r   r   r   .  s   r   c                   @   s    e Zd ZeejZdd ZdS )CudaConstModuleTemplatec                 C   s
   t tS r1   )r   r=   r0   r>   r   r   r   resolve_array_like:  s    z*CudaConstModuleTemplate.resolve_array_likeN)	r'   r(   r)   r   rA   r   r2   r-   r   r   r   r   r   r   6  s   r   c                   @   s    e Zd ZeejZdd ZdS )CudaLocalModuleTemplatec                 C   s
   t tS r1   )r   r=   r.   r>   r   r   r   r   B  s    z%CudaLocalModuleTemplate.resolve_arrayN)	r'   r(   r)   r   rA   r   r/   r-   r   r   r   r   r   r   >  s   r   c                   @   s   e Zd Zeej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S )CudaAtomicTemplatec                 C   s
   t tS r1   )r   r=   Cuda_atomic_addr>   r   r   r   resolve_addJ  s    zCudaAtomicTemplate.resolve_addc                 C   s
   t tS r1   )r   r=   Cuda_atomic_subr>   r   r   r   resolve_subM  s    zCudaAtomicTemplate.resolve_subc                 C   s
   t tS r1   )r   r=   Cuda_atomic_andr>   r   r   r   resolve_and_P  s    zCudaAtomicTemplate.resolve_and_c                 C   s
   t tS r1   )r   r=   Cuda_atomic_orr>   r   r   r   resolve_or_S  s    zCudaAtomicTemplate.resolve_or_c                 C   s
   t tS r1   )r   r=   Cuda_atomic_xorr>   r   r   r   resolve_xorV  s    zCudaAtomicTemplate.resolve_xorc                 C   s
   t tS r1   )r   r=   Cuda_atomic_incr>   r   r   r   resolve_incY  s    zCudaAtomicTemplate.resolve_incc                 C   s
   t tS r1   )r   r=   Cuda_atomic_decr>   r   r   r   resolve_dec\  s    zCudaAtomicTemplate.resolve_decc                 C   s
   t tS r1   )r   r=   Cuda_atomic_exchr>   r   r   r   resolve_exch_  s    zCudaAtomicTemplate.resolve_exchc                 C   s
   t tS r1   )r   r=   Cuda_atomic_maxr>   r   r   r   resolve_maxb  s    zCudaAtomicTemplate.resolve_maxc                 C   s
   t tS r1   )r   r=   Cuda_atomic_minr>   r   r   r   resolve_mine  s    zCudaAtomicTemplate.resolve_minc                 C   s
   t tS r1   )r   r=   Cuda_atomic_nanminr>   r   r   r   resolve_nanminh  s    z!CudaAtomicTemplate.resolve_nanminc                 C   s
   t tS r1   )r   r=   Cuda_atomic_nanmaxr>   r   r   r   resolve_nanmaxk  s    z!CudaAtomicTemplate.resolve_nanmaxc                 C   s
   t tS r1   )r   r=   r   r>   r   r   r   resolve_compare_and_swapn  s    z+CudaAtomicTemplate.resolve_compare_and_swapc                 C   s
   t tS r1   )r   r=   r   r>   r   r   r   resolve_casq  s    zCudaAtomicTemplate.resolve_casN)r'   r(   r)   r   rA   r   r   r-   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   F  s   r   c                   @   s  e Zd Zeej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 )>CudaFp16Templatec                 C   s
   t tS r1   )r   r=   	Cuda_haddr>   r   r   r   resolve_haddy  s    zCudaFp16Template.resolve_haddc                 C   s
   t tS r1   )r   r=   	Cuda_hsubr>   r   r   r   resolve_hsub|  s    zCudaFp16Template.resolve_hsubc                 C   s
   t tS r1   )r   r=   	Cuda_hmulr>   r   r   r   resolve_hmul  s    zCudaFp16Template.resolve_hmulc                 C   s   t S r1   )hdiv_devicer>   r   r   r   resolve_hdiv  s    zCudaFp16Template.resolve_hdivc                 C   s
   t tS r1   )r   r=   	Cuda_hnegr>   r   r   r   resolve_hneg  s    zCudaFp16Template.resolve_hnegc                 C   s
   t tS r1   )r   r=   	Cuda_habsr>   r   r   r   resolve_habs  s    zCudaFp16Template.resolve_habsc                 C   s
   t tS r1   )r   r=   rb   r>   r   r   r   resolve_hfma  s    zCudaFp16Template.resolve_hfmac                 C   s   t S r1   )hsin_devicer>   r   r   r   resolve_hsin  s    zCudaFp16Template.resolve_hsinc                 C   s   t S r1   )hcos_devicer>   r   r   r   resolve_hcos  s    zCudaFp16Template.resolve_hcosc                 C   s   t S r1   )hlog_devicer>   r   r   r   resolve_hlog  s    zCudaFp16Template.resolve_hlogc                 C   s   t S r1   )hlog10_devicer>   r   r   r   resolve_hlog10  s    zCudaFp16Template.resolve_hlog10c                 C   s   t S r1   )hlog2_devicer>   r   r   r   resolve_hlog2  s    zCudaFp16Template.resolve_hlog2c                 C   s   t S r1   )hexp_devicer>   r   r   r   resolve_hexp  s    zCudaFp16Template.resolve_hexpc                 C   s   t S r1   )hexp10_devicer>   r   r   r   resolve_hexp10  s    zCudaFp16Template.resolve_hexp10c                 C   s   t S r1   )hexp2_devicer>   r   r   r   resolve_hexp2  s    zCudaFp16Template.resolve_hexp2c                 C   s   t S r1   )hfloor_devicer>   r   r   r   resolve_hfloor  s    zCudaFp16Template.resolve_hfloorc                 C   s   t S r1   )hceil_devicer>   r   r   r   resolve_hceil  s    zCudaFp16Template.resolve_hceilc                 C   s   t S r1   )hsqrt_devicer>   r   r   r   resolve_hsqrt  s    zCudaFp16Template.resolve_hsqrtc                 C   s   t S r1   )hrsqrt_devicer>   r   r   r   resolve_hrsqrt  s    zCudaFp16Template.resolve_hrsqrtc                 C   s   t S r1   )hrcp_devicer>   r   r   r   resolve_hrcp  s    zCudaFp16Template.resolve_hrcpc                 C   s   t S r1   )hrint_devicer>   r   r   r   resolve_hrint  s    zCudaFp16Template.resolve_hrintc                 C   s   t S r1   )htrunc_devicer>   r   r   r   resolve_htrunc  s    zCudaFp16Template.resolve_htruncc                 C   s
   t tS r1   )r   r=   Cuda_heqr>   r   r   r   resolve_heq  s    zCudaFp16Template.resolve_heqc                 C   s
   t tS r1   )r   r=   Cuda_hner>   r   r   r   resolve_hne  s    zCudaFp16Template.resolve_hnec                 C   s
   t tS r1   )r   r=   Cuda_hger>   r   r   r   resolve_hge  s    zCudaFp16Template.resolve_hgec                 C   s
   t tS r1   )r   r=   Cuda_hgtr>   r   r   r   resolve_hgt  s    zCudaFp16Template.resolve_hgtc                 C   s
   t tS r1   )r   r=   Cuda_hler>   r   r   r   resolve_hle  s    zCudaFp16Template.resolve_hlec                 C   s
   t tS r1   )r   r=   Cuda_hltr>   r   r   r   resolve_hlt  s    zCudaFp16Template.resolve_hltc                 C   s
   t tS r1   )r   r=   	Cuda_hmaxr>   r   r   r   resolve_hmax  s    zCudaFp16Template.resolve_hmaxc                 C   s
   t tS r1   )r   r=   	Cuda_hminr>   r   r   r   resolve_hmin  s    zCudaFp16Template.resolve_hminN)&r'   r(   r)   r   rA   r   rc   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   u  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;S )<CudaModuleTemplatec                 C   s   t tjS r1   )r   rA   r   r;   r>   r   r   r   
resolve_cg  s    zCudaModuleTemplate.resolve_cgc                 C   s   t S r1   r   r>   r   r   r   resolve_threadIdx  s    z$CudaModuleTemplate.resolve_threadIdxc                 C   s   t S r1   r  r>   r   r   r   resolve_blockIdx  s    z#CudaModuleTemplate.resolve_blockIdxc                 C   s   t S r1   r  r>   r   r   r   resolve_blockDim  s    z#CudaModuleTemplate.resolve_blockDimc                 C   s   t S r1   r  r>   r   r   r   resolve_gridDim  s    z"CudaModuleTemplate.resolve_gridDimc                 C   s   t jS r1   r   r>   r   r   r   resolve_laneid  s    z!CudaModuleTemplate.resolve_laneidc                 C   s   t tjS r1   )r   rA   r   r+   r>   r   r   r   resolve_shared  s    z!CudaModuleTemplate.resolve_sharedc                 C   s
   t tS r1   )r   r=   rV   r>   r   r   r   resolve_popc  s    zCudaModuleTemplate.resolve_popcc                 C   s
   t tS r1   )r   r=   rf   r>   r   r   r   resolve_brev  s    zCudaModuleTemplate.resolve_brevc                 C   s
   t tS r1   )r   r=   rg   r>   r   r   r   resolve_clz  s    zCudaModuleTemplate.resolve_clzc                 C   s
   t tS r1   )r   r=   rh   r>   r   r   r   resolve_ffs  s    zCudaModuleTemplate.resolve_ffsc                 C   s
   t tS r1   )r   r=   r^   r>   r   r   r   resolve_fma  s    zCudaModuleTemplate.resolve_fmac                 C   s
   t tS r1   )r   r=   re   r>   r   r   r   resolve_cbrt  s    zCudaModuleTemplate.resolve_cbrtc                 C   s
   t tS r1   )r   r=   r3   r>   r   r   r   resolve_threadfence  s    z&CudaModuleTemplate.resolve_threadfencec                 C   s
   t tS r1   )r   r=   r6   r>   r   r   r   resolve_threadfence_block  s    z,CudaModuleTemplate.resolve_threadfence_blockc                 C   s
   t tS r1   )r   r=   r7   r>   r   r   r   resolve_threadfence_system  s    z-CudaModuleTemplate.resolve_threadfence_systemc                 C   s
   t tS r1   )r   r=   r8   r>   r   r   r   resolve_syncwarp  s    z#CudaModuleTemplate.resolve_syncwarpc                 C   s
   t tS r1   )r   r=   rK   r>   r   r   r   resolve_shfl_sync_intrinsic  s    z.CudaModuleTemplate.resolve_shfl_sync_intrinsicc                 C   s
   t tS r1   )r   r=   rP   r>   r   r   r   resolve_vote_sync_intrinsic  s    z.CudaModuleTemplate.resolve_vote_sync_intrinsicc                 C   s
   t tS r1   )r   r=   rQ   r>   r   r   r   resolve_match_any_sync  s    z)CudaModuleTemplate.resolve_match_any_syncc                 C   s
   t tS r1   )r   r=   rR   r>   r   r   r   resolve_match_all_sync  s    z)CudaModuleTemplate.resolve_match_all_syncc                 C   s
   t tS r1   )r   r=   rS   r>   r   r   r   resolve_activemask  s    z%CudaModuleTemplate.resolve_activemaskc                 C   s
   t tS r1   )r   r=   rU   r>   r   r   r   resolve_lanemask_lt  s    z&CudaModuleTemplate.resolve_lanemask_ltc                 C   s
   t tS r1   )r   r=   ri   r>   r   r   r   resolve_selp  s    zCudaModuleTemplate.resolve_selpc                 C   s
   t tS r1   )r   r=   r   r>   r   r   r   resolve_nanosleep   s    z$CudaModuleTemplate.resolve_nanosleepc                 C   s   t tjS r1   )r   rA   r   r   r>   r   r   r   resolve_atomic#  s    z!CudaModuleTemplate.resolve_atomicc                 C   s   t tjS r1   )r   rA   r   rc   r>   r   r   r   resolve_fp16&  s    zCudaModuleTemplate.resolve_fp16c                 C   s   t tjS r1   )r   rA   r   r2   r>   r   r   r   resolve_const)  s    z CudaModuleTemplate.resolve_constc                 C   s   t tjS r1   )r   rA   r   r/   r>   r   r   r   resolve_local,  s    z CudaModuleTemplate.resolve_localN)$r'   r(   r)   r   rA   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  )operatorZ
numba.corer   Znumba.core.typing.npydeclr   r   r   r   r   Znumba.core.typing.templatesr   r	   r
   r   r   r   Znumba.cuda.typesr   r   Znumba.core.typeconvr   Znumbar   Znumba.cuda.compilerr   registryrt   Zregister_attrrz   r   r*   r.   r0   r3   r6   r7   r8   r:   r<   rB   rI   rK   rP   rQ   rR   rS   rU   rV   r^   rb   re   rf   rg   rh   ri   rv   r{   r}   floatr~   r   r   r   r   rc   Zhaddr   addZCuda_addiaddZ	Cuda_iaddZhsubr   subZCuda_subisubZ	Cuda_isubZhmulr   mulZCuda_mulimulZ	Cuda_imulZhmaxr  Zhminr  Zhnegr   negZCuda_negZhabsr   absZCuda_absZheqr   eqhner   neZhger   geZhgtr   gthler   leZhltr   lttruedivitruedivr   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   ra   r`   rC   rT   rZ   r]   Zall_numba_typesr   Zunsigned_int_numba_typesr   r   r   maxr   minr   Znanmaxr   Znanminr   and_r   or_r   xorr   incr   decr   Zexchr   r   r   r   r   r   r   r   r   r   r  rA   funcr   r   r   r   <module>   sJ   	

			














   .^[