§
    1@Äj2�  ã                   ó°  — d dl Z d dlZd dlZd dlZd dlmZmZ d dlm	Z	 d dl
mZ d dlmZ d dlmZ d dlmZmZmZm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#„ Z0d$„ Z1d%„ Z2 ej3        d&¬'¦  «        d(„ ¦   «         Z4 ej3        d&¬'¦  «        d)„ ¦   «         Z5d*„ Z6d+„ Z7d,„ Z8d-„ Z9d.„ Z:d/„ Z;d0„ Z<d1„ Z=d2„ Z>d3„ Z?d4„ Z@d5„ ZAd6„ ZBd7„ ZCd8„ ZDd9„ ZEd:„ ZFd;„ ZGd<„ ZHd=„ ZId>„ ZJd?„ ZKd@„ ZLdA„ ZMdB„ ZNdC„ ZOdD„ ZPdE„ ZQdF„ ZRdG„ ZSdH„ ZTdI„ ZUdJ„ ZVdK„ ZW G dL„ dMe¦  «        ZXeYdNk    r ejZ        ¦   «          dS dS )Oé    N)ÚcudaÚint64)Úcompile_ptx)Úconfig)ÚTypingError)Úf2)ÚunittestÚCUDATestCaseÚskip_on_cudasimÚskip_unless_cc_53c                 ó2   — t           j        j        }|| d<   d S ©Nr   ©r   Ú	threadIdxÚx©ÚaryÚis     úXC:\Projects\Historical\venv\Lib\site-packages\numba/cuda/tests/cudapy/test_intrinsics.pyÚsimple_threadidxr      ó   € ÝŒÔ€AØ€Cˆ�F€F€Fó    c                 ó2   — t           j        j        }|| |<   d S ©Nr   r   s     r   Úfill_threadidxr      r   r   c                 óš   — t           j        j        }t           j        j        }t           j        j        }|dz   |dz   z  |dz   z  | |||f<   d S ©Né   )r   r   r   ÚyÚz)r   r   ÚjÚks       r   Úfill3d_threadidxr#      sI   € ÝŒÔ€AÝŒÔ€AÝŒÔ€Aà˜‘E˜a !™eÑ$¨¨A©Ñ.€Cˆˆ1ˆaˆ�L€L€Lr   c                 ó8   — t          j        d¦  «        }|| |<   d S r   ©r   Úgridr   s     r   Úsimple_grid1dr'       s   € ÝŒ	�!‰Œ€AØ€Cˆ�F€F€Fr   c                 óH   — t          j        d¦  «        \  }}||z   | ||f<   d S ©Né   r%   )r   r   r!   s      r   Úsimple_grid2dr+   %   s(   € ÝŒ9�Q‰<Œ<�D€A€qØ�A‘€Cˆˆ1ˆ�I€I€Ir   c                 óp   — t          j        d¦  «        }t          j        d¦  «        }|dk    r|| d<   d S d S ©Nr   r   ©r   r&   Úgridsize)r   r   r   s      r   Úsimple_gridsize1dr0   *   s;   € ÝŒ	�!‰Œ€AÝŒ�aÑÔ€AØˆA‚v€vØˆˆA‰ˆˆð €vr   c                 ó–   — t          j        d¦  «        \  }}t          j        d¦  «        \  }}|dk    r|dk    r|| d<   || d<   d S d S d S )Nr*   r   r   r.   )r   r   r!   r   r   s        r   Úsimple_gridsize2dr2   1   sX   € ÝŒ9�Q‰<Œ<�D€A€qÝŒ=˜ÑÔ�D€A€qØˆA‚v€v�!�q’&�&ØˆˆA‰ØˆˆA‰ˆˆð €v�&�&r   c                 ó8  — t          j        d¦  «        \  }}t           j        j        t           j        j        z  }t           j        j        t           j        j        z  }| j        \  }}t          |||¦  «        D ] }t          |||¦  «        D ]}||z   | ||f<   ŒŒ!d S r)   )r   r&   ÚgridDimr   ÚblockDimr   ÚshapeÚrange)	ÚcÚstartXÚstartYÚgridXÚgridYÚheightÚwidthr   r   s	            r   Úintrinsic_forloop_stepr?   9   sš   € Ý”Y˜q‘\”\�N€FˆFÝŒLŒN�Tœ]œ_Ñ,€EÝŒLŒN�Tœ]œ_Ñ,€EØ”G�M€FˆEå�6˜5 %Ñ(Ô(ð ð ˆÝ�v˜v uÑ-Ô-ð 	ð 	ˆAØ˜!‘eˆAˆa�ˆd‰GˆGð	ðð r   c                 ó4   — t          j        |¦  «        | d<   d S r   )r   Úpopc©r   r8   s     r   Úsimple_popcrC   D   ó   € ÝŒY�q‰\Œ\€Cˆ�F€F€Fr   c                 ó8   — t          j        |||¦  «        | d<   d S r   )r   Úfma©r   ÚaÚbr8   s       r   Ú
simple_fmarJ   H   s   € ÝŒX�a˜˜AÑÔ€Cˆ�F€F€Fr   c                 ód   — t           j                             |d         |d         ¦  «        | d<   d S r   ©r   Úfp16Úhadd©r   rH   rI   s      r   Úsimple_haddrP   L   ó'   € ÝŒY�^Š^˜A˜aœD ! A¤$Ñ'Ô'€Cˆ�F€F€Fr   c                 óL   — t           j                             ||¦  «        | d<   d S r   rL   rO   s      r   Úsimple_hadd_scalarrS   P   ó   € ÝŒY�^Š^˜A˜qÑ!Ô!€Cˆ�F€F€Fr   c                 ór   — t           j                             |d         |d         |d         ¦  «        | d<   d S r   ©r   rM   ÚhfmarG   s       r   Úsimple_hfmarX   T   s-   € ÝŒY�^Š^˜A˜aœD ! A¤$¨¨!¬Ñ-Ô-€Cˆ�F€F€Fr   c                 óN   — t           j                             |||¦  «        | d<   d S r   rV   rG   s       r   Úsimple_hfma_scalarrZ   X   s!   € ÝŒY�^Š^˜A˜q !Ñ$Ô$€Cˆ�F€F€Fr   c                 ód   — t           j                             |d         |d         ¦  «        | d<   d S r   ©r   rM   ÚhsubrO   s      r   Úsimple_hsubr^   \   rQ   r   c                 óL   — t           j                             ||¦  «        | d<   d S r   r\   rO   s      r   Úsimple_hsub_scalarr`   `   rT   r   c                 ód   — t           j                             |d         |d         ¦  «        | d<   d S r   ©r   rM   ÚhmulrO   s      r   Úsimple_hmulrd   d   rQ   r   c                 óL   — t           j                             ||¦  «        | d<   d S r   rb   rO   s      r   Úsimple_hmul_scalarrf   h   rT   r   c                 óL   — t           j                             ||¦  «        | d<   d S r   )r   rM   ÚhdivrO   s      r   Úsimple_hdiv_scalarri   l   rT   r   c                 ó®   — t          j        d¦  «        }|| j        k     r5||         }||         }t           j                             ||¦  «        | |<   d S d S r   )r   r&   ÚsizerM   rh   )r   Úarray_aÚarray_br   rH   rI   s         r   Úsimple_hdiv_kernelrn   p   sP   € ÝŒ	�!‰Œ€AØˆ3Œ8‚|€|Ø�AŒJˆØ�AŒJˆÝ”—’  1Ñ%Ô%ˆˆA‰ˆˆð €|r   c                 óV   — t           j                             |d         ¦  «        | d<   d S r   ©r   rM   Úhneg©r   rH   s     r   Úsimple_hnegrs   x   ó!   € ÝŒY�^Š^˜A˜aœDÑ!Ô!€Cˆ�F€F€Fr   c                 óJ   — t           j                             |¦  «        | d<   d S r   rp   rr   s     r   Úsimple_hneg_scalarrv   |   ó   € ÝŒY�^Š^˜AÑÔ€Cˆ�F€F€Fr   c                 óV   — t           j                             |d         ¦  «        | d<   d S r   ©r   rM   Úhabsrr   s     r   Úsimple_habsr{   €   rt   r   c                 óJ   — t           j                             |¦  «        | d<   d S r   ry   rr   s     r   Úsimple_habs_scalarr}   „   rw   r   c                 óL   — t           j                             ||¦  «        | d<   d S r   )r   rM   ÚheqrO   s      r   Úsimple_heq_scalarr€   ˆ   ó   € ÝŒY�]Š]˜1˜aÑ Ô €Cˆ�F€F€Fr   c                 óL   — t           j                             ||¦  «        | d<   d S r   )r   rM   ÚhnerO   s      r   Úsimple_hne_scalarr„   Œ   r�   r   c                 óL   — t           j                             ||¦  «        | d<   d S r   )r   rM   ÚhgerO   s      r   Úsimple_hge_scalarr‡   �   r�   r   c                 óL   — t           j                             ||¦  «        | d<   d S r   )r   rM   ÚhgtrO   s      r   Úsimple_hgt_scalarrŠ   ”   r�   r   c                 óL   — t           j                             ||¦  «        | d<   d S r   )r   rM   ÚhlerO   s      r   Úsimple_hle_scalarr�   ˜   r�   r   c                 óL   — t           j                             ||¦  «        | d<   d S r   ©r   rM   ÚhltrO   s      r   Úsimple_hlt_scalarr‘   œ   r�   r   T)Údevicec                 óB   — t           j                             | |¦  «        S r   r�   ©r   r   s     r   Ú
hlt_func_1r•       ó   € åŒ9�=Š=˜˜AÑÔÐr   c                 óB   — t           j                             | |¦  «        S r   r�   r”   s     r   Ú
hlt_func_2r˜   ¥   r–   r   c                 óL   — t          ||¦  «        ot          ||¦  «        | d<   d S r   )r•   r˜   ©ÚrrH   rI   r8   s       r   Útest_multiple_hcmp_1rœ   ª   s(   € å�a˜ÑÔÐ0¥
¨1¨aÑ 0Ô 0€A€a�D€D€Dr   c                 ól   — t          ||¦  «        ot          j                             ||¦  «        | d<   d S r   )r•   r   rM   r�   rš   s       r   Útest_multiple_hcmp_2rž   ¯   ó.   € å�a˜ÑÔÐ3¥¤	§¢¨a°Ñ 3Ô 3€A€a�D€D€Dr   c                 ól   — t          ||¦  «        ot          j                             ||¦  «        | d<   d S r   )r•   r   rM   r†   rš   s       r   Útest_multiple_hcmp_3r¡   ´   rŸ   r   c                 óŒ   — t           j                             ||¦  «        ot           j                             ||¦  «        | d<   d S r   r�   rš   s       r   Útest_multiple_hcmp_4r£   ¹   ó4   € åŒ9�=Š=˜˜AÑÔÐ6¥4¤9§=¢=°°AÑ#6Ô#6€A€a�D€D€Dr   c                 óŒ   — t           j                             ||¦  «        ot           j                             ||¦  «        | d<   d S r   )r   rM   r�   r†   rš   s       r   Útest_multiple_hcmp_5r¦   ¾   r¤   r   c                 óL   — t           j                             ||¦  «        | d<   d S r   )r   rM   ÚhmaxrO   s      r   Úsimple_hmax_scalarr©   Ã   rT   r   c                 óL   — t           j                             ||¦  «        | d<   d S r   )r   rM   ÚhminrO   s      r   Úsimple_hmin_scalarr¬   Ç   rT   r   c                 ó¨   — t          j        d¦  «        }|t          | ¦  «        k     r*t           j                             ||         ¦  «        | |<   d S d S r   )r   r&   ÚlenrM   Úhsin©r›   r   r   s      r   Úsimple_hsinr±   Ë   óD   € ÝŒ	�!‰Œ€Aà�3ˆq‰6Œ6‚z€zÝŒy�~Š~˜a œdÑ#Ô#ˆˆ!‰ˆˆð €zr   c                 ó¨   — t          j        d¦  «        }|t          | ¦  «        k     r*t           j                             ||         ¦  «        | |<   d S d S r   )r   r&   r®   rM   Úhcosr°   s      r   Úsimple_hcosrµ   Ò   r²   r   c                 ó¨   — t          j        d¦  «        }|t          | ¦  «        k     r*t           j                             ||         ¦  «        | |<   d S d S r   )r   r&   r®   rM   Úhlogr°   s      r   Úsimple_hlogr¸   Ù   r²   r   c                 ó¨   — t          j        d¦  «        }|t          | ¦  «        k     r*t           j                             ||         ¦  «        | |<   d S d S r   )r   r&   r®   rM   Úhlog2r°   s      r   Úsimple_hlog2r»   à   óD   € ÝŒ	�!‰Œ€Aà�3ˆq‰6Œ6‚z€zÝŒy�Š˜q œtÑ$Ô$ˆˆ!‰ˆˆð €zr   c                 ó¨   — t          j        d¦  «        }|t          | ¦  «        k     r*t           j                             ||         ¦  «        | |<   d S d S r   )r   r&   r®   rM   Úhlog10r°   s      r   Úsimple_hlog10r¿   ç   óF   € ÝŒ	�!‰Œ€Aà�3ˆq‰6Œ6‚z€zÝŒy×Ò  !¤Ñ%Ô%ˆˆ!‰ˆˆð €zr   c                 ó¨   — t          j        d¦  «        }|t          | ¦  «        k     r*t           j                             ||         ¦  «        | |<   d S d S r   )r   r&   r®   rM   Úhexpr°   s      r   Úsimple_hexprÃ   î   r²   r   c                 ó¨   — t          j        d¦  «        }|t          | ¦  «        k     r*t           j                             ||         ¦  «        | |<   d S d S r   )r   r&   r®   rM   Úhexp2r°   s      r   Úsimple_hexp2rÆ   õ   r¼   r   c                 ó¨   — t          j        d¦  «        }|t          | ¦  «        k     r*t           j                             ||         ¦  «        | |<   d S d S r   )r   r&   r®   rM   Úhsqrtr°   s      r   Úsimple_hsqrtrÉ   ü   r¼   r   c                 ó¨   — t          j        d¦  «        }|t          | ¦  «        k     r*t           j                             ||         ¦  «        | |<   d S d S r   )r   r&   r®   rM   Úhrsqrtr°   s      r   Úsimple_hrsqrtrÌ     sF   € åŒ	�!‰Œ€Aà�3ˆq‰6Œ6‚z€zÝŒy×Ò  !¤Ñ%Ô%ˆˆ!‰ˆˆð €zr   c                 ó   — | dz  S )Ng      à¿© )r   Údtypes     r   Únumpy_hrsqrtrÐ     s   € Ø�‰9Ðr   c                 ó¨   — t          j        d¦  «        }|t          | ¦  «        k     r*t           j                             ||         ¦  «        | |<   d S d S r   )r   r&   r®   rM   Úhceilr°   s      r   Úsimple_hceilrÓ     r¼   r   c                 ó¨   — t          j        d¦  «        }|t          | ¦  «        k     r*t           j                             ||         ¦  «        | |<   d S d S r   )r   r&   r®   rM   Úhfloorr°   s      r   Úsimple_hfloorrÖ     rÀ   r   c                 ó¨   — t          j        d¦  «        }|t          | ¦  «        k     r*t           j                             ||         ¦  «        | |<   d S d S r   )r   r&   r®   rM   Úhrcpr°   s      r   Úsimple_hrcprÙ     r²   r   c                 ó¨   — t          j        d¦  «        }|t          | ¦  «        k     r*t           j                             ||         ¦  «        | |<   d S d S r   )r   r&   r®   rM   Úhtruncr°   s      r   Úsimple_htruncrÜ   $  rÀ   r   c                 ó¨   — t          j        d¦  «        }|t          | ¦  «        k     r*t           j                             ||         ¦  «        | |<   d S d S r   )r   r&   r®   rM   Úhrintr°   s      r   Úsimple_hrintrß   +  r¼   r   c                 ó4   — t          j        |¦  «        | d<   d S r   )r   Úcbrtrr   s     r   Úsimple_cbrtrâ   2  rD   r   c                 ó4   — t          j        |¦  «        | d<   d S r   )r   ÚbrevrB   s     r   Úsimple_brevrå   6  rD   r   c                 ó4   — t          j        |¦  «        | d<   d S r   )r   ÚclzrB   s     r   Ú
simple_clzrè   :  ó   € ÝŒX�a‰[Œ[€Cˆ�F€F€Fr   c                 ó4   — t          j        |¦  «        | d<   d S r   )r   ÚffsrB   s     r   Ú
simple_ffsrì   >  ré   r   c                 ó*   — t          |¦  «        | d<   d S r   ©ÚroundrB   s     r   Úsimple_roundrð   B  s   € Ý�1‰XŒX€Cˆ�F€F€Fr   c                 ó,   — t          ||¦  «        | d<   d S r   rî   )r   r8   Úndigitss      r   Úsimple_round_toró   F  s   € Ý�1�gÑÔ€Cˆ�F€F€Fr   c                 óŠ   — t          j        d¦  «        }| |         dk    r|dz  dk    r||         | |<   d S d| |<   d S d| |<   d S )Nr   é   r*   r   é   é   r%   )rH   rI   r8   r   s       r   Úbranching_with_ifsrø   J  sT   € ÝŒ	�!‰Œ€Aàˆ„tˆa‚x€xØˆq‰5�AŠ:ˆ:Ø�Q”4ˆAˆa‰DˆDˆDàˆAˆa‰DˆDˆDàˆˆ!‰ˆˆr   c                 óº   — t          j        d¦  «        }t          j        |dz  dk    ||         d¦  «        }t          j        | |         dk    |d¦  «        | |<   d S )Nr   r*   r   rö   rõ   r÷   )r   r&   Úselp)rH   rI   r8   r   Úinners        r   Úbranching_with_selpsrü   V  sQ   € ÝŒ	�!‰Œ€AåŒI�a˜!‘e˜q’j ! A¤$¨Ñ+Ô+€EÝŒ9�Q�q”T˜A’X˜u aÑ(Ô(€A€a�D€D€Dr   c                 óL   — t          j        d¦  «        }t           j        | |<   d S r   )r   r&   Úlaneidr   s     r   Úsimple_laneidrÿ   ]  s   € ÝŒ	�!‰Œ€AÝŒ[€Cˆ�F€F€Fr   c                 ó$   — t           j        | d<   d S r   )r   Úwarpsize)r   s    r   Úsimple_warpsizer  b  s   € ÝŒ]€Cˆ�F€F€Fr   c                 ó.   — t          j        | ¦  «         d S r   r%   ©r   s    r   Únonliteral_gridr  f  s   € Ý„Iˆa�L„L€L€L€Lr   c                 ó.   — t          j        | ¦  «         d S r   )r   r/   r  s    r   Únonliteral_gridsizer  j  s   € Ý„M�!ÑÔÐÐÐr   c                   óô  ‡ — e Zd Zˆ fd„Zd„ Zd„ Zd„ Z ed¦  «        d„ ¦   «         Z ed¦  «        d„ ¦   «         Z	d„ Z
d	„ Zd
„ Z ed¦  «        d„ ¦   «         Z ed¦  «        d„ ¦   «         Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Zed„ ¦   «         Zed„ ¦   «         Z ed¦  «        d„ ¦   «         Zed„ ¦   «         Zed„ ¦   «         Z ed¦  «        d„ ¦   «         Zed„ ¦   «         Zed„ ¦   «         Z ed¦  «        d „ ¦   «         Z ed!„ ¦   «         Z!ed"„ ¦   «         Z" ed¦  «        d#„ ¦   «         Z#ed$„ ¦   «         Z$ed%„ ¦   «         Z%ed&„ ¦   «         Z&ed'„ ¦   «         Z' ed¦  «        d(„ ¦   «         Z(ed)„ ¦   «         Z)ed*„ ¦   «         Z* ed¦  «        d+„ ¦   «         Z+ed,„ ¦   «         Z,ed-„ ¦   «         Z-ed.„ ¦   «         Z.ed/„ ¦   «         Z/ed0„ ¦   «         Z0ed1„ ¦   «         Z1d2„ Z2d3„ Z3d4„ Z4 ed5¦  «        d6„ ¦   «         Z5d7„ Z6d8„ Z7d9„ Z8d:„ Z9 ed5¦  «        d;„ ¦   «         Z:d<„ Z;d=„ Z<d>„ Z=d?„ Z> ed5¦  «        d@„ ¦   «         Z?dA„ Z@dB„ ZAdC„ ZBdD„ ZCdE„ ZD edF¦  «        dG„ ¦   «         ZEdH„ ZFdI„ ZG edF¦  «        dJ„ ¦   «         ZHdK„ ZIˆ xZJS )LÚTestCudaIntrinsicc                 ó†   •— t          ¦   «                              ¦   «          t          j                             d¦  «         d S r   )ÚsuperÚsetUpÚnpÚrandomÚseed)ÚselfÚ	__class__s    €r   r  zTestCudaIntrinsic.setUpo  s.   ø€ Ý‰Œ�Š‰ŒˆÝ
Œ	�Š�qÑÔÐÐÐr   c                 óê   —  t          j        d¦  «        t          ¦  «        }t          j        dt          j        ¬¦  «        } |d         |¦  «         |                      |d         dk    ¦  «         d S )Núvoid(int32[:])r   ©rÏ   ©r   r   r   )r   Újitr   r  ÚonesÚint32Ú
assertTrue©r  Úcompiledr   s      r   Útest_simple_threadidxz'TestCudaIntrinsic.test_simple_threadidxs  sg   € Ø-•4”8Ð,Ñ-Ô-Õ.>Ñ?Ô?ˆÝŒg�a�rœxÐ(Ñ(Ô(ˆØˆ�Œ�sÑÔÐØ�Š˜˜Aœ !šÑ$Ô$Ð$Ð$Ð$r   c                 óJ  —  t          j        d¦  «        t          ¦  «        }d}t          j        |t          j        ¬¦  «        }t          j        |t          j        ¬¦  «        } |d|f         |¦  «         |                      t          j        ||k    ¦  «        ¦  «         d S )Nr  é
   r  r   )	r   r  r   r  r  r  Úaranger  Úall)r  r  ÚNr   Úexps        r   Útest_fill_threadidxz%TestCudaIntrinsic.test_fill_threadidxy  s‹   € Ø-•4”8Ð,Ñ-Ô-­nÑ=Ô=ˆØˆÝŒg�a�rœxÐ(Ñ(Ô(ˆÝŒi˜¥¤Ð*Ñ*Ô*ˆØˆ��A�Œ�sÑÔÐØ�Š�œ˜s cšzÑ*Ô*Ñ+Ô+Ð+Ð+Ð+r   c                 ó²   ‡‡‡— d\  ŠŠŠˆˆˆfd„}ˆˆˆfd„} |¦   «         } |¦   «         }|                       t          j        ||k    ¦  «        ¦  «         d S )N)rõ   é   é   c                  ó¾   •—  t          j        d¦  «        t          ¦  «        } t          j        ‰‰‰ft          j        ¬¦  «        } | d‰‰‰ff         |¦  «         |S )Nzvoid(int32[:,:,::1])r  r   )r   r  r#   r  Úzerosr  ©r  r   ÚXÚYÚZs     €€€r   Úc_contigousz<TestCudaIntrinsic.test_fill3d_threadidx.<locals>.c_contigous„  s^   ø€ Ø7•t”xÐ 6Ñ7Ô7Õ8HÑIÔIˆHÝ”(˜A˜q !˜9­B¬HÐ5Ñ5Ô5ˆCØ"ˆH�Q˜˜A˜q˜	�\Ô" 3Ñ'Ô'Ð'ØˆJr   c                  óâ   •—  t          j        d¦  «        t          ¦  «        } t          j        t          j        ‰‰‰ft          j        ¬¦  «        ¦  «        } | d‰‰‰ff         |¦  «         |S )Nzvoid(int32[::1,:,:])r  r   )r   r  r#   r  Úasfortranarrayr(  r  r)  s     €€€r   Úf_contigousz<TestCudaIntrinsic.test_fill3d_threadidx.<locals>.f_contigousŠ  si   ø€ Ø7•t”xÐ 6Ñ7Ô7Õ8HÑIÔIˆHÝÔ#¥B¤H¨a°°A¨Y½b¼hÐ$GÑ$GÔ$GÑHÔHˆCØ"ˆH�Q˜˜A˜q˜	�\Ô" 3Ñ'Ô'Ð'ØˆJr   )r  r  r   )r  r-  r0  Úc_resÚf_resr*  r+  r,  s        @@@r   Útest_fill3d_threadidxz'TestCudaIntrinsic.test_fill3d_threadidx�  s”   øøø€ Ø‰ˆˆ1ˆað	ð 	ð 	ð 	ð 	ð 	ð 	ð	ð 	ð 	ð 	ð 	ð 	ð 	ð �‘”ˆØ�‘”ˆØ�Š�œ˜u¨š~Ñ.Ô.Ñ/Ô/Ð/Ð/Ð/r   zCudasim does not check typesc                 ó²   — |                       t          d¦  «        5   t          j        d¦  «        t          ¦  «         d d d ¦  «         d S # 1 swxY w Y   d S ©NÚRequireLiteralValuezvoid(int32))ÚassertRaisesRegexr   r   r  r  ©r  s    r   Útest_nonliteral_grid_errorz,TestCudaIntrinsic.test_nonliteral_grid_error”  s—   € à×#Ò#¥KÐ1FÑGÔGð 	5ð 	5Ø#�DŒH�]Ñ#Ô#¥OÑ4Ô4Ð4ð	5ð 	5ð 	5ñ 	5ô 	5ð 	5ð 	5ð 	5ð 	5ð 	5ð 	5ð 	5øøøð 	5ð 	5ð 	5ð 	5ð 	5ð 	5ó   œ#AÁAÁAc                 ó²   — |                       t          d¦  «        5   t          j        d¦  «        t          ¦  «         d d d ¦  «         d S # 1 swxY w Y   d S r5  )r7  r   r   r  r  r8  s    r   Útest_nonliteral_gridsize_errorz0TestCudaIntrinsic.test_nonliteral_gridsize_error™  s˜   € à×#Ò#¥KÐ1FÑGÔGð 	9ð 	9Ø#�DŒH�]Ñ#Ô#Õ$7Ñ8Ô8Ð8ð	9ð 	9ð 	9ñ 	9ô 	9ð 	9ð 	9ð 	9ð 	9ð 	9ð 	9ð 	9øøøð 	9ð 	9ð 	9ð 	9ð 	9ð 	9r:  c                 ó>  —  t          j        d¦  «        t          ¦  «        }d\  }}||z  }t          j        |t          j        ¬¦  «        } |||f         |¦  «         |                      t          j        |t          j        |¦  «        k    ¦  «        ¦  «         d S )Núvoid(int32[::1])©r÷   é   r  )	r   r  r'   r  Úemptyr  r  r   r  )r  r  ÚntidÚnctaidÚnelemr   s         r   Útest_simple_grid1dz$TestCudaIntrinsic.test_simple_grid1dž  sŽ   € Ø/•4”8Ð.Ñ/Ô/µÑ>Ô>ˆØ‰ˆˆfØ�v‘ˆÝŒh�u¥B¤HÐ-Ñ-Ô-ˆØˆ�˜�Ô˜sÑ#Ô#Ð#Ø�Š�œ˜s¥b¤i°Ñ&6Ô&6Ò6Ñ7Ô7Ñ8Ô8Ð8Ð8Ð8r   c                 ó  —  t          j        d¦  «        t          ¦  «        }d}d}|d         |d         z  |d         |d         z  f}t          j        |t          j        ¬¦  «        }|                     ¦   «         } |||f         |¦  «         t          |j        d         ¦  «        D ])}t          |j        d         ¦  «        D ]}||z   |||f<   ŒŒ*|  	                    t          j
        ||k    ¦  «        ¦  «         d S ©Nzvoid(int32[:,::1])©rõ   r÷   ©r%  r&  r   r   r  )r   r  r+   r  rA  r  Úcopyr7   r6   r  r   )	r  r  rB  rC  r6   r   r"  r   r!   s	            r   Útest_simple_grid2dz$TestCudaIntrinsic.test_simple_grid2d¦  s   € Ø1•4”8Ð0Ñ1Ô1µ-Ñ@Ô@ˆØˆØˆØ�a”˜6 !œ9Ñ$ d¨1¤g°°q´	Ñ&9Ð:ˆÝŒh�u¥B¤HÐ-Ñ-Ô-ˆØ�hŠh‰jŒjˆØˆ�˜�Ô˜sÑ#Ô#Ð#å�s”y ”|Ñ$Ô$ð 	"ð 	"ˆAÝ˜3œ9 Qœ<Ñ(Ô(ð "ð "�Ø ™E��A�q�D‘	�	ð"ð 	�Š�œ˜s cšzÑ*Ô*Ñ+Ô+Ð+Ð+Ð+r   c                 óø   —  t          j        d¦  «        t          ¦  «        }d\  }}t          j        dt          j        ¬¦  «        } |||f         |¦  «         |                      |d         ||z  ¦  «         d S )Nr>  r?  r   r  r   )r   r  r0   r  r(  r  ÚassertEqual©r  r  rB  rC  r   s        r   Útest_simple_gridsize1dz(TestCudaIntrinsic.test_simple_gridsize1dµ  sy   € Ø/•4”8Ð.Ñ/Ô/Õ0AÑBÔBˆØ‰ˆˆfÝŒh�q¥¤Ð)Ñ)Ô)ˆØˆ�˜�Ô˜sÑ#Ô#Ð#Ø×Ò˜˜Qœ ¨$¡Ñ/Ô/Ð/Ð/Ð/r   zRequires too many threadsc                 óD  — t           j        d„ ¦   «         }t          j        dt          j        ¬¦  «        }t          j        dt          j        ¬¦  «        } |d         ||¦  «         |                      |d         d¦  «         |                      |d         d¦  «         d S )Nc                 ó6  — t          j        d¦  «        }t           j        j        t           j        j        z  t           j        j        z   }t          j        d¦  «        }t           j        j        t           j        j        z  }||k    rd| d<   ||k    rd|d<   d S d S r-   )r   r&   ÚblockIdxr   r5   r   r/   r4   )Ú
