+
    :i                     2   ^ RI Ht ^ RIt^ RIt^ RIHt ^ RIHt ^ RI	H
t
Ht ^ RIHt ^ RIHt ^ RIHtHt ^ RIHt ^ R	IHt ^R
IHt ^ RIHt ^ RIHtHtHt ^ RIH t H!t! ]
! 4       t"]"PF                  t#]"PH                  t%]"PL                  t&R t']%! ]PP                  ! ]4      R4      R 4       t)]%! ]PP                  ! ]4      R4      R 4       t*]%! ]PP                  ! ]4      R4      R 4       t+]%! ]PP                  ! ]4      R4      R 4       t,]%! ]PP                  ! ]4      R4      R 4       t-]%! ] R4      R 4       t.]%! ] R4      R 4       t/]%! ] R4      R 4       t0]#! ]Pb                  Pd                  ]Pf                  4      R 4       t4^ s5R  t6]#! ]Pn                  Pp                  ]Pr                  ]Pt                  4      R! 4       t;]#! ]Pn                  Pp                  ]Px                  ]Pt                  4      ]#! ]Pn                  Pp                  ]Pz                  ]Pt                  4      R" 4       4       t>]#! ]P~                  Pp                  ]Pr                  ]Pt                  4      R# 4       t@]#! ]P~                  Pp                  ]Px                  ]Pt                  4      ]#! ]P~                  Pp                  ]Pz                  ]Pt                  4      R$ 4       4       tA]#! ]P                  4      R% 4       tC]#! ]P                  4      R& 4       tE]#! ]P                  4      R' 4       tG]#! ]P                  4      R( 4       tI]#! ]P                  ]P                  4      R) 4       tK]#! ]P                  ]P                  ]P                  ]P                  ]P                  ]P                  4      ]#! ]P                  ]P                  ]P                  ]P                  ]P                  ]P                  4      ]#! ]P                  ]P                  ]P                  ]P                  ]P                  ]P                  4      ]#! ]P                  ]P                  ]P                  ]P                  ]P                  ]P                  4      R* 4       4       4       4       tP]#! ]P                  ]P                  ]P                  ]P                  4      R+ 4       tS]#! ]P                  ]P                  ]P                  4      ]#! ]P                  ]P                  ]P                  4      ]#! ]P                  ]P                  ]P                  4      ]#! ]P                  ]P                  ]P                  4      R, 4       4       4       4       tU]#! ]P                  ]P                  ]P                  4      ]#! ]P                  ]P                  ]P                  4      ]#! ]P                  ]P                  ]P                  4      ]#! ]P                  ]P                  ]P                  4      R- 4       4       4       4       tW]#! ]P                  4      R. 4       tY]#! ]P                  4      R/ 4       t[]#! ]P                  ]Pt                  4      R0 4       t]]#! ]P                  ]Pt                  ]Pt                  ]Pt                  4      R1 4       t_R2 t`]! ]P                  ]P                  4      R3 4       tc]! ]P                  ]P                  4      R4 4       tdR5 te]! ]P                  ]P                  4      R6 4       tg]! ]P                  ]P                  4      ]! ]Pr                  ]P                  4      R7 4       4       thR8 ti]i! ]P                  P                  R94       ]i! ]P                  R94       ]i! ]P                  R94       ]i! ]P                  P                  R:4       ]i! ]P                  R:4       ]i! ]P                  R:4       ]i! ]P                  P                  R;4       ]i! ]P                  R;4       ]i! ]P                  R;4       ]#! ]P                  P                  ]P                  4      R< 4       tu]#! ]P                  ]P                  4      R= 4       tw]#! ]P                  P                  ]P                  4      R> 4       ty]#! ]z]P                  4      R? 4       t{]#! ]P                  P                  ]P                  ]P                  ]P                  4      R@ 4       t}]#! ]P                  ]P                  ]P                  4      ]#! ]P                  ]P                  ]P                  4      RA 4       4       tRBtRC t]#! ]P                  EP                  ]P                  ]P                  4      ! ]! RD4      4       ]#! ]EP                  ]P                  ]P                  4      ! ]! RD4      4       ]#! ]P                  EP
                  ]P                  ]P                  4      ! ]! RE4      4       ]#! ]EP                  ]P                  ]P                  4      ! ]! RE4      4       ]#! ]P                  EP                  ]P                  ]P                  4      ! ]! RF4      4       ]#! ]EP                  ]P                  ]P                  4      ! ]! RF4      4       ]#! ]P                  EP                  ]P                  ]P                  4      ! ]! RG4      4       ]#! ]EP                  ]P                  ]P                  4      ! ]! RG4      4       ]#! ]P                  EP                  ]P                  ]P                  4      ! ]! RH4      4       ]#! ]EP                  ]P                  ]P                  4      ! ]! RH4      4       ]#! ]P                  EP                  ]P                  ]P                  4      ! ]! RI4      4       ]#! ]EP                  ]P                  ]P                  4      ! ]! RI4      4       RJ t]! ]P                  EP                   RKRG4       ]! ]P                  EP"                  RLRI4       ]EP$                  RM]EP&                  RN/t]#! ]EP*                  ]EP$                  4      ]#! ]EP*                  ]EP&                  4      RO 4       4       t]#! ]EP.                  ]EP0                  4      RP 4       t]#! ]EP.                  ]EP4                  4      RQ 4       t]#! ]EP8                  ]Pt                  4      RR 4       t]#! ]EP<                  ]P                  4      ]#! ]EP<                  ]EP0                  4      RS 4       4       t]#! ]EP<                  ]P                  4      ]#! ]EP<                  ]EP4                  4      RT 4       4       t]#! ]EPB                  ]Pt                  ]Pt                  ]Pt                  4      RU 4       t]#! ]]P                  ]P                  4      RV 4       t]#! ]]P                  ]P                  4      ]#! ]]P                  ]P                  4      ]#! ]]P                  ]P                  4      RW 4       4       4       t]#! ]]P                  ]P                  4      RX 4       t]#! ]]P                  ]P                  4      ]#! ]]P                  ]P                  4      ]#! ]]P                  ]P                  4      RY 4       4       4       t]#! ]]P                  4      ]#! ]]P                  4      RZ 4       4       t]#! ]]P                  ]P                  4      ]#! ]]P                  ]P                  4      R[ 4       4       tR\ t]EPZ                  R],          tR]]EPZ                  ,          t]#! ]EP`                  ]P                  4      ! ]! ]4      4       ]#! ]EP`                  ]P                  4      ! ]! ]4      4       ]#! ]EPb                  ]P                  4      ! ]! ]4      4       ]#! ]EPb                  ]P                  4      ! ]! ]4      4       R^ tR_ t]#! ]EPh                  P                  ]Pf                  ]EPj                  ]Pt                  4      ]#! ]EPh                  P                  ]Pf                  ]Pz                  ]Pt                  4      ]#! ]EPh                  P                  ]Pf                  ]Px                  ]Pt                  4      ]R` 4       4       4       4       t]#! ]EPh                  P                  ]Pf                  ]EPj                  ]Pt                  4      ]#! ]EPh                  P                  ]Pf                  ]Pz                  ]Pt                  4      ]#! ]EPh                  P                  ]Pf                  ]Px                  ]Pt                  4      ]Ra 4       4       4       4       t]#! ]EPh                  EPp                  ]Pf                  ]EPj                  ]Pt                  4      ]#! ]EPh                  EPp                  ]Pf                  ]Pz                  ]Pt                  4      ]#! ]EPh                  EPp                  ]Pf                  ]Px                  ]Pt                  4      ]Rb 4       4       4       4       t]#! ]EPh                  EPt                  ]Pf                  ]EPj                  ]Pt                  4      ]#! ]EPh                  EPt                  ]Pf                  ]Pz                  ]Pt                  4      ]#! ]EPh                  EPt                  ]Pf                  ]Px                  ]Pt                  4      ]Rc 4       4       4       4       tRd t]! ]EPh                  EPz                  Re4       ]! ]EPh                  EP|                  Rf4       ]! ]EPh                  EP~                  Rg4       ]#! ]EPh                  EP                  ]Pf                  ]EPj                  ]Pt                  4      ]#! ]EPh                  EP                  ]Pf                  ]Pz                  ]Pt                  4      ]#! ]EPh                  EP                  ]Pf                  ]Px                  ]Pt                  4      ]Rh 4       4       4       4       t]#! ]EPh                  EPF                  ]Pf                  ]EPj                  ]Pt                  4      ]#! ]EPh                  EPF                  ]Pf                  ]Px                  ]Pt                  4      ]#! ]EPh                  EPF                  ]Pf                  ]Pz                  ]Pt                  4      ]Ri 4       4       4       4       t]#! ]EPh                  EPL                  ]Pf                  ]EPj                  ]Pt                  4      ]#! ]EPh                  EPL                  ]Pf                  ]Px                  ]Pt                  4      ]#! ]EPh                  EPL                  ]Pf                  ]Pz                  ]Pt                  4      ]Rj 4       4       4       4       t]#! ]EPh                  EP                  ]Pf                  ]EPj                  ]Pt                  4      ]#! ]EPh                  EP                  ]Pf                  ]Px                  ]Pt                  4      ]#! ]EPh                  EP                  ]Pf                  ]Pz                  ]Pt                  4      ]Rk 4       4       4       4       t]#! ]EPh                  EP                  ]Pf                  ]EPj                  ]Pt                  4      ]#! ]EPh                  EP                  ]Pf                  ]Px                  ]Pt                  4      ]#! ]EPh                  EP                  ]Pf                  ]Pz                  ]Pt                  4      ]Rl 4       4       4       4       t]#! ]EPh                  EP                  ]Pf                  ]Pt                  ]Pt                  4      Rm 4       t]#! ]EPh                  EP                  ]Pf                  ]EPj                  ]Pt                  ]Pt                  4      ]#! ]EPh                  EP                  ]Pf                  ]Px                  ]Pt                  ]Pt                  4      ]#! ]EPh                  EP                  ]Pf                  ]Pz                  ]Pt                  ]Pt                  4      Rn 4       4       4       t]#! ]EP                  ]EP                  4      Ro 4       tRrRp lt]&! ]!4      Rq 4       t]! ]EP                  ! 4       ]#4       R# )s    )reduceN)ir)Registry
