§
    ƒÛgî˜  ã                   ó)  — 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!m"Z"  e
¦   «         Z#e#j$        Z$e#j%        Z&e#j'        Z'd„ Z( e& ej)        e¦  «        d¦  «        d„ ¦   «         Z* e& ej)        e¦  «        d¦  «        d„ ¦   «         Z+ e& ej)        e¦  «        d¦  «        d„ ¦   «         Z, e& ej)        e¦  «        d¦  «        d„ ¦   «         Z- e& ej)        e¦  «        d¦  «        d„ ¦   «         Z. e&e d¦  «        d„ ¦   «         Z/ e&e d¦  «        d„ ¦   «         Z0 e&e d¦  «        d„ ¦   «         Z1 e$ej2        j3        ¦  «        d „ ¦   «         Z4 e$d!e!¦  «        d"„ ¦   «         Z5 e$ej6        j7        ej8        ¦  «        d#„ ¦   «         Z9d a:d$„ Z; e$ej<        j=        ej>        ej?        ¦  «        d%„ ¦   «         Z@ e$ej<        j=        ejA        ej?        ¦  «         e$ej<        j=        ejB        ej?        ¦  «        d&„ ¦   «         ¦   «         ZC e$ejD        j=        ej>        ej?        ¦  «        d'„ ¦   «         ZE e$ejD        j=        ejA        ej?        ¦  «         e$ejD        j=        ejB        ej?        ¦  «        d(„ ¦   «         ¦   «         ZF e$ejG        ¦  «        d)„ ¦   «         ZH e$ejI        ¦  «        d*„ ¦   «         ZJ e$ejK        ¦  «        d+„ ¦   «         ZL e$ejM        ¦  «        d,„ ¦   «         ZN e$ejM        ejO        ¦  «        d-„ ¦   «         ZP e$ejQ        ejO        ejO        ejO        ejO        ejO        ¦  «         e$ejQ        ejO        ejO        ejR        ejO        ejO        ¦  «         e$ejQ        ejO        ejO        ejS        ejO        ejO        ¦  «         e$ejQ        ejO        ejO        ejT        ejO        ejO        ¦  «        d.„ ¦   «         ¦   «         ¦   «         ¦   «         ZU e$ejV        ejO        ejO        ejW        ¦  «        d/„ ¦   «         ZX e$ejY        ejO        ejO        ¦  «         e$ejY        ejO        ejR        ¦  «         e$ejY        ejO        ejS        ¦  «         e$ejY        ejO        ejT        ¦  «        d0„ ¦   «         ¦   «         ¦   «         ¦   «         ZZ e$ej[        ejO        ejO        ¦  «         e$ej[        ejO        ejR        ¦  «         e$ej[        ejO        ejS        ¦  «         e$ej[        ejO        ejT        ¦  «        d1„ ¦   «         ¦   «         ¦   «         ¦   «         Z\ e$ej]        ¦  «        d2„ ¦   «         Z^ e$ej_        ¦  «        d3„ ¦   «         Z` e$eja        ej?        ¦  «        d4„ ¦   «         Zb e$ejc        ej?        ej?        ej?        ¦  «        d5„ ¦   «         Zdd6„ Ze eejf        ejg        ¦  «        d7„ ¦   «         Zh eejg        ejf        ¦  «        d8„ ¦   «         Zid9„ Zj eejf        ejk        ¦  «        d:„ ¦   «         Zl eejk        ejf        ¦  «         eej>        ejf        ¦  «        d;„ ¦   «         ¦   «         Zmd<„ Zn enejo        jp        d=¦  «          enejq        d=¦  «          enejr        d=¦  «          enejo        js        d>¦  «          enejt        d>¦  «          eneju        d>¦  «          enejo        jv        d?¦  «          enejw        d?¦  «          enejx        d?¦  «          e$ejo        jy        ejf        ¦  «        d@„ ¦   «         Zz e$ej{        ejf        ¦  «        dA„ ¦   «         Z| e$ejo        j}        ejf        ¦  «        dB„ ¦   «         Z~ e$eejf        ¦  «        dC„ ¦   «         Z€ e$ejo        j�        ejf        ejf        ejf        ¦  «        dD„ ¦   «         Z‚ e$ejƒ        ejf        ejf        ¦  «         e$ej„        ejf        ejf        ¦  «        dE„ ¦   «         ¦   «         Z…dFZ†dG„ Z‡  e$ejo        jˆ        ejf        ejf        ¦  «         e‡dH¦  «        ¦  «           e$ej‰        ejf        ejf        ¦  «         e‡dH¦  «        ¦  «           e$ejo        jŠ        ejf        ejf        ¦  «         e‡dI¦  «        ¦  «           e$ej‹        ejf        ejf        ¦  «         e‡dI¦  «        ¦  «           e$ejo        jŒ        ejf        ejf        ¦  «         e‡dJ¦  «        ¦  «           e$ej�        ejf        ejf        ¦  «         e‡dJ¦  «        ¦  «           e$ejo        jŽ        ejf        ejf        ¦  «         e‡dK¦  «        ¦  «           e$ej�        ejf        ejf        ¦  «         e‡dK¦  «        ¦  «           e$ejo        j�        ejf        ejf        ¦  «         e‡dL¦  «        ¦  «           e$ej‘        ejf        ejf        ¦  «         e‡dL¦  «        ¦  «           e$ejo        j’        ejf        ejf        ¦  «         e‡dM¦  «        ¦  «           e$ej“        ejf        ejf        ¦  «         e‡dM¦  «        ¦  «         dN„ Z” e”ejo        j•        dOdK¦  «          e”ejo        j–        dPdM¦  «         ej—        dQej˜        dRiZ™ e$ejš        ej—        ¦  «         e$ejš        ej˜        ¦  «        dS„ ¦   «         ¦   «         Z› e$ejœ        ej�        ¦  «        dT„ ¦   «         Zž e$ejœ        ejŸ        ¦  «        dU„ ¦   «         Z  e$ej¡        ej?        ¦  «        dV„ ¦   «         Z¢ e$ej£        ejO        ¦  «         e$ej£        ej�        ¦  «        dW„ ¦   «         ¦   «         Z¤ e$ej£        ejR        ¦  «         e$ej£        ejŸ        ¦  «        dX„ ¦   «         ¦   «         Z¥ e$ej¦        ej?        ej?        ej?        ¦  «        dY„ ¦   «         Z§ e$e¨ejS        ejS        ¦  «        dZ„ ¦   «         Z© e$e¨ejT        ejS        ¦  «         e$e¨ejS        ejT        ¦  «         e$e¨ejT        ejT        ¦  «        d[„ ¦   «         ¦   «         ¦   «         Zª e$e«ejS        ejS        ¦  «        d\„ ¦   «         Z¬ e$e«ejT        ejS        ¦  «         e$e«ejS        ejT        ¦  «         e$e«ejT        ejT        ¦  «        d]„ ¦   «         ¦   «         ¦   «         Z­ e$e®ejS        ¦  «         e$e®ejT        ¦  «        d^„ ¦   «         ¦   «         Z¯ e$e®ejS        ejk        ¦  «         e$e®ejT        ejk        ¦  «        d_„ ¦   «         ¦   «         Z°d`„ Z±ej²        daz  Z³daej²        z  Z´  e$ejµ        ejS        ¦  «         e±e³¦  «        ¦  «           e$ejµ        ejT        ¦  «         e±e³¦  «        ¦  «           e$ej¶        ejS        ¦  «         e±e´¦  «        ¦  «           e$ej¶        ejT        ¦  «         e±e´¦  «        ¦  «         db„ Z·dc„ Z¸ e$ej¹        jq        ej8        ejº        ej?        ¦  «         e$ej¹        jq        ej8        ejB        ej?        ¦  «         e$ej¹        jq        ej8        ejA        ej?        ¦  «        e¸dd„ ¦   «         ¦   «         ¦   «         ¦   «         Z» e$ej¹        jt        ej8        ejº        ej?        ¦  «         e$ej¹        jt        ej8        ejB        ej?        ¦  «         e$ej¹        jt        ej8        ejA        ej?        ¦  «        e¸de„ ¦   «         ¦   «         ¦   «         ¦   «         Z¼ e$ej¹        j½        ej8        ejº        ej?        ¦  «         e$ej¹        j½        ej8        ejB        ej?        ¦  «         e$ej¹        j½        ej8        ejA        ej?        ¦  «        e¸df„ ¦   «         ¦   «         ¦   «         ¦   «         Z¾ e$ej¹        j¿        ej8        ejº        ej?        ¦  «         e$ej¹        j¿        ej8        ejB        ej?        ¦  «         e$ej¹        j¿        ej8        ejA        ej?        ¦  «        e¸dg„ ¦   «         ¦   «         ¦   «         ¦   «         ZÀdh„ ZÁ eÁej¹        jÂ        di¦  «          eÁej¹        jÃ        dj¦  «          eÁej¹        jÄ        dk¦  «          e$ej¹        jÅ        ej8        ejº        ej?        ¦  «         e$ej¹        jÅ        ej8        ejB        ej?        ¦  «         e$ej¹        jÅ        ej8        ejA        ej?        ¦  «        e¸dl„ ¦   «         ¦   «         ¦   «         ¦   «         ZÆ e$ej¹        j¨        ej8        ejº        ej?        ¦  «         e$ej¹        j¨        ej8        ejA        ej?        ¦  «         e$ej¹        j¨        ej8        ejB        ej?        ¦  «        e¸dm„ ¦   «         ¦   «         ¦   «         ¦   «         ZÇ e$ej¹        j«        ej8        ejº        ej?        ¦  «         e$ej¹        j«        ej8        ejA        ej?        ¦  «         e$ej¹        j«        ej8        ejB        ej?        ¦  «        e¸dn„ ¦   «         ¦   «         ¦   «         ¦   «         ZÈ e$ej¹        jÉ        ej8        ejº        ej?        ¦  «         e$ej¹        jÉ        ej8        ejA        ej?        ¦  «         e$ej¹        jÉ        ej8        ejB        ej?        ¦  «        e¸do„ ¦   «         ¦   «         ¦   «         ¦   «         ZÊ e$ej¹        jË        ej8        ejº        ej?        ¦  «         e$ej¹        jË        ej8        ejA        ej?        ¦  «         e$ej¹        jË        ej8        ejB        ej?        ¦  «        e¸dp„ ¦   «         ¦   «         ¦   «         ¦   «         ZÌ e$ej¹        jÍ        ej8        ej?        ej?        ¦  «        dq„ ¦   «         ZÎ e$ej¹        jÏ        ej8        ejº        ej?        ej?        ¦  «         e$ej¹        jÏ        ej8        ejA        ej?        ej?        ¦  «         e$ej¹        jÏ        ej8        ejB        ej?        ej?        ¦  «        dr„ ¦   «         ¦   «         ¦   «         ZÐ e$ejÑ        ejÒ        ¦  «        ds„ ¦   «         ZÓ	 dwdu„ZÔ e'e"¦  «        dv„ ¦   «         ZÕ e ejÖ        ¦   «         e$¦  «         dS )xé    )ÚreduceN)Úir)ÚRegistryÚ
lower_cast)Úparse_dtype)Úmodels)ÚtypesÚcgutils)Úufunc_db)Úregister_ufuncsé   )Únvvm)Úcuda)Ú	nvvmutilsÚstubsÚerrors)Údim3Ú
grid_groupÚ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        úI/granary/edev/quantum/lib/python3.11/site-packages/numba/cuda/cudaimpl.pyÚinitialize_dim3r      s]   € ÝÔ˜G V¨f¡_Ñ5Ô5€AÝÔ˜G V¨f¡_Ñ5Ô5€AÝÔ˜G V¨f¡_Ñ5Ô5€AÝÔ˜w¨¨A¨q¨	Ñ2Ô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_blockIdxr1   *   s   € å˜7 GÑ,Ô,Ð,r    ÚgridDimc                 ó"   — t          |d¦  «        S )NÚnctaidr$   r%   s       r   Úcuda_gridDimr5   /   s   € å˜7 HÑ-Ô-Ð-r    Úlaneidc                 ó,   — t          j        |d¦  «        S )Nr6   )r   r   r%   s       r   Úcuda_laneidr8   4   s   € åÔ˜w¨Ñ1Ô1Ð1r    r   c                 ó.   — |                      |d¦  «        S ©Nr   ©Úextract_valuer%   s       r   Údim3_xr=   9   ó   € à× Ò   qÑ)Ô)Ð)r    r   c                 ó.   — |                      |d¦  «        S ©Nr   r;   r%   s       r   Údim3_yrA   >   r>   r    r   c                 ó.   — |                      |d¦  «        S )Né   r;   r%   s       r   Údim3_zrD   C   r>   r    c                 ó¢   — |                       t          j        d¦  «        }|j        }|                     t          j        |¦  «        |f¦  «        S r@   )Úget_constantr	   Úint32ÚmoduleÚcallr   Ú declare_cudaCGGetIntrinsicHandle)r&   r   r'   r(   ÚoneÚlmods         r   Úcg_this_gridrM   H   sI   € à
×
Ò
�uœ{¨AÑ
.Ô
.€CØŒ>€DØ�<Š<ÝÔ2°4Ñ8Ô8Ø	ˆñô ð r    zGridGroup.syncc                 óª   — |                       t          j        d¦  «        }|j        }|                     t          j        |¦  «        g |¢|‘R ¦  «        S r:   )rF   r	   rG   rH   rI   r   Údeclare_cudaCGSynchronize)r&   r   r'   r(   ÚflagsrL   s         r   Úptx_sync_grouprQ   Q   sQ   € à× Ò ¥¤¨aÑ0Ô0€EØŒ>€DØ�<Š<ÝÔ+¨DÑ1Ô1Øˆ$ˆ�ˆˆñô ð r    c                 ó   — |d         S r:   © r%   s       r   Úcuda_const_array_likerT   \   s   € ð �Œ7€Nr    c                 óL   — t           dz  a d                     | 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_idrZ   f   s$   € õ �qÑ€OØ×Ò˜D¥/Ñ2Ô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_arrayrZ   r   ÚADDRSPACE_SHARED©r&   r   r'   r(   Úlengthr_   s         r   Úcuda_shared_array_integerrh   p   sX   € àŒX�aŒ[Ô&€FÝ˜œ œÑ$Ô$€EÝ˜' 7°6°)À5Ý&9¸.Ñ&IÔ&IÝ$(Ô$9Ø'+ð-ñ -ô -ð -r    c           	      ó¾   — d„ |j         d         D ¦   «         }t          |j         d         ¦  «        }t          | |||t          d¦  «        t          j        d¬¦  «        S )Nc                 ó   — g | ]	}|j         ‘Œ
S rS   ©rc   ©Ú.0Úss     r   ú
<listcomp>z+cuda_shared_array_tuple.<locals>.<listcomp>}   ó   € Ð4Ð4Ð4 !ˆaŒoÐ4Ð4Ð4r    r   r   r\   Tr]   )r(   r   rd   rZ   r   re   ©r&   r   r'   r(   r^   r_   s         r   Úcuda_shared_array_tuplerr   z   sc   € ð 5Ð4 s¤x°¤{Ð4Ñ4Ô4€EÝ˜œ œÑ$Ô$€EÝ˜' 7°%¸uÝ&9¸.Ñ&IÔ&IÝ$(Ô$9Ø'+ð-ñ -ô -ð -r    c           	      óœ   — |j         d         j        }t          |j         d         ¦  «        }t          | ||f|dt          j        d¬¦  «        S )Nr   r   Ú_cudapy_lmemFr]   )r(   rc   r   rd   r   ÚADDRSPACE_LOCALrf   s         r   Úcuda_local_array_integerrv   …   sP   € àŒX�aŒ[Ô&€FÝ˜œ œÑ$Ô$€EÝ˜' 7°6°)À5Ø&4Ý$(Ô$8Ø',ð.ñ .ô .ð .r    c           	      ó¤   — d„ |j         d         D ¦   «         }t          |j         d         ¦  «        }t          | |||dt          j        d¬¦  «        S )Nc                 ó   — g | ]	}|j         ‘Œ
S rS   rk   rl   s     r   ro   z(ptx_lmem_alloc_array.<locals>.<listcomp>’   rp   r    r   r   rt   Fr]   )r(   r   rd   r   ru   rq   s         r   Úptx_lmem_alloc_arrayry   �   s[   € ð 5Ð4 s¤x°¤{Ð4Ñ4Ô4€EÝ˜œ œÑ$Ô$€EÝ˜' 7°%¸uØ&4Ý$(Ô$8Ø',ð.ñ .ô .ð .r    c                 óè   — |rJ ‚d}|j         }t          j        t          j        ¦   «         d¦  «        }t	          j        |||¦  «        }|                     |d¦  «         |                      ¦   «         S )Nzllvm.nvvm.membar.ctarS   ©rH   r   ÚFunctionTypeÚVoidTyper
   Úget_or_insert_functionrI   Úget_dummy_value©r&   r   r'   r(   ÚfnamerL   ÚfntyÚsyncs           r   Úptx_threadfence_blockr„   š   óh   € à€O€O€OØ"€EØŒ>€DÝŒ?�2œ;™=œ=¨"Ñ-Ô-€DÝÔ)¨$°°eÑ<Ô<€DØ‡L‚L��rÑÔÐØ×"Ò"Ñ$Ô$Ð$r    c                 óè   — |rJ ‚d}|j         }t          j        t          j        ¦   «         d¦  «        }t	          j        |||¦  «        }|                     |d¦  «         |                      ¦   «         S )Nzllvm.nvvm.membar.sysrS   r{   r€   s           r   Úptx_threadfence_systemr‡   ¥   r…   r    c                 óè   — |rJ ‚d}|j         }t          j        t          j        ¦   «         d¦  «        }t	          j        |||¦  «        }|                     |d¦  «         |                      ¦   «         S )Nzllvm.nvvm.membar.glrS   r{   r€   s           r   Úptx_threadfence_devicer‰   °   sh   € à€O€O€OØ!€EØŒ>€DÝŒ?�2œ;™=œ=¨"Ñ-Ô-€DÝÔ)¨$°°eÑ<Ô<€DØ‡L‚L��rÑÔÐØ×"Ò"Ñ$Ô$Ð$r    c                 ó¤   — |                       t          j        d¦  «        }t          j        t          j        ¦  «        }t	          | |||g¦  «        S )Nl   ÿÿ )rF   r	   rG   ÚnoneÚptx_syncwarp_mask)r&   r   r'   r(   ÚmaskÚmask_sigs         r   Úptx_syncwarpr�   »   s@   € à×Ò¥¤¨ZÑ8Ô8€DÝŒz�%œ+Ñ&Ô&€HÝ˜W g¨x¸$¸Ñ@Ô@Ð@r    c                 ó  — d}|j         }t          j        t          j        ¦   «         t          j        d¦  «        f¦  «        }t          j        |||¦  «        }|                     ||¦  «         |                      ¦   «         S )Nzllvm.nvvm.bar.warp.syncé    )	rH   r   r|   r}   ÚIntTyper
   r~   rI   r   r€   s           r   rŒ   rŒ   Â   sj   € à%€EØŒ>€DÝŒ?�2œ;™=œ=­2¬:°b©>¬>Ð*;Ñ<Ô<€DÝÔ)¨$°°eÑ<Ô<€DØ‡L‚L��tÑÔÐØ×"Ò"Ñ$Ô$Ð$r    c           
      óH  — |\  }}}}}|j         d         }	|	t          j        v r-|                     |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—|                     ||||||f¦  «        }|	t          j        k    rj|                     |d¦  «        }|                     |d¦  «        }|                     |t	          j        ¦   «         ¦  «        }t          j        |||f¦  «        }�nä|                     |t	          j        d¦  «        ¦  «        }|                     ||                      t          j        d¦  «        ¦  «        }|                     |t	          j        d¦  «        ¦  «        }|                     ||||||f¦  «        }|                     ||||||f¦  «        }|                     |d¦  «        }|                     |d¦  «        }|                     |d¦  «        }|                     |t	          j        d¦  «        ¦  «        }|                     |t	          j        d¦  «        ¦  «        }|                     ||                      t          j        d¦  «        ¦  «        }|                     ||¦  «        }|	t          j        k    r'|                     |t	          j        ¦   «         ¦  «        }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.
    rC   zllvm.nvvm.shfl.sync.i32r‘   r   r   é@   )r(   r	   Úreal_domainÚbitcastr   r’   ÚbitwidthrH   r|   ÚLiteralStructTyper
   r~   rI   Úfloat32r<   Ú	FloatTypeÚmake_anonymous_structÚtruncÚlshrrF   Úi8ÚzextÚshlÚor_Úfloat64Ú
DoubleType)r&   r   r'   r(   r�   ÚmodeÚvalueÚindexÚclampÚ
value_typer�   rL   r‚   Ú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°*Ô2EÑ'FÔ'FÑGÔGˆØ%€EØŒ>€DÝŒ?Ý
Ô�bœj¨™nœn­b¬j¸©m¬mÐ<Ñ=Ô=ÝœZ¨™^œ^­R¬Z¸©^¬^½R¼ZÈ¹^¼^ÝœZ¨™^œ^­R¬Z¸©^¬^ð=ñô €Dõ
 Ô)¨$°°eÑ<Ô<€DØÔ˜bÒ Ð Ø�lŠl˜4 $¨¨e°U¸EÐ!BÑCÔCˆØ�œÒ&Ð&Ø×&Ò& s¨AÑ.Ô.ˆBØ×(Ò(¨¨aÑ0Ô0ˆDØ—’ ¥R¤\¡^¤^Ñ4Ô4ˆBÝÔ/°¸"¸d¸ÑDÔDˆCùà—’˜u¥b¤j°¡n¤nÑ5Ô5ˆØ—\’\ %¨×)=Ò)=½e¼hÈÑ)KÔ)KÑLÔLˆ
Ø—’˜z­2¬:°b©>¬>Ñ:Ô:ˆØ�|Š|˜D 4¨¨v°u¸eÐ"DÑEÔEˆØ�|Š|˜D 4¨¨v°u¸eÐ"DÑEÔEˆØ×#Ò# D¨!Ñ,Ô,ˆØ×#Ò# D¨!Ñ,Ô,ˆØ×$Ò$ T¨1Ñ-Ô-ˆØ—’˜c¥2¤:¨b¡>¤>Ñ2Ô2ˆØ—’˜c¥2¤:¨b¡>¤>Ñ2Ô2ˆØ—’˜V W×%9Ò%9½%¼(ÀBÑ%GÔ%GÑHÔHˆØ�[Š[˜ Ñ(Ô(ˆØ�œÒ&Ð&Ø—’ ¥R¤]¡_¤_Ñ5Ô5ˆBÝÔ+¨G°b¸$°ZÑ@Ô@ˆØ€Jr    c                 óx  — 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        |||¦  «        }|                     ||¦  «        S )Nzllvm.nvvm.vote.syncr‘   r   )rH   r   r|   r˜   r’   r
   r~   rI   )r&   r   r'   r(   r�   rL   r‚   r©   s           r   Úptx_vote_syncrº     s™   € à!€EØŒ>€DÝŒ?�2Ô/µ´¸B±´Ý13´¸A±´ð1@ñ Aô AåœJ r™NœN­B¬J°r©N¬N½B¼JÀq¹M¼MÐJñLô L€Dõ Ô)¨$°°eÑ<Ô<€DØ�<Š<˜˜dÑ#Ô#Ð#r    c                 óà  — |\  }}|j         d         j        }|j         d         t          j        v r(|                     |t          j        |¦  «        ¦  «        }d                     |¦  «        }|j        }t          j	        t          j        d¦  «        t          j        d¦  «        t          j        |¦  «        f¦  «        }	t          j        ||	|¦  «        }
|                     |
||f¦  «        S )Nr   zllvm.nvvm.match.any.sync.i{}r‘   )r(   r—   r	   r•   r–   r   r’   rW   rH   r|   r
   r~   rI   ©r&   r   r'   r(   r�   r¥   Úwidthr�   rL   r‚   r©   s              r   Úptx_match_any_syncr¾     sÂ   € ð
 �K€Dˆ%ØŒH�QŒKÔ €EØ
„x�„{•eÔ'Ð'Ð'Ø—’ ¥r¤z°%Ñ'8Ô'8Ñ9Ô9ˆØ*×1Ò1°%Ñ8Ô8€EØŒ>€DÝŒ?�2œ: b™>œ>­B¬J°r©N¬N½B¼JÀuÑ<MÔ<MÐ+NÑOÔO€DÝÔ)¨$°°eÑ<Ô<€DØ�<Š<˜˜t U˜mÑ,Ô,Ð,r    c                 ó,  — |\  }}|j         d         j        }|j         d         t          j        v r(|                     |t          j        |¦  «        ¦  «        }d                     |¦  «        }|j        }t          j	        t          j
        t          j        d¦  «        t          j        d¦  «        f¦  «        t          j        d¦  «        t          j        |¦  «        f¦  «        }	t          j        ||	|¦  «        }
