§
    ƒÛgÒ„  ã                   ó¤  — 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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"„ Z.d#„ Z/d$„ Z0 ej1        d%¬&¦  «        d'„ ¦   «         Z2 ej1        d%¬&¦  «        d(„ ¦   «         Z3d)„ Z4d*„ Z5d+„ Z6d,„ Z7d-„ Z8d.„ Z9d/„ Z:d0„ Z;d1„ Z<d2„ Z=d3„ Z>d4„ Z?d5„ Z@d6„ ZAd7„ ZBd8„ ZCd9„ ZDd:„ ZEd;„ ZFd<„ ZGd=„ ZHd>„ ZId?„ ZJd@„ ZKdA„ ZLdB„ ZMdC„ ZNdD„ ZOdE„ ZPdF„ ZQdG„ ZRdH„ ZSdI„ ZTdJ„ ZU G dK„ dLe¦  «        ZVeWdMk    r ejX        ¦   «          dS dS )Né    N)ÚcudaÚint64)Úcompile_ptx)ÚTypingError)Úf2)ÚunittestÚCUDATestCaseÚskip_on_cudasimÚskip_unless_cc_53c                 ó2   — t           j        j        }|| d<   d S ©Nr   ©r   Ú	threadIdxÚx©ÚaryÚis     ú]/granary/edev/quantum/lib/python3.11/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_gridsize1dr.   )   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_gridsize2dr0   0   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=   8   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   r6   s     r   Úsimple_popcrA   C   ó   € ÝŒY�q‰\Œ\€Cˆ�F€F€Fr   c                 ó8   — t          j        |||¦  «        | d<   d S r   )r   Úfma©r   ÚaÚbr6   s       r   Ú
