U
    9%e|\                     @   s  d dl Z d dlmZ d dlmZmZ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Z0eG d-d. d.eZ1eG d/d0 d0eZ2eG d1d2 d2eZ3eG d3d4 d4eZ4eG d5d6 d6eZ5eG d7d8 d8eZ6eG d9d: d:eZ7eG d;d< d<eZ8d=d> Z9d?d@ Z:dAdB Z;ee<G dCdD dDeZ=dEdF Z>dGdH Z?dIdJ Z@dKdL ZAe;ejBjCZDeAe jEZFeAe jGZHe;ejBjIZJeAe jKZLeAe jMZNe;ejBjOZPeAe jQZReAe jSZTe;ejBjUZVe;ejBjWZXe9ejBjYZZe:e j[Z\e9ejBj]Z^e:e_Z`e>ejBjaZbe@e jc e>ejBjdZee@e jf e>ejBjgZhe@e ji e>ejBjjZke@e jl e>ejBjmZne@e jo e>ejBjpZqe@e jr eAe js eAe jt dMdN ZudOdP ZveudQZweudRZxeudSZyeudTZzeudUZ{eudVZ|eudWZ}eudXZ~eudYZeudZZeud[Zeud\Zeud]Zeud^Zeud_Zevd`Zdadb ZejejejejejejfZejejejejfZejejfZeejjEeZeejjKe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 qe	D ]Zeee qe
D ]Zeee qdS )w    N)types)parse_dtypeparse_shaperegister_number_classesregister_numpy_ufunctrigonometric_functionscomparison_functionsbit_twiddling_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   R/var/www/html/Darija-Ai-API/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,   1   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   r0   6   s   r0   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   r2   ;   s   r2   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   r5   E   s   r5   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   r6   r7   r   r   r   r   r8   K   s   r8   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   r6   r7   r   r   r   r   r9   Q   s   r9   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   r6   i4r7   r   r   r   r   r:   W   s   r:   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   r7   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 r3   )r   Functionr<   r&   modr   r   r   resolve_this_gridg   s    z&CudaCgModuleTemplate.resolve_this_gridN)	r)   r*   r+   r   Moduler   r=   r/   rB   r   r   r   r   r>   c   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'   n   s    zCuda_grid_group_sync.genericNr)   r*   r+   r/   r'   r   r   r   r   rD   k   s   rD   c                   @   s   e Zd ZeZdd ZdS )GridGroup_attrsc                 C   s   t ttS r3   )r   ZBoundFunctionrD   r   r@   r   r   r   resolve_syncv   s    zGridGroup_attrs.resolve_syncN)r)   r*   r+   r   r/   rL   r   r   r   r   rK   r   s   rK   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!   r;   b1i8f4f8r7   r   r   r   r   rM   z   s<                   rM   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!   r;   rN   r7   r   r   r   r   rR      s     rR   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   r;   rO   rP   rQ   r7   r   r   r   r   rS      s   rS   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!   r;   rN   rO   rP   rQ   r7   r   r   r   r   rT      s   rT   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   uint32r7   r   r   r   r   rU      s   rU   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   rV   r7   r   r   r   r   rW      s   rW   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int16rE   int64uint8uint16rV   uint64r7   r   r   r   r   rX      s   rX   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+   rY   r   fmar/   r   r   float32float64r7   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   float16r7   r   r   r   r   rd      s   rd   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   rb   rc   r7   r   r   r   r   rg      s   rg   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   rV   r_   r7   r   r   r   r   rh      s   rh   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+   rY   r   Zclzr/   r   r   rZ   r[   rE   r\   r]   r^   rV   r_   r7   r   r   r   r   ri      s   ri   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+   rY   r   Zffsr/   r   r   rV   rZ   r[   rE   r\   r]   r^   r_   r7   r   r   r   r   rj      s   rj   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 r3   )AssertionErrorr   rc   rb   r[   r^   rE   rV   r\   r_   r   )r&   rH   rI   testabsupported_typesr   r   r   r'     s    
   zCuda_selp.genericN)r)   r*   r+   r   Zselpr/   r'   r   r   r   r   rk     s   rk   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   rf   r7   r   l_keyr   r   Cuda_fp16_unary)  s   rt   registerr   rs   rt   r   rr   r   _genfp16_unary(  s    rx   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   )rl   lenr   rf   r   rG   r   r   r   r'   6  s    z8_genfp16_unary_operator.<locals>.Cuda_fp16_unary.genericNrJ   r   rr   r   r   rt   2  s   rt   register_globalr   rw   r   rr   r   _genfp16_unary_operator1  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_binaryNrq   r   rr   r   r   Cuda_fp16_binary?  s   r~   ru   )rs   r~   r   rr   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 r3   )rl   r   rf   r   )r&   rH   rI   argr   r   r   r'   J  s    