|                     |
||f¦  «        S )Nr   zllvm.nvvm.match.all.sync.i{}r‘   )r(   r—   r	   r•   r–   r   r’   rW   rH   r|   r˜   r
   r~   rI   r¼   s              r   Úptx_match_all_syncrÀ     sí   € ð
 �K€Dˆ%ØŒH�QŒKÔ €EØ
„x�„{•eÔ'Ð'Ð'Ø—’ ¥r¤z°%Ñ'8Ô'8Ñ9Ô9ˆØ*×1Ò1°%Ñ8Ô8€EØŒ>€DÝŒ?�2Ô/µ´¸B±´Ý13´¸A±´ð1@ñ Aô AåœJ r™NœN­B¬J°uÑ,=Ô,=Ð>ñ@ô @€Dõ Ô)¨$°°eÑ<Ô<€DØ�<Š<˜˜t U˜mÑ,Ô,Ð,r    c                 ó¨   — t          j        t          j        t          j        d¦  «        g ¦  «        ddd¬¦  «        }|                     |g ¦  «        S )Nr‘   zactivemask.b32 $0;ú=rT©Úside_effect©r   Ú	InlineAsmr|   r’   rI   ©r&   r   r'   r(   Ú
activemasks        r   Úptx_activemaskrÉ   /  sL   € å”�bœo­b¬j¸©n¬n¸bÑAÔAØ2°DÀdðLñ Lô L€Jà�<Š<˜
 BÑ'Ô'Ð'r    c                 ó¨   — t          j        t          j        t          j        d¦  «        g ¦  «        ddd¬¦  «        }|                     |g ¦  «        S )Nr‘   zmov.u32 $0, %lanemask_lt;rÂ   TrÃ   rÅ   rÇ   s        r   Úptx_lanemask_ltrË   6  sL   € å”�bœo­b¬j¸©n¬n¸bÑAÔAØ9¸4Ø*.ð0ñ 0ô 0€Jð �<Š<˜
 BÑ'Ô'Ð'r    c                 ó8   — |                      |d         ¦  «        S r:   )Úctpopr%   s       r   Úptx_popcrÎ   >  s   € à�=Š=˜˜aœÑ!Ô!Ð!r    c                 ó   —  |j         |Ž S ©N)Úfmar%   s       r   Úptx_fmarÒ   C  s   € àˆ7Œ;˜ÐÐr    c                 ón   — 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ß   H  s[   € Ø \Ð2Ð2€Gð,Ø�xÔ Ð øÝð ,ð ,ð ,ØK°XÐKÐKÐKˆÝÔ& sÑ+Ô+Ð+ð,øøøs   ‡ �%4c                 ó4  — |j         |j         k    r|S t          |j         ¦  «        \  }}t          j        |                      |¦  «        t          j        d¦  «        g¦  «        }t          j        |d|› d�d|› d�¦  «        }|                     ||g¦  «        S )Né   zcvt.ú.f16 $0, $1;ú=ú,h)r—   rß   r   r|   Úget_value_typer’   rÆ   rI   ©	r&   r   ÚfromtyÚtotyÚvalÚtyÚ
constraintr‚   Úasms	            r   Úfloat16_to_float_castrí   R  s”   € à„˜$œ-Ò'Ð'Øˆ
å0°´Ñ?Ô?�N€Bˆ
åŒ?˜7×1Ò1°$Ñ7Ô7½"¼*ÀR¹.¼.Ð9IÑJÔJ€DÝ
Œ,�tÐ4 BÐ4Ð4Ð4Ð6H¸*Ð6HÐ6HÐ6HÑ
IÔ
I€CØ�<Š<˜˜c˜UÑ#Ô#Ð#r    c                 ó2  — |j         |j         k    r|S t          |j         ¦  «        \  }}t          j        t          j        d¦  «        |                      |¦  «        g¦  «        }t          j        |d|› d�d|› �¦  «        }|                     ||g¦  «        S )Nrá   úcvt.rn.f16.ú $0, $1;ú=h,)r—   rß   r   r|   r’   rå   rÆ   rI   ræ   s	            r   Úfloat_to_float16_castrò   ^  s‘   € à„˜$œ-Ò'Ð'Øˆ
å0°´ÑAÔA�N€Bˆ
åŒ?�2œ: b™>œ>¨G×,BÒ,BÀ6Ñ,JÔ,JÐ+KÑLÔL€DÝ
Œ,�tÐ7¨2Ð7Ð7Ð7Ð9K¸zÐ9KÐ9KÑ
LÔ
L€CØ�<Š<˜˜c˜UÑ#Ô#Ð#r    c                 ór   — 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ù   j  s_   € Ø˜C S¨cÐ3Ð3€Gð,Ø�xÔ Ð øÝð ,ð ,ð ,ØI°8ÐIÐIÐIˆÝÔ& sÑ+Ô+Ð+ð,øøøs   ‰ ‘%6c                 ó(  — |j         }t          |¦  «        }|j        rdnd}t          j        |                      |¦  «        t          j        d¦  «        g¦  «        }t          j        |d|› |› d�d|› d�¦  «        }	|                     |	|g¦  «        S )Nrn   Úurá   zcvt.rni.râ   rã   rä   )	r—   rù   Úsignedr   r|   rå   r’   rÆ   rI   ©
r&   r   rç   rè   ré   r—   rë   Ú
signednessr‚   rì   s
             r   Úfloat16_to_integer_castrÿ   t  sœ   € àŒ}€HÝ'¨Ñ1Ô1€JØœÐ,��¨€JåŒ?˜7×1Ò1°$Ñ7Ô7½"¼*ÀR¹.¼.Ð9IÑJÔJ€DÝ
Œ,�tØD *ÐD¨hÐDÐDÐDØ)˜:Ð)Ð)Ð)ñ+ô +€Cð �<Š<˜˜c˜UÑ#Ô#Ð#r    c                 ó&  — |j         }t          |¦  «        }|j        rdnd}t          j        t          j        d¦  «        |                      |¦  «        g¦  «        }t          j        |d|› |› d�d|› �¦  «        }	|                     |	|g¦  «        S )Nrn   rû   rá   rï   rð   rñ   )	r—   rù   rü   r   r|   r’   rå   rÆ   rI   rý   s
             r   Úinteger_to_float16_castr  �  s    € ð Œ€HÝ'¨Ñ1Ô1€JØœÐ.��¨3€JåŒ?�2œ: b™>œ>Ø#×2Ò2°6Ñ:Ô:Ð;ñ=ô =€Då