lower_cast)parse_dtype)models)typescgutils)ufunc_db)register_ufuncs)nvvm)cuda)	nvvmutilsstubserrors)dim3CUDADispatcherc                     \         P                  ! V R V,          4      p\         P                  ! V RV,          4      p\         P                  ! V RV,          4      p\        P                  ! WW434      # )z%s.xz%s.yz%s.z)r   	call_sregr
   pack_struct)builderprefixxyzs   &&   Q/var/www/html/photoedit/myenv/lib/python3.14/site-packages/numba/cuda/cudaimpl.pyinitialize_dim3r      sY    GVf_5AGVf_5AGVf_5AwA	22    	threadIdxc                     \        VR 4      # )tidr   contextr   sigargss   &&&&r   cuda_threadIdxr'       s    7E**r   blockDimc                     \        VR 4      # )ntidr"   r#   s   &&&&r   cuda_blockDimr+   %   s    7F++r   blockIdxc                     \        VR 4      # )ctaidr"   r#   s   &&&&r   cuda_blockIdxr/   *   s    7G,,r   gridDimc                     \        VR 4      # )nctaidr"   r#   s   &&&&r   cuda_gridDimr3   /   s    7H--r   laneidc                 0    \         P                  ! VR 4      # )r4   )r   r   r#   s   &&&&r   cuda_laneidr6   4   s    w11r   r   c                 &    VP                  V^ 4      # r   extract_valuer#   s   &&&&r   dim3_xr;   9         q))r   r   c                 &    VP                  V^4      # )   r9   r#   s   &&&&r   dim3_yr?   >   r<   r   r   c                 &    VP                  V^4      # )   r9   r#   s   &&&&r   dim3_zrB   C   r<   r   c                     V^ ,          # r8    r#   s   &&&&r   cuda_const_array_likerE   J   s     7Nr   c                H    \         ^,          s RP                  V \         4      # )zDue to bug with NVVM invalid internalizing of shared memory in the
