a
    DfY                     @   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 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'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Z0eG d/d0 d0eZ1eG d1d2 d2eZ2eG d3d4 d4eZ3d5d6 Z4d7d8 Z5d9d: Z6ee7G d;d< d<eZ8d=d> Z9d?d@ Z:dAdB Z;dCdD Z<e6ej=j>Z?e<e j@ZAe<e jBZCe6ej=jDZEe<e jFZGe<e jHZIe6ej=jJZKe<e jLZMe<e jNZOe6ej=jPZQe6ej=jRZSe4ej=jTZUe5e jVZWe4ej=jXZYe5eZZ[e9ej=j\Z]e;e j^ e9ej=j_Z`e;e ja e9ej=jbZce;e jd e9ej=jeZfe;e jg e9ej=jhZie;e jj e9ej=jkZle;e jm e<e jn e<e jo dEdF ZpdGdH ZqepdIZrepdJZsepdKZtepdLZuepdMZvepdNZwepdOZxepdPZyepdQZzepdRZ{epdSZ|epdTZ}epdUZ~epdVZepdWZeqdXZdYdZ ZejejejejejejfZejejejejfZejejfZeejj@eZeejjFeZeejjeZeejjeZeejjeZeejjeZeejjeZeejjeZeejjeZeejjeZeejjeZeejjeZeG d[d\ d\eZeG d]d^ d^eZeG d_d` d`eZeG dadb dbe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eee eD ]Zeee q>e	D ]Zeee qTe