grid_errorÚgridsize_errorÚi1Úi2Úgs1Úgs2s         r   Úfz,TestCudaIntrinsic.test_issue_9229.<locals>.fÀ  s~   € å”˜1‘”ˆBÝ””¥4¤=¤?Ñ2µT´^Ô5EÑEˆBÝ”- Ñ"Ô"ˆCÝ”-”/¥D¤L¤NÑ2ˆCØ�RŠxˆxØ !�
˜1‘Ø�cŠzˆzØ$%�˜qÑ!Ð!Ð!ð ˆzr   r   r  )i ‹Qé   r   )r   r  r  r(  Úuint64rM  )r  rY  rS  rT  s       r   Útest_issue_9229z!TestCudaIntrinsic.test_issue_9229¼  s¡   € õ 
Œð	&ð 	&ñ 
Œð	&õ ”X˜a¥r¤yÐ1Ñ1Ô1ˆ
Ýœ !­2¬9Ð5Ñ5Ô5ˆð 	ˆˆ-Ô˜ ^Ñ4Ô4Ð4à×Ò˜ Aœ¨Ñ*Ô*Ð*Ø×Ò˜¨Ô*¨AÑ.Ô.Ð.Ð.Ð.r   zTests PTX emissionc           	      óÞ  — t           d d …         t           t           d d …         f} t          j        |¦  «        t          ¦  «        } t          j        |¦  «        t          ¦  «        }d}d}t          j        ddt
          j         ¬¦  «        }|                     ¦   «         }d|d d…<   t          j        |t
          j         ¬¦  «        } ||df         |||¦  «         | 	                    |¦  «        }	|  
                    d	t          t          j        d