PTX output.  We can't mark shared memory to be internal. We have to
ensure unique name is generated for shared memory symbol.
z{0}_{1})_unique_smem_idformatnames   &r   _get_unique_smem_idrK   T   s!     qOD/22r   c           
          VP                   ^ ,          P                  p\        VP                   ^,          4      p\        WV3V\	        R4      \
        P                  RR7      # r   _cudapy_smemTshapedtypesymbol_name	addrspacecan_dynsizedr&   literal_valuer   _generic_arrayrK   r   ADDRSPACE_SHAREDr$   r   r%   r&   lengthrQ   s   &&&&  r   cuda_shared_array_integerr[   ^   sP    XXa[&&F$E'6)5&9.&I$($9$9'+- -r   c           
          VP                   ^ ,           Uu. uF  qDP                  NK  	  pp\        VP                   ^,          4      p\        WWV\	        R4      \
        P                  RR7      # u upi rM   rU   r$   r   r%   r&   srP   rQ   s   &&&&   r   cuda_shared_array_tupler_   h   s^     (+xx{4{!oo{E4$E'%&9.&I$($9$9'+- - 5s   A/c           
          VP                   ^ ,          P                  p\        VP                   ^,          4      p\        WV3VR\        P
                  RR7      # r   _cudapy_lmemFrO   r&   rV   r   rW   r   ADDRSPACE_LOCALrY   s   &&&&  r   cuda_local_array_integerre   s   sK    XXa[&&F$E'6)5&4$($8$8',. .r   c           
          VP                   ^ ,           Uu. uF  qDP                  NK  	  pp\        VP                   ^,          4      p\        WWVR\        P
                  RR7      # u upi ra   rc   r]   s   &&&&   r   ptx_lmem_alloc_arrayrg   }   sY     (+xx{4{!oo{E4$E'%&4$($8$8',. . 5s   A&c                     V'       d   Q hR pVP                   p\        P                  ! \        P                  ! 4       R4      p\        P
                  ! WVV4      pVP                  VR4       V P                  4       # )zllvm.nvvm.membar.ctarD   moduler   FunctionTypeVoidTyper
   get_or_insert_functioncallget_dummy_valuer$   r   r%   r&   fnamelmodfntysyncs   &&&&    r   ptx_threadfence_blockru      \    O8"E>>D??2;;="-D))$e<DLLr""$$r   c                     V'       d   Q hR pVP                   p\        P                  ! \        P                  ! 4       R4      p\        P
                  ! WVV4      pVP                  VR4       V P                  4       # )zllvm.nvvm.membar.sysrD   ri   rp   s   &&&&    r   ptx_threadfence_systemrx      rv   r   c                     V'       d   Q hR pVP                   p\        P                  ! \        P                  ! 4       R4      p\        P
                  ! WVV4      pVP                  VR4       V P                  4       # )zllvm.nvvm.membar.glrD   ri   rp   s   &&&&    r   ptx_threadfence_devicerz      s\    O8!E>>D??2;;="-D))$e<DLLr""$$r   c                     V P                  \        P                  R 4      p\        P                  ! \        P                  4      p\	        WWT.4      # )l    )get_constantr	   int32noneptx_syncwarp_mask)r$   r   r%   r&   maskmask_sigs   &&&&  r   ptx_syncwarpr      s9    Z8Dzz%++&HWx@@r   c                    R pVP                   p\        P                  ! \        P                  ! 4       \        P                  ! ^ 4      34      p\
        P                  ! WVV4      pVP                  Ws4       V P                  4       # )zllvm.nvvm.bar.warp.sync)	rj   r   rk   rl   IntTyper
   rm   rn   ro   rp   s   &&&&    r   r   r      s^    %E>>D??2;;=2::b>*;<D))$e<DLL""$$r   c           
     :   Vw  rErgpVP                   ^,          p	V	\        P                  9   d1   VP                  V\        P
                  ! V	P                  4      4      pRp
VP                  p\        P                  ! \        P                  ! \        P
                  ! ^ 4      \        P
                  ! ^4      34      \        P
                  ! ^ 4      \        P
                  ! ^ 4      \        P
                  ! ^ 4      \        P
                  ! ^ 4      \        P
                  ! ^ 4      34      p\        P                  ! WV