D ]Zeee qjdS )o    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)
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 ur||d ur|tj	||ddS d S )Nc                 S   s   g | ]}t |tj qS  )
isinstancer   IntegerLiteral).0sr   r   `/nfs/NAS7/SABIOD/METHODE/ermites/ermites_venv/lib/python3.9/site-packages/numba/cuda/cudadecl.py
<listcomp>#   s   z:Cuda_array_decl.generic.<locals>.typer.<locals>.<listcomp>C)dtypendimZlayout)
r   r   ZIntegerr   TupleZUniTupleanyr   r   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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   r<   ]   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_intrinsicr/   r   r   r    r;   r=   r7   r   r   r   r   rA   l   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_syncr/   r   r   r;   r>   r?   r@   r7   r   r   r   r   rB   s   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_syncr/   r   r   r    r;   r=   r>   r?   r@   r7   r   r   r   r   rC   ~   s   rC   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   rD      s   rD   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   rE   r7   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_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int16int32int64uint8uint16rE   uint64r7   r   r   r   r   rG      s   rG   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+   rH   r   fmar/   r   r   float32float64r7   r   r   r   r   rP      s
   rP   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   rT      s   rT   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   rR   rS   r7   r   r   r   r   rW      s   rW   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   rE   rO   r7   r   r   r   r   rX      s   rX   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+   rH   r   Zclzr/   r   r   rI   rJ   rK   rL   rM   rN   rE   rO   r7   r   r   r   r   rY      s   rY   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+   rH   r   Zffsr/   r   r   rE   rI   rJ   rK   rL   rM   rN   rO   r7   r   r   r   r   rZ      s   rZ   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sF||vrJd S t	||||S r3   )
r   rS   rR   rJ   rN   rK   rE   rL   rO   r   )r&   argskwstestabsupported_typesr   r   r   r'      s    
zCuda_selp.genericN)r)   r*   r+   r   Zselpr/   r'   r   r   r   r   r[      s   r[   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   rV   r7   r   l_keyr   r   Cuda_fp16_unary  s   re   registerr   rd   re   r   rc   r   _genfp16_unary  s    ri   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J t |dkr0|d tjkr0ttjtjS d S )N   r   )lenr   rV   r   )r&   r\   r]   r   r   r   r'     s    z8_genfp16_unary_operator.<locals>.Cuda_fp16_unary.genericNr)   r*   r+   r/   r'   r   rc   r   r   re     s   re   register_globalr   rh   r   rc   r   _genfp16_unary_operator  s    ro   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_binaryNrb   r   rc   r   r   Cuda_fp16_binary"  s   rp   rf   )rd   rp   r   rc   r   _genfp16_binary!  s    rq   c                   @   s   e Zd Zdd ZdS )Floatc                 C   s&   |rJ |\}|t jkr"t||S d S r3   )r   rV   r   )r&   r\   r]   argr   r   r   r'   -  s    
zFloat.genericNr(   r   r   r   r   rr   *  s   rr   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   r=   rV   r7   r   rc   r   r   Cuda_fp16_cmp7  s   rt   rf   )rd   rt   r   rc   r   _genfp16_binary_comparison6  s    ru   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J 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   rj   )
rk   r   rV   contextZcan_convertr   exactZpromotesafer   )r&   r\   r]   Zconvertible)rettyr   r   r'   S  s    

z9_fp16_binary_operator.<locals>.Cuda_fp16_operator.genericNrl   r   rd   rz   r   r   Cuda_fp16_operatorO  s   r|   rm   )rd   rz   r|   r   r{   r   _fp16_binary_operatorN  s    r}   c                 C   s   t | tjS r3   )r}   r   r=   opr   r   r   _genfp16_comparison_operatorm  s    r   c                 C   s   t | tjS r3   )r}   r   rV   r~   r   r   r   _genfp16_binary_operatorq  s    r   c                 C   s"   t d|  tjtjf}t|S NZ__numba_wrapper_r   r   rV   Function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J |\}}}|j  vr d S |jdkr>t|j |tj|j S |jdkrZt|j |||j S d S Nrj   )r   r   r   r   intp)r&   r\   r]   aryidxval)ra   r   r   r'     s    



z!_gen.<locals>.Cuda_atomic.genericNrl   r   rd   ra   r   r   Cuda_atomic  s   r   )rg   r   )rd   ra   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J |\}}}|j }|tv r8|jdkr8t||||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   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J |\}}}}|j }|tvr&d S |jdkrBt||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   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   voidrE   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   rK   r&   modr   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_array  s    z&CudaSharedModuleTemplate.resolve_arrayN)	r)   r*   r+   r   Moduler   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 r3   )r   r   r2   r   r   r   r   resolve_array_like  s    z*CudaConstModuleTemplate.resolve_array_likeN)	r)   r*   r+   r   r   r   r4   r/   r   r   r   r   r   r     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   '  s    z%CudaLocalModuleTemplate.resolve_arrayN)	r)   r*   r+   r   r   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_add/  s    zCudaAtomicTemplate.resolve_addc                 C   s
   t tS r3   )r   r   Cuda_atomic_subr   r   r   r   resolve_sub2  s    zCudaAtomicTemplate.resolve_subc                 C   s
   t tS r3   )r   r   Cuda_atomic_andr   r   r   r   resolve_and_5  s    zCudaAtomicTemplate.resolve_and_c                 C   s
   t tS r3   )r   r   Cuda_atomic_orr   r   r   r   resolve_or_8  s    zCudaAtomicTemplate.resolve_or_c                 C   s
   t tS r3   )r   r   Cuda_atomic_xorr   r   r   r   resolve_xor;  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_decA  s    zCudaAtomicTemplate.resolve_decc                 C   s
   t tS r3   )r   r   Cuda_atomic_exchr   r   r   r   resolve_exchD  s    zCudaAtomicTemplate.resolve_exchc                 C   s
   t tS r3   )r   r   Cuda_atomic_maxr   r   r   r   resolve_maxG  s    zCudaAtomicTemplate.resolve_maxc                 C   s
   t tS r3   )r   r   Cuda_atomic_minr   r   r   r   resolve_minJ  s    zCudaAtomicTemplate.resolve_minc                 C   s
   t tS r3   )r   r   Cuda_atomic_nanminr   r   r   r   resolve_nanminM  s    z!CudaAtomicTemplate.resolve_nanminc                 C   s
   t tS r3   )r   r   Cuda_atomic_nanmaxr   r   r   r   resolve_nanmaxP  s    z!CudaAtomicTemplate.resolve_nanmaxc                 C   s
   t tS r3   )r   r   r   r   r   r   r   resolve_compare_and_swapS  s    z+CudaAtomicTemplate.resolve_compare_and_swapc                 C   s
   t tS r3   )r   r   r   r   r   r   r   resolve_casV  s    zCudaAtomicTemplate.resolve_casN)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   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_hsuba  s    zCudaFp16Template.resolve_hsubc                 C   s
   t tS r3   )r   r   	Cuda_hmulr   r   r   r   resolve_hmuld  s    zCudaFp16Template.resolve_hmulc                 C   s   t S r3   )hdiv_devicer   r   r   r   resolve_hdivg  s    zCudaFp16Template.resolve_hdivc                 C   s
   t tS r3   )r   r   	Cuda_hnegr   r   r   r   resolve_hnegj  s    zCudaFp16Template.resolve_hnegc                 C   s
   t tS r3   )r   r   	Cuda_habsr   r   r   r   resolve_habsm  s    zCudaFp16Template.resolve_habsc                 C   s
   t tS r3   )r   r   rT   r   r   r   r   resolve_hfmap  s    zCudaFp16Template.resolve_hfmac                 C   s   t S r3   )hsin_devicer   r   r   r   resolve_hsins  s    zCudaFp16Template.resolve_hsinc                 C   s   t S r3   )hcos_devicer   r   r   r   resolve_hcosv  s    zCudaFp16Template.resolve_hcosc                 C   s   t S r3   )hlog_devicer   r   r   r   resolve_hlogy  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   r   r   rU   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   Z  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   r   r   Zcgr   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   r   r   r-   r   r   r   r   resolve_shared  s    z!CudaModuleTemplate.resolve_sharedc                 C   s
   t tS r3   )r   r   rG   r   r   r   r   resolve_popc  s    zCudaModuleTemplate.resolve_popcc                 C   s
   t tS r3   )r   r   rX   r   r   r   r   resolve_brev  s    zCudaModuleTemplate.resolve_brevc                 C   s
   t tS r3   )r   r   rY   r   r   r   r   resolve_clz  s    zCudaModuleTemplate.resolve_clzc                 C   s
   t tS r3   )r   r   rZ   r   r   r   r   resolve_ffs  s    zCudaModuleTemplate.resolve_ffsc                 C   s
   t tS r3   )r   r   rP   r   r   r   r   resolve_fma  s    zCudaModuleTemplate.resolve_fmac                 C   s
   t tS r3   )r   r   rW   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   r<   r   r   r   r   resolve_shfl_sync_intrinsic  s    z.CudaModuleTemplate.resolve_shfl_sync_intrinsicc                 C   s
   t tS r3   )r   r   rA   r   r   r   r   resolve_vote_sync_intrinsic  s    z.CudaModuleTemplate.resolve_vote_sync_intrinsicc                 C   s
   t tS r3   )r   r   rB   r   r   r   r   resolve_match_any_sync  s    z)CudaModuleTemplate.resolve_match_any_syncc                 C   s
   t tS r3   )r   r   rC   r   r   r   r   resolve_match_all_sync  s    z)CudaModuleTemplate.resolve_match_all_syncc                 C   s
   t tS r3   )r   r   rD   r   r   r   r   resolve_activemask  s    z%CudaModuleTemplate.resolve_activemaskc                 C   s
   t tS r3   )r   r   rF   r   r   r   r   resolve_lanemask_lt  s    z&CudaModuleTemplate.resolve_lanemask_ltc                 C   s
   t tS r3   )r   r   r[   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   r   r   r   r   r   r   r   resolve_atomic  s    z!CudaModuleTemplate.resolve_atomicc                 C   s   t tjS r3   )r   r   r   rU   r   r   r   r   resolve_fp16  s    zCudaModuleTemplate.resolve_fp16c                 C   s   t tjS r3   )r   r   r   r4   r   r   r   r   resolve_const  s    z CudaModuleTemplate.resolve_constc                 C   s   t tjS r3   )r   r   r   r1   r   r   r   r   resolve_local  s    z CudaModuleTemplate.resolve_localN)$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  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   Znumba.core.typeconvr   Znumbar   Znumba.cuda.compilerr   registryrg   Zregister_attrrn   r   r,   r0   r2   r5   r8   r9   r:   r<   rA   rB   rC   rD   rF   rG   rP   rT   rW   rX   rY   rZ   r[   ri   ro   rq   floatrr   ru   r}   r   r   rU   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   rS   rR   rK   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   r   funcr   r   r   r   <module>   s>  $ 	

			














.^[