Œ,�tØC ZÐC°ÐCÐCÐCØ)˜ZÐ)Ð)ñ+ô +€Cð �<Š<˜˜c˜UÑ#Ô#Ð#r    c                 óh   ‡— t          | t          j        t          j        ¦  «        ˆfd„¦   «         }d S )Nc                 óü   •— t          j        t          j        d¦  «        t          j        d¦  «        t          j        d¦  «        g¦  «        }t          j        |‰› d�d¦  «        }|                     ||¦  «        S )Nrá   z.f16 $0,$1,$2;ú=h,h,h©r   r|   r’   rÆ   rI   )r&   r   r'   r(   r‚   rì   Úops         €r   Úptx_fp16_binaryz*lower_fp16_binary.<locals>.ptx_fp16_binary‘  sh   ø€ åŒ�rœz¨"™~œ~Ý "¤
¨2¡¤µ´
¸2±´Ð?ñAô AˆåŒl˜4 BÐ!6Ð!6Ð!6¸ÑAÔAˆØ�|Š|˜C Ñ&Ô&Ð&r    ©Úlowerr	   Úfloat16)Úfnr  r  s    ` r   Úlower_fp16_binaryr  �  sA   ø€ Ý
ˆ2�uŒ}�eœmÑ,Ô,ð'ð 'ð 'ð 'ñ -Ô,ð'ð 'ð 'r    ÚaddÚsubÚmulc                 óÎ   — t          j        t          j        d¦  «        t          j        d¦  «        g¦  «        }t          j        |dd¦  «        }|                     ||¦  «        S )Nrá   zneg.f16 $0, $1;ú=h,hr  ©r&   r   r'   r(   r‚   rì   s         r   Úptx_fp16_hnegr  ¤  óN   € åŒ?�2œ: b™>œ>­B¬J°r©N¬NÐ+;Ñ<Ô<€DÝ
Œ,�tÐ.°Ñ
7Ô
7€CØ�<Š<˜˜TÑ"Ô"Ð"r    c                 ó&   — t          | |||¦  «        S rÐ   )r  r%   s       r   Úoperator_hnegr  «  ó   € å˜ '¨3°Ñ5Ô5Ð5r    c                 óÎ   — t          j        t          j        d¦  «        t          j        d¦  «        g¦  «        }t          j        |dd¦  «        }|                     ||¦  «        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¦  «        }|                     ||¦  «        S )Nrá   zfma.rn.f16 $0,$1,$2,$3;z=h,h,h,h)r   r’   r|   rÆ   rI   )r&   r   r'   r(   Úargtysr‚   rì   s          r   Úptx_hfmar  ¼  sg   € åŒj˜‰nŒn�bœj¨™nœn­b¬j¸©n¬nÐ=€FÝŒ?�2œ: b™>œ>¨6Ñ2Ô2€DÝ
Œ,�tÐ6¸
Ñ
CÔ
C€CØ�<Š<˜˜TÑ"Ô"Ð"r    c                 ó8   — d„ }|                       ||||¦  «        S )Nc                 óB   — t           j                             | |¦  «        S rÐ   )r   Úfp16Úhdiv)r   r   s     r   Úfp16_divzfp16_div_impl.<locals>.fp16_divÇ  s   € ÝŒy�~Š~˜a Ñ#Ô#Ð#r    ©Úcompile_internal)r&   r   r'   r(   r#  s        r   Úfp16_div_implr&  Ä  s-   € ð$ð $ð $ð ×#Ò# G¨X°s¸DÑAÔ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                               ‰	¬¦  «        d¦  «        }|                     ||¦  «        }|                      t          j	        d¦  «        }| 
                    |t          j        d¦  «        ¦  «        }|                     d||¦  «        S )Nrá   )r  r  r   z!=)r   r|   r’   rÆ   Ú	_fp16_cmprW   rI   rF   r	   Úint16r–   Úicmp_unsigned)
r&   r   r'   r(   r‚   rì   ÚresultÚzeroÚ
int_resultr  s
            €r   Úptx_fp16_comparisonz*_gen_fp16_cmp.<locals>.ptx_fp16_comparisonÕ  s²   ø€ ÝŒ�rœz¨"™~œ~µ´
¸2±´ÅÄ
È2ÁÄÐ/OÑPÔPˆÝŒl˜4¥×!1Ò!1°RÐ!1Ñ!8Ô!8¸(ÑCÔCˆØ—’˜c 4Ñ(Ô(ˆà×#Ò#¥E¤K°Ñ3Ô3ˆØ—_’_ V­R¬Z¸©^¬^Ñ<Ô<ˆ
Ø×$Ò$ T¨:°tÑ<Ô<Ð<r    rS   )r  r/  s   ` r   Ú_gen_fp16_cmpr0  Ô  s$   ø€ ð=ð =ð =ð =ð =ð Ðr    ÚeqÚneÚgeÚgtÚleÚltc                 óh   ‡— t          | t          j        t          j        ¦  «        ˆfd„¦   «         }d S )Nc                 ó€   •—  t          ‰¦  «        | |||¦  «        }|                     ||d         |d         ¦  «        S )Nr   r   )r0  Úselect)r&   r   r'   r(   Úchoicer  s        €r   Úptx_fp16_minmaxz*lower_fp16_minmax.<locals>.ptx_fp16_minmaxï  s>   ø€ à"•˜rÑ"Ô" 7¨G°S¸$Ñ?Ô?ˆØ�~Š~˜f d¨1¤g¨t°A¬wÑ7Ô7Ð7r    r  )r  r�   r  r;  s     ` r   Úlower_fp16_minmaxr<  î  sA   ø€ Ý
ˆ2�uŒ}�eœmÑ,Ô,ð8ð 8ð 8ð 8ñ -Ô,ð8ð 8ð 8r    ÚmaxÚminÚ
__nv_cbrtfÚ	__nv_cbrtc                 óæ   — |j         }t          |         }|                      |¦  «        }|j        }t	          j        ||g¦  «        }t          j        |||¦  «        }	|                     |	|¦  «        S rÐ   )	Úreturn_typeÚ
cbrt_funcsrå   rH   r   r|   r
   r~   rI   )