simple_fmarH   G   s   € ÝŒX�a˜˜AÑÔ€Cˆ�F€F€Fr   c                 ód   — t           j                             |d         |d         ¦  «        | d<   d S r   ©r   Úfp16Úhadd©r   rF   rG   s      r   Úsimple_haddrN   K   ó'   € ÝŒY�^Š^˜A˜aœD ! A¤$Ñ'Ô'€Cˆ�F€F€Fr   c                 óL   — t           j                             ||¦  «        | d<   d S r   rJ   rM   s      r   Úsimple_hadd_scalarrQ   O   ó   € ÝŒY�^Š^˜A˜qÑ!Ô!€Cˆ�F€F€Fr   c                 ór   — t           j                             |d         |d         |d         ¦  «        | d<   d S r   ©r   rK   ÚhfmarE   s       r   Úsimple_hfmarV   S   s-   € ÝŒY�^Š^˜A˜aœD ! A¤$¨¨!¬Ñ-Ô-€Cˆ�F€F€Fr   c                 óN   — t           j                             |||¦  «        | d<   d S r   rT   rE   s       r   Úsimple_hfma_scalarrX   W   s!   € ÝŒY�^Š^˜A˜q !Ñ$Ô$€Cˆ�F€F€Fr   c                 ód   — t           j                             |d         |d         ¦  «        | d<   d S r   ©r   rK   ÚhsubrM   s      r   Úsimple_hsubr\   [   rO   r   c                 óL   — t           j                             ||¦  «        | d<   d S r   rZ   rM   s      r   Úsimple_hsub_scalarr^   _   rR   r   c                 ód   — t           j                             |d         |d         ¦  «        | d<   d S r   ©r   rK   ÚhmulrM   s      r   Úsimple_hmulrb   c   rO   r   c                 óL   — t           j                             ||¦  «        | d<   d S r   r`   rM   s      r   Úsimple_hmul_scalarrd   g   rR   r   c                 óL   — t           j                             ||¦  «        | d<   d S r   )r   rK   ÚhdivrM   s      r   Úsimple_hdiv_scalarrg   k   rR   r   c                 ó®   — t          j        d¦  «        }|| j        k     r5||         }||         }t           j                             ||¦  «        | |<   d S d S r   )r   r%   ÚsizerK   rf   )r   Úarray_aÚarray_br   rF   rG   s         r   Úsimple_hdiv_kernelrl   o   sP   € ÝŒ	�!‰Œ€AØˆ3Œ8‚|€|Ø�AŒJˆØ�AŒJˆÝ”—’  1Ñ%Ô%ˆˆA‰ˆˆð €|r   c                 óV   — t           j                             |d         ¦  «        | d<   d S r   ©r   rK   Úhneg©r   rF   s     r   Úsimple_hnegrq   w   ó!   € ÝŒY�^Š^˜A˜aœDÑ!Ô!€Cˆ�F€F€Fr   c                 óJ   — t           j                             |¦  «        | d<   d S r   rn   rp   s     r   Úsimple_hneg_scalarrt   {   ó   € ÝŒY�^Š^˜AÑÔ€Cˆ�F€F€Fr   c                 óV   — t           j                             |d         ¦  «        | d<   d S r   ©r   rK   Úhabsrp   s     r   Úsimple_habsry      rr   r   c                 óJ   — t           j                             |¦  «        | d<   d S r   rw   rp   s     r   Úsimple_habs_scalarr{   ƒ   ru   r   c                 óL   — t           j                             ||¦  «        | d<   d S r   )r   rK   ÚheqrM   s      r   Úsimple_heq_scalarr~   ‡   ó   € ÝŒY�]Š]˜1˜aÑ Ô €Cˆ�F€F€Fr   c                 óL   — t           j                             ||¦  «        | d<   d S r   )r   rK   ÚhnerM   s      r   Úsimple_hne_scalarr‚   ‹   r   r   c                 óL   — t           j                             ||¦  «        | d<   d S r   )r   rK   ÚhgerM   s      r   Úsimple_hge_scalarr…   �   r   r   c                 óL   — t           j                             ||¦  «        | d<   d S r   )r   rK   ÚhgtrM   s      r   Úsimple_hgt_scalarrˆ   “   r   r   c                 óL   — t           j                             ||¦  «        | d<   d S r   )r   rK   ÚhlerM   s      r   Úsimple_hle_scalarr‹   —   r   r   c                 óL   — t           j                             ||¦  «        | d<   d S r   ©r   rK   ÚhltrM   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–   ©ÚrrF   rG   r6   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   rK   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   rK   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   rK   rŽ   r„   r˜   s       r   Útest_multiple_hcmp_5r¤   ½   r¢   r   c                 óL   — t           j                             ||¦  «        | d<   d S r   )r   rK   ÚhmaxrM   s      r   Úsimple_hmax_scalarr§   Â   rR   r   c                 óL   — t           j                             ||¦  «        | d<   d S r   )r   rK   ÚhminrM   s      r   Úsimple_hmin_scalarrª   Æ   rR   r   c                 ó¨   — t          j        d¦  «        }|t          | ¦  «        k     r*t           j                             ||         ¦  «        | |<   d S d S r   )r   r%   ÚlenrK   Ú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¬   rK   Úhcosr®   s      r   Úsimple_hcosr³   Ñ   r°   r   c                 ó¨   — t          j        d¦  «        }|t          | ¦  «        k     r*t           j                             ||         ¦  «        | |<   d S d S r   )r   r%   r¬   rK   Úhlogr®   s      r   Úsimple_hlogr¶   Ø   r°   r   c                 ó¨   — t          j        d¦  «        }|t          | ¦  «        k     r*t           j                             ||         ¦  «        | |<   d S d S r   )r   r%   r¬   rK   Ú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¬   rK   Ú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¬   rK   Úhexpr®   s      r   Úsimple_hexprÁ   í   r°   r   c                 ó¨   — t          j        d¦  «        }|t          | ¦  «        k     r*t           j                             ||         ¦  «        | |<   d S d S r   )r   r%   r¬   rK   Úhexp2r®   s      r   Úsimple_hexp2rÄ   ô   rº   r   c                 ó¨   — t          j        d¦  «        }|t          | ¦  «        k     r*t           j                             ||         ¦  «        | |<   d S d S r   )r   r%   r¬   rK   Úhsqrtr®   s      r   Úsimple_hsqrtrÇ   û   rº   r   c                 ó¨   — t          j        d¦  «        }|t          | ¦  «        k     r*t           j                             ||         ¦  «        | |<   d S d S r   )r   r%   r¬   rK   Ú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¬   rK   Úhceilr®   s      r   Úsimple_hceilrÑ     rº   r   c                 ó¨   — t          j        d¦  «        }|t          | ¦  «        k     r*t           j                             ||         ¦  «        | |<   d S d S r   )r   r%   r¬   rK   Úhfloorr®   s      r   Úsimple_hfloorrÔ     r¾   r   c                 ó¨   — t          j        d¦  «        }|t          | ¦  «        k     r*t           j                             ||         ¦  «        | |<   d S d S r   )r   r%   r¬   rK   Úhrcpr®   s      r   Úsimple_hrcpr×     r°   r   c                 ó¨   — t          j        d¦  «        }|t          | ¦  «        k     r*t           j                             ||         ¦  «        | |<   d S d S r   )r   r%   r¬   rK   Úhtruncr®   s      r   Úsimple_htruncrÚ   #  r¾   r   c                 ó¨   — t          j        d¦  «        }|t          | ¦  «        k     r*t           j                             ||         ¦  «        | |<   d S d S r   )r   r%   r¬   rK   Úhrintr®   s      r   Úsimple_hrintrÝ   *  rº   r   c                 ó4   — t          j        |¦  «        | d<   d S r   )r   Úcbrtrp   s     r   Úsimple_cbrtrà   1  rB   r   c                 ó4   — t          j        |¦  «        | d<   d S r   )r   Úbrevr@   s     r   Úsimple_brevrã   5  rB   r   c                 ó4   — t          j        |¦  «        | d<   d S r   )r   Úclzr@   s     r   Ú
simple_clzræ   9  ó   € ÝŒX�a‰[Œ[€Cˆ�F€F€Fr   c                 ó4   — t          j        |¦  «        | d<   d S r   )r   Úffsr@   s     r   Ú
simple_ffsrê   =  rç   r   c                 ó*   — t          |¦  «        | d<   d S r   ©Úroundr@   s     r   Úsimple_roundrî   A  s   € Ý�1‰XŒX€Cˆ�F€F€Fr   c                 ó,   — t          ||¦  «        | d<   d S r   rì   )r   r6   Úndigitss      r   Úsimple_round_torñ   E  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$   )rF   rG   r6   r   s       r   Úbranching_with_ifsrö   I  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)rF   rG   r6   r   Úinners        r   Úbranching_with_selpsrú   U  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   a  s   € ÝŒ]€Cˆ�F€F€Fr   c                 ó.   — t          j        | ¦  «         d S r   r$   ©r   s    r   Únonliteral_gridr  e  s   € Ý„Iˆa�L„L€L€L€Lr   c                 ó.   — t          j        | ¦  «         d S r   )r   r-   r  s    r   Únonliteral_gridsizer  i  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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d.„ ¦   «         Z/ed/„ ¦   «         Z0d0„ Z1d1„ Z2d2„ Z3 ed3¦  «        d4„ ¦   «         Z4d5„ Z5d6„ Z6d7„ Z7d8„ Z8 ed3¦  «        d9„ ¦   «         Z9d:„ Z:d;„ Z;d<„ Z<d=„ Z= ed3¦  «        d>„ ¦   «         Z>d?„ Z?d@„ Z@dA„ ZAdB„ ZBdC„ ZC edD¦  «        dE„ ¦   «         ZDdF„ ZEdG„ ZF edD¦  «        dH„ ¦   «         ZGdI„ ZHˆ xZIS )JÚTestCudaIntrinsicc                 ó†   •— t          ¦   «                              ¦   «          t          j                             d¦  «         d S r   )ÚsuperÚsetUpÚnpÚrandomÚseed)ÚselfÚ	__class__s    €r   r
  zTestCudaIntrinsic.setUpn  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_threadidxr  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_threadidxx  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+  r.  Ú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 r3  )r5  r   r   r  r  r6  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ð 	9r8  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  r?  r  Úcopyr5   r4   r  r  )	r  r  r@  rA  r4   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  r.   r  r&  r  ÚassertEqual©r  r  r@  rA  r   s        r   Útest_simple_gridsize1dz(TestCudaIntrinsic.test_simple_gridsize1d´  sy   € Ø/•4”8Ð.Ñ/Ô/Õ0AÑBÔBˆØ‰ˆˆfÝŒh�q¥¤Ð)Ñ)Ô)ˆØˆ�˜�Ô˜sÑ#Ô#Ð#Ø×Ò˜˜Qœ ¨$¡Ñ/Ô/Ð/Ð/Ð/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$  é   )r4   Ú
