§
    ƒÛg|\  ã                   ó.  — d dl Z d dlmZ d dlmZmZmZmZmZm	Z	m
Z
 d dlmZmZmZmZmZmZ d dlmZmZ d dlmZ d dlmZ d dlmZ  e¦   «         Zej        Zej        Zej        Z ee¦  «          G d	„ d
e¦  «        Ze G d„ de¦  «        ¦   «         Z e G d„ de¦  «        ¦   «         Z!e G d„ de¦  «        ¦   «         Z"e G d„ de¦  «        ¦   «         Z#e G d„ de¦  «        ¦   «         Z$e G d„ de¦  «        ¦   «         Z%e G d„ de¦  «        ¦   «         Z&e G d„ de¦  «        ¦   «         Z'e G d„ de¦  «        ¦   «         Z( G d„ de¦  «        Z)e G d„ d e¦  «        ¦   «         Z*e G d!„ d"e¦  «        ¦   «         Z+e G d#„ d$e¦  «        ¦   «         Z,e G d%„ d&e¦  «        ¦   «         Z-e G d'„ d(e¦  «        ¦   «         Z.e G d)„ d*e¦  «        ¦   «         Z/e G d+„ d,e¦  «        ¦   «         Z0e G d-„ d.e¦  «        ¦   «         Z1e G d/„ d0e¦  «        ¦   «         Z2e G d1„ d2e¦  «        ¦   «         Z3e G d3„ d4e¦  «        ¦   «         Z4e G d5„ d6e¦  «        ¦   «         Z5e G d7„ d8e¦  «        ¦   «         Z6e G d9„ d:e¦  «        ¦   «         Z7e G d;„ d<e¦  «        ¦   «         Z8d=„ Z9d>„ Z:d?„ Z; ee<¦  «         G d@„ dAe¦  «        ¦   «         Z=dB„ Z>dC„ Z?dD„ Z@dE„ ZA e;ejB        jC        ¦  «        ZD eAe jE        ¦  «        ZF eAe jG        ¦  «        ZH e;ejB        jI        ¦  «        ZJ eAe jK        ¦  «        ZL eAe jM        ¦  «        ZN e;ejB        jO        ¦  «        ZP eAe jQ        ¦  «        ZR eAe jS        ¦  «        ZT e;ejB        jU        ¦  «        ZV e;ejB        jW        ¦  «        ZX e9ejB        jY        ¦  «        ZZ e:e j[        ¦  «        Z\ e9ejB        j]        ¦  «        Z^ e:e_¦  «        Z` e>ejB        ja        ¦  «        Zb e@e jc        ¦  «          e>ejB        jd        ¦  «        Ze e@e jf        ¦  «          e>ejB        jg        ¦  «        Zh e@e ji        ¦  «          e>ejB        jj        ¦  «        Zk e@e jl        ¦  «          e>ejB        jm        ¦  «        Zn e@e jo        ¦  «          e>ejB        jp        ¦  «        Zq e@e jr        ¦  «          eAe js        ¦  «          eAe jt        ¦  «         dF„ ZudG„ Zv eudH¦  «        Zw eudI¦  «        Zx eudJ¦  «        Zy eudK¦  «        Zz eudL¦  «        Z{ eudM¦  «        Z| eudN¦  «        Z} eudO¦  «        Z~ eudP¦  «        Z eudQ¦  «        Z€ eudR¦  «        Z� eudS¦  «        Z‚ eudT¦  «        Zƒ eudU¦  «        Z„ eudV¦  «        Z… evdW¦  «        Z†dX„ Z‡ejˆ        ej‰        ejŠ        ej‹        ejŒ        ej�        fZŽejŠ        ej‹        ejŒ        ej�        fZ�ej‹        ej�        fZ� e‡ej‘        jE        eŽ¦  «        Z’ e‡ej‘        jK        eŽ¦  «        Z“ e‡ej‘        j”        eŽ¦  «        Z• e‡ej‘        j–        eŽ¦  «        Z— e‡ej‘        j˜        eŽ¦  «        Z™ e‡ej‘        jš        eŽ¦  «        Z› e‡ej‘        jœ        e�¦  «        Z� e‡ej‘        jž        e�¦  «        ZŸ e‡ej‘        j         e�¦  «        Z¡ e‡ej‘        j¢        e�¦  «        Z£ e‡ej‘        j¤        e�¦  «        Z¥ e‡ej‘        j¦        e�¦  «        Z§e G dY„ dZe¦  «        ¦   «         Z¨e G d[„ d\e¦  «        ¦   «         Z©e G d]„ d^e¦  «        ¦   «         Zªe G d_„ d`e¦  «        ¦   «         Z«e G da„ dbe¦  «        ¦   «         Z¬e G dc„ dde¦  «        ¦   «         Z­e G de„ dfe¦  «        ¦   «         Z®e G dg„ dhe¦  «        ¦   «         Z¯e G di„ dje¦  «        ¦   «         Z°e G dk„ dle¦  «        ¦   «         Z± ee ej²        e¦  «        ¦  «         eD ]Z³ ee³e¦  «         Œe	D ]Z³ ee³e¦  «         Œe
D ]Z³ ee³e¦  «         ŒdS )mé    N)Útypes)Úparse_dtypeÚparse_shapeÚregister_number_classesÚregister_numpy_ufuncÚtrigonometric_functionsÚcomparison_functionsÚbit_twiddling_functions)ÚAttributeTemplateÚConcreteTemplateÚAbstractTemplateÚCallableTemplateÚ	signatureÚRegistry)Údim3Ú
grid_group)Ú
Conversion)Úcuda)Ú declare_device_function_templatec                   ó   — e Zd Zd„ ZdS )ÚCuda_array_declc                 ó   — d„ }|S )Nc                 ór  — t          | t          j        ¦  «        rt          | t          j        ¦  «        sd S nDt          | t          j        t          j        f¦  «        rt          d„ | D ¦   «         ¦  «        rd S nd S t          | ¦  «        }t          |¦  «        }|�|�t          j	        ||d¬¦  «        S d S d S )Nc                 óD   — g | ]}t          |t          j        ¦  «         ‘ŒS © )Ú
isinstancer   ÚIntegerLiteral)Ú.0Úss     úI/granary/edev/quantum/lib/python3.11/site-packages/numba/cuda/cudadecl.pyú
<listcomp>z:Cuda_array_decl.generic.<locals>.typer.<locals>.<listcomp>#   s8   € ð (ð (ð (Øõ ' q­%Ô*>Ñ?Ô?Ð?ð (ð (ð (ó    ÚC)ÚdtypeÚndimÚlayout)
r   r   ÚIntegerr   ÚTupleÚUniTupleÚanyr   r   ÚArray)Úshaper$   r%   Únb_dtypes       r    Útyperz&Cuda_array_decl.generic.<locals>.typer   sÛ   € õ ˜%¥¤Ñ/Ô/ð Ý! %­Ô)=Ñ>Ô>ð  Ø˜4ð å˜E¥E¤Kµ´Ð#@ÑAÔAð Ýð (ð (Ø!&ð(ñ (ô (ñ )ô )ð  à˜4ð ð �tå˜uÑ%Ô%ˆDÝ" 5Ñ)Ô)ˆHØÐ#¨Ð(8Ý”{¨¸ÀSÐIÑIÔIÐIð $Ð#Ð(8Ð(8r"   r   ©Úselfr.   s     r    ÚgenericzCuda_array_decl.generic   s   € ð	Jð 	Jð 	Jð& ˆr"   N©Ú__name__Ú
__module__Ú__qualname__r1   r   r"   r    r   r      s#   € € € € € ðð ð ð ð r"   r   c                   ó&   — e Zd Zej        j        ZdS )ÚCuda_shared_arrayN)r3   r4   r5   r   ÚsharedÚarrayÚkeyr   r"   r    r7   r7   1   s   € € € € € à
Œ+Ô
€C€C€Cr"   r7   c                   ó&   — e Zd Zej        j        ZdS )ÚCuda_local_arrayN)r3   r4   r5   r   Úlocalr9   r:   r   r"   r    r<   r<   6   s   € € € € € à
Œ*Ô
€C€C€Cr"   r<   c                   ó,   — e Zd Zej        j        Zd„ ZdS )ÚCuda_const_array_likec                 ó   — d„ }|S )Nc                 ó   — | S ©Nr   )Úndarrays    r    r.   z,Cuda_const_array_like.generic.<locals>.typer@   s   € ØˆNr"   r   r/   s     r    r1   zCuda_const_array_like.generic?   s   € ð	ð 	ð 	àˆr"   N)r3   r4   r5   r   ÚconstÚ
array_liker:   r1   r   r"   r    r?   r?   ;   s-   € € € € € à
Œ*Ô
€Cðð ð ð ð r"   r?   c                   ó>   — e Zd Zej        Z eej        ¦  «        gZ	dS )ÚCuda_threadfence_deviceN)
r3   r4   r5   r   Úthreadfencer:   r   r   ÚnoneÚcasesr   r"   r    rG   rG   E   s*   € € € € € à
Ô
€CØˆY�u”zÑ"Ô"Ð#€E€E€Er"   rG   c                   ó>   — e Zd Zej        Z eej        ¦  «        gZ	dS )ÚCuda_threadfence_blockN)
r3   r4   r5   r   Úthreadfence_blockr:   r   r   rI   rJ   r   r"   r    rL   rL   K   s*   € € € € € à
Ô
 €CØˆY�u”zÑ"Ô"Ð#€E€E€Er"   rL   c                   ó>   — e Zd Zej        Z eej        ¦  «        gZ	dS )ÚCuda_threadfence_systemN)
