o
    8ήc(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                   @      e Zd Zdd ZdS )GridFunctionc                 C      dd }|S )Nc                 S   sX   t | tjst| | j}|dkrtj}n|dv r"ttj|}ntdt	|tjS )N   )      zargument can only be 1, 2, 3)

isinstancer   IntegerLiteralr   RequireLiteralValueliteral_valueint32UniTuple
ValueErrorr   )ndimvalrestype r    :/tmp/pip-target-vg8gfxp4/lib/python/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          r   c                   @      e Zd ZejZdS )	Cuda_gridN)r'   r(   r)   r   gridkeyr    r    r    r!   r,   #       
r,   c                   @   r+   )Cuda_gridsizeN)r'   r(   r)   r   gridsizer.   r    r    r    r!   r0   (   r/   r0   c                   @   r   )Cuda_array_declc                 C   r   )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 r    )r   r   r   ).0sr    r    r!   
<listcomp>7   s    z:Cuda_array_decl.generic.<locals>.typer.<locals>.<listcomp>C)dtyper   layout)
r   r   Integerr   Tupler   anyr   r   Array)shaper7   r   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!   r2   -   r*   r2   c                   @      e Zd ZejjZdS )Cuda_shared_arrayN)r'   r(   r)   r   sharedarrayr.   r    r    r    r!   r@   E       r@   c                   @   r?   )Cuda_local_arrayN)r'   r(   r)   r   localrB   r.   r    r    r    r!   rD   J   rC   rD   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"   T      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   const
array_liker.   r%   r    r    r    r!   rG   O       rG   c                   @      e Zd ZejZeejgZ	dS )Cuda_syncthreadsN)
r'   r(   r)   r   syncthreadsr.   r   r   nonecasesr    r    r    r!   rO   Y       rO   c                   @   "   e Zd ZejZeejejgZ	dS )Cuda_syncthreads_countN)
r'   r(   r)   r   syncthreads_countr.   r   r   i4rR   r    r    r    r!   rU   _       rU   c                   @   rT   )Cuda_syncthreads_andN)
r'   r(   r)   r   syncthreads_andr.   r   r   rW   rR   r    r    r    r!   rY   e   rX   rY   c                   @   rT   )Cuda_syncthreads_orN)
r'   r(   r)   r   syncthreads_orr.   r   r   rW   rR   r    r    r    r!   r[   k   rX   r[   c                   @   rN   )Cuda_threadfence_deviceN)
r'   r(   r)   r   threadfencer.   r   r   rQ   rR   r    r    r    r!   r]   q   rS   r]   c                   @   rN   )Cuda_threadfence_blockN)
r'   r(   r)   r   threadfence_blockr.   r   r   rQ   rR   r    r    r    r!   r_   w   rS   r_   c                   @   rN   )Cuda_threadfence_systemN)
r'   r(   r)   r   threadfence_systemr.   r   r   rQ   rR   r    r    r    r!   ra   }   rS   ra   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   syncwarpr.   r   r   rQ   rW   rR   r    r    r    r!   rc      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   cg	this_gridr.   r   r   rR   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 rH   )r   Functionre   r$   modr    r    r!   resolve_this_grid      
z&CudaCgModuleTemplate.resolve_this_gridN)	r'   r(   r)   r   Moduler   rf   r.   rn   r    r    r    r!   ri          ri   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   r   this)r$   argskwsr    r    r!   r%      s   zCuda_grid_group_sync.genericNr'   r(   r)   r.   r%   r    r    r    r!   rr      s    rr   c                   @   s   e Zd ZeZdd ZdS )GridGroup_attrsc                 C   s   t ttS rH   )r   BoundFunctionrr   r   rl   r    r    r!   resolve_sync      zGridGroup_attrs.resolve_syncN)r'   r(   r)   r   r.   rz   r    r    r    r!   rx      s    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   shfl_sync_intrinsicr.   r   r   r:   rW   b1i8f4f8rR   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   vote_sync_intrinsicr.   r   r   r:   rW   r~   rR   r    r    r    r!   r      s
    
r   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_syncr.   r   r   rW   r   r   r   rR   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   match_all_syncr.   r   r   r:   rW   r~   r   r   r   rR   r    r    r    r!   r      s    r   c                   @   rN   )Cuda_activemaskN)
r'   r(   r)   r   
activemaskr.   r   r   uint32rR   r    r    r    r!   r      rS   r   c                   @   rN   )Cuda_lanemask_ltN)
r'   r(   r)   r   lanemask_ltr.   r   r   r   rR   r    r    r    r!   r      rS   r   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   popcr.   r   r   int8int16r   int64uint8uint16r   uint64rR   r    r    r    r!   r          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)   r   r   fmar.   r   r   float32float64rR   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hfmar.   r   r   float16rR   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   cbrtr.   r   r   r   r   rR   r    r    r    r!   r     s
    r   c                   @   r   )	Cuda_brevN)r'   r(   r)   r   brevr.   r   r   r   r   rR   r    r    r    r!   r     s
    r   c                
   @   r   )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)   r   r   clzr.   r   r   r   r   r   r   r   r   r   r   rR   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)   r   r   ffsr.   r   r   r   r   r   r   r   r   r   r   rR   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 rH   )