|	¦  «        ¦  «        ¦  «         t
          j                             ||d¬¦  «         t          j        |t
          j         ¬¦  «        } ||df         |||¦  «         | 	                    |¦  «        }	|  
                    dt          t          j        d
|	¦  «        ¦  «        ¦  «         t
          j                             ||d¬¦  «         d S )Né    r&  é   )r6   Ú
fill_valuerÏ   r÷   r%  r  r   r*   z	\s+bra\s+Ú	branching)Úerr_msgr   rú   )r   r   r  rø   rü   r  ÚfullrJ  r  Úinspect_asmrM  r®   ÚreÚfindallÚtestingÚassert_array_equal)
r  ÚsigÚcu_branching_with_ifsÚcu_branching_with_selpsÚnrI   r8   ÚexpectedrH   Úptxs
             r   Ú	test_selpzTestCudaIntrinsic.test_selpÕ  s®  € å�Q�Q�QŒx�¥ a a a¤Ð)ˆØ -¥¤¨¡¤Õ.@Ñ AÔ AÐØ"/¥$¤(¨3¡-¤-Õ0DÑ"EÔ"EÐàˆØˆÝŒG˜"¨µ2´8Ð<Ñ<Ô<ˆà—6’6‘8”8ˆØˆ��!�‰åŒI�a�rœxÐ(Ñ(Ô(ˆØ#Ð˜a ˜dÔ# A q¨!Ñ,Ô,Ð,Ø#×/Ò/°Ñ4Ô4ˆØ×Ò˜�C¥¤
¨<¸Ñ =Ô =Ñ>Ô>Ñ?Ô?Ð?Ý
Œ
×%Ò% a¨¸;Ð%ÑGÔGÐGåŒI�a�rœxÐ(Ñ(Ô(ˆØ%Ð  1 Ô% a¨¨AÑ.Ô.Ð.Ø%×1Ò1°#Ñ6Ô6ˆØ×Ò˜�C¥¤
¨<¸Ñ =Ô =Ñ>Ô>Ñ?Ô?Ð?Ý
Œ
×%Ò% a¨¸6Ð%ÑBÔBÐBÐBÐBr   c                 ód  —  t          j        d¦  «        t          ¦  «        }d}d}t          j        dt          j        ¬¦  «        } |||f         |¦  «         |                      |d         |d         |d         z  ¦  «         |                      |d         |d         |d         z  ¦  «         d S )Nr>  rH  rI  r*   r  r   r   )r   r  r2   r  r(  r  rM  rN  s        r   Útest_simple_gridsize2dz(TestCudaIntrinsic.test_simple_gridsize2dî  sª   € Ø/•4”8Ð.Ñ/Ô/Õ0AÑBÔBˆØˆØˆÝŒh�q¥¤Ð)Ñ)Ô)ˆØˆ�˜�Ô˜sÑ#Ô#Ð#à×Ò˜˜Qœ ¨¤¨T°!¬WÑ!4Ñ5Ô5Ð5Ø×Ò˜˜Qœ ¨¤¨T°!¬WÑ!4Ñ5Ô5Ð5Ð5Ð5r   c           	      óZ  —  t          j        d¦  «        t          ¦  «        }d}d}|d         |d         z  |d         |d         z  f}t          j        |t          j        ¬¦  «        } |||f         |¦  «         |\  }}|j        \  }}	t          t          |d         ¦  «        t          |d         ¦  «        ¦  «        D ]i\  }
}||
z   ||z   }}t          ||	|¦  «        D ]H}t          |||¦  «        D ]4}|  	                    |||f         ||z   k    |||f         ||z   f¦  «         Œ5ŒIŒjd S rG  )
r   r  r?   r  rA  r  r6   Úzipr7   r  )r  r  rB  rC  r6   r   r;   r<   r=   r>   r   r!   r9   r:   r   r   s                   r   Útest_intrinsic_forloop_stepz-TestCudaIntrinsic.test_intrinsic_forloop_stepø  s]  € Ø1•4”8Ð0Ñ1Ô1Õ2HÑIÔIˆØˆØˆØ�a”˜6 !œ9Ñ$ d¨1¤g°°q´	Ñ&9Ð:ˆÝŒh�u¥B¤HÐ-Ñ-Ô-ˆàˆ�˜�Ô˜sÑ#Ô#Ð#à‰ˆˆuØœ	‰ˆ�Ý�˜d 1œg™œ­¨d°1¬g©¬Ñ7Ô7ð 	Lð 	L‰DˆAˆqØ" Q™Y¨°©	�FˆFÝ˜6 5¨%Ñ0Ô0ð Lð L�Ý˜v v¨uÑ5Ô5ð Lð L�AØ—O’O C¨¨1¨¤I°°Q±Ò$6¸¸QÀ¸T¼ÀAÈÁEÐ8JÑKÔKÐKÐKðLðLð	Lð 	Lr   c                 óü   — t           j        d„ ¦   «         }t          j        dt          j        ¬¦  «                             ddd¦  «        } |d         |¦  «         t          j                             |d¦  «         d S )Nc                 ó‚   — t          j        d¦  «        \  }}}t          j        d¦  «        \  }}}||z  |z  | |||f<   d S ©Nr÷   r.   )Úoutr   r   r    rH   rI   r8   s          r   Úfooz*TestCudaIntrinsic.test_3dgrid.<locals>.foo
  sE   € å”i ‘l”l‰GˆAˆq�!Ý”m AÑ&Ô&‰GˆAˆq�!Ø˜q™5 1™9ˆC��1�a�‰LˆLˆLr   iÙ  r  é	   )©r÷   r÷   r÷   r{  )r   r  r  r(  r  Úreshaperg  Úassert_equal)r  ry  Úarrs      r   Útest_3dgridzTestCudaIntrinsic.test_3dgrid	  s{   € Ý	Œð	%ð 	%ñ 
Œð	%õ
 Œh�v¥R¤XÐ.Ñ.Ô.×6Ò6°q¸!¸QÑ?Ô?ˆØ!ˆÐ Ô! #Ñ&Ô&Ð&å
Œ
×Ò  VÑ,Ô,Ð,Ð,Ð,r   c                 ó"  — t           j        d„ ¦   «         }d\  }}}t          j        ||z  |z  t          j        ¬¦  «                             |||¦  «        } |d         |¦  «         |                      t          j        |¦  «        ¦  «         d S )Nc                 óº  — t          j        d¦  «        \  }}}t          j        d¦  «        \  }}}|t           j        j        t           j        j        t           j        j        z  z   k    oq|t           j        j        t           j        j        t           j        j        z  z   k    o8|t           j        j        t           j        j        t           j        j        z  z   k    }|t           j        j        t           j	        j        z  k    oM|t           j        j        t           j	        j        z  k    o&|t           j        j        t           j	        j        z  k    }|o|| |||f<   d S rw  )
r   r&   r/   r   r   rR  r5   r   r    r4   )	rx  r   r   r    rH   rI   r8   Úgrid_is_rightÚgridsize_is_rights	            r   ry  z,TestCudaIntrinsic.test_3dgrid_2.<locals>.foo  s  € å”i ‘l”l‰GˆAˆq�!Ý”m AÑ&Ô&‰GˆAˆq�!à•T”^Ô%­¬¬½$¼-¼/Ñ(IÑIÒIð JØ•T”^Ô%­¬¬½$¼-¼/Ñ(IÑIÒIðJà•T”^Ô%­¬¬½$¼-¼/Ñ(IÑIÒIð ð
 "#¥d¤m¤o½¼¼Ñ&FÒ!Fð "GØ!"¥d¤m¤o½¼¼Ñ&FÒ!Fð"Gà!"¥d¤m¤o½¼¼Ñ&FÒ!Fð ð )Ð>Ð->ˆC��1�a�‰LˆLˆLr   )é   r&  é   r  ))rõ   r÷   r*   )r÷   r*   rõ   )r   r  r  r(  Úbool_r|  r  r   )r  ry  r   r   r    r~  s         r   Útest_3dgrid_2zTestCudaIntrinsic.test_3dgrid_2  s�   € Ý	Œð	?ð 	?ñ 
Œð	?ð (‰ˆˆ1ˆaÝŒh˜˜A™ ™	­"¬(Ð3Ñ3Ô3×;Ò;¸A¸qÀ!ÑDÔDˆØ!ˆÐ Ô! #Ñ&Ô&Ð&à�Š�œ˜s™œÑ$Ô$Ð$Ð$Ð$r   c                 óæ   —  t          j        d¦  «        t          ¦  «        }t          j        dt          j        ¬¦  «        } |d         |d¦  «         |                      |d         d¦  «         d S )Núvoid(int32[:], uint32)r   r  r  éð   r   rõ   ©r   r  rC   r  r(  r  rM  r  s      r   Útest_popc_u4zTestCudaIntrinsic.test_popc_u4*  sh   € Ø5•4”8Ð4Ñ5Ô5µkÑBÔBˆÝŒh�q¥¤Ð)Ñ)Ô)ˆØˆ�Œ�s˜DÑ!Ô!Ð!Ø×Ò˜˜Qœ Ñ#Ô#Ð#Ð#Ð#r   c                 óæ   —  t          j        d¦  «        t          ¦  «        }t          j        dt          j        ¬¦  «        } |d         |d¦  «         |                      |d         d¦  «         d S )Nzvoid(int32[:], uint64)r   r  r  l        @ r   rõ   r‹  r  s      r   Útest_popc_u8zTestCudaIntrinsic.test_popc_u80  sh   € Ø5•4”8Ð4Ñ5Ô5µkÑBÔBˆÝŒh�q¥¤Ð)Ñ)Ô)ˆØˆ�Œ�s˜NÑ+Ô+Ð+Ø×Ò˜˜Qœ Ñ#Ô#Ð#Ð#Ð#r   c                 óþ   —  t          j        d¦  «        t          ¦  «        }t          j        dt          j        ¬¦  «        } |d         |ddd¦  «         t          j                             |d         d	¦  «         d S )
Nzvoid(f4[:], f4, f4, f4)r   r  r  ç       @ç      @ç      @r   r  )r   r  rJ   r  r(  Úfloat32rg  Úassert_allcloser  s      r   Útest_fma_f4zTestCudaIntrinsic.test_fma_f46  ón   € Ø6•4”8Ð5Ñ6Ô6µzÑBÔBˆÝŒh�q¥¤
Ð+Ñ+Ô+ˆØˆ�Œ�s˜B  BÑ'Ô'Ð'Ý
Œ
×"Ò" 3 q¤6¨9Ñ5Ô5Ð5Ð5Ð5r   c                 óþ   —  t          j        d¦  «        t          ¦  «        }t          j        dt          j        ¬¦  «        } |d         |ddd¦  «         t          j                             |d         d	¦  «         d S )
Nzvoid(f8[:], f8, f8, f8)r   r  r  r�  r‘  r’  r   r  )r   r  rJ   r  r(  Úfloat64rg  r”  r  s      r   Útest_fma_f8zTestCudaIntrinsic.test_fma_f8<  r–  r   c                 ó†  —  t          j        d¦  «        t          ¦  «        }t          j        dt          j        ¬¦  «        }t          j        dgt          j        ¬¦  «        }t          j        dgt          j        ¬¦  «        } |d         |||¦  «         t          j                             |d         ||z   ¦  «         d S ©Núvoid(f2[:], f2[:], f2[:])r   r  r‘  r’  r  r   )	r   r  rP   r  r(  Úfloat16Úarrayrg  r”  ©r  r  r   Úarg1Úarg2s        r   Ú	test_haddzTestCudaIntrinsic.test_haddB  ó    € à8•4”8Ð7Ñ8Ô8½ÑEÔEˆÝŒh�q¥¤
Ð+Ñ+Ô+ˆÝŒx˜˜¥B¤JÐ/Ñ/Ô/ˆÝŒx˜˜¥B¤JÐ/Ñ/Ô/ˆØˆ�Œ�s˜D $Ñ'Ô'Ð'Ý
Œ
×"Ò" 3 q¤6¨4°$©;Ñ7Ô7Ð7Ð7Ð7r   c                 óV  —  t          j        d¦  «        t          ¦  «        }t          j        dt          j        ¬¦  «        }t          j        d¦  «        }t          j        d¦  «        } |d         |||¦  «         ||z   }t          j                             |d         |¦  «         d S )Núvoid(f2[:], f2, f2)r   r  çJØMû!	@r‘  r  r   )r   r  rS   r  r(  r�  rg  r”  ©r  r  r   r   r¡  Úrefs         r   Útest_hadd_scalarz"TestCudaIntrinsic.test_hadd_scalarK  s’   € à2•4”8Ð1Ñ2Ô2Õ3EÑFÔFˆÝŒh�q¥¤
Ð+Ñ+Ô+ˆÝŒz˜)Ñ$Ô$ˆÝŒz˜"‰~Œ~ˆØˆ�Œ�s˜D $Ñ'Ô'Ð'Ø�T‰kˆÝ
Œ
×"Ò" 3 q¤6¨3Ñ/Ô/Ð/Ð/Ð/r   z(Compilation unsupported in the simulatorc                 óž   — t           d d …         t           t           f}t          t          |d¬¦  «        \  }}|                      d|¦  «         d S )N©r%  r÷   ©Úcczadd.f16)r   r   rS   ÚassertIn©r  Úargsrn  Ú_s       r   Útest_hadd_ptxzTestCudaIntrinsic.test_hadd_ptxU  óH   € å�1�1�1”•r�2ˆˆÝÕ/°¸&ÐAÑAÔA‰ˆˆQØ�Š�i Ñ%Ô%Ð%Ð%Ð%r   c                 óÐ  —  t          j        d¦  «        t          ¦  «        }t          j        dt          j        ¬¦  «        }t          j        dgt          j        ¬¦  «        }t          j        dgt          j        ¬¦  «        }t          j        dgt          j        ¬¦  «        } |d         ||||¦  «         t          j                             |d         ||z  |z   ¦  «         d S )	Nz void(f2[:], f2[:], f2[:], f2[:])r   r  r�  r‘  r’  r  r   )	r   r  rX   r  r(  r�  rž  rg  r”  )r  r  r   r   r¡  Úarg3s         r   Ú	test_hfmazTestCudaIntrinsic.test_hfma[  s¿   € à?•4”8Ð>Ñ?Ô?ÅÑLÔLˆÝŒh�q¥¤
Ð+Ñ+Ô+ˆÝŒx˜˜¥B¤JÐ/Ñ/Ô/ˆÝŒx˜˜¥B¤JÐ/Ñ/Ô/ˆÝŒx˜˜¥B¤JÐ/Ñ/Ô/ˆØˆ�Œ�s˜D $¨Ñ-Ô-Ð-Ý
Œ
×"Ò" 3 q¤6¨4°$©;¸Ñ+=Ñ>Ô>Ð>Ð>Ð>r   c                 ó†  —  t          j        d¦  «        t          ¦  «        }t          j        dt          j        ¬¦  «        }t          j        d¦  «        }t          j        d¦  «        }t          j        d¦  «        } |d         ||||¦  «         ||z  |z   }t          j                             |d         |¦  «         d S )	Nzvoid(f2[:], f2, f2, f2)r   r  r�  r‘  r’  r  r   )r   r  rZ   r  r(  r�  rg  r”  )r  r  r   r   r¡  rµ  r¨  s          r   Útest_hfma_scalarz"TestCudaIntrinsic.test_hfma_scalare  s¤   € à6•4”8Ð5Ñ6Ô6Õ7IÑJÔJˆÝŒh�q¥¤
Ð+Ñ+Ô+ˆÝŒz˜"‰~Œ~ˆÝŒz˜"‰~Œ~ˆÝŒz˜"‰~Œ~ˆØˆ�Œ�s˜D $¨Ñ-Ô-Ð-Ø�T‰k˜DÑ ˆÝ
Œ
×"Ò" 3 q¤6¨3Ñ/Ô/Ð/Ð/Ð/r   c                 óª   — t           d d …         t           t           t           f}t          t          |d¬¦  «        \  }}|                      d|¦  «         d S )Nr«  r¬  z
fma.rn.f16)r   r   rZ   r®  r¯  s       r   Útest_hfma_ptxzTestCudaIntrinsic.test_hfma_ptxp  sK   € å�1�1�1”•r�2�rÐ"ˆÝÕ/°¸&ÐAÑAÔA‰ˆˆQØ�Š�l CÑ(Ô(Ð(Ð(Ð(r   c                 ó†  —  t          j        d¦  «        t          ¦  «        }t          j        dt          j        ¬¦  «        }t          j        dgt          j        ¬¦  «        }t          j        dgt          j        ¬¦  «        } |d         |||¦  «         t          j                             |d         ||z
  ¦  «         d S r›  )	r   r  r^   r  r(  r�  rž  rg  r”  rŸ  s        r   Ú	test_hsubzTestCudaIntrinsic.test_hsubv  r£  r   c                 óV  —  t          j        d¦  «        t          ¦  «        }t          j        dt          j        ¬¦  «        }t          j        d¦  «        }t          j        d¦  «        } |d         |||¦  «         ||z
  }t          j                             |d         |¦  «         d S ©Nr¥  r   r  r¦  g…ëQ¸ù?r  r   )r   r  r`   r  r(  r�  rg  r”  r§  s         r   Útest_hsub_scalarz"TestCudaIntrinsic.test_hsub_scalar  ó”   € à2•4”8Ð1Ñ2Ô2Õ3EÑFÔFˆÝŒh�q¥¤
Ð+Ñ+Ô+ˆÝŒz˜)Ñ$Ô$ˆÝŒz˜$ÑÔˆØˆ�Œ�s˜D $Ñ'Ô'Ð'Ø�T‰kˆÝ
Œ
×"Ò" 3 q¤6¨3Ñ/Ô/Ð/Ð/Ð/r   c                 óž   — t           d d …         t           t           f}t          t          |d¬¦  «        \  }}|                      d|¦  «         d S )Nr«  r¬  zsub.f16)r   r   r`   r®  r¯  s       r   Útest_hsub_ptxzTestCudaIntrinsic.test_hsub_ptx‰  r³  r   c                 ó„  —  t          j        ¦   «         t          ¦  «        }t          j        dt          j        ¬¦  «        }t          j        dgt          j        ¬¦  «        }t          j        dgt          j        ¬¦  «        } |d         |||¦  «         t          j                             |d         ||z  ¦  «         d S )Nr   r  r‘  r’  r  r   )	r   r  rd   r  r(  r�  rž  rg  r”  rŸ  s        r   Ú	test_hmulzTestCudaIntrinsic.test_hmul�  s›   € à•4”8‘:”:�kÑ*Ô*ˆÝŒh�q¥¤
Ð+Ñ+Ô+ˆÝŒx˜˜¥B¤JÐ/Ñ/Ô/ˆÝŒx˜˜¥B¤JÐ/Ñ/Ô/ˆØˆ�Œ�s˜D $Ñ'Ô'Ð'Ý
Œ
×"Ò" 3 q¤6¨4°$©;Ñ7Ô7Ð7Ð7Ð7r   c                 óV  —  t          j        d¦  «        t          ¦  «        }t          j        dt          j        ¬¦  «        }t          j        d¦  «        }t          j        d¦  «        } |d         |||¦  «         ||z  }t          j                             |d         |¦  «         d S r¾  )r   r  rf   r  r(  r�  rg  r”  r§  s         r   Útest_hmul_scalarz"TestCudaIntrinsic.test_hmul_scalar˜  rÀ  r   c                 óž   — t           d d …         t           t           f}t          t          |d¬¦  «        \  }}|                      d|¦  «         d S )Nr«  r¬  zmul.f16)r   r   rf   r®  r¯  s       r   Útest_hmul_ptxzTestCudaIntrinsic.test_hmul_ptx¢  r³  r   c                 óV  —  t          j        d¦  «        t          ¦  «        }t          j        dt          j        ¬¦  «        }t          j        d¦  «        }t          j        d¦  «        } |d         |||¦  «         ||z  }t          j                             |d         |¦  «         d S r¾  )r   r  ri   r  r(  r�  rg  r”  r§  s         r   Útest_hdiv_scalarz"TestCudaIntrinsic.test_hdiv_scalar¨  s”   € à2•4”8Ð1Ñ2Ô2Õ3EÑFÔFˆÝŒh�q¥¤
Ð+Ñ+Ô+ˆÝŒz˜)Ñ$Ô$ˆÝŒz˜$ÑÔˆàˆ�Œ�s˜D $Ñ'Ô'Ð'Ø�T‰kˆÝ
Œ
×"Ò" 3 q¤6¨3Ñ/Ô/Ð/Ð/Ð/r   c                 ó  —  t          j        d¦  «        t          ¦  «        }t          j                             ddd¬¦  «                             t          j        ¦  «        }t          j                             ddd¬¦  «                             t          j        ¦  «        }t          j        |t          j        ¬¦  «        } | 	                    |j
        ¦  «        |||¦  «         ||z  }t          j                             ||¦  «         d S )Nrœ  i  ÿÿéáÿ  iô  ©rk   r  )r   r  rn   r  r  ÚrandintÚastyper�  Ú
