Ë
    èÿæiä–  ã                   óÚ2  — d dl mZ d dlZd dlZd dlmZ d dlmZ d dl	m
Z
mZ d dlmZ d dlmZ d dlmZmZ d dlmZ d d	lmZ d
dlmZ d dlmZ d dlmZmZmZ d dlm Z m!Z!  e
«       Z"e"jF                  Z#e"jH                  Z%e"jL                  Z&d„ Z' e% ejP                  e«      d«      d„ «       Z) e% ejP                  e«      d«      d„ «       Z* e% ejP                  e«      d«      d„ «       Z+ e% ejP                  e«      d«      d„ «       Z, e% ejP                  e«      d«      d„ «       Z- e%e d«      d„ «       Z. e%e d«      d„ «       Z/ e%e d«      d„ «       Z0 e#ejb                  jd                  ejf                  «      d „ «       Z4d a5d!„ Z6 e#ejn                  jp                  ejr                  ejt                  «      d"„ «       Z; e#ejn                  jp                  ejx                  ejt                  «       e#ejn                  jp                  ejz                  ejt                  «      d#„ «       «       Z> e#ej~                  jp                  ejr                  ejt                  «      d$„ «       Z@ e#ej~                  jp                  ejx                  ejt                  «       e#ej~                  jp                  ejz                  ejt                  «      d%„ «       «       ZA e#ej„                  «      d&„ «       ZC e#ejˆ                  «      d'„ «       ZE e#ejŒ                  «      d(„ «       ZG e#ej�                  «      d)„ «       ZI e#ej�                  ej”                  «      d*„ «       ZK e#ej˜                  ej”                  ej”                  ej”                  ej”                  ej”                  «       e#ej˜                  ej”                  ej”                  ejš                  ej”                  ej”                  «       e#ej˜                  ej”                  ej”                  ejœ                  ej”                  ej”                  «       e#ej˜                  ej”                  ej”                  ejž                  ej”                  ej”                  «      d+„ «       «       «       «       ZP e#ej¢                  ej”                  ej”                  ej¤                  «      d,„ «       ZS e#ej¨                  ej”                  ej”                  «       e#ej¨                  ej”                  ejš                  «       e#ej¨                  ej”                  ejœ                  «       e#ej¨                  ej”                  ejž                  «      d-„ «       «       «       «       ZU e#ej¬                  ej”                  ej”                  «       e#ej¬                  ej”                  ejš                  «       e#ej¬                  ej”                  ejœ                  «       e#ej¬                  ej”                  ejž                  «      d.„ «       «       «       «       ZW e#ej°                  «      d/„ «       ZY e#ej´                  «      d0„ «       Z[ e#ej¸                  ejt                  «      d1„ «       Z] e#ej¼                  ejt                  ejt                  ejt                  «      d2„ «       Z_d3„ Z` eejÂ                  ejÄ                  «      d4„ «       Zc eejÄ                  ejÂ                  «      d5„ «       Zdd6„ Ze eejÂ                  ejÌ                  «      d7„ «       Zg eejÌ                  ejÂ                  «       eejr                  ejÂ                  «      d8„ «       «       Zhd9„ Zi eiejÔ                  jÖ                  d:«        eiejØ                  d:«        eiejÚ                  d:«        eiejÔ                  jÜ                  d;«        eiejÞ                  d;«        eiejà                  d;«        eiejÔ                  jâ                  d<«        eiejä                  d<«        eiejæ                  d<«        e#ejÔ                  jè                  ejÂ                  «      d=„ «       Zu e#ejì                  ejÂ                  «      d>„ «       Zw e#ejÔ                  jð                  ejÂ                  «      d?„ «       Zy e#ezejÂ                  «      d@„ «       Z{ e#ejÔ                  jø                  ejÂ                  ejÂ                  ejÂ                  «      dA„ «       Z} e#ejü                  ejÂ                  ejÂ                  «       e#ejþ                  ejÂ                  ejÂ                  «      dB„ «       «       Z€dCZ�dD„ Z‚  e#ejÔ                  �j                  ejÂ                  ejÂ                  «       e‚dE«      «         e#e�j                  ejÂ                  ejÂ                  «       e‚dE«      «         e#ejÔ                  �j
                  ejÂ                  ejÂ                  «       e‚dF«      «         e#e�j                  ejÂ                  ejÂ                  «       e‚dF«      «         e#ejÔ                  �j                  ejÂ                  ejÂ                  «       e‚dG«      «         e#e�j                  ejÂ                  ejÂ                  «       e‚dG«      «         e#ejÔ                  �j                  ejÂ                  ejÂ                  «       e‚dH«      «         e#e�j                  ejÂ                  ejÂ                  «       e‚dH«      «         e#ejÔ                  �j                  ejÂ                  ejÂ                  «       e‚dI«      «         e#e�j                  ejÂ                  ejÂ                  «       e‚dI«      «         e#ejÔ                  �j                  ejÂ                  ejÂ                  «       e‚dJ«      «         e#e�j                  ejÂ                  ejÂ                  «       e‚dJ«      «       dK„ Z� e�ejÔ                  �j                   dLdH«        e�ejÔ                  �j"                  dMdJ«       e�j$                  dNe�j&                  dOiZ” e#e�j*                  e�j$                  «       e#e�j*                  e�j&                  «      dP„ «       «       Z– e#e�j.                  e�j0                  «      dQ„ «       Z™ e#e�j.                  e�j4                  «      dR„ «       Z› e#e�j8                  ejt                  «      dS„ «       Z� e#e�j<                  ej”                  «       e#e�j<                  e�j0                  «      dT„ «       «       ZŸ e#e�j<                  ejš                  «       e#e�j<                  e�j4                  «      dU„ «       «       Z  e#e�jB                  ejt                  ejt                  ejt                  «      dV„ «       Z¢ e#e£ejœ                  ejœ                  «      dW„ «       Z¤ e#e£ejž                  ejœ                  «       e#e£ejœ                  ejž                  «       e#e£ejž                  ejž                  «      dX„ «       «       «       Z¥ e#e¦ejœ                  ejœ                  «      dY„ «       Z§ e#e¦ejž                  ejœ                  «       e#e¦ejœ                  ejž                  «       e#e¦ejž                  ejž                  «      dZ„ «       «       «       Z¨ e#e©ejœ                  «       e#e©ejž                  «      d[„ «       «       Zª e#e©ejœ                  ejÌ                  «       e#e©ejž                  ejÌ                  «      d\„ «       «       Z«d]„ Z¬e�jZ                  d^z  Z®d^e�jZ                  z  Z¯  e#e�j`                  ejœ                  «       e¬e®«      «         e#e�j`                  ejž                  «       e¬e®«      «         e#e�jb                  ejœ                  «       e¬e¯«      «         e#e�jb                  ejž                  «       e¬e¯«      «       d_„ Z²d`„ Z³ e#e�jh                  jØ                  ejf                  e�jj                  ejt                  «       e#e�jh                  jØ                  ejf                  ejz                  ejt                  «       e#e�jh                  jØ                  ejf                  ejx                  ejt                  «      e³da„ «       «       «       «       Z¶ e#e�jh                  jÞ                  ejf                  e�jj                  ejt                  «       e#e�jh                  jÞ                  ejf                  ejz                  ejt                  «       e#e�jh                  jÞ                  ejf                  ejx                  ejt                  «      e³db„ «       «       «       «       Z· e#e�jh                  �jp                  ejf                  e�jj                  ejt                  «       e#e�jh                  �jp                  ejf                  ejz                  ejt                  «       e#e�jh                  �jp                  ejf                  ejx                  ejt                  «      e³dc„ «       «       «       «       Z¹ e#e�jh                  �jt                  ejf                  e�jj                  ejt                  «       e#e�jh                  �jt                  ejf                  ejz                  ejt                  «       e#e�jh                  �jt                  ejf                  ejx                  ejt                  «      e³dd„ «       «       «       «       Z»de„ Z¼ e¼e�jh                  �jz                  df«        e¼e�jh                  �j|                  dg«        e¼e�jh                  �j~                  dh«        e#e�jh                  �j€                  ejf                  e�jj                  ejt                  «       e#e�jh                  �j€                  ejf                  ejz                  ejt                  «       e#e�jh                  �j€                  ejf                  ejx                  ejt                  «      e³di„ «       «       «       «       ZÁ e#e�jh                  �jF                  ejf                  e�jj                  ejt                  «       e#e�jh                  �jF                  ejf                  ejx                  ejt                  «       e#e�jh                  �jF                  ejf                  ejz                  ejt                  «      e³dj„ «       «       «       «       ZÂ e#e�jh                  �jL                  ejf                  e�jj                  ejt                  «       e#e�jh                  �jL                  ejf                  ejx                  ejt                  «       e#e�jh                  �jL                  ejf                  ejz                  ejt                  «      e³dk„ «       «       «       «       ZÃ e#e�jh                  �jˆ                  ejf                  e�jj                  ejt                  «       e#e�jh                  �jˆ                  ejf                  ejx                  ejt                  «       e#e�jh                  �jˆ                  ejf                  ejz                  ejt                  «      e³dl„ «       «       «       «       ZÅ e#e�jh                  �jŒ                  ejf                  e�jj                  ejt                  «       e#e�jh                  �jŒ                  ejf                  ejx                  ejt                  «       e#e�jh                  �jŒ                  ejf                  ejz                  ejt                  «      e³dm„ «       «       «       «       ZÇ e#e�jh                  �j�                  ejf                  ejt                  ejt                  «      dn„ «       ZÉ e#e�jh                  �j”                  ejf                  e�jj                  ejt                  ejt                  «       e#e�jh                  �j”                  ejf                  ejx                  ejt                  ejt                  «       e#e�jh                  �j”                  ejf                  ejz                  ejt                  ejt                  «      do„ «       «       «       ZË e#e�j˜                  e�jš                  «      dp„ «       ZÎ	 dsdq„ZÏ e&e!«      dr„ «       ZÐ e e�j¢                  «       e#«       y)té    )ÚreduceN)Úir)ÚRegistryÚ
