o
    i|\                     @   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 qGe	D ]Zeee qRe
D ]Zeee q]d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                   @      e Zd Zdd ZdS )Cuda_array_declc                 C      dd }|S )Nc                 S   s   t | tjrt | tjsd S nt | tjtjfr$tdd | D r#d S nd S t| }t|}|d ur>|d ur@tj	||ddS d S d S )Nc                 S   s   g | ]	}t |tj qS  )
isinstancer   IntegerLiteral).0sr   r   b/var/www/html/eduruby.in/lip-sync/lip-sync-env/lib/python3.10/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                   @      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       r/   c                   @   r.   )Cuda_local_arrayN)r+   r,   r-   r   localr1   r2   r   r   r   r   r4   6   r3   r4   c                   @      e Zd ZejjZdd ZdS )Cuda_const_array_likec                 C   r   )Nc                 S   s   | S Nr   )Zndarrayr   r   r   r&   @      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_liker2   r)   r   r   r   r   r7   ;       r7   c                   @      e Zd ZejZeejgZ	dS )Cuda_threadfence_deviceN)
r+   r,   r-   r   Zthreadfencer2   r   r   nonecasesr   r   r   r   r=   E       r=   c                   @   r<   )Cuda_threadfence_blockN)
r+   r,   r-   r   Zthreadfence_blockr2   r   r   r>   r?   r   r   r   r   rA   K   r@   rA   c                   @   r<   )Cuda_threadfence_systemN)
r+   r,   r-   r   Zthreadfence_systemr2   r   r   r>   r?   r   r   r   r   rB   Q   r@   rB   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syncwarpr2   r   r   r>   i4r?   r   r   r   r   rC   W   s    rC   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_gridr2   r   r   r?   r   r   r   r   rE   ]   s    rE   c                   @       e Zd ZeejZdd ZdS )CudaCgModuleTemplatec                 C   
   t tS r8   )r   FunctionrE   r(   modr   r   r   resolve_this_gridg      
z&CudaCgModuleTemplate.resolve_this_gridN)	r+   r,   r-   r   Moduler   rF   r2   rM   r   r   r   r   rH   c       rH   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-   r2   r)   r   r   r   r   rQ   k   s    rQ   c                   @   s   e Zd ZeZdd ZdS )GridGroup_attrsc                 C   s   t ttS r8   )r   ZBoundFunctionrQ   r   rK   r   r   r   resolve_syncv      zGridGroup_attrs.resolve_syncN)r+   r,   r-   r   r2   rY   r   r   r   r   rX   r       rX   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_intrinsicr2   r   r   r#   rD   b1i8f4f8r?   r   r   r   r   r\   z   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_intrinsicr2   r   r   r#   rD   r]   r?   r   r   r   r   ra      s
    
ra   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_syncr2   r   r   rD   r^   r_   r`   r?   r   r   r   r   rb      s    rb   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_syncr2   r   r   r#   rD   r]   r^   r_   r`   r?   r   r   r   r   rc      s    rc   c                   @   r<   )Cuda_activemaskN)
r+   r,   r-   r   Z
activemaskr2   r   r   uint32r?   r   r   r   r   rd      r@   rd   c                   @   r<   )Cuda_lanemask_ltN)
r+   r,   r-   r   Zlanemask_ltr2   r   r   re   r?   r   r   r   r   rf      r@   rf   c                
   @   z   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popcr2   r   r   int8int16rR   int64uint8uint16re   uint64r?   r   r   r   r   rh          rh   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-   ri   r   fmar2   r   r   float32float64r?   r   r   r   r   rq      s    rq   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hfmar2   r   r   float16r?   r   r   r   r   ru      s    ru   c                   @   .   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cbrtr2   r   r   rs   rt   r?   r   r   r   r   ry      s
    ry   c                   @   rx   )	Cuda_brevN)r+   r,   r-   r   Zbrevr2   r   r   re   ro   r?   r   r   r   r   rz      s
    rz   c                
   @   rg   )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-   ri   r   Zclzr2   r   r   rj   rk   rR   rl   rm   rn   re   ro   r?   r   r   r   r   r{      rp   r{   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-   ri   r   Zffsr2   r   r   re   rj   rk   rR   rl   rm   rn   ro   r?   r   r   r   r   r|      rp   r|   c                   @   s   e Zd ZejZdd ZdS )	Cuda_selpc                 C   sX   |rJ |\}}}t jt jt jt jt jt jt jt jf}||ks#||vr%d S t	||||S r8   )
