o
    TdV[                     @   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 qCd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                   @      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   f/home/ncw/WWW/www-new/content/articles/pi-bbp/venv/lib/python3.10/site-packages/numba/cuda/cudadecl.py
<listcomp>!   s    z:Cuda_array_decl.generic.<locals>.typer.<locals>.<listcomp>C)dtypendimlayout)
r   r   Integerr   TupleUniTupleanyr   r   Array)shaper   r   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   r1   /       r1   c                   @   r0   )Cuda_local_arrayN)r-   r.   r/   r   localr3   r4   r   r   r   r   r6   4   r5   r6   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   )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   const
array_liker4   r+   r   r   r   r   r9   9       r9   c                   @      e Zd ZejZeejgZ	dS )Cuda_threadfence_deviceN)
r-   r.   r/   r   threadfencer4   r   r   nonecasesr   r   r   r   rA   C       rA   c                   @   r@   )Cuda_threadfence_blockN)
r-   r.   r/   r   threadfence_blockr4   r   r   rC   rD   r   r   r   r   rF   I   rE   rF   c                   @   r@   )Cuda_threadfence_systemN)
r-   r.   r/   r   threadfence_systemr4   r   r   rC   rD   r   r   r   r   rH   O   rE   rH   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   syncwarpr4   r   r   rC   i4rD   r   r   r   r   rJ   U   s    rJ   c                   @   s   e Zd ZejjZeegZ	dS )Cuda_cg_this_gridN)
r-   r.   r/   r   cg	this_gridr4   r   r   rD   r   r   r   r   rM   [   s    rM   c                   @       e Zd ZeejZdd ZdS )CudaCgModuleTemplatec                 C   
   t tS r:   )r   FunctionrM   r*   modr   r   r   resolve_this_gride      
z&CudaCgModuleTemplate.resolve_this_gridN)	r-   r.   r/   r   Moduler   rN   r4   rV   r   r   r   r   rQ   a       rQ   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)recvr)r   r   int32thisr*   argskwsr   r   r   r+   l   s   zCuda_grid_group_sync.genericNr-   r.   r/   r4   r+   r   r   r   r   rZ   i   s    rZ   c                   @   s   e Zd ZeZdd ZdS )GridGroup_attrsc                 C   s   t ttS r:   )r   BoundFunctionrZ   r   rT   r   r   r   resolve_synct      zGridGroup_attrs.resolve_syncN)r-   r.   r/   r   r4   rd   r   r   r   r   rb   p       rb   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   shfl_sync_intrinsicr4   r   r   r"   rL   b1i8f4f8rD   r   r   r   r   rg   x   s    rg   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   vote_sync_intrinsicr4   r   r   r"   rL   ri   rD   r   r   r   r   rm      s
    
rm   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   match_any_syncr4   r   r   rL   rj   rk   rl   rD   r   r   r   r   ro      s    ro   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   match_all_syncr4   r   r   r"   rL   ri   rj   rk   rl   rD   r   r   r   r   rq      s    rq   c                   @   r@   )Cuda_activemaskN)
r-   r.   r/   r   
activemaskr4   r   r   uint32rD   r   r   r   r   rs      rE   rs   c                   @   r@   )Cuda_lanemask_ltN)
r-   r.   r/   r   lanemask_ltr4   r   r   ru   rD   r   r   r   r   rv      rE   rv   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   popcr4   r   r   int8int16r\   int64uint8uint16ru   uint64rD   r   r   r   r   ry          ry   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/   rz   r   fmar4   r   r   float32float64rD   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   fp16hfmar4   r   r   float16rD   r   r   r   r   r      s    r   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   cbrtr4   r   r   r   r   rD   r   r   r   r   r      s
    r   c                   @   r   )	Cuda_brevN)r-   r.   r/   r   brevr4   r   r   ru   r   rD   r   r   r   r   r      s
    r   c                
   @   rx   )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/   rz   r   clzr4   r   r   r|   r}   r\   r~   r   r   ru   r   rD   r   r   r   r   r      r   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/   rz   r   ffsr4   r   r   ru   r|   r}   r\   r~   r   r   r   rD   r   r   r   r   r      r   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 r:   )