r3   r4   r5   r   Úthreadfence_systemr:   r   r   rI   rJ   r   r"   r    rO   rO   Q   s*   € € € € € à
Ô
!€CØˆY�u”zÑ"Ô"Ð#€E€E€Er"   rO   c                   óh   — e Zd Zej        Z eej        ¦  «         eej        ej	        ¦  «        gZ
dS )ÚCuda_syncwarpN)r3   r4   r5   r   Úsyncwarpr:   r   r   rI   Úi4rJ   r   r"   r    rR   rR   W   s;   € € € € € à
Œ-€CØˆY�u”zÑ"Ô" I I¨e¬j¸%¼(Ñ$CÔ$CÐD€E€E€Er"   rR   c                   ó>   — e Zd Zej        j        Z ee¦  «        gZ	dS )ÚCuda_cg_this_gridN)
r3   r4   r5   r   ÚcgÚ	this_gridr:   r   r   rJ   r   r"   r    rV   rV   ]   s*   € € € € € à
Œ'Ô
€CØˆY�zÑ"Ô"Ð#€E€E€Er"   rV   c                   ó>   — e Zd Z ej        ej        ¦  «        Zd„ ZdS )ÚCudaCgModuleTemplatec                 ó4   — t          j        t          ¦  «        S rB   )r   ÚFunctionrV   ©r0   Úmods     r    Úresolve_this_gridz&CudaCgModuleTemplate.resolve_this_gridg   ó   € ÝŒ~Õ/Ñ0Ô0Ð0r"   N)	r3   r4   r5   r   ÚModuler   rW   r:   r_   r   r"   r    rZ   rZ   c   s6   € € € € € à
ˆ%Œ,�t”wÑ
Ô
€Cð1ð 1ð 1ð 1ð 1r"   rZ   c                   ó   — e Zd ZdZd„ ZdS )ÚCuda_grid_group_synczGridGroup.syncc                 óB   — t          t          j        | j        ¬¦  «        S )N)Úrecvr)r   r   Úint32Úthis©r0   ÚargsÚkwss      r    r1   zCuda_grid_group_sync.genericn   s   € Ý�œ¨D¬IÐ6Ñ6Ô6Ð6r"   N©r3   r4   r5   r:   r1   r   r"   r    rc   rc   k   s(   € € € € € Ø
€Cð7ð 7ð 7ð 7ð 7r"   rc   c                   ó   — e Zd ZeZd„ ZdS )ÚGridGroup_attrsc                 ó@   — t          j        t          t          ¦  «        S rB   )r   ÚBoundFunctionrc   r   r]   s     r    Úresolve_synczGridGroup_attrs.resolve_syncv   s   € ÝÔ"Õ#7½ÑDÔDÐDr"   N)r3   r4   r5   r   r:   rp   r   r"   r    rm   rm   r   s-   € € € € € à
€CðEð Eð Eð Eð Er"   rm   c                   ó0  — e Zd Zej        Z e ej        ej	        ej
        f¦  «        ej	        ej	        ej	        ej	        ej	        ¦  «         e ej        ej        ej
        f¦  «        ej	        ej	        ej        ej	        ej	        ¦  «         e ej        ej        ej
        f¦  «        ej	        ej	        ej        ej	        ej	        ¦  «         e ej        ej        ej
        f¦  «        ej	        ej	        ej        ej	        ej	        ¦  «        gZdS )ÚCuda_shfl_sync_intrinsicN)r3   r4   r5   r   Úshfl_sync_intrinsicr:   r   r   r(   rT   Úb1Úi8Úf4Úf8rJ   r   r"   r    rr   rr   z   s  € € € € € à
Ô
"€Càˆ	�+�%”+˜uœx¨¬Ð2Ñ3Ô3Ø”(˜EœH e¤h°´¸%¼(ñ	Dô 	Dàˆ	�+�%”+˜uœx¨¬Ð2Ñ3Ô3Ø”(˜EœH e¤h°´¸%¼(ñ	Dô 	Dàˆ	�+�%”+˜uœx¨¬Ð2Ñ3Ô3Ø”(˜EœH e¤h°´¸%¼(ñ	Dô 	Dàˆ	�+�%”+˜uœx¨¬Ð2Ñ3Ô3Ø”(˜EœH e¤h°´¸%¼(ñ	Dô 	Dð	€E€E€Er"   rr   c                   óŒ   — e Zd Zej        Z e ej        ej	        ej
        f¦  «        ej	        ej	        ej
        ¦  «        gZdS )ÚCuda_vote_sync_intrinsicN)r3   r4   r5   r   Úvote_sync_intrinsicr:   r   r   r(   rT   rt   rJ   r   r"   r    ry   ry   ‰   sP   € € € € € à
Ô
"€CØˆY�{�u”{ E¤H¨e¬hÐ#7Ñ8Ô8Ø”x ¤¨5¬8ñ5ô 5ð 6€E€E€Er"   ry   c                   óø   — e Zd Zej        Z eej        ej        ej        ¦  «         eej        ej        ej	        ¦  «         eej        ej        ej
        ¦  «         eej        ej        ej        ¦  «        gZdS )ÚCuda_match_any_syncN)r3   r4   r5   r   Úmatch_any_syncr:   r   r   rT   ru   rv   rw   rJ   r   r"   r    r|   r|   �   sy   € € € € € à