zFloat.genericNr(   r   r   r   r   r   G  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   rN   rf   r7   r   rr   r   r   Cuda_fp16_cmpT  s   r   ru   )rs   r   r   rr   r   _genfp16_binary_comparisonS  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   ry   )rl   rz   r   rf   contextZcan_convertr   exactZpromotesafer   )r&   rH   rI   Zconvertible)rettyr   r   r'   p  s    

z9_fp16_binary_operator.<locals>.Cuda_fp16_operator.genericNrJ   r   rs   r   r   r   Cuda_fp16_operatorl  s   r   r{   )rs   r   r   r   r   r   _fp16_binary_operatork  s    r   c                 C   s   t | tjS r3   )r   r   rN   opr   r   r   _genfp16_comparison_operator  s    r   c                 C   s   t | tjS r3   )r   r   rf   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   rf   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 Nry   )rl   r   r   r   r   intp)r&   rH   rI   aryidxval)rp   r   r   r'     s    



z!_gen.<locals>.Cuda_atomic.genericNrJ   r   rs   rp   r   r   Cuda_atomic  s   r   )rv   r   )rs   rp   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   )rl   r   integer_numba_typesr   r   )r&   rH   rI   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   )rl   r   r   r   r   r   r   )r&   rH   rI   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   voidrV   r7   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 r3   r   rE   r@   r   r   r   	resolve_x&  s    zDim3_attrs.resolve_xc                 C   s   t jS r3   r   r@   r   r   r   	resolve_y)  s    zDim3_attrs.resolve_yc                 C   s   t jS r3   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 r3   )r   r?   r,   r@   r   r   r   resolve_array4  s    z&CudaSharedModuleTemplate.resolve_arrayN)	r)   r*   r+   r   rC   r   r-   r/   r   r   r   r   r   r   0  s   r   c                   @   s    e Zd ZeejZdd ZdS )CudaConstModuleTemplatec                 C   s
   t tS r3   )r   r?   r2   r@   r   r   r   resolve_array_like<  s    z*CudaConstModuleTemplate.resolve_array_likeN)	r)   r*   r+   r   rC   r   r4   r/   r   r   r   r   r   r   8  s   r   c                   @   s    e Zd ZeejZdd ZdS )CudaLocalModuleTemplatec                 C   s
   t tS r3   )r   r?   r0   r@   r   r   r   r   D  s    z%CudaLocalModuleTemplate.resolve_arrayN)	r)   r*   r+   r   rC   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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 r3   )r   r?   Cuda_atomic_addr@   r   r   r   resolve_addL  s    zCudaAtomicTemplate.resolve_addc                 C   s
   t tS r3   )r   r?   Cuda_atomic_subr@   r   r   r   resolve_subO  s    zCudaAtomicTemplate.resolve_subc                 C   s
   t tS r3   )r   r?   Cuda_atomic_andr@   r   r   r   resolve_and_R  s    zCudaAtomicTemplate.resolve_and_c                 C   s
   t tS r3   )r   r?   Cuda_atomic_orr@   r   r   r   resolve_or_U  s    zCudaAtomicTemplate.resolve_or_c                 C   s
   t tS r3   )r   r?   Cuda_atomic_xorr@   r   r   r   resolve_xorX  s    zCudaAtomicTemplate.resolve_xorc                 C   s
   t tS r3   )r   r?   Cuda_atomic_incr@   r   r   r   resolve_inc[  s    zCudaAtomicTemplate.resolve_incc                 C   s
   t tS r3   )r   r?   Cuda_atomic_decr@   r   r   r   resolve_dec^  s    zCudaAtomicTemplate.resolve_decc                 C   s
   t tS r3   )r   r?   Cuda_atomic_exchr@   r   r   r   resolve_excha  s    zCudaAtomicTemplate.resolve_exchc                 C   s
   t tS r3   )r   r?   Cuda_atomic_maxr@   r   r   r   resolve_maxd  s    zCudaAtomicTemplate.resolve_maxc                 C   s
   t tS r3   )r   r?   Cuda_atomic_minr@   r   r   r   resolve_ming  s    zCudaAtomicTemplate.resolve_minc                 C   s
   t tS r3   )r   r?   Cuda_atomic_nanminr@   r   r   r   resolve_nanminj  s    z!CudaAtomicTemplate.resolve_nanminc                 C   s
   t tS r3   )r   r?   Cuda_atomic_nanmaxr@   r   r   r   resolve_nanmaxm  s    z!CudaAtomicTemplate.resolve_nanmaxc                 C   s
   t tS r3   )r   r?   r   r@   r   r   r   resolve_compare_and_swapp  s    z+CudaAtomicTemplate.resolve_compare_and_swapc                 C   s
   t tS r3   )r   r?   r   r@   r   r   r   resolve_cass  s    zCudaAtomicTemplate.resolve_casN)r)   r*   r+   r   rC   r   r   r/   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   H  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 r3   )r   r?   	Cuda_haddr@   r   r   r   resolve_hadd{  s    zCudaFp16Template.resolve_haddc                 C   s
   t tS r3   )r   r?   	Cuda_hsubr@   r   r   r   resolve_hsub~  s    zCudaFp16Template.resolve_hsubc                 C   s
   t tS r3   )r   r?   	Cuda_hmulr@   r   r   r   resolve_hmul  s    zCudaFp16Template.resolve_hmulc                 C   s   t S r3   )hdiv_devicer@   r   r   r   resolve_hdiv  s    zCudaFp16Template.resolve_hdivc                 C   s
   t tS r3   )r   r?   	Cuda_hnegr@   r   r   r   resolve_hneg  s    zCudaFp16Template.resolve_hnegc                 C   s
   t tS r3   )r   r?   	Cuda_habsr@   r   r   r   resolve_habs  s    zCudaFp16Template.resolve_habsc                 C   s
   t tS r3   )r   r?   rd   r@   r   r   r   resolve_hfma  s    zCudaFp16Template.resolve_hfmac                 C   s   t S r3   )hsin_devicer@   r   r   r   resolve_hsin  s    zCudaFp16Template.resolve_hsinc                 C   s   t S r3   )hcos_devicer@   r   r   r   resolve_hcos  s    zCudaFp16Template.resolve_hcosc                 C   s   t S r3   )hlog_devicer@   r   r   r   resolve_hlog  s    zCudaFp16Template.resolve_hlogc                 C   s   t S r3   )hlog10_devicer@   r   r   r   resolve_hlog10  s    zCudaFp16Template.resolve_hlog10c                 C   s   t S r3   )hlog2_devicer@   r   r   r   resolve_hlog2  s    zCudaFp16Template.resolve_hlog2c                 C   s   t S r3   )hexp_devicer@   r   r   r   resolve_hexp  s    zCudaFp16Template.resolve_hexpc                 C   s   t S r3   )hexp10_devicer@   r   r   r   resolve_hexp10  s    zCudaFp16Template.resolve_hexp10c                 C   s   t S r3   )hexp2_devicer@   r   r   r   resolve_hexp2  s    zCudaFp16Template.resolve_hexp2c                 C   s   t S r3   )hfloor_devicer@   r   r   r   resolve_hfloor  s    zCudaFp16Template.resolve_hfloorc                 C   s   t S r3   )hceil_devicer@   r   r   r   resolve_hceil  s    zCudaFp16Template.resolve_hceilc                 C   s   t S r3   )hsqrt_devicer@   r   r   r   resolve_hsqrt  s    zCudaFp16Template.resolve_hsqrtc                 C   s   t S r3   )hrsqrt_devicer@   r   r   r   resolve_hrsqrt  s    zCudaFp16Template.resolve_hrsqrtc                 C   s   t S r3   )hrcp_devicer@   r   r   r   resolve_hrcp  s    zCudaFp16Template.resolve_hrcpc                 C   s   t S r3   )hrint_devicer@   r   r   r   resolve_hrint  s    zCudaFp16Template.resolve_hrintc                 C   s   t S r3   )htrunc_devicer@   r   r   r   resolve_htrunc  s    zCudaFp16Template.resolve_htruncc                 C   s
   t tS r3   )r   r?   Cuda_heqr@   r   r   r   resolve_heq  s    zCudaFp16Template.resolve_heqc                 C   s
   t tS r3   )r   r?   Cuda_hner@   r   r   r   resolve_hne  s    zCudaFp16Template.resolve_hnec                 C   s
   t tS r3   )r   r?   Cuda_hger@   r   r   r   resolve_hge  s    zCudaFp16Template.resolve_hgec                 C   s
   t tS r3   )r   r?   Cuda_hgtr@   r   r   r   resolve_hgt  s    zCudaFp16Template.resolve_hgtc                 C   s
   t tS r3   )r   r?   Cuda_hler@   r   r   r   resolve_hle  s    zCudaFp16Template.resolve_hlec                 C   s
   t tS r3   )r   r?   Cuda_hltr@   r   r   r   resolve_hlt  s    zCudaFp16Template.resolve_hltc                 C   s
   t tS r3   )r   r?   	Cuda_hmaxr@   r   r   r   resolve_hmax  s    zCudaFp16Template.resolve_hmaxc                 C   s
   t tS r3   )r   r?   	Cuda_hminr@   r   r   r   resolve_hmin  s    zCudaFp16Template.resolve_hminN)&r)   r*   r+   r   rC   r   re   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   w  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 r3   )r   rC   r   r=   r@   r   r   r   