r   r   r   r}   r   r\   ru   r~   r   r   )r*   r_   r`   testabsupported_typesr   r   r   r+     s   
zCuda_selp.genericN)r-   r.   r/   r   selpr4   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/   r4   r   r   r   rD   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   r   r   r^   r   r   r   r+   4  s   z8_genfp16_unary_operator.<locals>.Cuda_fp16_unary.genericNra   r   r   r   r   r   0  rf   r   register_globalr
   r   r   r   r   _genfp16_unary_operator/  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 r:   )r   r   r   )r*   r_   r`   argr   r   r   r+   H  s
   

zFloat.genericNr,   r   r   r   r   r   E  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/   r4   r   r   ri   r   rD   r   r   r   r   Cuda_fp16_cmpR  s    r   r   )r   r   r   r   r   _genfp16_binary_comparisonQ  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   r   contextcan_convertr   exactpromotesafer   )r*   r_   r`   convertible)rettyr   r   r+   n  s"   

z9_fp16_binary_operator.<locals>.Cuda_fp16_operator.genericNra   r   r   r   r   r   Cuda_fp16_operatorj      r   r   )r   r   r   r   r   r   _fp16_binary_operatori  s   r   c                 C      t | tjS r:   )r   r   ri   opr   r   r   _genfp16_comparison_operator  re   r   c                 C   r   r:   )r   r   r   r   r   r   r   _genfp16_binary_operator  re   r   c                 C   s"   t d|  tjtjf}t|S N__numba_wrapper_r   r   r   rS   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   hsinhcoshloghlog10hlog2hexphexp10hexp2hsqrthrsqrthfloorhceilhrcphrinthtrunc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*   r_   r`   aryidxval)r   r   r   r+     s   



z!_gen.<locals>.Cuda_atomic.genericNra   r   r   r   r   r   Cuda_atomic  r   r   )r   r
   )r   r   r   r   r   r   _gen  s   r   c                   @   r8   )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*   r_   r`   r   oldr   dtyr   r   r   r+     s   
z$Cuda_atomic_compare_and_swap.genericN)r-   r.   r/   r   atomiccompare_and_swapr4   r+   r   r   r   r   r     r?   r   c                   @   r8   )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*   r_   r`   r   r   r   r   r   r   r   r   r+     s   