Ô
€Càˆ	�%”(˜EœH e¤hÑ/Ô/Øˆ	�%”(˜EœH e¤hÑ/Ô/Øˆ	�%”(˜EœH e¤hÑ/Ô/Øˆ	�%”(˜EœH e¤hÑ/Ô/ð	€E€E€Er"   r|   c            	       ó   — e Zd Zej        Z e ej        ej	        ej
        f¦  «        ej	        ej	        ¦  «         e ej        ej	        ej
        f¦  «        ej	        ej        ¦  «         e ej        ej	        ej
        f¦  «        ej	        ej        ¦  «         e ej        ej	        ej
        f¦  «        ej	        ej        ¦  «        gZdS )ÚCuda_match_all_syncN)r3   r4   r5   r   Úmatch_all_syncr:   r   r   r(   rT   rt   ru   rv   rw   rJ   r   r"   r    r   r   ›   sÅ   € € € € € à
Ô
€Càˆ	�+�%”+˜uœx¨¬Ð2Ñ3Ô3°U´X¸u¼xÑHÔHØˆ	�+�%”+˜uœx¨¬Ð2Ñ3Ô3°U´X¸u¼xÑHÔHØˆ	�+�%”+˜uœx¨¬Ð2Ñ3Ô3°U´X¸u¼xÑHÔHØˆ	�+�%”+˜uœx¨¬Ð2Ñ3Ô3°U´X¸u¼xÑHÔHð	€E€E€Er"   r   c                   ó>   — e Zd Zej        Z eej        ¦  «        gZ	dS )ÚCuda_activemaskN)
r3   r4   r5   r   Ú
activemaskr:   r   r   Úuint32rJ   r   r"   r    r‚   r‚   ¦   s)   € € € € € à
Œ/€CØˆY�u”|Ñ$Ô$Ð%€E€E€Er"   r‚   c                   ó>   — e Zd Zej        Z eej        ¦  «        gZ	dS )ÚCuda_lanemask_ltN)
r3   r4   r5   r   Úlanemask_ltr:   r   r   r„   rJ   r   r"   r    r†   r†   ¬   s*   € € € € € à
Ô
€CØˆY�u”|Ñ$Ô$Ð%€E€E€Er"   r†   c                   ót  — e Zd ZdZej        Z eej	        ej	        ¦  «         eej
        ej
        ¦  «         eej        ej        ¦  «         eej        ej        ¦  «         eej        ej        ¦  «         eej        ej        ¦  «         eej        ej        ¦  «         eej        ej        ¦  «        gZdS )Ú	Cuda_popcz†
    Supported types from `llvm.popc`
    [here](http://docs.nvidia.com/cuda/nvvm-ir-spec/index.html#bit-manipulations-intrinics)
    N)r3   r4   r5   Ú__doc__r   Úpopcr:   r   r   Úint8Úint16rf   Úint64Úuint8Úuint16r„   Úuint64rJ   r   r"   r    r‰   r‰   ²   sÀ   € € € € € ðð ð Œ)€Càˆ	�%”*˜eœjÑ)Ô)Øˆ	�%”+˜uœ{Ñ+Ô+Øˆ	�%”+˜uœ{Ñ+Ô+Øˆ	�%”+˜uœ{Ñ+Ô+Øˆ	�%”+˜uœ{Ñ+Ô+Øˆ	�%”, ¤Ñ-Ô-Øˆ	�%”, ¤Ñ-Ô-Øˆ	�%”, ¤Ñ-Ô-ð	€E€E€Er"   r‰   c                   ó¨   — e Zd ZdZej        Z eej	        ej	        ej	        ej	        ¦  «         eej
        ej
        ej
        ej
        ¦  «        gZdS )ÚCuda_fmaz‡
    Supported types from `llvm.fma`
    [here](https://docs.nvidia.com/cuda/nvvm-ir-spec/index.html#standard-c-library-intrinics)
    N)r3   r4   r5   rŠ   r   Úfmar:   r   r   Úfloat32Úfloat64rJ   r   r"   r    r“   r“   Å   s^   € € € € € ðð ð Œ(€Càˆ	�%”- ¤°´¸u¼}ÑMÔMØˆ	�%”- ¤°´¸u¼}ÑMÔMð€E€E€Er"   r“   c                   ól   — e Zd Zej        j        Z eej	        ej	        ej	        ej	        ¦  «        gZ
dS )Ú	Cuda_hfmaN)r3   r4   r5   r   Úfp16Úhfmar:   r   r   Úfloat16rJ   r   r"   r    r˜   r˜   Ò   s9   € € € € € à
Œ)Œ.€Càˆ	�%”- ¤°´¸u¼}ÑMÔMð€E€E€Er"   r˜   c                   ót   — e Zd Zej        Z eej        ej        ¦  «         eej	        ej	        ¦  «        gZ
dS )Ú	Cuda_cbrtN)r3   r4   r5   r   Úcbrtr:   r   r   r•   r–   rJ   r   r"   r    r�   r�   Ú   sD   € € € € € ð Œ)€Càˆ	�%”- ¤Ñ/Ô/Øˆ	�%”- ¤Ñ/Ô/ð€E€E€Er"   r�   c                   ót   — e Zd Zej        Z eej        ej        ¦  «         eej	        ej	        ¦  «        gZ
dS )Ú	Cuda_brevN)r3   r4   r5   r   Úbrevr:   r   r   r„   r‘   rJ   r   r"   r    r    r    ä   sB   € € € € € à
Œ)€Càˆ	�%”, ¤Ñ-Ô-Øˆ	�%”, ¤Ñ-Ô-ð€E€E€Er"   r    c                   ót  — e Zd ZdZej        Z eej	        ej	        ¦  «         eej
        ej
        ¦  «         eej        ej        ¦  «         eej        ej        ¦  «         eej        ej        ¦  «         eej        ej        ¦  «         eej        ej        ¦  «         eej        ej        ¦  «        gZdS )ÚCuda_clzz†
    Supported types from `llvm.ctlz`
    [here](http://docs.nvidia.com/cuda/nvvm-ir-spec/index.html#bit-manipulations-intrinics)
    N)r3   r4   r5   rŠ   r   Úclzr:   r   r   rŒ   r�   rf   rŽ   r�   r�   r„   r‘   rJ   r   r"   r    r£   r£   í   sÀ   € € € € € ðð ð Œ(€Càˆ	�%”*˜eœjÑ)Ô)Øˆ	�%”+˜uœ{Ñ+Ô+Øˆ	�%”+˜uœ{Ñ+Ô+Øˆ	�%”+˜uœ{Ñ+Ô+Øˆ	�%”+˜uœ{Ñ+Ô+Øˆ	�%”, ¤Ñ-Ô-Øˆ	�%”, ¤Ñ-Ô-Øˆ	�%”, ¤Ñ-Ô-ð	€E€E€Er"   r£   c                   ót  — e Zd ZdZej        Z eej	        ej
        ¦  «         eej	        ej        ¦  «         eej	        ej        ¦  «         eej	        ej        ¦  «         eej	        ej        ¦  «         eej	        ej        ¦  «         eej	        ej	        ¦  «         eej	        ej        ¦  «        gZdS )ÚCuda_ffsz†
    Supported types from `llvm.cttz`
    [here](http://docs.nvidia.com/cuda/nvvm-ir-spec/index.html#bit-manipulations-intrinics)
    N)r3   r4   r5   rŠ   r   Úffsr:   r   r   r„   rŒ   r�   rf   rŽ   r�   r�   r‘   rJ   r   r"   r    r¦   r¦      sÀ   € € € € € ðð ð Œ(€Càˆ	�%”, ¤
Ñ+Ô+Øˆ	�%”, ¤Ñ,Ô,Øˆ	�%”, ¤Ñ,Ô,Øˆ	�%”, ¤Ñ,Ô,Øˆ	�%”, ¤Ñ,Ô,Øˆ	�%”, ¤Ñ-Ô-Øˆ	�%”, ¤Ñ-Ô-Øˆ	�%”, ¤Ñ-Ô-ð	€E€E€Er"   r¦   c                   ó"   — e Zd Zej        Zd„ ZdS )Ú	Cuda_selpc                 ó  — |rJ ‚|\  }}}t           j        t           j        t           j        t           j        t           j        t           j        t           j        t           j        f}||k    s||vrd S t          ||||¦  «        S rB   )
r   r–   r•   r�   r�   rf   r„   rŽ   r‘   r   )r0   ri   rj   ÚtestÚaÚbÚsupported_typess          r    r1   zCuda_selp.generic  sq   € ØˆˆˆØ‰
ˆˆa�õ !œ=­%¬-Ý œ;­¬Ý œ;­¬Ý œ;­¬ð6ˆð
 �Š6ˆ6�Q˜oÐ-Ð-ØˆFå˜˜D ! QÑ'Ô'Ð'r"   N)r3   r4   r5   r   Úselpr:   r1   r   r"   r    r©   r©     s*   € € € € € à