r   rt   rs   rk   rn   rR   re   rl   ro   r   )r(   rU   rV   testabsupported_typesr   r   r   r)     s   
zCuda_selp.genericN)r+   r,   r-   r   Zselpr2   r)   r   r   r   r   r}     s    r}   c                       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-   r2   r   r   rw   r?   r   l_keyr   r   Cuda_fp16_unary)  s    r   registerr   r   r   r   r   r   _genfp16_unary(     r   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   s8   |rJ t |dkr|d tjkrttjtjS d S d S )N   r   )lenr   rw   r   rT   r   r   r   r)   6  s   z8_genfp16_unary_operator.<locals>.Cuda_fp16_unary.genericNrW   r   r   r   r   r   2  r[   r   register_globalr   r   r   r   r   _genfp16_unary_operator1  s   r   c                    r   )Nc                       s$   e Zd Z ZeejejejgZdS )z)_genfp16_binary.<locals>.Cuda_fp16_binaryNr   r   r   r   r   Cuda_fp16_binary?  s    r   r   )r   r   r   r   r   _genfp16_binary>  r   r   c                   @   r   )Floatc                 C   s&   |rJ |\}|t jkrt||S d S r8   )r   rw   r   )r(   rU   rV   argr   r   r   r)   J  s
   

zFloat.genericNr*   r   r   r   r   r   G  s    r   c                    r   )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-   r2   r   r   r]   rw   r?   r   r   r   r   Cuda_fp16_cmpT  s    r   r   )r   r   r   r   r   _genfp16_binary_comparisonS  s   r   c                    s"   t  G  fdddt}|S )Nc                          e Zd Z ZfddZdS )z1_fp16_binary_operator.<locals>.Cuda_fp16_operatorc                    s   |rJ t |dkrM|d tjks|d tjkrO|d tjkr+| j|d |d }n| j|d |d }|tjksE|tjksE|tjkrQt	 tjtjS d S d S d S )N   r   r   )
r   r   rw   contextZcan_convertr   exactZpromotesafer   )r(   rU   rV   Zconvertible)rettyr   r   r)   p  s   



z9_fp16_binary_operator.<locals>.Cuda_fp16_operator.genericNrW   r   r   r   r   r   Cuda_fp16_operatorl      r   r   )r   r   r   r   r   r   _fp16_binary_operatork  s   r   c                 C      t | tjS r8   )r   r   r]   opr   r   r   _genfp16_comparison_operator  rZ   r   c                 C   r   r8   )r   r   rw   r   r   r   r   _genfp16_binary_operator  rZ   r   c                 C   s"   t d|  tjtjf}t|S NZ__numba_wrapper_r   r   rw   rJ   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                       r   )z_gen.<locals>.Cuda_atomicc                    s^   |rJ |\}}}|j  vrd S |jdkrt|j |tj|j S |jdkr-t|j |||j S d S Nr   )r    r!   r   r   intp)r(   rU   rV   aryidxval)r   r   r   r)     s   