lower_cast)Úparse_dtype)Úmodels)ÚtypesÚcgutils)Úufunc_db)Úregister_ufuncsé   )Únvvm)Úcuda)Ú	nvvmutilsÚstubsÚerrors)Údim3ÚCUDADispatcherc                 óÊ   — t        j                  | d|z  «      }t        j                  | d|z  «      }t        j                  | d|z  «      }t        j                  | |||f«      S )Nz%s.xz%s.yz%s.z)r   Ú	call_sregr
   Úpack_struct)ÚbuilderÚprefixÚxÚyÚzs        úh/Volumes/fast/ai/experiments/voice-extract-mac/.venv/lib/python3.12/site-packages/numba/cuda/cudaimpl.pyÚinitialize_dim3r      s]   € Ü×Ñ˜G V¨f¡_Ó5€AÜ×Ñ˜G V¨f¡_Ó5€AÜ×Ñ˜G V¨f¡_Ó5€AÜ×Ñ˜w¨¨A¨q¨	Ó2Ð2ó    Ú	threadIdxc                 ó   — t        |d«      S )NÚtid©r   ©Úcontextr   ÚsigÚargss       r   Úcuda_threadIdxr(       s   € ä˜7 EÓ*Ð*r   ÚblockDimc                 ó   — t        |d«      S )NÚntidr#   r$   s       r   Úcuda_blockDimr,   %   s   € ä˜7 FÓ+Ð+r   ÚblockIdxc                 ó   — t        |d«      S )NÚctaidr#   r$   s       r   Úcuda_blockIdxr0   *   s   € ä˜7 GÓ,Ð,r   ÚgridDimc                 ó   — t        |d«      S )NÚnctaidr#   r$   s       r   Úcuda_gridDimr4   /   s   € ä˜7 HÓ-Ð-r   Úlaneidc                 ó.   — t        j                  |d«      S )Nr5   )r   r   r$   s       r   Úcuda_laneidr7   4   s   € ä×Ñ˜w¨Ó1Ð1r   r   c                 ó&   — |j                  |d«      S ©Nr   ©Úextract_valuer$   s       r   Údim3_xr<   9   ó   € à× Ñ   qÓ)Ð)r   r   c                 ó&   — |j                  |d«      S )Nr   r:   r$   s       r   Údim3_yr?   >   r=   r   r   c                 ó&   — |j                  |d«      S )Né   r:   r$   s       r   Údim3_zrB   C   r=   r   c                 ó   — |d   S r9   © r$   s       r   Úcuda_const_array_likerE   J   s   € ð �‰7€Nr   c                 ó@   — t         dz  a dj                  | t         «      S )zÍDue 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.
    r   z{0}_{1})Ú_unique_smem_idÚformat©Únames    r   Ú_get_unique_smem_idrK   T   s!   € ô �qÑ€OØ×Ñ˜D¤/Ó2Ð2r   c           	      ó¸   — |j                   d   j                  }t        |j                   d   «      }t        | ||f|t	        d«      t
        j                  d¬«      S ©Nr   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[   ^   sR   € à�X‰X�a‰[×&Ñ&€FÜ˜Ÿ™ ™Ó$€EÜ˜' 7°6°)À5Ü&9¸.Ó&IÜ$(×$9Ñ$9Ø'+ô-ð -r   c           	      óÚ   — |j                   d   D �cg c]  }|j                  ‘Œ }}t        |j                   d   «      }t        | |||t	        d«      t
        j                  d¬«      S c c}w rM   rU   ©r%   r   r&   r'   ÚsrP   rQ   s          r   Úcuda_shared_array_tupler_   h   sd   € ð (+§x¡x°¢{Ó4¡{ !ˆa�o‹o {€EÐ4Ü˜Ÿ™ ™Ó$€EÜ˜' 7°%¸uÜ&9¸.Ó&IÜ$(×$9Ñ$9Ø'+ô-ð -ùò 5s   ’A(c           	      ó¦   — |j                   d   j                  }t        |j                   d   «      }t        | ||f|dt        j
                  d¬«      S ©Nr   r   Ú_cudapy_lmemFrO   ©r'   rV   r   rW   r   ÚADDRSPACE_LOCALrY   s         r   Úcuda_local_array_integerre   s   sM   € à�X‰X�a‰[×&Ñ&€FÜ˜Ÿ™ ™Ó$€EÜ˜' 7°6°)À5Ø&4Ü$(×$8Ñ$8Ø',ô.ð .r   c           	      óÈ   — |j                   d   D �cg c]  }|j                  ‘Œ }}t        |j                   d   «      }t        | |||dt        j
                  d¬«      S c c}w ra   rc   r]   s          r   Úptx_lmem_alloc_arrayrg   }   s_   € ð (+§x¡x°¢{Ó4¡{ !ˆa�o‹o {€EÐ4Ü˜Ÿ™ ™Ó$€EÜ˜' 7°%¸uØ&4Ü$(×$8Ñ$8Ø',ô.ð .ùò 5s   ’Ac                 óè   — |rJ ‚d}|j                   }t        j                  t        j                  «       d«      }t	        j
                  |||«      }|j                  |d«       | j                  «       S )Nz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   ˆ   ó^   € á€Oˆ8Ø"€EØ�>‰>€DÜ�?‰?œ2Ÿ;™;›=¨"Ó-€DÜ×)Ñ)¨$°°eÓ<€DØ‡L�L��rÔØ×"Ñ"Ó$Ð$r   c                 óè   — |rJ ‚d}|j                   }t        j                  t        j                  «       d«      }t	        j
                  |||«      }|j                  |d«       | j                  «       S )Nzllvm.nvvm.membar.sysrD   ri   rp   s           r   Úptx_threadfence_systemrx   “   rv   r   c                 óè   — |rJ ‚d}|j                   }t        j                  t        j                  «       d«      }t	        j
                  |||«      }|j                  |d«       | j                  «       S )Nzllvm.nvvm.membar.glrD   ri   rp   s           r   Úptx_threadfence_devicerz   ž   s^   € á€Oˆ8Ø!€EØ�>‰>€DÜ�?‰?œ2Ÿ;™;›=¨"Ó-€DÜ×)Ñ)¨$°°eÓ<€DØ‡L�L��rÔØ×"Ñ"Ó$Ð$r   c                 ó¦   — | j                  t        j                  d«      }t        j                  t        j                  «      }t	        | |||g«      S )Nl   ÿÿ )Úget_constantr	   Úint32ÚnoneÚptx_syncwarp_mask)r%   r   r&   r'   ÚmaskÚmask_sigs         r   Úptx_syncwarpr‚   ©   s=   € à×Ñ¤§¡¨ZÓ8€DÜ�z‰zœ%Ÿ+™+Ó&€HÜ˜W g¨x¸$¸Ó@Ð@r   c                 ó  — d}|j                   }t        j                  t        j                  «       t        j                  d«      f«      }t        j                  |||«      }|j                  ||«       | j                  «       S )Nzllvm.nvvm.bar.warp.syncé    )	rj   r   rk   rl   ÚIntTyper
   rm   rn   ro   rp   s           r   r   r   °   sb   € à%€EØ�>‰>€DÜ�?‰?œ2Ÿ;™;›=¬2¯:©:°b«>Ð*;Ó<€DÜ×)Ñ)¨$°°eÓ<€DØ‡L�L��tÔØ×"Ñ"Ó$Ð$r   c           
      ó  — |\  }}}}}|j                   d   }	|	t        j                  v r/|j                  |t	        j
                  |	j                  «      «      }d}
