
    
fY                        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
e          Ze G d de                      Ze G d de                      Z e G d de                      Z!e G d de                      Z"e G d de                      Z#e G d de                      Z$e G d de                      Z%e G d de                      Z&e G d de                      Z'e G d de                      Z(e G d d e                      Z)e G d! d"e                      Z*e G d# d$e                      Z+e G d% d&e                      Z,e G d' d(e                      Z-e G d) d*e                      Z.e G d+ d,e                      Z/e G d- d.e                      Z0e G d/ d0e                      Z1e G d1 d2e                      Z2e G d3 d4e                      Z3d5 Z4d6 Z5d7 Z6 ee7           G d8 d9e                      Z8d: Z9d; Z:d< Z;d= Z< e6ej=        j>                  Z? e<e j@                  ZA e<e jB                  ZC e6ej=        jD                  ZE e<e jF                  ZG e<e jH                  ZI e6ej=        jJ                  ZK e<e jL                  ZM e<e jN                  ZO e6ej=        jP                  ZQ e6ej=        jR                  ZS e4ej=        jT                  ZU e5e jV                  ZW e4ej=        jX                  ZY e5eZ          Z[ e9ej=        j\                  Z] e;e j^                    e9ej=        j_                  Z` e;e ja                    e9ej=        jb                  Zc e;e jd                    e9ej=        je                  Zf e;e jg                    e9ej=        jh                  Zi e;e jj                    e9ej=        jk                  Zl e;e jm                    e<e jn                    e<e jo                   d> Zpd? Zq epd@          Zr epdA          Zs epdB          Zt epdC          Zu epdD          Zv epdE          Zw epdF          Zx epdG          Zy epdH          Zz epdI          Z{ epdJ          Z| epdK          Z} epdL          Z~ epdM          Z epdN          Z eqdO          ZdP Zej        ej        ej        ej        ej        ej        fZej        ej        ej        ej        fZej        ej        fZ eej        j@        e          Z eej        jF        e          Z eej        j        e          Z eej        j        e          Z eej        j        e          Z eej        j        e          Z eej        j        e          Z eej        j        e          Z eej        j        e          Z eej        j        e          Z eej        j        e          Z eej        j        e          Ze G dQ dRe                      Ze G dS dTe                      Ze G dU dVe                      Ze G dW dXe                      Ze G dY dZe                      Ze G d[ d\e                      Ze G d] d^e                      Ze G d_ d`e                      Ze G da dbe                      Ze G dc dde                      Z ee ej        e                     eD ]Z eee           e	D ]Z eee           e
D ]Z eee           dS )e    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                       e Zd Zd ZdS )Cuda_array_declc                     d }|S )Nc                 r   t          | t          j                  rt          | t          j                  sd S nDt          | t          j        t          j        f          rt          d | D                       rd S nd S t          |           }t          |          }||t          j	        ||d          S d S d S )Nc                 D    g | ]}t          |t          j                   S  )
isinstancer   IntegerLiteral).0ss     3lib/python3.11/site-packages/numba/cuda/cudadecl.py
<listcomp>z:Cuda_array_decl.generic.<locals>.typer.<locals>.<listcomp>#   s8     ( ( ( 'q%*>??? ( ( (    C)dtypendimlayout)
r   r   Integerr   TupleUniTupleanyr   r   Array)shaper$   r%   nb_dtypes       r    typerz&Cuda_array_decl.generic.<locals>.typer   s     %// !%)=>>  4 EEK#@AA  ( (!&( ( ( ) )  4  tu%%D"5))H#(8{SIIII $#(8(8r"   r   selfr.   s     r    genericzCuda_array_decl.generic   s    	J 	J 	J& r"   N__name__
__module____qualname__r1   r   r"   r    r   r      s#            r"   r   c                   &    e Zd Zej        j        ZdS )Cuda_shared_arrayN)r3   r4   r5   r   sharedarraykeyr   r"   r    r7   r7   1   s        
+
CCCr"   r7   c                   &    e Zd Zej        j        ZdS )Cuda_local_arrayN)r3   r4   r5   r   localr9   r:   r   r"   r    r<   r<   6   s        
*
CCCr"   r<   c                   ,    e Zd Zej        j        Zd ZdS )Cuda_const_array_likec                     d }|S )Nc                     | S Nr   )ndarrays    r    r.   z,Cuda_const_array_like.generic.<locals>.typer@   s    Nr"   r   r/   s     r    r1   zCuda_const_array_like.generic?   s    	 	 	r"   N)r3   r4   r5   r   const
array_liker:   r1   r   r"   r    r?   r?   ;   s-        
*
C    r"   r?   c                   >    e Zd Zej        Z eej                  gZ	dS )Cuda_threadfence_deviceN)
r3   r4   r5   r   threadfencer:   r   r   nonecasesr   r"   r    rG   rG   E   s*        