zeros_likeÚforallrk   rg  r”  )r  r  Úarry1Úarry2r   r¨  s         r   Ú	test_hdivzTestCudaIntrinsic.test_hdiv³  sÔ   € à8•4”8Ð7Ñ8Ô8Õ9KÑLÔLˆÝ”	×!Ò! &¨%°cÐ!Ñ:Ô:×AÒAÅ"Ä*ÑMÔMˆÝ”	×!Ò! &¨%°cÐ!Ñ:Ô:×AÒAÅ"Ä*ÑMÔMˆÝŒm˜E­¬Ð4Ñ4Ô4ˆà!ˆ�Š˜œÑ!Ô! # u¨eÑ4Ô4Ð4Ø�e‰mˆÝ
Œ
×"Ò" 3¨Ñ,Ô,Ð,Ð,Ð,r   c                 ó>  —  t          j        d¦  «        t          ¦  «        }t          j        dt          j        ¬¦  «        }t          j        dgt          j        ¬¦  «        } |d         ||¦  «         t          j                             |d         | ¦  «         d S )Núvoid(f2[:], f2[:])r   r  r‘  r  r   )	r   r  rs   r  r(  r�  rž  rg  r”  ©r  r  r   r   s       r   Ú	test_hnegzTestCudaIntrinsic.test_hneg¾  s„   € à1•4”8Ð0Ñ1Ô1µ+Ñ>Ô>ˆÝŒh�q¥¤
Ð+Ñ+Ô+ˆÝŒx˜˜¥B¤JÐ/Ñ/Ô/ˆØˆ�Œ�s˜DÑ!Ô!Ð!Ý
Œ
×"Ò" 3 q¤6¨D¨5Ñ1Ô1Ð1Ð1Ð1r   c                 ó(  —  t          j        d¦  «        t          ¦  «        }t          j        dt          j        ¬¦  «        }t          j        d¦  «        } |d         ||¦  «         | }t          j                             |d         |¦  «         d S )Núvoid(f2[:], f2)r   r  r¦  r  r   )r   r  rv   r  r(  r�  rg  r”  ©r  r  r   r   r¨  s        r   Útest_hneg_scalarz"TestCudaIntrinsic.test_hneg_scalarÆ  s�   € à.•4”8Ð-Ñ.Ô.Õ/AÑBÔBˆÝŒh�q¥¤
Ð+Ñ+Ô+ˆÝŒz˜)Ñ$Ô$ˆØˆ�Œ�s˜DÑ!Ô!Ð!ØˆeˆÝ
Œ
×"Ò" 3 q¤6¨3Ñ/Ô/Ð/Ð/Ð/r   c                 ó’   — t           d d …         t           f}t          t          |d¬¦  «        \  }}|                      d|¦  «         d S )Nr«  r¬  zneg.f16)r   r   rv   r®  r¯  s       r   Útest_hneg_ptxzTestCudaIntrinsic.test_hneg_ptxÏ  óF   € å�1�1�1”•rˆ{ˆÝÕ/°¸&ÐAÑAÔA‰ˆˆQØ�Š�i Ñ%Ô%Ð%Ð%Ð%r   c                 óT  —  t          j        ¦   «         t          ¦  «        }t          j        dt          j        ¬¦  «        }t          j        dgt          j        ¬¦  «        } |d         ||¦  «         t          j                             |d         t          |¦  «        ¦  «         d S )Nr   r  ç      Àr  r   )
r   r  r{   r  r(  r�  rž  rg  r”  Úabsr×  s       r   Ú	test_habszTestCudaIntrinsic.test_habsÕ  sƒ   € à•4”8‘:”:�kÑ*Ô*ˆÝŒh�q¥¤
Ð+Ñ+Ô+ˆÝŒx˜˜¥R¤ZÐ0Ñ0Ô0ˆØˆ�Œ�s˜DÑ!Ô!Ð!Ý
Œ
×"Ò" 3 q¤6­3¨t©9¬9Ñ5Ô5Ð5Ð5Ð5r   c                 ó@  —  t          j        d¦  «        t          ¦  «        }t          j        dt          j        ¬¦  «        }t          j        d¦  «        } |d         ||¦  «         t          |¦  «        }t          j                             |d         |¦  «         d S )NrÚ  r   r  gJØMû!	Àr  r   )	r   r  r}   r  r(  r�  râ  rg  r”  rÛ  s        r   Útest_habs_scalarz"TestCudaIntrinsic.test_habs_scalarÝ  s…   € à.•4”8Ð-Ñ.Ô.Õ/AÑBÔBˆÝŒh�q¥¤
Ð+Ñ+Ô+ˆÝŒz˜*Ñ%Ô%ˆØˆ�Œ�s˜DÑ!Ô!Ð!Ý�$‰iŒiˆÝ
Œ
×"Ò" 3 q¤6¨3Ñ/Ô/Ð/Ð/Ð/r   c                 ó’   — t           d d …         t           f}t          t          |d¬¦  «        \  }}|                      d|¦  «         d S )Nr«  r¬  zabs.f16)r   r   r}   r®  r¯  s       r   Útest_habs_ptxzTestCudaIntrinsic.test_habs_ptxæ  rß  r   c                 ó&  — t           t          t          t          t          t
          t          t          t          t          t          t          f}t          t          f}t          j        t          j        t          j        t          j        t          j        t          j        t          j        t          j        t          j        t          j        t          j        t4          f}t          j        t          j        f}d}t          j                             d¦  «         t          j                             dd|¬¦  «                              t          j!        ¦  «        }t          j"        |¦  «        }tF          j$        rdnd}tK          ||¦  «        D ]�\  }	}
