
    shZ                     <   S SK r S SKJr  S SKJrJrJrJrJrJ	r	J
r
Jr  S SKJrJrJrJrJrJr  S SKJr  S SKJr  S SKJr  S SKJr  \" 5       r\R8                  r\R:                  r\R<                  r\" \5         " S	 S
\5      r\ " S S\5      5       r \ " S S\5      5       r!\ " S S\5      5       r"\ " S S\5      5       r#\ " S S\5      5       r$\ " S S\5      5       r%\ " S S\5      5       r&\ " S S\5      5       r'\ " S S\5      5       r(\ " S S\5      5       r)\ " S S \5      5       r*\ " S! S"\5      5       r+\ " S# S$\5      5       r,\ " S% S&\5      5       r-\ " S' S(\5      5       r.\ " S) S*\5      5       r/\ " S+ S,\5      5       r0\ " S- S.\5      5       r1\ " S/ S0\5      5       r2\ " S1 S2\5      5       r3\ " S3 S4\5      5       r4S5 r5S6 r6S7 r7\" \85       " S8 S9\5      5       r9S: r:S; r;S< r<S= r=\7" \R|                  R~                  5      r@\=" \ R                  5      rB\=" \ R                  5      rD\7" \R|                  R                  5      rF\=" \ R                  5      rH\=" \ R                  5      rJ\7" \R|                  R                  5      rL\=" \ R                  5      rN\=" \ R                  5      rP\7" \R|                  R                  5      rR\7" \R|                  R                  5      rT\5" \R|                  R                  5      rV\6" \ R                  5      rX\5" \R|                  R                  5      rZ\6" \[5      r\\:" \R|                  R                  5      r^\<" \ R                  5        \:" \R|                  R                  5      ra\<" \ R                  5        \:" \R|                  R                  5      rd\<" \ R                  5        \:" \R|                  R                  5      rg\<" \ R                  5        \:" \R|                  R                  5      rj\<" \ R                  5        \:" \R|                  R                  5      rm\<" \ R                  5        \=" \ R                  5        \=" \ R                  5        S> rqS? rr\q" S@5      rs\q" SA5      rt\q" SB5      ru\q" SC5      rv\q" SD5      rw\q" SE5      rx\q" SF5      ry\q" SG5      rz\q" SH5      r{\q" SI5      r|\q" SJ5      r}\q" SK5      r~\q" SL5      r\q" SM5      r\q" SN5      r\r" SO5      rSP r\GR                  \GR
                  \GR                  \GR                  \GR                  \GR                  4r\GR                  \GR                  \GR                  \GR                  4r\GR                  \GR                  4r\" \GR                  R                  \5      r\" \GR                  R                  \5      r\" \GR                  GR                   \5      r\" \GR                  GR$                  \5      r\" \GR                  GR(                  \5      r\" \GR                  GR,                  \5      r\" \GR                  GR0                  \5      r\" \GR                  GR4                  \5      r\" \GR                  GR8                  \5      r\" \GR                  GR<                  \5      r\" \GR                  GR@                  \5      r\" \GR                  GRD                  \5      r\ " SQ SR\5      5       r\ " SS ST\5      5       r\ " SU SV\5      5       r\ " SW SX\5      5       r\ " SY SZ\5      5       r\ " S[ S\\5      5       r\ " S] S^\5      5       r\ " S_ S`\5      5       r\ " Sa Sb\5      5       r\ " Sc Sd\5      5       r\" \\GR\                  " \5      5        \ H  r\" \\5        M     \	 H  r\" \\5        M     \ H  r\" \\5        M     \
 H  r\Se;   d  M  \" \\5        M     g)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                       \ rS rSrS rSrg)Cuda_array_decl   c                     S nU$ )Nc           	         [        U [        R                  5      (       a!  [        U [        R                  5      (       d  g Op[        U [        R                  [        R
                  45      (       a@  [        U  Vs/ s H"  n[        U[        R                  5      (       + PM$     sn5      (       a  g Og [        U 5      n[        U5      nUb  Ub  [        R                  " XCSS9$ g g s  snf )NC)dtypendimlayout)

isinstancer   IntegerIntegerLiteralTupleUniTupleanyr   r   Array)shaper   sr   nb_dtypes        g/Users/tiagomarins/Projetos/claudeai/copy_bank/venv/lib/python3.13/site-packages/numba/cuda/cudadecl.pytyper&Cuda_array_decl.generic.<locals>.typer   s     %//!%)=)=>> ?EEKK#@AA!&(!&A 'q%*>*>??!&( ) )) u%D"5)H#(8{{SII )9#(s   9)C$ selfr+   s     r*   genericCuda_array_decl.generic   s    	J&     r-   N__name__
__module____qualname____firstlineno__r0   __static_attributes__r-   r2   r*   r   r      s    r2   r   c                   @    \ rS rSr\R
                  R                  rSrg)Cuda_shared_array2   r-   N)	r4   r5   r6   r7   r   sharedarraykeyr8   r-   r2   r*   r:   r:   2   s    
++