|j                  }t	        j                  t	        j                  t	        j
                  d«      t	        j
                  d«      f«      t	        j
                  d«      t	        j
                  d«      t	        j
                  d«      t	        j
                  d«      t	        j
                  d«      f«      }t        j                  |||
«      }|	j                  dk(  rŒ|j                  ||||||f«      }|	t        j                  k(  r`|j                  |d«      }|j                  |d«      }|j                  |t	        j                  «       «      }t        j                   |||f«      }|S |j#                  |t	        j
                  d«      «      }|j%                  || j'                  t        j(                  d«      «      }|j#                  |t	        j
                  d«      «      }|j                  ||||||f«      }|j                  ||||||f«      }|j                  |d«      }|j                  |d«      }|j                  |d«      }|j+                  |t	        j
                  d«      «      }|j+                  |t	        j
                  d«      «      }|j-                  || j'                  t        j(                  d«      «      }|j/                  ||«      }|	t        j0                  k(  r$|j                  |t	        j2                  «       «      }t        j                   |||f«      }|S )a�  
    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.
    rA   zllvm.nvvm.shfl.sync.i32r„   r   r   é@   )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ˆ$��u˜eØ—‘˜!‘€JØ”U×&Ñ&Ñ&Ø—‘ ¤r§z¡z°*×2EÑ2EÓ'FÓGˆØ%€EØ�>‰>€DÜ�?‰?Ü
×ÑœbŸj™j¨›n¬b¯j©j¸«mÐ<Ó=ÜŸZ™Z¨›^¬R¯Z©Z¸«^¼R¿Z¹ZÈ»^ÜŸZ™Z¨›^¬R¯Z©Z¸«^ð=ó€Dô
 ×)Ñ)¨$°°eÓ<€DØ×Ñ˜bÒ Ø�l‰l˜4 $¨¨e°U¸EÐ!BÓCˆØœŸ™Ò&Ø×&Ñ& s¨AÓ.ˆBØ×(Ñ(¨¨aÓ0ˆDØ—‘ ¤R§\¡\£^Ó4ˆBÜ×/Ñ/°¸"¸d¸ÓDˆCð" €Jð —‘˜u¤b§j¡j°£nÓ5ˆØ—\‘\ %¨×)=Ñ)=¼e¿h¹hÈÓ)KÓLˆ
Ø—‘˜z¬2¯:©:°b«>Ó:ˆØ�|‰|˜D 4¨¨v°u¸eÐ"DÓEˆØ�|‰|˜D 4¨¨v°u¸eÐ"DÓEˆØ×#Ñ# D¨!Ó,ˆØ×#Ñ# D¨!Ó,ˆØ×$Ñ$ T¨1Ó-ˆØ—‘˜c¤2§:¡:¨b£>Ó2ˆØ—‘˜c¤2§:¡:¨b£>Ó2ˆØ—‘˜V W×%9Ñ%9¼%¿(¹(ÀBÓ%GÓHˆØ�[‰[˜ Ó(ˆØœŸ™Ò&Ø—‘ ¤R§]¡]£_Ó5ˆBÜ×+Ñ+¨G°b¸$°ZÓ@ˆØ€Jr   c                 óŠ  — d}|j                   }t        j                  t        j                  t        j                  d«      t        j                  d«      f«      t        j                  d«      t        j                  d«      t        j                  d«      f«      }t        j                  |||«      }|j                  ||«      S )Nzllvm.nvvm.vote.syncr„   r   )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äŸJ™J r›N¬B¯J©J°r«N¼B¿J¹JÀq»MÐJóL€Dô ×)Ñ)¨$°°eÓ<€DØ�<‰<˜˜dÓ#Ð#r   c                 óø  — |\  }}|j                   d   j                  }|j                   d   t        j                  v r%|j	                  |t        j                  |«      «      }dj                  |«      }|j                  }t        j                  t        j                  d«      t        j                  d«      t        j                  |«      f«      }	t        j                  ||	|«      }
|j                  |
||f«      S )Nr   zllvm.nvvm.match.any.sync.i{}r„   )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¿   € ð
 �K€Dˆ%Ø�H‰H�Q‰K× Ñ €EØ
‡x�x��{”e×'Ñ'Ñ'Ø—‘ ¤r§z¡z°%Ó'8Ó9ˆØ*×1Ñ1°%Ó8€EØ�>‰>€DÜ�?‰?œ2Ÿ:™: b›>¬B¯J©J°r«N¼B¿J¹JÀuÓ<MÐ+NÓO€DÜ×)Ñ)¨$°°eÓ<€DØ�<‰<˜˜t U˜mÓ,Ð,r   c                 óH  — |\  }}|j                   d   j                  }|j                   d   t        j                  v r%|j	                  |t        j                  |«      «      }dj                  |«      }|j                  }t        j                  t        j                  t        j                  d«      t        j                  d«      f«      t        j                  d«      t        j                  |«      f«      }	t        j                  ||	|«      }
|j                  |
||f«      S )Nr   zllvm.nvvm.match.all.sync.i{}r„   )r'   rŠ   r	   rˆ   r‰   r   r…   rH   rj   rk   r‹   r
   rm   rn   r¯   s              r   Úptx_match_all_syncr³     sä   € ð
 �K€Dˆ%Ø�H‰H�Q‰K× Ñ €EØ
‡x�x��{”e×'Ñ'Ñ'Ø—‘ ¤r§z¡z°%Ó'8Ó9ˆØ*×1Ñ1°%Ó8€EØ�>‰>€DÜ�?‰?œ2×/Ñ/´·±¸B³Ü13·±¸A³ð1@ó AäŸJ™J r›N¬B¯J©J°uÓ,=Ð>ó@€Dô ×)Ñ)¨$°°eÓ<€DØ�<‰<˜˜t U˜mÓ,Ð,r   c                 ó¦   — t        j                  t        j                  t        j                  d«      g «      ddd¬«      }|j	                  |g «      S )Nr„   zactivemask.b32 $0;ú=rT©Úside_effect©r   Ú	InlineAsmrk   r…   rn   ©r%   r   r&   r'   Ú
activemasks        r   Úptx_activemaskr¼     s>   € ä—‘œbŸo™o¬b¯j©j¸«n¸bÓAØ2°DÀdôL€Jà�<‰<˜
 BÓ'Ð'r   c                 ó¦   — t        j                  t        j                  t        j                  d«      g «      ddd¬«      }|j	                  |g «      S )Nr„   zmov.u32 $0, %lanemask_lt;rµ   Tr¶   r¸   rº   s        r   Úptx_lanemask_ltr¾   $  s@   € ä—‘œbŸo™o¬b¯j©j¸«n¸bÓAØ9¸4Ø*.ô0€Jð �<‰<˜
 BÓ'Ð'r   c                 ó*   — |j                  |d   «      S r9   )Úctpopr$   s       r   Úptx_popcrÁ   ,  s   € à�=‰=˜˜a™Ó!Ð!r   c                 ó    —  |j                   |Ž S ©N)Úfmar$   s       r   Úptx_fmarÅ   1  s   € àˆ7�;‰;˜ÐÐr   c                 óh   — dddœ}	 ||    S # t         $ r d| › d�}t        j                  |«      ‚w xY w)N)Úf32Úf)Úf64Úd)r„   r‡   z$Conversion between float16 and floatú unsupported©ÚKeyErrorr   ÚCudaLoweringError©rŠ   ÚtypemapÚmsgs      r   Úfloat16_float_ty_constraintrÒ   6  sJ   € Ø \Ñ2€Gð,Ø�xÑ Ð øÜò ,Ø4°X°J¸lÐKˆÜ×&Ñ& sÓ+Ð+ð,ús   ‡ Œ%1c                 ó>  — |j                   |j                   k(  r|S t        |j                   «      \  }}t        j                  | j	                  |«      t        j
                  d«      g«      }t        j                  |d|› d�d|› d�«      }|j                  ||g«      S )Né   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°·±Ó?�N€Bˆ
ä�?‰?˜7×1Ñ1°$Ó7¼"¿*¹*ÀR».Ð9IÓJ€DÜ
�,‰,�t˜t B 4 |Ð4¸¸*¸ÀRÐ6HÓ
I€CØ�<‰<˜˜c˜UÓ#Ð#r   c                 ó<  — |j                   |j                   k(  r|S t        |j                   «      \  }}t        j                  t        j                  d«      | j                  |«      g«      }t        j                  |d|› d�d|› �«      }|j                  ||g«      S )NrÔ   úcvt.rn.f16.ú $0, $1;ú=h,)rŠ   rÒ   r   rk   r…   rØ   r¹   rn   rÙ   s	            r   Úfloat_to_float16_castrå   L  s„   € à‡�˜$Ÿ-™-Ò'Øˆ
ä0°·±ÓA�N€Bˆ
ä�?‰?œ2Ÿ:™: b›>¨G×,BÑ,BÀ6Ó,JÐ+KÓL€DÜ
�,‰,�t˜{¨2¨$¨hÐ7¸3¸z¸lÐ9KÓ
L€CØ�<‰<˜˜c˜UÓ#Ð#r   c                 ól   — dddddœ}	 ||    S # t         $ r d| › d�}t        j                  |«      ‚w xY w)NÚcÚhÚrÚl)é   rÔ   r„   r‡   z"Conversion between float16 and intrË   rÌ   rÏ   s      r   Úfloat16_int_constraintrì   X  sN   € Ø˜C S¨cÑ3€Gð,Ø�xÑ Ð øÜò ,Ø2°8°*¸LÐIˆÜ×&Ñ& sÓ+Ð+ð,ús   ‰ Ž%3c                 ó*  — |j                   }t        |«      }|j                  rdnd}t        j                  | j                  |«      t        j                  d«      g«      }t        j                  |d|› |› d�d|› d�«      }	|j                  |	|g«      S )Nr^   ÚurÔ   z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Ü'¨Ó1€JØŸš‘¨€Jä�?‰?˜7×1Ñ1°$Ó7¼"¿*¹*ÀR».Ð9IÓJ€DÜ
�,‰,�tØ! * ¨h¨Z°|ÐDØ˜:˜, bÐ)ó+€Cð �<‰<˜˜c˜UÓ#Ð#r   c                 ó(  — |j                   }t        |«      }|j                  rdnd}t        j                  t        j
                  d«      | j                  |«      g«      }t        j                  |d|› |› d�d|› �«      }	|j                  |	|g«      S )Nr^   rî   rÔ   râ   rã   rä   )	rŠ   rì   rï   r   rk   r…   rØ   r¹   rn   rð   s
             r   Úinteger_to_float16_castrô   o  sŒ   € ð �‰€HÜ'¨Ó1€JØŸš‘¨3€Jä�?‰?œ2Ÿ:™: b›>Ø#×2Ñ2°6Ó:Ð;ó=€Dä
�,‰,�tØ$ Z L°°
¸(ÐCØ˜Z˜LÐ)ó+€Cð �<‰<˜˜c˜UÓ#Ð#r   c                 óh   ‡— t        | t        j                  t        j                  «      ˆfd„«       }y )Nc                 óþ   •— t        j                  t        j                  d«      t        j                  d«      t        j                  d«      g«      }t        j                  |‰› d�d«      }|j	                  ||«      S )NrÔ   z.f16 $0,$1,$2;ú=h,h,h©r   rk   r…   r¹   rn   )r%   r   r&   r'   rs   rß   Úops         €r   Úptx_fp16_binaryz*lower_fp16_binary.<locals>.ptx_fp16_binary  s^   ø€ ä�‰œrŸz™z¨"›~Ü "§
¡
¨2£´·
±
¸2³Ð?óAˆä�l‰l˜4 B 4 ~Ð!6¸ÓAˆØ�|‰|˜C Ó&Ð&r   ©Úlowerr	   Úfloat16)Úfnrù   rú   s    ` r   Úlower_fp16_binaryrÿ   ~  s&   ø€ Ü
ˆ2Œu�}‰}œeŸm™mÓ,ó'ó -ñ'r   ÚaddÚsubÚmulc                 óÎ   — t        j                  t        j                  d«      t        j                  d«      g«      }t        j                  |dd«      }|j	                  ||«      S )NrÔ   zneg.f16 $0, $1;ú=h,hrø   ©r%   r   r&   r'   rs   rß   s         r   Úptx_fp16_hnegr  ’  óI   € ä�?‰?œ2Ÿ:™: b›>¬B¯J©J°r«NÐ+;Ó<€DÜ
�,‰,�tÐ.°Ó
7€CØ�<‰<˜˜TÓ"Ð"r   c                 ó   — t        | |||«      S rÃ   )r  r$   s       r   Úoperator_hnegr	  ™  ó   € ä˜ '¨3°Ó5Ð5r   c                 óÎ   — t        j                  t        j                  d«      t        j                  d«      g«      }t        j                  |dd«      }|j	                  ||«      S )NrÔ   zabs.f16 $0, $1;r  rø   r  s         r   Úptx_fp16_habsr  ž  r  r   c                 ó   — t        | |||«      S rÃ   )r  r$   s       r   Úoperator_habsr  ¥  r
  r   c                 ó"  — t        j                  d«      t        j                  d«      t        j                  d«      g}t        j                  t        j                  d«      |«      }t        j                  |dd«      }|j	                  ||«      S )NrÔ   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  ª  sb   € ä�j‰j˜‹nœbŸj™j¨›n¬b¯j©j¸«nÐ=€FÜ�?‰?œ2Ÿ:™: b›>¨6Ó2€DÜ
