U
    ,d(I                     @   sp  d dl mZmZ d dl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 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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'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 d0d1 d1e	Z,eG d2d3 d3e	Z-eG d4d5 d5e	Z.eG d6d7 d7e	Z/eG d8d9 d9e	Z0eG d:d; d;e	Z1eG d<d= d=e	Z2eG d>d? d?e	Z3eG d@dA dAe	Z4eG dBdC dCe	Z5eG dDdE dEe	Z6eG dFdG dGe
Z7dHdI Z8dJdK Z9dLdM Z:e9ej;j<Z=e9ej;j>Z?e9ej;j@ZAe9ej;jBZCe9ej;jDZEe8ej;jFZGe8ej;jHZIe:ej;jJZKe:ej;jLZMe:ej;jNZOe:ej;jPZQe:ej;jRZSe:ej;jTZUdNdO ZVejWejXejYejZej[ej\fZ]ejYejZej[ej\fZ^ejZej\fZ_eVej`jae]ZbeVej`jce]ZdeVej`jee]ZfeVej`jge]ZheVej`jie]ZjeVej`jke]ZleVej`jme^ZneVej`joe^ZpeVej`jqe^ZreVej`jse_ZteVej`jue_ZveVej`jwe^ZxeG dPdQ dQe
ZyeG dRdS dSe	ZzeG dTdU dUeZ{eG dVdW dWeZ|eG dXdY dYeZ}eG dZd[ d[eZ~eG d\d] d]eZeG d^d_ d_eZeG d`da daeZeeee dbS )c    )typeserrors)parse_dtypeparse_shaperegister_number_classes)AttributeTemplateConcreteTemplateAbstractTemplateCallableTemplate	signatureRegistry)dim3
grid_group)cudac                   @   s   e Zd Zdd ZdS )GridFunctionc                 C   s   dd }|S )Nc                 S   sX   t | tjst| | j}|dkr,tj}n |dkrDttj|}ntdt	|tjS )N   )      zargument can only be 1, 2, 3)

isinstancer   IntegerLiteralr   ZRequireLiteralValueZliteral_valueint32UniTuple
ValueErrorr   )ndimvalrestype r   7/tmp/pip-unpacked-wheel-eu7e0c37/numba/cuda/cudadecl.pytyper   s    
z#GridFunction.generic.<locals>.typerr   selfr   r   r   r   generic   s    zGridFunction.genericN__name__
__module____qualname__r!   r   r   r   r   r      s   r   c                   @   s   e Zd ZejZdS )	Cuda_gridN)r#   r$   r%   r   Zgridkeyr   r   r   r   r&   #   s   r&   c                   @   s   e Zd ZejZdS )Cuda_gridsizeN)r#   r$   r%   r   Zgridsizer'   r   r   r   r   r(   (   s   r(   c                   @   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 r   )r   r   r   ).0sr   r   r   
<listcomp>7   s   z:Cuda_array_decl.generic.<locals>.typer.<locals>.<listcomp>C)dtyper   Zlayout)
r   r   ZIntegerr   Tupler   anyr   r   Array)shaper.   r   Znb_dtyper   r   r   r   /   s    z&Cuda_array_decl.generic.<locals>.typerr   r   r   r   r   r!   .   s    zCuda_array_decl.genericNr"   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arrayr'   r   r   r   r   r3   E   s   r3   c                   @   s   e Zd ZejjZdS )Cuda_local_arrayN)r#   r$   r%   r   localr5   r'   r   r   r   r   r6   J   s   r6   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   T   s    z,Cuda_const_array_like.generic.<locals>.typerr   r   r   r   r   r!   S   s    zCuda_const_array_like.genericN)r#   r$   r%   r   constZ
array_liker'   r!   r   r   r   r   r8   O   s   r8   c                   @   s   e Zd ZejZeejgZ	dS )Cuda_syncthreadsN)
r#   r$   r%   r   Zsyncthreadsr'   r   r   nonecasesr   r   r   r   r;   Y   s   r;   c                   @   s"   e Zd ZejZeejejgZ	dS )Cuda_syncthreads_countN)
r#   r$   r%   r   Zsyncthreads_countr'   r   r   i4r=   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_syncthreads_andN)
r#   r$   r%   r   Zsyncthreads_andr'   r   r   r?   r=   r   r   r   r   r@   e   s   r@   c                   @   s"   e Zd ZejZeejejgZ	dS )Cuda_syncthreads_orN)
r#   r$   r%   r   Zsyncthreads_orr'   r   r   r?   r=   r   r   r   r   rA   k   s   rA   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   r<   r=   r   r   r   r   rB   q   s   rB   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   r<   r=   r   r   r   r   rC   w   s   rC   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   r<   r=   r   r   r   r   rD   }   s   rD   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   r<   r?   r=   r   r   r   r   rE      s   rE   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   r=   r   r   r   r   rF      s   rF   c                   @   s    e Zd ZeejZdd ZdS )CudaCgModuleTemplatec                 C   s
   t tS r9   )r   FunctionrF   r    modr   r   r   resolve_this_grid   s    z&CudaCgModuleTemplate.resolve_this_gridN)	r#   r$   r%   r   Moduler   rG   r'   rL   r   r   r   r   rH      s   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   r   this)r    argskwsr   r   r   r!      s    zCuda_grid_group_sync.genericNr#   r$   r%   r'   r!   r   r   r   r   rN      s   rN   c                   @   s   e Zd ZeZdd ZdS )GridGroup_attrsc                 C   s   t ttS r9   )r   ZBoundFunctionrN   r   rJ   r   r   r   resolve_sync   s    zGridGroup_attrs.resolve_syncN)r#   r$   r%   r   r'   rT   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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f8r=   r   r   r   r   rU      s<                   rU   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?   rV   r=   r   r   r   r   rZ      s     rZ   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?   rW   rX   rY   r=   r   r   r   r   r[      s   r[   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?   rV   rW   rX   rY   r=   r   r   r   r   r\      s   r\   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   uint32r=   r   r   r   r   r]      s   r]   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   r^   r=   r   r   r   r   r_      s   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_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int16r   int64uint8uint16r^   uint64r=   r   r   r   r   r`      s   r`   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%   ra   r   fmar'   r   r   float32float64r=   r   r   r   r   rh      s
   rh   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   float16r=   r   r   r   r   rl      s   rl   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   rj   rk   r=   r   r   r   r   ro     s   ro   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   r^   rg   r=   r   r   r   r   rp     s   rp   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%   ra   r   Zclzr'   r   r   rb   rc   r   rd   re   rf   r^   rg   r=   r   r   r   r   rq     s   rq   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%   ra   r   Zffsr'   r   r   r^   rb   rc   r   rd   re   rf   rg   r=   r   r   r   r   rr   ,  s   rr   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 r9   )AssertionErrorr   rk   rj   rc   rf   r   r^   rd   rg   r   )r    rP   rQ   testabsupported_typesr   r   r   r!   C  s    
   zCuda_selp.genericN)r#   r$   r%   r   Zselpr'   r!   r   r   r   r   rs   ?  s   rs   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   rn   r=   r   l_keyr   r   Cuda_fp16_unaryU  s   r|   registerr   )r{   r|   r   rz   r   _genfp16_unaryT  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_binaryNry   r   rz   r   r   Cuda_fp16_binary^  s   r   r}   )r{   r   r   rz   r   _genfp16_binary]  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   rV   rn   r=   r   rz   r   r   Cuda_fp16_cmpg  s   r   r}   )r{   r   r   rz   r   _genfp16_binary_comparisonf  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 )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 Nr   )rt   r.   r   r   r   Zintp)r    rP   rQ   aryidxr   )rx   r   r   r!     s    



z!_gen.<locals>.Cuda_atomic.genericNrR   r   r{   rx   r   r   Cuda_atomic  s   r   )r~   r	   )r{   rx   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   )rt   r.   integer_numba_typesr   r   )r    rP   rQ   r   oldr   Z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Zeejej	gZ
dS )Cuda_nanosleepN)r#   r$   r%   r   Z	nanosleepr'   r   r   voidr^   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   s   t jS r9   r   r   rJ   r   r   r   	resolve_x  s    zDim3_attrs.resolve_xc                 C   s   t jS r9   r   rJ   r   r   r   	resolve_y  s    zDim3_attrs.resolve_yc                 C   s   t jS r9   r   rJ   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 r9   )r   rI   r3   rJ   r   r   r   resolve_array  s    z&CudaSharedModuleTemplate.resolve_arrayN)	r#   r$   r%   r   rM   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 )CudaConstModuleTemplatec                 C   s
   t tS r9   )r   rI   r8   rJ   r   r   r   resolve_array_like  s    z*CudaConstModuleTemplate.resolve_array_likeN)	r#   r$   r%   r   rM   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 )CudaLocalModuleTemplatec                 C   s
   t tS r9   )r   rI   r6   rJ   r   r   r   r     s    z%CudaLocalModuleTemplate.resolve_arrayN)	r#   r$   r%   r   rM   r   r7   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S )CudaAtomicTemplatec                 C   s
   t tS r9   )r   rI   Cuda_atomic_addrJ   r   r   r   resolve_add  s    zCudaAtomicTemplate.resolve_addc                 C   s
   t tS r9   )r   rI   Cuda_atomic_subrJ   r   r   r   resolve_sub  s    zCudaAtomicTemplate.resolve_subc                 C   s
   t tS r9   )r   rI   Cuda_atomic_andrJ   r   r   r   resolve_and_  s    zCudaAtomicTemplate.resolve_and_c                 C   s
   t tS r9   )r   rI   Cuda_atomic_orrJ   r   r   r   resolve_or_  s    zCudaAtomicTemplate.resolve_or_c                 C   s
   t tS r9   )r   rI   Cuda_atomic_xorrJ   r   r   r   resolve_xor  s    zCudaAtomicTemplate.resolve_xorc                 C   s
   t tS r9   )r   rI   Cuda_atomic_incrJ   r   r   r   resolve_inc  s    zCudaAtomicTemplate.resolve_incc                 C   s
   t tS r9   )r   rI   Cuda_atomic_decrJ   r   r   r   resolve_dec  s    zCudaAtomicTemplate.resolve_decc                 C   s
   t tS r9   )r   rI   Cuda_atomic_exchrJ   r   r   r   resolve_exch  s    zCudaAtomicTemplate.resolve_exchc                 C   s
   t tS r9   )r   rI   Cuda_atomic_maxrJ   r   r   r   resolve_max  s    zCudaAtomicTemplate.resolve_maxc                 C   s
   t tS r9   )r   rI   Cuda_atomic_minrJ   r   r   r   resolve_min  s    zCudaAtomicTemplate.resolve_minc                 C   s
   t tS r9   )r   rI   Cuda_atomic_nanminrJ   r   r   r   resolve_nanmin  s    z!CudaAtomicTemplate.resolve_nanminc                 C   s
   t tS r9   )r   rI   Cuda_atomic_nanmaxrJ   r   r   r   resolve_nanmax
  s    z!CudaAtomicTemplate.resolve_nanmaxc                 C   s
   t tS r9   )r   rI   r   rJ   r   r   r   resolve_compare_and_swap  s    z+CudaAtomicTemplate.resolve_compare_and_swapN)r#   r$   r%   r   rM   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S )CudaFp16Templatec                 C   s
   t tS r9   )r   rI   	Cuda_haddrJ   r   r   r   resolve_hadd  s    zCudaFp16Template.resolve_haddc                 C   s
   t tS r9   )r   rI   	Cuda_hsubrJ   r   r   r   resolve_hsub  s    zCudaFp16Template.resolve_hsubc                 C   s
   t tS r9   )r   rI   	Cuda_hmulrJ   r   r   r   resolve_hmul  s    zCudaFp16Template.resolve_hmulc                 C   s
   t tS r9   )r   rI   	Cuda_hnegrJ   r   r   r   resolve_hneg  s    zCudaFp16Template.resolve_hnegc                 C   s
   t tS r9   )r   rI   	Cuda_habsrJ   r   r   r   resolve_habs!  s    zCudaFp16Template.resolve_habsc                 C   s
   t tS r9   )r   rI   rl   rJ   r   r   r   resolve_hfma$  s    zCudaFp16Template.resolve_hfmac                 C   s
   t tS r9   )r   rI   Cuda_heqrJ   r   r   r   resolve_heq'  s    zCudaFp16Template.resolve_heqc                 C   s
   t tS r9   )r   rI   Cuda_hnerJ   r   r   r   resolve_hne*  s    zCudaFp16Template.resolve_hnec                 C   s
   t tS r9   )r   rI   Cuda_hgerJ   r   r   r   resolve_hge-  s    zCudaFp16Template.resolve_hgec                 C   s
   t tS r9   )r   rI   Cuda_hgtrJ   r   r   r   resolve_hgt0  s    zCudaFp16Template.resolve_hgtc                 C   s
   t tS r9   )r   rI   Cuda_hlerJ   r   r   r   resolve_hle3  s    zCudaFp16Template.resolve_hlec                 C   s
   t tS r9   )r   rI   Cuda_hltrJ   r   r   r   resolve_hlt6  s    zCudaFp16Template.resolve_hltc                 C   s
   t tS r9   )r   rI   	Cuda_hmaxrJ   r   r   r   resolve_hmax9  s    zCudaFp16Template.resolve_hmaxc                 C   s
   t tS r9   )r   rI   	Cuda_hminrJ   r   r   r   resolve_hmin<  s    zCudaFp16Template.resolve_hminN)r#   r$   r%   r   rM   r   rm   r'   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r     s   r   c                   @   s6  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;d< Z$d=d> Z%d?d@ Z&dAdB Z'dCdD Z(dEdF Z)dGdH Z*dIS )JCudaModuleTemplatec                 C   s
   t tS r9   )r   rI   r&   rJ   r   r   r   resolve_gridD  s    zCudaModuleTemplate.resolve_gridc                 C   s
   t tS r9   )r   rI   r(   rJ   r   r   r   resolve_gridsizeG  s    z#CudaModuleTemplate.resolve_gridsizec                 C   s   t tjS r9   )r   rM   r   rG   rJ   r   r   r   
resolve_cgJ  s    zCudaModuleTemplate.resolve_cgc                 C   s   t S r9   r   rJ   r   r   r   resolve_threadIdxM  s    z$CudaModuleTemplate.resolve_threadIdxc                 C   s   t S r9   r   rJ   r   r   r   resolve_blockIdxP  s    z#CudaModuleTemplate.resolve_blockIdxc                 C   s   t S r9   r   rJ   r   r   r   resolve_blockDimS  s    z#CudaModuleTemplate.resolve_blockDimc                 C   s   t S r9   r   rJ   r   r   r   resolve_gridDimV  s    z"CudaModuleTemplate.resolve_gridDimc                 C   s   t jS r9   r   rJ   r   r   r   resolve_warpsizeY  s    z#CudaModuleTemplate.resolve_warpsizec                 C   s   t jS r9   r   rJ   r   r   r   resolve_laneid\  s    z!CudaModuleTemplate.resolve_laneidc                 C   s   t tjS r9   )r   rM   r   r4   rJ   r   r   r   resolve_shared_  s    z!CudaModuleTemplate.resolve_sharedc                 C   s
   t tS r9   )r   rI   r`   rJ   r   r   r   resolve_popcb  s    zCudaModuleTemplate.resolve_popcc                 C   s
   t tS r9   )r   rI   rp   rJ   r   r   r   resolve_breve  s    zCudaModuleTemplate.resolve_brevc                 C   s
   t tS r9   )r   rI   rq   rJ   r   r   r   resolve_clzh  s    zCudaModuleTemplate.resolve_clzc                 C   s
   t tS r9   )r   rI   rr   rJ   r   r   r   resolve_ffsk  s    zCudaModuleTemplate.resolve_ffsc                 C   s
   t tS r9   )r   rI   rh   rJ   r   r   r   resolve_fman  s    zCudaModuleTemplate.resolve_fmac                 C   s
   t tS r9   )r   rI   ro   rJ   r   r   r   resolve_cbrtq  s    zCudaModuleTemplate.resolve_cbrtc                 C   s
   t tS r9   )r   rI   r;   rJ   r   r   r   resolve_syncthreadst  s    z&CudaModuleTemplate.resolve_syncthreadsc                 C   s
   t tS r9   )r   rI   r>   rJ   r   r   r   resolve_syncthreads_countw  s    z,CudaModuleTemplate.resolve_syncthreads_countc                 C   s
   t tS r9   )r   rI   r@   rJ   r   r   r   resolve_syncthreads_andz  s    z*CudaModuleTemplate.resolve_syncthreads_andc                 C   s
   t tS r9   )r   rI   rA   rJ   r   r   r   resolve_syncthreads_or}  s    z)CudaModuleTemplate.resolve_syncthreads_orc                 C   s
   t tS r9   )r   rI   rB   rJ   r   r   r   resolve_threadfence  s    z&CudaModuleTemplate.resolve_threadfencec                 C   s
   t tS r9   )r   rI   rC   rJ   r   r   r   resolve_threadfence_block  s    z,CudaModuleTemplate.resolve_threadfence_blockc                 C   s
   t tS r9   )r   rI   rD   rJ   r   r   r   resolve_threadfence_system  s    z-CudaModuleTemplate.resolve_threadfence_systemc                 C   s
   t tS r9   )r   rI   rE   rJ   r   r   r   resolve_syncwarp  s    z#CudaModuleTemplate.resolve_syncwarpc                 C   s
   t tS r9   )r   rI   rU   rJ   r   r   r   resolve_shfl_sync_intrinsic  s    z.CudaModuleTemplate.resolve_shfl_sync_intrinsicc                 C   s
   t tS r9   )r   rI   rZ   rJ   r   r   r   resolve_vote_sync_intrinsic  s    z.CudaModuleTemplate.resolve_vote_sync_intrinsicc                 C   s
   t tS r9   )r   rI   r[   rJ   r   r   r   resolve_match_any_sync  s    z)CudaModuleTemplate.resolve_match_any_syncc                 C   s
   t tS r9   )r   rI   r\   rJ   r   r   r   resolve_match_all_sync  s    z)CudaModuleTemplate.resolve_match_all_syncc                 C   s
   t tS r9   )r   rI   r]   rJ   r   r   r   resolve_activemask  s    z%CudaModuleTemplate.resolve_activemaskc                 C   s
   t tS r9   )r   rI   r_   rJ   r   r   r   resolve_lanemask_lt  s    z&CudaModuleTemplate.resolve_lanemask_ltc                 C   s
   t tS r9   )r   rI   rs   rJ   r   r   r   resolve_selp  s    zCudaModuleTemplate.resolve_selpc                 C   s
   t tS r9   )r   rI   r   rJ   r   r   r   resolve_nanosleep  s    z$CudaModuleTemplate.resolve_nanosleepc                 C   s   t tjS r9   )r   rM   r   r   rJ   r   r   r   resolve_atomic  s    z!CudaModuleTemplate.resolve_atomicc                 C   s   t tjS r9   )r   rM   r   rm   rJ   r   r   r   resolve_fp16  s    zCudaModuleTemplate.resolve_fp16c                 C   s   t tjS r9   )r   rM   r   r:   rJ   r   r   r   resolve_const  s    z CudaModuleTemplate.resolve_constc                 C   s   t tjS r9   )r   rM   r   r7   rJ   r   r   r   resolve_local  s    z CudaModuleTemplate.resolve_localN)+r#   r$   r%   r   rM   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   r   r   @  sJ   
r   N)Z
numba.corer   r   Znumba.core.typing.npydeclr   r   r   Znumba.core.typing.templatesr   r   r	   r
   r   r   Znumba.cuda.typesr   r   Znumbar   registryr~   Zregister_attrZregister_globalr   r&   r(   r)   r3   r6   r8   r;   r>   r@   rA   rB   rC   rD   rE   rF   rH   rN   rS   rU   rZ   r[   r\   r]   r_   r`   rh   rl   ro   rp   rq   rr   rs   r   r   r   rm   Zhaddr   Zhsubr   Zhmulr   Zhmaxr   Zhminr   Zhnegr   Zhabsr   Zheqr   hner   Zhger   Zhgtr   hler   Zhltr   r   rk   rj   r   r^   rd   rg   Zall_numba_typesr   Zunsigned_int_numba_typesr   addr   subr   maxr   minr   Znanmaxr   Znanminr   and_r   or_r   xorr   incr   decr   Zexchr   r   r   r   r   r   r   r   r   r   rM   r   r   r   r   <module>   s   	

			   +.p