r&   r   r'   r(   rê   r�   ÚftyrL   r‚   r  s
             r   Úptx_cbrtrE    si   € ð 
Œ€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¦  «        }|                     ||¦  «        S )Nr‘   Ú	__nv_brev©r
   r~   rH   r   r|   r’   rI   ©r&   r   r'   r(   r  s        r   Úptx_brev_u4rJ    sV   € õ
 
Ô	'ØŒÝ
Œ�œ
 2™œ­¬°B©¬Ð(9Ñ:Ô:Øñ
ô 
€Bð �<Š<˜˜DÑ!Ô!Ð!r    c           	      óÔ   — t          j        |j        t          j        t          j        d¦  «        t          j        d¦  «        f¦  «        d¦  «        }|                     ||¦  «        S )Nr”   Ú__nv_brevllrH  rI  s        r   Úptx_brev_u8rM    sV   € õ
 
Ô	'ØŒÝ
Œ�œ
 2™œ­¬°B©¬Ð(9Ñ:Ô:Øñ
ô 
€Bð �<Š<˜˜DÑ!Ô!Ð!r    c                 óv   — |                      |d         |                      t          j        d¦  «        ¦  «        S r:   )ÚctlzrF   r	   Úbooleanr%   s       r   Úptx_clzrQ  '  s4   € à�<Š<ØˆQŒØ×Ò�Uœ]¨AÑ.Ô.ñ0ô 0ð 0r    c           	      óÔ   — t          j        |j        t          j        t          j        d¦  «        t          j        d¦  «        f¦  «        d¦  «        }|                     ||¦  «        S )Nr‘   Ú__nv_ffsrH  rI  s        r   Ú
ptx_ffs_32rT  .  sV   € õ 
Ô	'ØŒÝ
Œ�œ
 2™œ­¬°B©¬Ð(9Ñ:Ô:Øñ
ô 
€Bð �<Š<˜˜DÑ!Ô!Ð!r    c           	      óÔ   — t          j        |j        t          j        t          j        d¦  «        t          j        d¦  «        f¦  «        d¦  «        }|                     ||¦  «        S )Nr‘   r”   Ú
__nv_ffsllrH  rI  s        r   Ú
ptx_ffs_64rW  8  sV   € õ 
Ô	'ØŒÝ
Œ�œ
 2™œ­¬°B©¬Ð(9Ñ:Ô:Øñ
ô 
€Bð �<Š<˜˜DÑ!Ô!Ð!r    c                 ó<   — |\  }}}|                      |||¦  «        S rÐ   )r9  )r&   r   r'   r(   ÚtestÚaÚbs          r   Úptx_selpr\  B  s#   € à�J€Dˆ!ˆQØ�>Š>˜$  1Ñ%Ô%Ð%r    c           	      óô   — t          j        |j        t          j        t          j        ¦   «         t          j        ¦   «         t          j        ¦   «         f¦  «        d¦  «        }|                     ||¦  «        S )NÚ
__nv_fmaxf©r
   r~   rH   r   r|   rš   rI   rI  s        r   Ú