�,‰,�tÐ6¸
Ó
C€CØ�<‰<˜˜TÓ"Ð"r   c                 ó0   — d„ }| j                  ||||«      S )Nc                 óB   — t         j                  j                  | |«      S rÃ   )r   Úfp16Úhdiv)r   r   s     r   Úfp16_divzfp16_div_impl.<locals>.fp16_divµ  s   € Ü�y‰y�~‰~˜a Ó#Ð#r   ©Úcompile_internal)r%   r   r&   r'   r  s        r   Úfp16_div_implr  ²  s    € ò$ð ×#Ñ# G¨X°s¸DÓ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                 ó   ‡ — ˆ fd„}|S )Nc                 óÐ  •— t        j                  t        j                  d«      t        j                  d«      t        j                  d«      g«      }t        j                  |t        j                  ‰	¬«      d«      }|j                  ||«      }| j                  t        j                  d«      }|j                  |t        j                  d«      «      }|j                  d||«      S )NrÔ   )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_comparisonz*_gen_fp16_cmp.<locals>.ptx_fp16_comparisonÃ  s£   ø€ Ü�‰œrŸz™z¨"›~´·
±
¸2³ÄÇ
Á
È2ÃÐ/OÓPˆÜ�l‰l˜4¤×!1Ñ!1°RÐ!1Ó!8¸(ÓCˆØ—‘˜c 4Ó(ˆà×#Ñ#¤E§K¡K°Ó3ˆØ—_‘_ V¬R¯Z©Z¸«^Ó<ˆ
Ø×$Ñ$ T¨:°tÓ<Ð<r   rD   )rù   r"  s   ` r   Ú_gen_fp16_cmpr#  Â  s   ø€ ô=ð Ðr   ÚeqÚneÚgeÚgtÚleÚltc                 óh   ‡— t        | t        j                  t        j                  «      ˆfd„«       }y )Nc                 ó^   •—  t        ‰«      | |||«      }|j                  ||d   |d   «      S )Nr   r   )r#  Úselect)r%   r   r&   r'   Úchoicerù   s        €r   Úptx_fp16_minmaxz*lower_fp16_minmax.<locals>.ptx_fp16_minmaxÝ  s5   ø€ à"”˜rÓ" 7¨G°S¸$Ó?ˆØ�~‰~˜f d¨1¡g¨t°A©wÓ7Ð7r   rû   )rþ   rq   rù   r.  s     ` r   Úlower_fp16_minmaxr/  Ü  s&   ø€ Ü
ˆ2Œu�}‰}œeŸm™mÓ,ó8ó -ñ8r   ÚmaxÚminÚ
__nv_cbrtfÚ	__nv_cbrtc                 óæ   — |j                   }t        |   }| j                  |«      }|j                  }t	        j
                  ||g«      }t        j                  |||«      }	|j                  |	|«      S rÃ   )	Úreturn_typeÚ
cbrt_funcsrØ   rj   r   rk   r
   rm   rn   )
r%   r   r&   r'   rÝ   rq   Úftyrr   rs   rþ   s
             r   Úptx_cbrtr8  ñ  sf   € ð 
�‰€BÜ�r‰N€EØ
×
 Ñ
  Ó