4      pV	P                  ^ 8X  d   VP                  WWVWx34      pV	\        P                  8X  dc   VP                  V^ 4      pVP                  V^4      pVP                  V\        P                  ! 4       4      p\        P                   ! VVV34      pV# VP#                  V\        P
                  ! ^ 4      4      pVP%                  W`P'                  \        P(                  ^ 4      4      pVP#                  V\        P
                  ! ^ 4      4      pVP                  WVVWx34      pVP                  WVVWx34      pVP                  V^ 4      pVP                  V^ 4      pVP                  V^4      pVP+                  V\        P
                  ! ^@4      4      pVP+                  V\        P
                  ! ^@4      4      pVP-                  VV P'                  \        P(                  ^ 4      4      pVP/                  VV4      pV	\        P0                  8X  d&   VP                  V\        P2                  ! 4       4      p\        P                   ! WV34      pV# )aw  
The NVVM intrinsic for shfl only supports i32, but the cuda intrinsic
function supports both 32 and 64 bit ints and floats, so for feature parity,
i64, f32, and f64 are implemented. Floats by way of bitcasting the float to
an int, then shuffling, then bitcasting back. And 64-bit values by packing
them into 2 32bit values, shuffling thoose, and then packing back together.
zllvm.nvvm.shfl.sync.i32)r&   r	   real_domainbitcastr   r   bitwidthrj   rk   LiteralStructTyper
   rm   rn   float32r:   	FloatTypemake_anonymous_structtrunclshrr|   i8zextshlor_float64
DoubleType)r$   r   r%   r&   r   modevalueindexclamp
value_typerq   rr   rs   funcretrvpredfvvalue1
value_lshrvalue2ret1ret2rv1rv2rv1_64rv2_64rv_shls   &&&&                        r   ptx_shfl_sync_i32r      s     '+#De!JU&&&rzz*2E2E'FG%E>>D??
bjjnbjjm<=ZZ^RZZ^RZZ^ZZ^RZZ^=D
 ))$e<Db ll4U!BC&&&sA.B((a0DR\\^4B//"dDC" J ubjjn5\\%)=)=ehh)KL
z2::b>:||Dvu"DE||Dvu"DE##D!,##D!,$$T1-c2::b>2c2::b>2VW%9%9%((B%GH[[(&R]]_5B++G$Z@Jr   c                    R pVP                   p\        P                  ! \        P                  ! \        P                  ! ^ 4      \        P                  ! ^4      34      \        P                  ! ^ 4      \        P                  ! ^ 4      \        P                  ! ^4      34      p\
        P                  ! WVV4      pVP                  Ws4      # )zllvm.nvvm.vote.sync)rj   r   rk   r   r   r
   rm   rn   )r$   r   r%   r&   rq   rr   rs   r   s   &&&&    r   ptx_vote_syncr      s    !E>>D??2//B13A1@ AJJrNBJJrNBJJqMJLD ))$e<D<<##r   c                    Vw  rEVP                   ^,          P                  pVP                   ^,          \        P                  9   d'   VP	                  V\
        P                  ! V4      4      pRP                  V4      pVP                  p\
        P                  ! \
        P                  ! ^ 4      \
        P                  ! ^ 4      \
        P                  ! V4      34      p	\        P                  ! WV4      p
VP                  WV34      # )r>   zllvm.nvvm.match.any.sync.i{})r&   r   r	   r   r   r   r   rH   rj   rk   r
   rm   rn   r$   r   r%   r&   r   r   widthrq   rr   rs   r   s   &&&&       r   ptx_match_any_syncr      s    
 KDHHQK  E
xx{e'''rzz%'89*11%8E>>D??2::b>BJJrNBJJu<M+NOD))$e<D<<Um,,r   c                 h   Vw  rEVP                   ^,          P                  pVP                   ^,          \        P                  9   d'   VP	                  V\
        P                  ! V4      4      pRP                  V4      pVP                  p\
        P                  ! \
        P                  ! \
        P                  ! ^ 4      \
        P                  ! ^4      34      \
        P                  ! ^ 4      \
        P                  ! V4      34      p	\        P                  ! WV4      p
VP                  WV34      # )r>   zllvm.nvvm.match.all.sync.i{})r&   r   r	   r   r   r   r   rH   rj   rk   r   r
   rm   rn   r   s   &&&&       r   ptx_match_all_syncr     s    
 KDHHQK  E
xx{e'''rzz%'89*11%8E>>D??2//B13A1@ AJJrNBJJu,=>@D ))$e<D<<Um,,r   c                     \         P                  ! \         P                  ! \         P                  ! ^ 4      . 4      RRRR7      pVP	                  V. 4      # )    zactivemask.b32 $0;=rTside_effectr   	InlineAsmrk   r   rn   r$   r   r%   r&   
activemasks   &&&& r   ptx_activemaskr     s>    boobjjnbA2DdLJ<<
B''r   c                     \         P                  ! \         P                  ! \         P                  ! ^ 4      . 4      RRRR7      pVP	                  V. 4      # )r   zmov.u32 $0, %lanemask_lt;r   Tr   r   r   s   &&&& r   ptx_lanemask_ltr   $  s@    boobjjnbA94*.0J <<
B''r   c                 2    VP                  V^ ,          4      # r8   )ctpopr#   s   &&&&r   ptx_popcr   ,  s    ==a!!r   c                 "    VP                   ! V!  # N)fmar#   s   &&&&r   ptx_fmar   1  s    ;;r   c                 v    ^ R^@R/p W,          #   \          d    RT  R2p\        P                  ! T4      hi ; i)r   z$Conversion between float16 and float unsupported)f32f)f64dKeyErrorr   CudaLoweringErrorr   typemapmsgs   &  r   float16_float_ty_constraintr   6  sL    <\2G,   ,4XJlK&&s++,s    (8c                 D   VP                   VP                   8X  d   V# \        VP                   4      w  rV\        P                  ! V P	                  V4      \        P
                  ! ^4      .4      p\        P                  ! VRV R2RV R24      pVP                  W.4      # )   zcvt..f16 $0, $1;=,h)r   r   r   rk   get_value_typer   r   rn   	r$   r   fromtytotyvalty
constraintrs   asms	   &&&&&    r   float16_to_float_castr   @  s    $--'
0?NB??711$7"**R.9IJD
,,ttB4|4*R6H
IC<<U##r   c                 B   VP                   VP                   8X  d   V# \        VP                   4      w  rV\        P                  ! \        P                  ! ^4      V P                  V4      .4      p\        P                  ! VRV R2RV 24      pVP                  W.4      # )r   cvt.rn.f16. $0, $1;=h,)r   r   r   rk   r   r   r   rn   r   s	   &&&&&    r   float_to_float16_castr   L  s    $--'
0ANB??2::b>G,B,B6,J+KLD
,,t{2$h73zl9K
LC<<U##r   c                 ~    ^R^R^ R^@R/p W,          #   \          d    RT  R2p\        P                  ! T4      hi ; i)   chrlz"Conversion between float16 and intr   r   r   s   &  r   float16_int_constraintr   X  sT    3CS"c3G,   ,28*LI&&s++,s    (<c                 :   VP                   p\        V4      pVP                  '       d   R MRp\        P                  ! V P                  V4      \        P                  ! ^4      .4      p\        P                  ! VRV V R2RV R24      p	VP                  W.4      # )r^   uzcvt.rni.r   r   r   )	r   r   signedr   rk   r   r   r   rn   
r$   r   r   r   r   r   r   
signednessrs   r   s
   &&&&&     r   float16_to_integer_castr   b  s    }}H'1JJ??711$7"**R.9IJD
,,t!*hZ|D:,b)+C <<U##r   c                 8   VP                   p\        V4      pVP                  '       d   R MRp\        P                  ! \        P
                  ! ^4      V P                  V4      .4      p\        P                  ! VRV V R2RV 24      p	VP                  W.4      # )r^   r   r   r   r   )	r   r   r   r   rk   r   r   r   rn   r   s
   &&&&&     r   integer_to_float16_castr   o  s     H'1J3J??2::b>#226:;=D
,,t$ZL
(CZL)+C <<U##r   c                 l   a \        V \        P                  \        P                  4      V3R  l4       pR# )c                   < \         P                  ! \         P                  ! ^4      \         P                  ! ^4      \         P                  ! ^4      .4      p\         P                  ! VS R2R4      pVP	                  WS4      # )r   z.f16 $0,$1,$2;=h,h,hr   rk   r   r   rn   )r$   r   r%   r&   rs   r   ops   &&&&  r   ptx_fp16_binary*lower_fp16_binary.<locals>.ptx_fp16_binary  s\    rzz"~ "

2

2?All4B4~!6A||C&&r   Nlowerr	   float16)fnr   r   s   &f r   lower_fp16_binaryr   ~  s&    
2u}}emm,' -'r   addsubmulc                     \         P                  ! \         P                  ! ^4      \         P                  ! ^4      .4      p\         P                  ! VRR4      pVP	                  WS4      # )r   zneg.f16 $0, $1;=h,hr   r$   r   r%   r&   rs   r   s   &&&&  r   ptx_fp16_hnegr    G    ??2::b>BJJrN+;<D
,,t.
7C<<""r   c                     \        WW#4      # r   )r  r#   s   &&&&r   operator_hnegr	        355r   c                     \         P                  ! \         P                  ! ^4      \         P                  ! ^4      .4      p\         P                  ! VRR4      pVP	                  WS4      # )r   zabs.f16 $0, $1;r  r   r  s   &&&&  r   ptx_fp16_habsr    r  r   c                     \        WW#4      # r   )r  r#   s   &&&&r   operator_habsr    r
  r   c                 ,   \         P                  ! ^4      \         P                  ! ^4      \         P                  ! ^4      .p\         P                  ! \         P                  ! ^4      V4      p\         P                  ! VRR4      pVP	                  Wc4      # )r   zfma.rn.f16 $0,$1,$2,$3;z=h,h,h,h)r   r   rk   r   rn   )r$   r   r%   r&   argtysrs   r   s   &&&&   r   ptx_hfmar    s`    jjnbjjnbjjn=F??2::b>62D