|  &                    |
¬¦  «        5   tO          j(        d¦  «        |	¦  «        }	 |	d|f         ||¦  «          |
|t          j!        ¬	¦  «        }t          j)         *                    |||¬
¦  «         d d d ¦  «         n# 1 swxY w Y   Œžt          j                             dd|¬¦  «                              t          j!        ¦  «        }tK          ||¦  «        D ]�\  }	}
|  &                    |
¬¦  «        5   tO          j(        d¦  «        |	¦  «        }	 |	d|f         ||¦  «          |
|t          j!        ¬	¦  «        }t          j)         *                    |||¬
¦  «         d d d ¦  «         n# 1 swxY w Y   Œžd S )Nr^  r   rÌ  rÍ  gü©ñÒMbP?gH¯¼šò×z>©ÚfnrÖ  r  )Úrtolr  )+r±   rµ   r¸   r»   r¿   rÉ   rÓ   rÖ   rÙ   rÜ   rß   rÌ   rÃ   rÆ   r  ÚsinÚcosÚlogÚlog2Úlog10ÚsqrtÚceilÚfloorÚ
reciprocalÚtruncÚrintrÐ   r"  Úexp2r  r  rÎ  rÏ  r�  rÐ  r   ÚENABLE_CUDASIMrs  ÚsubTestr   r  rg  r”  )r  ÚkernelsÚexp_kernelsÚexpected_functionsÚexpected_exp_functionsr!  r   r›   rë  Úkernelrê  rm  Úx2s                r   Útest_fp16_intrinsics_commonz-TestCudaIntrinsic.test_fp16_intrinsics_commonì  s#  € å¥Ý¥­mÝ¥­}Ý¥­|Ý ð	"ˆõ
 #¥LÐ1ˆÝ œf¥b¤fÝ œf¥b¤g­r¬xÝ œg¥r¤wµ´Ý œm­R¬Xµr´wÝ*ð	,Ðõ
 #%¤&­"¬'Ð!2Ðð ˆÝ