Œ)€Cð(ð (ð (ð (ð (r"   r©   c                 óL   ‡ — t            G ˆ fd„dt          ¦  «        ¦   «         }|S )Nc                   óB   •— e Zd Z” Z eej        ej        ¦  «        gZdS )ú'_genfp16_unary.<locals>.Cuda_fp16_unaryN©r3   r4   r5   r:   r   r   r›   rJ   ©Úl_keys   €r    ÚCuda_fp16_unaryr²   )  s,   ø€ € € € € àˆØ�˜5œ=¨%¬-Ñ8Ô8Ð9ˆˆˆr"   r¶   ©Úregisterr   ©rµ   r¶   s   ` r    Ú_genfp16_unaryrº   (  sJ   ø€ Ýð:ð :ð :ð :ð :ð :ð :Õ*ñ :ô :ñ „Xð:ð Ðr"   c                 ó\   ‡ — t          ‰ ¦  «         G ˆ fd„dt          ¦  «        ¦   «         }|S )Nc                   ó   •— e Zd Z” Zd„ ZdS )ú0_genfp16_unary_operator.<locals>.Cuda_fp16_unaryc                 ó¬   — |rJ ‚t          |¦  «        dk    r:|d         t          j        k    r&t          t          j        t          j        ¦  «        S d S d S )Né   r   )Úlenr   r›   r   rh   s      r    r1   z8_genfp16_unary_operator.<locals>.Cuda_fp16_unary.generic6  sK   € ØˆNˆNˆNÝ�4‰yŒy˜AŠ~ˆ~ $ q¤'­U¬]Ò":Ð":Ý ¥¤µ´Ñ>Ô>Ð>ð ˆ~Ð":Ð":r"   Nrk   r´   s   €r    r¶   r½   2  s)   ø€ € € € € àˆð	?ð 	?ð 	?ð 	?ð 	?r"   r¶   ©Úregister_globalr   r¹   s   ` r    Ú_genfp16_unary_operatorrÃ   1  sS   ø€ Ý�UÑÔð?ð ?ð ?ð ?ð ?ð ?ð ?Õ*ñ ?ô ?ñ Ôð?ð Ðr"   c                 óL   ‡ — t            G ˆ fd„dt          ¦  «        ¦   «         }|S )Nc                   óN   •— e Zd Z” Z eej        ej        ej        ¦  «        gZdS )ú)_genfp16_binary.<locals>.Cuda_fp16_binaryNr³   r´   s   €r    ÚCuda_fp16_binaryrÆ   ?  s0   ø€ € € € € àˆØ�˜5œ=¨%¬-¸¼ÑGÔGÐHˆˆˆr"   rÇ   r·   )rµ   rÇ   s   ` r    Ú_genfp16_binaryrÈ   >  sT   ø€ ÝðIð Ið Ið Ið Ið Ið IÕ+ñ Iô Iñ „XðIð Ðr"   c                   ó   — e Zd Zd„ ZdS )ÚFloatc                 óV   — |rJ ‚|\  }|t           j        k    rt          ||¦  «        S d S rB   )r   r›   r   )r0   ri   rj   Úargs       r    r1   zFloat.genericJ  s7   € Øˆˆˆà‰ˆà•%”-ÒÐÝ˜S #Ñ&Ô&Ð&ð  Ðr"   Nr2   r   r"   r    rÊ   rÊ   G  s#   € € € € € ð'ð 'ð 'ð 'ð 'r"   rÊ   c                 óL   ‡ — t            G ˆ fd„dt          ¦  «        ¦   «         }|S )Nc                   óN   •— e Zd Z” Z eej        ej        ej        ¦  «        gZdS )ú1_genfp16_binary_comparison.<locals>.Cuda_fp16_cmpN)	r3   r4   r5   r:   r   r   rt   r›   rJ   r´   s   €r    ÚCuda_fp16_cmprÏ   T  s4   ø€ € € € € àˆð ˆI�e”h ¤¨u¬}Ñ=Ô=ð
ˆˆˆr"   rÐ   r·   )rµ   rÐ   s   ` r    Ú_genfp16_binary_comparisonrÑ   S  sJ   ø€ Ýð
ð 
ð 
ð 
ð 
ð 
ð 
Õ(ñ 
ô 
ñ „Xð
ð Ðr"   c                 ó`   ‡ ‡— t          ‰ ¦  «         G ˆ ˆfd„dt          ¦  «        ¦   «         }|S )Nc                   ó   •— e Zd Z” Zˆfd„ZdS )ú1_fp16_binary_operator.<locals>.Cuda_fp16_operatorc                 ó
  •— |rJ ‚t          |¦  «        dk    ræ|d         t          j        k    s|d         t          j        k    r¼|d         t          j        k    r(| j                             |d         |d         ¦  «        }n'| j                             |d         |d         ¦  «        }|t
          j        k    s |t
          j        k    s|t
          j        k    r)t          ‰t          j        t          j        ¦  «        S d S d S d S )Né   r   r¿   )
rÀ   r   r›   ÚcontextÚcan_convertr   ÚexactÚpromoteÚsafer   )r0   ri   rj   ÚconvertibleÚrettys       €r    r1   z9_fp16_binary_operator.<locals>.Cuda_fp16_operator.genericp  sã   ø€ ØˆNˆNˆNå�4‰yŒy˜AŠ~ˆ~Ø˜!”W¥¤Ò-Ð-°°a´½E¼MÒ1IÐ1IØ˜”G�uœ}Ò,Ð,Ø"&¤,×":Ò":¸4À¼7ÀDÈÄGÑ"LÔ"L�K�Kà"&¤,×":Ò":¸4À¼7ÀDÈÄGÑ"LÔ"L�Kð  ¥:Ô#3Ò3Ð3Ø¥:Ô#5Ò5Ð5Ø¥:¤?Ò2Ð2Ý$ U­E¬M½5¼=ÑIÔIÐIð% ˆ~Ø1IÐ1Ið  3Ð2r"   Nrk   )rµ   rÝ   s   €€r    ÚCuda_fp16_operatorrÔ   l  s:   ø€ € € € € àˆð	Jð 	Jð 	Jð 	Jð 	Jð 	Jð 	Jr"   rÞ   rÁ   )rµ   rÝ   rÞ   s   `` r    Ú_fp16_binary_operatorrß   k  sd   øø€ Ý�UÑÔðJð Jð Jð Jð Jð Jð Jð JÕ-ñ Jô Jñ ÔðJð4 Ðr"   c                 ó6   — t          | t          j        ¦  «        S rB   )rß   r   rt   ©Úops    r    Ú_genfp16_comparison_operatorrã   Š  s   € Ý  ¥U¤XÑ.Ô.Ð.r"   c                 ó6   — t          | t          j        ¦  «        S rB   )rß   r   r›   rá   s    r    Ú_genfp16_binary_operatorrå   Ž  s   € Ý  ¥U¤]Ñ3Ô3Ð3r"   c                 ó|   — t          d| › �t          j        t          j        f¦  «        }t          j        |¦  «        S ©NÚ__numba_wrapper_©r   r   r›   r\   ©ÚfnameÚdecls     r    Ú_resolve_wrapped_unaryrí   ±  s;   € Ý+Ð,F¸uÐ,FÐ,FÝ,1¬MÝ-2¬]Ð,<ñ>ô >€Dõ Œ>˜$ÑÔÐr"   c                 ó’   — t          d| › �t          j        t          j        t          j        f¦  «        }t          j        |¦  «        S rç   ré   rê   s     r    Ú_resolve_wrapped_binaryrï   ¸  sA   € Ý+Ð,F¸uÐ,FÐ,FÝ,1¬MÝ-2¬]½E¼MÐ,KñMô M€Dõ Œ>˜$ÑÔÐr"   ÚhsinÚhcosÚhlogÚhlog10Úhlog2ÚhexpÚhexp10Úhexp2ÚhsqrtÚhrsqrtÚhfloorÚhceilÚhrcpÚhrintÚhtruncÚhdivc                 óP   ‡ ‡— t            G ˆ ˆfd„dt          ¦  «        ¦   «         }|S )Nc                   ó   •— e Zd Z” Zˆfd„ZdS )ú_gen.<locals>.Cuda_atomicc                 óâ   •— |rJ ‚|\  }}}|j         ‰vrd S |j        dk    r&t          |j         |t          j        |j         ¦  «        S |j        dk    rt          |j         |||j         ¦  «        S d S ©Nr¿   )r$   r%   r   r   Úintp)r0   ri   rj   ÚaryÚidxÚvalr®   s         €r    r1   z!_gen.<locals>.Cuda_atomic.generic×  sx   ø€ ØˆNˆNˆNØ ‰MˆC��càŒy Ð/Ð/Ø�àŒx˜1Š}ˆ}Ý  ¤¨Cµ´¸S¼YÑGÔGÐGØ”˜A’�Ý  ¤¨C°°c´iÑ@Ô@Ð@ð �r"   Nrk   )rµ   r®   s   €€r    ÚCuda_atomicr  Ó  s:   ø€ € € € € àˆð
	Að 
	Að 
	Að 
	Að 
	Að 
	Að 
	Ar"   r	  )r¸   r   )rµ   r®   r	  s   `` r    Ú_genr
  Ò  s[   øø€ ÝðAð Að Að Að Að Að Að AÕ&ñ Aô Añ „XðAð Ðr"   c                   ó,   — e Zd Zej        j        Zd„ ZdS )ÚCuda_atomic_compare_and_swapc                 óx   — |rJ ‚|\  }}}|j         }|t          v r|j        dk    rt          ||||¦  «        S d S d S r  )r$   Úinteger_numba_typesr%   r   )r0   ri   rj   r  Úoldr  Údtys          r    r1   z$Cuda_atomic_compare_and_swap.generic   sS   € ØˆˆˆØ‰ˆˆS�#ØŒiˆàÕ%Ð%Ð%¨#¬(°aª-¨-Ý˜S # s¨CÑ0Ô0Ð0ð &Ð%¨-¨-r"   N)r3   r4   r5   r   ÚatomicÚcompare_and_swapr:   r1   r   r"   r    r  r  ü  s-   € € € € € à