,,t6

CC<<""r   c                 ,    R  pV P                  WW#4      # )c                 @    \         P                  P                  W4      # r   )r   fp16hdiv)r   r   s   &&r   fp16_divfp16_div_impl.<locals>.fp16_div  s    yy~~a##r   compile_internal)r$   r   r%   r&   r  s   &&&& r   fp16_div_implr    s    $ ##GsAAr   z{{
          .reg .pred __$$f16_cmp_tmp;
          setp.{op}.f16 __$$f16_cmp_tmp, $1, $2;
          selp.u16 $0, 1, 0, __$$f16_cmp_tmp;
        }}c                    a  V 3R  lpV# )c                   < \         P                  ! \         P                  ! ^4      \         P                  ! ^4      \         P                  ! ^4      .4      p\         P                  ! V\        P                  S	R7      R4      pVP                  WS4      pV P                  \        P                  ^ 4      pVP                  V\         P                  ! ^4      4      pVP                  RW4      # )r   )r   r   z!=)r   rk   r   r   	_fp16_cmprH   rn   r|   r	   int16r   icmp_unsigned)
r$   r   r%   r&   rs   r   resultzero
int_resultr   s
   &&&&     r   ptx_fp16_comparison*_gen_fp16_cmp.<locals>.ptx_fp16_comparison  s    rzz"~

2

2/OPll4!1!1R!1!8(Cc(##EKK3__VRZZ^<
$$T:<<r   rD   )r   r#  s   f r   _gen_fp16_cmpr%    s    = r   eqnegegtleltc                 l   a \        V \        P                  \        P                  4      V3R  l4       pR# )c                 h   < \        S4      ! WW#4      pVP                  WC^ ,          V^,          4      # r8   )r%  select)r$   r   r%   r&   choicer   s   &&&& r   ptx_fp16_minmax*lower_fp16_minmax.<locals>.ptx_fp16_minmax  s-    r"7S?~~f1gtAw77r   Nr   )r   rq   r   r0  s   &&f r   lower_fp16_minmaxr2    s&    
2u}}emm,8 -8r   maxmin
__nv_cbrtf	__nv_cbrtc                     VP                   p\        V,          pV P                  V4      pVP                  p\        P
                  ! Wf.4      p\        P                  ! WxV4      p	VP                  W4      # r   )	return_type
cbrt_funcsr   rj   r   rk   r
   rm   rn   )