resolve_cg  s    zCudaModuleTemplate.resolve_cgc                 C   s   t S r3   r   r@   r   r   r   resolve_threadIdx  s    z$CudaModuleTemplate.resolve_threadIdxc                 C   s   t S r3   r
  r@   r   r   r   resolve_blockIdx  s    z#CudaModuleTemplate.resolve_blockIdxc                 C   s   t S r3   r
  r@   r   r   r   resolve_blockDim  s    z#CudaModuleTemplate.resolve_blockDimc                 C   s   t S r3   r
  r@   r   r   r   resolve_gridDim  s    z"CudaModuleTemplate.resolve_gridDimc                 C   s   t jS r3   r   r@   r   r   r   resolve_laneid  s    z!CudaModuleTemplate.resolve_laneidc                 C   s   t tjS r3   )r   rC   r   r-   r@   r   r   r   resolve_shared  s    z!CudaModuleTemplate.resolve_sharedc                 C   s
   t tS r3   )r   r?   rX   r@   r   r   r   resolve_popc  s    zCudaModuleTemplate.resolve_popcc                 C   s
   t tS r3   )r   r?   rh   r@   r   r   r   resolve_brev  s    zCudaModuleTemplate.resolve_brevc                 C   s
   t tS r3   )r   r?   ri   r@   r   r   r   resolve_clz  s    zCudaModuleTemplate.resolve_clzc                 C   s
   t tS r3   )r   r?   rj   r@   r   r   r   resolve_ffs  s    zCudaModuleTemplate.resolve_ffsc                 C   s
   t tS r3   )r   r?   r`   r@   r   r   r   resolve_fma  s    zCudaModuleTemplate.resolve_fmac                 C   s
   t tS r3   )r   r?   rg   r@   r   r   r   resolve_cbrt  s    zCudaModuleTemplate.resolve_cbrtc                 C   s
   t tS r3   )r   r?   r5   r@   r   r   r   resolve_threadfence  s    z&CudaModuleTemplate.resolve_threadfencec                 C   s
   t tS r3   )r   r?   r8   r@   r   r   r   resolve_threadfence_block  s    z,CudaModuleTemplate.resolve_threadfence_blockc                 C   s
   t tS r3   )r   r?   r9   r@   r   r   r   resolve_threadfence_system  s    z-CudaModuleTemplate.resolve_threadfence_systemc                 C   s
   t tS r3   )r   r?   r:   r@   r   r   r   resolve_syncwarp
  s    z#CudaModuleTemplate.resolve_syncwarpc                 C   s
   t tS r3   )r   r?   rM   r@   r   r   r   resolve_shfl_sync_intrinsic  s    z.CudaModuleTemplate.resolve_shfl_sync_intrinsicc                 C   s
   t tS r3   )r   r?   rR   r@   r   r   r   resolve_vote_sync_intrinsic  s    z.CudaModuleTemplate.resolve_vote_sync_intrinsicc                 C   s
   t tS r3   )r   r?   rS   r@   r   r   r   resolve_match_any_sync  s    z)CudaModuleTemplate.resolve_match_any_syncc                 C   s
   t tS r3   )r   r?   rT   r@   r   r   r   resolve_match_all_sync  s    z)CudaModuleTemplate.resolve_match_all_syncc                 C   s
   t tS r3   )r   r?   rU   r@   r   r   r   resolve_activemask  s    z%CudaModuleTemplate.resolve_activemaskc                 C   s
   t tS r3   )r   r?   rW   r@   r   r   r   resolve_lanemask_lt  s    z&CudaModuleTemplate.resolve_lanemask_ltc                 C   s
   t tS r3   )r   r?   rk   r@   r   r   r   resolve_selp  s    zCudaModuleTemplate.resolve_selpc                 C   s
   t tS r3   )r   r?   r   r@   r   r   r   resolve_nanosleep"  s    z$CudaModuleTemplate.resolve_nanosleepc                 C   s   t tjS r3   )r   rC   r   r   r@   r   r   r   resolve_atomic%  s    z!CudaModuleTemplate.resolve_atomicc                 C   s   t tjS r3   )r   rC   r   re   r@   r   r   r   resolve_fp16(  s    zCudaModuleTemplate.resolve_fp16c                 C   s   t tjS r3   )r   rC   r   r4   r@   r   r   r   resolve_const+  s    z CudaModuleTemplate.resolve_constc                 C   s   t tjS r3   )r   rC   r   r1   r@   r   r   r   resolve_local.  s    z CudaModuleTemplate.resolve_localN)$r)   r*   r+   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    s<   
r  )operatorZ
numba.corer   Znumba.core.typing.npydeclr   r   r   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   registryrv   Zregister_attrr|   r   r,   r0   r2   r5   r8   r9   r:   r<   r>   rD   rK   rM   rR   rS   rT   rU   rW   rX   r`   rd   rg   rh   ri   rj   rk   rx   r}   r   floatr   r   r   r   r   re   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   rc   rb   rE   rV   r\   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  rC   funcr   r   r   r   <module>   sR  $ 	

			














   .^[