Œ+Ô
&€Cð1ð 1ð 1ð 1ð 1r"   r  c                   ó,   — e Zd Zej        j        Zd„ ZdS )ÚCuda_atomic_casc                 óÌ   — |rJ ‚|\  }}}}|j         }|t          vrd S |j        dk    rt          ||t          j        ||¦  «        S |j        dk    rt          |||||¦  «        S d S r  )r$   r  r%   r   r   r  )r0   ri   rj   r  r  r  r  r  s           r    r1   zCuda_atomic_cas.generic  s|   € ØˆˆˆØ!ÑˆˆS�#�sØŒiˆàÕ)Ð)Ð)ØˆFàŒ8�qŠ=ˆ=Ý˜S #¥u¤z°3¸Ñ<Ô<Ð<ØŒX˜Š\ˆ\Ý˜S # s¨C°Ñ5Ô5Ð5ð ˆ\r"   N)r3   r4   r5   r   r  Úcasr:   r1   r   r"   r    r  r  	  s,   € € € € € à
Œ+Œ/€Cð6ð 6ð 6ð 6ð 6r"   r  c                   óJ   — e Zd Zej        Z eej        ej	        ¦  «        gZ
dS )ÚCuda_nanosleepN)r3   r4   r5   r   Ú	nanosleepr:   r   r   Úvoidr„   rJ   r   r"   r    r  r    s-   € € € € € à
Œ.€CàˆY�u”z 5¤<Ñ0Ô0Ð1€E€E€Er"   r  c                   ó$   — e Zd ZeZd„ Zd„ Zd„ ZdS )Ú
Dim3_attrsc                 ó   — t           j        S rB   ©r   rf   r]   s     r    Ú	resolve_xzDim3_attrs.resolve_x&  ó
   € ÝŒ{Ðr"   c                 ó   — t           j        S rB   r  r]   s     r    Ú	resolve_yzDim3_attrs.resolve_y)  r   r"   c                 ó   — t           j        S rB   r  r]   s     r    Ú	resolve_zzDim3_attrs.resolve_z,  r   r"   N)r3   r4   r5   r   r:   r  r"  r$  r   r"   r    r  r  "  sF   € € € € € à