r$   r   r%   r&   r   rq   ftyrr   rs   r   s
   &&&&      r   ptx_cbrtr;    s`     
BrNE

 
 
$C>>D??3&D		'	'E	:B<<!!r   c           	          \         P                  ! VP                  \        P                  ! \        P
                  ! ^ 4      \        P
                  ! ^ 4      34      R4      pVP                  WC4      # )r   	__nv_brevr
   rm   rj   r   rk   r   rn   r$   r   r%   r&   r   s   &&&& r   ptx_brev_u4r@    sP    
 
	'	'


2B(9:
B <<!!r   c           	          \         P                  ! VP                  \        P                  ! \        P
                  ! ^@4      \        P
                  ! ^@4      34      R4      pVP                  WC4      # )@   __nv_brevllr>  r?  s   &&&& r   ptx_brev_u8rD  	  sP    
 
	'	'


2B(9:
B <<!!r   c                 p    VP                  V^ ,          V P                  \        P                  ^ 4      4      # r8   )ctlzr|   r	   booleanr#   s   &&&&r   ptx_clzrH    s.    <<QU]]A.0 0r   c           	          \         P                  ! VP                  \        P                  ! \        P
                  ! ^ 4      \        P
                  ! ^ 4      34      R4      pVP                  WC4      # )r   __nv_ffsr>  r?  s   &&&& r   
ptx_ffs_32rK    sP     
	'	'


2B(9:
B <<!!r   c           	          \         P                  ! VP                  \        P                  ! \        P
                  ! ^ 4      \        P
                  ! ^@4      34      R4      pVP                  WC4      # )r   
__nv_ffsllr>  r?  s   &&&& r   
ptx_ffs_64rN  &  sP     
	'	'


2B(9:
B <<!!r   c                 0    Vw  rEpVP                  WEV4      # r   )r.  )r$   r   r%   r&   testabs   &&&&   r   ptx_selprS  0  s    JDQ>>$1%%r   c           	         \         P                  ! VP                  \        P                  ! \        P
                  ! 4       \        P
                  ! 4       \        P
                  ! 4       34      R 4      pVP                  WC4      # )
__nv_fmaxfr
   rm   rj   r   rk   r   rn   r?  s   &&&& r   
ptx_max_f4rW  6  X    		'	'
LLN\\^R\\^,	. 	
B <<!!r   c           
         \         P                  ! VP                  \        P                  ! \        P
                  ! 4       \        P
                  ! 4       \        P
                  ! 4       34      R 4      pVP                  VV P                  W^ ,          VP                  ^ ,          \        P                  4      V P                  W^,          VP                  ^,          \        P                  4      .4      # )	__nv_fmaxr
   rm   rj   r   rk   r   rn   castr&   r	   doubler?  s   &&&& r   
ptx_max_f8r^  A       
	'	'
MMO]]_bmmo.	0 	
B <<W1gsxx{ELLAW1gsxx{ELLA  r   c           	         \         P                  ! VP                  \        P                  ! \        P
                  ! 4       \        P
                  ! 4       \        P
                  ! 4       34      R 4      pVP                  WC4      # )
__nv_fminfrV  r?  s   &&&& r   
ptx_min_f4rb  R  rX  r   c           
         \         P                  ! VP                  \        P                  ! \        P
                  ! 4       \        P
                  ! 4       \        P
                  ! 4       34      R 4      pVP                  VV P                  W^ ,          VP                  ^ ,          \        P                  4      V P                  W^,          VP                  ^,          \        P                  4      .4      # )	__nv_fminr[  r?  s   &&&& r   
ptx_min_f8re  ]  r_  r   c           	      T   \         P                  ! VP                  \        P                  ! \        P
                  ! ^@4      \        P                  ! 4       34      R4      pVP                  VV P                  W^ ,          VP                  ^ ,          \        P                  4      .4      # )rB  __nv_llrint)r
   rm   rj   r   rk   r   r   rn   r\  r&   r	   r]  r?  s   &&&& r   	ptx_roundrh  n  s}     
	'	'
JJrN]]_	  	
B <<W1gsxx{ELLA  r   c                 ,    R  pV P                  WW#4      # )c                    \         P                  ! V 4      '       g   \         P                  ! V 4      '       d   V # V^ 8  dT   V^8  d   RV^,
          ,          pRpMRV,          pRpW,          V,          p\         P                  ! V4      '       d   V # MRV) ,          pW,          p\        V4      p\         P                  ! WE,
          4      R8X  d   R\        VR,          4      ,          pV^ 8  d   VX,          V,          pV# WR,          pV# )r   g      $@gMDg      ?g      ?g       @)mathisinfisnanroundfabs)r   ndigitspow1pow2r   r   s   &&    r   round_ndigits$round_to_impl.<locals>.round_ndigits  s    ::a==DJJqMMHa<| "-wT!Azz!}}  WH%DA!HIIae#eAGn$Aa<TT!A  IAr   r  )r$   r   r%   r&   rs  s   &&&& r   round_to_implru    s    B ##GCHHr   c                    a  V 3R  lpV# )c                 t   < VP                   w  pV P                  VS4      pVP                  WS^ ,          4      # r8   )r&   r|   fmul)r$   r   r%   r&   argtyfactorconsts   &&&&  r   implgen_deg_rad.<locals>.impl  s3    %%eU3||FG,,r   rD   )r{  r|  s   f r   gen_deg_radr~    s    - Kr   g     f@c           
        V\         P                  9   d   \         P                  ! V^R7      pV.pM!\        P                  ! W\        V4      R7      p\        W&4       UUu. uF&  w  rxV P                  WV\         P                  4      NK(  	  pppVP                  p	W8w  d   \        RV	: RV: 24      hVP                  \        V4      8w  d(   \        RVP                  \        V4      3,          4      hW&3# u uppi )z,
Convert integer indices into tuple of intp
)rQ   count)r  zexpect z	 but got z#indexing %d-D array with %d-D index)r	   integer_domainUniTupler
   unpack_tuplelenzipr\  intprQ   	TypeErrorndim)
r$   r   indtyindsarytyvaltyindicestirQ   s
   &&&&&&    r   _normalize_indicesr    s     $$$U!4&&&wCJGu.0.41 ||G5::6.  0 KKE~%?@@zzSZ=SZ01 2 	2 >0s   !,C9c                    a  V 3R  lpV# )c           	         < VP                   w  rEpVw  rxp	VP                  p
\        WWXWF4      w  r[V P                  V4      ! WV4      p\        P
                  ! WWLVR R7      pS! WWV	4      # )T
wraparound)r&   rQ   r  
make_arrayr
   get_item_pointer)r$   r   r%   r&   r  r  r  aryr  r   rQ   r  laryptrdispatch_fns   &&&&          r   imp_atomic_dispatcher.<locals>.imp  sy    !hhe3+Ge,1: !!%(3?&&wg268 7U==r   rD   )r  r  s   f r   _atomic_dispatcherr    s    > Jr   c                 H   V\         P                  8X  d4   VP                  pVP                  \        P
                  ! V4      W434      # V\         P                  8X  d4   VP                  pVP                  \        P                  ! V4      W434      # VP                  R W4R4      # )r   	monotonic)	r	   r   rj   rn   r   declare_atomic_add_float32r   declare_atomic_add_float64
atomic_rmwr$   r   rQ   r  r   rr   s   &&&&& r   ptx_atomic_add_tupler        
 ~~||I@@F J( 	(	%--	~~||I@@F J( 	( !!%;??r   c                 H   V\         P                  8X  d4   VP                  pVP                  \        P
                  ! V4      W434      # V\         P                  8X  d4   VP                  pVP                  \        P                  ! V4      W434      # VP                  R W4R4      # )r  r  )	r	   r   rj   rn   r   declare_atomic_sub_float32r   declare_atomic_sub_float64r  r  s   &&&&& r   ptx_atomic_subr    r  r   c                     V\         P                  P                  9   dE   VP                  pVP                  p\        \        R V 24      pVP                  V! V4      W434      # \        RV R24      h)declare_atomic_inc_intzUnimplemented atomic inc with  array	r   cudadeclunsigned_int_numba_typesr   rj   getattrr   rn   r  r$   r   rQ   r  r   bwrr   r   s   &&&&&   r   ptx_atomic_incr    h    
 666^^~~Y"8 =>||BtHsj118vFGGr   c                     V\         P                  P                  9   dE   VP                  pVP                  p\        \        R V 24      pVP                  V! V4      W434      # \        RV R24      h)declare_atomic_dec_intzUnimplemented atomic dec with r  r  r  s   &&&&&   r   ptx_atomic_decr    r  r   c                    a \         V3R  l4       p\        P                  \        P                  \        P                  3 F3  p\        V \        P                  V\        P                  4      ! V4       K5  	  R# )c                    < V\         P                  P                  9   d   VP                  SW4R 4      # \	        RS RV R24      h)r  zUnimplemented atomic z with r  r   r  integer_numba_typesr  r  )r$   r   rQ   r  r   r   s   &&&&&r   impl_ptx_atomic+ptx_atomic_bitwise.<locals>.impl_ptx_atomic  sE    T]]667%%b#K@@3B4veWFKLLr   N)r  r	   r  r  Tupler   ArrayAny)stubr   r  r   s   &f  r   ptx_atomic_bitwiser    sQ    M M zz5>>5;;7dEKKUYY/@ 8r   andorxorc                     V\         P                  P                  9   d   VP                  R W4R4      # \	        RV R24      h)xchgr  zUnimplemented atomic exch with r  r  )r$   r   rQ   r  r   s   &&&&&r   ptx_atomic_exchr  /  s@    
 223!!&#K@@9%GHHr   c                    VP                   pV\        P                  8X  d(   VP                  \        P
                  ! V4      W434      # V\        P                  8X  d(   VP                  \        P                  ! V4      W434      # V\        P                  \        P                  39   d   VP                  R W4RR7      # V\        P                  \        P                  39   d   VP                  RW4RR7      # \        RV,          4      hr3  r  orderingumaxz&Unimplemented atomic max with %s array)rj   r	   r   rn   r   declare_atomic_max_float64r   declare_atomic_max_float32r}   int64r  uint32uint64r  r  s   &&&&& r   ptx_atomic_maxr  :      
 >>D||I@@F J( 	(	%--	||I@@F J( 	(	5;;,	,!!%K!HH	5<<.	.!!&#[!II@5HIIr   c                    VP                   pV\        P                  8X  d(   VP                  \        P
                  ! V4      W434      # V\        P                  8X  d(   VP                  \        P                  ! V4      W434      # V\        P                  \        P                  39   d   VP                  R W4RR7      # V\        P                  \        P                  39   d   VP                  RW4RR7      # \        RV,          4      hr4  r  r  uminz&Unimplemented atomic min with %s array)rj   r	   r   rn   r   declare_atomic_min_float64r   declare_atomic_min_float32r}   r  r  r  r  r  r  s   &&&&& r   ptx_atomic_minr  N  r  r   c                    VP                   pV\        P                  8X  d(   VP                  \        P
                  ! V4      W434      # V\        P                  8X  d(   VP                  \        P                  ! V4      W434      # V\        P                  \        P                  39   d   VP                  R W4RR7      # V\        P                  \        P                  39   d   VP                  RW4RR7      # \        RV,          4      hr  )rj   r	   r   rn   r   declare_atomic_nanmax_float64r   declare_atomic_nanmax_float32r}   r  r  r  r  r  r  s   &&&&& r   ptx_atomic_nanmaxr  b      
 >>D||ICCDI J( 	(	%--	||ICCDI J( 	(	5;;,	,!!%K!HH	5<<.	.!!&#[!II@5HIIr   c                    VP                   pV\        P                  8X  d(   VP                  \        P
                  ! V4      W434      # V\        P                  8X  d(   VP                  \        P                  ! V4      W434      # V\        P                  \        P                  39   d   VP                  R W4RR7      # V\        P                  \        P                  39   d   VP                  RW4RR7      # \        RV,          4      hr  )rj   r	   r   rn   r   declare_atomic_nanmin_float64r   declare_atomic_nanmin_float32r}   r  r  r  r  r  r  s   &&&&& r   ptx_atomic_nanminr  v  r  r   c                 6   VP                  VP                  ^ ,          \        P                  VP                  ^,          VP                  ^,          4      pV^ ,          V P	                  \        P                  ^ 4      V^,          V^,          3p\        WW#4      # r8   )r8  r&   r	   r  r|   ptx_atomic_casr#   s   &&&&r   ptx_atomic_compare_and_swapr    sj    
//#((1+uzz388A;
LCGW))%**a8$q'47KD'C66r   c           	         VP                   w  rErgVw  rr\        WWYVV4      w  r\V P                  V4      ! WV4      p\        P                  ! WWMVR R7      pVP
                  \        P                  P                  9   d<   VP                  pVP
                  P                  p\        P                  ! WVWV4      # \        RVP
                  ,          4      h)Tr  z&Unimplemented atomic cas with %s array)r&   r  r  r
   r  rQ   r   r  r  rj   r   r   atomic_cmpxchgr  )r$   r   r%   r&   r  r  oldtyr  r  r  oldr   r  r  r  rr   r   s   &&&&             r   r  r    s     "%E%Cs'%u(-/NE e$Ws;D

"
"7U'.24C {{t}}889~~;;''''x3OO@5;;NOOr   c                     \         P                  ! \         P                  ! \         P                  ! 4       \         P                  ! ^ 4      .4      RRRR7      pV^ ,          pVP                  WE.4       R# )r   znanosleep.u32 $0;r   Tr   N)r   r   rk   rl   r   rn   )r$   r   r%   r&   	nanosleepnss   &&&&  r   ptx_nanosleepr    sM    R__R[[]RZZ^<LM0#4II	aBLLD!r   c           
      .	   \        \        P                  V^4      pV^ 8*  ;'       d    T;'       d    \        V4      ^8H  pV^ 8:  d   V'       g   \	        R4      hV P
                  V,          p	\        V\        P                  \        P                  34      ;'       g6    \        V	\        P                  4      ;'       g    V\        P                  8H  p
V\        P                  9  d   V
'       g   \        RV,          4      hV P                  V4      p\         P"                  ! W4      pV\$        P&                  8X  d   \(        P*                  ! WVR7      pMVP,                  p\(        P.                  ! WVV4      pV P1                  V4      p^V^,
          P3                  4       ,          Vn        V'       d	   RVn        M*\         P8                  ! V\         P:                  4      Vn        VP?                  V\         P@                  ! \         PB                  ! ^4      4      R4      p\D        PF                  ! \$        PH                  ! 4       PJ                  4      pV P                  V4      pVPM                  V4      pTp. p\O        \Q        V4      4       F   w  ppVPS                  V4       VV,          pK"  	  \Q        V4       Uu. uF  pVNK  	  ppV Uu. uF#  pV PU                  \        PV                  V4      NK%  	  ppV'       d   \         PX                  ! \         PZ                  ! \         PB                  ! ^ 4      . 4      RRRR	7      pVP]                  VP_                  V. 4      \         PB                  ! ^@4      4      pV PU                  \        PV                  V4      pVPa                  VV4      .pM/V Uu. uF#  pV PU                  \        PV                  V4      NK%  	  pp\        V4      p\        Pb                  ! VVR
R7      pV Pe                  V4      ! W4      p V Pg                  V VPi                  VV Pj                  Pl                  4      VVV PU                  \        PV                  V4      RR7       V Po                  4       # u upi u upi u upi )r>   zarray length <= 0zunsupported type: %srI   externalgenericzmov.u32 $0, %dynamic_smem_size;r   Tr   C)rQ   r  layoutN)datarP   stridesitemsizememinfo)8r   operatorr  r  
ValueErrordata_model_manager
isinstancer	   RecordBooleanr   StructModelr   number_domainr  get_data_typer   	ArrayTyper   rd   r
   alloca_oncerj   add_global_variableget_abi_sizeof
bit_lengthalignlinkageConstant	UndefinedinitializeraddrspacecastPointerTyper   llcreate_target_dataNVVMdata_layoutget_abi_size	enumeratereversedappendr|   r  r   rk   r   rn   udivr  r  populate_arrayr   r  type	_getvalue)!r$   r   rP   rQ   rR   rS   rT   	elemcountdynamic_smem
data_modelother_supported_typelldtypelarytydataptrrr   gvmemr  
targetdatar  
laststriderstridesr  lastsizer^   r  kstridesget_dynshared_sizedynsmem_size	kitemsizekshaper  r  r  s!   &&&&&&&                          r   rW   rW     s   x||UA.I >FFlFFs5zQLA~l,-- ++E2J55<<78 	" 	"j&"4"45	" 	"EMM! 
 E'''0D.677##E*G\\'-FD((( %%gKH~~ ++D+,57 &&w/ EAI2244&EM !#FBLL AE ''r~~bjjm/L(13 &&tyy{'>'>?J##E*G##J/H JH %18
#h
 2 #8,-,Qq,G-=DEW$$UZZ3WHE 
  \\"//"**R."*M*K*.DB ||GLL1CR$H$&JJrN4 ((X>	,,|Y78?DEu!'&&uzz15uE u:DKKe$s;E


U
#G
5C3 ' G!'#+$+$8$8X$N#'  ) ==?A .E$ Fs   R)R9)Rc                 "    V P                  4       # r   )ro   )r$   r   r   pyvals   &&&&r   cuda_dispatcher_constr,    s    ""$$r   )F)	functoolsr   r  rk  llvmliter   llvmlite.bindingbindingr  numba.core.imputilsr   r   numba.core.typing.npydeclr   numba.core.datamodelr   
numba.corer	   r
   numba.npr   numba.np.npyimplr   cudadrvr   numbar   
numba.cudar   r   r   numba.cuda.typesr   r   registryr   lower_getattr
lower_attrlower_constantr   Moduler'   r+   r/   r3   r6   r;   r?   rB   r{  
array_liker  rE   rG   rK   sharedarrayIntegerLiteralr  r[   r  r  r_   localre   rg   threadfence_blockru   threadfence_systemrx   threadfencerz   syncwarpr   i4r   shfl_sync_intrinsicr   f4f8r   vote_sync_intrinsicrG  r   match_any_syncr   match_all_syncr   r   r   lanemask_ltr   popcr   r   r   r   r   Floatr   r   r   Integerr   r   r   r  haddr   iaddhsubr  isubhmulr  imulhnegr  negr	  habsr  absr  hfmar  truedivitruedivr  r  r%  heqr&  hner'  hger(  hgtr)  hler*  hltr+  r2  hmaxhminr   r   r9  cbrtr;  brevu4r@  u8rD  clzrH  ffsrK  rN  selprS  r3  rW  r^  r4  rb  re  rn  rh  ru  r~  pi_deg2rad_rad2degradiansdegreesr  r  atomicr  r  r  incr  decr  r  and_r   r  exchr  r  r  nanmaxr  nanminr  compare_and_swapr  casr  r  r  r  rW   r,  
get_ufuncsrD   r   r   <module>r     s(        4 1 ' %  ,   / / 1:##
((3 ELL,+ -+ ELL
+, ,, ELL
+- ,- ELL	*. +. ELL)2 *2 D#* * D#* * D#* * tzzekk* + 3 t{{%..		:- ;- t{{%++uyy1t{{%..%))4- 5 2- tzz--uyy9. :. tzzeii0tzz3. 4 1. u%  % u % !% u% % u~~A A u~~uxx % !% u  %((EHHehhxxu  %((EHHehhxxu  %((EHHehhxxu  %((EHHehhxx++\ u  %((EHHemmD$ E$ uUXXuxx0uUXXuxx0uUXXuxx0uUXXuxx0	- 1 1 1 1	- uUXXuxx0uUXXuxx0uUXXuxx0uUXXuxx0- 1 1 1 1- u( ( u( ( uzz599" " uyy%))UYY		2 3, EMM5;;'$ ($ EKK'$ ($, EMM5==)	$ *	$ EMM5==)E  %--0
$ 1 *
$' %**//5 ) (,, & (-- ' %**//5 ) (,, & (-- ' %**//5 ) (,, & (-- ' uzz&# '# x||U]]#6 $6 uzz&# '# sEMM6 6 uzzu}}emmD# E# x6x%--7B 8 7B		 ejjnnemmU]] 3M$4G H hkk5==%-- 0t1D E ejjnnemmU]] 3M$4G H hkk5==%-- 0t1D E ejjnnemmU]] 3M$4G H hkk5==%-- 0t1D E ejjnnemmU]] 3M$4G H hkk5==%-- 0t1D E ejjnnemmU]] 3M$4G H hkk5==%-- 0t1D E ejjnnemmU]] 3M$4G H hkk5==%-- 0t1D E8 %**//5$ / %**//5$ / 
MM<	MM;
 uzz5==!uzz5==!" " "" uzz588" " uzz588" " uyy%))0 0 uyy%((uyy%(("  " uyy%((uyy%(("  " uzz599eii3& 4&
 sEHHehh"  " sEHHehhsEHHehhsEHHehh       sEHHehh"  " sEHHehhsEHHehhsEHHehh       uehhuehh	  	  uehh&uehh&"I ' '"IJ 77T>$''> dllEHH k(3 4 dllEHH k(3 4 dllEHH k(3 4 dllEHH k(3 4.$ u||ejj%))<u||enneii@u||ekk599=
@  > A =
@ u||ejj%))<u||enneii@u||ekk599=
@  > A =
@ u||ejj%))<u||enneii@u||ekk599=H  > A =H u||ejj%))<u||enneii@u||ekk599=H  > A =H	A 5<<$$e , 5<<##T * 5<<##U + u||%++uzz599=u||%++u~~uyyAu||%++u{{EII>I  ? B >I u||ejj%))<u||ekk599=u||enneii@J  A > =J  u||ejj%))<u||ekk599=u||enneii@J  A > =J  u||EKKUYY?u||EKKeii@u||EKKCJ  D A @J  u||EKKUYY?u||EKKeii@u||EKKCJ  D A @J  u||$$ekk599eiiH7 I7 u||ejj%))UYYGu||ekk599eiiHu||enneiiKP L I HP* u%" &"aH %  % ##%u -r   