Cr2   r:   c                   @    \ rS rSr\R
                  R                  rSrg)Cuda_local_array7   r-   N)	r4   r5   r6   r7   r   localr=   r>   r8   r-   r2   r*   r@   r@   7   s    
**

Cr2   r@   c                   F    \ rS rSr\R
                  R                  rS rSr	g)Cuda_const_array_like<   c                     S nU$ )Nc                     U $ Nr-   )ndarrays    r*   r+   ,Cuda_const_array_like.generic.<locals>.typerA   s    Nr2   r-   r.   s     r*   r0   Cuda_const_array_like.generic@   s    	r2   r-   N)
r4   r5   r6   r7   r   const
array_liker>   r0   r8   r-   r2   r*   rD   rD   <   s    
**

Cr2   rD   c                   R    \ rS rSr\R
                  r\" \R                  5      /r
Srg)Cuda_threadfence_deviceF   r-   N)r4   r5   r6   r7   r   threadfencer>   r   r   nonecasesr8   r-   r2   r*   rO   rO   F   s    


Cuzz"#Er2   rO   c                   R    \ rS rSr\R
                  r\" \R                  5      /r
Srg)Cuda_threadfence_blockL   r-   N)r4   r5   r6   r7   r   threadfence_blockr>   r   r   rR   rS   r8   r-   r2   r*   rU   rU   L   s    

 
 Cuzz"#Er2   rU   c                   R    \ rS rSr\R
                  r\" \R                  5      /r
Srg)Cuda_threadfence_systemR   r-   N)r4   r5   r6   r7   r   threadfence_systemr>   r   r   rR   rS   r8   r-   r2   r*   rY   rY   R   s    

!
!Cuzz"#Er2   rY   c                       \ rS rSr\R
                  r\" \R                  5      \" \R                  \R                  5      /rSrg)Cuda_syncwarpX   r-   N)r4   r5   r6   r7   r   syncwarpr>   r   r   rR   i4rS   r8   r-   r2   r*   r]   r]   X   s-    
--Cuzz"Iejj%(($CDEr2   r]   c                   P   \ rS rSr\R
                  r\" \R                  " \R                  \R                  45      \R                  \R                  \R                  \R                  \R                  5      \" \R                  " \R                  \R                  45      \R                  \R                  \R                  \R                  \R                  5      \" \R                  " \R                  \R                  45      \R                  \R                  \R                  \R                  \R                  5      \" \R                  " \R                  \R                  45      \R                  \R                  \R                  \R                  \R                  5      /rSrg)Cuda_shfl_sync_intrinsic^   r-   N)r4   r5   r6   r7   r   shfl_sync_intrinsicr>   r   r   r#   r`   b1i8f4f8rS   r8   r-   r2   r*   rb   rb   ^   s   

"
"C%++uxx23((EHHehh%((	D%++uxx23((EHHehh%((	D%++uxx23((EHHehh%((	D%++uxx23((EHHehh%((	D	Er2   rb   c                       \ rS rSr\R
                  r\" \R                  " \R                  \R                  45      \R                  \R                  \R                  5      /rSrg)Cuda_vote_sync_intrinsicm   r-   N)r4   r5   r6   r7   r   vote_sync_intrinsicr>   r   r   r#   r`   re   rS   r8   r-   r2   r*   rj   rj   m   sG    

"
"Cu{{EHHehh#78xx5885 6Er2   rj   c                   h   \ rS rSr\R
                  r\" \R                  \R                  \R                  5      \" \R                  \R                  \R                  5      \" \R                  \R                  \R                  5      \" \R                  \R                  \R                  5      /rSrg)Cuda_match_any_synct   r-   N)r4   r5   r6   r7   r   match_any_syncr>   r   r   r`   rf   rg   rh   rS   r8   r-   r2   r*   rn   rn   t   st    


