§
    1@Äj¾]  ã                   ó¸  — 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
mZ d dlmZmZmZmZmZmZ d dl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(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¦  «        ¦   «         Z0e G d-„ d.e¦  «        ¦   «         Z1e G d/„ d0e¦  «        ¦   «         Z2e G d1„ d2e¦  «        ¦   «         Z3e G d3„ d4e¦  «        ¦   «         Z4d5„ Z5d6„ Z6d7„ Z7 ee8¦  «         G d8„ d9e¦  «        ¦   «         Z9d:„ Z:d;„ Z;d<„ Z<d=„ Z= e7ej>        j?        ¦  «        Z@ e=e jA        ¦  «        ZB e=e jC        ¦  «        ZD e7ej>        jE        ¦  «        ZF e=e jG        ¦  «        ZH e=e jI        ¦  «        ZJ e7ej>        jK        ¦  «        ZL e=e jM        ¦  «        ZN e=e jO        ¦  «        ZP e7ej>        jQ        ¦  «        ZR e7ej>        jS        ¦  «        ZT e5ej>        jU        ¦  «        ZV e6e jW        ¦  «        ZX e5ej>        jY        ¦  «        ZZ e6e[¦  «        Z\ e:ej>        j]        ¦  «        Z^ e<e j_        ¦  «          e:ej>        j`        ¦  «        Za e<e jb        ¦  «          e:ej>        jc        ¦  «        Zd e<e je        ¦  «          e:ej>        jf        ¦  «        Zg e<e jh        ¦  «          e:ej>        ji        ¦  «        Zj e<e jk        ¦  «          e:ej>        jl        ¦  «        Zm e<e jn        ¦  «          e=e jo        ¦  «          e=e jp        ¦  «         d>„ Zqd?„ Zr eqd@¦  «        Zs eqdA¦  «        Zt eqdB¦  «        Zu eqdC¦  «        Zv eqdD¦  «        Zw eqdE¦  «        Zx eqdF¦  «        Zy eqdG¦  «        Zz eqdH¦  «        Z{ eqdI¦  «        Z| eqdJ¦  «        Z} eqdK¦  «        Z~ eqdL¦  «        Z eqdM¦  «        Z€ eqdN¦  «        Z� erdO¦  «        Z‚dP„ 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�        jA        eŠ¦  «        ZŽ eƒej�        jG        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 dQ„ dRe¦  «        ¦   «         Z¤e G dS„ dTe¦  «        ¦   «         Z¥e G dU„ dVe¦  «        ¦   «         Z¦e G dW„ dXe¦  «        ¦   «         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e ej®        e¦  «        ¦  «         eD ]Z¯ ee¯e¦  «         Œe	D ]Z¯ ee¯e¦  «         ŒeD ]Z¯ ee¯e¦  «         Œe
D ]Z¯e¯dev r ee¯e¦  «         ŒdS )fé    N)Útypes)Úparse_dtypeÚparse_shapeÚregister_number_classesÚregister_numpy_ufuncÚtrigonometric_functionsÚcomparison_functionsÚmath_operationsÚbit_twiddling_functions)ÚAttributeTemplateÚConcreteTemplateÚAbstractTemplateÚCallableTemplateÚ	signatureÚRegistry©Údim3)Ú
Conversion)Úcuda)Ú declare_device_function_templatec                   ó   — 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     úDC:\Projects\Historical\venv\Lib\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__r2   r   r#   r!   r   r      s#   € € € € € ðð ð ð ð r#   r   c                   ó&   — e Zd Zej        j        ZdS )ÚCuda_shared_arrayN)r4   r5   r6   r   ÚsharedÚarrayÚkeyr   r#   r!   r8   r8   2   s   € € € € € à
Œ+Ô
€C€C€Cr#   r8   c                   ó&   — e Zd Zej        j        ZdS )ÚCuda_local_arrayN)r4   r5   r6   r   Úlocalr:   r;   r   r#   r!   r=   r=   7   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>.typerA   s   € ØˆNr#   r   r0   s     r!   r2   zCuda_const_array_like.generic@   s   € ð	ð 	ð 	àˆr#   N)r4   r5   r6   r   ÚconstÚ
array_liker;   r2   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)
r4   r5   r6   r   Úthreadfencer;   r   r   ÚnoneÚcasesr   r#   r!   rH   rH   F   s*   € € € € € à
Ô
€CØˆY�u”zÑ"Ô"Ð#€E€E€Er#   rH   c                   ó>   — e Zd Zej        Z eej        ¦  «        gZ	dS )ÚCuda_threadfence_blockN)
r4   r5   r6   r   Úthreadfence_blockr;   r   r   rJ   rK   r   r#   r!   rM   rM   L   s*   € € € € € à
Ô
 €CØˆY�u”zÑ"Ô"Ð#€E€E€Er#   rM   c                   ó>   — e Zd Zej        Z eej        ¦  «        gZ	dS )ÚCuda_threadfence_systemN)
r4   r5   r6   r   Úthreadfence_systemr;   r   r   rJ   rK   r   r#   r!   rP   rP   R   s*   € € € € € à
Ô
!€CØˆY�u”zÑ"Ô"Ð#€E€E€Er#   rP   c                   óh   — e Zd Zej        Z eej        ¦  «         eej        ej	        ¦  «        gZ
dS )ÚCuda_syncwarpN)r4   r5   r6   r   Úsyncwarpr;   r   r   rJ   Úi4rK   r   r#   r!   rS   rS   X   s;   € € € € € à
Œ-€CØˆY�u”zÑ"Ô" I I¨e¬j¸%¼(Ñ$CÔ$CÐD€E€E€Er#   rS   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)r4   r5   r6   r   Úshfl_sync_intrinsicr;   r   r   r)   rU   Úb1Úi8Úf4Úf8rK   r   r#   r!   rW   rW   ^   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#   rW   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)r4   r5   r6   r   Úvote_sync_intrinsicr;   r   r   r)   rU   rY   rK   r   r#   r!   r^   r^   m   sP   € € € € € à
Ô
"€CØˆY�{�u”{ E¤H¨e¬hÐ#7Ñ8Ô8Ø”x ¤¨5¬8ñ5ô 5ð 6€E€E€Er#   r^   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)r4   r5   r6   r   Úmatch_any_syncr;   r   r   rU   rZ   r[   r\   rK   r   r#   r!   ra   ra   t   sy   € € € € € à
Ô
€Càˆ	�%”(˜EœH e¤hÑ/Ô/Øˆ	�%”(˜EœH e¤hÑ/Ô/Øˆ	�%”(˜EœH e¤hÑ/Ô/Øˆ	�%”(˜EœH e¤hÑ/Ô/ð	€E€E€Er#   ra   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)r4   r5   r6   r   Úmatch_all_syncr;   r   r   r)   rU   rY   rZ   r[   r\   rK   r   r#   r!   rd   rd      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#   rd   c                   ó>   — e Zd Zej        Z eej        ¦  «        gZ	dS )ÚCuda_activemaskN)
r4   r5   r6   r   Ú
activemaskr;   r   r   Úuint32rK   r   r#   r!   rg   rg   Š   s)   € € € € € à
Œ/€CØˆY�u”|Ñ$Ô$Ð%€E€E€Er#   rg   c                   ó>   — e Zd Zej        Z eej        ¦  «        gZ	dS )ÚCuda_lanemask_ltN)
r4   r5   r6   r   Úlanemask_ltr;   r   r   ri   rK   r   r#   r!   rk   rk   �   s*   € € € € € à
Ô
€CØˆY�u”|Ñ$Ô$Ð%€E€E€Er#   rk   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)r4   r5   r6   Ú__doc__r   Úpopcr;   r   r   Úint8Úint16Úint32Úint64Úuint8Úuint16ri   Úuint64rK   r   r#   r!   rn   rn   –   sÀ   € € € € € ðð ð Œ)€Càˆ	�%”*˜eœjÑ)Ô)Øˆ	�%”+˜uœ{Ñ+Ô+Øˆ	�%”+˜uœ{Ñ+Ô+Øˆ	�%”+˜uœ{Ñ+Ô+Øˆ	�%”+˜uœ{Ñ+Ô+Øˆ	�%”, ¤Ñ-Ô-Øˆ	�%”, ¤Ñ-Ô-Øˆ	�%”, ¤Ñ-Ô-ð	€E€E€Er#   rn   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)r4   r5   r6   ro   r   Úfmar;   r   r   Úfloat32Úfloat64rK   r   r#   r!   ry   ry   ©   s^   € € € € € ðð ð Œ(€Càˆ	�%”- ¤°´¸u¼}ÑMÔMØˆ	�%”- ¤°´¸u¼}ÑMÔMð€E€E€Er#   ry   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)r4   r5   r6   r   Úfp16Úhfmar;   r   r   Úfloat16rK   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)r4   r5   r6   r   Úcbrtr;   r   r   r{   r|   rK   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)r4   r5   r6   r   Úbrevr;   r   r   ri   rw   rK   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)r4   r5   r6   ro   r   Úclzr;   r   r   rq   rr   rs   rt   ru   rv   ri   rw   rK   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)r4   r5   r6   ro   r   Úffsr;   r   r   ri   rq   rr   rs   rt   ru   rv   rw   rK   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 rC   )
r   r|   r{   rr   rv   rs   ri   rt   rw   r   )r1   ÚargsÚkwsÚtestÚaÚbÚsupported_typess          r!   r2   zCuda_selp.genericû   sq   € ØˆˆˆwØ‰
ˆˆa�õ !œ=­%¬-Ý œ;­¬Ý œ;­¬Ý œ;­¬ð6ˆð
 �Š6ˆ6�Q˜oÐ-Ð-ØˆFå˜˜D ! QÑ'Ô'Ð'r#   N)r4   r5   r6   r   Úselpr;   r2   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©r4   r5   r6   r;   r   r   r�   rK   ©Ú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   )r1   r‘   r’   s      r!   r2   z8_genfp16_unary_operator.<locals>.Cuda_fp16_unary.generic  sK   € ØˆNˆN�7Ý�4‰yŒy˜AŠ~ˆ~ $ q¤'­U¬]Ò":Ð":Ý ¥¤µ´Ñ>Ô>Ð>ð ˆ~Ð":Ð":r#   N©r4   r5   r6   r;   r2   rœ   s   €r!   rž   r¥     s)   ø€ € € € € àˆð	?ð 	?ð 	?ð 	?ð 	?r#   rž   ©Úregister_globalr   r¡   s   ` r!   Ú_genfp16_unary_operatorr¬     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 rC   )r   r�   r   )r1   r‘   r’   Úargs       r!   r2   zFloat.generic.  s7   € Øˆˆˆwà‰ˆà•%”-ÒÐÝ˜S #Ñ&Ô&Ð&ð  Ðr#   Nr3   r   r#   r!   r³   r³   +  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)	r4   r5   r6   r;   r   r   rY   r�   rK   rœ   s   €r!   ÚCuda_fp16_cmpr¸   8  s4   ø€ € € € € àˆð ˆI�e”h ¤¨u¬}Ñ=Ô=ð