Œ	�Š�qÑÔÐÝŒI×Ò˜a ¨QÐÑ/Ô/×6Ò6µr´zÑBÔBˆÝŒM˜!ÑÔˆÝÔ,Ð6ˆtˆt°$ˆÝ˜gÐ'9Ñ:Ô:ð 	Cð 	C‰JˆF�BØ—’ �Ñ$Ô$ð Cð CØ7�œÐ"6Ñ7Ô7¸Ñ?Ô?�Ø��q˜�s”˜A˜qÑ!Ô!Ð!Ø˜2˜a¥r¤zÐ2Ñ2Ô2�Ý”
×*Ò*¨1¨h¸TÐ*ÑBÔBÐBð	Cð Cð Cñ Cô Cð Cð Cð Cð Cð Cð Cøøøð Cð Cð Cð Cøõ ŒY×Ò˜q "¨1ÐÑ-Ô-×4Ò4µR´ZÑ@Ô@ˆÝ˜kÐ+AÑBÔBð 	Cð 	C‰JˆF�BØ—’ �Ñ$Ô$ð Cð CØ7�œÐ"6Ñ7Ô7¸Ñ?Ô?�Ø��q˜�s”˜A˜rÑ"Ô"Ð"Ø˜2˜b­¬
Ð3Ñ3Ô3�Ý”
×*Ò*¨1¨h¸TÐ*ÑBÔBÐBð	Cð Cð Cñ Cô Cð Cð Cð Cð Cð Cð Cøøøð Cð Cð Cð Cøð	Cð 	Cs&   Æ!A+HÈH	ÈH	ÊA+LÌL		ÌL		c                 óŽ  — t          j        ¦   «         d„ ¦   «         }d}t          j                             d¦  «         t          j                             |¦  «                             t          j        ¦  «        }t          j        |¦  «        } |d|f         ||¦  «         t          j	         
                    |d|z  ¦  «         d S )Nc                 ó¨   — t          j        d¦  «        }|t          | ¦  «        k     r*t           j                             ||         ¦  «        | |<   d S d S r   )r   r&   r®   rM   Úhexp10r°   s      r   Úhexp10_vectorsz5TestCudaIntrinsic.test_hexp10.<locals>.hexp10_vectors  sF   € å”	˜!‘”ˆAà•3�q‘6”6ŠzˆzÝ”y×'Ò'¨¨!¬Ñ-Ô-��!‘��ð ˆzr   r^  r   r  )r   r  r  r  r  ÚrandrÏ  r�  rÐ  rg  r”  )r  r  r!  r   r›   s        r   Útest_hexp10zTestCudaIntrinsic.test_hexp10  s­   € å	Œ‰Œð	.ð 	.ñ 
Œð	.ð ˆÝ
Œ	�Š�qÑÔÐÝŒI�NŠN˜1ÑÔ×$Ò$¥R¤ZÑ0Ô0ˆÝŒM˜!ÑÔˆð 	ˆ�q˜!�tÔ˜Q Ñ"Ô"Ð"Ý
Œ
×"Ò" 1 b¨A¡gÑ.Ô.Ð.Ð.Ð.r   c                 óú  — t           t          t          t          t          t
          f}t          j        t          j        t          j	        t          j
        t          j        t          j        f}t          ||¦  «        D �]~\  }}|                      |¬¦  «        5   t          j        d¦  «        |¦  «        }t#          j        dt"          j        ¬¦  «        }t#          j        dt"          j        ¬¦  «        }t#          j        d¦  «        }t#          j        d¦  «        }	t#          j        d¦  «        }
 |d         ||	|	¦  «          ||	|	¦  «        }|                      ||d	         ¦  «          |d         ||	|
¦  «          ||	|
¦  «        }|                      ||d	         ¦  «          |d         ||	|¦  «          ||	|¦  «        }|                      ||d	         ¦  «         d d d ¦  «         n# 1 swxY w Y   �Œ€d S )
N)Úopzvoid(b1[:], f2, f2)r   r  r*   r÷   rõ   r  r   )r€   r„   r‡   rŠ   r�   r‘   ÚoperatorÚeqÚneÚgeÚgtÚleÚltrs  rù  r   r  r  r(  r†  r�  rM  )r  ÚfnsÚopsrê  r  rþ  rm  Úgotr¡  rµ  Úarg4s              r   Útest_fp16_comparisonz&TestCudaIntrinsic.test_fp16_comparison#  s  € å Õ"3Õ5FÝ Õ"3Õ5FðHˆåŒ{�HœK­¬ÝŒ{�HœK­¬ð6ˆõ ˜#˜s‘m”mð 	3ñ 	3‰FˆB�Ø—’ �Ñ$Ô$ð 3ð 3Ø8�œÐ"7Ñ8Ô8¸Ñ<Ô<�åœ8 A­R¬XÐ6Ñ6Ô6�Ý”h˜q­¬Ð1Ñ1Ô1�Ý”z !‘}”}�Ý”z !‘}”}�Ý”z !‘}”}�ð ��t”˜S $¨Ñ-Ô-Ð-Ø˜2˜d D™>œ>�Ø× Ò  ¨3¨q¬6Ñ2Ô2Ð2ð ��t”˜S $¨Ñ-Ô-Ð-Ø˜2˜d D™>œ>�Ø× Ò  ¨3¨q¬6Ñ2Ô2Ð2ð ��t”˜S $¨Ñ-Ô-Ð-Ø˜2˜d D™>œ>�Ø× Ò  ¨3¨q¬6Ñ2Ô2Ð2ð-3ð 3ð 3ñ 3ô 3ð 3ð 3ð 3ð 3ð 3ð 3øøøð 3ð 3ð 3ð 3ùð	3ð 	3s   ÂEG.Ç.G2	Ç5G2	c                 óü  — t           t          t          t          t          f}|D ]Ø}|                      |¬¦  «        5   t          j        d¦  «        |¦  «        }t          j	        dt          j
        ¬¦  «        }t          j        d¦  «        }t          j        d¦  «        }t          j        d¦  «        } |d         ||||¦  «         |                      |d	         ¦  «         d d d ¦  «         n# 1 swxY w Y   ŒÙd S )
Nré  zvoid(b1[:], f2, f2, f2)r   r  r�  r‘  r’  r  r   )rœ   rž   r¡   r£   r¦   rù  r   r  r  r(  r†  r�  r  )r  Ú	functionsrê  r  r   r   r¡  rµ  s           r   Ú!test_multiple_float16_comparisonsz3TestCudaIntrinsic.test_multiple_float16_comparisonsC  s$  € å)Ý)Ý)Ý)Ý)ð	+ˆ	ð
 ð 	(ð 	(ˆBØ—’ �Ñ$Ô$ð (ð (Ø>�4œ8Ð$=Ñ>Ô>¸rÑBÔB�Ý”h˜q­¬Ð1Ñ1Ô1�Ý”z "‘~”~�Ý”z "‘~”~�Ý”z "‘~”~�Ø�˜”˜s D¨$°Ñ5Ô5Ð5Ø—’  A¤Ñ'Ô'Ð'ð(ð (ð (ñ (ô (ð (ð (ð (ð (ð (ð (øøøð (ð (ð (ð (øð	(ð 	(s   »B)C0Ã0C4	Ã7C4	c                 óæ  —  t          j        d¦  «        t          ¦  «        }t          j        dt          j        ¬¦  «        }t          j        d¦  «        }t          j        d¦  «        } |d         |||¦  «         t          j                             |d         |¦  «         t          j        d¦  «        } |d         |||¦  «         t          j                             |d         |¦  «         d S ©	Nr¥  r   r  r‘  r’  r  r   g      @)r   r  r©   r  r(  r�  rg  r”  rŸ  s        r   Ú	test_hmaxzTestCudaIntrinsic.test_hmaxT  óÈ   € à2•4”8Ð1Ñ2Ô2Õ3EÑFÔFˆÝŒh�q¥¤
Ð+Ñ+Ô+ˆÝŒz˜"‰~Œ~ˆÝŒz˜"‰~Œ~ˆØˆ�Œ�s˜D $Ñ'Ô'Ð'Ý
Œ
×"Ò" 3 q¤6¨4Ñ0Ô0Ð0ÝŒz˜"‰~Œ~ˆØˆ�Œ�s˜D $Ñ'Ô'Ð'Ý
Œ
×"Ò" 3 q¤6¨4Ñ0Ô0Ð0Ð0Ð0r   c                 óæ  —  t          j        d¦  «        t          ¦  «        }t          j        dt          j        ¬¦  «        }t          j        d¦  «        }t          j        d¦  «        } |d         |||¦  «         t          j                             |d         |¦  «         t          j        d¦  «        } |d         |||¦  «         t          j                             |d         |¦  «         d S r  )r   r  r¬   r  r(  r�  rg  r”  rŸ  s        r   Ú	test_hminzTestCudaIntrinsic.test_hmin`  r  r   c                 ó  —  t          j        d¦  «        t          ¦  «        }t          j        dt          j        ¬¦  «        }d} |d         ||¦  «         t          j                             |d         |dz  ¦  «         d S )Nzvoid(float32[:], float32)r   r  r�  r  r   çUUUUUUÕ?)r   r  râ   r  r(  r“  rg  r”  ©r  r  r   Úcbrt_args       r   Útest_cbrt_f32zTestCudaIntrinsic.test_cbrt_f32l  ót   € Ø8•4”8Ð7Ñ8Ô8½ÑEÔEˆÝŒh�q¥¤
Ð+Ñ+Ô+ˆØˆØˆ�Œ�s˜HÑ%Ô%Ð%Ý
Œ
×"Ò" 3 q¤6¨8¸Ñ+>Ñ?Ô?Ð?Ð?Ð?r   c                 ó  —  t          j        d¦  «        t          ¦  «        }t          j        dt          j        ¬¦  «        }d} |d         ||¦  «         t          j                             |d         |dz  ¦  «         d S )Nzvoid(float64[:], float64)r   r  g      @r  r   r  )r   r  râ   r  r(  r˜  rg  r”  r   s       r   Útest_cbrt_f64zTestCudaIntrinsic.test_cbrt_f64s  r#  r   c                 óæ   —  t          j        d¦  «        t          ¦  «        }t          j        dt          j        ¬¦  «        } |d         |d¦  «         |                      |d         d¦  «         d S )Nzvoid(uint32[:], uint32)r   r  r  ið0  r   i  )r   r  rå   r  r(  Úuint32rM  r  s      r   Útest_brev_u4zTestCudaIntrinsic.test_brev_u4z  sh   € Ø6•4”8Ð5Ñ6Ô6µ{ÑCÔCˆÝŒh�q¥¤	Ð*Ñ*Ô*ˆØˆ�Œ�s˜JÑ'Ô'Ð'Ø×Ò˜˜Qœ Ñ,Ô,Ð,Ð,Ð,r   z.only get given a Python "int", assumes 32 bitsc                 óæ   —  t          j        d¦  «        t          ¦  «        }t          j        dt          j        ¬¦  «        } |d         |d¦  «         |                      |d         d¦  «         d S )Nzvoid(uint64[:], uint64)r   r  r  l   ð0  ÀC r   l       `x)r   r  rå   r  r(  r[  rM  r  s      r   Útest_brev_u8zTestCudaIntrinsic.test_brev_u8€  sj   € à6•4”8Ð5Ñ6Ô6µ{ÑCÔCˆÝŒh�q¥¤	Ð*Ñ*Ô*ˆØˆ�Œ�sÐ.Ñ/Ô/Ð/Ø×Ò˜˜QœÐ!3Ñ4Ô4Ð4Ð4Ð4r   c                 óæ   —  t          j        d¦  «        t          ¦  «        }t          j        dt          j        ¬¦  «        } |d         |d¦  «         |                      |d         d¦  «         d S )Núvoid(int32[:], int32)r   r  r  é   r   é   ©r   r  rè   r  r(  r  rM  r  s      r   Útest_clz_i4zTestCudaIntrinsic.test_clz_i4‡  sh   € Ø4•4”8Ð3Ñ4Ô4µZÑ@Ô@ˆÝŒh�q¥¤Ð)Ñ)Ô)ˆØˆ�Œ�s˜JÑ'Ô'Ð'Ø×Ò˜˜Qœ Ñ$Ô$Ð$Ð$Ð$r   c                 óæ   —  t          j        d¦  «        t          ¦  «        }t          j        dt          j        ¬¦  «        } |d         |d¦  «         |                      |d         d¦  «         dS )	aº  
        Although the CUDA Math API
        (http://docs.nvidia.com/cuda/cuda-math-api/group__CUDA__MATH__INTRINSIC__INT.html)
        only says int32 & int64 arguments are supported in C code, the LLVM
        IR input supports i8, i16, i32 & i64 (LLVM doesn't have a concept of
        unsigned integers, just unsigned operations on integers).
        http://docs.nvidia.com/cuda/nvvm-ir-spec/index.html#bit-manipulations-intrinics
        r‰  r   r  r  r-  r   r.  Nr/  r  s      r   Útest_clz_u4zTestCudaIntrinsic.test_clz_u4�  sj   € ð 6•4”8Ð4Ñ5Ô5µjÑAÔAˆÝŒh�q¥¤Ð)Ñ)Ô)ˆØˆ�Œ�s˜JÑ'Ô'Ð'Ø×Ò˜˜Qœ Ñ$Ô$Ð$Ð$Ð$r   c                 óæ   —  t          j        d¦  «        t          ¦  «        }t          j        dt          j        ¬¦  «        } |d         |d¦  «         |                      |d         d¦  «         d S ©Nr,  r   r  r  l   ÿÿ r   r/  r  s      r   Útest_clz_i4_1sz TestCudaIntrinsic.test_clz_i4_1s›  óh   € Ø4•4”8Ð3Ñ4Ô4µZÑ@Ô@ˆÝŒh�q¥¤Ð)Ñ)Ô)ˆØˆ�Œ�s˜JÑ'Ô'Ð'Ø×Ò˜˜Qœ Ñ#Ô#Ð#Ð#Ð#r   c                 óè   —  t          j        d¦  «        t          ¦  «        }t          j        dt          j        ¬¦  «        } |d         |d¦  «         |                      |d         dd¦  «         d S )Nr,  r   r  r  r   r^  úCUDA semanticsr/  r  s      r   Útest_clz_i4_0sz TestCudaIntrinsic.test_clz_i4_0s¡  sk   € Ø4•4”8Ð3Ñ4Ô4µZÑ@Ô@ˆÝŒh�q¥¤Ð)Ñ)Ô)ˆØˆ�Œ�s˜CÑ Ô Ð Ø×Ò˜˜Qœ Ð%5Ñ6Ô6Ð6Ð6Ð6r   c                 óæ   —  t          j        d¦  «        t          ¦  «        }t          j        dt          j        ¬¦  «        } |d         |d¦  «         |                      |d         d¦  «         d S )Núvoid(int32[:], int64)r   r  r  é   r   é/   r/  r  s      r   Útest_clz_i8zTestCudaIntrinsic.test_clz_i8§  si   € à4•4”8Ð3Ñ4Ô4µZÑ@Ô@ˆÝŒh�q¥¤Ð)Ñ)Ô)ˆØˆ�Œ�sÐ-Ñ.Ô.Ð.Ø×Ò˜˜Qœ Ñ$Ô$Ð$Ð$Ð$r   c                 óB  —  t          j        d¦  «        t          ¦  «        }t          j        dt          j        ¬¦  «        } |d         |d¦  «         |                      |d         d¦  «          |d         |d¦  «         |                      |d         d	¦  «         d S )
