
    J/PhZ                        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
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                      Z0e G d- d.e                      Z1e G d/ d0e                      Z2e G d1 d2e                      Z3e G d3 d4e                      Z4d5 Z5d6 Z6d7 Z7 ee8           G d8 d9e                      Z9d: Z:d; Z;d< Z<d= Z= e7ej>        j?                  Z@ e=e jA                  ZB e=e jC                  ZD e7ej>        jE                  ZF e=e jG                  ZH e=e jI                  ZJ e7ej>        jK                  ZL e=e jM                  ZN e=e jO                  ZP e7ej>        jQ                  ZR e7ej>        jS                  ZT e5ej>        jU                  ZV e6e jW                  ZX e5ej>        jY                  ZZ e6e[          Z\ e:ej>        j]                  Z^ e<e j_                    e:ej>        j`                  Za e<e jb                    e:ej>        jc                  Zd e<e je                    e:ej>        jf                  Zg e<e jh                    e:ej>        ji                  Zj e<e jk                    e:ej>        jl                  Zm e<e jn                    e=e jo                    e=e jp                   d> Zqd? Zr eqd@          Zs eqdA          Zt eqdB          Zu eqdC          Zv eqdD          Zw eqdE          Zx eqdF          Zy eqdG          Zz eqdH          Z{ eqdI          Z| eqdJ          Z} eqdK          Z~ eqdL          Z eqdM          Z eqdN          Z erdO          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        jA        e          Z eej        jG        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           e
D ]Zedev r eee           dS )f    N)types)parse_dtypeparse_shaperegister_number_classesregister_numpy_ufunctrigonometric_functionscomparison_functionsmath_operations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     S/var/www/html/test/jupyter/venv/lib/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__r2   r   r#   r!   r   r      s#            r#   r   c                   &    e Zd Zej        j        ZdS )Cuda_shared_arrayN)r4   r5   r6   r   sharedarraykeyr   r#   r!   r8   r8   2   s        
+
CCCr#   r8   c                   &    e Zd Zej        j        ZdS )Cuda_local_arrayN)r4   r5   r6   r   localr:   r;   r   r#   r!   r=   r=   7   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>.typerA   s    Nr#   r   r0   s     r!   r2   zCuda_const_array_like.generic@   s    	 	 	r#   N)r4   r5   r6   r   const
array_liker;   r2   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)
r4   r5   r6   r   threadfencer;   r   r   nonecasesr   r#   r!   rH   rH   F   s*        

CYuz""#EEEr#   rH   c                   >    e Zd Zej        Z eej                  gZ	dS )Cuda_threadfence_blockN)
r4   r5   r6   r   threadfence_blockr;   r   r   rJ   rK   r   r#   r!   rM   rM   L   s*        

 CYuz""#EEEr#   rM   c                   >    e Zd Zej        Z eej                  gZ	dS )Cuda_threadfence_systemN)
r4   r5   r6   r   threadfence_systemr;   r   r   rJ   rK   r   r#   r!   rP   rP   R   s*        

!CYuz""#EEEr#   rP   c                   h    e Zd Zej        Z eej                   eej        ej	                  gZ
dS )Cuda_syncwarpN)r4   r5   r6   r   syncwarpr;   r   r   rJ   i4rK   r   r#   r!   rS   rS   X   s;        
-CYuz""IIej%($C$CDEEEr#   rS   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)r4   r5   r6   r   shfl_sync_intrinsicr;   r   r   r)   rU   b1i8f4f8rK   r   r#   r!   rW   rW   ^   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#   rW   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)r4   r5   r6   r   vote_sync_intrinsicr;   r   r   r)   rU   rY   rK   r   r#   r!   r^   r^   m   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)r4   r5   r6   r   match_any_syncr;   r   r   rU   rZ   r[   r\   rK   r   r#   r!   ra   ra   t   sy        

C	%(EHeh//	%(EHeh//	%(EHeh//	%(EHeh//	EEEr#   ra   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)r4   r5   r6   r   match_all_syncr;   r   r   r)   rU   rY   rZ   r[   r\   rK   r   r#   r!   rd   rd      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#   rd   c                   >    e Zd Zej        Z eej                  gZ	dS )Cuda_activemaskN)
r4   r5   r6   r   
activemaskr;   r   r   uint32rK   r   r#   r!   rg   rg      s)        
/CYu|$$%EEEr#   rg   c                   >    e Zd Zej        Z eej                  gZ	dS )Cuda_lanemask_ltN)
r4   r5   r6   r   lanemask_ltr;   r   r   ri   rK   r   r#   r!   rk   rk      s*        

CYu|$$%EEEr#   rk   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)r4   r5   r6   __doc__r   popcr;   r   r   int8int16int32int64uint8uint16ri   uint64rK   r   r#   r!   rn   rn      s          )C	%*ej))	%+u{++	%+u{++	%+u{++	%+u{++	%,--	%,--	%,--	EEEr#   rn   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)r4   r5   r6   ro   r   fmar;   r   r   float32float64rK   r   r#   r!   ry   ry      s^          (C	%-u}MM	%-u}MMEEEr#   ry   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)r4   r5   r6   r   fp16hfmar;   r   r   float16rK   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)r4   r5   r6   r   cbrtr;   r   r   r{   r|   rK   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)r4   r5   r6   r   brevr;   r   r   ri   rw   rK   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)r4   r5   r6   ro   r   clzr;   r   r   rq   rr   rs   rt   ru   rv   ri   rw   rK   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)r4   r5   r6   ro   r   ffsr;   r   r   ri   rq   rr   rs   rt   ru   rv   rw   rK   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 rC   )
r   r|   r{   rr   rv   rs   ri   rt   rw   r   )r1   argskwstestabsupported_typess          r!   r2   zCuda_selp.generic   sq    
a !=%- ; ; ;6
 66Qo--FD!Q'''r#   N)r4   r5   r6   r   selpr;   r2   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r4   r5   r6   r;   r   r   r   rK   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   )r1   r   r   s      r!   r2   z8_genfp16_unary_operator.<locals>.Cuda_fp16_unary.generic  sK    NNN4yyA~~$q'U]":": >>> ~":":r#   Nr4   r5   r6   r;   r2   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 rC   )r   r   r   )r1   r   r   args       r!   r2   zFloat.generic.  s7    %-S#&&&  r#   Nr3   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)	r4   r5   r6   r;   r   r   rY   r   rK   r   s   r!   Cuda_fp16_cmpr   8  s4         Iehu}==
r#   r   r   )r   r   s   ` r!   _genfp16_binary_comparisonr   7  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   )r1   r   r   convertiblerettys       r!   r2   z9_fp16_binary_operator.<locals>.Cuda_fp16_operator.genericT  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   P  s:        	J 	J 	J 	J 	J 	J 	Jr#   r   r   )r   r   r   s   `` r!   _fp16_binary_operatorr   O  sd    UJ J J J J J J J- J J J4 r#   c                 6    t          | t          j                  S rC   )r   r   rY   ops    r!   _genfp16_comparison_operatorr   n  s     UX...r#   c                 6    t          | t          j                  S rC   )r   r   r   r   s    r!   _genfp16_binary_operatorr   r  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)r1   r   r   aryidxvalr   s         r!   r2   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   )r1   r   r   r   oldr   dtys          r!   r2   z$Cuda_atomic_compare_and_swap.generic  sS    S#i%%%#(a--S#sC000 &%--r#   N)r4   r5   r6   r   atomiccompare_and_swapr;   r2   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   )r1   r   r   r   r   r   r   r   s           r!   r2   zCuda_atomic_cas.generic  s|    !S#si)))F8q==S#uz3<<<X\\S#sC555 \r#   N)r4   r5   r6   r   r   casr;   r2   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)r4   r5   r6   r   	nanosleepr;   r   r   voidri   rK   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 rC   r   rs   r1   mods     r!   	resolve_xzDim3_attrs.resolve_x
  
    {r#   c                     t           j        S rC   r  r	  s     r!   	resolve_yzDim3_attrs.resolve_y  r  r#   c                     t           j        S rC   r  r	  s     r!   	resolve_zzDim3_attrs.resolve_z  r  r#   N)r4   r5   r6   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 rC   )r   r   r8   r	  s     r!   resolve_arrayz&CudaSharedModuleTemplate.resolve_array  s    ~/000r#   N)	r4   r5   r6   r   Moduler   r9   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 rC   )r   r   r@   r	  s     r!   resolve_array_likez*CudaConstModuleTemplate.resolve_array_like   s    ~3444r#   N)	r4   r5   r6   r   r  r   rE   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 rC   )r   r   r=   r	  s     r!   r  z%CudaLocalModuleTemplate.resolve_array(      ~.///r#   N)	r4   r5   r6   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 rC   )r   r   Cuda_atomic_addr	  s     r!   resolve_addzCudaAtomicTemplate.resolve_add0      ~o...r#   c                 4    t          j        t                    S rC   )r   r   Cuda_atomic_subr	  s     r!   resolve_subzCudaAtomicTemplate.resolve_sub3  r#  r#   c                 4    t          j        t                    S rC   )r   r   Cuda_atomic_andr	  s     r!   resolve_and_zCudaAtomicTemplate.resolve_and_6  r#  r#   c                 4    t          j        t                    S rC   )r   r   Cuda_atomic_orr	  s     r!   resolve_or_zCudaAtomicTemplate.resolve_or_9      ~n---r#   c                 4    t          j        t                    S rC   )r   r   Cuda_atomic_xorr	  s     r!   resolve_xorzCudaAtomicTemplate.resolve_xor<  r#  r#   c                 4    t          j        t                    S rC   )r   r   Cuda_atomic_incr	  s     r!   resolve_inczCudaAtomicTemplate.resolve_inc?  r#  r#   c                 4    t          j        t                    S rC   )r   r   Cuda_atomic_decr	  s     r!   resolve_deczCudaAtomicTemplate.resolve_decB  r#  r#   c                 4    t          j        t                    S rC   )r   r   Cuda_atomic_exchr	  s     r!   resolve_exchzCudaAtomicTemplate.resolve_exchE  r  r#   c                 4    t          j        t                    S rC   )r   r   Cuda_atomic_maxr	  s     r!   resolve_maxzCudaAtomicTemplate.resolve_maxH  r#  r#   c                 4    t          j        t                    S rC   )r   r   Cuda_atomic_minr	  s     r!   resolve_minzCudaAtomicTemplate.resolve_minK  r#  r#   c                 4    t          j        t                    S rC   )r   r   Cuda_atomic_nanminr	  s     r!   resolve_nanminz!CudaAtomicTemplate.resolve_nanminN      ~0111r#   c                 4    t          j        t                    S rC   )r   r   Cuda_atomic_nanmaxr	  s     r!   resolve_nanmaxz!CudaAtomicTemplate.resolve_nanmaxQ  rC  r#   c                 4    t          j        t                    S rC   )r   r   r   r	  s     r!   resolve_compare_and_swapz+CudaAtomicTemplate.resolve_compare_and_swapT  s    ~:;;;r#   c                 4    t          j        t                    S rC   )r   r   r   r	  s     r!   resolve_caszCudaAtomicTemplate.resolve_casW  r#  r#   N)r4   r5   r6   r   r  r   r   r;   r"  r&  r)  r,  r0  r3  r6  r9  r<  r?  rB  rF  rH  rJ  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 rC   )r   r   	Cuda_haddr	  s     r!   resolve_haddzCudaFp16Template.resolve_hadd_      ~i(((r#   c                 4    t          j        t                    S rC   )r   r   	Cuda_hsubr	  s     r!   resolve_hsubzCudaFp16Template.resolve_hsubb  rP  r#   c                 4    t          j        t                    S rC   )r   r   	Cuda_hmulr	  s     r!   resolve_hmulzCudaFp16Template.resolve_hmule  rP  r#   c                     t           S rC   )hdiv_devicer	  s     r!   resolve_hdivzCudaFp16Template.resolve_hdivh      r#   c                 4    t          j        t                    S rC   )r   r   	Cuda_hnegr	  s     r!   resolve_hnegzCudaFp16Template.resolve_hnegk  rP  r#   c                 4    t          j        t                    S rC   )r   r   	Cuda_habsr	  s     r!   resolve_habszCudaFp16Template.resolve_habsn  rP  r#   c                 4    t          j        t                    S rC   )r   r   r~   r	  s     r!   resolve_hfmazCudaFp16Template.resolve_hfmaq  rP  r#   c                     t           S rC   )hsin_devicer	  s     r!   resolve_hsinzCudaFp16Template.resolve_hsint  rZ  r#   c                     t           S rC   )hcos_devicer	  s     r!   resolve_hcoszCudaFp16Template.resolve_hcosw  rZ  r#   c                     t           S rC   )hlog_devicer	  s     r!   resolve_hlogzCudaFp16Template.resolve_hlogz  rZ  r#   c                     t           S rC   )hlog10_devicer	  s     r!   resolve_hlog10zCudaFp16Template.resolve_hlog10}      r#   c                     t           S rC   )hlog2_devicer	  s     r!   resolve_hlog2zCudaFp16Template.resolve_hlog2      r#   c                     t           S rC   )hexp_devicer	  s     r!   resolve_hexpzCudaFp16Template.resolve_hexp  rZ  r#   c                     t           S rC   )hexp10_devicer	  s     r!   resolve_hexp10zCudaFp16Template.resolve_hexp10  ro  r#   c                     t           S rC   )hexp2_devicer	  s     r!   resolve_hexp2zCudaFp16Template.resolve_hexp2  rs  r#   c                     t           S rC   )hfloor_devicer	  s     r!   resolve_hfloorzCudaFp16Template.resolve_hfloor  ro  r#   c                     t           S rC   )hceil_devicer	  s     r!   resolve_hceilzCudaFp16Template.resolve_hceil  rs  r#   c                     t           S rC   )hsqrt_devicer	  s     r!   resolve_hsqrtzCudaFp16Template.resolve_hsqrt  rs  r#   c                     t           S rC   )hrsqrt_devicer	  s     r!   resolve_hrsqrtzCudaFp16Template.resolve_hrsqrt  ro  r#   c                     t           S rC   )hrcp_devicer	  s     r!   resolve_hrcpzCudaFp16Template.resolve_hrcp  rZ  r#   c                     t           S rC   )hrint_devicer	  s     r!   resolve_hrintzCudaFp16Template.resolve_hrint  rs  r#   c                     t           S rC   )htrunc_devicer	  s     r!   resolve_htrunczCudaFp16Template.resolve_htrunc  ro  r#   c                 4    t          j        t                    S rC   )r   r   Cuda_heqr	  s     r!   resolve_heqzCudaFp16Template.resolve_heq      ~h'''r#   c                 4    t          j        t                    S rC   )r   r   Cuda_hner	  s     r!   resolve_hnezCudaFp16Template.resolve_hne  r  r#   c                 4    t          j        t                    S rC   )r   r   Cuda_hger	  s     r!   resolve_hgezCudaFp16Template.resolve_hge  r  r#   c                 4    t          j        t                    S rC   )r   r   Cuda_hgtr	  s     r!   resolve_hgtzCudaFp16Template.resolve_hgt  r  r#   c                 4    t          j        t                    S rC   )r   r   Cuda_hler	  s     r!   resolve_hlezCudaFp16Template.resolve_hle  r  r#   c                 4    t          j        t                    S rC   )r   r   Cuda_hltr	  s     r!   resolve_hltzCudaFp16Template.resolve_hlt  r  r#   c                 4    t          j        t                    S rC   )r   r   	Cuda_hmaxr	  s     r!   resolve_hmaxzCudaFp16Template.resolve_hmax  rP  r#   c                 4    t          j        t                    S rC   )r   r   	Cuda_hminr	  s     r!   resolve_hminzCudaFp16Template.resolve_hmin  rP  r#   N)&r4   r5   r6   r   r  r   r   r;   rO  rS  rV  rY  r]  r`  rb  re  rh  rk  rn  rr  rv  ry  r|  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r   r#   r!   rL  rL  [  s       
%,ty
!
!C) ) )) ) )) ) )  ) ) )) ) )) ) )                              ( ( (( ( (( ( (( ( (( ( (( ( () ) )) ) ) ) )r#   rL  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 rC   )r   r  r   cgr	  s     r!   
resolve_cgzCudaModuleTemplate.resolve_cg  s    |DG$$$r#   c                     t           S rC   r   r	  s     r!   resolve_threadIdxz$CudaModuleTemplate.resolve_threadIdx      r#   c                     t           S rC   r   r	  s     r!   resolve_blockIdxz#CudaModuleTemplate.resolve_blockIdx  r  r#   c                     t           S rC   r   r	  s     r!   resolve_blockDimz#CudaModuleTemplate.resolve_blockDim  r  r#   c                     t           S rC   r   r	  s     r!   resolve_gridDimz"CudaModuleTemplate.resolve_gridDim  r  r#   c                     t           j        S rC   r  r	  s     r!   resolve_laneidz!CudaModuleTemplate.resolve_laneid  r  r#   c                 >    t          j        t          j                  S rC   )r   r  r   r9   r	  s     r!   resolve_sharedz!CudaModuleTemplate.resolve_shared      |DK(((r#   c                 4    t          j        t                    S rC   )r   r   rn   r	  s     r!   resolve_popczCudaModuleTemplate.resolve_popc  rP  r#   c                 4    t          j        t                    S rC   )r   r   r   r	  s     r!   resolve_brevzCudaModuleTemplate.resolve_brev  rP  r#   c                 4    t          j        t                    S rC   )r   r   r   r	  s     r!   resolve_clzzCudaModuleTemplate.resolve_clz  r  r#   c                 4    t          j        t                    S rC   )r   r   r   r	  s     r!   resolve_ffszCudaModuleTemplate.resolve_ffs  r  r#   c                 4    t          j        t                    S rC   )r   r   ry   r	  s     r!   resolve_fmazCudaModuleTemplate.resolve_fma  r  r#   c                 4    t          j        t                    S rC   )r   r   r   r	  s     r!   resolve_cbrtzCudaModuleTemplate.resolve_cbrt  rP  r#   c                 4    t          j        t                    S rC   )r   r   rH   r	  s     r!   resolve_threadfencez&CudaModuleTemplate.resolve_threadfence      ~5666r#   c                 4    t          j        t                    S rC   )r   r   rM   r	  s     r!   resolve_threadfence_blockz,CudaModuleTemplate.resolve_threadfence_block  s    ~4555r#   c                 4    t          j        t                    S rC   )r   r   rP   r	  s     r!   resolve_threadfence_systemz-CudaModuleTemplate.resolve_threadfence_system  r  r#   c                 4    t          j        t                    S rC   )r   r   rS   r	  s     r!   resolve_syncwarpz#CudaModuleTemplate.resolve_syncwarp  s    ~m,,,r#   c                 4    t          j        t                    S rC   )r   r   rW   r	  s     r!   resolve_shfl_sync_intrinsicz.CudaModuleTemplate.resolve_shfl_sync_intrinsic      ~6777r#   c                 4    t          j        t                    S rC   )r   r   r^   r	  s     r!   resolve_vote_sync_intrinsicz.CudaModuleTemplate.resolve_vote_sync_intrinsic  r  r#   c                 4    t          j        t                    S rC   )r   r   ra   r	  s     r!   resolve_match_any_syncz)CudaModuleTemplate.resolve_match_any_sync      ~1222r#   c                 4    t          j        t                    S rC   )r   r   rd   r	  s     r!   resolve_match_all_syncz)CudaModuleTemplate.resolve_match_all_sync  r  r#   c                 4    t          j        t                    S rC   )r   r   rg   r	  s     r!   resolve_activemaskz%CudaModuleTemplate.resolve_activemask  r#  r#   c                 4    t          j        t                    S rC   )r   r   rk   r	  s     r!   resolve_lanemask_ltz&CudaModuleTemplate.resolve_lanemask_lt   r  r#   c                 4    t          j        t                    S rC   )r   r   r   r	  s     r!   resolve_selpzCudaModuleTemplate.resolve_selp  rP  r#   c                 4    t          j        t                    S rC   )r   r   r  r	  s     r!   resolve_nanosleepz$CudaModuleTemplate.resolve_nanosleep  r-  r#   c                 >    t          j        t          j                  S rC   )r   r  r   r   r	  s     r!   resolve_atomicz!CudaModuleTemplate.resolve_atomic	  r  r#   c                 >    t          j        t          j                  S rC   )r   r  r   r   r	  s     r!   resolve_fp16zCudaModuleTemplate.resolve_fp16  s    |DI&&&r#   c                 >    t          j        t          j                  S rC   )r   r  r   rE   r	  s     r!   resolve_constz CudaModuleTemplate.resolve_const      |DJ'''r#   c                 >    t          j        t          j                  S rC   )r   r  r   r>   r	  s     r!   resolve_localz CudaModuleTemplate.resolve_local  r  r#   N)$r4   r5   r6   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  )loglog2log10)operator
numba.corer   numba.core.typing.npydeclr   r   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   r8   r=   r@   rH   rM   rP   rS   rW   r^   ra   rd   rg   rk   rn   ry   r~   r   r   r   r   r   r   r   r   floatr   r   r   r   r   r   haddrN  addCuda_addiadd	Cuda_iaddhsubrR  subCuda_subisub	Cuda_isubhmulrU  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   rd  rg  rj  rm  rq  ru  rx  r{  r  r  r~  r  r  r  r  rX  r   r|   r{   rs   ri   rt   rw   all_numba_typesr   unsigned_int_numba_typesr   r!  r%  maxr;  minr>  nanmaxrE  nanminrA  and_r(  or_r+  xorr/  incr2  decr5  exchr8  r   r   r  r  r  r  r  r  rL  r  r  funcr   r#   r!   <module>r.     sx         @ @ @ @ @ @ @ @ @ @ @ @ @ @ @ @ @ @ @ @> > > > > > > > > > > > > > > > " ! ! ! ! ! * * * * * *       @ @ @ @ @ @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//// 4 4D'''T?3334 4r#   