ptx_max_f4r`  H  ó`   € å	Ô	'ØŒÝ
ŒÝŒL‰NŒNÝŒ\‰^Œ^�Rœ\™^œ^Ð,ñ	.ô 	.ð 	ñ
ô 
€Bð �<Š<˜˜DÑ!Ô!Ð!r    c           
      ó¼  — t          j        |j        t          j        t          j        ¦   «         t          j        ¦   «         t          j        ¦   «         f¦  «        d¦  «        }|                     ||                      ||d         |j        d         t          j
        ¦  «        |                      ||d         |j        d         t          j
        ¦  «        g¦  «        S )NÚ	__nv_fmaxr   r   ©r
   r~   rH   r   r|   r£   rI   Úcastr(   r	   ÚdoublerI  s        r   Ú
ptx_max_f8rg  S  ó­   € õ 
Ô	'ØŒÝ
ŒÝŒM‰OŒOÝŒ]‰_Œ_�bœm™oœoÐ.ñ	0ô 	0ð 	ñ
ô 
€Bð �<Š<˜Ø�Š�W˜d 1œg s¤x°¤{µE´LÑAÔAØ�Š�W˜d 1œg s¤x°¤{µE´LÑAÔAðñ ô ð r    c           	      óô   — t          j        |j        t          j        t          j        ¦   «         t          j        ¦   «         t          j        ¦   «         f¦  «        d¦  «        }|                     ||¦  «        S )NÚ
__nv_fminfr_  rI  s        r   Ú
ptx_min_f4rk  d  ra  r    c           
      ó¼  — t          j        |j        t          j        t          j        ¦   «         t          j        ¦   «         t          j        ¦   «         f¦  «        d¦  «        }|                     ||                      ||d         |j        d         t          j
        ¦  «        |                      ||d         |j        d         t          j
        ¦  «        g¦  «        S )NÚ	__nv_fminr   r   rd  rI  s        r   Ú
ptx_min_f8rn  o  rh  r    c           	      ó6  — t          j        |j        t          j        t          j        d¦  «        t          j        ¦   «         f¦  «        d¦  «        }|                     ||                      ||d         |j	        d         t          j        ¦  «        g¦  «        S )Nr”   Ú__nv_llrintr   )r
   r~   rH   r   r|   r’   r£   rI   re  r(   r	   rf  rI  s        r   Ú	ptx_roundrq  €  s†   € õ 
Ô	'ØŒÝ
ŒÝŒJ�r‰NŒNÝŒ]‰_Œ_Ðñ	 ô 	 ð 	ñ
ô 
€Bð �<Š<˜Ø�Š�W˜d 1œg s¤x°¤{µE´LÑAÔAðñ ô ð r    c                 ó8   — d„ }|                       ||||¦  «        S )Nc                 ó�  — t          j        | ¦  «        st          j        | ¦  «        r| S |dk    r7|dk    rd|dz
  z  }d}nd|z  }d}| |z  |z  }t          j        |¦  «        r| S nd| z  }| |z  }t          |¦  «        }t          j        ||z
  ¦  «        dk    rdt          |dz  ¦  «        z  }|dk    r	||z  |z  }n||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 q™MœMð 	ØˆHà�aŠ<ˆ<Ø˜Š|ˆ|ð  ¨"¡Ñ-�Ø��à˜w‘�Ø�Ø�T‘˜TÑ!ˆAÝŒz˜!‰}Œ}ð Ø�ðð ˜W˜HÑ%ˆDØ�D‘ˆAå�!‰HŒHˆÝŒI�a˜!‘eÑÔ Ò#Ð#à•e˜A ™G‘n”nÑ$ˆAà�aŠ<ˆ<Ø�T‘˜TÑ!ˆAˆAà�‰IˆAàˆr    r$  )r&   r   r'   r(   r}  s        r   Úround_to_implr~  ’  s.   € ðð ð ðB ×#Ò# G¨]¸CÀÑHÔHÐHr    c                 ó   ‡ — ˆ fd„}|S )Nc                 óz   •— |j         \  }|                      |‰¦  «        }|                     ||d         ¦  «        S r:   )r(   rF   Úfmul)r&   r   r'   r(   ÚargtyÚfactorÚconsts         €r   Úimplzgen_deg_rad.<locals>.implº  s9   ø€ Ø”‰ˆØ×%Ò% e¨UÑ3Ô3ˆØ�|Š|˜F D¨¤GÑ,Ô,Ð,r    rS   )r„  r…  s   ` r   Úgen_deg_radr†  ¹  s#   ø€ ð-ð -ð -ð -ð -ð €Kr    g     €f@c                 ó   ‡ ‡— |t           j        v rt          j        |d¬¦  «        }|g}n$t          j        ‰|t          |¦  «        ¬¦  «        }ˆˆ fd„t          ||¦  «        D ¦   «         }|j        }||k    rt          d|›d|›�¦  «        ‚|j	        t          |¦  «        k    r&t          d|j	        t          |¦  «        fz  ¦  «        ‚||fS )z4
    Convert integer indices into tuple of intp
    r   )r_   Úcount)rˆ  c                 óZ   •— g | ]'\  }}‰                      ‰||t          j        ¦  «        ‘Œ(S rS   )re  r	   Úintp)rm   ÚtÚir   r&   s      €€r   ro   z&_normalize_indices.<locals>.<listcomp>Ò  sA   ø€ ð 0ð 0ð 0Ù�1�að �|Š|˜G Q¨­5¬:Ñ6Ô6ð 0ð 0ð 0r    zexpect z	 but got z#indexing %d-D array with %d-D index)
r	   Úinteger_domainÚUniTupler
   Úunpack_tupleÚlenÚzipr_   Ú	TypeErrorÚndim)r&   r   ÚindtyÚindsÚarytyÚvaltyÚindicesr_   s   ``      r   Ú_normalize_indicesr™  É  sû   øø€ ð •Ô$Ð$Ð$Ý” U°!Ð4Ñ4Ô4ˆØ�&ˆˆåÔ& w°½CÀ¹J¼JÐGÑGÔGˆð0ð 0ð 0ð 0ð 0Ý˜u gÑ.Ô.ð0ñ 0ô 0€Gð ŒK€EØ�‚~€~Ýˆi°%°%°%¸¸Ð?Ñ@Ô@Ð@à„z•S˜‘Z”ZÒÐÝÐ=Øœ¥S¨¡Z¤ZÐ0ñ1ñ 2ô 2ð 	2ð �'ˆ>Ðr    c                 ó   ‡ — ˆ fd„}|S )Nc                 óô   •— |j         \  }}}|\  }}}	|j        }