$€CØ�>‰>€DÜ�?‰?˜3  Ó&€DÜ	×	'Ñ	'¨¨d°EÓ	:€BØ�<‰<˜˜DÓ!Ð!r   c           	      óÞ   — t        j                  |j                  t        j                  t        j
                  d«      t        j
                  d«      f«      d«      }|j                  ||«      S )Nr„   Ú	__nv_brev©r
   rm   rj   r   rk   r…   rn   ©r%   r   r&   r'   rþ   s        r   Úptx_brev_u4r=  ý  sR   € ô
 
×	'Ñ	'Ø�‰Ü
�‰œŸ
™
 2›¬¯©°B«Ð(9Ó:Øó
€Bð �<‰<˜˜DÓ!Ð!r   c           	      óÞ   — t        j                  |j                  t        j                  t        j
                  d«      t        j
                  d«      f«      d«      }|j                  ||«      S )Nr‡   Ú__nv_brevllr;  r<  s        r   Úptx_brev_u8r@  	  sR   € ô
 
×	'Ñ	'Ø�‰Ü
�‰œŸ
™
 2›¬¯©°B«Ð(9Ó:Øó
€Bð �<‰<˜˜DÓ!Ð!r   c                 óh   — |j                  |d   | j                  t        j                  d«      «      S r9   )Úctlzr|   r	   Úbooleanr$   s       r   Úptx_clzrD    s.   € à�<‰<ØˆQ‰Ø×ÑœUŸ]™]¨AÓ.ó0ð 0r   c           	      óÞ   — t        j                  |j                  t        j                  t        j
                  d«      t        j
                  d«      f«      d«      }|j                  ||«      S )Nr„   Ú__nv_ffsr;  r<  s        r   Ú
ptx_ffs_32rG    sR   € ô 
×	'Ñ	'Ø�‰Ü
�‰œŸ
™
 2›¬¯©°B«Ð(9Ó:Øó
€Bð �<‰<˜˜DÓ!Ð!r   c           	      óÞ   — t        j                  |j                  t        j                  t        j
                  d«      t        j
                  d«      f«      d«      }|j                  ||«      S )Nr„   r‡   Ú
__nv_ffsllr;  r<  s        r   Ú
ptx_ffs_64rJ  &  sR   € ô 
×	'Ñ	'Ø�‰Ü
�‰œŸ
™
 2›¬¯©°B«Ð(9Ó:Øó
€Bð �<‰<˜˜DÓ!Ð!r   c                 ó4   — |\  }}}|j                  |||«      S rÃ   )r,  )r%   r   r&   r'   ÚtestÚaÚbs          r   Úptx_selprO  0  s    € à�J€Dˆ!ˆQØ�>‰>˜$  1Ó%Ð%r   c           	      ó   — t        j                  |j                  t        j                  t        j
                  «       t        j
                  «       t        j
                  «       f«      d«      }|j                  ||«      S )NÚ
__nv_fmaxf©r
   rm   rj   r   rk   r�   rn   r<  s        r   Ú
ptx_max_f4rS  6  óZ   € ä	×	'Ñ	'Ø�‰Ü
�‰Ü�L‰L‹NÜ�\‰\‹^œRŸ\™\›^Ð,ó	.ð 	ó
€Bð �<‰<˜˜DÓ!Ð!r   c           
      óÄ  — t        j                  |j                  t        j                  t        j
                  «       t        j
                  «       t        j
                  «       f«      d«      }|j                  || j                  ||d   |j                  d   t        j                  «      | j                  ||d   |j                  d   t        j                  «      g«      S )NÚ	__nv_fmaxr   r   ©r
   rm   rj   r   rk   r–   rn   Úcastr'   r	   Údoubler<  s        r   Ú
ptx_max_f8rZ  A  ó§   € ô 
×	'Ñ	'Ø�‰Ü
�‰Ü�M‰M‹OÜ�]‰]‹_œbŸm™m›oÐ.ó	0ð 	ó
€Bð �<‰<˜Ø�‰�W˜d 1™g s§x¡x°¡{´E·L±LÓAØ�‰�W˜d 1™g s§x¡x°¡{´E·L±LÓAðó ð r   c           	      ó   — t        j                  |j                  t        j                  t        j
                  «       t        j
                  «       t        j
                  «       f«      d«      }|j                  ||«      S )NÚ
__nv_fminfrR  r<  s        r   Ú
ptx_min_f4r^  R  rT  r   c           
      óÄ  — t        j                  |j                  t        j                  t        j
                  «       t        j
                  «       t        j
                  «       f«      d«      }|j                  || j                  ||d   |j                  d   t        j                  «      | j                  ||d   |j                  d   t        j                  «      g«      S )NÚ	__nv_fminr   r   rW  r<  s        r   Ú
ptx_min_f8ra  ]  r[  r   c           	      ó>  — t        j                  |j                  t        j                  t        j
                  d«      t        j                  «       f«      d«      }|j                  || j                  ||d   |j                  d   t        j                  «      g«      S )Nr‡   Ú__nv_llrintr   )r
   rm   rj   r   rk   r…   r–   rn   rX  r'   r	   rY  r<  s        r   Ú	ptx_roundrd  n  s   € ô 
×	'Ñ	'Ø�‰Ü
�‰Ü�J‰J�r‹NÜ�]‰]‹_Ðó	 ð 	ó
€Bð �<‰<˜Ø�‰�W˜d 1™g s§x¡x°¡{´E·L±LÓAðó ð r   c                 ó0   — d„ }| j                  ||||«      S )Nc                 ó€  — t        j                  | «      st        j                  | «      r| S |dk\  r6|dkD  rd|dz
  z  }d}nd|z  }d}| |z  |z  }t        j                  |«      r| S d| z  }| |z  }t        |«      }t        j                  ||z
  «      dk(  rdt        |dz  «      z  }|dk\  r
|z  |z  }|S ||z  }|S )Nr   é   g      $@g’ÕMÏð€Dg      ð?g      à?g       @)ÚmathÚisinfÚisnanÚroundÚfabs)r   ÚndigitsÚpow1Úpow2r   r   s         r   Úround_ndigitsz$round_to_impl.<locals>.round_ndigitsƒ  sÞ   € Ü�:‰:�aŒ=œDŸJ™J qœMØˆHà�aŠ<Ø˜Š|ð  ¨"¡Ñ-�Ø‘à˜w‘�Ø�Ø�T‘˜TÑ!ˆAÜ�z‰z˜!Œ}Ø�ð ˜W˜HÑ%ˆDØ�D‘ˆAä�!‹HˆÜ�I‰I�a˜!‘eÓ Ò#à”e˜A ™G“nÑ$ˆAà�aŠ<Ø�T‘˜TÑ!ˆAð ˆð �‰IˆAàˆr   r  )r%   r   r&   r'   rp  s        r   Úround_to_implrq  €  s!   € òðB ×#Ñ# G¨]¸CÀÓHÐHr   c                 ó   ‡ — ˆ fd„}|S )Nc                 ón   •— |j                   \  }| j                  |‰«      }|j                  ||d   «      S r9   )r'   r|   Úfmul)r%   r   r&   r'   ÚargtyÚfactorÚconsts         €r   Úimplzgen_deg_rad.<locals>.impl¨  s5   ø€ Ø—‘‰ˆØ×%Ñ% e¨UÓ3ˆØ�|‰|˜F D¨¡GÓ,Ð,r   rD   )rw  rx  s   ` r   Úgen_deg_radry  §  s   ø€ ô-ð €Kr   g     €f@c           
      óè  — |t         j                  v rt        j                  |d¬«      }|g}n!t        j                  ||t        |«      ¬«      }t        ||«      D ��cg c]'  \  }}| j                  |||t         j                  «      ‘Œ) }}}|j                  }	|	|k7  rt        d|	›d|›�«      ‚|j                  t        |«      k7  r#t        d|j                  t        |«      fz  «      ‚||fS c c}}w )z4
    Convert integer indices into tuple of intp
    r   )rQ   Úcount)r{  zexpect z	 but got z#indexing %d-D array with %d-D index)r	   Úinteger_domainÚUniTupler
   Úunpack_tupleÚlenÚziprX  ÚintprQ   Ú	TypeErrorÚndim)
r%   r   ÚindtyÚindsÚarytyÚvaltyÚindicesÚtÚirQ   s
             r   Ú_normalize_indicesr‹  ·  sç   € ð ”×$Ñ$Ñ$Ü—‘ U°!Ô4ˆØ�&‰ä×&Ñ& w°¼CÀ»JÔGˆä˜u gÔ.ô0Ù.‘4�1�að �|‰|˜G Q¨¬5¯:©:Õ6Ø.ð ñ 0ð �K‰K€EØ�‚~Ý²%¹Ð?Ó@Ð@à‡z�z”S˜“ZÒÜÐ=ØŸ™¤S¨£ZÐ0ñ1ó 2ð 	2ð �'ˆ>Ðùó0s   Á,C.c                 ó   ‡ — ˆ fd„}|S )Nc                 óî   •— |j                   \  }}}|\  }}}	|j                  }