CYuz""#EEEr"   rG   c                   >    e Zd Zej        Z eej                  gZ	dS )Cuda_threadfence_blockN)
r3   r4   r5   r   threadfence_blockr:   r   r   rI   rJ   r   r"   r    rL   rL   K   s*        

 CYuz""#EEEr"   rL   c                   >    e Zd Zej        Z eej                  gZ	dS )Cuda_threadfence_systemN)
r3   r4   r5   r   threadfence_systemr:   r   r   rI   rJ   r   r"   r    rO   rO   Q   s*        

!CYuz""#EEEr"   rO   c                   h    e Zd Zej        Z eej                   eej        ej	                  gZ
dS )Cuda_syncwarpN)r3   r4   r5   r   syncwarpr:   r   r   rI   i4rJ   r   r"   r    rR   rR   W   s;        
-CYuz""IIej%($C$CDEEEr"   rR   c                   0   e Zd Zej        Z e ej        ej	        ej
        f          ej	        ej	        ej	        ej	        ej	                   e ej        ej        ej
        f          ej	        ej	        ej        ej	        ej	                   e ej        ej        ej
        f          ej	        ej	        ej        ej	        ej	                   e ej        ej        ej
        f          ej	        ej	        ej        ej	        ej	                  gZdS )Cuda_shfl_sync_intrinsicN)r3   r4   r5   r   shfl_sync_intrinsicr:   r   r   r(   rT   b1i8f4f8rJ   r   r"   r    rV   rV   ]   s       

"C	+%+ux233(EHeh%(	D 	D	+%+ux233(EHeh%(	D 	D	+%+ux233(EHeh%(	D 	D	+%+ux233(EHeh%(	D 	D	EEEr"   rV   c                       e Zd Zej        Z e ej        ej	        ej
        f          ej	        ej	        ej
                  gZdS )Cuda_vote_sync_intrinsicN)r3   r4   r5   r   vote_sync_intrinsicr:   r   r   r(   rT   rX   rJ   r   r"   r    r]   r]   l   sP        

"CY{u{EHeh#788x585 5 6EEEr"   r]   c                       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)r3   r4   r5   r   match_any_syncr:   r   r   rT   rY   rZ   r[   rJ   r   r"   r    r`   r`   s   sy        

C	%(EHeh//	%(EHeh//	%(EHeh//	%(EHeh//	EEEr"   r`   c            	          e Zd Zej        Z e ej        ej	        ej
        f          ej	        ej	                   e ej        ej	        ej
        f          ej	        ej                   e ej        ej	        ej
        f          ej	        ej                   e ej        ej	        ej
        f          ej	        ej                  gZdS )Cuda_match_all_syncN)r3   r4   r5   r   match_all_syncr:   r   r   r(   rT   rX   rY   rZ   r[   rJ   r   r"   r    rc   rc   ~   s        

C	+%+ux233UXuxHH	+%+ux233UXuxHH	+%+ux233UXuxHH	+%+ux233UXuxHH	EEEr"   rc   c                   >    e Zd Zej        Z eej                  gZ	dS )Cuda_activemaskN)
r3   r4   r5   r   
activemaskr:   r   r   uint32rJ   r   r"   r    rf   rf      s)        
/CYu|$$%EEEr"   rf   c                   >    e Zd Zej        Z eej                  gZ	dS )Cuda_lanemask_ltN)
r3   r4   r5   r   lanemask_ltr:   r   r   rh   rJ   r   r"   r    rj   rj      s*        

CYu|$$%EEEr"   rj   c                   t   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)r3   r4   r5   __doc__r   popcr:   r   r   int8int16int32int64uint8uint16rh   uint64rJ   r   r"   r    rm   rm      s          )C	%*ej))	%+u{++	%+u{++	%+u{++	%+u{++	%,--	%,--	%,--	EEEr"   rm   c                       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)r3   r4   r5   rn   r   fmar:   r   r   float32float64rJ   r   r"   r    rx   rx      s^          (C	%-u}MM	%-u}MMEEEr"   rx   c                   l    e Zd Zej        j        Z eej	        ej	        ej	        ej	                  gZ
dS )	Cuda_hfmaN)r3   r4   r5   r   fp16hfmar:   r   r   float16rJ   r   r"   r    r}   r}      s9        
).C	%-u}MMEEEr"   r}   c                   t    e Zd Zej        Z eej        ej                   eej	        ej	                  gZ
dS )	Cuda_cbrtN)r3   r4   r5   r   cbrtr:   r   r   rz   r{   rJ   r   r"   r    r   r      sD         )C	%-//	%-//EEEr"   r   c                   t    e Zd Zej        Z eej        ej                   eej	        ej	                  gZ
dS )	Cuda_brevN)r3   r4   r5   r   brevr:   r   r   rh   rv   rJ   r   r"   r    r   r      sB        
)C	%,--	%,--EEEr"   r   c                   t   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)r3   r4   r5   rn   r   clzr:   r   r   rp   rq   rr   rs   rt   ru   rh   rv   rJ   r   r"   r    r   r      s          (C	%*ej))	%+u{++	%+u{++	%+u{++	%+u{++	%,--	%,--	%,--	EEEr"   r   c                   t   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)r3   r4   r5   rn   r   ffsr:   r   r   rh   rp   rq   rr   rs   rt   ru   rv   rJ   r   r"   r    r   r      s          (C	%,
++	%,,,	%,,,	%,,,	%,,,	%,--	%,--	%,--	EEEr"   r   c                   "    e Zd Zej        Zd ZdS )	Cuda_selpc                    |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 rB   )
r   r{   rz   rq   ru   rr   rh   rs   rv   r   )r0   argskwstestabsupported_typess          r    r1   zCuda_selp.generic   sq    
a !=%- ; ; ;6
 66Qo--FD!Q'''r"   N)r3   r4   r5   r   selpr:   r1   r   r"   r    r   r      s*        
)C( ( ( ( (r"   r   c                 L     t            G  fddt                                }|S )Nc                   B    e Zd Z Z eej        ej                  gZdS )'_genfp16_unary.<locals>.Cuda_fp16_unaryNr3   r4   r5   r:   r   r   r   rJ   l_keys   r    Cuda_fp16_unaryr     s,        5=%-889r"   r   registerr   r   r   s   ` r    _genfp16_unaryr     sJ    : : : : : : :* : : X: r"   c                 \     t                      G  fddt                                }|S )Nc                       e Zd Z Zd ZdS )0_genfp16_unary_operator.<locals>.Cuda_fp16_unaryc                     |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   )r0   r   r   s      r    r1   z8_genfp16_unary_operator.<locals>.Cuda_fp16_unary.generic  sK    NNN4yyA~~$q'U]":": >>> ~":":r"   Nr3   r4   r5   r:   r1   r   s   r    r   r     s)        	? 	? 	? 	? 	?r"   r   register_globalr   r   s   ` r    _genfp16_unary_operatorr     sS    U? ? ? ? ? ? ?* ? ? ? r"   c                 L     t            G  fddt                                }|S )Nc                   N    e Zd Z Z eej        ej        ej                  gZdS ))_genfp16_binary.<locals>.Cuda_fp16_binaryNr   r   s   r    Cuda_fp16_binaryr   "  s0        5=%-GGHr"   r   r   )r   r   s   ` r    _genfp16_binaryr   !  sT    I I I I I I I+ I I XI r"   c                       e Zd Zd ZdS )Floatc                 V    |rJ |\  }|t           j        k    rt          ||          S d S rB   )r   r   r   )r0   r   r   args       r    r1   zFloat.generic-  s7    %-S#&&&  r"   Nr2   r   r"   r    r   r   *  s#        ' ' ' ' 'r"   r   c                 L     t            G  fddt                                }|S )Nc                   N    e Zd Z Z eej        ej        ej                  gZdS )1_genfp16_binary_comparison.<locals>.Cuda_fp16_cmpN)	r3   r4   r5   r:   r   r   rX   r   rJ   r   s   r    Cuda_fp16_cmpr   7  s4         Iehu}==
r"   r   r   )r   r   s   ` r    _genfp16_binary_comparisonr   6  sJ    
 
 
 
 
 
 
( 
 
 X
 r"   c                 `     t                      G  fddt                                }|S )Nc                       e Zd Z ZfdZdS )1_fp16_binary_operator.<locals>.Cuda_fp16_operatorc                 
   |rJ t          |          dk    r|d         t          j        k    s|d         t          j        k    r|d         t          j        k    r(| j                            |d         |d                   }n'| j                            |d         |d                   }|t
          j        k    s |t
          j        k    s|t
          j        k    r)t          t          j        t          j                  S d S d S d S )N   r   r   )
r   r   r   contextcan_convertr   exactpromotesafer   )r0   r   r   convertiblerettys       r    r1   z9_fp16_binary_operator.<locals>.Cuda_fp16_operator.genericS  s    NNN4yyA~~!W--aEM1I1IGu},,"&,":":47DG"L"LKK"&,":":47DG"L"LK  :#333:#555:?22$UEM5=III% ~1I1I  32r"   Nr   )r   r   s   r    Cuda_fp16_operatorr   O  s:        	J 	J 	J 	J 	J 	J 	Jr"   r   r   )r   r   r   s   `` r    _fp16_binary_operatorr   N  sd    UJ J J J J J J J- J J J4 r"   c                 6    t          | t          j                  S rB   )r   r   rX   ops    r    _genfp16_comparison_operatorr   m  s     UX...r"   c                 6    t          | t          j                  S rB   )r   r   r   r   s    r    _genfp16_binary_operatorr   q  s     U]333r"   c                 |    t          d|  t          j        t          j        f          }t          j        |          S N__numba_wrapper_r   r   r   Functionfnamedecls     r    _resolve_wrapped_unaryr     s;    +,Fu,F,F,1M-2],<> >D >$r"   c                     t          d|  t          j        t          j        t          j        f          }t          j        |          S r   r   r   s     r    _resolve_wrapped_binaryr     sA    +,Fu,F,F,1M-2]EM,KM MD >$r"   hsinhcoshloghlog10hlog2hexphexp10hexp2hsqrthrsqrthfloorhceilhrcphrinthtrunchdivc                 P     t            G  fddt                                }|S )Nc                       e Zd Z ZfdZdS )_gen.<locals>.Cuda_atomicc                     |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)r0   r   r   aryidxvalr   s         r    r1   z!_gen.<locals>.Cuda_atomic.generic  sx    NNN MCcy//x1}} CSYGGGA Cci@@@ r"   Nr   )r   r   s   r    Cuda_atomicr     s:        
	A 
	A 
	A 
	A 
	A 
	A 
	Ar"   r   )r   r   )r   r   r   s   `` r    _genr     s[    A A A A A A A A& A A XA r"   c                   ,    e Zd Zej        j        Zd ZdS )Cuda_atomic_compare_and_swapc                 x    |rJ |\  }}}|j         }|t          v r|j        dk    rt          ||||          S d S d S r   )r$   integer_numba_typesr%   r   )r0   r   r   r   oldr   dtys          r    r1   z$Cuda_atomic_compare_and_swap.generic  sS    S#i%%%#(a--S#sC000 &%--r"   N)r3   r4   r5   r   atomiccompare_and_swapr:   r1   r   r"   r    r   r     s-        
+
&C1 1 1 1 1r"   r   c                   ,    e Zd Zej        j        Zd ZdS )Cuda_atomic_casc                     |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   )r0   r   r   r   r   r   r   r   s           r    r1   zCuda_atomic_cas.generic  s|    !S#si)))F8q==S#uz3<<<X\\S#sC555 \r"   N)r3   r4   r5   r   r   casr:   r1   r   r"   r    r   r     s,        
+/C6 6 6 6 6r"   r   c                   J    e Zd Zej        Z eej        ej	                  gZ
dS )Cuda_nanosleepN)r3   r4   r5   r   	nanosleepr:   r   r   voidrh   rJ   r   r"   r    r  r    s-        
.CYuz5<001EEEr"   r  c                   $    e Zd ZeZd Zd Zd ZdS )
Dim3_attrsc                     t           j        S rB   r   rr   r0   mods     r    	resolve_xzDim3_attrs.resolve_x	  
    {r"   c                     t           j        S rB   r  r  s     r    	resolve_yzDim3_attrs.resolve_y  r  r"   c                     t           j        S rB   r  r  s     r    	resolve_zzDim3_attrs.resolve_z  r  r"   N)r3   r4   r5   r   r:   r
  r  r  r   r"   r    r  r    sF        
C        r"   r  c                   >    e Zd Z ej        ej                  Zd ZdS )CudaSharedModuleTemplatec                 4    t          j        t                    S rB   )r   r   r7   r  s     r    resolve_arrayz&CudaSharedModuleTemplate.resolve_array  s    ~/000r"   N)	r3   r4   r5   r   Moduler   r8   r:   r  r   r"   r    r  r    s6        
%,t{
#
#C1 1 1 1 1r"   r  c                   >    e Zd Z ej        ej                  Zd ZdS )CudaConstModuleTemplatec                 4    t          j        t                    S rB   )r   r   r?   r  s     r    resolve_array_likez*CudaConstModuleTemplate.resolve_array_like  s    ~3444r"   N)	r3   r4   r5   r   r  r   rD   r:   r  r   r"   r    r  r    s6        
%,tz
"
"C5 5 5 5 5r"   r  c                   >    e Zd Z ej        ej                  Zd ZdS )CudaLocalModuleTemplatec                 4    t          j        t                    S rB   )r   r   r<   r  s     r    r  z%CudaLocalModuleTemplate.resolve_array'      ~.///r"   N)	r3   r4   r5   r   r  r   r=   r:   r  r   r"   r    r  r  #  s6        
%,tz
"
"C0 0 0 0 0r"   r  c                       e Zd Z ej        ej                  Zd Zd Z	d Z
d Zd Zd Zd Zd Zd	 Zd
 Zd Zd Zd Zd ZdS )CudaAtomicTemplatec                 4    t          j        t                    S rB   )r   r   Cuda_atomic_addr  s     r    resolve_addzCudaAtomicTemplate.resolve_add/      ~o...r"   c                 4    t          j        t                    S rB   )r   r   Cuda_atomic_subr  s     r    resolve_subzCudaAtomicTemplate.resolve_sub2  r"  r"   c                 4    t          j        t                    S rB   )r   r   Cuda_atomic_andr  s     r    resolve_and_zCudaAtomicTemplate.resolve_and_5  r"  r"   c                 4    t          j        t                    S rB   )r   r   Cuda_atomic_orr  s     r    resolve_or_zCudaAtomicTemplate.resolve_or_8      ~n---r"   c                 4    t          j        t                    S rB   )r   r   Cuda_atomic_xorr  s     r    resolve_xorzCudaAtomicTemplate.resolve_xor;  r"  r"   c                 4    t          j        t                    S rB   )r   r   Cuda_atomic_incr  s     r    resolve_inczCudaAtomicTemplate.resolve_inc>  r"  r"   c                 4    t          j        t                    S rB   )r   r   Cuda_atomic_decr  s     r    resolve_deczCudaAtomicTemplate.resolve_decA  r"  r"   c                 4    t          j        t                    S rB   )r   r   Cuda_atomic_exchr  s     r    resolve_exchzCudaAtomicTemplate.resolve_exchD  r  r"   c                 4    t          j        t                    S rB   )r   r   Cuda_atomic_maxr  s     r    resolve_maxzCudaAtomicTemplate.resolve_maxG  r"  r"   c                 4    t          j        t                    S rB   )r   r   Cuda_atomic_minr  s     r    resolve_minzCudaAtomicTemplate.resolve_minJ  r"  r"   c                 4    t          j        t                    S rB   )r   r   Cuda_atomic_nanminr  s     r    resolve_nanminz!CudaAtomicTemplate.resolve_nanminM      ~0111r"   c                 4    t          j        t                    S rB   )r   r   Cuda_atomic_nanmaxr  s     r    resolve_nanmaxz!CudaAtomicTemplate.resolve_nanmaxP  rB  r"   c                 4    t          j        t                    S rB   )r   r   r   r  s     r    resolve_compare_and_swapz+CudaAtomicTemplate.resolve_compare_and_swapS  s    ~:;;;r"   c                 4    t          j        t                    S rB   )r   r   r   r  s     r    resolve_caszCudaAtomicTemplate.resolve_casV  r"  r"   N)r3   r4   r5   r   r  r   r   r:   r!  r%  r(  r+  r/  r2  r5  r8  r;  r>  rA  rE  rG  rI  r   r"   r    r  r  +  s        
%,t{
#
#C/ / // / // / /. . ./ / // / // / /0 0 0/ / // / /2 2 22 2 2< < </ / / / /r"   r  c                       e Zd Z ej        ej                  Zd Zd Z	d Z
d Zd Zd Zd Zd Zd	 Zd
 Zd Zd Zd Zd Zd Zd Zd Zd Zd Zd Zd Zd Zd Zd Zd Z d Z!d Z"d Z#d Z$d Z%dS ) CudaFp16Templatec                 4    t          j        t                    S rB   )r   r   	Cuda_haddr  s     r    resolve_haddzCudaFp16Template.resolve_hadd^      ~i(((r"   c                 4    t          j        t                    S rB   )r   r   	Cuda_hsubr  s     r    resolve_hsubzCudaFp16Template.resolve_hsuba  rO  r"   c                 4    t          j        t                    S rB   )r   r   	Cuda_hmulr  s     r    resolve_hmulzCudaFp16Template.resolve_hmuld  rO  r"   c                     t           S rB   )hdiv_devicer  s     r    resolve_hdivzCudaFp16Template.resolve_hdivg      r"   c                 4    t          j        t                    S rB   )r   r   	Cuda_hnegr  s     r    resolve_hnegzCudaFp16Template.resolve_hnegj  rO  r"   c                 4    t          j        t                    S rB   )r   r   	Cuda_habsr  s     r    resolve_habszCudaFp16Template.resolve_habsm  rO  r"   c                 4    t          j        t                    S rB   )r   r   r}   r  s     r    resolve_hfmazCudaFp16Template.resolve_hfmap  rO  r"   c                     t           S rB   )hsin_devicer  s     r    resolve_hsinzCudaFp16Template.resolve_hsins  rY  r"   c                     t           S rB   )hcos_devicer  s     r    resolve_hcoszCudaFp16Template.resolve_hcosv  rY  r"   c                     t           S rB   )hlog_devicer  s     r    resolve_hlogzCudaFp16Template.resolve_hlogy  rY  r"   c                     t           S rB   )hlog10_devicer  s     r    resolve_hlog10zCudaFp16Template.resolve_hlog10|      r"   c                     t           S rB   )hlog2_devicer  s     r    resolve_hlog2zCudaFp16Template.resolve_hlog2      r"   c                     t           S rB   )hexp_devicer  s     r    resolve_hexpzCudaFp16Template.resolve_hexp  rY  r"   c                     t           S rB   )hexp10_devicer  s     r    resolve_hexp10zCudaFp16Template.resolve_hexp10  rn  r"   c                     t           S rB   )hexp2_devicer  s     r    resolve_hexp2zCudaFp16Template.resolve_hexp2  rr  r"   c                     t           S rB   )hfloor_devicer  s     r    resolve_hfloorzCudaFp16Template.resolve_hfloor  rn  r"   c                     t           S rB   )hceil_devicer  s     r    resolve_hceilzCudaFp16Template.resolve_hceil  rr  r"   c                     t           S rB   )hsqrt_devicer  s     r    resolve_hsqrtzCudaFp16Template.resolve_hsqrt  rr  r"   c                     t           S rB   )hrsqrt_devicer  s     r    resolve_hrsqrtzCudaFp16Template.resolve_hrsqrt  rn  r"   c                     t           S rB   )hrcp_devicer  s     r    resolve_hrcpzCudaFp16Template.resolve_hrcp  rY  r"   c                     t           S rB   )hrint_devicer  s     r    resolve_hrintzCudaFp16Template.resolve_hrint  rr  r"   c                     t           S rB   )htrunc_devicer  s     r    resolve_htrunczCudaFp16Template.resolve_htrunc  rn  r"   c                 4    t          j        t                    S rB   )r   r   Cuda_heqr  s     r    resolve_heqzCudaFp16Template.resolve_heq      ~h'''r"   c                 4    t          j        t                    S rB   )r   r   Cuda_hner  s     r    resolve_hnezCudaFp16Template.resolve_hne  r  r"   c                 4    t          j        t                    S rB   )r   r   Cuda_hger  s     r    resolve_hgezCudaFp16Template.resolve_hge  r  r"   c                 4    t          j        t                    S rB   )r   r   Cuda_hgtr  s     r    resolve_hgtzCudaFp16Template.resolve_hgt  r  r"   c                 4    t          j        t                    S rB   )r   r   Cuda_hler  s     r    resolve_hlezCudaFp16Template.resolve_hle  r  r"   c                 4    t          j        t                    S rB   )r   r   Cuda_hltr  s     r    resolve_hltzCudaFp16Template.resolve_hlt  r  r"   c                 4    t          j        t                    S rB   )r   r   	Cuda_hmaxr  s     r    resolve_hmaxzCudaFp16Template.resolve_hmax  rO  r"   c                 4    t          j        t                    S rB   )r   r   	Cuda_hminr  s     r    resolve_hminzCudaFp16Template.resolve_hmin  rO  r"   N)&r3   r4   r5   r   r  r   r~   r:   rN  rR  rU  rX  r\  r_  ra  rd  rg  rj  rm  rq  ru  rx  r{  r~  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r   r"   r    rK  rK  Z  s       
%,ty
!
!C) ) )) ) )) ) )  ) ) )) ) )) ) )                              ( ( (( ( (( ( (( ( (( ( (( ( () ) )) ) ) ) )r"   rK  c                       e Zd Z ej        e          Zd Zd Zd Z	d Z
d Zd Zd Zd Zd	 Zd
 Zd Zd Zd Zd Zd Zd Zd Zd Zd Zd Zd Zd Zd Zd Zd Zd Z d Z!d Z"d Z#dS )CudaModuleTemplatec                 >    t          j        t          j                  S rB   )r   r  r   cgr  s     r    
resolve_cgzCudaModuleTemplate.resolve_cg  s    |DG$$$r"   c                     t           S rB   r   r  s     r    resolve_threadIdxz$CudaModuleTemplate.resolve_threadIdx      r"   c                     t           S rB   r   r  s     r    resolve_blockIdxz#CudaModuleTemplate.resolve_blockIdx  r  r"   c                     t           S rB   r   r  s     r    resolve_blockDimz#CudaModuleTemplate.resolve_blockDim  r  r"   c                     t           S rB   r   r  s     r    resolve_gridDimz"CudaModuleTemplate.resolve_gridDim  r  r"   c                     t           j        S rB   r  r  s     r    resolve_laneidz!CudaModuleTemplate.resolve_laneid  r  r"   c                 >    t          j        t          j                  S rB   )r   r  r   r8   r  s     r    resolve_sharedz!CudaModuleTemplate.resolve_shared      |DK(((r"   c                 4    t          j        t                    S rB   )r   r   rm   r  s     r    resolve_popczCudaModuleTemplate.resolve_popc  rO  r"   c                 4    t          j        t                    S rB   )r   r   r   r  s     r    resolve_brevzCudaModuleTemplate.resolve_brev  rO  r"   c                 4    t          j        t                    S rB   )r   r   r   r  s     r    resolve_clzzCudaModuleTemplate.resolve_clz  r  r"   c                 4    t          j        t                    S rB   )r   r   r   r  s     r    resolve_ffszCudaModuleTemplate.resolve_ffs  r  r"   c                 4    t          j        t                    S rB   )r   r   rx   r  s     r    resolve_fmazCudaModuleTemplate.resolve_fma  r  r"   c                 4    t          j        t                    S rB   )r   r   r   r  s     r    resolve_cbrtzCudaModuleTemplate.resolve_cbrt  rO  r"   c                 4    t          j        t                    S rB   )r   r   rG   r  s     r    resolve_threadfencez&CudaModuleTemplate.resolve_threadfence      ~5666r"   c                 4    t          j        t                    S rB   )r   r   rL   r  s     r    resolve_threadfence_blockz,CudaModuleTemplate.resolve_threadfence_block  s    ~4555r"   c                 4    t          j        t                    S rB   )r   r   rO   r  s     r    resolve_threadfence_systemz-CudaModuleTemplate.resolve_threadfence_system  r  r"   c                 4    t          j        t                    S rB   )r   r   rR   r  s     r    resolve_syncwarpz#CudaModuleTemplate.resolve_syncwarp  s    ~m,,,r"   c                 4    t          j        t                    S rB   )r   r   rV   r  s     r    resolve_shfl_sync_intrinsicz.CudaModuleTemplate.resolve_shfl_sync_intrinsic      ~6777r"   c                 4    t          j        t                    S rB   )r   r   r]   r  s     r    resolve_vote_sync_intrinsicz.CudaModuleTemplate.resolve_vote_sync_intrinsic  r  r"   c                 4    t          j        t                    S rB   )r   r   r`   r  s     r    resolve_match_any_syncz)CudaModuleTemplate.resolve_match_any_sync      ~1222r"   c                 4    t          j        t                    S rB   )r   r   rc   r  s     r    resolve_match_all_syncz)CudaModuleTemplate.resolve_match_all_sync  r  r"   c                 4    t          j        t                    S rB   )r   r   rf   r  s     r    resolve_activemaskz%CudaModuleTemplate.resolve_activemask  r"  r"   c                 4    t          j        t                    S rB   )r   r   rj   r  s     r    resolve_lanemask_ltz&CudaModuleTemplate.resolve_lanemask_lt  r  r"   c                 4    t          j        t                    S rB   )r   r   r   r  s     r    resolve_selpzCudaModuleTemplate.resolve_selp  rO  r"   c                 4    t          j        t                    S rB   )r   r   r  r  s     r    resolve_nanosleepz$CudaModuleTemplate.resolve_nanosleep  r,  r"   c                 >    t          j        t          j                  S rB   )r   r  r   r   r  s     r    resolve_atomicz!CudaModuleTemplate.resolve_atomic  r  r"   c                 >    t          j        t          j                  S rB   )r   r  r   r~   r  s     r    resolve_fp16zCudaModuleTemplate.resolve_fp16  s    |DI&&&r"   c                 >    t          j        t          j                  S rB   )r   r  r   rD   r  s     r    resolve_constz CudaModuleTemplate.resolve_const      |DJ'''r"   c                 >    t          j        t          j                  S rB   )r   r  r   r=   r  s     r    resolve_localz CudaModuleTemplate.resolve_local  r  r"   N)$r3   r4   r5   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       
%,t

C% % %          ) ) )) ) )) ) )( ( (( ( (( ( () ) )7 7 76 6 67 7 7- - -8 8 88 8 83 3 33 3 3/ / /0 0 0) ) ). . .) ) )' ' '( ( (( ( ( ( (r"   r  )operator
numba.corer   numba.core.typing.npydeclr   r   r   r   r   r	   r
   numba.core.typing.templatesr   r   r   r   r   r   numba.cuda.typesr   numba.core.typeconvr   numbar   numba.cuda.compilerr   registryr   register_attrr   r   r7   r<   r?   rG   rL   rO   rR   rV   r]   r`   rc   rf   rj   rm   rx   r}   r   r   r   r   r   r   r   r   floatr   r   r   r   r   r~   haddrM  addCuda_addiadd	Cuda_iaddhsubrQ  subCuda_subisub	Cuda_isubhmulrT  mulCuda_mulimul	Cuda_imulhmaxr  hminr  hnegr[  negCuda_neghabsr^  absCuda_absheqr  eqhner  nehger  gehgtr  gthler  lehltr  lttruedivitruedivr   r   rc  rf  ri  rl  rp  rt  rw  rz  r  r  r}  r  r  r  r  rW  r   r{   rz   rr   rh   rs   rv   all_numba_typesr   unsigned_int_numba_typesr   r   r$  maxr:  minr=  nanmaxrD  nanminr@  and_r'  or_r*  xorr.  incr1  decr4  exchr7  r   r   r  r  r  r  r  r  rK  r  r  funcr   r"   r    <module>r*     s=         @ @ @ @ @ @ @ @ @ @ @ @ @ @ @ @ @ @> > > > > > > > > > > > > > > > " ! ! ! ! ! * * * * * *       @ @ @ @ @ @8::&*   ( ( (    &   0 
       
 
       
 
    ,   
 
$ $ $ $ $. $ $ 
$
 
$ $ $ $ $- $ $ 
$
 
$ $ $ $ $. $ $ 
$
 
E E E E E$ E E 
E
 
    /   
 
6 6 6 6 6/ 6 6 
6 
    *   
 
    *   
 
& & & & && & & 
&
 
& & & & &' & & 
&
 
        
$ 
	 	 	 	 	 	 	 
	 
        
 
        
 
        
 
       
$ 
       
$ 
( ( ( ( (  ( ( 
((  
 
 
   ' ' ' ' ' ' ' '  0  >/ / /4 4 4 ODIN++	##HL11$$X]33	ODIN++	##HL11$$X]33	ODIN++	##HL11$$X]33	ODIN++	ODIN++	N49>**	""8<00N49>**	""3''%%dim44  X[ ) ) )%%dim44  X[ ) ) )%%dim44  X[ ) ) )%%dim44  X[ ) ) )%%dim44  X[ ) ) )%%dim44  X[ ) ) )  ) * * *  * + + +           %$V,,$$V,,$$V,,&&x00%%g..$$V,,&&x00%%g..%%g..&&x00&&x00%%g..$$V,,%%g..&&x00%%f--  & =%-;;. {EL{EL2  "L%,7 $t{88$t{88$t{88$t{88T$+,o>> T$+,o>> $t{')<==dko':;;$t{(;<<$t{(@AA$t{(@AA4(*=>>  
	1 	1 	1 	1 	1#3 	1 	1 
	1 
6 6 6 6 6& 6 6 
6" 
2 2 2 2 2% 2 2 
2 
 
 
 
 
" 
 
 
 1 1 1 1 10 1 1 1 5 5 5 5 5/ 5 5 5 0 0 0 0 0/ 0 0 0 +/ +/ +/ +/ +/* +/ +/ +/\ [) [) [) [) [)( [) [) [)| X( X( X( X( X(* X( X( X(v lel4(( ) ) )
 $ 0 0D////  0 0D////# 0 0D////0 0r"   