t          | |||||¦  «        \  }} |                      |¦  «        | ||¦  «        }t	          j        | ||||d¬¦  «        } ‰| ||
||	¦  «        S )NT©Ú
wraparound)r(   r_   r™  Ú
make_arrayr
   Úget_item_pointer)r&   r   r'   r(   r–  r”  r—  Úaryr•  ré   r_   r˜  ÚlaryÚptrÚdispatch_fns                 €r   Úimpz_atomic_dispatcher.<locals>.impá  s¤   ø€ à!œhÑˆˆu�eØ‰ˆˆT�3Ø”ˆå+¨G°W¸eÀTØ,1°5ñ:ô :‰ˆˆwð )ˆw×!Ò! %Ñ(Ô(¨°'¸3Ñ?Ô?ˆÝÔ& w°¸ÀÀgØ26ð8ñ 8ô 8ˆð ˆ{˜7 G¨U°C¸Ñ=Ô=Ð=r    rS   )r£  r¤  s   ` r   Ú_atomic_dispatcherr¥  à  s#   ø€ ð>ð >ð >ð >ð >ð €Jr    c                 ó6  — |t           j        k    r1|j        }|                     t	          j        |¦  «        ||f¦  «        S |t           j        k    r1|j        }|                     t	          j        |¦  «        ||f¦  «        S |                     d||d¦  «        S )Nr  Ú	monotonic)	r	   r™   rH   rI   r   Údeclare_atomic_add_float32r¢   Údeclare_atomic_add_float64Ú
atomic_rmw©r&   r   r_   r¢  ré   rL   s         r   Úptx_atomic_add_tupler¬  ò  óœ   € ð
 •”ÒÐØŒ~ˆØ�|Š|�IÔ@ÀÑFÔFØ  #˜Jñ(ô (ð 	(à	•%”-Ò	Ð	ØŒ~ˆØ�|Š|�IÔ@ÀÑFÔFØ  #˜Jñ(ô (ð 	(ð ×!Ò! %¨¨c°;Ñ?Ô?Ð?r    c                 ó6  — |t           j        k    r1|j        }|                     t	          j        |¦  «        ||f¦  «        S |t           j        k    r1|j        }|                     t	          j        |¦  «        ||f¦  «        S |                     d||d¦  «        S )Nr  r§  )	r	   r™   rH   rI   r   Údeclare_atomic_sub_float32r¢   Údeclare_atomic_sub_float64rª  r«  s         r   Úptx_atomic_subr±    r­  r    c                 óÜ   — |t           j        j        v rG|j        }|j        }t          t          d|› �¦  «        }|                      ||¦  «        ||f¦  «        S t          d|› d�¦  «        ‚)NÚdeclare_atomic_inc_intzUnimplemented atomic inc with ú array©	r   ÚcudadeclÚunsigned_int_numba_typesr—   rH   Úgetattrr   rI   r’  ©r&   r   r_   r¢  ré   ÚbwrL   r  s           r   Úptx_atomic_incr»    ót   € ð
 •”Ô6Ð6Ð6ØŒ^ˆØŒ~ˆÝ•YÐ =¸Ð =Ð =Ñ>Ô>ˆØ�|Š|˜B˜B˜t™HœH s¨C jÑ1Ô1Ð1åÐF¸ÐFÐFÐFÑGÔGÐGr    c                 óÜ   — |t           j        j        v rG|j        }|j        }t          t          d|› �¦  «        }|                      ||¦  «        ||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 ]1} t          | t          j        |t          j        ¦  «        |¦  «         Œ2d S )Nc                 ó†   •— |t           j        j        v r|                     ‰||d¦  «        S t	          d‰› d|› d�¦  «        ‚)Nr§  zUnimplemented atomic z with r´  ©r   r¶  Úinteger_numba_typesrª  r’  )r&   r   r_   r¢  ré   r  s        €r   Úimpl_ptx_atomicz+ptx_atomic_bitwise.<locals>.impl_ptx_atomic1  sO   ø€ à•T”]Ô6Ð7Ð7Ø×%Ò% b¨#¨s°KÑ@Ô@Ð@åÐK°BÐKÐK¸eÐKÐKÐKÑLÔLÐLr    )r¥  r	   rŠ  rŽ  ÚTupler	  ÚArrayÚAny)Ústubr  rÄ  rê   s    `  r   Úptx_atomic_bitwiserÉ  0  s}   ø€ ÝðMð Mð Mð Mñ ÔðMõ Œz�5œ>­5¬;Ð7ð Að AˆØ/�ˆd•E”K ¥U¤YÑ/Ô/°Ñ@Ô@Ð@Ð@ðAð Ar    ÚandÚorÚxorc                 ó~   — |t           j        j        v r|                     d||d¦  «        S t	          d|› d�¦  «        ‚)NÚxchgr§  zUnimplemented atomic exch with r´  rÂ  )r&   r   r_   r¢  ré   s        r   Úptx_atomic_exchrÏ  A  sH   € ð
 •”Ô2Ð3Ð3Ø×!Ò! &¨#¨s°KÑ@Ô@Ð@åÐG¸%ÐGÐGÐGÑHÔHÐHr    c                 óè  — |j         }|t          j        k    r*|                     t	          j        |¦  «        ||f¦  «        S |t          j        k    r*|                     t	          j        |¦  «        ||f¦  «        S |t          j        t          j	        fv r| 
                    d||d¬¦  «        S |t          j        t          j        fv r| 
                    d||d¬¦  «        S t          d|z  ¦  «        ‚©Nr=  r§  ©ÚorderingÚumaxz&Unimplemented atomic max with %s array)rH   r	   r¢   rI   r   Údeclare_atomic_max_float64r™   Údeclare_atomic_max_float32rG   Úint64rª  Úuint32Úuint64r’  r«  s         r   Úptx_atomic_maxrÚ  L  óï   € ð
 Œ>€DØ•”ÒÐØ�|Š|�IÔ@ÀÑFÔFØ  #˜Jñ(ô (ð 	(à	•%”-Ò	Ð	Ø�|Š|�IÔ@ÀÑFÔFØ  #˜Jñ(ô (ð 	(à	•5”;¥¤Ð,Ð	,Ð	,Ø×!Ò! %¨¨c¸KÐ!ÑHÔHÐHØ	•5”<¥¤Ð.Ð	.Ð	.Ø×!Ò! &¨#¨s¸[Ð!ÑIÔIÐIåÐ@À5ÑHÑIÔIÐIr    c                 óè  — |j         }|t          j        k    r*|                     t	          j        |¦  «        ||f¦  «        S |t          j        k    r*|                     t	          j        |¦  «        ||f¦  «        S |t          j        t          j	        fv r| 
                    d||d¬¦  «        S |t          j        t          j        fv r| 
                    d||d¬¦  «        S t          d|z  ¦  «        ‚©Nr>  r§  rÒ  Úuminz&Unimplemented atomic min with %s array)rH   r	   r¢   rI   r   Údeclare_atomic_min_float64r™   Údeclare_atomic_min_float32rG   r×  rª  rØ  rÙ  r’  r«  s         r   Úptx_atomic_minrá  `  rÛ  r    c                 óè  — |j         }|t          j        k    r*|                     t	          j        |¦  «        ||f¦  «        S |t          j        k    r*|                     t	          j        |¦  «        ||f¦  «        S |t          j        t          j	        fv r| 
                    d||d¬¦  «        S |t          j        t          j        fv r| 
                    d||d¬¦  «        S t          d|z  ¦  «        ‚rÑ  )rH   r	   r¢   rI   r   Údeclare_atomic_nanmax_float64r™   Údeclare_atomic_nanmax_float32rG   r×  rª  rØ  rÙ  r’  r«  s         r   Úptx_atomic_nanmaxrå  t  óï   € ð
 Œ>€DØ•”ÒÐØ�|Š|�IÔCÀDÑIÔIØ  #˜Jñ(ô (ð 	(à	•%”-Ò	Ð	Ø�|Š|�IÔCÀDÑIÔIØ  #˜Jñ(ô (ð 	(à	•5”;¥¤Ð,Ð	,Ð	,Ø×!Ò! %¨¨c¸KÐ!ÑHÔHÐHØ	•5”<¥¤Ð.Ð	.Ð	.Ø×!Ò! &¨#¨s¸[Ð!ÑIÔIÐIåÐ@À5ÑHÑIÔIÐIr    c                 óè  — |j         }|t          j        k    r*|                     t	          j        |¦  «        ||f¦  «        S |t          j        k    r*|                     t	          j        |¦  «        ||f¦  «        S |t          j        t          j	        fv r| 
                    d||d¬¦  «        S |t          j        t          j        fv r| 
                    d||d¬¦  «        S t          d|z  ¦  «        ‚rÝ  )rH   r	   r¢   rI   r   Údeclare_atomic_nanmin_float64r™   Údeclare_atomic_nanmin_float32rG   r×  rª  rØ  rÙ  r’  r«  s         r   Úptx_atomic_nanminrê  ˆ  ræ  r    c                 ó  — |                      |j        d         t          j        |j        d         |j        d         ¦  «        }|d         |                      t          j        d¦  «        |d         |d         f}t          | |||¦  «        S )Nr   r   rC   )rB  r(   r	   rŠ  rF   Úptx_atomic_casr%   s       r   Úptx_atomic_compare_and_swaprí  œ  sm   € à
�/Š/˜#œ( 1œ+¥u¤z°3´8¸A´;ÀÄÈÄÑ
LÔ
L€CØ�ŒG�W×)Ò)­%¬*°aÑ8Ô8¸$¸q¼'À4ÈÄ7ÐK€DÝ˜' 7¨C°Ñ6Ô6Ð6r    c                 ó€  — |j         \  }}}}|\  }}	}
}t          | |||	||¦  «        \  }} |                      |¦  «        | ||¦  «        }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Ÿ  r_   r   r¶  rÃ  rH   r—   r   Úatomic_cmpxchgr’  )r&   r   r'   r(   r–  r”  Úoldtyr—  r   r•  Úoldré   r˜  r¡  r¢  rL   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ñ 4ô 4€Cð „{•t”}Ô8Ð9Ð9ØŒ~ˆØ”;Ô'ˆÝÔ'¨°°xÀÀcÈ3ÑOÔOÐOåÐ@À5Ä;ÑNÑOÔOÐOr    c                 óâ   — t          j        t          j        t          j        ¦   «         t          j        d¦  «        g¦  «        ddd¬¦  «        }|d         }|                     ||g¦  «         d S )Nr‘   znanosleep.u32 $0;rö   TrÃ   r   )r   rÆ   r|   r}   r’   rI   )r&   r   r'   r(   Ú	nanosleepÚnss         r   Úptx_nanosleeprõ  »  sf   € å”�Rœ_­R¬[©]¬]½R¼ZÈ¹^¼^Ð<LÑMÔMØ0°#À4ðIñ Iô I€Ià	ˆaŒ€BØ‡L‚L�˜R˜DÑ!Ô!Ð!Ð!Ð!r    Fc           
      ó  ‡ — t          t          j        |d¦  «        }|dk    o|ot          |¦  «        dk    }|dk    r|st	          d¦  «        ‚‰ j        |         }	t          |t          j        t          j	        f¦  «        p)t          |	t          j        ¦  «        p|t          j        k    }
|t          j        vr|
st          d|z  ¦  «        ‚‰                      |¦  «        }t!          j        ||¦  «        }|t$          j        k    rt)          j        |||¬¦  «        }n»|j        }t)          j        ||||¦  «        }‰                      |¦  «        }d|dz
                       ¦   «         z  |_        |rd|_        n$t!          j        |t           j        ¦  «        |_        |                     |t!          j         t!          j!        d¦  «        ¦  «        d¦  «        }tE          j#        t%          j$        ¦   «         j%        ¦  «        }‰                      |¦  «        }| &                    |¦  «        }|}g }tO          tQ          |¦  «        ¦  «        D ]\  }}| )                    |¦  «         ||z  }Œ d	„ tQ          |¦  «        D ¦   «         }ˆ fd
„|D ¦   «         }|r±t!          j*        t!          j+        t!          j!        d¦  «        g ¦  «        ddd¬¦  «        }| ,                    | -                    |g ¦  «        t!          j!        d¦  «        ¦  «        }‰  .                    t          j/        |¦  «        }| 0                    ||¦  «        g}nˆ fd„|D ¦   «         }t          |¦  «        }t          j1        ||d¬¦  «        } ‰  2                    |¦  «        ‰ |¦  «        }‰  3                    || 4                    ||j5        j6        ¦  «        ||‰  .                    t          j/        |¦  «        d ¬¦  «         | 7                    ¦   «         S )Nr   r   zarray length <= 0zunsupported type: %srX   Úexternalrø   Úgenericc                 ó   — g | ]}|‘ŒS rS   rS   rl   s     r   ro   z"_generic_array.<locals>.<listcomp>  s   € Ð-Ð-Ð-�QˆqÐ-Ð-Ð-r    c                 óP   •— g | ]"}‰                      t          j        |¦  «        ‘Œ#S rS   ©rF   r	   rŠ  ©rm   rn   r&   s     €r   ro   z"_generic_array.<locals>.<listcomp>  s+   ø€ ÐEÐEÐE¸�×$Ò$¥U¤Z°Ñ3Ô3ÐEÐEÐEr    r‘   zmov.u32 $0, %dynamic_smem_size;rÂ   TrÃ   r”   c                 óP   •— g | ]"}‰                      t          j        |¦  «        ‘Œ#S rS   rû  rü  s     €r   ro   z"_generic_array.<locals>.<listcomp>  s+   ø€ ÐEÐEÐE¸!�'×&Ò&¥u¤z°1Ñ5Ô5ÐEÐEÐEr    ÚC)r_   r“  Úlayout)Údatar^   ÚstridesÚitemsizeÚmeminfo)8r   Úoperatorr  r�  Ú
ValueErrorÚdata_model_managerÚ
isinstancer	   ÚRecordÚBooleanr   ÚStructModelr
  Únumber_domainr’  Úget_data_typer   Ú	ArrayTyper   ru   r
   Úalloca_oncerH   Ú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Ÿ   rI   rF   rŠ  ÚudivrÆ  rž  Úpopulate_arrayr–   r   ÚtypeÚ	_getvalue) r&   r   r^   r_   r`   ra   rb   Ú	elemcountÚdynamic_smemÚ