fill_valuerÍ   rõ   r#  r  r   r)   z	\s+bra\s+Ú	branching)Úerr_msgr   rø   )r   r   r  rö   rú   r  ÚfullrH  r  Úinspect_asmrK  r¬   ÚreÚfindallÚtestingÚassert_array_equal)
r  ÚsigÚcu_branching_with_ifsÚcu_branching_with_selpsÚnrG   r6   ÚexpectedrF   Ú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<  rF  rG  r)   r  r   r   )r   r  r0   r  r&  r  rK  rL  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 rE  )
r   r  r=   r  r?  r  r4   Úzipr5   r  )r  r  r@  rA  r4   r   r9   r:   r;   r<   r   r    r7   r8   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   rF   rG   r6   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õ   rl  )r   r  r  r&  r  ÚreshaperX  Úassert_equal)r  rj  Ú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 rh  )
r   r%   r-   r   r   ÚblockIdxr3   r   r   r2   )	ri  r   r   r   rF   rG   r6   Úgrid_is_rightÚgridsize_is_rights	            r   rj  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_rm  r  r  )r  rj  r   r   r   ro  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  rA   r  r&  r  ÚassertEqualsr  s      r   Útest_popc_u4zTestCudaIntrinsic.test_popc_u4  sh   € Ø5•4”8Ð4Ñ5Ô5µkÑBÔBˆÝŒh�q¥¤Ð)Ñ)Ô)ˆØˆ�Œ�s˜DÑ!Ô!Ð!Ø×Ò˜#˜aœ& !Ñ$Ô$Ð$Ð$Ð$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_u8  sh   € Ø5•4”8Ð4Ñ5Ô5µkÑBÔBˆÝŒh�q¥¤Ð)Ñ)Ô)ˆØˆ�Œ�s˜NÑ+Ô+Ð+Ø×Ò˜#˜aœ& !Ñ$Ô$Ð$Ð$Ð$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  rH   r  r&  Úfloat32rX  Úassert_allcloser  s      r   Útest_fma_f4zTestCudaIntrinsic.test_fma_f4  ó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  rH   r  r&  Úfloat64rX  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  rN   r  r&  Úfloat16ÚarrayrX  r‡  ©r  r  r   Úarg1Úarg2s        r   Ú	test_haddzTestCudaIntrinsic.test_hadd(  ó    € à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  rQ   r  r&  r�  rX  r‡  ©r  r  r   r“  r”  Úrefs         r   Útest_hadd_scalarz"TestCudaIntrinsic.test_hadd_scalar1  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   rQ   ÚassertIn©r  Úargsr_  Ú_s       r   Útest_hadd_ptxzTestCudaIntrinsic.test_hadd_ptx;  ó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  rV   r  r&  r�  r‘  rX  r‡  )r  r  r   r“  r”  Úarg3s         r   Ú	test_hfmazTestCudaIntrinsic.test_hfmaA  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  rX   r  r&  r�  rX  r‡  )r  r  r   r“  r”  r¨  r›  s          r   Útest_hfma_scalarz"TestCudaIntrinsic.test_hfma_scalarK  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   rX   r¡  r¢  s       r   Útest_hfma_ptxzTestCudaIntrinsic.test_hfma_ptxV  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‘  rX  r‡  r’  s        r   Ú	test_hsubzTestCudaIntrinsic.test_hsub\  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�  rX  r‡  rš  s         r   Útest_hsub_scalarz"TestCudaIntrinsic.test_hsub_scalare  ó”   € à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_ptxo  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  rb   r  r&  r�  r‘  rX  r‡  r’  s        r   Ú	test_hmulzTestCudaIntrinsic.test_hmulu  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  rd   r  r&  r�  rX  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   rd   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  rg   r  r&  r�  rX  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ô  ©ri   r  )r   r  rl   r  r  ÚrandintÚastyper�  Ú
zeros_likeÚforallri   rX  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  rq   r  r&  r�  r‘  rX  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  rt   r  r&  r�  rX  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   rt   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  ry   r  r&  r�  r‘  rX  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Õ  rX  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"        |¦  «        }tG          ||¦  «        D ]›\  }}	|  $                    |	¬¦  «        5   tK          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!        ¦  «        }tG          ||¦  «        D ]›\  }}	|  $                    |	¬¦  «        5   tK          j&        d¦  «        |¦  «        } |d|f         ||¦  «          |	|t          j!        ¬¦  «        }
t          j'         (                    ||
¦  «         d d d ¦  «         n# 1 swxY w Y   Œœd S )	NrO  r   r¿  rÀ  ©ÚfnrÉ  r  r  ))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Ã  rd  ÚsubTestr   r  rX  r‡  )r  ÚkernelsÚexp_kernelsÚexpected_functionsÚexpected_exp_functionsr  r   r™   ÚkernelrÝ  r^  Ú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˜!ÑÔˆÝ˜gÐ'9Ñ:Ô:ð 	8ð 	8‰JˆF�BØ—’ �Ñ$Ô$ð 8ð 8Ø7�œÐ"6Ñ7Ô7¸Ñ?Ô?�Ø��q˜�s”˜A˜qÑ!Ô!Ð!Ø˜2˜a¥r¤zÐ2Ñ2Ô2�Ý”
×*Ò*¨1¨hÑ7Ô7Ð7ð	8ð 8ð 8ñ 8ô 8ð 8ð 8ð 8ð 8ð 8ð 8øøøð 8ð 8ð 8ð 8øõ ŒY×Ò˜q "¨1ÐÑ-Ô-×4Ò4µR´ZÑ@Ô@ˆÝ˜kÐ+AÑBÔBð 	8ð 	8‰JˆF�BØ—’ �Ñ$Ô$ð 8ð 8Ø7�œÐ"6Ñ7Ô7¸Ñ?Ô?�Ø��q˜�s”˜A˜rÑ"Ô"Ð"Ø˜2˜b­¬
Ð3Ñ3Ô3�Ý”
×*Ò*¨1¨hÑ7Ô7Ð7ð	8ð 8ð 8ñ 8ô 8ð 8ð 8ð 8ð 8ð 8ð 8øøøð 8ð 8ð 8ð 8øð	8ð 	8s&   ÆA)HÈH
	ÈH
	É<A)K1Ë1K5	Ë8K5	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¬   rK   Úhexp10r®   s      r   Úhexp10_vectorsz5TestCudaIntrinsic.test_hexp10.<locals>.hexp10_vectors÷  sF   € å”	˜!‘”ˆAà•3�q‘6”6ŠzˆzÝ”y×'Ò'¨¨!¬Ñ-Ô-��!‘��ð ˆzr   rO  r   r  )r   r  r  r  r  ÚrandrÂ  r�  rÃ  rX  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Últrd  rê  r   r  r  r&  Úbool8r�  rK  )r  ÚfnsÚopsrÝ  rù  rï  r^  Ú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_comparisons(  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�  rX  r‡  r’  s        r   Ú	test_hmaxzTestCudaIntrinsic.test_hmax9  óÈ   € à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�  rX  r‡  r’  s        r   Ú	test_hminzTestCudaIntrinsic.test_hminE  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†  rX  r‡  ©r  r  r   Úcbrt_args       r   Útest_cbrt_f32zTestCudaIntrinsic.test_cbrt_f32Q  ó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‹  rX  r‡  r  s       r   Útest_cbrt_f64zTestCudaIntrinsic.test_cbrt_f64X  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&  Úuint32r~  r  s      r   Útest_brev_u4zTestCudaIntrinsic.test_brev_u4_  sh   € Ø6•4”8Ð5Ñ6Ô6µ{ÑCÔCˆÝŒh�q¥¤	Ð*Ñ*Ô*ˆØˆ�Œ�s˜JÑ'Ô'Ð'Ø×Ò˜#˜aœ& *Ñ-Ô-Ð-Ð-Ð-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&  Úuint64r~  r  s      r   Útest_brev_u8zTestCudaIntrinsic.test_brev_u8e  sj   € à6•4”8Ð5Ñ6Ô6µ{ÑCÔCˆÝŒh�q¥¤	Ð*Ñ*Ô*ˆØˆ�Œ�sÐ.Ñ/Ô/Ð/Ø×Ò˜#˜aœ&Ð"4Ñ5Ô5Ð5Ð5Ð5r   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  r~  r  s      r   Útest_clz_i4zTestCudaIntrinsic.test_clz_i4l  sh   € Ø4•4”8Ð3Ñ4Ô4µZÑ@Ô@ˆÝŒh�q¥¤Ð)Ñ)Ô)ˆØˆ�Œ�s˜JÑ'Ô'Ð'Ø×Ò˜#˜aœ& "Ñ%Ô%Ð%Ð%Ð%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_u4r  sj   € ð 6•4”8Ð4Ñ5Ô5µjÑAÔAˆÝŒh�q¥¤Ð)Ñ)Ô)ˆØˆ�Œ�s˜JÑ'Ô'Ð'Ø×Ò˜#˜aœ& "Ñ%Ô%Ð%Ð%Ð%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Ñ'Ô'Ð'Ø×Ò˜#˜aœ& !Ñ$Ô$Ð$Ð$Ð$r   c                 óè   —  t          j        d¦  «        t          ¦  «        }t          j        dt          j        ¬¦  «        } |d         |d¦  «         |                      |d         dd¦  «         d S )Nr  r   r  r  r   rO  ú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Ñ Ô Ð Ø×Ò˜#˜aœ& "Ð&6Ñ7Ô7Ð7Ð7Ð7r   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Ð-Ñ.Ô.Ð.Ø×Ò˜#˜aœ& "Ñ%Ô%Ð%Ð%Ð%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   é   ì        rO  ©r   r  rê   r  r&  r  r~  r  s      r   Útest_ffs_i4zTestCudaIntrinsic.test_ffs_i4“  s˜   € Ø4•4”8Ð3Ñ4Ô4µZÑ@Ô@ˆÝŒh�q¥¤Ð)Ñ)Ô)ˆØˆ�Œ�s˜JÑ'Ô'Ð'Ø×Ò˜#˜aœ& "Ñ%Ô%Ð%Øˆ�Œ�s˜JÑ'Ô'Ð'Ø×Ò˜#˜aœ& "Ñ%Ô%Ð%Ð%Ð%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   r3  r4  rO  r5  r  s      r   Útest_ffs_u4zTestCudaIntrinsic.test_ffs_u4›  s˜   € Ø5•4”8Ð4Ñ5Ô5µjÑAÔAˆÝŒh�q¥¤Ð)Ñ)Ô)ˆØˆ�Œ�s˜JÑ'Ô'Ð'Ø×Ò˜#˜aœ& "Ñ%Ô%Ð%Øˆ�Œ�s˜JÑ'Ô'Ð'Ø×Ò˜#˜aœ& "Ñ%Ô%Ð%Ð%Ð%r   c                 óæ   —  t          j        d¦  «        t          ¦  «        }t          j        dt          j        ¬¦  «        } |d         |d¦  «         |                      |d         d¦  «         d S r'  r5  r  s      r   Útest_ffs_i4_1sz TestCudaIntrinsic.test_ffs_i4_1s£  r)  r   c                 óæ   —  t          j        d¦  «        t          ¦  «        }t          j        dt          j        ¬¦  «        } |d         |d¦  «         |                      |d         d¦  «         d S )Nr  r   r  r  r   r5  r  s      r   Útest_ffs_i4_0sz TestCudaIntrinsic.test_ffs_i4_0s©  sh   € Ø4•4”8Ð3Ñ4Ô4µZÑ@Ô@ˆÝŒh�q¥¤Ð)Ñ)Ô)ˆØˆ�Œ�s˜CÑ Ô Ð Ø×Ò˜#˜aœ& !Ñ$Ô$Ð$Ð$Ð$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   rP  l        é!   r5  r  s      r   Útest_ffs_i8zTestCudaIntrinsic.test_ffs_i8¯  s™   € à4•4”8Ð3Ñ4Ô4µZÑ@Ô@ˆÝŒh�q¥¤Ð)Ñ)Ô)ˆØˆ�Œ�sÐ-Ñ.Ô.Ð.Ø×Ò˜#˜aœ& "Ñ%Ô%Ð%Øˆ�Œ�s˜KÑ(Ô(Ð(Ø×Ò˜#˜aœ& "Ñ%Ô%Ð%Ð%Ð%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)   rO  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   rO  r+  )r   r  r   r  r&  r  r~  r  s      r   Útest_simple_warpsizez&TestCudaIntrinsic.test_simple_warpsizeÀ  si   € Ø-•4”8Ð,Ñ-Ô-­oÑ>Ô>ˆÝŒh�q¥¤Ð)Ñ)Ô)ˆØˆ�Œ�sÑÔÐØ×Ò˜#˜aœ& "Ð&6Ñ7Ô7Ð7Ð7Ð7r   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   r~  rí   ©r  r  r   r   s       r   Útest_round_f4zTestCudaIntrinsic.test_round_f4Æ  ó„   € Ø6•4”8Ð5Ñ6Ô6µ|ÑDÔDˆÝŒh�q¥¤Ð)Ñ)Ô)ˆà@ð 	0ð 	0ˆAØˆH�TŒN˜3 Ñ"Ô"Ð"Ø×Ò˜c !œf¥e¨A¡h¤hÑ/Ô/Ð/Ð/ð	0ð 	0r   c                 ó
  —  t          j        d¦  «        t          ¦  «        }t          j        dt          j        ¬¦  «        }dD ]=} |d         ||¦  «         |                      |d         t          |¦  «        ¦  «         Œ>d S )Nzvoid(int64[:], float64)r   r  rG  r  r   rH  rI  s       r   Útest_round_f8zTestCudaIntrinsic.test_round_f8Î  rK  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  é{   rO  )éûÿÿÿéüÿÿÿéýÿÿÿéþÿÿÿéÿÿÿÿ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ÚdigitsrW  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 )NrO  r   r  i,  r  r   )	r   r  rñ   r  r&  r†  ÚfinfoÚmaxrK  ©r  r  r   rW  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 )