ˆˆˆr#   r¹   rŸ   )r�   r¹   s   ` r!   Ú_genfp16_binary_comparisonrº   7  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   )r1   r‘   r’   ÚconvertibleÚrettys       €r!   r2   z9_fp16_binary_operator.<locals>.Cuda_fp16_operator.genericT  sã   ø€ ØˆNˆN�7å�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#   Nr©   )r�   rÆ   s   €€r!   ÚCuda_fp16_operatorr½   P  s:   ø€ € € € € àˆð	Jð 	Jð 	Jð 	Jð 	Jð 	Jð 	Jr#   rÇ   rª   )r�   rÆ   rÇ   s   `` r!   Ú_fp16_binary_operatorrÈ   O  sd   øø€ Ý�UÑÔðJð Jð Jð Jð Jð Jð Jð JÕ-ñ Jô Jñ ÔðJð4 Ðr#   c                 ó6   — t          | t          j        ¦  «        S rC   )rÈ   r   rY   ©Úops    r!   Ú_genfp16_comparison_operatorrÌ   n  s   € Ý  ¥U¤XÑ.Ô.Ð.r#   c                 ó6   — t          | t          j        ¦  «        S rC   )rÈ   r   r�   rÊ   s    r!   Ú_genfp16_binary_operatorrÎ   r  s   € Ý  ¥U¤]Ñ3Ô3Ð3r#   c                 ó|   — t          d| › �t          j        t          j        f¦  «        }t          j        |¦  «        S ©NÚ__numba_wrapper_©r   r   r�   ÚFunction©Ú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)r1   r‘   r’   ÚaryÚidxÚvalr–   s         €r!   r2   z!_gen.<locals>.Cuda_atomic.generic»  sx   ø€ ØˆNˆN�7Ø ‰MˆC��càŒy Ð/Ð/Ø�àŒx˜1Š}ˆ}Ý  ¤¨Cµ´¸S¼YÑGÔGÐGØ”˜A’�Ý  ¤¨C°°c´iÑ@Ô@Ð@ð �r#   Nr©   )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   )r1   r‘   r’   rð   Úoldrò   Údtys          r!   r2   z$Cuda_atomic_compare_and_swap.genericä  sS   € ØˆˆˆwØ‰ˆˆS�#ØŒiˆàÕ%Ð%Ð%¨#¬(°aª-¨-Ý˜S # s¨CÑ0Ô0Ð0ð &Ð%¨-¨-r#   N)r4   r5   r6   r   ÚatomicÚcompare_and_swapr;   r2   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ï   )r1   r‘   r’   rð   rñ   rù   rò   rú   s           r!   r2   zCuda_atomic_cas.genericñ  s|   € ØˆˆˆwØ!ÑˆˆS�#�sØŒiˆàÕ)Ð)Ð)ØˆFàŒ8�qŠ=ˆ=Ý˜S #¥u¤z°3¸Ñ<Ô<Ð<ØŒX˜Š\ˆ\Ý˜S # s¨C°Ñ5Ô5Ð5ð ˆ\r#   N)r4   r5   r6   r   rû   Úcasr;   r2   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)r4   r5   r6   r   Ú	nanosleepr;   r   r   Úvoidri   rK   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 rC   ©r   rs   ©r1   Úmods     r!   Ú	resolve_xzDim3_attrs.resolve_x
  ó
   € ÝŒ{Ðr#   c                 ó   — t           j        S rC   r  r	  s     r!   Ú	resolve_yzDim3_attrs.resolve_y  r  r#   c                 ó   — t           j        S rC   r  r	  s     r!   Ú	resolve_zzDim3_attrs.resolve_z  r  r#   N)r4   r5   r6   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 rC   )r   rÓ   r8   r	  s     r!   Úresolve_arrayz&CudaSharedModuleTemplate.resolve_array  s   € ÝŒ~Õ/Ñ0Ô0Ð0r#   N)	r4   r5   r6   r   ÚModuler   r9   r;   r  r   r#   r!   r  r    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 rC   )r   rÓ   r@   r	  s     r!   Úresolve_array_likez*CudaConstModuleTemplate.resolve_array_like   s   € ÝŒ~Õ3Ñ4Ô4Ð4r#   N)	r4   r5   r6   r   r  r   rE   r;   r  r   r#   r!   r  r    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 rC   )r   rÓ   r=   r	  s     r!   r  z%CudaLocalModuleTemplate.resolve_array(  ó   € ÝŒ~Õ.Ñ/Ô/Ð/r#   N)	r4   r5   r6   r   r  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 rC   )r   rÓ   ÚCuda_atomic_addr	  s     r!   Úresolve_addzCudaAtomicTemplate.resolve_add0  ó   € ÝŒ~�oÑ.Ô.Ð.r#   c                 ó4   — t          j        t          ¦  «        S rC   )r   rÓ   ÚCuda_atomic_subr	  s     r!   Úresolve_subzCudaAtomicTemplate.resolve_sub3  r#  r#   c                 ó4   — t          j        t          ¦  «        S rC   )r   rÓ   ÚCuda_atomic_andr	  s     r!   Úresolve_and_zCudaAtomicTemplate.resolve_and_6  r#  r#   c                 ó4   — t          j        t          ¦  «        S rC   )r   rÓ   ÚCuda_atomic_orr	  s     r!   Úresolve_or_zCudaAtomicTemplate.resolve_or_9  ó   € ÝŒ~�nÑ-Ô-Ð-r#   c                 ó4   — t          j        t          ¦  «        S rC   )r   rÓ   ÚCuda_atomic_xorr	  s     r!   Úresolve_xorzCudaAtomicTemplate.resolve_xor<  r#  r#   c                 ó4   — t          j        t          ¦  «        S rC   )r   rÓ   ÚCuda_atomic_incr	  s     r!   Úresolve_inczCudaAtomicTemplate.resolve_inc?  r#  r#   c                 ó4   — t          j        t          ¦  «        S rC   )r   rÓ   ÚCuda_atomic_decr	  s     r!   Úresolve_deczCudaAtomicTemplate.resolve_decB  r#  r#   c                 ó4   — t          j        t          ¦  «        S rC   )r   rÓ   ÚCuda_atomic_exchr	  s     r!   Úresolve_exchzCudaAtomicTemplate.resolve_exchE  r  r#   c                 ó4   — t          j        t          ¦  «        S rC   )r   rÓ   ÚCuda_atomic_maxr	  s     r!   Úresolve_maxzCudaAtomicTemplate.resolve_maxH  r#  r#   c                 ó4   — t          j        t          ¦  «        S rC   )r   rÓ   ÚCuda_atomic_minr	  s     r!   Úresolve_minzCudaAtomicTemplate.resolve_minK  r#  r#   c                 ó4   — t          j        t          ¦  «        S rC   )r   rÓ   ÚCuda_atomic_nanminr	  s     r!   Úresolve_nanminz!CudaAtomicTemplate.resolve_nanminN  ó   € ÝŒ~Õ0Ñ1Ô1Ð1r#   c                 ó4   — t          j        t          ¦  «        S rC   )r   rÓ   ÚCuda_atomic_nanmaxr	  s     r!   Úresolve_nanmaxz!CudaAtomicTemplate.resolve_nanmaxQ  rC  r#   c                 ó4   — t          j        t          ¦  «        S rC   )r   rÓ   rö   r	  s     r!   Úresolve_compare_and_swapz+CudaAtomicTemplate.resolve_compare_and_swapT  s   € ÝŒ~Õ:Ñ;Ô;Ð;r#   c                 ó4   — t          j        t          ¦  «        S rC   )r   rÓ   rþ   r	  s     r!   Úresolve_caszCudaAtomicTemplate.resolve_casW  r#  r#   N)r4   r5   r6   r   r  r   rû   r;   r"  r&  r)  r,  r0  r3  r6  r9  r<  r?  rB  rF  rH  rJ  r   r#   r!   r  r  ,  sù   € € € € € à
ˆ%Œ,�t”{Ñ
#Ô
#€Cð/ð /ð /ð/ð /ð /ð/ð /ð /ð.ð .ð .ð/ð /ð /ð/ð /ð /ð/ð /ð /ð0ð 0ð 0ð/ð /ð /ð/ð /ð /ð2ð 2ð 2ð2ð 2ð 2ð<ð <ð <ð/ð /ð /ð /ð /r#   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„ 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 rC   )r   rÓ   Ú	Cuda_haddr	  s     r!   Úresolve_haddzCudaFp16Template.resolve_hadd_  ó   € ÝŒ~�iÑ(Ô(Ð(r#   c                 ó4   — t          j        t          ¦  «        S rC   )r   rÓ   Ú	Cuda_hsubr	  s     r!   Úresolve_hsubzCudaFp16Template.resolve_hsubb  rP  r#   c                 ó4   — t          j        t          ¦  «        S rC   )r   rÓ   Ú	Cuda_hmulr	  s     r!   Úresolve_hmulzCudaFp16Template.resolve_hmule  rP  r#   c                 ó   — t           S rC   )Úhdiv_devicer	  s     r!   Úresolve_hdivzCudaFp16Template.resolve_hdivh  ó   € ÝÐr#   c                 ó4   — t          j        t          ¦  «        S rC   )r   rÓ   Ú	Cuda_hnegr	  s     r!   Úresolve_hnegzCudaFp16Template.resolve_hnegk  rP  r#   c                 ó4   — t          j        t          ¦  «        S rC   )r   rÓ   Ú	Cuda_habsr	  s     r!   Úresolve_habszCudaFp16Template.resolve_habsn  rP  r#   c                 ó4   — t          j        t          ¦  «        S rC   )r   rÓ   r~   r	  s     r!   Úresolve_hfmazCudaFp16Template.resolve_hfmaq  rP  r#   c                 ó   — t           S rC   )Úhsin_devicer	  s     r!   Úresolve_hsinzCudaFp16Template.resolve_hsint  rZ  r#   c                 ó   — t           S rC   )Úhcos_devicer	  s     r!   Úresolve_hcoszCudaFp16Template.resolve_hcosw  rZ  r#   c                 ó   — t           S rC   )Úhlog_devicer	  s     r!   Úresolve_hlogzCudaFp16Template.resolve_hlogz  rZ  r#   c                 ó   — t           S rC   )Úhlog10_devicer	  s     r!   Úresolve_hlog10zCudaFp16Template.resolve_hlog10}  ó   € ÝÐr#   c                 ó   — t           S rC   )Úhlog2_devicer	  s     r!   Úresolve_hlog2zCudaFp16Template.resolve_hlog2€  ó   € ÝÐr#   c                 ó   — t           S rC   )Úhexp_devicer	  s     r!   Úresolve_hexpzCudaFp16Template.resolve_hexpƒ  rZ  r#   c                 ó   — t           S rC   )Úhexp10_devicer	  s     r!   Úresolve_hexp10zCudaFp16Template.resolve_hexp10†  ro  r#   c                 ó   — t           S rC   )Úhexp2_devicer	  s     r!   Úresolve_hexp2zCudaFp16Template.resolve_hexp2‰  rs  r#   c                 ó   — t           S rC   )Úhfloor_devicer	  s     r!   Úresolve_hfloorzCudaFp16Template.resolve_hfloorŒ  ro  r#   c                 ó   — t           S rC   )Úhceil_devicer	  s     r!   Úresolve_hceilzCudaFp16Template.resolve_hceil�  rs  r#   c                 ó   — t           S rC   )Úhsqrt_devicer	  s     r!   Úresolve_hsqrtzCudaFp16Template.resolve_hsqrt’  rs  r#   c                 ó   — t           S rC   )Úhrsqrt_devicer	  s     r!   Úresolve_hrsqrtzCudaFp16Template.resolve_hrsqrt•  ro  r#   c                 ó   — t           S rC   )Úhrcp_devicer	  s     r!   Úresolve_hrcpzCudaFp16Template.resolve_hrcp˜  rZ  r#   c                 ó   — t           S rC   )Úhrint_devicer	  s     r!   Úresolve_hrintzCudaFp16Template.resolve_hrint›  rs  r#   c                 ó   — t           S rC   )Úhtrunc_devicer	  s     r!   Úresolve_htrunczCudaFp16Template.resolve_htruncž  ro  r#   c                 ó4   — t          j        t          ¦  «        S rC   )r   rÓ   ÚCuda_heqr	  s     r!   Úresolve_heqzCudaFp16Template.resolve_heq¡  ó   € ÝŒ~�hÑ'Ô'Ð'r#   c                 ó4   — t          j        t          ¦  «        S rC   )r   rÓ   ÚCuda_hner	  s     r!   Úresolve_hnezCudaFp16Template.resolve_hne¤  r•  r#   c                 ó4   — t          j        t          ¦  «        S rC   )r   rÓ   ÚCuda_hger	  s     r!   Úresolve_hgezCudaFp16Template.resolve_hge§  r•  r#   c                 ó4   — t          j        t          ¦  «        S rC   )r   rÓ   ÚCuda_hgtr	  s     r!   Úresolve_hgtzCudaFp16Template.resolve_hgtª  r•  r#   c                 ó4   — t          j        t          ¦  «        S rC   )r   rÓ   ÚCuda_hler	  s     r!   Úresolve_hlezCudaFp16Template.resolve_hle­  r•  r#   c                 ó4   — t          j        t          ¦  «        S rC   )r   rÓ   ÚCuda_hltr	  s     r!   Úresolve_hltzCudaFp16Template.resolve_hlt°  r•  r#   c                 ó4   — t          j        t          ¦  «        S rC   )r   rÓ   Ú	Cuda_hmaxr	  s     r!   Úresolve_hmaxzCudaFp16Template.resolve_hmax³  rP  r#   c                 ó4   — t          j        t          ¦  «        S rC   )r   rÓ   Ú	Cuda_hminr	  s     r!   Úresolve_hminzCudaFp16Template.resolve_hmin¶  rP  r#   N)&r4   r5   r6   r   r  r   r   r;   rO  rS  rV  rY  r]  r`  rb  re  rh  rk  rn  rr  rv  ry  r|  r  r‚  r…  rˆ  r‹  rŽ  r‘  r”  r˜  r›  rž  r¡  r¤  r§  rª  r   r#   r!   rL  rL  [  sé  € € € € € à
ˆ%Œ,�t”yÑ
!Ô
!€Cð)ð )ð )ð)ð )ð )ð)ð )ð )ðð ð ð)ð )ð )ð)ð )ð )ð)ð )ð )ðð ð ðð ð ðð ð ðð ð ðð ð ðð ð ðð ð ðð ð ðð ð ðð ð ðð ð ðð ð ðð ð ðð ð ðð ð ð(ð (ð (ð(ð (ð (ð(ð (ð (ð(ð (ð (ð(ð (ð (ð(ð (ð (ð)ð )ð )ð)ð )ð )ð )ð )r#   rL  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 rC   )r   r  r   Úcgr	  s     r!   Ú
resolve_cgzCudaModuleTemplate.resolve_cg¾  s   € ÝŒ|�DœGÑ$Ô$Ð$r#   c                 ó   — t           S rC   r   r	  s     r!   Úresolve_threadIdxz$CudaModuleTemplate.resolve_threadIdxÁ  ó   € Ýˆr#   c                 ó   — t           S rC   r   r	  s     r!   Úresolve_blockIdxz#CudaModuleTemplate.resolve_blockIdxÄ  r²  r#   c                 ó   — t           S rC   r   r	  s     r!   Úresolve_blockDimz#CudaModuleTemplate.resolve_blockDimÇ  r²  r#   c                 ó   — t           S rC   r   r	  s     r!   Úresolve_gridDimz"CudaModuleTemplate.resolve_gridDimÊ  r²  r#   c                 ó   — t           j        S rC   r  r	  s     r!   Úresolve_laneidz!CudaModuleTemplate.resolve_laneidÍ  r  r#   c                 ó>   — t          j        t          j        ¦  «        S rC   )r   r  r   r9   r	  s     r!   Úresolve_sharedz!CudaModuleTemplate.resolve_sharedÐ  ó   € ÝŒ|�DœKÑ(Ô(Ð(r#   c                 ó4   — t          j        t          ¦  «        S rC   )r   rÓ   rn   r	  s     r!   Úresolve_popczCudaModuleTemplate.resolve_popcÓ  rP  r#   c                 ó4   — t          j        t          ¦  «        S rC   )r   rÓ   r†   r	  s     r!   Úresolve_brevzCudaModuleTemplate.resolve_brevÖ  rP  r#   c                 ó4   — t          j        t          ¦  «        S rC   )r   rÓ   r‰   r	  s     r!   Úresolve_clzzCudaModuleTemplate.resolve_clzÙ  r•  r#   c                 ó4   — t          j        t          ¦  «        S rC   )r   rÓ   rŒ   r	  s     r!   Úresolve_ffszCudaModuleTemplate.resolve_ffsÜ  r•  r#   c                 ó4   — t          j        t          ¦  «        S rC   )r   rÓ   ry   r	  s     r!   Úresolve_fmazCudaModuleTemplate.resolve_fmaß  r•  r#   c                 ó4   — t          j        t          ¦  «        S rC   )r   rÓ   rƒ   r	  s     r!   Úresolve_cbrtzCudaModuleTemplate.resolve_cbrtâ  rP  r#   c                 ó4   — t          j        t          ¦  «        S rC   )r   rÓ   rH   r	  s     r!   Úresolve_threadfencez&CudaModuleTemplate.resolve_threadfenceå  ó   € ÝŒ~Õ5Ñ6Ô6Ð6r#   c                 ó4   — t          j        t          ¦  «        S rC   )r   rÓ   rM   r	  s     r!   Úresolve_threadfence_blockz,CudaModuleTemplate.resolve_threadfence_blockè  s   € ÝŒ~Õ4Ñ5Ô5Ð5r#   c                 ó4   — t          j        t          ¦  «        S rC   )r   rÓ   rP   r	  s     r!   Úresolve_threadfence_systemz-CudaModuleTemplate.resolve_threadfence_systemë  rÌ  r#   c                 ó4   — t          j        t          ¦  «        S rC   )r   rÓ   rS   r	  s     r!   Úresolve_syncwarpz#CudaModuleTemplate.resolve_syncwarpî  s   € ÝŒ~�mÑ,Ô,Ð,r#   c                 ó4   — t          j        t          ¦  «        S rC   )r   rÓ   rW   r	  s     r!   Úresolve_shfl_sync_intrinsicz.CudaModuleTemplate.resolve_shfl_sync_intrinsicñ  ó   € ÝŒ~Õ6Ñ7Ô7Ð7r#   c                 ó4   — t          j        t          ¦  «        S rC   )r   rÓ   r^   r	  s     r!   Úresolve_vote_sync_intrinsicz.CudaModuleTemplate.resolve_vote_sync_intrinsicô  rÕ  r#   c                 ó4   — t          j        t          ¦  «        S rC   )r   rÓ   ra   r	  s     r!   Úresolve_match_any_syncz)CudaModuleTemplate.resolve_match_any_sync÷  ó   € ÝŒ~Õ1Ñ2Ô2Ð2r#   c                 ó4   — t          j        t          ¦  «        S rC   )r   rÓ   rd   r	  s     r!   Úresolve_match_all_syncz)CudaModuleTemplate.resolve_match_all_syncú  rÚ  r#   c                 ó4   — t          j        t          ¦  «        S rC   )r   rÓ   rg   r	  s     r!   Úresolve_activemaskz%CudaModuleTemplate.resolve_activemaský  r#  r#   c                 ó4   — t          j        t          ¦  «        S rC   )r   rÓ   rk   r	  s     r!   Úresolve_lanemask_ltz&CudaModuleTemplate.resolve_lanemask_lt   r  r#   c                 ó4   — t          j        t          ¦  «        S rC   )r   rÓ   r�   r	  s     r!   Úresolve_selpzCudaModuleTemplate.resolve_selp  rP  r#   c                 ó4   — t          j        t          ¦  «        S rC   )r   rÓ   r  r	  s     r!   Úresolve_nanosleepz$CudaModuleTemplate.resolve_nanosleep  r-  r#   c                 ó>   — t          j        t          j        ¦  «        S rC   )r   r  r   rû   r	  s     r!   Úresolve_atomicz!CudaModuleTemplate.resolve_atomic	  r½  r#   c                 ó>   — t          j        t          j        ¦  «        S rC   )r   r  r   r   r	  s     r!   Úresolve_fp16zCudaModuleTemplate.resolve_fp16  s   € ÝŒ|�DœIÑ&Ô&Ð&r#   c                 ó>   — t          j        t          j        ¦  «        S rC   )r   r  r   rE   r	  s     r!   Úresolve_constz CudaModuleTemplate.resolve_const  ó   € ÝŒ|�DœJÑ'Ô'Ð'r#   c                 ó>   — t          j        t          j        ¦  «        S rC   )r   r  r   r>   r	  s     r!   Úresolve_localz CudaModuleTemplate.resolve_local  rë  r#   N)$r4   r5   r6   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!   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¬  )ÚlogÚlog2Úlog10)°ÚoperatorÚ
numba.corer   Únumba.core.typing.npydeclr   r   r   r   r   r	   r
   r   Únumba.core.typing.templatesr   r   r   r   r   r   Únumba.cuda.typesr   Únumba.core.typeconvr   Únumbar   Únumba.cuda.compilerr   Úregistryr    Úregister_attrr«   r   r8   r=   r@   rH   rM   rP   rS   rW   r^   ra   rd   rg   rk   rn   ry   r~   rƒ   r†   r‰   rŒ   r�   r¢   r¬   r±   Úfloatr³   rº   rÈ   rÌ   rÎ   r   ÚhaddrN  ÚaddÚCuda_addÚiaddÚ	Cuda_iaddÚhsubrR  ÚsubÚCuda_subÚisubÚ	Cuda_isubÚhmulrU  ÚmulÚCuda_mulÚimulÚ	Cuda_imulÚhmaxr¦  Úhminr©  Úhnegr\  ÚnegÚCuda_negÚhabsr_  ÚabsÚCuda_absÚheqr“  ÚeqÚhner—  ÚneÚhgerš  ÚgeÚhgtr�  ÚgtÚhler   ÚleÚhltr£  ÚltÚtruedivÚitruedivr×   rÙ   rd  rg  rj  rm  rq  ru  rx  r{  r„  r‡  r~  r�  rŠ  r�  r�  rX  rô   r|   r{   rs   ri   rt   rw   Úall_numba_typesrø   Úunsigned_int_numba_typesrû   r!  r%  Úmaxr;  Úminr>  ÚnanmaxrE  ÚnanminrA  Úand_r(  Úor_r+  Úxorr/  Úincr2  Údecr5  Úexchr8  rö   rþ   r  r  r  r  r  r  rL  r¬  r  Úfuncr   r#   r!   ú<module>r.     sx  ðØ €€€Ø Ð Ð Ð Ð Ð ð@ð @ð @ð @ð @ð @ð @ð @ð @ð @ð @ð @ð @ð @ð @ð @ð @ð @ð @ð @ð>ð >ð >ð >ð >ð >ð >ð >ð >ð >ð >ð >ð >ð >ð >ð >ð "Ð !Ð !Ð !Ð !Ð !Ø *Ð *Ð *Ð *Ð *Ð *Ø Ð Ð Ð Ð Ð Ø @Ð @Ð @Ð @Ð @Ð @àˆ8‰:Œ:€ØÔ€ØÔ&€ØÔ*€à Ð ˜Ñ (Ô (Ð (ðð ð ð ð Ð&ñ ô ð ð0 
ðð ð ð ð ˜ñ ô ñ 
„ðð 
ðð ð ð ð �ñ ô ñ 
„ðð 
ðð ð ð ð Ð,ñ ô ñ 
„ðð 
ð$ð $ð $ð $ð $Ð.ñ $ô $ñ 
„ð$ð
 
ð$ð $ð $ð $ð $Ð-ñ $ô $ñ 
„ð$ð
 
ð$ð $ð $ð $ð $Ð.ñ $ô $ñ 
„ð$ð
 
ð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ØÐ˜˜Ñ/Ô/Ð/Ð/àð 4ð 4€DØÐ'Ð'Ð'ØÐ˜T ?Ñ3Ô3Ð3øð4ð 4r#   