Nr,  r   r  r  r-  r   é   ì        r^  ©r   r  rì   r  r(  r  rM  r  s      r   Útest_ffs_i4zTestCudaIntrinsic.test_ffs_i4®  s˜   € Ø4•4”8Ð3Ñ4Ô4µZÑ@Ô@ˆÝŒh�q¥¤Ð)Ñ)Ô)ˆØˆ�Œ�s˜JÑ'Ô'Ð'Ø×Ò˜˜Qœ Ñ$Ô$Ð$Øˆ�Œ�s˜JÑ'Ô'Ð'Ø×Ò˜˜Qœ Ñ$Ô$Ð$Ð$Ð$r   c                 óB  —  t          j        d¦  «        t          ¦  «        }t          j        dt          j        ¬¦  «        } |d         |d¦  «         |                      |d         d¦  «          |d         |d¦  «         |                      |d         d	¦  «         d S )
Nr‰  r   r  r  r-  r   r@  rA  r^  rB  r  s      r   Útest_ffs_u4zTestCudaIntrinsic.test_ffs_u4¶  s˜   € Ø5•4”8Ð4Ñ5Ô5µjÑAÔAˆÝŒh�q¥¤Ð)Ñ)Ô)ˆØˆ�Œ�s˜JÑ'Ô'Ð'Ø×Ò˜˜Qœ Ñ$Ô$Ð$Øˆ�Œ�s˜JÑ'Ô'Ð'Ø×Ò˜˜Qœ Ñ$Ô$Ð$Ð$Ð$r   c                 óæ   —  t          j        d¦  «        t          ¦  «        }t          j        dt          j        ¬¦  «        } |d         |d¦  «         |                      |d         d¦  «         d S r4  rB  r  s      r   Útest_ffs_i4_1sz TestCudaIntrinsic.test_ffs_i4_1s¾  r6  r   c                 óæ   —  t          j        d¦  «        t          ¦  «        }t          j        dt          j        ¬¦  «        } |d         |d¦  «         |                      |d         d¦  «         d S )Nr,  r   r  r  r   rB  r  s      r   Útest_ffs_i4_0sz TestCudaIntrinsic.test_ffs_i4_0sÄ  sh   € Ø4•4”8Ð3Ñ4Ô4µZÑ@Ô@ˆÝŒh�q¥¤Ð)Ñ)Ô)ˆØˆ�Œ�s˜CÑ Ô Ð Ø×Ò˜˜Qœ Ñ#Ô#Ð#Ð#Ð#r   c                 óB  —  t          j        d¦  «        t          ¦  «        }t          j        dt          j        ¬¦  «        } |d         |d¦  «         |                      |d         d¦  «          |d         |d¦  «         |                      |d         d	¦  «         d S )
Nr;  r   r  r  r<  r   r_  l        é!   rB  r  s      r   Útest_ffs_i8zTestCudaIntrinsic.test_ffs_i8Ê  s™   € à4•4”8Ð3Ñ4Ô4µZÑ@Ô@ˆÝŒh�q¥¤Ð)Ñ)Ô)ˆØˆ�Œ�sÐ-Ñ.Ô.Ð.Ø×Ò˜˜Qœ Ñ$Ô$Ð$Øˆ�Œ�s˜KÑ(Ô(Ð(Ø×Ò˜˜Qœ Ñ$Ô$Ð$Ð$Ð$r   c                 ó|  —  t          j        d¦  «        t          ¦  «        }d}t          j        |dz  t          j        ¬¦  «        }t          j        t          j        dt          j        ¬¦  «        |¦  «        } |d|dz  f         |¦  «         |                      t          j	        ||k    ¦  «        ¦  «         d S )Nr  r*   r^  r  r   )
r   r  rÿ   r  r(  r  Útiler  r  r   )r  r  Úcountr   r"  s        r   Útest_simple_laneidz$TestCudaIntrinsic.test_simple_laneidÓ  s    € Ø-•4”8Ð,Ñ-Ô-­mÑ<Ô<ˆØˆÝŒh�u˜r‘z­¬Ð2Ñ2Ô2ˆÝŒg•b”i ­"¬(Ð3Ñ3Ô3°UÑ;Ô;ˆØˆ��E˜B‘J�Ô Ñ$Ô$Ð$Ø�Š�œ˜s cšzÑ*Ô*Ñ+Ô+Ð+Ð+Ð+r   c                 óæ   —  t          j        d¦  «        t          ¦  «        }t          j        dt          j        ¬¦  «        } |d         |¦  «         |                      |d         dd¦  «         d S )Nr  r   r  r  r   r^  r8  )r   r  r  r  r(  r  rM  r  s      r   Útest_simple_warpsizez&TestCudaIntrinsic.test_simple_warpsizeÛ  si   € Ø-•4”8Ð,Ñ-Ô-­oÑ>Ô>ˆÝŒh�q¥¤Ð)Ñ)Ô)ˆØˆ�Œ�sÑÔÐØ×Ò˜˜Qœ Ð%5Ñ6Ô6Ð6Ð6Ð6r   c                 ó
  —  t          j        d¦  «        t          ¦  «        }t          j        dt          j        ¬¦  «        }dD ]=} |d         ||¦  «         |                      |d         t          |¦  «        ¦  «         Œ>d S )Nzvoid(int64[:], float32)r   r  ©rá  g      Àg      Àg      ø¿g      ø?g      @g      @g      @r  r   ©r   r  rð   r  r(  r   rM  rï   ©r  r  r   r   s       r   Útest_round_f4zTestCudaIntrinsic.test_round_f4á  ó„   € Ø6•4”8Ð5Ñ6Ô6µ|ÑDÔDˆÝŒh�q¥¤Ð)Ñ)Ô)ˆà@ð 	/ð 	/ˆAØˆH�TŒN˜3 Ñ"Ô"Ð"Ø×Ò˜S œV¥U¨1¡X¤XÑ.Ô.Ð.Ð.ð	/ð 	/r   c                 ó
  —  t          j        d¦  «        t          ¦  «        }t          j        dt          j        ¬¦  «        }dD ]=} |d         ||¦  «         |                      |d         t          |¦  «        ¦  «         Œ>d S )Nzvoid(int64[:], float64)r   r  rT  r  r   rU  rV  s       r   Útest_round_f8zTestCudaIntrinsic.test_round_f8é  rX  r   c           	      óê  —  t          j        d¦  «        t          ¦  «        }t          j        dt          j        ¬¦  «        }t          j                             d¦  «         t          j                             d¦  «                             t          j        ¦  «        }t          j	        |t          j
        t          j        t          j         t          j        g¦  «        f¦  «         d}t          j        ||¦  «        D ]s\  }}|                      ||¬¦  «        5   |d         |||¦  «         |                      |d	         t#          ||¦  «        d