NrO  r   r  gìQ¸…ëÕ?rõ   r  r   rX  rY  )r   r  rñ   r  r&  r†  r`  rí   rh  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  rP  rO  )rQ  rR  rS  rT  rU  r   r   r)   rõ   ró   r#  rV  r  r   ÚexactrY  g`¢8„p=Ö<é   Údouble)r   r  rñ   r  r&  r‹  r  r  r[  r‘  r\  r]  r^  r_  rê  r`  rí   ra  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 )Nrm  r   r  rv  r  r   )	r   r  rñ   r  r&  r‹  rf  rg  rK  rh  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 )
Nrm  r   r  g\�Âõ(\á?rõ   r  r   rp  rY  )r   r  rñ   r  r&  r‹  r`  rí   rh  s        r   Útest_round_to_f8_halfwayz*TestCudaIntrinsic.test_round_to_f8_halfway,  s…   € Ø?•4”8Ð>Ñ?Ô?ÅÑPÔPˆÝŒh�q¥¤
Ð+Ñ+Ô+ˆð ˆØˆØˆ�Œ�s˜C Ñ)Ô)Ð)Ø×Ò  A¤­¨c°7Ñ(;Ô(;À(ÐÑKÔKÐKÐKÐKr   )JÚ__name__Ú
__module__Ú__qualname__r
  r  r!  r1  r
   r7  r:  rC  rI  rM  r`  rb  re  rp  ry  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,  r1  r6  r8  r:  r<  r?  rC  rE  rJ  rM  rd  ri  rk  rq  rs  ru  Ú__classcell__)r  s   @r   r  r  m  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ð €_Ð)Ñ*Ô*ð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Ô@ð&ð
 ð 8ð  8ñ Ôð 8ðD ð/ð /ñ Ôð/ð$ ð3ð 3ñ Ôð3ð> ð(ð (ñ Ôð(ð  ð	1ð 	1ñ Ôð	1ð ð	1ð 	1ñ Ôð	1ð@ð @ð @ð@ð @ð @ð.ð .ð .ð €_ÐEÑFÔFð6ð 6ñ GÔFð6ð&ð &ð &ð&ð &ð &ð%ð %ð %ð8ð 8ð 8ð €_ÐEÑFÔFð&ð &ñ GÔFð&ð&ð &ð &ð&ð &ð &ð%ð %ð %ð%ð %ð %ð €_ÐEÑFÔFð&ð &ñ GÔFð&ð,ð ,ð ,ð8ð 8ð 8ð0ð 0ð 0ð0ð 0ð 0ð7ð 7ð 7ð4 €_Ð;Ñ<Ô<ð&ð &ñ =Ô<ð&ðLð Lð Lð3ð 3ð 3ð. €_Ð;Ñ<Ô<ð
&ð 
&ñ =Ô<ð
&ð	Lð 	Lð 	Lð 	Lð 	Lð 	Lð 	Lr   r  Ú__main__)Yr^  Únumpyr  rú  rV  Únumbar   r   Ú
numba.cudar   Únumba.core.errorsr   Únumba.core.typesr   Únumba.cuda.testingr   r	   r
   r   r   r   r"   r&   r*   r.   r0   r=   rA   rH   rN   rQ   rV   rX   r\   r^   rb   rd   rg   rl   rq   rt   ry   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  rv  ÚmainrÌ   r   r   ú<module>r‚     s0  ðØ Ð Ð Ð Ø Ð Ð Ð Ø €€€Ø 	€	€	€	Ø Ð Ð Ð Ð Ð Ð Ð Ø "Ð "Ð "Ð "Ð "Ð "Ø )Ð )Ð )Ð )Ð )Ð )Ø Ð Ð Ð Ð Ð ð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ð
"ð "ð "ð"ð "ð "ð$ð $ð $ð$ð $ð $ð$ð $ð $ð%ð %ð %ð&ð &ð &ð$ð $ð $ð%ð %ð %ð%ð %ð %ð&ð &ð &ðð ð ð%ð %ð %ð&ð &ð &ð$ð $ð $ð&ð &ð &ð%ð %ð %ðð ð ðð ð ðð ð ðð ð ðð ð ðð ð ð	ð 	ð 	ð)ð )ð )ðð ð ð
ð ð ðð ð ðð ð ðHLð HLð HLð HLð HL˜ñ HLô HLð HLðV ˆzÒÐØ€H„M�O„O€O€O€Oð Ðr   