U
    hâËdV[  ã                   @   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\ƒZ€esd]ƒZ�esd^ƒZ‚esd_ƒZƒetd`ƒZ„dadb„ Z…ej†ej‡ejˆej‰ejŠej‹fZŒejˆej‰ejŠej‹fZ�ej‰ej‹fZŽe…ej�jCeŒƒZ�e…ej�jIeŒƒZ‘e…ej�j’eŒƒZ“e…ej�j”eŒƒZ•e…ej�j–eŒƒZ—e…ej�j˜eŒƒZ™e…ej�jše�ƒZ›e…ej�jœe�ƒZ�e…ej�jže�ƒZŸe…ej�j eŽƒZ¡e…ej�j¢eŽƒZ£e…ej�j¤e�ƒZ¥eG dcdd„ ddeƒƒZ¦eG dedf„ dfeƒƒZ§eG dgdh„ dheƒƒZ¨eG didj„ dje
ƒƒZ©eG dkdl„ dle
ƒƒZªeG dmdn„ dne
ƒƒZ«eG dodp„ dpe
ƒƒZ¬eG dqdr„ dre
ƒƒZ­eG dsdt„ dte
ƒƒZ®eG dudv„ dve
ƒƒZ¯eee °e¡ƒ eD ]Z±ee±eƒ �q†d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ršt
ˆ 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   	

			














  þ ÿ.^[