¬¦  «         d d d ¦  «         n# 1 swxY w Y   Œtd S )Nú void(float32[:], float32, int32)r   r  é{   r^  )éûÿÿÿéüÿÿÿéýÿÿÿéþÿÿÿéÿÿÿÿr   r   r*   r÷   rõ   r%  rö   ©Úvalrò   r  r   Úsingle©Úprec)r   r  ró   r  r(  r“  r  r  rÏ  Úconcatenaterž  ÚinfÚnanÚ	itertoolsÚproductrù  ÚassertPreciseEqualrï   ©r  r  r   ÚvalsÚdigitsrd  rò   s          r   Útest_round_to_f4z"TestCudaIntrinsic.test_round_to_f4ñ  s‚  € Ø?•4”8Ð>Ñ?Ô?ÅÑPÔPˆÝŒh�q¥¤
Ð+Ñ+Ô+ˆÝ
Œ	�Š�sÑÔÐÝŒy×Ò Ñ#Ô#×*Ò*­2¬:Ñ6Ô6ˆÝ
Œ˜�bœh­¬µ´°½¼Ð'@ÑAÔAÐBÑCÔCÐCð
ˆõ &Ô-¨d°FÑ;Ô;ð 	7ð 	7‰LˆC�Ø—’ #¨w�Ñ7Ô7ð 7ð 7Ø�˜”˜s C¨Ñ1Ô1Ð1Ø×'Ò'¨¨A¬µ°c¸7Ñ0CÔ0CØ-5ð (ñ 7ô 7ð 7ð7ð 7ð 7ñ 7ô 7ð 7ð 7ð 7ð 7ð 7ð 7øøøð 7ð 7ð 7ð 7øð	7ð 	7s   ÄA E'Å'E+	Å.E+	z$Overflow behavior differs on CPythonc                 ó2  —  t          j        d¦  «        t          ¦  «        }t          j        dt          j        ¬¦  «        }t          j        t          j        ¦  «        j        }d} |d         |||¦  «         |                      |d         |¦  «         d S )Nr\  r   r  i,  r  r   )	r   r  ró   r  r(  r“  ÚfinfoÚmaxrM  ©r  r  r   rd  rò   s        r   Útest_round_to_f4_overflowz+TestCudaIntrinsic.test_round_to_f4_overflow  sˆ   € ð @•4”8Ð>Ñ?Ô?ÅÑPÔPˆÝŒh�q¥¤
Ð+Ñ+Ô+ˆÝŒh•r”zÑ"Ô"Ô&ˆð ˆØˆ�Œ�s˜C Ñ)Ô)Ð)Ø×Ò˜˜Qœ Ñ%Ô%Ð%Ð%Ð%r   c                 ó  —  t          j        d¦  «        t          ¦  «        }t          j        dt          j        ¬¦  «        }d}d} |d         |||¦  «         |                      |d         t          ||¦  «        d¬	¦  «         d S )
Nr\  r   r  gìQ¸…ëÕ?r÷   r  r   re  rf  )r   r  ró   r  r(  r“  rm  rï   ru  s        r   Útest_round_to_f4_halfwayz*TestCudaIntrinsic.test_round_to_f4_halfway  s…   € Ø?•4”8Ð>Ñ?Ô?ÅÑPÔPˆÝŒh�q¥¤
Ð+Ñ+Ô+ˆð ˆØˆØˆ�Œ�s˜C Ñ)Ô)Ð)Ø×Ò  A¤­¨c°7Ñ(;Ô(;À(ÐÑKÔKÐKÐKÐKr   c           	      ó–  —  t          j        d¦  «        t          ¦  «        }t          j        dt          j        ¬¦  «        }t          j                             d¦  «         t          j                             d¦  «        }t          j        |t          j	        t          j
        t          j
         t          j        g¦  «        f¦  «         d}t          j        ||¦  «        D ]s\  }}|                      ||¬¦  «        5   |d         |||¦  «         |                      |d	         t!          ||¦  «        d
¬¦  «         d d d ¦  «         n# 1 swxY w Y   Œtd}d}|                      ||¬¦  «        5   |d         |||¦  «         |                      |d	         t!          ||¦  «        d¬¦  «         d d d ¦  «         d S # 1 swxY w Y   d S )Nú void(float64[:], float64, int32)r   r  r]  r^  )r^  r_  r`  ra  rb  r   r   r*   r÷   rõ   r%  rc  r  r   Úexactrf  g`¢8„p=Ö<é   Údouble)r   r  ró   r  r(  r˜  r  r  rh  rž  ri  rj  rk  rl  rù  rm  rï   rn  s          r   Útest_round_to_f8z"TestCudaIntrinsic.test_round_to_f8#  s4  € Ø?•4”8Ð>Ñ?Ô?ÅÑPÔPˆÝŒh�q¥¤
Ð+Ñ+Ô+ˆÝ
Œ	�Š�sÑÔÐÝŒy×Ò Ñ#Ô#ˆÝ
Œ˜�bœh­¬µ´°½¼Ð'@ÑAÔAÐBÑCÔCÐCØ7ˆå%Ô-¨d°FÑ;Ô;ð 	6ð 	6‰LˆC�Ø—’ #¨w�Ñ7Ô7ð 6ð 6Ø�˜”˜s C¨Ñ1Ô1Ð1Ø×'Ò'¨¨A¬µ°c¸7Ñ0CÔ0CØ-4ð (ñ 6ô 6ð 6ð6ð 6ð 6ñ 6ô 6ð 6ð 6ð 6ð 6ð 6ð 6øøøð 6ð 6ð 6ð 6øð +ˆØˆØ�\Š\˜c¨7ˆ\Ñ3Ô3ð 	3ð 	3ØˆH�TŒN˜3  WÑ-Ô-Ð-Ø×#Ò# C¨¤F­E°#°wÑ,?Ô,?Ø)1ð $ñ 3ô 3ð 3ð	3ð 	3ð 	3ñ 	3ô 	3ð 	3ð 	3ð 	3ð 	3ð 	3ð 	3ð 	3øøøð 	3ð 	3ð 	3ð 	3ð 	3ð 	3s&   Ã>A E
Å
E	ÅE	Å1A F>Æ>GÇGc                 ó2  —  t          j        d¦  «        t          ¦  «        }t          j        dt          j        ¬¦  «        }t          j        t          j        ¦  «        j        }d} |d         |||¦  «         |                      |d         |¦  «         d S )Nrz  r   r  r„  r  r   )	r   r  ró   r  r(  r˜  rs  rt  rM  ru  s        r   Útest_round_to_f8_overflowz+TestCudaIntrinsic.test_round_to_f8_overflow:  sˆ   € ð @•4”8Ð>Ñ?Ô?ÅÑPÔPˆÝŒh�q¥¤
Ð+Ñ+Ô+ˆÝŒh•r”zÑ"Ô"Ô&ˆð ˆØˆ�Œ�s˜C Ñ)Ô)Ð)Ø×Ò˜˜Qœ Ñ%Ô%Ð%Ð%Ð%r   c                 ó  —  t          j        d¦  «        t          ¦  «        }t          j        dt          j        ¬¦  «        }d}d} |d         |||¦  «         |                      |d         t          ||¦  «        d¬	¦  «         d S )
Nrz  r   r  g\�Âõ(\á?r÷   r  r   r}  rf  )r   r  ró   r  r(  r˜  rm  rï   ru  s        r   Útest_round_to_f8_halfwayz*TestCudaIntrinsic.test_round_to_f8_halfwayG  s…   € Ø?•4”8Ð>Ñ?Ô?ÅÑPÔPˆÝŒh�q¥¤
Ð+Ñ+Ô+ˆð ˆØˆØˆ�Œ�s˜C Ñ)Ô)Ð)Ø×Ò  A¤­¨c°7Ñ(;Ô(;À(ÐÑKÔKÐKÐKÐKr   )KÚ__name__Ú
__module__Ú__qualname__r  r  r#  r3  r   r9  r<  rE  rK  rO  r\  ro  rq  rt  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*  r0  r2  r5  r9  r>  rC  rE  rG  rI  rL  rP  rR  rW  rZ  rq  rv  rx  r~  r€  r‚  Ú__classcell__)r  s   @r   r	  r	  n  sÃ  ø€ € € € € ðð ð ð ð ð%ð %ð %ð,ð ,ð ,ð0ð 0ð 0ð& €_Ð3Ñ4Ô4ð5ð 5ñ 5Ô4ð5ð €_Ð3Ñ4Ô4ð9ð 9ñ 5Ô4ð9ð9ð 9ð 9ð,ð ,ð ,ð0ð 0ð 0ð €_Ð0Ñ1Ô1ð/ð /ñ 2Ô1ð/ð0 €_Ð)Ñ*Ô*ðCð Cñ +Ô*ðCð06ð 6ð 6ðLð Lð Lð"
-ð 
-ð 
-ð%ð %ð %ð*$ð $ð $ð$ð $ð $ð6ð 6ð 6ð6ð 6ð 6ð ð8ð 8ñ Ôð8ð ð0ð 0ñ Ôð0ð €_Ð?Ñ@Ô@ð&ð &ñ AÔ@ð&ð
 ð?ð ?ñ Ôð?ð ð0ð 0ñ Ôð0ð €_Ð?Ñ@Ô@ð)ð )ñ AÔ@ð)ð
 ð8ð 8ñ Ôð8ð ð0ð 0ñ Ôð0ð €_Ð?Ñ@Ô@ð&ð &ñ AÔ@ð&ð
 ð8ð 8ñ Ôð8ð ð0ð 0ñ Ôð0ð €_Ð?Ñ@Ô@ð&ð &ñ AÔ@ð&ð
 ð0ð 0ñ Ôð0ð ð-ð -ñ Ôð-ð ð2ð 2ñ Ôð2ð ð0ð 0ñ Ôð0ð €_Ð?Ñ@Ô@ð&ð &ñ AÔ@ð&ð
 ð6ð 6ñ Ôð6ð ð0ð 0ñ Ôð0ð €_Ð?Ñ@Ô@ð&ð &ñ AÔ@ð&ð
 ð!Cð !Cñ Ôð!CðF ð/ð /ñ Ôð/ð$ ð3ð 3ñ Ôð3ð> ð(ð (ñ Ôð(ð  ð	1ð 	1ñ Ôð	1ð ð	1ð 	1ñ Ôð	1ð@ð @ð @ð@ð @ð @ð-ð -ð -ð €_ÐEÑFÔFð5ð 5ñ GÔFð5ð%ð %ð %ð%ð %ð %ð$ð $ð $ð7ð 7ð 7ð €_ÐEÑFÔFð%ð %ñ GÔFð%ð%ð %ð %ð%ð %ð %ð$ð $ð $ð$ð $ð $ð €_ÐEÑFÔFð%ð %ñ GÔFð%ð,ð ,ð ,ð7ð 7ð 7ð/ð /ð /ð/ð /ð /ð7ð 7ð 7ð4 €_Ð;Ñ<Ô<ð&ð &ñ =Ô<ð&ðLð Lð Lð3ð 3ð 3ð. €_Ð;Ñ<Ô<ð
&ð 
&ñ =Ô<ð
&ð	Lð 	Lð 	Lð 	Lð 	Lð 	Lð 	Lr   r	  Ú__main__)[rk  Únumpyr  r	  re  Únumbar   r   Ú
numba.cudar   Ú
numba.corer   Únumba.core.errorsr   Únumba.core.typesr   Únumba.cuda.testingr	   r
   r   r   r   r   r#   r'   r+   r0   r2   r?   rC   rJ   rP   rS   rX   rZ   r^   r`   rd   rf   ri   rn   rs   rv   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ð   ró   rø   rü   rÿ   r  r  r  r	  rƒ  ÚmainrÎ   r   r   ú<module>r�     sB  ðØ Ð Ð Ð Ø Ð Ð Ð Ø €€€Ø 	€	€	€	Ø Ð Ð Ð Ð Ð Ð Ð Ø "Ð "Ð "Ð "Ð "Ð "Ø Ð Ð Ð Ð Ð Ø )Ð )Ð )Ð )Ð )Ð )Ø Ð Ð Ð Ð Ð ð3ð 3ð 3ð 3ð 3ð 3ð 3ð 3ð 3ð 3ð 3ð 3ðð ð ð
ð ð ð
/ð /ð /ðð ð ð
ð ð ð
ð ð ðð ð ðð ð ðð ð ðð ð ð(ð (ð (ð"ð "ð "ð.ð .ð .ð%ð %ð %ð(ð (ð (ð"ð "ð "ð(ð (ð (ð"ð "ð "ð"ð "ð "ð&ð &ð &ð"ð "ð "ðð ð ð"ð "ð "ðð ð ð!ð !ð !ð!ð !ð !ð!ð !ð !ð!ð !ð !ð!ð !ð !ð!ð !ð !ð 
€„�ÐÑÔðð ñ Ôðð 
€„�ÐÑÔðð ñ Ôðð1ð 1ð 1ð
4ð 4ð 4ð
4ð 4ð 4ð
7ð 7ð 7ð
7ð 7ð 7ð
"ð "ð "ð"ð "ð "ð$ð $ð $ð$ð $ð $ð$ð $ð $ð%ð %ð %ð&ð &ð &ð$ð $ð $ð%ð %ð %ð%ð %ð %ð&ð &ð &ðð ð ð%ð %ð %ð&ð &ð &ð$ð $ð $ð&ð &ð &ð%ð %ð %ðð ð ðð ð ðð ð ðð ð ðð ð ðð ð ð	ð 	ð 	ð)ð )ð )ðð ð ð
ð ð ðð ð ðð ð ðbLð bLð bLð bLð bL˜ñ bLô bLð bLðJ ˆzÒÐØ€H„M�O„O€O€O€Oð Ðr   