t        | |||||«      \  }} | j                  |«      | ||«      }t	        j
                  | ||||d¬«      } ‰| ||
||	«      S )NT©Ú
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   Úimpz_atomic_dispatcher.<locals>.impÏ  s‘   ø€ à!Ÿh™hÑˆˆu�eØ‰ˆˆT�3Ø—‘ˆä+¨G°W¸eÀTØ,1°5ó:‰ˆˆwð )ˆw×!Ñ! %Ó(¨°'¸3Ó?ˆÜ×&Ñ& w°¸ÀÀgØ26ô8ˆñ ˜7 G¨U°C¸Ó=Ð=r   rD   )r•  r–  s   ` r   Ú_atomic_dispatcherr—  Î  s   ø€ ô>ð €Jr   c                 óB  — |t         j                  k(  r3|j                  }|j                  t	        j
                  |«      ||f«      S |t         j                  k(  r3|j                  }|j                  t	        j                  |«      ||f«      S |j                  d||d«      S )Nr   Ú	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ó(ð 	(ð ×!Ñ! %¨¨c°;Ó?Ð?r   c                 óB  — |t         j                  k(  r3|j                  }|j                  t	        j
                  |«      ||f«      S |t         j                  k(  r3|j                  }|j                  t	        j                  |«      ||f«      S |j                  d||d«      S )Nr  r™  )	r	   rŒ   rj   rn   r   Údeclare_atomic_sub_float32r•   Údeclare_atomic_sub_float64rœ  r�  s         r   Úptx_atomic_subr£  ñ  rŸ  r   c                 óâ   — |t         j                  j                  v rE|j                  }|j                  }t        t        d|› �«      }|j                   ||«      ||f«      S t        d|› d�«      ‚)NÚ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­    ój   € ð
 ”—‘×6Ñ6Ñ6Ø�^‰^ˆØ�~‰~ˆÜ”YÐ"8¸¸Ð =Ó>ˆØ�|‰|™B˜t›H s¨C jÓ1Ð1äÐ8¸¸¸vÐFÓGÐGr   c                 óâ   — |t         j                  j                  v rE|j                  }|j                  }t        t        d|› �«      }|j                   ||«      ||f«      S t        d|› d�«      ‚)NÚdeclare_atomic_dec_intzUnimplemented atomic dec with r¦  r§  r«  s           r   Úptx_atomic_decr±    r®  r   c                 óê   ‡— t         ˆfd„«       }t        j                  t        j                  t        j                  fD ]2  } t        | t        j                  |t        j                  «      |«       Œ4 y )Nc                 óˆ   •— |t         j                  j                  v r|j                  ‰||d«      S t	        d‰› d|› d�«      ‚)Nr™  zUnimplemented atomic z with r¦  ©r   r¨  Úinteger_numba_typesrœ  r‚  )r%   r   rQ   r”  rÜ   rù   s        €r   Úimpl_ptx_atomicz+ptx_atomic_bitwise.<locals>.impl_ptx_atomic  sG   ø€ à”T—]‘]×6Ñ6Ñ7Ø×%Ñ% b¨#¨s°KÓ@Ð@äÐ3°B°4°v¸e¸WÀFÐKÓLÐLr   )r—  r	   r�  r}  ÚTuplerü   ÚArrayÚAny)Ústubrù   r¶  rÝ   s    `  r   Úptx_atomic_bitwiser»    sS   ø€ ÜóMó ðMô �z‰zœ5Ÿ>™>¬5¯;©;Ó7ˆØ/Œˆd”E—K‘K ¤U§Y¡YÓ/°Õ@ñ 8r   ÚandÚorÚxorc                 ó€   — |t         j                  j                  v r|j                  d||d«      S t	        d|› d�«      ‚)NÚxchgr™  zUnimplemented atomic exch with r¦  r´  )r%   r   rQ   r”  rÜ   s        r   Úptx_atomic_exchrÁ  /  sB   € ð
 ”—‘×2Ñ2Ñ3Ø×!Ñ! &¨#¨s°KÓ@Ð@äÐ9¸%¸ÀÐGÓHÐHr   c                 óú  — |j                   }|t        j                  k(  r'|j                  t	        j
                  |«      ||f«      S |t        j                  k(  r'|j                  t	        j                  |«      ||f«      S |t        j                  t        j                  fv r|j                  d||d¬«      S |t        j                  t        j                  fv r|j                  d||d¬«      S t        d|z  «      ‚©Nr0  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—;‘;¤§¡Ð,Ñ	,Ø×!Ñ! %¨¨c¸KÐ!ÓHÐHØ	”5—<‘<¤§¡Ð.Ñ	.Ø×!Ñ! &¨#¨s¸[Ð!ÓIÐIäÐ@À5ÑHÓIÐIr   c                 óú  — |j                   }|t        j                  k(  r'|j                  t	        j
                  |«      ||f«      S |t        j                  k(  r'|j                  t	        j                  |«      ||f«      S |t        j                  t        j                  fv r|j                  d||d¬«      S |t        j                  t        j                  fv r|j                  d||d¬«      S t        d|z  «      ‚©Nr1  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                 óú  — |j                   }|t        j                  k(  r'|j                  t	        j
                  |«      ||f«      S |t        j                  k(  r'|j                  t	        j                  |«      ||f«      S |t        j                  t        j                  fv r|j                  d||d¬«      S |t        j                  t        j                  fv r|j                  d||d¬«      S t        d|z  «      ‚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ÑCÀDÓIØ  #˜Jó(ð 	(à	”%—-‘-Ò	Ø�|‰|œI×CÑCÀDÓIØ  #˜Jó(ð 	(à	”5—;‘;¤§¡Ð,Ñ	,Ø×!Ñ! %¨¨c¸KÐ!ÓHÐHØ	”5—<‘<¤§¡Ð.Ñ	.Ø×!Ñ! &¨#¨s¸[Ð!ÓIÐIäÐ@À5ÑHÓIÐIr   c                 óú  — |j                   }|t        j                  k(  r'|j                  t	        j
                  |«      ||f«      S |t        j                  k(  r'|j                  t	        j                  |«      ||f«      S |t        j                  t        j                  fv r|j                  d||d¬«      S |t        j                  t        j                  fv r|j                  d||d¬«      S t        d|z  «      ‚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                 ó
  — |j                  |j                  d   t        j                  |j                  d   |j                  d   «      }|d   | j	                  t        j                  d«      |d   |d   f}t        | |||«      S )Nr   r   rA   )r5  r'   r	   r�  r|   Úptx_atomic_casr$   s       r   Úptx_atomic_compare_and_swaprß  Š  sn   € à
�/‰/˜#Ÿ(™( 1™+¤u§z¡z°3·8±8¸A±;ÀÇÁÈÁÓ
L€CØ�‰G�W×)Ñ)¬%¯*©*°aÓ8¸$¸q¹'À4ÈÁ7ÐK€DÜ˜' 7¨C°Ó6Ð6r   c                 ó´  — |j                   \  }}}}|\  }}	}
}t        | |||	||«      \  }} | j                  |«      | ||«      }t        j                  | ||||d¬«      }|j
                  t        j                  j                  v r<|j                  }|j
                  j                  }t        j                  |||||
|«      S t        d|j
                  z  «      ‚)NTrŽ  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ˆ5�%˜ØÑ€Cˆˆs�Cä'¨°¸%ÀÀuØ(-ó/�N€Eˆ7ð %ˆ7×Ñ˜eÓ$ W¨g°sÓ;€DÜ
×
"Ñ
" 7¨G°U¸DÀ'Ø.2ô4€Cð ‡{�{”t—}‘}×8Ñ8Ñ9Ø�~‰~ˆØ—;‘;×'Ñ'ˆÜ×'Ñ'¨°°xÀÀcÈ3ÓOÐOäÐ@À5Ç;Á;ÑNÓOÐOr   c                 óÚ   — t        j                  t        j                  t        j                  «       t        j                  d«      g«      ddd¬«      }|d   }|j                  ||g«       y )Nr„   znanosleep.u32 $0;ré   Tr¶   r   )r   r¹   rk   rl   r…   rn   )r%   r   r&   r'   Ú	nanosleepÚnss         r   Úptx_nanosleeprç  ©  sO   € ä—‘œRŸ_™_¬R¯[©[«]¼R¿Z¹ZÈ»^Ð<LÓMØ0°#À4ôI€Ià	ˆa‰€BØ‡L�L�˜R˜DÕ!r   c           
      óv  — t        t        j                  |d«      }|dk  xr |xr t        |«      dk(  }|dk  r|st	        d«      ‚| j
                  |   }	t        |t        j                  t        j                  f«      xs/ t        |	t        j                  «      xs |t        j                  k(  }