r   r   r   r   r   r   r   r   r   r   )r$   ru   rv   testabsupported_typesr    r    r!   r%   C  s   
zCuda_selp.genericN)r'   r(   r)   r   selpr.   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)   r.   r   r   r   rR   r    l_keyr    r!   Cuda_fp16_unaryU  s    r   registerr   )r   r   r    r   r!   _genfp16_unaryT     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   )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~   r   rR   r    r   r    r!   Cuda_fp16_cmpg  s    r   r   )r   r   r    r   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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   )r7   r   r   r   intp)r$   ru   rv   aryidxr   )r   r    r!   r%     s   



z!_gen.<locals>.Cuda_atomic.genericNrw   r    r   r   r    r!   Cuda_atomic  s    r   )r   r	   )r   r   r   r    r   r!   _gen  s   r   c                   @   rF   )Cuda_atomic_compare_and_swapc                 C   s@   |rJ |\}}}|j }|tv r|jdkrt||||S d S d S r   )r7   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   atomiccompare_and_swapr.   r%   r    r    r    r!   r     rM   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   	nanosleepr.   r   r   voidr   rR   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 rH   r   r   rl   r    r    r!   	resolve_x     zDim3_attrs.resolve_xc                 C   r   rH   r   rl   r    r    r!   	resolve_y  r   zDim3_attrs.resolve_yc                 C   r   rH   r   rl   r    r    r!   	resolve_z  r   zDim3_attrs.resolve_zN)r'   r(   r)   r   r.   r   r   r   r    r    r    r!   r     s
    r   c                   @   rh   )CudaSharedModuleTemplatec                 C   rj   rH   )r   rk   r@   rl   r    r    r!   resolve_array  ro   z&CudaSharedModuleTemplate.resolve_arrayN)	r'   r(   r)   r   rp   r   rA   r.   r   r    r    r    r!   r     rq   r   c                   @   rh   )CudaConstModuleTemplatec                 C   rj   rH   )r   rk   rG   rl   r    r    r!   resolve_array_like  ro   z*CudaConstModuleTemplate.resolve_array_likeN)	r'   r(   r)   r   rp   r   rK   r.   r   r    r    r    r!   r     rq   r   c                   @   rh   )CudaLocalModuleTemplatec                 C   rj   rH   )r   rk   rD   rl   r    r    r!   r     ro   z%CudaLocalModuleTemplate.resolve_arrayN)	r'   r(   r)   r   rp   r   rE   r.   r   r    r    r    r!   r     rq   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   rj   rH   )r   rk   Cuda_atomic_addrl   r    r    r!   resolve_add  ro   zCudaAtomicTemplate.resolve_addc                 C   rj   rH   )r   rk   Cuda_atomic_subrl   r    r    r!   resolve_sub  ro   zCudaAtomicTemplate.resolve_subc                 C   rj   rH   )r   rk   Cuda_atomic_andrl   r    r    r!   resolve_and_  ro   zCudaAtomicTemplate.resolve_and_c                 C   rj   rH   )r   rk   Cuda_atomic_orrl   r    r    r!   resolve_or_  ro   zCudaAtomicTemplate.resolve_or_c                 C   rj   rH   )r   rk   Cuda_atomic_xorrl   r    r    r!   resolve_xor  ro   zCudaAtomicTemplate.resolve_xorc                 C   rj   rH   )r   rk   Cuda_atomic_incrl   r    r    r!   resolve_inc  ro   zCudaAtomicTemplate.resolve_incc                 C   rj   rH   )r   rk   Cuda_atomic_decrl   r    r    r!   resolve_dec  ro   zCudaAtomicTemplate.resolve_decc                 C   rj   rH   )r   rk   Cuda_atomic_exchrl   r    r    r!   resolve_exch  ro   zCudaAtomicTemplate.resolve_exchc                 C   rj   rH   )r   rk   Cuda_atomic_maxrl   r    r    r!   resolve_max  ro   zCudaAtomicTemplate.resolve_maxc                 C   rj   rH   )r   rk   Cuda_atomic_minrl   r    r    r!   resolve_min  ro   zCudaAtomicTemplate.resolve_minc                 C   rj   rH   )r   rk   Cuda_atomic_nanminrl   r    r    r!   resolve_nanmin  ro   z!CudaAtomicTemplate.resolve_nanminc                 C   rj   rH   )r   rk   Cuda_atomic_nanmaxrl   r    r    r!   resolve_nanmax
  ro   z!CudaAtomicTemplate.resolve_nanmaxc                 C   rj   rH   )r   rk   r   rl   r    r    r!   resolve_compare_and_swap  ro   z+CudaAtomicTemplate.resolve_compare_and_swapN)r'   r(   r)   r   rp   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   rj   rH   )r   rk   	Cuda_haddrl   r    r    r!   resolve_hadd  ro   zCudaFp16Template.resolve_haddc                 C   rj   rH   )r   rk   	Cuda_hsubrl   r    r    r!   resolve_hsub  ro   zCudaFp16Template.resolve_hsubc                 C   rj   rH   )r   rk   	Cuda_hmulrl   r    r    r!   resolve_hmul  ro   zCudaFp16Template.resolve_hmulc                 C   rj   rH   )r   rk   	Cuda_hnegrl   r    r    r!   resolve_hneg  ro   zCudaFp16Template.resolve_hnegc                 C   rj   rH   )r   rk   	Cuda_habsrl   r    r    r!   resolve_habs!  ro   zCudaFp16Template.resolve_habsc                 C   rj   rH   )r   rk   r   rl   r    r    r!   resolve_hfma$  ro   zCudaFp16Template.resolve_hfmac                 C   rj   rH   )r   rk   Cuda_heqrl   r    r    r!   resolve_heq'  ro   zCudaFp16Template.resolve_heqc                 C   rj   rH   )r   rk   Cuda_hnerl   r    r    r!   resolve_hne*  ro   zCudaFp16Template.resolve_hnec                 C   rj   rH   )r   rk   Cuda_hgerl   r    r    r!   resolve_hge-  ro   zCudaFp16Template.resolve_hgec                 C   rj   rH   )r   rk   Cuda_hgtrl   r    r    r!   resolve_hgt0  ro   zCudaFp16Template.resolve_hgtc                 C   rj   rH   )r   rk   Cuda_hlerl   r    r    r!   resolve_hle3  ro   zCudaFp16Template.resolve_hlec                 C   rj   rH   )r   rk   Cuda_hltrl   r    r    r!   resolve_hlt6  ro   zCudaFp16Template.resolve_hltc                 C   rj   rH   )r   rk   	Cuda_hmaxrl   r    r    r!   resolve_hmax9  ro   zCudaFp16Template.resolve_hmaxc                 C   rj   rH   )r   rk   	Cuda_hminrl   r    r    r!   resolve_hmin<  ro   zCudaFp16Template.resolve_hminN)r'   r(   r)   r   rp   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                   @   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   rj   rH   )r   rk   r,   rl   r    r    r!   resolve_gridD  ro   zCudaModuleTemplate.resolve_gridc                 C   rj   rH   )r   rk   r0   rl   r    r    r!   resolve_gridsizeG  ro   z#CudaModuleTemplate.resolve_gridsizec                 C      t tjS rH   )r   rp   r   rf   rl   r    r    r!   