C%((EHHehh/%((EHHehh/%((EHHehh/%((EHHehh/	Er2   rn   c            	       H   \ rS rSr\R
                  r\" \R                  " \R                  \R                  45      \R                  \R                  5      \" \R                  " \R                  \R                  45      \R                  \R                  5      \" \R                  " \R                  \R                  45      \R                  \R                  5      \" \R                  " \R                  \R                  45      \R                  \R                  5      /rSrg)Cuda_match_all_sync   r-   N)r4   r5   r6   r7   r   match_all_syncr>   r   r   r#   r`   re   rf   rg   rh   rS   r8   r-   r2   r*   rr   rr      s    


C%++uxx23UXXuxxH%++uxx23UXXuxxH%++uxx23UXXuxxH%++uxx23UXXuxxH	Er2   rr   c                   R    \ rS rSr\R
                  r\" \R                  5      /r
Srg)Cuda_activemask   r-   N)r4   r5   r6   r7   r   
activemaskr>   r   r   uint32rS   r8   r-   r2   r*   rv   rv      s    
//Cu||$%Er2   rv   c                   R    \ rS rSr\R
                  r\" \R                  5      /r
Srg)Cuda_lanemask_lt   r-   N)r4   r5   r6   r7   r   lanemask_ltr>   r   r   ry   rS   r8   r-   r2   r*   r{   r{      s    


Cu||$%Er2   r{   c                      \ rS rSrSr\R                  r\" \	R                  \	R                  5      \" \	R                  \	R                  5      \" \	R                  \	R                  5      \" \	R                  \	R                  5      \" \	R                  \	R                  5      \" \	R                  \	R                  5      \" \	R                   \	R                   5      \" \	R"                  \	R"                  5      /rSrg)	Cuda_popc   zz
Supported types from `llvm.popc`
[here](http://docs.nvidia.com/cuda/nvvm-ir-spec/index.html#bit-manipulations-intrinics)
r-   N)r4   r5   r6   r7   __doc__r   popcr>   r   r   int8int16int32int64uint8uint16ry   uint64rS   r8   r-   r2   r*   r   r      s     ))C%**ejj)%++u{{+%++u{{+%++u{{+%++u{{+%,,-%,,-%,,-	Er2   r   c                       \ rS rSrSr\R                  r\" \	R                  \	R                  \	R                  \	R                  5      \" \	R                  \	R                  \	R                  \	R                  5      /rSrg)Cuda_fma   z{
Supported types from `llvm.fma`
[here](https://docs.nvidia.com/cuda/nvvm-ir-spec/index.html#standard-c-library-intrinics)
r-   N)r4   r5   r6   r7   r   r   fmar>   r   r   float32float64rS   r8   r-   r2   r*   r   r      sU     ((C%--u}}M%--u}}MEr2   r   c                       \ rS rSr\R
                  R                  r\" \	R                  \	R                  \	R                  \	R                  5      /rSrg)	Cuda_hfma   r-   N)r4   r5   r6   r7   r   fp16hfmar>   r   r   float16rS   r8   r-   r2   r*   r   r      s4    
))..C%--u}}MEr2   r   c                       \ rS rSr\R
                  r\" \R                  \R                  5      \" \R                  \R                  5      /rSrg)	Cuda_cbrt   r-   N)r4   r5   r6   r7   r   cbrtr>   r   r   r   r   rS   r8   r-   r2   r*   r   r      s8     ))C%--/%--/Er2   r   c                       \ rS rSr\R
                  r\" \R                  \R                  5      \" \R                  \R                  5      /rSrg)	Cuda_brev   r-   N)r4   r5   r6   r7   r   brevr>   r   r   ry   r   rS   r8   r-   r2   r*   r   r      s6    
))C%,,-%,,-Er2   r   c                      \ rS rSrSr\R                  r\" \	R                  \	R                  5      \" \	R                  \	R                  5      \" \	R                  \	R                  5      \" \	R                  \	R                  5      \" \	R                  \	R                  5      \" \	R                  \	R                  5      \" \	R                   \	R                   5      \" \	R"                  \	R"                  5      /rSrg)Cuda_clz   zz
Supported types from `llvm.ctlz`
[here](http://docs.nvidia.com/cuda/nvvm-ir-spec/index.html#bit-manipulations-intrinics)
r-   N)r4   r5   r6   r7   r   r   clzr>   r   r   r   r   r   r   r   r   ry   r   rS   r8   r-   r2   r*   r   r      s     ((C%**ejj)%++u{{+%++u{{+%++u{{+%++u{{+%,,-%,,-%,,-	Er2   r   c                      \ rS rSrSr\R                  r\" \	R                  \	R                  5      \" \	R                  \	R                  5      \" \	R                  \	R                  5      \" \	R                  \	R                  5      \" \	R                  \	R                  5      \" \	R                  \	R                   5      \" \	R                  \	R                  5      \" \	R                  \	R"                  5      /rSrg)Cuda_ffs   zz
Supported types from `llvm.cttz`
[here](http://docs.nvidia.com/cuda/nvvm-ir-spec/index.html#bit-manipulations-intrinics)
r-   N)r4   r5   r6   r7   r   r   ffsr>   r   r   ry   r   r   r   r   r   r   r   rS   r8   r-   r2   r*   r   r      s     ((C%,,

+%,,,%,,,%,,,%,,,%,,-%,,-%,,-	Er2   r   c                   2    \ rS rSr\R
                  rS rSrg)	Cuda_selp   c                 @   U(       a   eUu  p4n[         R                  [         R                  [         R                  [         R                  [         R
                  [         R                  [         R                  [         R                  4nXE:w  d  XF;  a  g [        XCXD5      $ rH   )
r   r   r   r   r   r   ry   r   r   r   )r/   argskwstestabsupported_typess          r*   r0   Cuda_selp.generic   sk    w
 !==%-- ;; ;; ;;6
 6Q-!''r2   r-   N)	r4   r5   r6   r7   r   selpr>   r0   r8   r-   r2   r*   r   r      s    
))C(r2   r   c                 >   ^  [          " U 4S jS[        5      5       nU$ )Nc                   X   > \ rS rSr Y r\" \R                  \R                  5      /rSr	g)'_genfp16_unary.<locals>.Cuda_fp16_unaryi  r-   N
r4   r5   r6   r7   r>   r   r   r   rS   r8   l_keys   r*   Cuda_fp16_unaryr     s    5==%--89r2   r   registerr   r   r   s   ` r*   _genfp16_unaryr     s%    :* : : r2   c                 H   ^  [        T 5       " U 4S jS[        5      5       nU$ )Nc                   "   > \ rS rSr Y rS rSrg)0_genfp16_unary_operator.<locals>.Cuda_fp16_unaryi  c                     U(       a   e[        U5      S:X  a@  US   [        R                  :X  a(  [        [        R                  [        R                  5      $ g g )N   r   )lenr   r   r   )r/   r   r   s      r*   r0   8_genfp16_unary_operator.<locals>.Cuda_fp16_unary.generic  s?    N74yA~$q'U]]": >> #;~r2   r-   Nr4   r5   r6   r7   r>   r0   r8   r   s   r*   r   r     s    	?r2   r   register_globalr   r   s   ` r*   _genfp16_unary_operatorr     s*    U?* ? ? r2   c                 >   ^  [          " U 4S jS[        5      5       nU$ )Nc                   n   > \ rS rSr Y r\" \R                  \R                  \R                  5      /rSr	g))_genfp16_binary.<locals>.Cuda_fp16_binaryi#  r-   Nr   r   s   r*   Cuda_fp16_binaryr   #  s%    5==%--GHr2   r   r   )r   r   s   ` r*   _genfp16_binaryr   "  s(    I+ I I r2   c                       \ rS rSrS rSrg)Floati+  c                 \    U(       a   eUu  nU[         R                  :X  a  [        X35      $ g rH   )r   r   r   )r/   r   r   args       r*   r0   Float.generic.  s,    w%--S&&  r2   r-   Nr3   r-   r2   r*   r   r   +  s    'r2   r   c                 >   ^  [          " U 4S jS[        5      5       nU$ )Nc                   n   > \ rS rSr Y r\" \R                  \R                  \R                  5      /r	Sr
g)1_genfp16_binary_comparison.<locals>.Cuda_fp16_cmpi8  r-   N)r4   r5   r6   r7   r>   r   r   re   r   rS   r8   r   s   r*   Cuda_fp16_cmpr   8  s)     ehhu}}=
r2   r   r   )r   r   s   ` r*   _genfp16_binary_comparisonr   7  s%    
( 
 
 r2   c                 L   ^ ^ [        T 5       " U U4S jS[        5      5       nU$ )Nc                   (   > \ rS rSr Y rU4S jrSrg)1_fp16_binary_operator.<locals>.Cuda_fp16_operatoriP  c                   > U(       a   e[        U5      S:X  a  US   [        R                  :X  d  US   [        R                  :X  a  US   [        R                  :X  a#  U R                  R	                  US   US   5      nO"U R                  R	                  US   US   5      nU[
        R                  :X  d(  U[
        R                  :X  d  U[
        R                  :X  a)  [        T[        R                  [        R                  5      $ g g g )N   r   r   )
r   r   r   contextcan_convertr   exactpromotesafer   )r/   r   r   convertiblerettys       r*   r0   9_fp16_binary_operator.<locals>.Cuda_fp16_operator.genericT  s    N74yA~!W-aEMM1IGu}},"&,,":":47DG"LK"&,,":":47DG"LK  :#3#33:#5#55:??2$UEMM5==II 3! 2J r2   r-   Nr   )r   r   s   r*   Cuda_fp16_operatorr   P  s    	J 	Jr2   r   r   )r   r   r   s   `` r*   _fp16_binary_operatorr   O  s3    UJ J- J J4 r2   c                 6    [        U [        R                  5      $ rH   )r   r   re   ops    r*   _genfp16_comparison_operatorr   n  s     UXX..r2   c                 6    [        U [        R                  5      $ rH   )r   r   r   r   s    r*   _genfp16_binary_operatorr   r  s     U]]33r2   c                     [        SU  3[        R                  [        R                  45      n[        R                  " U5      $ N__numba_wrapper_r   r   r   Functionfnamedecls     r*   _resolve_wrapped_unaryr     s8    +.>ug,F,1MM-2]],<>D >>$r2   c                     [        SU  3[        R                  [        R                  [        R                  45      n[        R                  " U5      $ r   r   r   s     r*   _resolve_wrapped_binaryr     s?    +.>ug,F,1MM-2]]EMM,KMD >>$r2   hsinhcoshloghlog10hlog2hexphexp10hexp2hsqrthrsqrthfloorhceilhrcphrinthtrunchdivc                 B   ^ ^ [          " U U4S jS[        5      5       nU$ )Nc                   (   > \ rS rSr Y rU4S jrSrg)_gen.<locals>.Cuda_atomici  c                 &  > U(       a   eUu  p4nUR                   T;  a  g UR                  S:X  a0  [        UR                   U[        R                  UR                   5      $ UR                  S:  a!  [        UR                   X4UR                   5      $ g Nr   )r   r   r   r   intp)r/   r   r   aryidxvalr   s         r*   r0   !_gen.<locals>.Cuda_atomic.generic  sq    N7 MCcyy/xx1} CSYYGGA Ccii@@ r2   r-   Nr   )r   r   s   r*   Cuda_atomicr	    s    
	A 
	Ar2   r  )r   r   )r   r   r  s   `` r*   _genr    s.    A A& A A r2   c                   F    \ rS rSr\R
                  R                  rS rSr	g)Cuda_atomic_compare_and_swapi  c                     U(       a   eUu  p4nUR                   nU[        ;   a  UR                  S:X  a  [        XcXf5      $ g g r  )r   integer_numba_typesr   r   )r/   r   r   r  oldr  dtys          r*   r0   $Cuda_atomic_compare_and_swap.generic  sC    w#ii%%#((a-Ss00 +8%r2   r-   N)
r4   r5   r6   r7   r   atomiccompare_and_swapr>   r0   r8   r-   r2   r*   r  r    s    
++
&
&C1r2   r  c                   F    \ rS rSr\R
                  R                  rS rSr	g)Cuda_atomic_casi  c                     U(       a   eUu  p4pVUR                   nU[        ;  a  g UR                  S:X  a  [        Xs[        R
                  Xw5      $ UR                  S:  a  [        XsXGU5      $ g r  )r   r  r   r   r   r  )r/   r   r   r  r  r  r  r  s           r*   r0   Cuda_atomic_cas.generic  se    w!#ii))88q=Suzz3<<XX\Ss55 r2   r-   N)
r4   r5   r6   r7   r   r  casr>   r0   r8   r-   r2   r*   r  r    s    
++//C6r2   r  c                   h    \ rS rSr\R
                  r\" \R                  \R                  5      /rSrg)Cuda_nanosleepi  r-   N)r4   r5   r6   r7   r   	nanosleepr>   r   r   voidry   rS   r8   r-   r2   r*   r"  r"    s"    
..Cuzz5<<01Er2   r"  c                   *    \ rS rSr\rS rS rS rSr	g)
Dim3_attrsi  c                 "    [         R                  $ rH   r   r   r/   mods     r*   	resolve_xDim3_attrs.resolve_x
      {{r2   c                 "    [         R                  $ rH   r(  r)  s     r*   	resolve_yDim3_attrs.resolve_y  r-  r2   c                 "    [         R                  $ rH   r(  r)  s     r*   	resolve_zDim3_attrs.resolve_z  r-  r2   r-   N)
r4   r5   r6   r7   r   r>   r+  r/  r2  r8   r-   r2   r*   r&  r&    s    
Cr2   r&  c                   R    \ rS rSr\R
                  " \R                  5      rS r	Sr
g)CudaSharedModuleTemplatei  c                 6    [         R                  " [        5      $ rH   )r   r   r:   r)  s     r*   resolve_array&CudaSharedModuleTemplate.resolve_array  s    ~~/00r2   r-   N)r4   r5   r6   r7   r   Moduler   r<   r>   r7  r8   r-   r2   r*   r5  r5    s    
,,t{{
#C1r2   r5  c                   R    \ rS rSr\R
                  " \R                  5      rS r	Sr
g)CudaConstModuleTemplatei  c                 6    [         R                  " [        5      $ rH   )r   r   rD   r)  s     r*   resolve_array_like*CudaConstModuleTemplate.resolve_array_like   s    ~~344r2   r-   N)r4   r5   r6   r7   r   r9  r   rL   r>   r=  r8   r-   r2   r*   r;  r;    s    
,,tzz
"C5r2   r;  c                   R    \ rS rSr\R
                  " \R                  5      rS r	Sr
g)CudaLocalModuleTemplatei$  c                 6    [         R                  " [        5      $ rH   )r   r   r@   r)  s     r*   r7  %CudaLocalModuleTemplate.resolve_array(      ~~.//r2   r-   N)r4   r5   r6   r7   r   r9  r   rB   r>   r7  r8   r-   r2   r*   r@  r@  $  s    
,,tzz
"C0r2   r@  c                       \ rS rSr\R
                  " \R                  5      rS r	S r
S rS rS rS rS rS	 rS
 rS rS rS rS rS rSrg)CudaAtomicTemplatei,  c                 6    [         R                  " [        5      $ rH   )r   r   Cuda_atomic_addr)  s     r*   resolve_addCudaAtomicTemplate.resolve_add0      ~~o..r2   c                 6    [         R                  " [        5      $ rH   )r   r   Cuda_atomic_subr)  s     r*   resolve_subCudaAtomicTemplate.resolve_sub3  rJ  r2   c                 6    [         R                  " [        5      $ rH   )r   r   Cuda_atomic_andr)  s     r*   resolve_and_CudaAtomicTemplate.resolve_and_6  rJ  r2   c                 6    [         R                  " [        5      $ rH   )r   r   Cuda_atomic_orr)  s     r*   resolve_or_CudaAtomicTemplate.resolve_or_9      ~~n--r2   c                 6    [         R                  " [        5      $ rH   )r   r   Cuda_atomic_xorr)  s     r*   resolve_xorCudaAtomicTemplate.resolve_xor<  rJ  r2   c                 6    [         R                  " [        5      $ rH   )r   r   Cuda_atomic_incr)  s     r*   resolve_incCudaAtomicTemplate.resolve_inc?  rJ  r2   c                 6    [         R                  " [        5      $ rH   )r   r   Cuda_atomic_decr)  s     r*   resolve_decCudaAtomicTemplate.resolve_decB  rJ  r2   c                 6    [         R                  " [        5      $ rH   )r   r   Cuda_atomic_exchr)  s     r*   resolve_exchCudaAtomicTemplate.resolve_exchE  rC  r2   c                 6    [         R                  " [        5      $ rH   )r   r   Cuda_atomic_maxr)  s     r*   resolve_maxCudaAtomicTemplate.resolve_maxH  rJ  r2   c                 6    [         R                  " [        5      $ rH   )r   r   Cuda_atomic_minr)  s     r*   resolve_minCudaAtomicTemplate.resolve_minK  rJ  r2   c                 6    [         R                  " [        5      $ rH   )r   r   Cuda_atomic_nanminr)  s     r*   resolve_nanmin!CudaAtomicTemplate.resolve_nanminN      ~~011r2   c                 6    [         R                  " [        5      $ rH   )r   r   Cuda_atomic_nanmaxr)  s     r*   resolve_nanmax!CudaAtomicTemplate.resolve_nanmaxQ  rt  r2   c                 6    [         R                  " [        5      $ rH   )r   r   r  r)  s     r*   resolve_compare_and_swap+CudaAtomicTemplate.resolve_compare_and_swapT  s    ~~:;;r2   c                 6    [         R                  " [        5      $ rH   )r   r   r  r)  s     r*   resolve_casCudaAtomicTemplate.resolve_casW  rJ  r2   r-   N)r4   r5   r6   r7   r   r9  r   r  r>   rH  rM  rQ  rU  rZ  r^  rb  rf  rj  rn  rr  rw  rz  r}  r8   r-   r2   r*   rE  rE  ,  sZ    
,,t{{
#C///.///0//22</r2   rE  c                       \ rS rSr\R
                  " \R                  5      rS r	S r
S rS rS rS rS rS	 rS
 rS rS rS rS rS rS rS rS rS rS rS rS rS rS rS r S r!S r"S r#S r$S r%S r&S r'g!)"CudaFp16Templatei[  c                 6    [         R                  " [        5      $ rH   )r   r   	Cuda_haddr)  s     r*   resolve_haddCudaFp16Template.resolve_hadd_      ~~i((r2   c                 6    [         R                  " [        5      $ rH   )r   r   	Cuda_hsubr)  s     r*   resolve_hsubCudaFp16Template.resolve_hsubb  r  r2   c                 6    [         R                  " [        5      $ rH   )r   r   	Cuda_hmulr)  s     r*   resolve_hmulCudaFp16Template.resolve_hmule  r  r2   c                     [         $ rH   )hdiv_devicer)  s     r*   resolve_hdivCudaFp16Template.resolve_hdivh      r2   c                 6    [         R                  " [        5      $ rH   )r   r   	Cuda_hnegr)  s     r*   resolve_hnegCudaFp16Template.resolve_hnegk  r  r2   c                 6    [         R                  " [        5      $ rH   )r   r   	Cuda_habsr)  s     r*   resolve_habsCudaFp16Template.resolve_habsn  r  r2   c                 6    [         R                  " [        5      $ rH   )r   r   r   r)  s     r*   resolve_hfmaCudaFp16Template.resolve_hfmaq  r  r2   c                     [         $ rH   )hsin_devicer)  s     r*   resolve_hsinCudaFp16Template.resolve_hsint  r  r2   c                     [         $ rH   )hcos_devicer)  s     r*   resolve_hcosCudaFp16Template.resolve_hcosw  r  r2   c                     [         $ rH   )hlog_devicer)  s     r*   resolve_hlogCudaFp16Template.resolve_hlogz  r  r2   c                     [         $ rH   )hlog10_devicer)  s     r*   resolve_hlog10CudaFp16Template.resolve_hlog10}      r2   c                     [         $ rH   )hlog2_devicer)  s     r*   resolve_hlog2CudaFp16Template.resolve_hlog2      r2   c                     [         $ rH   )hexp_devicer)  s     r*   resolve_hexpCudaFp16Template.resolve_hexp  r  r2   c                     [         $ rH   )hexp10_devicer)  s     r*   resolve_hexp10CudaFp16Template.resolve_hexp10  r  r2   c                     [         $ rH   )hexp2_devicer)  s     r*   resolve_hexp2CudaFp16Template.resolve_hexp2  r  r2   c                     [         $ rH   )hfloor_devicer)  s     r*   resolve_hfloorCudaFp16Template.resolve_hfloor  r  r2   c                     [         $ rH   )hceil_devicer)  s     r*   resolve_hceilCudaFp16Template.resolve_hceil  r  r2   c                     [         $ rH   )hsqrt_devicer)  s     r*   resolve_hsqrtCudaFp16Template.resolve_hsqrt  r  r2   c                     [         $ rH   )hrsqrt_devicer)  s     r*   resolve_hrsqrtCudaFp16Template.resolve_hrsqrt  r  r2   c                     [         $ rH   )hrcp_devicer)  s     r*   resolve_hrcpCudaFp16Template.resolve_hrcp  r  r2   c                     [         $ rH   )hrint_devicer)  s     r*   resolve_hrintCudaFp16Template.resolve_hrint  r  r2   c                     [         $ rH   )htrunc_devicer)  s     r*   resolve_htruncCudaFp16Template.resolve_htrunc  r  r2   c                 6    [         R                  " [        5      $ rH   )r   r   Cuda_heqr)  s     r*   resolve_heqCudaFp16Template.resolve_heq      ~~h''r2   c                 6    [         R                  " [        5      $ rH   )r   r   Cuda_hner)  s     r*   resolve_hneCudaFp16Template.resolve_hne  r  r2   c                 6    [         R                  " [        5      $ rH   )r   r   Cuda_hger)  s     r*   resolve_hgeCudaFp16Template.resolve_hge  r  r2   c                 6    [         R                  " [        5      $ rH   )r   r   Cuda_hgtr)  s     r*   resolve_hgtCudaFp16Template.resolve_hgt  r  r2   c                 6    [         R                  " [        5      $ rH   )r   r   Cuda_hler)  s     r*   resolve_hleCudaFp16Template.resolve_hle  r  r2   c                 6    [         R                  " [        5      $ rH   )r   r   Cuda_hltr)  s     r*   resolve_hltCudaFp16Template.resolve_hlt  r  r2   c                 6    [         R                  " [        5      $ rH   )r   r   	Cuda_hmaxr)  s     r*   resolve_hmaxCudaFp16Template.resolve_hmax  r  r2   c                 6    [         R                  " [        5      $ rH   )r   r   	Cuda_hminr)  s     r*   resolve_hminCudaFp16Template.resolve_hmin  r  r2   r-   N)(r4   r5   r6   r7   r   r9  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  r8   r-   r2   r*   r  r  [  s    
,,tyy
!C))))))(((((())r2   r  c                       \ rS rSr\R
                  " \5      rS rS r	S r
S rS rS rS rS	 rS
 rS rS rS rS rS rS rS rS rS rS rS rS rS rS rS rS r S r!S r"S r#S r$Sr%g )!CudaModuleTemplatei  c                 J    [         R                  " [        R                  5      $ rH   )r   r9  r   cgr)  s     r*   
resolve_cgCudaModuleTemplate.resolve_cg  s    ||DGG$$r2   c                     [         $ rH   r   r)  s     r*   resolve_threadIdx$CudaModuleTemplate.resolve_threadIdx      r2   c                     [         $ rH   r   r)  s     r*   resolve_blockIdx#CudaModuleTemplate.resolve_blockIdx  r  r2   c                     [         $ rH   r   r)  s     r*   resolve_blockDim#CudaModuleTemplate.resolve_blockDim  r  r2   c                     [         $ rH   r   r)  s     r*   resolve_gridDim"CudaModuleTemplate.resolve_gridDim  r  r2   c                 "    [         R                  $ rH   r(  r)  s     r*   resolve_laneid!CudaModuleTemplate.resolve_laneid  r-  r2   c                 J    [         R                  " [        R                  5      $ rH   )r   r9  r   r<   r)  s     r*   resolve_shared!CudaModuleTemplate.resolve_shared      ||DKK((r2   c                 6    [         R                  " [        5      $ rH   )r   r   r   r)  s     r*   resolve_popcCudaModuleTemplate.resolve_popc  r  r2   c                 6    [         R                  " [        5      $ rH   )r   r   r   r)  s     r*   resolve_brevCudaModuleTemplate.resolve_brev  r  r2   c                 6    [         R                  " [        5      $ rH   )r   r   r   r)  s     r*   resolve_clzCudaModuleTemplate.resolve_clz  r  r2   c                 6    [         R                  " [        5      $ rH   )r   r   r   r)  s     r*   resolve_ffsCudaModuleTemplate.resolve_ffs  r  r2   c                 6    [         R                  " [        5      $ rH   )r   r   r   r)  s     r*   resolve_fmaCudaModuleTemplate.resolve_fma  r  r2   c                 6    [         R                  " [        5      $ rH   )r   r   r   r)  s     r*   resolve_cbrtCudaModuleTemplate.resolve_cbrt  r  r2   c                 6    [         R                  " [        5      $ rH   )r   r   rO   r)  s     r*   resolve_threadfence&CudaModuleTemplate.resolve_threadfence      ~~566r2   c                 6    [         R                  " [        5      $ rH   )r   r   rU   r)  s     r*   resolve_threadfence_block,CudaModuleTemplate.resolve_threadfence_block  s    ~~455r2   c                 6    [         R                  " [        5      $ rH   )r   r   rY   r)  s     r*   resolve_threadfence_system-CudaModuleTemplate.resolve_threadfence_system  r,  r2   c                 6    [         R                  " [        5      $ rH   )r   r   r]   r)  s     r*   resolve_syncwarp#CudaModuleTemplate.resolve_syncwarp  s    ~~m,,r2   c                 6    [         R                  " [        5      $ rH   )r   r   rb   r)  s     r*   resolve_shfl_sync_intrinsic.CudaModuleTemplate.resolve_shfl_sync_intrinsic      ~~677r2   c                 6    [         R                  " [        5      $ rH   )r   r   rj   r)  s     r*   resolve_vote_sync_intrinsic.CudaModuleTemplate.resolve_vote_sync_intrinsic  r9  r2   c                 6    [         R                  " [        5      $ rH   )r   r   rn   r)  s     r*   resolve_match_any_sync)CudaModuleTemplate.resolve_match_any_sync      ~~122r2   c                 6    [         R                  " [        5      $ rH   )r   r   rr   r)  s     r*   resolve_match_all_sync)CudaModuleTemplate.resolve_match_all_sync  r@  r2   c                 6    [         R                  " [        5      $ rH   )r   r   rv   r)  s     r*   resolve_activemask%CudaModuleTemplate.resolve_activemask  rJ  r2   c                 6    [         R                  " [        5      $ rH   )r   r   r{   r)  s     r*   resolve_lanemask_lt&CudaModuleTemplate.resolve_lanemask_lt   rC  r2   c                 6    [         R                  " [        5      $ rH   )r   r   r   r)  s     r*   resolve_selpCudaModuleTemplate.resolve_selp  r  r2   c                 6    [         R                  " [        5      $ rH   )r   r   r"  r)  s     r*   resolve_nanosleep$CudaModuleTemplate.resolve_nanosleep  rW  r2   c                 J    [         R                  " [        R                  5      $ rH   )r   r9  r   r  r)  s     r*   resolve_atomic!CudaModuleTemplate.resolve_atomic	  r  r2   c                 J    [         R                  " [        R                  5      $ rH   )r   r9  r   r   r)  s     r*   resolve_fp16CudaModuleTemplate.resolve_fp16  s    ||DII&&r2   c                 J    [         R                  " [        R                  5      $ rH   )r   r9  r   rL   r)  s     r*   resolve_const CudaModuleTemplate.resolve_const      ||DJJ''r2   c                 J    [         R                  " [        R                  5      $ rH   )r   r9  r   rB   r)  s     r*   resolve_local CudaModuleTemplate.resolve_local  rY  r2   r-   N)&r4   r5   r6   r7   r   r9  r   r>   r  r  r  r  r  r  r  r  r  r  r!  r$  r'  r*  r.  r1  r4  r7  r;  r>  rB  rE  rH  rK  rN  rQ  rT  rW  r[  r8   r-   r2   r*   r  r    s    
,,t
C%)))((()767-8833/0).)'((r2   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   r:   r@   rD   rO   rU   rY   r]   rb   rj   rn   rr   rv   r{   r   r   r   r   r   r   r   r   r   r   r   floatr   r   r   r   r   r   haddr  addCuda_addiadd	Cuda_iaddhsubr  subCuda_subisub	Cuda_isubhmulr  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   r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r   r   r   ry   r   r   all_numba_typesr  unsigned_int_numba_typesr  rG  rL  maxri  minrm  nanmaxrv  nanminrq  and_rP  or_rT  xorrY  incr]  decra  exchre  r  r  r"  r&  r5  r;  r@  rE  r  r  r9  funcr-   r2   r*   <module>r     sq    @ @ @> > " *  @:&&**  (& 0 
  
 
  
 
,  
 
$. $ 
$
 
$- $ 
$
 
$. $ 
$
 
E$ E 
E
 
/  
 
6/ 6 
6 
*  
 
*  
 
&& & 
&
 
&' & 
&
 
   
$ 
	 	 
	 
   
 
   
 
   
 
  
$ 
  
$ 
(  ( 
((
 ' ' '0>/4 DIINN+	#HLL1$X]]3	DIINN+	#HLL1$X]]3	DIINN+	#HLL1$X]]3	DIINN+	DIINN+	499>>*	"8<<0499>>*	"3'%diimm4 X[[ )%diimm4 X[[ )%diimm4 X[[ )%diimm4 X[[ )%diimm4 X[[ )%diimm4 X[[ ) )) * ** +   %V,$V,$V,&x0%g.$V,&x0%g.%g.&x0&x0%g.$V,%g.&x0%f-& ==%--;;;;. {{ELL{{ELL2  "LL%,,7 t{{8t{{8t{{8t{{8$++,,o> $++,,o> t{{'')<=dkkoo':;t{{(;<t{{(@At{{(@A((*=>  
	1#3 	1 
	1 
6& 6 
6" 
2% 2 
2 
" 
 
 10 1 1 5/ 5 5 0/ 0 0 +/* +/ +/\ [)( [) [)| X(* X( X(v ell4( )
 $D/ $ !D/ ! $D/ $ D''T?3 r2   