|t        j                  vr|
st        d|z  «      ‚| j                  |«      }t!        j"                  ||«      }|t$        j&                  k(  rt)        j*                  |||¬«      }n¼|j,                  }t)        j.                  ||||«      }| j1                  |«      }d|dz
  j3                  «       z  |_        |rd|_        n)t!        j8                  |t         j:                  «      |_        |j?                  |t!        j@                  t!        jB                  d«      «      d«      }tE        jF                  t%        jH                  «       jJ                  «      }| j                  |«      }|jM                  |«      }|}g }tO        tQ        |«      «      D ]  \  }}|jS                  |«       ||z  }Œ tQ        |«      D �cg c]  }|‘Œ }}|D �cg c]"  }| jU                  t        jV                  |«      ‘Œ$ }}|r©t!        jX                  t!        jZ                  t!        jB                  d	«      g «      d
dd¬«      }|j]                  |j_                  |g «      t!        jB                  d«      «      }| jU                  t        jV                  |«      }|ja                  ||«      g}n-|D �cg c]"  }| jU                  t        jV                  |«      ‘Œ$ }}t        |«      }t        jb                  ||d¬«      } | je                  |«      | |«      } | jg                  | |ji                  || jj                  jl                  «      ||| jU                  t        jV                  |«      d ¬«       | jo                  «       S c c}w c c}w c c}w )Nr   r   zarray length <= 0zunsupported type: %srI   Úexternalrë   Úgenericr„   zmov.u32 $0, %dynamic_smem_size;rµ   Tr¶   r‡   ÚC)rQ   rƒ  Úlayout)Ú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—|‘| U¨AÓ.€Ið  ‘>ÒF lÒF´s¸5³zÀQ±€LØ�A‚~™lÜÐ,Ó-Ð-ð ×+Ñ+¨EÑ2€Jä�5œ5Ÿ<™<¬¯©Ð7Ó8ò 	"Ü�j¤&×"4Ñ"4Ó5ò	"à”E—M‘MÑ!ð ð
 ”E×'Ñ'Ñ'Ñ0DÜÐ.°Ñ6Ó7Ð7à×#Ñ# EÓ*€GÜ�\‰\˜' 9Ó-€Fà”D×(Ñ(Ò(ô ×%Ñ% g¨v¸KÔH‰à�~‰~ˆô ×+Ñ+¨D°&¸+Ø,5ó7ˆð ×&Ñ& wÓ/ˆð ˜E A™I×2Ñ2Ó4Ñ4ˆŒáØ&ˆE�Mô !#§¡¨F´B·L±LÓ AˆEÔð ×'Ñ'¨¬r¯~©~¼b¿j¹jÈ»mÓ/LØ(1ó3ˆô ×&Ñ&¤t§y¡y£{×'>Ñ'>Ó?€JØ×#Ñ# EÓ*€GØ×#Ñ# JÓ/€Hð €JØ€HÜ ¤¨%£Ö1‰ˆˆ8Ø�‰˜
Ô#Ø�hÑ‰
ð 2ô # 8Ô,Ó-Ñ,�QŠqÐ,€GÐ-Ù=DÓE¹W¸�×$Ñ$¤U§Z¡Z°Õ3¸W€HÐEñ ô
  Ÿ\™\¬"¯/©/¼"¿*¹*ÀR».È"Ó*MØ*KØ*.¸DôBÐð —|‘| G§L¡LÐ1CÀRÓ$HÜ$&§J¡J¨r£Nó4ˆð ×(Ñ(¬¯©°XÓ>ˆ	Ø—,‘,˜|¨YÓ7Ð8‰á?DÓE¹u¸!�'×&Ñ&¤u§z¡z°1Õ5¸uˆÐEô ˆu‹:€DÜ�K‰K˜e¨$°sÔ;€EØ
#ˆ'×
Ñ
˜UÓ
# G¨WÓ
5€Cà×Ñ˜3Ø '§¡°¸¿¹¿¹Ó GØ!'Ø#+Ø$+×$8Ñ$8¼¿¹ÀXÓ$NØ#'ð ô )ð �=‰=‹?ÐùòA .ùÚEùò$ Fs   É8	P,Ê'P1Í'P6c                 ó"   — | j                  «       S rÃ   )ro   )r%   r   rÝ   Úpyvals       r   Úcuda_dispatcher_constr%    s   € à×"Ñ"Ó$Ð$r   )F)ÒÚ	functoolsr   rñ  rh  Ú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,   r0   r4   r7   r<   r?   rB   rw  Ú
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_intrinsicrC  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)  r/  ÚhmaxÚhminrŒ   r•   r6  Úcbrtr8  ÚbrevÚu4r=  Úu8r@  ÚclzrD  ÚffsrG  rJ  ÚselprO  r0  rS  rZ  r1  r^  ra  rk  rd  rq  ry  Ú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>rx     sZ  ðÝ Û Û å Ý ç 4Ý 1Ý 'ß %Ý Ý ,Ý Ý ß /Ñ /ß 1á‹:€Ø�‰€Ø×#Ñ#€
Ø×(Ñ(€ò3ñ ˆLˆE�L‰L˜Ó Ó,ñ+ó -ð+ñ ˆLˆE�L‰L˜Ó 
Ó+ñ,ó ,ð,ñ ˆLˆE�L‰L˜Ó 
Ó+ñ-ó ,ð-ñ ˆLˆE�L‰L˜Ó 	Ó*ñ.ó +ð.ñ ˆLˆE�L‰L˜Ó Ó)ñ2ó *ð2ñ ˆD�#Óñ*ó ð*ñ ˆD�#Óñ*ó ð*ñ ˆD�#Óñ*ó ð*ñ €t‡z�z×Ñ˜eŸk™kÓ*ñó +ðð €ò3ñ €t‡{�{×Ñ˜%×.Ñ.°·	±	Ó:ñ-ó ;ð-ñ €t‡{�{×Ñ˜%Ÿ+™+ u§y¡yÓ1Ù€t‡{�{×Ñ˜%Ÿ.™.¨%¯)©)Ó4ñ-ó 5ó 2ð-ñ €t‡z�z×Ñ˜×-Ñ-¨u¯y©yÓ9ñ.ó :ð.ñ €t‡z�z×Ñ˜Ÿ™ e§i¡iÓ0Ù€t‡z�z×Ñ˜Ÿ™¨¯©Ó3ñ.ó 4ó 1ð.ñ €u×ÑÓñ%ó  ð%ñ €u×ÑÓ ñ%ó !ð%ñ €u×ÑÓñ%ó ð%ñ €u‡~�~ÓñAó ðAñ €u‡~�~�u—x‘xÓ ñ%ó !ð%ñ €u× Ñ  %§(¡(¨E¯H©H°e·h±hÀÇÁØ‡x�xóá€u× Ñ  %§(¡(¨E¯H©H°e·h±hÀÇÁØ‡x�xóá€u× Ñ  %§(¡(¨E¯H©H°e·h±hÀÇÁØ‡x�xóá€u× Ñ  %§(¡(¨E¯H©H°e·h±hÀÇÁØ‡x�xóñ+óóóóð+ñ\ €u× Ñ  %§(¡(¨E¯H©H°e·m±mÓDñ$ó Eð$ñ €u×Ñ˜UŸX™X u§x¡xÓ0Ù€u×Ñ˜UŸX™X u§x¡xÓ0Ù€u×Ñ˜UŸX™X u§x¡xÓ0Ù€u×Ñ˜UŸX™X u§x¡xÓ0ñ	-ó 1ó 1ó 1ó 1ð	-ñ €u×Ñ˜UŸX™X u§x¡xÓ0Ù€u×Ñ˜UŸX™X u§x¡xÓ0Ù€u×Ñ˜UŸX™X u§x¡xÓ0Ù€u×Ñ˜UŸX™X u§x¡xÓ0ñ-ó 1ó 1ó 1ó 1ð-ñ €u×ÑÓñ(ó ð(ñ €u×ÑÓñ(ó ð(ñ €u‡z�z�5—9‘9Óñ"ó ð"ñ €u‡y�y�%—)‘)˜UŸY™Y¨¯	©	Ó2ñó 3ðò,ñ ˆE�M‰M˜5Ÿ;™;Ó'ñ$ó (ð$ñ ˆE�K‰K˜Ÿ™Ó'ñ$ó (ð$ò,ñ ˆE�M‰M˜5Ÿ=™=Ó)ñ	$ó *ð	$ñ ˆE�M‰M˜5Ÿ=™=Ó)ÙˆE× Ñ  %§-¡-Ó0ñ
$ó 1ó *ð
$ò'ñ �%—*‘*—/‘/ 5Ô )Ù �(—,‘, Ô &Ù �(—-‘- Ô 'Ù �%—*‘*—/‘/ 5Ô )Ù �(—,‘, Ô &Ù �(—-‘- Ô 'Ù �%—*‘*—/‘/ 5Ô )Ù �(—,‘, Ô &Ù �(—-‘- Ô 'ñ €u‡z�z‡�˜Ÿ™Ó&ñ#ó 'ð#ñ €x‡|�|�U—]‘]Ó#ñ6ó $ð6ñ €u‡z�z‡�˜Ÿ™Ó&ñ#ó 'ð#ñ €sˆE�M‰MÓñ6ó ð6ñ €u‡z�z‡�˜Ÿ™ u§}¡}°e·m±mÓDñ#ó Eð#ñ €x×Ñ˜Ÿ™¨¯©Ó6Ù€x×Ñ˜%Ÿ-™-¨¯©Ó7ñBó 8ó 7ðBð€	ò	ð 4�€e‡j�j‡n‚n�e—m‘m U§]¡]Ó 3±MÀ$Ó4GÔ HØ 0�€h‡k‚k�5—=‘= %§-¡-Ó 0±¸tÓ1DÔ EØ 3�€e‡j�j‡n‚n�e—m‘m U§]¡]Ó 3±MÀ$Ó4GÔ HØ 0�€h‡k‚k�5—=‘= %§-¡-Ó 0±¸tÓ1DÔ EØ 3�€e‡j�j‡n‚n�e—m‘m U§]¡]Ó 3±MÀ$Ó4GÔ HØ 0�€h‡k‚k�5—=‘= %§-¡-Ó 0±¸tÓ1DÔ EØ 3�€e‡j�j‡n‚n�e—m‘m U§]¡]Ó 3±MÀ$Ó4GÔ HØ 0�€h‡k‚k�5—=‘= %§-¡-Ó 0±¸tÓ1DÔ EØ 3�€e‡j�j‡n‚n�e—m‘m U§]¡]Ó 3±MÀ$Ó4GÔ HØ 0�€h‡k‚k�5—=‘= %§-¡-Ó 0±¸tÓ1DÔ EØ 3�€e‡j�j‡n‚n�e—m‘m U§]¡]Ó 3±MÀ$Ó4GÔ HØ 0�€h‡k‚k�5—=‘= %§-¡-Ó 0±¸tÓ1DÔ Eò8ñ �%—*‘*—/’/ 5¨$Ô /Ù �%—*‘*—/’/ 5¨$Ô /ð 
‡M‚M�<Ø	‡M‚M�;ð€
ñ €u‡z‚z�5—=’=Ó!Ù€u‡z‚z�5—=’=Ó!ñ"ó "ó "ð"ñ €u‡z‚z�5—8’8Óñ"ó ð"ñ €u‡z‚z�5—8’8Óñ"ó ð"ñ €u‡y‚y�%—)‘)Óñ0ó ð0ñ €u‡y‚y�%—(‘(ÓÙ€u‡y‚y�%—(’(Óñ"ó ó ð"ñ €u‡y‚y�%—(‘(ÓÙ€u‡y‚y�%—(’(Óñ"ó ó ð"ñ €u‡z‚z�5—9‘9˜eŸi™i¨¯©Ó3ñ&ó 4ð&ñ
 €sˆE�H‰H�e—h‘hÓñ"ó  ð"ñ €sˆE�H‰H�e—h‘hÓÙ€sˆE�H‰H�e—h‘hÓÙ€sˆE�H‰H�e—h‘hÓñó  ó  ó  ðñ €sˆE�H‰H�e—h‘hÓñ"ó  ð"ñ €sˆE�H‰H�e—h‘hÓÙ€sˆE�H‰H�e—h‘hÓÙ€sˆE�H‰H�e—h‘hÓñó  ó  ó  ðñ €uˆe�h‰hÓÙ€uˆe�h‰hÓñ	ó ó ð	ñ  €uˆe�h‰h˜Ÿ™Ó&Ù€uˆe�h‰h˜Ÿ™Ó&ñ"Ió 'ó 'ð"IòJð �7Š7�T‰>€Ø�$—'’'‰>€Ø �€d‡l‚l�E—H‘HÓ ™k¨(Ó3Ô 4Ø �€d‡l‚l�E—H‘HÓ ™k¨(Ó3Ô 4Ø �€d‡l‚l�E—H‘HÓ ™k¨(Ó3Ô 4Ø �€d‡l‚l�E—H‘HÓ ™k¨(Ó3Ô 4òò.ñ$ €u‡|‚|×Ñ˜Ÿ™ e§j¢j°%·)±)Ó<Ù€u‡|‚|×Ñ˜Ÿ™ e§n¡n°e·i±iÓ@Ù€u‡|‚|×Ñ˜Ÿ™ e§k¡k°5·9±9Ó=Øñ
@ó ó >ó Aó =ð
@ñ €u‡|‚|×Ñ˜Ÿ™ e§j¢j°%·)±)Ó<Ù€u‡|‚|×Ñ˜Ÿ™ e§n¡n°e·i±iÓ@Ù€u‡|‚|×Ñ˜Ÿ™ e§k¡k°5·9±9Ó=Øñ
@ó ó >ó Aó =ð
@ñ €u‡|‚|×Ò˜Ÿ™ e§j¢j°%·)±)Ó<Ù€u‡|‚|×Ò˜Ÿ™ e§n¡n°e·i±iÓ@Ù€u‡|‚|×Ò˜Ÿ™ e§k¡k°5·9±9Ó=ØñHó ó >ó Aó =ðHñ €u‡|‚|×Ò˜Ÿ™ e§j¢j°%·)±)Ó<Ù€u‡|‚|×Ò˜Ÿ™ e§n¡n°e·i±iÓ@Ù€u‡|‚|×Ò˜Ÿ™ e§k¡k°5·9±9Ó=ØñHó ó >ó Aó =ðHò	Añ �5—<’<×$Ò$ eÔ ,Ù �5—<’<×#Ò# TÔ *Ù �5—<’<×#Ò# UÔ +ñ €u‡|‚|×Ò˜%Ÿ+™+ u§z¢z°5·9±9Ó=Ù€u‡|‚|×Ò˜%Ÿ+™+ u§~¡~°u·y±yÓAÙ€u‡|‚|×Ò˜%Ÿ+™+ u§{¡{°E·I±IÓ>ØñIó ó ?ó Bó >ðIñ €u‡|‚|×Ò˜Ÿ™ e§j¢j°%·)±)Ó<Ù€u‡|‚|×Ò˜Ÿ™ e§k¡k°5·9±9Ó=Ù€u‡|‚|×Ò˜Ÿ™ e§n¡n°e·i±iÓ@ØñJó ó Aó >ó =ðJñ  €u‡|‚|×Ò˜Ÿ™ e§j¢j°%·)±)Ó<Ù€u‡|‚|×Ò˜Ÿ™ e§k¡k°5·9±9Ó=Ù€u‡|‚|×Ò˜Ÿ™ e§n¡n°e·i±iÓ@ØñJó ó Aó >ó =ðJñ  €u‡|‚|×Ò˜EŸK™K¨¯ª°U·Y±YÓ?Ù€u‡|‚|×Ò˜EŸK™K¨¯©°e·i±iÓ@Ù€u‡|‚|×Ò˜EŸK™K¨¯©¸¿¹ÓCØñJó ó Dó Aó @ðJñ  €u‡|‚|×Ò˜EŸK™K¨¯ª°U·Y±YÓ?Ù€u‡|‚|×Ò˜EŸK™K¨¯©°e·i±iÓ@Ù€u‡|‚|×Ò˜EŸK™K¨¯©¸¿¹ÓCØñJó ó Dó Aó @ðJñ  €u‡|‚|×$Ò$ e§k¡k°5·9±9¸e¿i¹iÓHñ7ó Ið7ñ €u‡|‚|×Ò˜Ÿ™ e§j¢j°%·)±)¸U¿Y¹YÓGÙ€u‡|‚|×Ò˜Ÿ™ e§k¡k°5·9±9¸e¿i¹iÓHÙ€u‡|‚|×Ò˜Ÿ™ e§n¡n°e·i±iÀÇÁÓKñPó Ló Ió HðPñ* €u‡‚˜ŸšÓ%ñ"ó &ð"ð !&óañH �Óñ%ó  ð%ñ Ð#�×#Ò#Ó% uÕ -r   