€Cðð ð ðð ð ðð ð ð ð r"   r  c                   ó>   — e Zd Z ej        ej        ¦  «        Zd„ ZdS )ÚCudaSharedModuleTemplatec                 ó4   — t          j        t          ¦  «        S rB   )r   r\   r7   r]   s     r    Úresolve_arrayz&CudaSharedModuleTemplate.resolve_array4  r`   r"   N)	r3   r4   r5   r   ra   r   r8   r:   r(  r   r"   r    r&  r&  0  s6   € € € € € à
ˆ%Œ,�t”{Ñ
#Ô
#€Cð1ð 1ð 1ð 1ð 1r"   r&  c                   ó>   — e Zd Z ej        ej        ¦  «        Zd„ ZdS )ÚCudaConstModuleTemplatec                 ó4   — t          j        t          ¦  «        S rB   )r   r\   r?   r]   s     r    Úresolve_array_likez*CudaConstModuleTemplate.resolve_array_like<  s   € ÝŒ~Õ3Ñ4Ô4Ð4r"   N)	r3   r4   r5   r   ra   r   rD   r:   r,  r   r"   r    r*  r*  8  s6   € € € € € à
ˆ%Œ,�t”zÑ
"Ô
"€Cð5ð 5ð 5ð 5ð 5r"   r*  c                   ó>   — e Zd Z ej        ej        ¦  «        Zd„ ZdS )ÚCudaLocalModuleTemplatec                 ó4   — t          j        t          ¦  «        S rB   )r   r\   r<   r]   s     r    r(  z%CudaLocalModuleTemplate.resolve_arrayD  ó   € ÝŒ~Õ.Ñ/Ô/Ð/r"   N)	r3   r4   r5   r   ra   r   r=   r:   r(  r   r"   r    r.  r.  @  s6   € € € € € à
ˆ%Œ,�t”zÑ
"Ô
"€Cð0ð 0ð 0ð 0ð 0r"   r.  c                   óŒ   — e Zd Z ej        ej        ¦  «        Zd„ Zd„ Z	d„ Z
d„ Zd„ Zd„ Zd„ Zd„ Zd	„ Zd
„ Zd„ Zd„ Zd„ Zd„ ZdS )ÚCudaAtomicTemplatec                 ó4   — t          j        t          ¦  «        S rB   )r   r\   ÚCuda_atomic_addr]   s     r    Úresolve_addzCudaAtomicTemplate.resolve_addL  ó   € ÝŒ~�oÑ.Ô.Ð.r"   c                 ó4   — t          j        t          ¦  «        S rB   )r   r\   ÚCuda_atomic_subr]   s     r    Úresolve_subzCudaAtomicTemplate.resolve_subO  r6  r"   c                 ó4   — t          j        t          ¦  «        S rB   )r   r\   ÚCuda_atomic_andr]   s     r    Úresolve_and_zCudaAtomicTemplate.resolve_and_R  r6  r"   c                 ó4   — t          j        t          ¦  «        S rB   )r   r\   ÚCuda_atomic_orr]   s     r    Úresolve_or_zCudaAtomicTemplate.resolve_or_U  ó   € ÝŒ~�nÑ-Ô-Ð-r"   c                 ó4   — t          j        t          ¦  «        S rB   )r   r\   ÚCuda_atomic_xorr]   s     r    Úresolve_xorzCudaAtomicTemplate.resolve_xorX  r6  r"   c                 ó4   — t          j        t          ¦  «        S rB   )r   r\   ÚCuda_atomic_incr]   s     r    Úresolve_inczCudaAtomicTemplate.resolve_inc[  r6  r"   c                 ó4   — t          j        t          ¦  «        S rB   )r   r\   ÚCuda_atomic_decr]   s     r    Úresolve_deczCudaAtomicTemplate.resolve_dec^  r6  r"   c                 ó4   — t          j        t          ¦  «        S rB   )r   r\   ÚCuda_atomic_exchr]   s     r    Úresolve_exchzCudaAtomicTemplate.resolve_excha  r0  r"   c                 ó4   — t          j        t          ¦  «        S rB   )r   r\   ÚCuda_atomic_maxr]   s     r    Úresolve_maxzCudaAtomicTemplate.resolve_maxd  r6  r"   c                 ó4   — t          j        t          ¦  «        S rB   )r   r\   ÚCuda_atomic_minr]   s     r    Úresolve_minzCudaAtomicTemplate.resolve_ming  r6  r"   c                 ó4   — t          j        t          ¦  «        S rB   )r   r\   ÚCuda_atomic_nanminr]   s     r    Úresolve_nanminz!CudaAtomicTemplate.resolve_nanminj  ó   € ÝŒ~Õ0Ñ1Ô1Ð1r"   c                 ó4   — t          j        t          ¦  «        S rB   )r   r\   ÚCuda_atomic_nanmaxr]   s     r    Úresolve_nanmaxz!CudaAtomicTemplate.resolve_nanmaxm  rV  r"   c                 ó4   — t          j        t          ¦  «        S rB   )r   r\   r  r]   s     r    Úresolve_compare_and_swapz+CudaAtomicTemplate.resolve_compare_and_swapp  s   € ÝŒ~Õ:Ñ;Ô;Ð;r"   c                 ó4   — t          j        t          ¦  «        S rB   )r   r\   r  r]   s     r    Úresolve_caszCudaAtomicTemplate.resolve_cass  r6  r"   N)r3   r4   r5   r   ra   r   r  r:   r5  r9  r<  r?  rC  rF  rI  rL  rO  rR  rU  rY  r[  r]  r   r"   r    r2  r2  H  sù   € € € € € à
ˆ%Œ,�t”{Ñ
#Ô
#€Cð/ð /ð /ð/ð /ð /ð/ð /ð /ð.ð .ð .ð/ð /ð /ð/ð /ð /ð/ð /ð /ð0ð 0ð 0ð/ð /ð /ð/ð /ð /ð2ð 2ð 2ð2ð 2ð 2ð<ð <ð <ð/ð /ð /ð /ð /r"   r2  c                   óì   — e Zd Z ej        ej        ¦  «        Zd„ Zd„ Z	d„ Z
d„ Zd„ Zd„ Zd„ Zd„ Zd	„ Zd
„ Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Z d„ Z!d„ Z"d„ Z#d„ Z$d„ Z%dS ) ÚCudaFp16Templatec                 ó4   — t          j        t          ¦  «        S rB   )r   r\   Ú	Cuda_haddr]   s     r    Úresolve_haddzCudaFp16Template.resolve_hadd{  ó   € ÝŒ~�iÑ(Ô(Ð(r"   c                 ó4   — t          j        t          ¦  «        S rB   )r   r\   Ú	Cuda_hsubr]   s     r    Úresolve_hsubzCudaFp16Template.resolve_hsub~  rc  r"   c                 ó4   — t          j        t          ¦  «        S rB   )r   r\   Ú	Cuda_hmulr]   s     r    Úresolve_hmulzCudaFp16Template.resolve_hmul�  rc  r"   c                 ó   — t           S rB   )Úhdiv_devicer]   s     r    Úresolve_hdivzCudaFp16Template.resolve_hdiv„  ó   € ÝÐr"   c                 ó4   — t          j        t          ¦  «        S rB   )r   r\   Ú	Cuda_hnegr]   s     r    Úresolve_hnegzCudaFp16Template.resolve_hneg‡  rc  r"   c                 ó4   — t          j        t          ¦  «        S rB   )r   r\   Ú	Cuda_habsr]   s     r    Úresolve_habszCudaFp16Template.resolve_habsŠ  rc  r"   c                 ó4   — t          j        t          ¦  «        S rB   )r   r\   r˜   r]   s     r    Úresolve_hfmazCudaFp16Template.resolve_hfma�  rc  r"   c                 ó   — t           S rB   )Úhsin_devicer]   s     r    Úresolve_hsinzCudaFp16Template.resolve_hsin�  rm  r"   c                 ó   — t           S rB   )Úhcos_devicer]   s     r    Úresolve_hcoszCudaFp16Template.resolve_hcos“  rm  r"   c                 ó   — t           S rB   )Úhlog_devicer]   s     r    Úresolve_hlogzCudaFp16Template.resolve_hlog–  rm  r"   c                 ó   — t           S rB   )Úhlog10_devicer]   s     r    Úresolve_hlog10zCudaFp16Template.resolve_hlog10™  ó   € ÝÐr"   c                 ó   — t           S rB   )Úhlog2_devicer]   s     r    Úresolve_hlog2zCudaFp16Template.resolve_hlog2œ  ó   € ÝÐr"   c                 ó   — t           S rB   )Úhexp_devicer]   s     r    Úresolve_hexpzCudaFp16Template.resolve_hexpŸ  rm  r"   c                 ó   — t           S rB   )Úhexp10_devicer]   s     r    Úresolve_hexp10zCudaFp16Template.resolve_hexp10¢  r‚  r"   c                 ó   — t           S rB   )Úhexp2_devicer]   s     r    Úresolve_hexp2zCudaFp16Template.resolve_hexp2¥  r†  r"   c                 ó   — t           S rB   )Úhfloor_devicer]   s     r    Úresolve_hfloorzCudaFp16Template.resolve_hfloor¨  r‚  r"   c                 ó   — t           S rB   )Úhceil_devicer]   s     r    Úresolve_hceilzCudaFp16Template.resolve_hceil«  r†  r"   c                 ó   — t           S rB   )Úhsqrt_devicer]   s     r    Úresolve_hsqrtzCudaFp16Template.resolve_hsqrt®  r†  r"   c                 ó   — t           S rB   )Úhrsqrt_devicer]   s     r    Úresolve_hrsqrtzCudaFp16Template.resolve_hrsqrt±  r‚  r"   c                 ó   — t           S rB   )Úhrcp_devicer]   s     r    Úresolve_hrcpzCudaFp16Template.resolve_hrcp´  rm  r"   c                 ó   — t           S rB   )Úhrint_devicer]   s     r    Úresolve_hrintzCudaFp16Template.resolve_hrint·  r†  r"   c                 ó   — t           S rB   )Úhtrunc_devicer]   s     r    Úresolve_htrunczCudaFp16Template.resolve_htruncº  r‚  r"   c                 ó4   — t          j        t          ¦  «        S rB   )r   r\   ÚCuda_heqr]   s     r    Úresolve_heqzCudaFp16Template.resolve_heq½  ó   € ÝŒ~�hÑ'Ô'Ð'r"   c                 ó4   — t          j        t          ¦  «        S rB   )r   r\   ÚCuda_hner]   s     r    Úresolve_hnezCudaFp16Template.resolve_hneÀ  r¨  r"   c                 ó4   — t          j        t          ¦  «        S rB   )r   r\   ÚCuda_hger]   s     r    Úresolve_hgezCudaFp16Template.resolve_hgeÃ  r¨  r"   c                 ó4   — t          j        t          ¦  «        S rB   )r   r\   ÚCuda_hgtr]   s     r    Úresolve_hgtzCudaFp16Template.resolve_hgtÆ  r¨  r"   c                 ó4   — t          j        t          ¦  «        S rB   )r   r\   ÚCuda_hler]   s     r    Úresolve_hlezCudaFp16Template.resolve_hleÉ  r¨  r"   c                 ó4   — t          j        t          ¦  «        S rB   )r   r\   ÚCuda_hltr]   s     r    Úresolve_hltzCudaFp16Template.resolve_hltÌ  r¨  r"   c                 ó4   — t          j        t          ¦  «        S rB   )r   r\   Ú	Cuda_hmaxr]   s     r    Úresolve_hmaxzCudaFp16Template.resolve_hmaxÏ  rc  r"   c                 ó4   — t          j        t          ¦  «        S rB   )r   r\   Ú	Cuda_hminr]   s     r    Úresolve_hminzCudaFp16Template.resolve_hminÒ  rc  r"   N)&r3   r4   r5   r   ra   r   r™   r:   rb  rf  ri  rl  rp  rs  ru  rx  r{  r~  r�  r…  r‰  rŒ  r�  r’  r•  r˜  r›  rž  r¡  r¤  r§  r«  r®  r±  r´  r·  rº  r½  r   r"   r    r_  r_  w  sé  € € € € € à
ˆ%Œ,�t”yÑ
!Ô
!€Cð)ð )ð )ð)ð )ð )ð)ð )ð )ðð ð ð)ð )ð )ð)ð )ð )ð)ð )ð )ðð ð ðð ð ðð ð ðð ð ðð ð ðð ð ðð ð ðð ð ðð ð ðð ð ðð ð ðð ð ðð ð ðð ð ðð ð ð(ð (ð (ð(ð (ð (ð(ð (ð (ð(ð (ð (ð(ð (ð (ð(ð (ð (ð)ð )ð )ð)ð )ð )ð )ð )r"   r_  c                   óÜ   — e Zd Z ej        e¦  «        Zd„ Zd„ Zd„ Z	d„ Z
d„ Zd„ Zd„ Zd„ Zd	„ Zd
„ Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Z d„ Z!d„ Z"d„ Z#dS )ÚCudaModuleTemplatec                 ó>   — t          j        t          j        ¦  «        S rB   )r   ra   r   rW   r]   s     r    Ú
resolve_cgzCudaModuleTemplate.resolve_cgÚ  s   € ÝŒ|�DœGÑ$Ô$Ð$r"   c                 ó   — t           S rB   ©r   r]   s     r    Úresolve_threadIdxz$CudaModuleTemplate.resolve_threadIdxÝ  ó   € Ýˆr"   c                 ó   — t           S rB   rÃ  r]   s     r    Úresolve_blockIdxz#CudaModuleTemplate.resolve_blockIdxà  rÅ  r"   c                 ó   — t           S rB   rÃ  r]   s     r    Úresolve_blockDimz#CudaModuleTemplate.resolve_blockDimã  rÅ  r"   c                 ó   — t           S rB   rÃ  r]   s     r    Úresolve_gridDimz"CudaModuleTemplate.resolve_gridDimæ  rÅ  r"   c                 ó   — t           j        S rB   r  r]   s     r    Úresolve_laneidz!CudaModuleTemplate.resolve_laneidé  r   r"   c                 ó>   — t          j        t          j        ¦  «        S rB   )r   ra   r   r8   r]   s     r    Úresolve_sharedz!CudaModuleTemplate.resolve_sharedì  ó   € ÝŒ|�DœKÑ(Ô(Ð(r"   c                 ó4   — t          j        t          ¦  «        S rB   )r   r\   r‰   r]   s     r    Úresolve_popczCudaModuleTemplate.resolve_popcï  rc  r"   c                 ó4   — t          j        t          ¦  «        S rB   )r   r\   r    r]   s     r    Úresolve_brevzCudaModuleTemplate.resolve_brevò  rc  r"   c                 ó4   — t          j        t          ¦  «        S rB   )r   r\   r£   r]   s     r    Úresolve_clzzCudaModuleTemplate.resolve_clzõ  r¨  r"   c                 ó4   — t          j        t          ¦  «        S rB   )r   r\   r¦   r]   s     r    Úresolve_ffszCudaModuleTemplate.resolve_ffsø  r¨  r"   c                 ó4   — t          j        t          ¦  «        S rB   )r   r\   r“   r]   s     r    Úresolve_fmazCudaModuleTemplate.resolve_fmaû  r¨  r"   c                 ó4   — t          j        t          ¦  «        S rB   )r   r\   r�   r]   s     r    Úresolve_cbrtzCudaModuleTemplate.resolve_cbrtþ  rc  r"   c                 ó4   — t          j        t          ¦  «        S rB   )r   r\   rG   r]   s     r    Úresolve_threadfencez&CudaModuleTemplate.resolve_threadfence  ó   € ÝŒ~Õ5Ñ6Ô6Ð6r"   c                 ó4   — t          j        t          ¦  «        S rB   )r   r\   rL   r]   s     r    Úresolve_threadfence_blockz,CudaModuleTemplate.resolve_threadfence_block  s   € ÝŒ~Õ4Ñ5Ô5Ð5r"   c                 ó4   — t          j        t          ¦  «        S rB   )r   r\   rO   r]   s     r    Úresolve_threadfence_systemz-CudaModuleTemplate.resolve_threadfence_system  rß  r"   c                 ó4   — t          j        t          ¦  «        S rB   )r   r\   rR   r]   s     r    Úresolve_syncwarpz#CudaModuleTemplate.resolve_syncwarp
  s   € ÝŒ~�mÑ,Ô,Ð,r"   c                 ó4   — t          j        t          ¦  «        S rB   )r   r\   rr   r]   s     r    Úresolve_shfl_sync_intrinsicz.CudaModuleTemplate.resolve_shfl_sync_intrinsic  ó   € ÝŒ~Õ6Ñ7Ô7Ð7r"   c                 ó4   — t          j        t          ¦  «        S rB   )r   r\   ry   r]   s     r    Úresolve_vote_sync_intrinsicz.CudaModuleTemplate.resolve_vote_sync_intrinsic  rè  r"   c                 ó4   — t          j        t          ¦  «        S rB   )r   r\   r|   r]   s     r    Úresolve_match_any_syncz)CudaModuleTemplate.resolve_match_any_sync  ó   € ÝŒ~Õ1Ñ2Ô2Ð2r"   c                 ó4   — t          j        t          ¦  «        S rB   )r   r\   r   r]   s     r    Úresolve_match_all_syncz)CudaModuleTemplate.resolve_match_all_sync  rí  r"   c                 ó4   — t          j        t          ¦  «        S rB   )r   r\   r‚   r]   s     r    Úresolve_activemaskz%CudaModuleTemplate.resolve_activemask  r6  r"   c                 ó4   — t          j        t          ¦  «        S rB   )r   r\   r†   r]   s     r    Úresolve_lanemask_ltz&CudaModuleTemplate.resolve_lanemask_lt  r0  r"   c                 ó4   — t          j        t          ¦  «        S rB   )r   r\   r©   r]   s     r    Úresolve_selpzCudaModuleTemplate.resolve_selp  rc  r"   c                 ó4   — t          j        t          ¦  «        S rB   )r   r\   r  r]   s     r    Úresolve_nanosleepz$CudaModuleTemplate.resolve_nanosleep"  r@  r"   c                 ó>   — t          j        t          j        ¦  «        S rB   )r   ra   r   r  r]   s     r    Úresolve_atomicz!CudaModuleTemplate.resolve_atomic%  rÐ  r"   c                 ó>   — t          j        t          j        ¦  «        S rB   )r   ra   r   r™   r]   s     r    Úresolve_fp16zCudaModuleTemplate.resolve_fp16(  s   € ÝŒ|�DœIÑ&Ô&Ð&r"   c                 ó>   — t          j        t          j        ¦  «        S rB   )r   ra   r   rD   r]   s     r    Úresolve_constz CudaModuleTemplate.resolve_const+  ó   € ÝŒ|�DœJÑ'Ô'Ð'r"   c                 ó>   — t          j        t          j        ¦  «        S rB   )r   ra   r   r=   r]   s     r    Úresolve_localz CudaModuleTemplate.resolve_local.  rþ  r"   N)$r3   r4   r5   r   ra   r   r:   rÁ  rÄ  rÇ  rÉ  rË  rÍ  rÏ  rÒ  rÔ  rÖ  rØ  rÚ  rÜ  rÞ  rá  rã  rå  rç  rê  rì  rï  rñ  ró  rõ  r÷  rù  rû  rý  r   r   r"   r    r¿  r¿  Ö  sØ  € € € € € à