data_modelÚother_supported_typeÚlldtypeÚlarytyÚdataptrrL   Úgvmemr  Ú
targetdatar  Ú
laststrideÚrstridesrŒ  Úlastsizer  ÚkstridesÚget_dynshared_sizeÚdynsmem_sizeÚ	kitemsizeÚkshaper“  r–  r   s    `                               r   rd   rd   Æ  së  ø€ å•x”| U¨AÑ.Ô.€Ið  ’>ÐF lÐFµs¸5±z´zÀQ²€LØ�A‚~€~˜l€~ÝÐ,Ñ-Ô-Ð-ð Ô+¨EÔ2€Jå�5�5œ<­¬Ð7Ñ8Ô8ð 	"Ý�j¥&Ô"4Ñ5Ô5ð	"à•E”MÒ!ð ð
 •EÔ'Ð'Ð'Ð0DÐ'ÝÐ.°Ñ6Ñ7Ô7Ð7à×#Ò# EÑ*Ô*€GÝŒ\˜' 9Ñ-Ô-€Fà•DÔ(Ò(Ð(õ Ô% g¨v¸KÐHÑHÔHˆˆàŒ~ˆõ Ô+¨D°&¸+Ø,5ñ7ô 7ˆð ×&Ò& wÑ/Ô/ˆð ˜E A™I×2Ò2Ñ4Ô4Ñ4ˆŒàð 		BØ&ˆEŒMˆMõ !#¤¨FµB´LÑ AÔ AˆEÔð ×'Ò'¨­r¬~½b¼jÈ¹m¼mÑ/LÔ/LØ(1ñ3ô 3ˆõ Ô&¥t¤y¡{¤{Ô'>Ñ?Ô?€JØ×#Ò# EÑ*Ô*€GØ×#Ò# JÑ/Ô/€Hð €JØ€HÝ ¥¨%¡¤Ñ1Ô1ð ð ‰ˆˆ8Ø�Š˜
Ñ#Ô#Ð#Ø�hÑˆ
ˆ
Ø-Ð-�( 8Ñ,Ô,Ð-Ñ-Ô-€GØEÐEÐEÐE¸WÐEÑEÔE€Hð ð Fõ
  œ\­"¬/½"¼*ÀR¹.¼.È"Ñ*MÔ*MØ*KØ*.¸DðBñ Bô BÐð —|’| G§L¢LÐ1CÀRÑ$HÔ$HÝ$&¤J¨r¡N¤Nñ4ô 4ˆð ×(Ò(­¬°XÑ>Ô>ˆ	Ø—,’,˜|¨YÑ7Ô7Ð8ˆˆàEÐEÐEÐE¸uÐEÑEÔEˆõ ˆu‰:Œ:€DÝŒK˜e¨$°sÐ;Ñ;Ô;€EØ
#ˆ'×
Ò
˜UÑ
#Ô
# G¨WÑ
5Ô
5€Cà×Ò˜3Ø '§¢°¸¼¼Ñ GÔ GØ!'Ø#+Ø$+×$8Ò$8½¼ÀXÑ$NÔ$NØ#'ð ñ )ô )ð )ð �=Š=‰?Œ?Ðr    c                 ó*   — |                       ¦   «         S rÐ   )r   )r&   r   rê   Úpyvals       r   Úcuda_dispatcher_constr8  *  s   € à×"Ò"Ñ$Ô$Ð$r    )F)×Ú	functoolsr   r  ru  Ú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   r   Úregistryr	  Úlower_getattrÚ
lower_attrÚlower_constantr   ÚModuler)   r-   r1   r5   r8   r=   rA   rD   ÚcgÚ	this_gridrM   rQ   r„  Ú
array_likerÆ  rT   rV   rZ   ÚsharedÚarrayÚIntegerLiteralrÇ  rh   rÅ  rŽ  rr   Úlocalrv   ry   Úthreadfence_blockr„   Úthreadfence_systemr‡   Úthreadfencer‰   Úsyncwarpr�   Úi4rŒ   Úshfl_sync_intrinsicrž   Úf4Úf8r¸   Úvote_sync_intrinsicrP  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)  r0  Úheqr1  Úhner2  Úhger3  Úhgtr4  Úhler5  Úhltr6  r<  ÚhmaxÚhminr™   r¢   rC  ÚcbrtrE  ÚbrevÚu4rJ  Úu8rM  ÚclzrQ  ÚffsrT  rW  Úselpr\  r=  r`  rg  r>  rk  rn  rx  rq  r~  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õ  rd   r8  Ú
get_ufuncsrS   r    r   ú<module>r�     sõ  ðØ Ð Ð Ð Ð Ð Ø €€€Ø €€€à Ð Ð Ð Ð Ð Ø Ð Ð Ð Ð Ð à 4Ð 4Ð 4Ð 4Ð 4Ð 4Ð 4Ð 4Ø 1Ð 1Ð 1Ð 1Ð 1Ð 1Ø 'Ð 'Ð 'Ð 'Ð 'Ð 'Ø %Ð %Ð %Ð %Ð %Ð %Ð %Ð %Ø Ð Ð Ð Ð Ð Ø ,Ð ,Ð ,Ð ,Ð ,Ð ,Ø Ð Ð Ð Ð Ð Ø Ð Ð Ð Ð Ð Ø /Ð /Ð /Ð /Ð /Ð /Ð /Ð /Ð /Ð /Ø =Ð =Ð =Ð =Ð =Ð =Ð =Ð =Ð =Ð =àˆ8‰:Œ:€ØŒ€ØÔ#€
ØÔ(€ð3ð 3ð 3ð €ˆLˆEŒL˜ÑÔ Ñ,Ô,ð+ð +ñ -Ô,ð+ð €ˆLˆEŒL˜ÑÔ 
Ñ+Ô+ð,ð ,ñ ,Ô+ð,ð €ˆLˆEŒL˜ÑÔ 
Ñ+Ô+ð-ð -ñ ,Ô+ð-ð €ˆLˆEŒL˜ÑÔ 	Ñ*Ô*ð.ð .ñ +Ô*ð.ð €ˆLˆEŒL˜ÑÔ Ñ)Ô)ð2ð 2ñ *Ô)ð2ð €ˆD�#ÑÔð*ð *ñ Ôð*ð €ˆD�#ÑÔð*ð *ñ Ôð*ð €ˆD�#ÑÔð*ð *ñ Ôð*ð €€t„wÔÑÔðð ñ Ôðð €Ð˜Ñ$Ô$ðð ñ %Ô$ðð €€t„zÔ˜eœkÑ*Ô*ðð ñ +Ô*ðð €ð3ð 3ð 3ð €€t„{Ô˜%Ô.°´	Ñ:Ô:ð-ð -ñ ;Ô:ð-ð €€t„{Ô˜%œ+ u¤yÑ1Ô1Ø€€t„{Ô˜%œ.¨%¬)Ñ4Ô4ð-ð -ñ 5Ô4ñ 2Ô1ð-ð €€t„zÔ˜Ô-¨u¬yÑ9Ô9ð.ð .ñ :Ô9ð.ð €€t„zÔ˜œ e¤iÑ0Ô0Ø€€t„zÔ˜œ¨¬Ñ3Ô3ð.ð .ñ 4Ô3ñ 1Ô0ð.ð €€uÔÑÔð%ð %ñ  Ôð%ð €€uÔÑ Ô ð%ð %ñ !Ô ð%ð €€uÔÑÔð%ð %ñ Ôð%ð €€u„~ÑÔðAð Añ ÔðAð €€u„~�u”xÑ Ô ð%ð %ñ !Ô ð%ð €€uÔ  %¤(¨E¬H°e´hÀÄØ„xñô à€€uÔ  %¤(¨E¬H°e´hÀÄØ„xñô à€€uÔ  %¤(¨E¬H°e´hÀÄØ„xñô à€€uÔ  %¤(¨E¬H°e´hÀÄØ„xñô ð+ð +ñô ñô ñô ñô ð+ð\ €€uÔ  %¤(¨E¬H°e´mÑDÔDð$ð $ñ EÔDð$ð €€uÔ˜UœX u¤xÑ0Ô0Ø€€uÔ˜UœX u¤xÑ0Ô0Ø€€uÔ˜UœX u¤xÑ0Ô0Ø€€uÔ˜UœX u¤xÑ0Ô0ð	-ð 	-ñ 1Ô0ñ 1Ô0ñ 1Ô0ñ 1Ô0ð	-ð €€uÔ˜UœX u¤xÑ0Ô0Ø€€uÔ˜UœX u¤xÑ0Ô0Ø€€uÔ˜UœX u¤xÑ0Ô0Ø€€uÔ˜UœX u¤xÑ0Ô0ð-ð -ñ 1Ô0ñ 1Ô0ñ 1Ô0ñ 1Ô0ð-ð €€uÔÑÔð(ð (ñ Ôð(ð €€uÔÑÔð(ð (ñ Ôð(ð €€u„z�5”9ÑÔð"ð "ñ Ôð"ð €€u„y�%”)˜UœY¨¬	Ñ2Ô2ðð ñ 3Ô2ðð,ð ,ð ,ð €ˆEŒM˜5œ;Ñ'Ô'ð$ð $ñ (Ô'ð$ð €ˆEŒK˜œÑ'Ô'ð$ð $ñ (Ô'ð$ð,ð ,ð ,ð €ˆEŒM˜5œ=Ñ)Ô)ð	$ð 	$ñ *Ô)ð	$ð €ˆEŒM˜5œ=Ñ)Ô)Ø€ˆEÔ  %¤-Ñ0Ô0ð
$ð 
$ñ 1Ô0ñ *Ô)ð
$ð'ð 'ð 'ð Ð �%”*”/ 5Ñ )Ô )Ð )Ø Ð �(”, Ñ &Ô &Ð &Ø Ð �(”- Ñ 'Ô 'Ð 'Ø Ð �%”*”/ 5Ñ )Ô )Ð )Ø Ð �(”, Ñ &Ô &Ð &Ø Ð �(”- Ñ 'Ô 'Ð 'Ø Ð �%”*”/ 5Ñ )Ô )Ð )Ø Ð �(”, Ñ &Ô &Ð &Ø Ð �(”- Ñ 'Ô 'Ð 'ð €€u„z„˜œÑ&Ô&ð#ð #ñ 'Ô&ð#ð €€x„|�U”]Ñ#Ô#ð6ð 6ñ $Ô#ð6ð €€u„z„˜œÑ&Ô&ð#ð #ñ 'Ô&ð#ð €€sˆEŒMÑÔð6ð 6ñ Ôð6ð €€u„z„˜œ u¤}°e´mÑDÔDð#ð #ñ EÔDð#ð €€xÔ˜œ¨¬Ñ6Ô6Ø€€xÔ˜%œ-¨¬Ñ7Ô7ðBð Bñ 8Ô7ñ 7Ô6ðBð€	ð	ð 	ð 	ð 4€€€e„j„n�e”m U¤]Ñ 3Ô 3°M°MÀ$Ñ4GÔ4GÑ HÔ HÐ HØ 0€€€h„k�5”= %¤-Ñ 0Ô 0°°¸tÑ1DÔ1DÑ EÔ EÐ EØ 3€€€e„j„n�e”m U¤]Ñ 3Ô 3°M°MÀ$Ñ4GÔ4GÑ HÔ HÐ HØ 0€€€h„k�5”= %¤-Ñ 0Ô 0°°¸tÑ1DÔ1DÑ EÔ EÐ EØ 3€€€e„j„n�e”m U¤]Ñ 3Ô 3°M°MÀ$Ñ4GÔ4GÑ HÔ HÐ HØ 0€€€h„k�5”= %¤-Ñ 0Ô 0°°¸tÑ1DÔ1DÑ EÔ EÐ EØ 3€€€e„j„n�e”m U¤]Ñ 3Ô 3°M°MÀ$Ñ4GÔ4GÑ HÔ HÐ HØ 0€€€h„k�5”= %¤-Ñ 0Ô 0°°¸tÑ1DÔ1DÑ EÔ EÐ EØ 3€€€e„j„n�e”m U¤]Ñ 3Ô 3°M°MÀ$Ñ4GÔ4GÑ HÔ HÐ HØ 0€€€h„k�5”= %¤-Ñ 0Ô 0°°¸tÑ1DÔ1DÑ EÔ EÐ EØ 3€€€e„j„n�e”m U¤]Ñ 3Ô 3°M°MÀ$Ñ4GÔ4GÑ HÔ HÐ HØ 0€€€h„k�5”= %¤-Ñ 0Ô 0°°¸tÑ1DÔ1DÑ EÔ EÐ Eð8ð 8ð 8ð Ð �%”*”/ 5¨$Ñ /Ô /Ð /Ø Ð �%”*”/ 5¨$Ñ /Ô /Ð /ð 
„M�<Ø	„M�;ð€
ð €€u„z�5”=Ñ!Ô!Ø€€u„z�5”=Ñ!Ô!ð"ð "ñ "Ô!ñ "Ô!ð"ð €€u„z�5”8ÑÔð"ð "ñ Ôð"ð €€u„z�5”8ÑÔð"ð "ñ Ôð"ð €€u„y�%”)ÑÔð0ð 0ñ Ôð0ð €€u„y�%”(ÑÔØ€€u„y�%”(ÑÔð"ð "ñ Ôñ Ôð"ð €€u„y�%”(ÑÔØ€€u„y�%”(ÑÔð"ð "ñ Ôñ Ôð"ð €€u„z�5”9˜eœi¨¬Ñ3Ô3ð&ð &ñ 4Ô3ð&ð
 €€sˆEŒH�e”hÑÔð"ð "ñ  Ôð"ð €€sˆEŒH�e”hÑÔØ€€sˆEŒH�e”hÑÔØ€€sˆEŒH�e”hÑÔðð ñ  Ôñ  Ôñ  Ôðð €€sˆEŒH�e”hÑÔð"ð "ñ  Ôð"ð €€sˆEŒH�e”hÑÔØ€€sˆEŒH�e”hÑÔØ€€sˆEŒH�e”hÑÔðð ñ  Ôñ  Ôñ  Ôðð €€uˆeŒhÑÔØ€€uˆeŒhÑÔð	ð 	ñ Ôñ Ôð	ð  €€uˆeŒh˜œÑ&Ô&Ø€€uˆeŒh˜œÑ&Ô&ð"Ið "Iñ 'Ô&ñ 'Ô&ð"IðJð ð ð Œ7�T‰>€Ø�$”'‰>€Ø €€€d„l�E”HÑ Ô ˜k˜k¨(Ñ3Ô3Ñ 4Ô 4Ð 4Ø €€€d„l�E”HÑ Ô ˜k˜k¨(Ñ3Ô3Ñ 4Ô 4Ð 4Ø €€€d„l�E”HÑ Ô ˜k˜k¨(Ñ3Ô3Ñ 4Ô 4Ð 4Ø €€€d„l�E”HÑ Ô ˜k˜k¨(Ñ3Ô3Ñ 4Ô 4Ð 4ðð ð ð.ð ð ð$ €€u„|Ô˜œ e¤j°%´)Ñ<Ô<Ø€€u„|Ô˜œ e¤n°e´iÑ@Ô@Ø€€u„|Ô˜œ e¤k°5´9Ñ=Ô=Øð
@ð 
@ñ Ôñ >Ô=ñ AÔ@ñ =Ô<ð
@ð €€u„|Ô˜œ e¤j°%´)Ñ<Ô<Ø€€u„|Ô˜œ e¤n°e´iÑ@Ô@Ø€€u„|Ô˜œ e¤k°5´9Ñ=Ô=Øð
@ð 
@ñ Ôñ >Ô=ñ AÔ@ñ =Ô<ð
@ð €€u„|Ô˜œ e¤j°%´)Ñ<Ô<Ø€€u„|Ô˜œ e¤n°e´iÑ@Ô@Ø€€u„|Ô˜œ e¤k°5´9Ñ=Ô=ØðHð Hñ Ôñ >Ô=ñ AÔ@ñ =Ô<ðHð €€u„|Ô˜œ e¤j°%´)Ñ<Ô<Ø€€u„|Ô˜œ e¤n°e´iÑ@Ô@Ø€€u„|Ô˜œ e¤k°5´9Ñ=Ô=ØðHð Hñ Ôñ >Ô=ñ AÔ@ñ =Ô<ðHð	Að 	Að 	Að Ð �5”<Ô$ eÑ ,Ô ,Ð ,Ø Ð �5”<Ô# TÑ *Ô *Ð *Ø Ð �5”<Ô# UÑ +Ô +Ð +ð €€u„|Ô˜%œ+ u¤z°5´9Ñ=Ô=Ø€€u„|Ô˜%œ+ u¤~°u´yÑAÔAØ€€u„|Ô˜%œ+ u¤{°E´IÑ>Ô>ØðIð Iñ Ôñ ?Ô>ñ BÔAñ >Ô=ðIð €€u„|Ô˜œ e¤j°%´)Ñ<Ô<Ø€€u„|Ô˜œ e¤k°5´9Ñ=Ô=Ø€€u„|Ô˜œ e¤n°e´iÑ@Ô@ØðJð Jñ Ôñ AÔ@ñ >Ô=ñ =Ô<ðJð  €€u„|Ô˜œ e¤j°%´)Ñ<Ô<Ø€€u„|Ô˜œ e¤k°5´9Ñ=Ô=Ø€€u„|Ô˜œ e¤n°e´iÑ@Ô@ØðJð Jñ Ôñ AÔ@ñ >Ô=ñ =Ô<ðJð  €€u„|Ô˜EœK¨¬°U´YÑ?Ô?Ø€€u„|Ô˜EœK¨¬°e´iÑ@Ô@Ø€€u„|Ô˜EœK¨¬¸¼ÑCÔCØðJð Jñ Ôñ DÔCñ AÔ@ñ @Ô?ðJð  €€u„|Ô˜EœK¨¬°U´YÑ?Ô?Ø€€u„|Ô˜EœK¨¬°e´iÑ@Ô@Ø€€u„|Ô˜EœK¨¬¸¼ÑCÔCØðJð Jñ Ôñ DÔCñ AÔ@ñ @Ô?ðJð  €€u„|Ô$ e¤k°5´9¸e¼iÑHÔHð7ð 7ñ IÔHð7ð €€u„|Ô˜œ e¤j°%´)¸U¼YÑGÔGØ€€u„|Ô˜œ e¤k°5´9¸e¼iÑHÔHØ€€u„|Ô˜œ e¤n°e´iÀÄÑKÔKðPð Pñ LÔKñ IÔHñ HÔGðPð* €€u„˜œÑ%Ô%ð"ð "ñ &Ô%ð"ð !&ðað að að aðH €�ÑÔð%ð %ñ  Ôð%ð €Ð#�Ô#Ñ%Ô% uÑ -Ô -Ð -Ð -Ð -r    