zCuda_atomic_cas.genericN)r-   r.   r/   r   r   casr4   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   	nanosleepr4   r   r   voidru   rD   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 r:   r   r\   rT   r   r   r   	resolve_x$     zDim3_attrs.resolve_xc                 C   r   r:   r   rT   r   r   r   	resolve_y'  r   zDim3_attrs.resolve_yc                 C   r   r:   r   rT   r   r   r   	resolve_z*  r   zDim3_attrs.resolve_zN)r-   r.   r/   r   r4   r   r   r   r   r   r   r   r      s
    r   c                   @   rP   )CudaSharedModuleTemplatec                 C   rR   r:   )r   rS   r1   rT   r   r   r   resolve_array2  rW   z&CudaSharedModuleTemplate.resolve_arrayN)	r-   r.   r/   r   rX   r   r2   r4   r   r   r   r   r   r   .  rY   r   c                   @   rP   )CudaConstModuleTemplatec                 C   rR   r:   )r   rS   r9   rT   r   r   r   resolve_array_like:  rW   z*CudaConstModuleTemplate.resolve_array_likeN)	r-   r.   r/   r   rX   r   r=   r4   r   r   r   r   r   r   6  rY   r   c                   @   rP   )CudaLocalModuleTemplatec                 C   rR   r:   )r   rS   r6   rT   r   r   r   r   B  rW   z%CudaLocalModuleTemplate.resolve_arrayN)	r-   r.   r/   r   rX   r   r7   r4   r   r   r   r   r   r   >  rY   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   rR   r:   )r   rS   Cuda_atomic_addrT   r   r   r   resolve_addJ  rW   zCudaAtomicTemplate.resolve_addc                 C   rR   r:   )r   rS   Cuda_atomic_subrT   r   r   r   resolve_subM  rW   zCudaAtomicTemplate.resolve_subc                 C   rR   r:   )r   rS   Cuda_atomic_andrT   r   r   r   resolve_and_P  rW   zCudaAtomicTemplate.resolve_and_c                 C   rR   r:   )r   rS   Cuda_atomic_orrT   r   r   r   resolve_or_S  rW   zCudaAtomicTemplate.resolve_or_c                 C   rR   r:   )r   rS   Cuda_atomic_xorrT   r   r   r   resolve_xorV  rW   zCudaAtomicTemplate.resolve_xorc                 C   rR   r:   )r   rS   Cuda_atomic_incrT   r   r   r   resolve_incY  rW   zCudaAtomicTemplate.resolve_incc                 C   rR   r:   )r   rS   Cuda_atomic_decrT   r   r   r   resolve_dec\  rW   zCudaAtomicTemplate.resolve_decc                 C   rR   r:   )r   rS   Cuda_atomic_exchrT   r   r   r   resolve_exch_  rW   zCudaAtomicTemplate.resolve_exchc                 C   rR   r:   )r   rS   Cuda_atomic_maxrT   r   r   r   resolve_maxb  rW   zCudaAtomicTemplate.resolve_maxc                 C   rR   r:   )r   rS   Cuda_atomic_minrT   r   r   r   resolve_mine  rW   zCudaAtomicTemplate.resolve_minc                 C   rR   r:   )r   rS   Cuda_atomic_nanminrT   r   r   r   resolve_nanminh  rW   z!CudaAtomicTemplate.resolve_nanminc                 C   rR   r:   )r   rS   Cuda_atomic_nanmaxrT   r   r   r   resolve_nanmaxk  rW   z!CudaAtomicTemplate.resolve_nanmaxc                 C   rR   r:   )r   rS   r   rT   r   r   r   resolve_compare_and_swapn  rW   z+CudaAtomicTemplate.resolve_compare_and_swapc                 C   rR   r:   )r   rS   r   rT   r   r   r   resolve_casq  rW   zCudaAtomicTemplate.resolve_casN)r-   r.   r/   r   rX   r   r   r4   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   rR   r:   )r   rS   	Cuda_haddrT   r   r   r   resolve_haddy  rW   zCudaFp16Template.resolve_haddc                 C   rR   r:   )r   rS   	Cuda_hsubrT   r   r   r   resolve_hsub|  rW   zCudaFp16Template.resolve_hsubc                 C   rR   r:   )r   rS   	Cuda_hmulrT   r   r   r   resolve_hmul  rW   zCudaFp16Template.resolve_hmulc                 C      t S r:   )hdiv_devicerT   r   r   r   resolve_hdiv  r<   zCudaFp16Template.resolve_hdivc                 C   rR   r:   )r   rS   	Cuda_hnegrT   r   r   r   resolve_hneg  rW   zCudaFp16Template.resolve_hnegc                 C   rR   r:   )r   rS   	Cuda_habsrT   r   r   r   resolve_habs  rW   zCudaFp16Template.resolve_habsc                 C   rR   r:   )r   rS   r   rT   r   r   r   resolve_hfma  rW   zCudaFp16Template.resolve_hfmac                 C   r  r:   )hsin_devicerT   r   r   r   resolve_hsin  r<   zCudaFp16Template.resolve_hsinc                 C   r  r:   )hcos_devicerT   r   r   r   resolve_hcos  r<   zCudaFp16Template.resolve_hcosc                 C   r  r:   )hlog_devicerT   r   r   r   resolve_hlog  r<   zCudaFp16Template.resolve_hlogc                 C   r  r:   )hlog10_devicerT   r   r   r   resolve_hlog10  r<   zCudaFp16Template.resolve_hlog10c                 C   r  r:   )hlog2_devicerT   r   r   r   resolve_hlog2  r<   zCudaFp16Template.resolve_hlog2c                 C   r  r:   )hexp_devicerT   r   r   r   resolve_hexp  r<   zCudaFp16Template.resolve_hexpc                 C   r  r:   )hexp10_devicerT   r   r   r   resolve_hexp10  r<   zCudaFp16Template.resolve_hexp10c                 C   r  r:   )hexp2_devicerT   r   r   r   resolve_hexp2  r<   zCudaFp16Template.resolve_hexp2c                 C   r  r:   )hfloor_devicerT   r   r   r   resolve_hfloor  r<   zCudaFp16Template.resolve_hfloorc                 C   r  r:   )hceil_devicerT   r   r   r   resolve_hceil  r<   zCudaFp16Template.resolve_hceilc                 C   r  r:   )hsqrt_devicerT   r   r   r   resolve_hsqrt  r<   zCudaFp16Template.resolve_hsqrtc                 C   r  r:   )hrsqrt_devicerT   r   r   r   resolve_hrsqrt  r<   zCudaFp16Template.resolve_hrsqrtc                 C   r  r:   )hrcp_devicerT   r   r   r   resolve_hrcp  r<   zCudaFp16Template.resolve_hrcpc                 C   r  r:   )hrint_devicerT   r   r   r   resolve_hrint  r<   zCudaFp16Template.resolve_hrintc                 C   r  r:   )htrunc_devicerT   r   r   r   resolve_htrunc  r<   zCudaFp16Template.resolve_htruncc                 C   rR   r:   )r   rS   Cuda_heqrT   r   r   r   resolve_heq  rW   zCudaFp16Template.resolve_heqc                 C   rR   r:   )r   rS   Cuda_hnerT   r   r   r   resolve_hne  rW   zCudaFp16Template.resolve_hnec                 C   rR   r:   )r   rS   Cuda_hgerT   r   r   r   resolve_hge  rW   zCudaFp16Template.resolve_hgec                 C   rR   r:   )r   rS   Cuda_hgtrT   r   r   r   resolve_hgt  rW   zCudaFp16Template.resolve_hgtc                 C   rR   r:   )r   rS   Cuda_hlerT   r   r   r   resolve_hle  rW   zCudaFp16Template.resolve_hlec                 C   rR   r:   )r   rS   Cuda_hltrT   r   r   r   resolve_hlt  rW   zCudaFp16Template.resolve_hltc                 C   rR   r:   )r   rS   	Cuda_hmaxrT   r   r   r   resolve_hmax  rW   zCudaFp16Template.resolve_hmaxc                 C   rR   r:   )r   rS   	Cuda_hminrT   r   r   r   resolve_hmin  rW   zCudaFp16Template.resolve_hminN)&r-   r.   r/   r   rX   r   r   r4   r  r  r  r  r  r   r!  r#  r%  r'  r)  r+  r-  r/  r1  r3  r5  r7  r9  r;  r=  r?  rA  rC  rE  rG  rI  rK  rM  rO  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      t tjS r:   )r   rX   r   rN   rT   r   r   r   