ˆ%Œ,�tÑ
Ô
€Cð%ð %ð %ðð ð ðð ð ðð ð ðð ð ðð ð ð)ð )ð )ð)ð )ð )ð)ð )ð )ð(ð (ð (ð(ð (ð (ð(ð (ð (ð)ð )ð )ð7ð 7ð 7ð6ð 6ð 6ð7ð 7ð 7ð-ð -ð -ð8ð 8ð 8ð8ð 8ð 8ð3ð 3ð 3ð3ð 3ð 3ð/ð /ð /ð0ð 0ð 0ð)ð )ð )ð.ð .ð .ð)ð )ð )ð'ð 'ð 'ð(ð (ð (ð(ð (ð (ð (ð (r"   r¿  )´ÚoperatorÚ
numba.corer   Únumba.core.typing.npydeclr   r   r   r   r   r	   r
   Únumba.core.typing.templatesr   r   r   r   r   r   Únumba.cuda.typesr   r   Únumba.core.typeconvr   Únumbar   Únumba.cuda.compilerr   Úregistryr¸   Úregister_attrrÂ   r   r7   r<   r?   rG   rL   rO   rR   rV   rZ   rc   rm   rr   ry   r|   r   r‚   r†   r‰   r“   r˜   r�   r    r£   r¦   r©   rº   rÃ   rÈ   ÚfloatrÊ   rÑ   rß   rã   rå   r™   Úhaddra  ÚaddÚCuda_addÚiaddÚ	Cuda_iaddÚhsubre  ÚsubÚCuda_subÚisubÚ	Cuda_isubÚhmulrh  ÚmulÚCuda_mulÚimulÚ	Cuda_imulÚhmaxr¹  Úhminr¼  Úhnegro  ÚnegÚCuda_negÚhabsrr  ÚabsÚCuda_absÚheqr¦  ÚeqÚhnerª  ÚneÚhger­  ÚgeÚhgtr°  ÚgtÚhler³  ÚleÚhltr¶  ÚltÚtruedivÚitruedivrí   rï   rw  rz  r}  r€  r„  rˆ  r‹  rŽ  r—  rš  r‘  r”  r�  r   r£  rk  r
  r–   r•   rf   r„   rŽ   r‘   Úall_numba_typesr  Úunsigned_int_numba_typesr  r4  r8  ÚmaxrN  ÚminrQ  ÚnanmaxrX  ÚnanminrT  Úand_r;  Úor_r>  ÚxorrB  ÚincrE  ÚdecrH  ÚexchrK  r  r  r  r  r&  r*  r.  r2  r_  r¿  ra   Úfuncr   r"   r    ú<module>r>     s  ðØ €€€Ø Ð Ð Ð Ð Ð ð@ð @ð @ð @ð @ð @ð @ð @ð @ð @ð @ð @ð @ð @ð @ð @ð @ð @ð>ð >ð >ð >ð >ð >ð >ð >ð >ð >ð >ð >ð >ð >ð >ð >ð .Ð -Ð -Ð -Ð -Ð -Ð -Ð -Ø *Ð *Ð *Ð *Ð *Ð *Ø Ð Ð Ð Ð Ð Ø @Ð @Ð @Ð @Ð @Ð @àˆ8‰:Œ:€ØÔ€ØÔ&€ØÔ*€à Ð ˜Ñ (Ô (Ð (ðð ð ð ð Ð&ñ ô ð ð0 