resolve_cgJ  r{   zCudaModuleTemplate.resolve_cgc                 C      t S rH   r   rl   r    r    r!   resolve_threadIdxM  rJ   z$CudaModuleTemplate.resolve_threadIdxc                 C   r  rH   r  rl   r    r    r!   resolve_blockIdxP  rJ   z#CudaModuleTemplate.resolve_blockIdxc                 C   r  rH   r  rl   r    r    r!   resolve_blockDimS  rJ   z#CudaModuleTemplate.resolve_blockDimc                 C   r  rH   r  rl   r    r    r!   resolve_gridDimV  rJ   z"CudaModuleTemplate.resolve_gridDimc                 C   r   rH   r   rl   r    r    r!   resolve_warpsizeY  r   z#CudaModuleTemplate.resolve_warpsizec                 C   r   rH   r   rl   r    r    r!   resolve_laneid\  r   z!CudaModuleTemplate.resolve_laneidc                 C   r  rH   )r   rp   r   rA   rl   r    r    r!   resolve_shared_  r{   z!CudaModuleTemplate.resolve_sharedc                 C   rj   rH   )r   rk   r   rl   r    r    r!   resolve_popcb  ro   zCudaModuleTemplate.resolve_popcc                 C   rj   rH   )r   rk   r   rl   r    r    r!   resolve_breve  ro   zCudaModuleTemplate.resolve_brevc                 C   rj   rH   )r   rk   r   rl   r    r    r!   resolve_clzh  ro   zCudaModuleTemplate.resolve_clzc                 C   rj   rH   )r   rk   r   rl   r    r    r!   resolve_ffsk  ro   zCudaModuleTemplate.resolve_ffsc                 C   rj   rH   )r   rk   r   rl   r    r    r!   resolve_fman  ro   zCudaModuleTemplate.resolve_fmac                 C   rj   rH   )r   rk   r   rl   r    r    r!   resolve_cbrtq  ro   zCudaModuleTemplate.resolve_cbrtc                 C   rj   rH   )r   rk   rO   rl   r    r    r!   resolve_syncthreadst  ro   z&CudaModuleTemplate.resolve_syncthreadsc                 C   rj   rH   )r   rk   rU   rl   r    r    r!   resolve_syncthreads_countw  ro   z,CudaModuleTemplate.resolve_syncthreads_countc                 C   rj   rH   )r   rk   rY   rl   r    r    r!   resolve_syncthreads_andz  ro   z*CudaModuleTemplate.resolve_syncthreads_andc                 C   rj   rH   )r   rk   r[   rl   r    r    r!   resolve_syncthreads_or}  ro   z)CudaModuleTemplate.resolve_syncthreads_orc                 C   rj   rH   )r   rk   r]   rl   r    r    r!   resolve_threadfence  ro   z&CudaModuleTemplate.resolve_threadfencec                 C   rj   rH   )r   rk   r_   rl   r    r    r!   resolve_threadfence_block  ro   z,CudaModuleTemplate.resolve_threadfence_blockc                 C   rj   rH   )r   rk   ra   rl   r    r    r!   resolve_threadfence_system  ro   z-CudaModuleTemplate.resolve_threadfence_systemc                 C   rj   rH   )r   rk   rc   rl   r    r    r!   resolve_syncwarp  ro   z#CudaModuleTemplate.resolve_syncwarpc                 C   rj   rH   )r   rk   r|   rl   r    r    r!   resolve_shfl_sync_intrinsic  ro   z.CudaModuleTemplate.resolve_shfl_sync_intrinsicc                 C   rj   rH   )r   rk   r   rl   r    r    r!   resolve_vote_sync_intrinsic  ro   z.CudaModuleTemplate.resolve_vote_sync_intrinsicc                 C   rj   rH   )r   rk   r   rl   r    r    r!   resolve_match_any_sync  ro   z)CudaModuleTemplate.resolve_match_any_syncc                 C   rj   rH   )r   rk   r   rl   r    r    r!   resolve_match_all_sync  ro   z)CudaModuleTemplate.resolve_match_all_syncc                 C   rj   rH   )r   rk   r   rl   r    r    r!   resolve_activemask  ro   z%CudaModuleTemplate.resolve_activemaskc                 C   rj   rH   )r   rk   r   rl   r    r    r!   resolve_lanemask_lt  ro   z&CudaModuleTemplate.resolve_lanemask_ltc                 C   rj   rH   )r   rk   r   rl   r    r    r!   resolve_selp  ro   zCudaModuleTemplate.resolve_selpc                 C   rj   rH   )r   rk   r   rl   r    r    r!   resolve_nanosleep  ro   z$CudaModuleTemplate.resolve_nanosleepc                 C   r  rH   )r   rp   r   r   rl   r    r    r!   resolve_atomic  r{   z!CudaModuleTemplate.resolve_atomicc                 C   r  rH   )r   rp   r   r   rl   r    r    r!   resolve_fp16  r{   zCudaModuleTemplate.resolve_fp16c                 C   r  rH   )r   rp   r   rK   rl   r    r    r!   resolve_const  r{   z CudaModuleTemplate.resolve_constc                 C   r  rH   )r   rp   r   rE   rl   r    r    r!   resolve_local  r{   z CudaModuleTemplate.resolve_localN)+r'   r(   r)   r   rp   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/  r0  r1  r2  r3  r4  r5  r    r    r    r!   r  @  sL    
r  N)
numba.corer   r   numba.core.typing.npydeclr   r   r   numba.core.typing.templatesr   r   r	   r
   r   r   numba.cuda.typesr   r   numbar   registryr   register_attrregister_globalr   r,   r0   r2   r@   rD   rG   rO   rU   rY   r[   r]   r_   ra   rc   re   ri   rr   rx   r|   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   haddr   hsubr   hmulr   hmaxr
  hminr  hnegr   habsr   heqr   hner   hger  hgtr  hler  hltr  r   r   r   r   r   r   r   all_numba_typesr   unsigned_int_numba_typesr   addr   subr   maxr   minr   nanmaxr   nanminr   and_r   or_r   xorr   incr   decr   exchr   r   r   r   r   r   r   r   r   r  rp   r    r    r    r!   <module>   s    	

			+.p