resolve_cg  re   zCudaModuleTemplate.resolve_cgc                 C   r  r:   r   rT   r   r   r   resolve_threadIdx  r<   z$CudaModuleTemplate.resolve_threadIdxc                 C   r  r:   rS  rT   r   r   r   resolve_blockIdx  r<   z#CudaModuleTemplate.resolve_blockIdxc                 C   r  r:   rS  rT   r   r   r   resolve_blockDim  r<   z#CudaModuleTemplate.resolve_blockDimc                 C   r  r:   rS  rT   r   r   r   resolve_gridDim  r<   z"CudaModuleTemplate.resolve_gridDimc                 C   r   r:   r   rT   r   r   r   resolve_laneid  r   z!CudaModuleTemplate.resolve_laneidc                 C   rQ  r:   )r   rX   r   r2   rT   r   r   r   resolve_shared  re   z!CudaModuleTemplate.resolve_sharedc                 C   rR   r:   )r   rS   ry   rT   r   r   r   resolve_popc  rW   zCudaModuleTemplate.resolve_popcc                 C   rR   r:   )r   rS   r   rT   r   r   r   resolve_brev  rW   zCudaModuleTemplate.resolve_brevc                 C   rR   r:   )r   rS   r   rT   r   r   r   resolve_clz  rW   zCudaModuleTemplate.resolve_clzc                 C   rR   r:   )r   rS   r   rT   r   r   r   resolve_ffs  rW   zCudaModuleTemplate.resolve_ffsc                 C   rR   r:   )r   rS   r   rT   r   r   r   resolve_fma  rW   zCudaModuleTemplate.resolve_fmac                 C   rR   r:   )r   rS   r   rT   r   r   r   resolve_cbrt  rW   zCudaModuleTemplate.resolve_cbrtc                 C   rR   r:   )r   rS   rA   rT   r   r   r   resolve_threadfence  rW   z&CudaModuleTemplate.resolve_threadfencec                 C   rR   r:   )r   rS   rF   rT   r   r   r   resolve_threadfence_block  rW   z,CudaModuleTemplate.resolve_threadfence_blockc                 C   rR   r:   )r   rS   rH   rT   r   r   r   resolve_threadfence_system  rW   z-CudaModuleTemplate.resolve_threadfence_systemc                 C   rR   r:   )r   rS   rJ   rT   r   r   r   resolve_syncwarp  rW   z#CudaModuleTemplate.resolve_syncwarpc                 C   rR   r:   )r   rS   rg   rT   r   r   r   resolve_shfl_sync_intrinsic  rW   z.CudaModuleTemplate.resolve_shfl_sync_intrinsicc                 C   rR   r:   )r   rS   rm   rT   r   r   r   resolve_vote_sync_intrinsic  rW   z.CudaModuleTemplate.resolve_vote_sync_intrinsicc                 C   rR   r:   )r   rS   ro   rT   r   r   r   resolve_match_any_sync  rW   z)CudaModuleTemplate.resolve_match_any_syncc                 C   rR   r:   )r   rS   rq   rT   r   r   r   resolve_match_all_sync  rW   z)CudaModuleTemplate.resolve_match_all_syncc                 C   rR   r:   )r   rS   rs   rT   r   r   r   resolve_activemask  rW   z%CudaModuleTemplate.resolve_activemaskc                 C   rR   r:   )r   rS   rv   rT   r   r   r   resolve_lanemask_lt  rW   z&CudaModuleTemplate.resolve_lanemask_ltc                 C   rR   r:   )r   rS   r   rT   r   r   r   resolve_selp  rW   zCudaModuleTemplate.resolve_selpc                 C   rR   r:   )r   rS   r   rT   r   r   r   resolve_nanosleep   rW   z$CudaModuleTemplate.resolve_nanosleepc                 C   rQ  r:   )r   rX   r   r   rT   r   r   r   resolve_atomic#  re   z!CudaModuleTemplate.resolve_atomicc                 C   rQ  r:   )r   rX   r   r   rT   r   r   r   resolve_fp16&  re   zCudaModuleTemplate.resolve_fp16c                 C   rQ  r:   )r   rX   r   r=   rT   r   r   r   resolve_const)  re   z CudaModuleTemplate.resolve_constc                 C   rQ  r:   )r   rX   r   r7   rT   r   r   r   resolve_local,  re   z CudaModuleTemplate.resolve_localN)$r-   r.   r/   r   rX   r   r4   rR  rT  rU  rV  rW  rX  rY  rZ  r[  r\  r]  r^  r_  r`  ra  rb  rc  rd  re  rf  rg  rh  ri  rj  rk  rl  rm  rn  ro  r   r   r   r   rP    s>    
rP  )operator
numba.corer   numba.core.typing.npydeclr   r   r   r   r   numba.core.typing.templatesr   r	   r
   r   r   r   numba.cuda.typesr   r   numba.core.typeconvr   numbar   numba.cuda.compilerr   registryr   register_attrr   r   r1   r6   r9   rA   rF   rH   rJ   rM   rQ   rZ   rb   rg   rm   ro   rq   rs   rv   ry   r   r   r   r   r   r   r   r   r   r   floatr   r   r   r   r   r   haddr  addCuda_addiadd	Cuda_iaddhsubr  subCuda_subisub	Cuda_isubhmulr  mulCuda_mulimul	Cuda_imulhmaxrL  hminrN  hnegr  negCuda_neghabsr  absCuda_absheqr@  eqhnerB  nehgerD  gehgtrF  gthlerH  lehltrJ  lttruedivitruedivr   r   r"  r$  r&  r(  r*  r,  r.  r0  r6  r8  r2  r4  r:  r<  r>  r  r   r   r   r\   ru   r~   r   all_numba_typesr   unsigned_int_numba_typesr   r   r   maxr	  minr  nanmaxr  nanminr  and_r   or_r   xorr  incr  decr  exchr  r   r   r   r   r   r   r   r   r  rP  rX   funcr   r   r   r   <module>   sH    	

			














.^[