ðð ð ð ð ˜ñ ô ñ 
„ðð 
ðð ð ð ð �ñ ô ñ 
„ðð 
ðð ð ð ð Ð,ñ ô ñ 
„ðð 
ð$ð $ð $ð $ð $Ð.ñ $ô $ñ 
„ð$ð
 
ð$ð $ð $ð $ð $Ð-ñ $ô $ñ 
„ð$ð
 
ð$ð $ð $ð $ð $Ð.ñ $ô $ñ 
„ð$ð
 
ðEð Eð Eð Eð EÐ$ñ Eô Eñ 
„ðEð
 
ð$ð $ð $ð $ð $Ð(ñ $ô $ñ 
„ð$ð
 ð1ð 1ð 1ð 1ð 1Ð,ñ 1ô 1ñ „ð1ð7ð 7ð 7ð 7ð 7Ð+ñ 7ô 7ð 7ð ðEð Eð Eð Eð EÐ'ñ Eô Eñ „ðEð 
ðð ð ð ð Ð/ñ ô ñ 
„ðð 
ð6ð 6ð 6ð 6ð 6Ð/ñ 6ô 6ñ 
„ð6ð 
ðð ð ð ð Ð*ñ ô ñ 
„ðð 
ðð ð ð ð Ð*ñ ô ñ 
„ðð 
ð&ð &ð &ð &ð &Ð&ñ &ô &ñ 
„ð&ð
 
ð&ð &ð &ð &ð &Ð'ñ &ô &ñ 
„ð&ð
 
ðð ð ð ð Ð ñ ô ñ 
„ðð$ 
ð	ð 	ð 	ð 	ð 	Ðñ 	ô 	ñ 
„ð	ð 
ðð ð ð ð Ð ñ ô ñ 
„ðð 
ðð ð ð ð Ð ñ ô ñ 
„ðð 
ðð ð ð ð Ð ñ ô ñ 
„ðð 
ðð ð ð ð Ðñ ô ñ 
„ðð$ 
ðð ð ð ð Ðñ ô ñ 
„ðð$ 
ð(ð (ð (ð (ð (Ð ñ (ô (ñ 
„ð(ð(ð ð ð
ð 
ð 
ðð ð ð €�ÑÔð'ð 'ð 'ð 'ð 'Ðñ 'ô 'ñ Ôð'ðð ð ð0ð ð ð>/ð /ð /ð4ð 4ð 4ð ˆO˜DœIœNÑ+Ô+€	Ø#Ð# H¤LÑ1Ô1€Ø$Ð$ X¤]Ñ3Ô3€	ØˆO˜DœIœNÑ+Ô+€	Ø#Ð# H¤LÑ1Ô1€Ø$Ð$ X¤]Ñ3Ô3€	ØˆO˜DœIœNÑ+Ô+€	Ø#Ð# H¤LÑ1Ô1€Ø$Ð$ X¤]Ñ3Ô3€	ØˆO˜DœIœNÑ+Ô+€	ØˆO˜DœIœNÑ+Ô+€	ØˆN˜4œ9œ>Ñ*Ô*€	Ø"Ð" 8¤<Ñ0Ô0€ØˆN˜4œ9œ>Ñ*Ô*€	Ø"Ð" 3Ñ'Ô'€Ø%Ð% d¤i¤mÑ4Ô4€Ø Ð ˜Xœ[Ñ )Ô )Ð )Ø%Ð% d¤i¤mÑ4Ô4€Ø Ð ˜Xœ[Ñ )Ô )Ð )Ø%Ð% d¤i¤mÑ4Ô4€Ø Ð ˜Xœ[Ñ )Ô )Ð )Ø%Ð% d¤i¤mÑ4Ô4€Ø Ð ˜Xœ[Ñ )Ô )Ð )Ø%Ð% d¤i¤mÑ4Ô4€Ø Ð ˜Xœ[Ñ )Ô )Ð )Ø%Ð% d¤i¤mÑ4Ô4€Ø Ð ˜Xœ[Ñ )Ô )Ð )Ø Ð ˜Ô)Ñ *Ô *Ð *Ø Ð ˜Ô*Ñ +Ô +Ð +ð ð  ð  ð ð  ð  ð %Ð$ VÑ,Ô,€Ø$Ð$ VÑ,Ô,€Ø$Ð$ VÑ,Ô,€Ø&Ð& xÑ0Ô0€Ø%Ð% gÑ.Ô.€Ø$Ð$ VÑ,Ô,€Ø&Ð& xÑ0Ô0€Ø%Ð% gÑ.Ô.€Ø%Ð% gÑ.Ô.€Ø&Ð& xÑ0Ô0€Ø&Ð& xÑ0Ô0€Ø%Ð% gÑ.Ô.€Ø$Ð$ VÑ,Ô,€Ø%Ð% gÑ.Ô.€Ø&Ð& xÑ0Ô0€Ø%Ð% fÑ-Ô-€ðð ð ð& ”= %¤-Ø”; ¤Ø”; ¤ð.€ð ”{ E¤LØ”{ E¤Lð2Ð ð "œL¨%¬,Ð7Ð à�$�t”{”¨Ñ8Ô8€Ø�$�t”{”¨Ñ8Ô8€Ø�$�t”{”¨Ñ8Ô8€Ø�$�t”{”¨Ñ8Ô8€Ø�T˜$œ+Ô,¨oÑ>Ô>Ð Ø�T˜$œ+Ô,¨oÑ>Ô>Ð Ø�$�t”{Ô'Ð)<Ñ=Ô=€Ø��d”k”oÐ':Ñ;Ô;€Ø�$�t”{”Ð(;Ñ<Ô<€Ø�$�t”{”Ð(@ÑAÔA€Ø�$�t”{”Ð(@ÑAÔA€Ø�4˜œÔ(Ð*=Ñ>Ô>Ð ð 
ð	1ð 	1ð 	1ð 	1ð 	1Ð#3ñ 	1ô 	1ñ 
„ð	1ð 
ð6ð 6ð 6ð 6ð 6Ð&ñ 6ô 6ñ 
„ð6ð" 
ð2ð 2ð 2ð 2ð 2Ð%ñ 2ô 2ñ 
„ð2ð ð
ð 
ð 
ð 
ð 
Ð"ñ 
ô 
ñ „ð
ð ð1ð 1ð 1ð 1ð 1Ð0ñ 1ô 1ñ „ð1ð ð5ð 5ð 5ð 5ð 5Ð/ñ 5ô 5ñ „ð5ð ð0ð 0ð 0ð 0ð 0Ð/ñ 0ô 0ñ „ð0ð ð+/ð +/ð +/ð +/ð +/Ð*ñ +/ô +/ñ „ð+/ð\ ð[)ð [)ð [)ð [)ð [)Ð(ñ [)ô [)ñ „ð[)ð| ðX(ð X(ð X(ð X(ð X(Ð*ñ X(ô X(ñ „ðX(ðv €��l�e”l 4Ñ(Ô(Ñ )Ô )Ð )ð
 $ð 0ð 0€DØÐ˜˜Ñ/Ô/Ð/Ð/à ð 0ð 0€DØÐ˜˜Ñ/Ô/Ð/Ð/à#ð 0ð 0€DØÐ˜˜Ñ/Ô/Ð/Ð/ð0ð 0r"   