z!_gen.<locals>.Cuda_atomic.genericNrW   r   r   r   r   r   Cuda_atomic  r   r   )r   r   )r   r   r   r   r   r   _gen  s   r   c                   @   r6   )Cuda_atomic_compare_and_swapc                 C   s@   |rJ |\}}}|j }|tv r|jdkrt||||S d S d S r   )r    integer_numba_typesr!   r   )r(   rU   rV   r   oldr   dtyr   r   r   r)      s   
z$Cuda_atomic_compare_and_swap.genericN)r+   r,   r-   r   atomicZcompare_and_swapr2   r)   r   r   r   r   r     r;   r   c                   @   r6   )Cuda_atomic_casc                 C   s`   |rJ |\}}}}|j }|tvrd S |jdkr!t||tj||S |jdkr.t|||||S d S r   )r    r   r!   r   r   r   )r(   rU   rV   r   r   r   r   r   r   r   r   r)     s   

zCuda_atomic_cas.genericN)r+   r,   r-   r   r   Zcasr2   r)   r   r   r   r   r   	  r;   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	nanosleepr2   r   r   voidre   r?   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      t jS r8   r   rR   rK   r   r   r   	resolve_x&     zDim3_attrs.resolve_xc                 C   r   r8   r   rK   r   r   r   	resolve_y)  r   zDim3_attrs.resolve_yc                 C   r   r8   r   rK   r   r   r   	resolve_z,  r   zDim3_attrs.resolve_zN)r+   r,   r-   r   r2   r   r   r   r   r   r   r   r   "  s
    r   c                   @   rG   )CudaSharedModuleTemplatec                 C   rI   r8   )r   rJ   r/   rK   r   r   r   resolve_array4  rN   z&CudaSharedModuleTemplate.resolve_arrayN)	r+   r,   r-   r   rO   r   r0   r2   r   r   r   r   r   r   0  rP   r   c                   @   rG   )CudaConstModuleTemplatec                 C   rI   r8   )r   rJ   r7   rK   r   r   r   resolve_array_like<  rN   z*CudaConstModuleTemplate.resolve_array_likeN)	r+   r,   r-   r   rO   r   r:   r2   r   r   r   r   r   r   8  rP   r   c                   @   rG   )CudaLocalModuleTemplatec                 C   rI   r8   )r   rJ   r4   rK   r   r   r   r   D  rN   z%CudaLocalModuleTemplate.resolve_arrayN)	r+   r,   r-   r   rO   r   r5   r2   r   r   r   r   r   r   @  rP   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   rI   r8   )r   rJ   Cuda_atomic_addrK   r   r   r   resolve_addL  rN   zCudaAtomicTemplate.resolve_addc                 C   rI   r8   )r   rJ   Cuda_atomic_subrK   r   r   r   resolve_subO  rN   zCudaAtomicTemplate.resolve_subc                 C   rI   r8   )r   rJ   Cuda_atomic_andrK   r   r   r   resolve_and_R  rN   zCudaAtomicTemplate.resolve_and_c                 C   rI   r8   )r   rJ   Cuda_atomic_orrK   r   r   r   resolve_or_U  rN   zCudaAtomicTemplate.resolve_or_c                 C   rI   r8   )r   rJ   Cuda_atomic_xorrK   r   r   r   resolve_xorX  rN   zCudaAtomicTemplate.resolve_xorc                 C   rI   r8   )r   rJ   Cuda_atomic_incrK   r   r   r   resolve_inc[  rN   zCudaAtomicTemplate.resolve_incc                 C   rI   r8   )r   rJ   Cuda_atomic_decrK   r   r   r   resolve_dec^  rN   zCudaAtomicTemplate.resolve_decc                 C   rI   r8   )r   rJ   Cuda_atomic_exchrK   r   r   r   resolve_excha  rN   zCudaAtomicTemplate.resolve_exchc                 C   rI   r8   )r   rJ   Cuda_atomic_maxrK   r   r   r   resolve_maxd  rN   zCudaAtomicTemplate.resolve_maxc                 C   rI   r8   )r   rJ   Cuda_atomic_minrK   r   r   r   resolve_ming  rN   zCudaAtomicTemplate.resolve_minc                 C   rI   r8   )r   rJ   Cuda_atomic_nanminrK   r   r   r   resolve_nanminj  rN   z!CudaAtomicTemplate.resolve_nanminc                 C   rI   r8   )r   rJ   Cuda_atomic_nanmaxrK   r   r   r   resolve_nanmaxm  rN   z!CudaAtomicTemplate.resolve_nanmaxc                 C   rI   r8   )r   rJ   r   rK   r   r   r   resolve_compare_and_swapp  rN   z+CudaAtomicTemplate.resolve_compare_and_swapc                 C   rI   r8   )r   rJ   r   rK   r   r   r   resolve_cass  rN   zCudaAtomicTemplate.resolve_casN)r+   r,   r-   r   rO   r   r   r2   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   rI   r8   )r   rJ   	Cuda_haddrK   r   r   r   resolve_hadd{  rN   zCudaFp16Template.resolve_haddc                 C   rI   r8   )r   rJ   	Cuda_hsubrK   r   r   r   resolve_hsub~  rN   zCudaFp16Template.resolve_hsubc                 C   rI   r8   )r   rJ   	Cuda_hmulrK   r   r   r   resolve_hmul  rN   zCudaFp16Template.resolve_hmulc                 C      t S r8   )hdiv_devicerK   r   r   r   resolve_hdiv  r9   zCudaFp16Template.resolve_hdivc                 C   rI   r8   )r   rJ   	Cuda_hnegrK   r   r   r   resolve_hneg  rN   zCudaFp16Template.resolve_hnegc                 C   rI   r8   )r   rJ   	Cuda_habsrK   r   r   r   resolve_habs  rN   zCudaFp16Template.resolve_habsc                 C   rI   r8   )r   rJ   ru   rK   r   r   r   resolve_hfma  rN   zCudaFp16Template.resolve_hfmac                 C   r   r8   )hsin_devicerK   r   r   r   resolve_hsin  r9   zCudaFp16Template.resolve_hsinc                 C   r   r8   )hcos_devicerK   r   r   r   resolve_hcos  r9   zCudaFp16Template.resolve_hcosc                 C   r   r8   )hlog_devicerK   r   r   r   resolve_hlog  r9   zCudaFp16Template.resolve_hlogc                 C   r   r8   )hlog10_devicerK   r   r   r   resolve_hlog10  r9   zCudaFp16Template.resolve_hlog10c                 C   r   r8   )hlog2_devicerK   r   r   r   resolve_hlog2  r9   zCudaFp16Template.resolve_hlog2c                 C   r   r8   )hexp_devicerK   r   r   r   resolve_hexp  r9   zCudaFp16Template.resolve_hexpc                 C   r   r8   )hexp10_devicerK   r   r   r   resolve_hexp10  r9   zCudaFp16Template.resolve_hexp10c                 C   r   r8   )hexp2_devicerK   r   r   r   resolve_hexp2  r9   zCudaFp16Template.resolve_hexp2c                 C   r   r8   )hfloor_devicerK   r   r   r   resolve_hfloor  r9   zCudaFp16Template.resolve_hfloorc                 C   r   r8   )hceil_devicerK   r   r   r   resolve_hceil  r9   zCudaFp16Template.resolve_hceilc                 C   r   r8   )hsqrt_devicerK   r   r   r   resolve_hsqrt  r9   zCudaFp16Template.resolve_hsqrtc                 C   r   r8   )hrsqrt_devicerK   r   r   r   resolve_hrsqrt  r9   zCudaFp16Template.resolve_hrsqrtc                 C   r   r8   )hrcp_devicerK   r   r   r   resolve_hrcp  r9   zCudaFp16Template.resolve_hrcpc                 C   r   r8   )hrint_devicerK   r   r   r   resolve_hrint  r9   zCudaFp16Template.resolve_hrintc                 C   r   r8   )htrunc_devicerK   r   r   r   resolve_htrunc  r9   zCudaFp16Template.resolve_htruncc                 C   rI   r8   )r   rJ   Cuda_heqrK   r   r   r   resolve_heq  rN   zCudaFp16Template.resolve_heqc                 C   rI   r8   )r   rJ   Cuda_hnerK   r   r   r   resolve_hne  rN   zCudaFp16Template.resolve_hnec                 C   rI   r8   )r   rJ   Cuda_hgerK   r   r   r   resolve_hge  rN   zCudaFp16Template.resolve_hgec                 C   rI   r8   )r   rJ   Cuda_hgtrK   r   r   r   resolve_hgt  rN   zCudaFp16Template.resolve_hgtc                 C   rI   r8   )r   rJ   Cuda_hlerK   r   r   r   resolve_hle  rN   zCudaFp16Template.resolve_hlec                 C   rI   r8   )r   rJ   Cuda_hltrK   r   r   r   resolve_hlt  rN   zCudaFp16Template.resolve_hltc                 C   rI   r8   )r   rJ   	Cuda_hmaxrK   r   r   r   resolve_hmax  rN   zCudaFp16Template.resolve_hmaxc                 C   rI   r8   )r   rJ   	Cuda_hminrK   r   r   r   resolve_hmin  rN   zCudaFp16Template.resolve_hminN)&r+   r,   r-   r   rO   r   rv   r2   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      t tjS r8   )r   rO   r   rF   rK   r   r   r   
resolve_cg  rZ   zCudaModuleTemplate.resolve_cgc                 C   r   r8   r   rK   r   r   r   resolve_threadIdx  r9   z$CudaModuleTemplate.resolve_threadIdxc                 C   r   r8   r$  rK   r   r   r   resolve_blockIdx  r9   z#CudaModuleTemplate.resolve_blockIdxc                 C   r   r8   r$  rK   r   r   r   resolve_blockDim  r9   z#CudaModuleTemplate.resolve_blockDimc                 C   r   r8   r$  rK   r   r   r   resolve_gridDim  r9   z"CudaModuleTemplate.resolve_gridDimc                 C   r   r8   r   rK   r   r   r   resolve_laneid  r   z!CudaModuleTemplate.resolve_laneidc                 C   r"  r8   )r   rO   r   r0   rK   r   r   r   resolve_shared  rZ   z!CudaModuleTemplate.resolve_sharedc                 C   rI   r8   )r   rJ   rh   rK   r   r   r   resolve_popc  rN   zCudaModuleTemplate.resolve_popcc                 C   rI   r8   )r   rJ   rz   rK   r   r   r   resolve_brev  rN   zCudaModuleTemplate.resolve_brevc                 C   rI   r8   )r   rJ   r{   rK   r   r   r   resolve_clz  rN   zCudaModuleTemplate.resolve_clzc                 C   rI   r8   )r   rJ   r|   rK   r   r   r   resolve_ffs  rN   zCudaModuleTemplate.resolve_ffsc                 C   rI   r8   )r   rJ   rq   rK   r   r   r   resolve_fma  rN   zCudaModuleTemplate.resolve_fmac                 C   rI   r8   )r   rJ   ry   rK   r   r   r   resolve_cbrt  rN   zCudaModuleTemplate.resolve_cbrtc                 C   rI   r8   )r   rJ   r=   rK   r   r   r   resolve_threadfence  rN   z&CudaModuleTemplate.resolve_threadfencec                 C   rI   r8   )r   rJ   rA   rK   r   r   r   resolve_threadfence_block  rN   z,CudaModuleTemplate.resolve_threadfence_blockc                 C   rI   r8   )r   rJ   rB   rK   r   r   r   resolve_threadfence_system  rN   z-CudaModuleTemplate.resolve_threadfence_systemc                 C   rI   r8   )r   rJ   rC   rK   r   r   r   resolve_syncwarp
  rN   z#CudaModuleTemplate.resolve_syncwarpc                 C   rI   r8   )r   rJ   r\   rK   r   r   r   resolve_shfl_sync_intrinsic  rN   z.CudaModuleTemplate.resolve_shfl_sync_intrinsicc                 C   rI   r8   )r   rJ   ra   rK   r   r   r   resolve_vote_sync_intrinsic  rN   z.CudaModuleTemplate.resolve_vote_sync_intrinsicc                 C   rI   r8   )r   rJ   rb   rK   r   r   r   resolve_match_any_sync  rN   z)CudaModuleTemplate.resolve_match_any_syncc                 C   rI   r8   )r   rJ   rc   rK   r   r   r   resolve_match_all_sync  rN   z)CudaModuleTemplate.resolve_match_all_syncc                 C   rI   r8   )r   rJ   rd   rK   r   r   r   resolve_activemask  rN   z%CudaModuleTemplate.resolve_activemaskc                 C   rI   r8   )r   rJ   rf   rK   r   r   r   resolve_lanemask_lt  rN   z&CudaModuleTemplate.resolve_lanemask_ltc                 C   rI   r8   )r   rJ   r}   rK   r   r   r   resolve_selp  rN   zCudaModuleTemplate.resolve_selpc                 C   rI   r8   )r   rJ   r   rK   r   r   r   resolve_nanosleep"  rN   z$CudaModuleTemplate.resolve_nanosleepc                 C   r"  r8   )r   rO   r   r   rK   r   r   r   resolve_atomic%  rZ   z!CudaModuleTemplate.resolve_atomicc                 C   r"  r8   )r   rO   r   rv   rK   r   r   r   resolve_fp16(  rZ   zCudaModuleTemplate.resolve_fp16c                 C   r"  r8   )r   rO   r   r:   rK   r   r   r   resolve_const+  rZ   z CudaModuleTemplate.resolve_constc                 C   r"  r8   )r   rO   r   r5   rK   r   r   r   resolve_local.  rZ   z CudaModuleTemplate.resolve_localN)$r+   r,   r-   r   rO   r   r2   r#  r%  r&  r'  r(  r)  r*  r+  r,  r-  r.  r/  r0  r1  r2  r3  r4  r5  r6  r7  r8  r9  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   registryr   Zregister_attrr   r   r/   r4   r7   r=   rA   rB   rC   rE   rH   rQ   rX   r\   ra   rb   rc   rd   rf   rh   rq   ru   ry   rz   r{   r|   r}   r   r   r   floatr   r   r   r   r   rv   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   rt   rs   rR   re   rl   ro   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!  rO   funcr   r   r   r   <module>   sP   $ 	

			














.^[