§
    ƒÛgQä  ã                   ó^  — d dl Zd dlmZ d dlmZmZmZmZm	Z	 d dl
mZmZmZ d dlmZ  ej        d¬¦  «        d„ ¦   «         Z ej        d¬¦  «        d	„ ¦   «         Z ej        d¬¦  «        d
„ ¦   «         Z ej        d¬¦  «        d„ ¦   «         Z ej        d¬¦  «        d„ ¦   «         Z ej        d¬¦  «        d„ ¦   «         Z ej        d¬¦  «        d„ ¦   «         Z ej        d¬¦  «        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)„ Z2d*„ Z3d+„ Z4d,„ Z5d-„ Z6d.„ Z7d/„ Z8d0„ Z9d1„ Z:d2„ Z;d3„ Z<d4„ Z=d5„ Z>d6„ Z?d7„ Z@d8„ ZAd9„ ZBd:„ ZCd;„ ZDd<„ ZEd=„ ZFd>„ ZGd?„ ZHd@„ ZIdA„ ZJdB„ ZKdC„ ZLdD„ ZMdE„ ZNdF„ ZOdG„ ZPdH„ ZQdI„ ZRdJ„ ZSdK„ ZTdL„ ZUdM„ ZVdN„ ZWdO„ ZXdP„ ZYdQ„ ZZ eZdR¦  «        \  Z[Z\Z]Z^ eZdS¦  «        \  Z_Z`ZaZb eZdT¦  «        \  ZcZdZeZf eZdU¦  «        \  ZgZhZiZjdV„ ZkdW„ ZldX„ Zm G dY„ dZe¦  «        Zneod[k    r ejp        ¦   «          dS dS )\é    N)Údedent)ÚcudaÚuint32Úuint64Úfloat32Úfloat64)ÚunittestÚCUDATestCaseÚcc_X_or_above)ÚconfigT)Údevicec                 ó    — t          | ¦  «        S ©N)r   ©Únums    úZ/granary/edev/quantum/lib/python3.11/site-packages/numba/cuda/tests/cudapy/test_atomics.pyÚatomic_cast_to_uint64r   	   s   € å�#‰;Œ;Ðó    c                 ó    — t          | ¦  «        S r   )Úintr   s    r   Úatomic_cast_to_intr      s   € åˆs‰8Œ8€Or   c                 ó   — | S r   © r   s    r   Úatomic_cast_noner      s   € à€Jr   c	                 ó$  — t           j        j        }	t           j                             ||¦  «        }
||
|	<   t          j        ¦   «           |||	         |z  ¦  «        }|r||z
  } ||
||¦  «         t          j        ¦   «          |
|	         | |	<   d S r   ©r   Ú	threadIdxÚxÚsharedÚarrayÚsyncthreads)ÚaryÚidxÚop2Ú	ary_dtypeÚary_nelementsÚ
binop_funcÚ	cast_funcÚinitializerÚneg_idxÚtidÚsmÚbins               r   Úatomic_binary_1dim_sharedr.      s›   € õ Œ.Ô
€CÝ	Œ×	Ò	˜=¨)Ñ	4Ô	4€BØ€B€s�GÝÔÑÔÐØ
ˆ)�C˜”H˜}Ñ,Ñ
-Ô
-€CØð "Ø�MÑ!ˆØ€Jˆr�3˜ÑÔÐÝÔÑÔÐØ�#Œw€Cˆ�H€H€Hr   c                 ó"  — t           j        j        }t           j                             ||¦  «        }| |         ||<   t          j        ¦   «           |||         |z  ¦  «        }	 |||	|¦  «         t          j        ¦   «          ||         | |<   d S r   r   )
r"   r#   r$   r%   r&   r'   r(   r+   r,   r-   s
             r   Úatomic_binary_1dim_shared2r0   (   s�   € õ Œ.Ô
€CÝ	Œ×	Ò	˜=¨)Ñ	4Ô	4€BØ�#Œh€B€s�GÝÔÑÔÐØ
ˆ)�C˜”H˜}Ñ,Ñ
-Ô
-€CØ€Jˆr�3˜ÑÔÐÝÔÑÔÐØ�#Œw€Cˆ�H€H€Hr   c                 óŽ  — t           j        j        }t           j        j        }t           j                             ||¦  «        }	| ||f         |	||f<   t          j        ¦   «          | ||¦  «        f}
|r"|
d         |d         z
  |
d         |d         z
  f}
 ||	|
|¦  «         t          j        ¦   «          |	||f         | ||f<   d S ©Nr   é   )r   r   r   Úyr   r    r!   )r"   r$   r%   Ú	ary_shaper'   Úy_cast_funcr*   ÚtxÚtyr,   r-   s              r   Úatomic_binary_2dim_sharedr9   5   sÍ   € õ 
ŒÔ	€BÝ	ŒÔ	€BÝ	Œ×	Ò	˜9 iÑ	0Ô	0€BØ�R˜�V”€B€rˆ2€v�JÝÔÑÔÐØˆ{ˆ{˜2‰ŒÐ
€CØð =Ø�1Œv˜	 !œÑ$ c¨!¤f¨y¸¬|Ñ&;Ð<ˆØ€Jˆr�3˜ÑÔÐÝÔÑÔÐØ�R˜�V”*€CˆˆBˆ�K€K€Kr   c                 óÚ   — t           j        j        }t           j        j        }| ||¦  «        f}|r,|d         | j        d         z
  |d         | j        d         z
  f} || ||¦  «         d S r2   )r   r   r   r4   Úshape)r"   r$   r'   r6   r*   r7   r8   r-   s           r   Úatomic_binary_2dim_globalr<   E   sr   € å	ŒÔ	€BÝ	ŒÔ	€BØˆ{ˆ{˜2‰ŒÐ
€CØð =Ø�1Œv˜œ	 !œÑ$ c¨!¤f¨s¬y¸¬|Ñ&;Ð<ˆØ€Jˆs�C˜ÑÔÐÐÐr   c                 ó€   — t           j        j        }t          ||         |z  ¦  «        }|r||z
  } || ||¦  «         d S r   )r   r   r   r   )r"   r#   r&   r$   r'   r*   r+   r-   s           r   Úatomic_binary_1dim_globalr>   O   sP   € õ Œ.Ô
€CÝ
ˆc�#Œh˜Ñ&Ñ
'Ô
'€CØð "Ø�MÑ!ˆØ€Jˆs�C˜ÑÔÐÐÐr   c                 óf   — t          | | dt          dt          j        j        t
          dd¦	  «	         d S ©Nr3   é    r   F©r.   r   r   ÚatomicÚaddr   ©r"   s    r   Ú
atomic_addrF   Y   ó;   € Ý˜c 3¨­6°2Ý"œkœoÕ/?ÀÀEñKô Kð Kð Kð Kr   c                 óf   — t          | | dt          dt          j        j        t
          dd¦	  «	         d S )Nr3   rA   r   TrB   rE   s    r   Úatomic_add_wraprI   ^   s;   € Ý˜c 3¨­6°2Ý"œkœoÕ/?ÀÀDñJô Jð Jð Jð Jr   c           	      ób   — t          | dt          dt          j        j        t
          d¦  «         d S ©Nr3   ©é   é   F©r9   r   r   rC   rD   r   rE   s    r   Úatomic_add2rP   c   ó7   € Ý˜c 1¥f¨fÝ"œkœoÕ/?ÀñHô Hð Hð Hð Hr   c           	      ób   — t          | dt          dt          j        j        t
          d¦  «         d S )Nr3   rL   TrO   rE   s    r   Úatomic_add2_wraprS   h   s7   € Ý˜c 1¥f¨fÝ"œkœoÕ/?ÀñGô Gð Gð Gð Gr   c           	      ób   — t          | dt          dt          j        j        t
          d¦  «         d S rK   )r9   r   r   rC   rD   r   rE   s    r   Úatomic_add3rU   m   ó7   € Ý˜c 1¥f¨fÝ"œkœoÕ/DÀeñMô Mð Mð Mð Mr   c                 óf   — t          | | dt          dt          j        j        t
          dd¦	  «	         d S ©Nç      ð?rA   ç        F©r.   r   r   rC   rD   r   rE   s    r   Úatomic_add_floatr\   r   ó;   € Ý˜c 3¨­W°bÝ"œkœoÕ/AÀ3ÈñOô Oð Oð Oð Or   c                 óf   — t          | | dt          dt          j        j        t
          dd¦	  «	         d S ©NrY   rA   rZ   Tr[   rE   s    r   Úatomic_add_float_wrapr`   w   s;   € Ý˜c 3¨­W°bÝ"œkœoÕ/AÀ3ÈñNô Nð Nð Nð Nr   c           	      ób   — t          | dt          dt          j        j        t
          d¦  «         d S ©NrY   rL   F©r9   r   r   rC   rD   r   rE   s    r   Úatomic_add_float_2rd   |   ó7   € Ý˜c 3­°Ý"œkœoÕ/?ÀñHô Hð Hð Hð Hr   c           	      ób   — t          | dt          dt          j        j        t
          d¦  «         d S ©NrY   rL   Trc   rE   s    r   Úatomic_add_float_2_wraprh   �   ó7   € Ý˜c 3­°Ý"œkœoÕ/?ÀñGô Gð Gð Gð Gr   c           	      ób   — t          | dt          dt          j        j        t
          d¦  «         d S rb   )r9   r   r   rC   rD   r   rE   s    r   Úatomic_add_float_3rk   †   ó7   € Ý˜c 3­°Ý"œkœoÕ/DÀeñMô Mð Mð Mð Mr   c                 óL   — t          || ddt          j        j        d¦  «         d S ©NrA   rY   F©r>   r   rC   rD   ©r#   r"   s     r   Úatomic_add_double_globalrq   ‹   ó$   € Ý˜c 3¨¨Cµ´´À%ÑHÔHÐHÐHÐHr   c                 óL   — t          || ddt          j        j        d¦  «         d S )NrA   rY   Tro   rp   s     r   Úatomic_add_double_global_wraprt   �   s$   € Ý˜c 3¨¨Cµ´´À$ÑGÔGÐGÐGÐGr   c                 óT   — t          | dt          j        j        t          d¦  «         d S ©Nr3   F©r<   r   rC   rD   r   rE   s    r   Úatomic_add_double_global_2rx   “   s#   € Ý˜c 1¥d¤k¤oÕ7GÈÑOÔOÐOÐOÐOr   c                 óT   — t          | dt          j        j        t          d¦  «         d S )Nr3   Trw   rE   s    r   Úatomic_add_double_global_2_wraprz   —   s#   € Ý˜c 1¥d¤k¤oÕ7GÈÑNÔNÐNÐNÐNr   c                 óT   — t          | dt          j        j        t          d¦  «         d S rv   )r<   r   rC   rD   r   rE   s    r   Úatomic_add_double_global_3r|   ›   s.   € Ý˜c 1¥d¤k¤oÕ7LØ#ñ%ô %ð %ð %ð %r   c                 óf   — t          || dt          dt          j        j        t
          dd¦	  «	         d S rX   ©r.   r   r   rC   rD   r   rp   s     r   Úatomic_add_doubler       ó;   € Ý˜c 3¨­W°bÝ"œkœoÕ/?ÀÀeñMô Mð Mð Mð Mr   c                 óf   — t          || dt          dt          j        j        t
          dd¦	  «	         d S r_   r~   rp   s     r   Úatomic_add_double_wrapr‚   ¥   s;   € Ý˜c 3¨­W°bÝ"œkœoÕ/?ÀÀdñLô Lð Lð Lð Lr   c           	      ób   — t          | dt          dt          j        j        t
          d¦  «         d S rb   ©r9   r   r   rC   rD   r   rE   s    r   Úatomic_add_double_2r…   ª   re   r   c           	      ób   — t          | dt          dt          j        j        t
          d¦  «         d S rg   r„   rE   s    r   Úatomic_add_double_2_wrapr‡   ¯   ri   r   c           	      ób   — t          | dt          dt          j        j        t
          d¦  «         d S rb   )r9   r   r   rC   rD   r   rE   s    r   Úatomic_add_double_3r‰   ´   rl   r   c                 óf   — t          | | dt          dt          j        j        t
          dd¦	  «	         d S r@   )r.   r   r   rC   Úsubr   rE   s    r   Ú
atomic_subrŒ   ¹   rG   r   c           	      ób   — t          | dt          dt          j        j        t
          d¦  «         d S rK   )r9   r   r   rC   r‹   r   rE   s    r   Úatomic_sub2rŽ   ¾   rQ   r   c           	      ób   — t          | dt          dt          j        j        t
          d¦  «         d S rK   )r9   r   r   rC   r‹   r   rE   s    r   Úatomic_sub3r�   Ã   rV   r   c                 óf   — t          | | dt          dt          j        j        t
          dd¦	  «	         d S rX   )r.   r   r   rC   r‹   r   rE   s    r   Úatomic_sub_floatr’   È   r]   r   c           	      ób   — t          | dt          dt          j        j        t
          d¦  «         d S rb   )r9   r   r   rC   r‹   r   rE   s    r   Úatomic_sub_float_2r”   Í   re   r   c           	      ób   — t          | dt          dt          j        j        t
          d¦  «         d S rb   )r9   r   r   rC   r‹   r   rE   s    r   Úatomic_sub_float_3r–   Ò   rl   r   c                 óf   — t          || dt          dt          j        j        t
          dd¦	  «	         d S rX   )r.   r   r   rC   r‹   r   rp   s     r   Úatomic_sub_doubler˜   ×   r€   r   c           	      ób   — t          | dt          dt          j        j        t
          d¦  «         d S rb   )r9   r   r   rC   r‹   r   rE   s    r   Úatomic_sub_double_2rš   Ü   re   r   c           	      ób   — t          | dt          dt          j        j        t
          d¦  «         d S rb   ©r9   r   r   rC   r‹   r   rE   s    r   Úatomic_sub_double_3r�   á   rl   r   c                 óL   — t          || ddt          j        j        d¦  «         d S rn   )r>   r   rC   r‹   rp   s     r   Úatomic_sub_double_globalrŸ   æ   rr   r   c                 óT   — t          | dt          j        j        t          d¦  «         d S )NrY   F)r<   r   rC   r‹   r   rE   s    r   Úatomic_sub_double_global_2r¡   ê   s.   € Ý˜c 3­¬¬Õ9IØ#ñ%ô %ð %ð %ð %r   c           	      ób   — t          | dt          dt          j        j        t
          d¦  «         d S rb   rœ   rE   s    r   Úatomic_sub_double_global_3r£   ï   rl   r   c                 óf   — t          | | |t          dt          j        j        t
          dd¦	  «	         d S )NrA   r3   F)r.   r   r   rC   Úand_r   ©r"   r$   s     r   Ú
atomic_andr§   ô   s<   € Ý˜c 3¨­V°RÝ"œkÔ.Õ0@À!ÀUñLô Lð Lð Lð Lr   c           	      ób   — t          | |t          dt          j        j        t
          d¦  «         d S ©NrL   F)r9   r   r   rC   r¥   r   r¦   s     r   Úatomic_and2rª   ù   ó8   € Ý˜c 3­°Ý"œkÔ.Õ0@À%ñIô Ið Ið Ið Ir   c           	      ób   — t          | |t          dt          j        j        t
          d¦  «         d S r©   )r9   r   r   rC   r¥   r   r¦   s     r   Úatomic_and3r­   þ   s8   € Ý˜c 3­°Ý"œkÔ.Õ0EÀuñNô Nð Nð Nð Nr   c                 óL   — t          || d|t          j        j        d¦  «         d S ©NrA   F)r>   r   rC   r¥   ©r#   r"   r$   s      r   Úatomic_and_globalr±     ó%   € Ý˜c 3¨¨Cµ´Ô1AÀ5ÑIÔIÐIÐIÐIr   c                 óT   — t          | |t          j        j        t          d¦  «         d S ©NF)r<   r   rC   r¥   r   r¦   s     r   Úatomic_and_global_2rµ     s.   € Ý˜c 3­¬Ô(8Ý.°ñ7ô 7ð 7ð 7ð 7r   c                 óf   — t          | | |t          dt          j        j        t
          dd¦	  «	         d S ©NrA   r   F)r.   r   r   rC   Úor_r   r¦   s     r   Ú	atomic_orr¹     ó;   € Ý˜c 3¨­V°RÝ"œkœoÕ/?ÀÀEñKô Kð Kð Kð Kr   c           	      ób   — t          | |t          dt          j        j        t
          d¦  «         d S r©   )r9   r   r   rC   r¸   r   r¦   s     r   Ú
atomic_or2r¼     ó7   € Ý˜c 3­°Ý"œkœoÕ/?ÀñHô Hð Hð Hð Hr   c           	      ób   — t          | |t          dt          j        j        t
          d¦  «         d S r©   )r9   r   r   rC   r¸   r   r¦   s     r   Ú
atomic_or3r¿     ó7   € Ý˜c 3­°Ý"œkœoÕ/DÀeñMô Mð Mð Mð Mr   c                 óL   — t          || d|t          j        j        d¦  «         d S r¯   )r>   r   rC   r¸   r°   s      r   Úatomic_or_globalrÂ     rr   r   c                 óT   — t          | |t          j        j        t          d¦  «         d S r´   )r<   r   rC   r¸   r   r¦   s     r   Úatomic_or_global_2rÄ     ó-   € Ý˜c 3­¬¬Ý.°ñ7ô 7ð 7ð 7ð 7r   c                 óf   — t          | | |t          dt          j        j        t
          dd¦	  «	         d S r·   )r.   r   r   rC   Úxorr   r¦   s     r   Ú
atomic_xorrÈ   $  rº   r   c           	      ób   — t          | |t          dt          j        j        t
          d¦  «         d S r©   )r9   r   r   rC   rÇ   r   r¦   s     r   Úatomic_xor2rÊ   )  r½   r   c           	      ób   — t          | |t          dt          j        j        t
          d¦  «         d S r©   )r9   r   r   rC   rÇ   r   r¦   s     r   Úatomic_xor3rÌ   .  rÀ   r   c                 óL   — t          || d|t          j        j        d¦  «         d S r¯   )r>   r   rC   rÇ   r°   s      r   Úatomic_xor_globalrÎ   3  rr   r   c                 óT   — t          | |t          j        j        t          d¦  «         d S r´   )r<   r   rC   rÇ   r   r¦   s     r   Úatomic_xor_global_2rÐ   7  rÅ   r   c           	      ób   — t          | ||t          dt          j        j        t
          ¦  «         d S ©NrA   )r0   r   r   rC   Úincr   ©r"   r#   r$   s      r   Úatomic_inc32rÕ   <  ó7   € Ý˜s C¨­f°bÝ#œ{œÕ0@ñBô Bð Bð Bð Br   c           	      ób   — t          | ||t          dt          j        j        t
          ¦  «         d S rÒ   )r0   r   r   rC   rÓ   r   rÔ   s      r   Úatomic_inc64rØ   A  ó7   € Ý˜s C¨­f°bÝ#œ{œÕ0BñDô Dð Dð Dð Dr   c           	      ób   — t          | |t          dt          j        j        t
          d¦  «         d S r©   )r9   r   r   rC   rÓ   r   r¦   s     r   Úatomic_inc2_32rÛ   F  r½   r   c           	      ób   — t          | |t          dt          j        j        t
          d¦  «         d S r©   )r9   r   r   rC   rÓ   r   r¦   s     r   Úatomic_inc2_64rÝ   K  r½   r   c           	      ób   — t          | |t          dt          j        j        t
          d¦  «         d S r©   )r9   r   r   rC   rÓ   r   r¦   s     r   Úatomic_inc3rß   P  rÀ   r   c                 óL   — t          || d|t          j        j        d¦  «         d S r¯   )r>   r   rC   rÓ   r°   s      r   Úatomic_inc_globalrá   U  rr   r   c                 óT   — t          | |t          j        j        t          d¦  «         d S r´   )r<   r   rC   rÓ   r   r¦   s     r   Úatomic_inc_global_2rã   Y  rÅ   r   c           	      ób   — t          | ||t          dt          j        j        t
          ¦  «         d S rÒ   )r0   r   r   rC   Údecr   rÔ   s      r   Úatomic_dec32ræ   ^  rÖ   r   c           	      ób   — t          | ||t          dt          j        j        t
          ¦  «         d S rÒ   )r0   r   r   rC   rå   r   rÔ   s      r   Úatomic_dec64rè   c  rÙ   r   c           	      ób   — t          | |t          dt          j        j        t
          d¦  «         d S r©   )r9   r   r   rC   rå   r   r¦   s     r   Úatomic_dec2_32rê   h  r½   r   c           	      ób   — t          | |t          dt          j        j        t
          d¦  «         d S r©   )r9   r   r   rC   rå   r   r¦   s     r   Úatomic_dec2_64rì   m  r½   r   c           	      ób   — t          | |t          dt          j        j        t
          d¦  «         d S r©   )r9   r   r   rC   rå   r   r¦   s     r   Úatomic_dec3rî   r  rÀ   r   c                 óL   — t          || d|t          j        j        d¦  «         d S r¯   )r>   r   rC   rå   r°   s      r   Úatomic_dec_globalrð   w  rr   r   c                 óT   — t          | |t          j        j        t          d¦  «         d S r´   )r<   r   rC   rå   r   r¦   s     r   Úatomic_dec_global_2rò   {  rÅ   r   c           	      ób   — t          | ||t          dt          j        j        t
          ¦  «         d S rÒ   )r0   r   r   rC   Úexchr   rÔ   s      r   Úatomic_exchrõ   €  s8   € Ý˜s C¨­f°bÝ#œ{Ô/Õ1AñCô Cð Cð Cð Cr   c           	      ób   — t          | |t          dt          j        j        t
          d¦  «         d S r©   )r9   r   r   rC   rô   r   r¦   s     r   Úatomic_exch2r÷   …  r«   r   c           	      ób   — t          | |t          dt          j        j        t
          d¦  «         d S r©   )r9   r   r   rC   rô   r   r¦   s     r   Úatomic_exch3rù   Š  r«   r   c                 óL   — t          || d|t          j        j        d¦  «         d S r¯   )r>   r   rC   rô   r°   s      r   Úatomic_exch_globalrû   �  r²   r   c                 óÐ   — t          d¦  «                             | ¬¦  «        }i }t          |t          t          t
          dœ|¦  «         |d         |d         |d         |d         fS )Na  
    def atomic(res, ary):
        tx = cuda.threadIdx.x
        bx = cuda.blockIdx.x
        {func}(res, 0, ary[tx, bx])

    def atomic_double_normalizedindex(res, ary):
        tx = cuda.threadIdx.x
        bx = cuda.blockIdx.x
        {func}(res, 0, ary[tx, uint64(bx)])

    def atomic_double_oneindex(res, ary):
        tx = cuda.threadIdx.x
        {func}(res, 0, ary[tx])

    def atomic_double_shared(res, ary):
        tid = cuda.threadIdx.x
        smary = cuda.shared.array(32, float64)
        smary[tid] = ary[tid]
        smres = cuda.shared.array(1, float64)
        if tid == 0:
            smres[0] = res[0]
        cuda.syncthreads()
        {func}(smres, 0, smary[tid])
        cuda.syncthreads()
        if tid == 0:
            res[0] = smres[0]
    )Úfunc)r   r   r   rC   Úatomic_double_normalizedindexÚatomic_double_oneindexÚatomic_double_shared)r   ÚformatÚexecr   r   r   )rý   ÚfnsÚlds      r   Úgen_atomic_extreme_funcsr  “  sy   € å
ð ñ 	ô 	÷6 
Š�TˆÑ	Ô	ð7 ð8 
€BÝˆ•t­½6ÐBÐBÀBÑGÔGÐGØˆxŒL˜"Ð<Ô=ØÐ'Ô(¨"Ð-CÔ*DðFð Fr   zcuda.atomic.maxzcuda.atomic.minzcuda.atomic.nanmaxzcuda.atomic.nanminc                 ó¬   — t          j        d¦  «        }|| j        k     r4t           j                             | |d …         |||         ¦  «        ||<   d S d S ©Nr3   )r   ÚgridÚsizerC   Úcompare_and_swap©ÚresÚoldr"   Úfill_valÚgids        r   Úatomic_compare_and_swapr  Ã  sN   € Ý
Œ)�A‰,Œ,€CØ
ˆSŒX‚~€~Ý”;×/Ò/°°C°D°D´	¸8ÀSÈÄXÑNÔNˆˆC‰ˆˆð €~r   c                 óž   — t          j        d¦  «        }|| j        k     r-t           j                             | ||||         ¦  «        ||<   d S d S r  )r   r  r	  rC   Úcasr  s        r   Úatomic_cas_1dimr  É  sF   € Ý
Œ)�A‰,Œ,€CØ
ˆSŒX‚~€~Ý”;—?’? 3¨¨X°s¸3´xÑ@Ô@ˆˆC‰ˆˆð €~r   c                 óè   — t          j        d¦  «        }|d         | j        d         k     rD|d         | j        d         k     r/t           j                             | ||||         ¦  «        ||<   d S d S d S )Né   r   r3   )r   r  r;   rC   r  r  s        r   Úatomic_cas_2dimr  Ï  sk   € Ý
Œ)�A‰,Œ,€CØ
ˆ1„v�”	˜!”ÒÐ  Q¤¨#¬)°A¬,Ò!6Ð!6Ý”;—?’? 3¨¨X°s¸3´xÑ@Ô@ˆˆC‰ˆˆð ÐÐ!6Ð!6r   c                   óŽ  ‡ — e Zd Zˆ fd„Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Z	d–d	„Z
d
„ Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Z d „ Z!d!„ Z"d"„ Z#d#„ Z$d$„ Z%d%„ Z&d&„ Z'd'„ Z(d(„ Z)d)„ Z*d*„ Z+d+„ Z,d,„ Z-d-„ Z.d.„ Z/d/„ Z0d0„ Z1d1„ Z2d2„ Z3d3„ Z4d4„ Z5d5„ Z6d6„ Z7d7„ Z8d8„ Z9d9„ Z:d:„ Z;d;„ Z<d<„ Z=d=„ Z>d>„ Z?d?„ Z@d@„ ZAdA„ ZBdB„ ZCdC„ ZDdD„ ZEdE„ ZFdF„ ZGdG„ ZHdH„ ZIdI„ ZJdJ„ ZKdK„ ZLdL„ ZMdM„ ZNdN„ ZOdO„ ZPdP„ ZQdQ„ ZRdR„ ZSdS„ ZTdT„ ZUdU„ ZVdV„ ZWdW„ ZXdX„ ZYdY„ ZZdZ„ Z[d[„ Z\d\„ Z]d]„ Z^d^„ Z_d_„ Z`d`„ Zada„ Zbdb„ Zcd—dd„Zdde„ Zedf„ Zfdg„ Zgdh„ Zhdi„ Zidj„ Zjdk„ Zkdl„ Zldm„ Zmdn„ Zndo„ Zodp„ Zpdq„ Zqdr„ Zrds„ Zsdt„ Ztdu„ Zudv„ Zvdw„ Zwdx„ Zxdy„ Zydz„ Zzd{„ 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•ˆ xZ–S )˜ÚTestCudaAtomicsc                 ó†   •— t          ¦   «                              ¦   «          t          j                             d¦  «         d S )Nr   )ÚsuperÚsetUpÚnpÚrandomÚseed)ÚselfÚ	__class__s    €r   r  zTestCudaAtomics.setUpÖ  s.   ø€ Ý‰Œ�Š‰ŒˆÝ
Œ	�Š�qÑÔÐÐÐr   c                 óæ  — t           j                             ddd¬¦  «                             t           j        ¦  «        }|                     ¦   «         }|                     ¦   «         } t          j        d¦  «        t          ¦  «        } |d         |¦  «          t          j        d¦  «        t          ¦  «        } |d         |¦  «         t          j
        dt           j        ¬¦  «        }t          |j        ¦  «        D ]}|||         xx         dz  cc<   Œ|                      t          j        ||k    ¦  «        ¦  «         |                      t          j        ||k    ¦  «        ¦  «         d S ©Nr   rA   ©r	  zvoid(uint32[:])©r3   rA   ©Údtyper3   )r  r  ÚrandintÚastyper   Úcopyr   ÚjitrF   rI   ÚzerosÚranger	  Ú
assertTrueÚall)r  r"   Úary_wrapÚorigÚcuda_atomic_addÚcuda_atomic_add_wrapÚgoldÚis           r   Útest_atomic_addzTestCudaAtomics.test_atomic_addÚ  s:  € ÝŒi×Ò  2¨BÐÑ/Ô/×6Ò6µr´yÑAÔAˆØ—8’8‘:”:ˆØ�xŠx‰zŒzˆà5�$œ(Ð#4Ñ5Ô5µjÑAÔAˆØˆ˜Ô˜sÑ#Ô#Ð#à:�tœxÐ(9Ñ:Ô:½?ÑKÔKÐØ#Ð˜UÔ# HÑ-Ô-Ð-åŒx˜¥"¤)Ð,Ñ,Ô,ˆÝ�t”yÑ!Ô!ð 	ð 	ˆAØ��a”ˆMˆMŒM˜QÑˆMˆM‰MˆMà�Š�œ˜s dš{Ñ+Ô+Ñ,Ô,Ð,Ø�Š�œ˜x¨4Ò/Ñ0Ô0Ñ1Ô1Ð1Ð1Ð1r   c                 ó€  — t           j                             ddd¬¦  «                             t           j        ¦  «                             dd¦  «        }|                     ¦   «         }|                     ¦   «         } t          j        d¦  «        t          ¦  «        } |d         |¦  «          t          j        d¦  «        t          ¦  «        } |d         |¦  «         |                      t          j        ||dz   k    ¦  «        ¦  «         |                      t          j        ||dz   k    ¦  «        ¦  «         d S ©	Nr   rA   r#  rM   rN   zvoid(uint32[:,:])©r3   rL   r3   )r  r  r'  r(  r   Úreshaper)  r   r*  rP   rS   r-  r.  )r  r"   r/  r0  Úcuda_atomic_add2Úcuda_atomic_add2_wraps         r   Útest_atomic_add2z TestCudaAtomics.test_atomic_add2ì  s
  € ÝŒi×Ò  2¨BÐÑ/Ô/×6Ò6µr´yÑAÔA×IÒIÈ!ÈQÑOÔOˆØ—8’8‘:”:ˆØ�xŠx‰zŒzˆà8�4œ8Ð$7Ñ8Ô8½ÑEÔEÐØ#Ð˜Ô# CÑ(Ô(Ð(à =¥¤Ð)<Ñ =Ô =Õ>NÑ OÔ OÐØ(Ð˜iÔ(¨Ñ2Ô2Ð2à�Š�œ˜s d¨Q¡hšÑ/Ô/Ñ0Ô0Ð0Ø�Š�œ˜x¨4°!©8Ò3Ñ4Ô4Ñ5Ô5Ð5Ð5Ð5r   c                 ó–  — t           j                             ddd¬¦  «                             t           j        ¦  «                             dd¦  «        }|                     ¦   «         } t          j        d¦  «        t          ¦  «        } |d         |¦  «         |  
                    t          j        ||dz   k    ¦  «        ¦  «         d S r7  )r  r  r'  r(  r   r9  r)  r   r*  rU   r-  r.  ©r  r"   r0  Úcuda_atomic_add3s       r   Útest_atomic_add3z TestCudaAtomics.test_atomic_add3ú  s¦   € ÝŒi×Ò  2¨BÐÑ/Ô/×6Ò6µr´yÑAÔA×IÒIÈ!ÈQÑOÔOˆØ�xŠx‰zŒzˆØ8�4œ8Ð$7Ñ8Ô8½ÑEÔEÐØ#Ð˜Ô# CÑ(Ô(Ð(à�Š�œ˜s d¨Q¡hšÑ/Ô/Ñ0Ô0Ð0Ð0Ð0r   c                 ó   — t           j                             ddd¬¦  «                             t           j        ¦  «        }|                     ¦   «         }|                     ¦   «                              t           j        ¦  «        } t          j        d¦  «        t          ¦  «        } |d         |¦  «          t          j        d¦  «        t          ¦  «        } |d         |¦  «         t          j        dt           j        ¬¦  «        }t          |j        ¦  «        D ]}|||         xx         dz  cc<   Œ|                      t          j        ||k    ¦  «        ¦  «         |                      t          j        ||k    ¦  «        ¦  «         d S ©Nr   rA   r#  zvoid(float32[:])r$  r%  rY   )r  r  r'  r(  r   r)  Úintpr   r*  r\   r`   r+  r   r,  r	  r-  r.  )r  r"   r/  r0  Úcuda_atomic_add_floatÚadd_float_wrapr3  r4  s           r   Útest_atomic_add_floatz%TestCudaAtomics.test_atomic_add_float  sL  € ÝŒi×Ò  2¨BÐÑ/Ô/×6Ò6µr´zÑBÔBˆØ—8’8‘:”:ˆØ�xŠx‰zŒz× Ò ¥¤Ñ)Ô)ˆà <¥¤Ð);Ñ <Ô <Õ=MÑ NÔ NÐØ$Ð˜eÔ$ SÑ)Ô)Ð)à5�œÐ"4Ñ5Ô5Õ6KÑLÔLˆØˆ�uÔ˜hÑ'Ô'Ð'åŒx˜¥"¤)Ð,Ñ,Ô,ˆÝ�t”yÑ!Ô!ð 	!ð 	!ˆAØ��a”ˆMˆMŒM˜SÑ ˆMˆM‰MˆMà�Š�œ˜s dš{Ñ+Ô+Ñ,Ô,Ð,Ø�Š�œ˜x¨4Ò/Ñ0Ô0Ñ1Ô1Ð1Ð1Ð1r   c                 ó€  — t           j                             ddd¬¦  «                             t           j        ¦  «                             dd¦  «        }|                     ¦   «         }|                     ¦   «         } t          j        d¦  «        t          ¦  «        } |d         |¦  «          t          j        d¦  «        t          ¦  «        } |d         |¦  «         |                      t          j        ||dz   k    ¦  «        ¦  «         |                      t          j        ||dz   k    ¦  «        ¦  «         d S ©	Nr   rA   r#  rM   rN   zvoid(float32[:,:])r8  r3   )r  r  r'  r(  r   r9  r)  r   r*  rd   rh   r-  r.  )r  r"   r/  r0  r:  Úcuda_func_wraps         r   Útest_atomic_add_float_2z'TestCudaAtomics.test_atomic_add_float_2  s	  € ÝŒi×Ò  2¨BÐÑ/Ô/×6Ò6µr´zÑBÔB×JÒJÈ1ÈaÑPÔPˆØ—8’8‘:”:ˆØ�xŠx‰zŒzˆà9�4œ8Ð$8Ñ9Ô9Õ:LÑMÔMÐØ#Ð˜Ô# CÑ(Ô(Ð(à7�œÐ"6Ñ7Ô7Õ8OÑPÔPˆØ!ˆ�yÔ! (Ñ+Ô+Ð+à�Š�œ˜s d¨Q¡hšÑ/Ô/Ñ0Ô0Ð0Ø�Š�œ˜x¨4°!©8Ò3Ñ4Ô4Ñ5Ô5Ð5Ð5Ð5r   c                 ó–  — t           j                             ddd¬¦  «                             t           j        ¦  «                             dd¦  «        }|                     ¦   «         } t          j        d¦  «        t          ¦  «        } |d         |¦  «         |  
                    t          j        ||dz   k    ¦  «        ¦  «         d S rH  )r  r  r'  r(  r   r9  r)  r   r*  rk   r-  r.  r>  s       r   Útest_atomic_add_float_3z'TestCudaAtomics.test_atomic_add_float_3"  s§   € ÝŒi×Ò  2¨BÐÑ/Ô/×6Ò6µr´zÑBÔB×JÒJÈ1ÈaÑPÔPˆØ�xŠx‰zŒzˆØ9�4œ8Ð$8Ñ9Ô9Õ:LÑMÔMÐØ#Ð˜Ô# CÑ(Ô(Ð(à�Š�œ˜s d¨Q¡hšÑ/Ô/Ñ0Ô0Ð0Ð0Ð0r   Tc                 ó´  — t           j        rd S t          t          |                     ¦   «                              ¦   «         ¦  «        ¦  «        }t          dd¦  «        rIt          j         	                    ¦   «         dk    rd}nd}|r|› d�}|  
                    |› d�|¦  «         d S |r|  
                    d|¦  «         d S |  
                    d	|¦  «         d S )
Né   r   )é   r3   ÚredÚatomz.sharedz.add.f64zatom.shared.cas.b64zatom.cas.b64)r   ÚENABLE_CUDASIMÚnextÚiterÚinspect_asmÚvaluesr   r   ÚruntimeÚget_versionÚassertIn)r  Úkernelr   ÚasmÚinsts        r   ÚassertCorrectFloat64Atomicsz+TestCudaAtomics.assertCorrectFloat64Atomics*  sð   € ÝÔ ð 	ØˆFõ •4˜×*Ò*Ñ,Ô,×3Ò3Ñ5Ô5Ñ6Ô6Ñ7Ô7ˆÝ˜˜AÑÔð 	3ÝŒ|×'Ò'Ñ)Ô)¨GÒ3Ð3ð ��à�àð (ØÐ'Ð'Ð'�à�MŠM˜TÐ+Ð+Ð+¨SÑ1Ô1Ð1Ð1Ð1àð 3Ø—’Ð3°SÑ9Ô9Ð9Ð9Ð9à—’˜n¨cÑ2Ô2Ð2Ð2Ð2r   c                 ó  — t           j                             dddt           j        ¬¦  «        }t          j        dt           j        ¦  «        }|                     ¦   «         } t          j        d¦  «        t          ¦  «        } |d         ||¦  «          t          j        d¦  «        t          ¦  «        } |d         ||¦  «         t          j        dt           j        ¬¦  «        }t          |j        ¦  «        D ]}|||         xx         dz  cc<   Œt           j                             ||¦  «         t           j                             ||¦  «         |                      |¦  «         |                      |¦  «         d S ©Nr   rA   ©r	  r&  úvoid(int64[:], float64[:])r$  r%  rY   )r  r  r'  Úint64r+  r   r)  r   r*  r   r‚   r   r,  r	  ÚtestingÚassert_equalr]  )r  r#   r"   r/  Úcuda_fnÚwrap_fnr3  r4  s           r   Útest_atomic_add_doublez&TestCudaAtomics.test_atomic_add_doubleC  sK  € ÝŒi×Ò  2¨Bµb´hÐÑ?Ô?ˆÝŒh�r�2œ:Ñ&Ô&ˆØ—8’8‘:”:ˆà8•$”(Ð7Ñ8Ô8Õ9JÑKÔKˆØˆ�Œ�s˜CÑ Ô Ð à8•$”(Ð7Ñ8Ô8Õ9OÑPÔPˆØˆ�Œ�s˜HÑ%Ô%Ð%åŒx˜¥"¤)Ð,Ñ,Ô,ˆÝ�s”x‘”ð 	 ð 	 ˆAØ��Q”ˆLˆLŒL˜CÑˆLˆL‰LˆLå
Œ
×Ò  TÑ*Ô*Ð*Ý
Œ
×Ò ¨$Ñ/Ô/Ð/Ø×(Ò(¨Ñ1Ô1Ð1Ø×(Ò(¨Ñ1Ô1Ð1Ð1Ð1r   c                 ó¨  — t           j                             ddd¬¦  «                             t           j        ¦  «                             dd¦  «        }|                     ¦   «         }|                     ¦   «         } t          j        d¦  «        t          ¦  «        } |d         |¦  «          t          j        d¦  «        t          ¦  «        } |d         |¦  «         t           j                             ||dz   ¦  «         t           j                             ||dz   ¦  «         |                      |¦  «         |                      |¦  «         d S ©	Nr   rA   r#  rM   rN   úvoid(float64[:,:])r8  r3   )r  r  r'  r(  r   r9  r)  r   r*  r…   r‡   rc  rd  r]  )r  r"   r/  r0  re  Úcuda_fn_wraps         r   Útest_atomic_add_double_2z(TestCudaAtomics.test_atomic_add_double_2W  s  € ÝŒi×Ò  2¨BÐÑ/Ô/×6Ò6µr´zÑBÔB×JÒJÈ1ÈaÑPÔPˆØ—8’8‘:”:ˆØ�xŠx‰zŒzˆà0•$”(Ð/Ñ0Ô0Õ1DÑEÔEˆØˆ�	Ô˜3ÑÔÐà5•t”xÐ 4Ñ5Ô5Õ6NÑOÔOˆØˆ�YÔ Ñ)Ô)Ð)å
Œ
×Ò  T¨A¡XÑ.Ô.Ð.Ý
Œ
×Ò ¨$°©(Ñ3Ô3Ð3Ø×(Ò(¨Ñ1Ô1Ð1Ø×(Ò(¨Ñ6Ô6Ð6Ð6Ð6r   c                 óª  — t           j                             ddd¬¦  «                             t           j        ¦  «                             dd¦  «        }|                     ¦   «         } t          j        d¦  «        t          ¦  «        } |d         |¦  «         t           j
                             ||dz   ¦  «         |                      |¦  «         d S ri  )r  r  r'  r(  r   r9  r)  r   r*  r‰   rc  rd  r]  ©r  r"   r0  Ú	cuda_funcs       r   Útest_atomic_add_double_3z(TestCudaAtomics.test_atomic_add_double_3g  s±   € ÝŒi×Ò  2¨BÐÑ/Ô/×6Ò6µr´zÑBÔB×JÒJÈ1ÈaÑPÔPˆØ�xŠx‰zŒzˆØ2•D”HÐ1Ñ2Ô2Õ3FÑGÔGˆ	Øˆ	�)Ô˜SÑ!Ô!Ð!å
Œ
×Ò  T¨A¡XÑ.Ô.Ð.Ø×(Ò(¨Ñ3Ô3Ð3Ð3Ð3r   c                 ó  — t           j                             dddt           j        ¬¦  «        }t          j        dt           j        ¦  «        }|                     ¦   «         }d} t          j        |¦  «        t          ¦  «        } t          j        |¦  «        t          ¦  «        } |d         ||¦  «          |d         ||¦  «         t          j        dt           j        ¬¦  «        }t          |j        ¦  «        D ]}|||         xx         dz  cc<   Œt           j                             ||¦  «         t           j                             ||¦  «         |                      |d¬	¦  «         |                      |d¬	¦  «         d S )
Nr   rA   r`  ra  r$  r%  rY   F©r   )r  r  r'  rb  r+  r   r)  r   r*  rq   rt   r   r,  r	  rc  rd  r]  )	r  r#   r"   r/  Úsigro  Úwrap_cuda_funcr3  r4  s	            r   Útest_atomic_add_double_globalz-TestCudaAtomics.test_atomic_add_double_globalp  sV  € ÝŒi×Ò  2¨Bµb´hÐÑ?Ô?ˆÝŒh�r�2œ:Ñ&Ô&ˆØ—8’8‘:”:ˆà*ˆØ!•D”H˜S‘M”MÕ":Ñ;Ô;ˆ	Ø&�œ #™œÕ'DÑEÔEˆàˆ	�%Ô˜˜cÑ"Ô"Ð"Øˆ�uÔ˜c 8Ñ,Ô,Ð,åŒx˜¥"¤)Ð,Ñ,Ô,ˆÝ�s”x‘”ð 	 ð 	 ˆAØ��Q”ˆLˆLŒL˜CÑˆLˆL‰LˆLå
Œ
×Ò  TÑ*Ô*Ð*Ý
Œ
×Ò ¨$Ñ/Ô/Ð/Ø×(Ò(¨¸5Ð(ÑAÔAÐAØ×(Ò(¨ÀÐ(ÑFÔFÐFÐFÐFr   c                 ó´  — t           j                             ddd¬¦  «                             t           j        ¦  «                             dd¦  «        }|                     ¦   «         }|                     ¦   «         }d} t          j        |¦  «        t          ¦  «        } t          j        |¦  «        t          ¦  «        } |d         |¦  «          |d         |¦  «         t           j                             ||dz   ¦  «         t           j                             ||dz   ¦  «         |                      |d	¬
¦  «         |                      |d	¬
¦  «         d S ©Nr   rA   r#  rM   rN   rj  r8  r3   Frr  )r  r  r'  r(  r   r9  r)  r   r*  rx   rz   rc  rd  r]  )r  r"   r/  r0  rs  ro  rt  s          r   Útest_atomic_add_double_global_2z/TestCudaAtomics.test_atomic_add_double_global_2…  s'  € ÝŒi×Ò  2¨BÐÑ/Ô/×6Ò6µr´zÑBÔB×JÒJÈ1ÈaÑPÔPˆØ—8’8‘:”:ˆØ�xŠx‰zŒzˆà"ˆØ!•D”H˜S‘M”MÕ"<Ñ=Ô=ˆ	Ø&�œ #™œÕ'FÑGÔGˆàˆ	�)Ô˜SÑ!Ô!Ð!Ø!ˆ�yÔ! (Ñ+Ô+Ð+å
Œ
×Ò  T¨A¡XÑ.Ô.Ð.Ý
Œ
×Ò ¨$°©(Ñ3Ô3Ð3Ø×(Ò(¨¸5Ð(ÑAÔAÐAØ×(Ò(¨ÀÐ(ÑFÔFÐFÐFÐFr   c                 ó®  — t           j                             ddd¬¦  «                             t           j        ¦  «                             dd¦  «        }|                     ¦   «         } t          j        d¦  «        t          ¦  «        } |d         |¦  «         t           j
                             ||dz   ¦  «         |                      |d	¬
¦  «         d S rw  )r  r  r'  r(  r   r9  r)  r   r*  r|   rc  rd  r]  rn  s       r   Útest_atomic_add_double_global_3z/TestCudaAtomics.test_atomic_add_double_global_3–  s¶   € ÝŒi×Ò  2¨BÐÑ/Ô/×6Ò6µr´zÑBÔB×JÒJÈ1ÈaÑPÔPˆØ�xŠx‰zŒzˆØ2•D”HÐ1Ñ2Ô2Õ3MÑNÔNˆ	Øˆ	�)Ô˜SÑ!Ô!Ð!å
Œ
×Ò  T¨A¡XÑ.Ô.Ð.Ø×(Ò(¨¸5Ð(ÑAÔAÐAÐAÐAr   c                 ó  — t           j                             ddd¬¦  «                             t           j        ¦  «        }|                     ¦   «         } t          j        d¦  «        t          ¦  «        } |d         |¦  «         t          j	        dt           j        ¬¦  «        }t          |j        ¦  «        D ]}|||         xx         dz  cc<   Œ|                      t          j        ||k    ¦  «        ¦  «         d S r"  )r  r  r'  r(  r   r)  r   r*  rŒ   r+  r,  r	  r-  r.  )r  r"   r0  Úcuda_atomic_subr3  r4  s         r   Útest_atomic_subzTestCudaAtomics.test_atomic_subŸ  sÛ   € ÝŒi×Ò  2¨BÐÑ/Ô/×6Ò6µr´yÑAÔAˆØ�xŠx‰zŒzˆØ5�$œ(Ð#4Ñ5Ô5µjÑAÔAˆØˆ˜Ô˜sÑ#Ô#Ð#åŒx˜¥"¤)Ð,Ñ,Ô,ˆÝ�t”yÑ!Ô!ð 	ð 	ˆAØ��a”ˆMˆMŒM˜QÑˆMˆM‰MˆMà�Š�œ˜s dš{Ñ+Ô+Ñ,Ô,Ð,Ð,Ð,r   c                 ó–  — t           j                             ddd¬¦  «                             t           j        ¦  «                             dd¦  «        }|                     ¦   «         } t          j        d¦  «        t          ¦  «        } |d         |¦  «         |  
                    t          j        ||dz
  k    ¦  «        ¦  «         d S r7  )r  r  r'  r(  r   r9  r)  r   r*  rŽ   r-  r.  ©r  r"   r0  Úcuda_atomic_sub2s       r   Útest_atomic_sub2z TestCudaAtomics.test_atomic_sub2«  ó¦   € ÝŒi×Ò  2¨BÐÑ/Ô/×6Ò6µr´yÑAÔA×IÒIÈ!ÈQÑOÔOˆØ�xŠx‰zŒzˆØ8�4œ8Ð$7Ñ8Ô8½ÑEÔEÐØ#Ð˜Ô# CÑ(Ô(Ð(Ø�Š�œ˜s d¨Q¡hšÑ/Ô/Ñ0Ô0Ð0Ð0Ð0r   c                 ó–  — t           j                             ddd¬¦  «                             t           j        ¦  «                             dd¦  «        }|                     ¦   «         } t          j        d¦  «        t          ¦  «        } |d         |¦  «         |  
                    t          j        ||dz
  k    ¦  «        ¦  «         d S r7  )r  r  r'  r(  r   r9  r)  r   r*  r�   r-  r.  ©r  r"   r0  Úcuda_atomic_sub3s       r   Útest_atomic_sub3z TestCudaAtomics.test_atomic_sub3²  r‚  r   c                 ó<  — t           j                             ddd¬¦  «                             t           j        ¦  «        }|                     ¦   «                              t           j        ¦  «        } t          j        d¦  «        t          ¦  «        } |d         |¦  «         t          j
        dt           j        ¬¦  «        }t          |j        ¦  «        D ]}|||         xx         dz  cc<   Œ|                      t          j        ||k    ¦  «        ¦  «         d S rB  )r  r  r'  r(  r   r)  rC  r   r*  r’   r+  r,  r	  r-  r.  )r  r"   r0  Úcuda_atomic_sub_floatr3  r4  s         r   Útest_atomic_sub_floatz%TestCudaAtomics.test_atomic_sub_float¹  sî   € ÝŒi×Ò  2¨BÐÑ/Ô/×6Ò6µr´zÑBÔBˆØ�xŠx‰zŒz× Ò ¥¤Ñ)Ô)ˆØ <¥¤Ð);Ñ <Ô <Õ=MÑ NÔ NÐØ$Ð˜eÔ$ SÑ)Ô)Ð)åŒx˜¥"¤*Ð-Ñ-Ô-ˆÝ�t”yÑ!Ô!ð 	!ð 	!ˆAØ��a”ˆMˆMŒM˜SÑ ˆMˆM‰MˆMà�Š�œ˜s dš{Ñ+Ô+Ñ,Ô,Ð,Ð,Ð,r   c                 ó–  — t           j                             ddd¬¦  «                             t           j        ¦  «                             dd¦  «        }|                     ¦   «         } t          j        d¦  «        t          ¦  «        } |d         |¦  «         |  
                    t          j        ||dz
  k    ¦  «        ¦  «         d S rH  )r  r  r'  r(  r   r9  r)  r   r*  r”   r-  r.  r  s       r   Útest_atomic_sub_float_2z'TestCudaAtomics.test_atomic_sub_float_2Å  ó§   € ÝŒi×Ò  2¨BÐÑ/Ô/×6Ò6µr´zÑBÔB×JÒJÈ1ÈaÑPÔPˆØ�xŠx‰zŒzˆØ9�4œ8Ð$8Ñ9Ô9Õ:LÑMÔMÐØ#Ð˜Ô# CÑ(Ô(Ð(Ø�Š�œ˜s d¨Q¡hšÑ/Ô/Ñ0Ô0Ð0Ð0Ð0r   c                 ó–  — t           j                             ddd¬¦  «                             t           j        ¦  «                             dd¦  «        }|                     ¦   «         } t          j        d¦  «        t          ¦  «        } |d         |¦  «         |  
                    t          j        ||dz
  k    ¦  «        ¦  «         d S rH  )r  r  r'  r(  r   r9  r)  r   r*  r–   r-  r.  r„  s       r   Útest_atomic_sub_float_3z'TestCudaAtomics.test_atomic_sub_float_3Ì  rŒ  r   c                 óà  — t           j                             dddt           j        ¬¦  «        }t          j        dt           j        ¦  «        } t          j        d¦  «        t          ¦  «        } |d         ||¦  «         t          j        dt           j        ¬¦  «        }t          |j
        ¦  «        D ]}|||         xx         dz  cc<   Œt           j                             ||¦  «         d S r_  )r  r  r'  rb  r+  r   r   r*  r˜   r,  r	  rc  rd  )r  r#   r"   ro  r3  r4  s         r   Útest_atomic_sub_doublez&TestCudaAtomics.test_atomic_sub_doubleÓ  sÎ   € ÝŒi×Ò  2¨Bµb´hÐÑ?Ô?ˆÝŒh�r�2œ:Ñ&Ô&ˆØ:•D”HÐ9Ñ:Ô:Õ;LÑMÔMˆ	Øˆ	�%Ô˜˜cÑ"Ô"Ð"åŒx˜¥"¤*Ð-Ñ-Ô-ˆÝ�s”x‘”ð 	 ð 	 ˆAØ��Q”ˆLˆLŒL˜CÑˆLˆL‰LˆLå
Œ
×Ò  TÑ*Ô*Ð*Ð*Ð*r   c                 ó€  — t           j                             ddd¬¦  «                             t           j        ¦  «                             dd¦  «        }|                     ¦   «         } t          j        d¦  «        t          ¦  «        } |d         |¦  «         t           j
                             ||dz
  ¦  «         d S ri  )r  r  r'  r(  r   r9  r)  r   r*  rš   rc  rd  rn  s       r   Útest_atomic_sub_double_2z(TestCudaAtomics.test_atomic_sub_double_2ß  ó�   € ÝŒi×Ò  2¨BÐÑ/Ô/×6Ò6µr´zÑBÔB×JÒJÈ1ÈaÑPÔPˆØ�xŠx‰zŒzˆØ2•D”HÐ1Ñ2Ô2Õ3FÑGÔGˆ	Øˆ	�)Ô˜SÑ!Ô!Ð!Ý
Œ
×Ò  T¨A¡XÑ.Ô.Ð.Ð.Ð.r   c                 ó€  — t           j                             ddd¬¦  «                             t           j        ¦  «                             dd¦  «        }|                     ¦   «         } t          j        d¦  «        t          ¦  «        } |d         |¦  «         t           j
                             ||dz
  ¦  «         d S ri  )r  r  r'  r(  r   r9  r)  r   r*  r�   rc  rd  rn  s       r   Útest_atomic_sub_double_3z(TestCudaAtomics.test_atomic_sub_double_3æ  r“  r   c                 óä  — t           j                             dddt           j        ¬¦  «        }t          j        dt           j        ¦  «        }d} t          j        |¦  «        t          ¦  «        } |d         ||¦  «         t          j        dt           j        ¬¦  «        }t          |j
        ¦  «        D ]}|||         xx         dz  cc<   Œt           j                             ||¦  «         d S r_  )r  r  r'  rb  r+  r   r   r*  rŸ   r,  r	  rc  rd  )r  r#   r"   rs  ro  r3  r4  s          r   Útest_atomic_sub_double_globalz-TestCudaAtomics.test_atomic_sub_double_globalí  sÐ   € ÝŒi×Ò  2¨Bµb´hÐÑ?Ô?ˆÝŒh�r�2œ:Ñ&Ô&ˆØ*ˆØ!•D”H˜S‘M”MÕ":Ñ;Ô;ˆ	Øˆ	�%Ô˜˜cÑ"Ô"Ð"åŒx˜¥"¤*Ð-Ñ-Ô-ˆÝ�s”x‘”ð 	 ð 	 ˆAØ��Q”ˆLˆLŒL˜CÑˆLˆL‰LˆLå
Œ
×Ò  TÑ*Ô*Ð*Ð*Ð*r   c                 ó€  — t           j                             ddd¬¦  «                             t           j        ¦  «                             dd¦  «        }|                     ¦   «         } t          j        d¦  «        t          ¦  «        } |d         |¦  «         t           j
                             ||dz
  ¦  «         d S ri  )r  r  r'  r(  r   r9  r)  r   r*  r¡   rc  rd  rn  s       r   Útest_atomic_sub_double_global_2z/TestCudaAtomics.test_atomic_sub_double_global_2ú  ó�   € ÝŒi×Ò  2¨BÐÑ/Ô/×6Ò6µr´zÑBÔB×JÒJÈ1ÈaÑPÔPˆØ�xŠx‰zŒzˆØ2•D”HÐ1Ñ2Ô2Õ3MÑNÔNˆ	Øˆ	�)Ô˜SÑ!Ô!Ð!Ý
Œ
×Ò  T¨A¡XÑ.Ô.Ð.Ð.Ð.r   c                 ó€  — t           j                             ddd¬¦  «                             t           j        ¦  «                             dd¦  «        }|                     ¦   «         } t          j        d¦  «        t          ¦  «        } |d         |¦  «         t           j
                             ||dz
  ¦  «         d S ri  )r  r  r'  r(  r   r9  r)  r   r*  r£   rc  rd  rn  s       r   Útest_atomic_sub_double_global_3z/TestCudaAtomics.test_atomic_sub_double_global_3  rš  r   c                 ó*  — t           j                             d¦  «        }t           j                             ddd¬¦  «                             t           j        ¦  «        }|                     ¦   «         } t          j        d¦  «        t          ¦  «        } |d         ||¦  «         |                     ¦   «         }t          |j
        ¦  «        D ]}|||         xx         |z  cc<   Œ|                      t          j        ||k    ¦  «        ¦  «         d S )Néô  r   rA   r#  úvoid(uint32[:], uint32)r$  )r  r  r'  r(  r   r)  r   r*  r§   r,  r	  r-  r.  ©r  Ú
rand_constr"   r0  ro  r3  r4  s          r   Útest_atomic_andzTestCudaAtomics.test_atomic_and  sé   € Ý”Y×&Ò& sÑ+Ô+ˆ
ÝŒi×Ò  2¨BÐÑ/Ô/×6Ò6µr´yÑAÔAˆØ�xŠx‰zŒzˆØ7•D”HÐ6Ñ7Ô7½
ÑCÔCˆ	Øˆ	�%Ô˜˜jÑ)Ô)Ð)à�xŠx‰zŒzˆÝ�t”yÑ!Ô!ð 	(ð 	(ˆAØ��a”ˆMˆMŒM˜ZÑ'ˆMˆM‰MˆMà�Š�œ˜s dš{Ñ+Ô+Ñ,Ô,Ð,Ð,Ð,r   c                 óÖ  — t           j                             d¦  «        }t           j                             ddd¬¦  «                             t           j        ¦  «                             dd¦  «        }|                     ¦   «         } t          j        d¦  «        t          ¦  «        } |d         ||¦  «         |  
                    t          j        |||z  k    ¦  «        ¦  «         d S ©	Nrž  r   rA   r#  rM   rN   úvoid(uint32[:,:], uint32)r8  )r  r  r'  r(  r   r9  r)  r   r*  rª   r-  r.  ©r  r¡  r"   r0  Úcuda_atomic_and2s        r   Útest_atomic_and2z TestCudaAtomics.test_atomic_and2  ó¿   € Ý”Y×&Ò& sÑ+Ô+ˆ
ÝŒi×Ò  2¨BÐÑ/Ô/×6Ò6µr´yÑAÔA×IÒIÈ!ÈQÑOÔOˆØ�xŠx‰zŒzˆØ@�4œ8Ð$?Ñ@Ô@ÅÑMÔMÐØ#Ð˜Ô# C¨Ñ4Ô4Ð4Ø�Š�œ˜s d¨ZÑ&7Ò7Ñ8Ô8Ñ9Ô9Ð9Ð9Ð9r   c                 óÖ  — t           j                             d¦  «        }t           j                             ddd¬¦  «                             t           j        ¦  «                             dd¦  «        }|                     ¦   «         } t          j        d¦  «        t          ¦  «        } |d         ||¦  «         |  
                    t          j        |||z  k    ¦  «        ¦  «         d S r¤  )r  r  r'  r(  r   r9  r)  r   r*  r­   r-  r.  ©r  r¡  r"   r0  Úcuda_atomic_and3s        r   Útest_atomic_and3z TestCudaAtomics.test_atomic_and3  r©  r   c                 ó(  — t           j                             d¦  «        }t           j                             dddt           j        ¬¦  «        }t           j                             dddt           j        ¬¦  «        }d} t	          j        |¦  «        t          ¦  «        } |d         |||¦  «         |                     ¦   «         }t          |j	        ¦  «        D ]}|||         xx         |z  cc<   Œt           j
                             ||¦  «         d S ©Nrž  r   rA   r`  zvoid(int32[:], int32[:], int32)r$  )r  r  r'  Úint32r   r*  r±   r)  r,  r	  rc  rd  ©r  r¡  r#   r"   rs  ro  r3  r4  s           r   Útest_atomic_and_globalz&TestCudaAtomics.test_atomic_and_global%  së   € Ý”Y×&Ò& sÑ+Ô+ˆ
ÝŒi×Ò  2¨Bµb´hÐÑ?Ô?ˆÝŒi×Ò  2¨Bµb´hÐÑ?Ô?ˆØ/ˆØ!•D”H˜S‘M”MÕ"3Ñ4Ô4ˆ	Øˆ	�%Ô˜˜c :Ñ.Ô.Ð.à�xŠx‰zŒzˆÝ�s”x‘”ð 	'ð 	'ˆAØ��Q”ˆLˆLŒL˜JÑ&ˆLˆL‰LˆLå
Œ
×Ò  TÑ*Ô*Ð*Ð*Ð*r   c                 óÀ  — t           j                             d¦  «        }t           j                             ddd¬¦  «                             t           j        ¦  «                             dd¦  «        }|                     ¦   «         } t          j        d¦  «        t          ¦  «        } |d         ||¦  «         t           j
                             |||z  ¦  «         d S r¤  )r  r  r'  r(  r   r9  r)  r   r*  rµ   rc  rd  ©r  r¡  r"   r0  ro  s        r   Útest_atomic_and_global_2z(TestCudaAtomics.test_atomic_and_global_23  óµ   € Ý”Y×&Ò& sÑ+Ô+ˆ
ÝŒi×Ò  2¨BÐÑ/Ô/×6Ò6µr´yÑAÔA×IÒIÈ!ÈQÑOÔOˆØ�xŠx‰zŒzˆØ9•D”HÐ8Ñ9Ô9Õ:MÑNÔNˆ	Øˆ	�)Ô˜S *Ñ-Ô-Ð-Ý
Œ
×Ò  T¨JÑ%6Ñ7Ô7Ð7Ð7Ð7r   c                 óB  — t           j                             d¦  «        }t           j                             ddd¬¦  «                             t           j        ¦  «        }|                     ¦   «         } t          j        d¦  «        t          ¦  «        } |d         ||¦  «         t          j	        dt           j        ¬¦  «        }t          |j        ¦  «        D ]}|||         xx         |z  cc<   Œ|                      t          j        ||k    ¦  «        ¦  «         d S ©Nrž  r   rA   r#  rŸ  r$  r%  )r  r  r'  r(  r   r)  r   r*  r¹   r+  r,  r	  r-  r.  r   s          r   Útest_atomic_orzTestCudaAtomics.test_atomic_or;  sò   € Ý”Y×&Ò& sÑ+Ô+ˆ
ÝŒi×Ò  2¨BÐÑ/Ô/×6Ò6µr´yÑAÔAˆØ�xŠx‰zŒzˆØ7•D”HÐ6Ñ7Ô7½	ÑBÔBˆ	Øˆ	�%Ô˜˜jÑ)Ô)Ð)åŒx˜¥"¤)Ð,Ñ,Ô,ˆÝ�t”yÑ!Ô!ð 	(ð 	(ˆAØ��a”ˆMˆMŒM˜ZÑ'ˆMˆM‰MˆMà�Š�œ˜s dš{Ñ+Ô+Ñ,Ô,Ð,Ð,Ð,r   c                 óÖ  — t           j                             d¦  «        }t           j                             ddd¬¦  «                             t           j        ¦  «                             dd¦  «        }|                     ¦   «         } t          j        d¦  «        t          ¦  «        } |d         ||¦  «         |  
                    t          j        |||z  k    ¦  «        ¦  «         d S r¤  )r  r  r'  r(  r   r9  r)  r   r*  r¼   r-  r.  r¦  s        r   Útest_atomic_or2zTestCudaAtomics.test_atomic_or2H  ó¿   € Ý”Y×&Ò& sÑ+Ô+ˆ
ÝŒi×Ò  2¨BÐÑ/Ô/×6Ò6µr´yÑAÔA×IÒIÈ!ÈQÑOÔOˆØ�xŠx‰zŒzˆØ@�4œ8Ð$?Ñ@Ô@ÅÑLÔLÐØ#Ð˜Ô# C¨Ñ4Ô4Ð4Ø�Š�œ˜s d¨ZÑ&7Ò7Ñ8Ô8Ñ9Ô9Ð9Ð9Ð9r   c                 óÖ  — t           j                             d¦  «        }t           j                             ddd¬¦  «                             t           j        ¦  «                             dd¦  «        }|                     ¦   «         } t          j        d¦  «        t          ¦  «        } |d         ||¦  «         |  
                    t          j        |||z  k    ¦  «        ¦  «         d S r¤  )r  r  r'  r(  r   r9  r)  r   r*  r¿   r-  r.  r«  s        r   Útest_atomic_or3zTestCudaAtomics.test_atomic_or3P  r¼  r   c                 ó(  — t           j                             d¦  «        }t           j                             dddt           j        ¬¦  «        }t           j                             dddt           j        ¬¦  «        }d} t	          j        |¦  «        t          ¦  «        } |d         |||¦  «         |                     ¦   «         }t          |j	        ¦  «        D ]}|||         xx         |z  cc<   Œt           j
                             ||¦  «         d S r¯  )r  r  r'  r°  r   r*  rÂ   r)  r,  r	  rc  rd  r±  s           r   Útest_atomic_or_globalz%TestCudaAtomics.test_atomic_or_globalX  së   € Ý”Y×&Ò& sÑ+Ô+ˆ
ÝŒi×Ò  2¨Bµb´hÐÑ?Ô?ˆÝŒi×Ò  2¨Bµb´hÐÑ?Ô?ˆØ/ˆØ!•D”H˜S‘M”MÕ"2Ñ3Ô3ˆ	Øˆ	�%Ô˜˜c :Ñ.Ô.Ð.à�xŠx‰zŒzˆÝ�s”x‘”ð 	'ð 	'ˆAØ��Q”ˆLˆLŒL˜JÑ&ˆLˆL‰LˆLå
Œ
×Ò  TÑ*Ô*Ð*Ð*Ð*r   c                 óÀ  — t           j                             d¦  «        }t           j                             ddd¬¦  «                             t           j        ¦  «                             dd¦  «        }|                     ¦   «         } t          j        d¦  «        t          ¦  «        } |d         ||¦  «         t           j
                             |||z  ¦  «         d S r¤  )r  r  r'  r(  r   r9  r)  r   r*  rÄ   rc  rd  r´  s        r   Útest_atomic_or_global_2z'TestCudaAtomics.test_atomic_or_global_2f  sµ   € Ý”Y×&Ò& sÑ+Ô+ˆ
ÝŒi×Ò  2¨BÐÑ/Ô/×6Ò6µr´yÑAÔA×IÒIÈ!ÈQÑOÔOˆØ�xŠx‰zŒzˆØ9•D”HÐ8Ñ9Ô9Õ:LÑMÔMˆ	Øˆ	�)Ô˜S *Ñ-Ô-Ð-Ý
Œ
×Ò  T¨JÑ%6Ñ7Ô7Ð7Ð7Ð7r   c                 óB  — t           j                             d¦  «        }t           j                             ddd¬¦  «                             t           j        ¦  «        }|                     ¦   «         } t          j        d¦  «        t          ¦  «        } |d         ||¦  «         t          j	        dt           j        ¬¦  «        }t          |j        ¦  «        D ]}|||         xx         |z  cc<   Œ|                      t          j        ||k    ¦  «        ¦  «         d S r¸  )r  r  r'  r(  r   r)  r   r*  rÈ   r+  r,  r	  r-  r.  r   s          r   Útest_atomic_xorzTestCudaAtomics.test_atomic_xorn  sò   € Ý”Y×&Ò& sÑ+Ô+ˆ
ÝŒi×Ò  2¨BÐÑ/Ô/×6Ò6µr´yÑAÔAˆØ�xŠx‰zŒzˆØ7•D”HÐ6Ñ7Ô7½
ÑCÔCˆ	Øˆ	�%Ô˜˜jÑ)Ô)Ð)åŒx˜¥"¤)Ð,Ñ,Ô,ˆÝ�t”yÑ!Ô!ð 	(ð 	(ˆAØ��a”ˆMˆMŒM˜ZÑ'ˆMˆM‰MˆMà�Š�œ˜s dš{Ñ+Ô+Ñ,Ô,Ð,Ð,Ð,r   c                 óÖ  — t           j                             d¦  «        }t           j                             ddd¬¦  «                             t           j        ¦  «                             dd¦  «        }|                     ¦   «         } t          j        d¦  «        t          ¦  «        } |d         ||¦  «         |  
                    t          j        |||z  k    ¦  «        ¦  «         d S r¤  )r  r  r'  r(  r   r9  r)  r   r*  rÊ   r-  r.  )r  r¡  r"   r0  Úcuda_atomic_xor2s        r   Útest_atomic_xor2z TestCudaAtomics.test_atomic_xor2{  r©  r   c                 óÖ  — t           j                             d¦  «        }t           j                             ddd¬¦  «                             t           j        ¦  «                             dd¦  «        }|                     ¦   «         } t          j        d¦  «        t          ¦  «        } |d         ||¦  «         |  
                    t          j        |||z  k    ¦  «        ¦  «         d S r¤  )r  r  r'  r(  r   r9  r)  r   r*  rÌ   r-  r.  )r  r¡  r"   r0  Úcuda_atomic_xor3s        r   Útest_atomic_xor3z TestCudaAtomics.test_atomic_xor3ƒ  r©  r   c                 ó(  — t           j                             d¦  «        }t           j                             dddt           j        ¬¦  «        }t           j                             dddt           j        ¬¦  «        }|                     ¦   «         }d} t          j        |¦  «        t          ¦  «        } |d         |||¦  «         t          |j	        ¦  «        D ]}|||         xx         |z  cc<   Œt           j
                             ||¦  «         d S r¯  )r  r  r'  r°  r)  r   r*  rÎ   r,  r	  rc  rd  )r  r¡  r#   r"   r3  rs  ro  r4  s           r   Útest_atomic_xor_globalz&TestCudaAtomics.test_atomic_xor_global‹  së   € Ý”Y×&Ò& sÑ+Ô+ˆ
ÝŒi×Ò  2¨Bµb´hÐÑ?Ô?ˆÝŒi×Ò  2¨Bµb´hÐÑ?Ô?ˆØ�xŠx‰zŒzˆØ/ˆØ!•D”H˜S‘M”MÕ"3Ñ4Ô4ˆ	Øˆ	�%Ô˜˜c :Ñ.Ô.Ð.å�s”x‘”ð 	'ð 	'ˆAØ��Q”ˆLˆLŒL˜JÑ&ˆLˆL‰LˆLå
Œ
×Ò  TÑ*Ô*Ð*Ð*Ð*r   c                 óÀ  — t           j                             d¦  «        }t           j                             ddd¬¦  «                             t           j        ¦  «                             dd¦  «        }|                     ¦   «         } t          j        d¦  «        t          ¦  «        } |d         ||¦  «         t           j
                             |||z  ¦  «         d S r¤  )r  r  r'  r(  r   r9  r)  r   r*  rÐ   rc  rd  r´  s        r   Útest_atomic_xor_global_2z(TestCudaAtomics.test_atomic_xor_global_2™  r¶  r   c                 óä   — t           j                             d|¬¦  «        }t           j                             ddd¬¦  «                             |¦  «        }t          j        d|¬¦  «        }|||fS )NrA   r%  r   r#  )r  r  r'  r(  Úarange)r  r&  ÚrconstÚraryÚary_idxs        r   Úinc_dec_1dim_setupz"TestCudaAtomics.inc_dec_1dim_setup¡  sg   € Ý”×"Ò" 2¨eÐ"Ñ4Ô4ˆÝŒy× Ò   B¨RÐ Ñ0Ô0×7Ò7¸Ñ>Ô>ˆÝ”)˜B eÐ,Ñ,Ô,ˆØ�t˜WÐ$Ð$r   c                 óÞ   — t           j                             d|¬¦  «        }t           j                             ddd¬¦  «                             |¦  «                             dd¦  «        }||fS )NrA   r%  r   r#  rM   rN   )r  r  r'  r(  r9  )r  r&  rÑ  rÒ  s       r   Úinc_dec_2dim_setupz"TestCudaAtomics.inc_dec_2dim_setup§  s`   € Ý”×"Ò" 2¨UÐ"Ñ3Ô3ˆÝŒy× Ò   B¨RÐ Ñ0Ô0×7Ò7¸Ñ>Ô>×FÒFÀqÈ!ÑLÔLˆØ�tˆ|Ðr   c           	      ó  — |                      ¦   «         } t          j        |¦  «        |¦  «        }	 |	||f         |||¦  «         t          j                             |t          j        ||k    d|dz   ¦  «        ¦  «         d S r2   ©r)  r   r*  r  rc  rd  Úwhere©
r  r"   r#   rÑ  rs  ÚnblocksÚblksizerý   r0  ro  s
             r   Úcheck_inc_indexzTestCudaAtomics.check_inc_index¬  óx   € Ø�xŠx‰zŒzˆØ!•D”H˜S‘M”M $Ñ'Ô'ˆ	Ø#ˆ	�'˜7Ð"Ô# C¨¨fÑ5Ô5Ð5Ý
Œ
×Ò ¥R¤X¨d°fªn¸aÀÈÁÑ%JÔ%JÑKÔKÐKÐKÐKr   c           	      ó  — |                      ¦   «         } t          j        |¦  «        |¦  «        }	 |	||f         |||¦  «         t          j                             |t          j        ||k    d|dz   ¦  «        ¦  «         d S r2   rØ  rÚ  s
             r   Úcheck_inc_index2z TestCudaAtomics.check_inc_index2²  rÞ  r   c           	      ó  — |                      ¦   «         } t          j        |¦  «        |¦  «        } |||f         ||¦  «         t          j                             |t          j        ||k    d|dz   ¦  «        ¦  «         d S r2   rØ  ©	r  r"   rÑ  rs  rÛ  rÜ  rý   r0  ro  s	            r   Ú	check_inczTestCudaAtomics.check_inc¸  sv   € Ø�xŠx‰zŒzˆØ!•D”H˜S‘M”M $Ñ'Ô'ˆ	Ø#ˆ	�'˜7Ð"Ô# C¨Ñ0Ô0Ð0Ý
Œ
×Ò ¥R¤X¨d°fªn¸aÀÈÁÑ%JÔ%JÑKÔKÐKÐKÐKr   c           	      ó’   — |                       t          j        ¬¦  «        \  }}}d}|                      ||||ddt          ¦  «         d S ©Nr%  ú"void(uint32[:], uint32[:], uint32)r3   rA   )rÔ  r  r   rÝ  rÕ   ©r  r¡  r"   r#   rs  s        r   Útest_atomic_inc_32z"TestCudaAtomics.test_atomic_inc_32¾  óL   € Ø#×6Ò6½R¼YÐ6ÑGÔGÑˆ
�C˜Ø2ˆØ×Ò˜S # z°3¸¸2½|ÑLÔLÐLÐLÐLr   c           	      ó’   — |                       t          j        ¬¦  «        \  }}}d}|                      ||||ddt          ¦  «         d S ©Nr%  z"void(uint64[:], uint64[:], uint64)r3   rA   )rÔ  r  r   rÝ  rØ   rç  s        r   Útest_atomic_inc_64z"TestCudaAtomics.test_atomic_inc_64Ã  ré  r   c                 óŒ   — |                       t          j        ¦  «        \  }}d}|                      |||ddt          ¦  «         d S ©Nr¥  r3   rL   )rÖ  r  r   rã  rÛ   ©r  r¡  r"   rs  s       r   Útest_atomic_inc2_32z#TestCudaAtomics.test_atomic_inc2_32È  óB   € Ø×1Ò1µ"´)Ñ<Ô<‰ˆ
�CØ)ˆØ�Š�s˜J¨¨Q°µ~ÑFÔFÐFÐFÐFr   c                 óŒ   — |                       t          j        ¦  «        \  }}d}|                      |||ddt          ¦  «         d S ©Núvoid(uint64[:,:], uint64)r3   rL   )rÖ  r  r   rã  rÝ   rï  s       r   Útest_atomic_inc2_64z#TestCudaAtomics.test_atomic_inc2_64Í  rñ  r   c                 óŒ   — |                       t          j        ¦  «        \  }}d}|                      |||ddt          ¦  «         d S rî  )rÖ  r  r   rã  rß   rï  s       r   Útest_atomic_inc3z TestCudaAtomics.test_atomic_inc3Ò  óB   € Ø×1Ò1µ"´)Ñ<Ô<‰ˆ
�CØ)ˆØ�Š�s˜J¨¨Q°µ{ÑCÔCÐCÐCÐCr   c           	      ó’   — |                       t          j        ¬¦  «        \  }}}d}|                      ||||ddt          ¦  «         d S rå  )rÔ  r  r   rà  rá   rç  s        r   Útest_atomic_inc_global_32z)TestCudaAtomics.test_atomic_inc_global_32×  óW   € Ø#×6Ò6½R¼YÐ6ÑGÔGÑˆ
�C˜Ø2ˆØ×Ò˜c 3¨
°C¸¸BÝ/ñ	1ô 	1ð 	1ð 	1ð 	1r   c           	      ó’   — |                       t          j        ¬¦  «        \  }}}d}|                      ||||ddt          ¦  «         d S rë  )rÔ  r  r   rà  rá   rç  s        r   Útest_atomic_inc_global_64z)TestCudaAtomics.test_atomic_inc_global_64Ý  rû  r   c                 óŒ   — |                       t          j        ¦  «        \  }}d}|                      |||ddt          ¦  «         d S rî  )rÖ  r  r   rã  rã   rï  s       r   Útest_atomic_inc_global_2_32z+TestCudaAtomics.test_atomic_inc_global_2_32ã  óC   € Ø×1Ò1µ"´)Ñ<Ô<‰ˆ
�CØ)ˆØ�Š�s˜J¨¨Q°Õ7JÑKÔKÐKÐKÐKr   c                 óŒ   — |                       t          j        ¦  «        \  }}d}|                      |||ddt          ¦  «         d S ró  )rÖ  r  r   rã  rã   rï  s       r   Útest_atomic_inc_global_2_64z+TestCudaAtomics.test_atomic_inc_global_2_64è  r   r   c                 ó8  — |                      ¦   «         } t          j        |¦  «        |¦  «        }	 |	||f         |||¦  «         t          j                             |t          j        |dk    |t          j        ||k    ||dz
  ¦  «        ¦  «        ¦  «         d S r2   rØ  rÚ  s
             r   Úcheck_dec_indexzTestCudaAtomics.check_dec_indexí  ó¨   € Ø�xŠx‰zŒzˆØ!•D”H˜S‘M”M $Ñ'Ô'ˆ	Ø#ˆ	�'˜7Ð"Ô# C¨¨fÑ5Ô5Ð5Ý
Œ
×Ò ¥R¤X¨d°aªi¸Ý.0¬h°t¸f²}Ø7=Ø7;¸a±xñ/Aô /Añ&Bô &Bñ 	Cô 	Cð 	Cð 	Cð 	Cr   c                 ó8  — |                      ¦   «         } t          j        |¦  «        |¦  «        }	 |	||f         |||¦  «         t          j                             |t          j        |dk    |t          j        ||k    ||dz
  ¦  «        ¦  «        ¦  «         d S r2   rØ  rÚ  s
             r   Úcheck_dec_index2z TestCudaAtomics.check_dec_index2ö  r  r   c                 ó6  — |                      ¦   «         } t          j        |¦  «        |¦  «        } |||f         ||¦  «         t          j                             |t          j        |dk    |t          j        ||k    ||dz
  ¦  «        ¦  «        ¦  «         d S r2   rØ  râ  s	            r   Ú	check_deczTestCudaAtomics.check_decÿ  s¦   € Ø�xŠx‰zŒzˆØ!•D”H˜S‘M”M $Ñ'Ô'ˆ	Ø#ˆ	�'˜7Ð"Ô# C¨Ñ0Ô0Ð0Ý
Œ
×Ò ¥R¤X¨d°aªi¸Ý.0¬h°t¸f²}Ø7=Ø7;¸a±xñ/Aô /Añ&Bô &Bñ 	Cô 	Cð 	Cð 	Cð 	Cr   c           	      ó’   — |                       t          j        ¬¦  «        \  }}}d}|                      ||||ddt          ¦  «         d S rå  )rÔ  r  r   r  ræ   rç  s        r   Útest_atomic_dec_32z"TestCudaAtomics.test_atomic_dec_32  ré  r   c           	      ó’   — |                       t          j        ¬¦  «        \  }}}d}|                      ||||ddt          ¦  «         d S rë  )rÔ  r  r   r  rè   rç  s        r   Útest_atomic_dec_64z"TestCudaAtomics.test_atomic_dec_64  ré  r   c                 óŒ   — |                       t          j        ¦  «        \  }}d}|                      |||ddt          ¦  «         d S rî  )rÖ  r  r   r	  rê   rï  s       r   Útest_atomic_dec2_32z#TestCudaAtomics.test_atomic_dec2_32  rñ  r   c                 óŒ   — |                       t          j        ¦  «        \  }}d}|                      |||ddt          ¦  «         d S ró  )rÖ  r  r   r	  rì   rï  s       r   Útest_atomic_dec2_64z#TestCudaAtomics.test_atomic_dec2_64  rñ  r   c                 óŒ   — |                       t          j        ¦  «        \  }}d}|                      |||ddt          ¦  «         d S rî  )rÖ  r  r   r	  rî   rï  s       r   Útest_atomic_dec3_newz$TestCudaAtomics.test_atomic_dec3_new  rø  r   c           	      ó’   — |                       t          j        ¬¦  «        \  }}}d}|                      ||||ddt          ¦  «         d S rå  )rÔ  r  r   r  rð   rç  s        r   Útest_atomic_dec_global_32z)TestCudaAtomics.test_atomic_dec_global_32!  rû  r   c           	      ó’   — |                       t          j        ¬¦  «        \  }}}d}|                      ||||ddt          ¦  «         d S rë  )rÔ  r  r   r  rð   rç  s        r   Útest_atomic_dec_global_64z)TestCudaAtomics.test_atomic_dec_global_64'  rû  r   c                 óŒ   — |                       t          j        ¦  «        \  }}d}|                      |||ddt          ¦  «         d S rî  )rÖ  r  r   r	  rò   rï  s       r   Útest_atomic_dec_global2_32z*TestCudaAtomics.test_atomic_dec_global2_32-  r   r   c                 óŒ   — |                       t          j        ¦  «        \  }}d}|                      |||ddt          ¦  «         d S ró  )rÖ  r  r   r	  rò   rï  s       r   Útest_atomic_dec_global2_64z*TestCudaAtomics.test_atomic_dec_global2_642  r   r   c                 óÆ  — t           j                             ddt           j        ¬¦  «        }t           j                             ddd¬¦  «                             t           j        ¦  «        }t          j        dt           j        ¬¦  «        } t          j        d¦  «        t          ¦  «        } |d         |||¦  «         t           j	         
                    ||¦  «         d S )	Né2   éd   r%  r   rA   r#  ræ  r$  )r  r  r'  r   r(  rÐ  r   r*  rõ   rc  rd  )r  r¡  r"   r#   ro  s        r   Útest_atomic_exchz TestCudaAtomics.test_atomic_exch7  s³   € Ý”Y×&Ò& r¨3µb´iÐ&Ñ@Ô@ˆ
ÝŒi×Ò  2¨BÐÑ/Ô/×6Ò6µr´yÑAÔAˆÝŒi˜¥"¤)Ð,Ñ,Ô,ˆàB•D”HÐAÑBÔBÅ;ÑOÔOˆ	Øˆ	�%Ô˜˜c :Ñ.Ô.Ð.å
Œ
×Ò  ZÑ0Ô0Ð0Ð0Ð0r   c                 ó¬  — t           j                             ddt           j        ¬¦  «        }t           j                             ddd¬¦  «                             t           j        ¦  «                             dd¦  «        } t          j        d	¦  «        t          ¦  «        } |d
         ||¦  «         t           j	         
                    ||¦  «         d S )Nr  r  r%  r   rA   r#  rM   rN   r¥  r8  )r  r  r'  r   r(  r9  r   r*  r÷   rc  rd  ©r  r¡  r"   ro  s       r   Útest_atomic_exch2z!TestCudaAtomics.test_atomic_exch2A  ó«   € Ý”Y×&Ò& r¨3µb´iÐ&Ñ@Ô@ˆ
ÝŒi×Ò  2¨BÐÑ/Ô/×6Ò6µr´yÑAÔA×IÒIÈ!ÈQÑOÔOˆà9•D”HÐ8Ñ9Ô9½,ÑGÔGˆ	Øˆ	�)Ô˜S *Ñ-Ô-Ð-Ý
Œ
×Ò  ZÑ0Ô0Ð0Ð0Ð0r   c                 ó¬  — t           j                             ddt           j        ¬¦  «        }t           j                             ddd¬¦  «                             t           j        ¦  «                             dd¦  «        } t          j        d	¦  «        t          ¦  «        } |d
         ||¦  «         t           j	         
                    ||¦  «         d S )Nr  r  r%  r   rA   r#  rM   rN   rô  r8  )r  r  r'  r   r(  r9  r   r*  rù   rc  rd  r!  s       r   Útest_atomic_exch3z!TestCudaAtomics.test_atomic_exch3I  r#  r   c                 ó¦  — t           j                             ddt           j        ¬¦  «        }t          j        dt           j        ¬¦  «        }t           j                             dddt           j        ¬¦  «        }d} t          j        |¦  «        t          ¦  «        } |d         |||¦  «         t           j         	                    ||¦  «         d S )	Nr  r  r%  rA   r   r`  ræ  r$  )
r  r  r'  r   rÐ  r   r*  rû   rc  rd  )r  r¡  r#   r"   rs  ro  s         r   Útest_atomic_exch_globalz'TestCudaAtomics.test_atomic_exch_globalQ  sª   € Ý”Y×&Ò& r¨3µb´iÐ&Ñ@Ô@ˆ
ÝŒi˜¥"¤)Ð,Ñ,Ô,ˆÝŒi×Ò  2¨Bµb´iÐÑ@Ô@ˆà2ˆØ!•D”H˜S‘M”MÕ"4Ñ5Ô5ˆ	Øˆ	�%Ô˜˜c :Ñ.Ô.Ð.Ý
Œ
×Ò  ZÑ0Ô0Ð0Ð0Ð0r   c                 ód  — t           j                             ||d¬¦  «                             |¦  «        }t          j        d|j        ¬¦  «        }t          j        t          ¦  «        } |d         ||¦  «         t          j	        |¦  «        }t           j
                             ||¦  «         d S )N©rA   rA   r#  r3   r%  )r  r  r'  r(  r+  r&  r   r*  Ú
atomic_maxÚmaxrc  rd  ©r  r&  ÚloÚhiÚvalsr  ro  r3  s           r   Úcheck_atomic_maxz TestCudaAtomics.check_atomic_max[  s“   € ÝŒy× Ò   R¨hÐ Ñ7Ô7×>Ò>¸uÑEÔEˆÝŒh�q ¤
Ð+Ñ+Ô+ˆÝ”H�ZÑ(Ô(ˆ	Øˆ	�&Ô˜#˜tÑ$Ô$Ð$ÝŒv�d‰|Œ|ˆÝ
Œ
×Ò  TÑ*Ô*Ð*Ð*Ð*r   c                 óJ   — |                       t          j        dd¬¦  «         d S ©Né ÿÿéÿÿ  ©r&  r-  r.  )r0  r  r°  ©r  s    r   Útest_atomic_max_int32z%TestCudaAtomics.test_atomic_max_int32c  ó%   € Ø×Ò¥B¤H°¸EÐÑBÔBÐBÐBÐBr   c                 óJ   — |                       t          j        dd¬¦  «         d S ©Nr   r4  r5  )r0  r  r   r6  s    r   Útest_atomic_max_uint32z&TestCudaAtomics.test_atomic_max_uint32f  ó%   € Ø×Ò¥B¤I°!¸ÐÑ>Ô>Ð>Ð>Ð>r   c                 óJ   — |                       t          j        dd¬¦  «         d S r2  )r0  r  rb  r6  s    r   Útest_atomic_max_int64z%TestCudaAtomics.test_atomic_max_int64i  r8  r   c                 óJ   — |                       t          j        dd¬¦  «         d S r:  )r0  r  r   r6  s    r   Útest_atomic_max_uint64z&TestCudaAtomics.test_atomic_max_uint64l  r<  r   c                 óJ   — |                       t          j        dd¬¦  «         d S r2  )r0  r  r   r6  s    r   Útest_atomic_max_float32z'TestCudaAtomics.test_atomic_max_float32o  ó%   € Ø×Ò¥B¤J°6¸eÐÑDÔDÐDÐDÐDr   c                 óJ   — |                       t          j        dd¬¦  «         d S r2  )r0  r  r   r6  s    r   Útest_atomic_max_doublez&TestCudaAtomics.test_atomic_max_doubler  rC  r   c                 ó’  — t           j                             ddd¬¦  «                             t           j        ¦  «        }t          j        dt           j        ¦  «        } t          j        d¦  «        t          ¦  «        } |d         ||¦  «         t          j	        |¦  «        }t           j
                             ||¦  «         d S ©Nr   r4  r)  r#  r3   úvoid(float64[:], float64[:,:]))r  r  r'  r(  r   r+  r   r*  Ú!atomic_max_double_normalizedindexr+  rc  rd  ©r  r/  r  ro  r3  s        r   Ú&test_atomic_max_double_normalizedindexz6TestCudaAtomics.test_atomic_max_double_normalizedindexu  s¢   € ÝŒy× Ò   E°Ð Ñ9Ô9×@Ò@ÅÄÑLÔLˆÝŒh�q�"œ*Ñ%Ô%ˆØ>•D”HÐ=Ñ>Ô>Ý-ñ/ô /ˆ	àˆ	�&Ô˜#˜tÑ$Ô$Ð$åŒv�d‰|Œ|ˆÝ
Œ
×Ò  TÑ*Ô*Ð*Ð*Ð*r   c                 ó’  — t           j                             ddd¬¦  «                             t           j        ¦  «        }t          j        dt           j        ¦  «        } t          j        d¦  «        t          ¦  «        } |d         ||¦  «         t          j	        |¦  «        }t           j
                             ||¦  «         d S ©Nr   é€   rA   r#  r3   úvoid(float64[:], float64[:])r$  )r  r  r'  r(  r   r+  r   r*  Úatomic_max_double_oneindexr+  rc  rd  rJ  s        r   Útest_atomic_max_double_oneindexz/TestCudaAtomics.test_atomic_max_double_oneindex  s¢   € ÝŒy× Ò   C¨bÐ Ñ1Ô1×8Ò8½¼ÑDÔDˆÝŒh�q�"œ*Ñ%Ô%ˆØ<•D”HÐ;Ñ<Ô<Ý&ñ(ô (ˆ	àˆ	�%Ô˜˜dÑ#Ô#Ð#åŒv�d‰|Œ|ˆÝ
Œ
×Ò  TÑ*Ô*Ð*Ð*Ð*r   c                 óf  — t           j                             ||d¬¦  «                             |¦  «        }t          j        dg|j        ¬¦  «        }t          j        t          ¦  «        } |d         ||¦  «         t          j	        |¦  «        }t           j
                             ||¦  «         d S )Nr)  r#  r4  r%  )r  r  r'  r(  r    r&  r   r*  Ú
atomic_minÚminrc  rd  r,  s           r   Úcheck_atomic_minz TestCudaAtomics.check_atomic_min‰  s•   € ÝŒy× Ò   R¨hÐ Ñ7Ô7×>Ò>¸uÑEÔEˆÝŒh˜�w d¤jÐ1Ñ1Ô1ˆÝ”H�ZÑ(Ô(ˆ	Øˆ	�&Ô˜#˜tÑ$Ô$Ð$åŒv�d‰|Œ|ˆÝ
Œ
×Ò  TÑ*Ô*Ð*Ð*Ð*r   c                 óJ   — |                       t          j        dd¬¦  «         d S r2  )rU  r  r°  r6  s    r   Útest_atomic_min_int32z%TestCudaAtomics.test_atomic_min_int32’  r8  r   c                 óJ   — |                       t          j        dd¬¦  «         d S r:  )rU  r  r   r6  s    r   Útest_atomic_min_uint32z&TestCudaAtomics.test_atomic_min_uint32•  r<  r   c                 óJ   — |                       t          j        dd¬¦  «         d S r2  )rU  r  rb  r6  s    r   Útest_atomic_min_int64z%TestCudaAtomics.test_atomic_min_int64˜  r8  r   c                 óJ   — |                       t          j        dd¬¦  «         d S r:  )rU  r  r   r6  s    r   Útest_atomic_min_uint64z&TestCudaAtomics.test_atomic_min_uint64›  r<  r   c                 óJ   — |                       t          j        dd¬¦  «         d S r2  )rU  r  r   r6  s    r   Útest_atomic_min_floatz%TestCudaAtomics.test_atomic_min_floatž  rC  r   c                 óJ   — |                       t          j        dd¬¦  «         d S r2  )rU  r  r   r6  s    r   Útest_atomic_min_doublez&TestCudaAtomics.test_atomic_min_double¡  rC  r   c                 ó˜  — t           j                             ddd¬¦  «                             t           j        ¦  «        }t          j        dt           j        ¦  «        dz  } t          j        d¦  «        t          ¦  «        } |d         ||¦  «         t          j	        |¦  «        }t           j
                             ||¦  «         d S rG  )r  r  r'  r(  r   Úonesr   r*  Ú!atomic_min_double_normalizedindexrT  rc  rd  rJ  s        r   Ú&test_atomic_min_double_normalizedindexz6TestCudaAtomics.test_atomic_min_double_normalizedindex¤  s§   € ÝŒy× Ò   E°Ð Ñ9Ô9×@Ò@ÅÄÑLÔLˆÝŒg�a�œÑ$Ô$ uÑ,ˆØ>•D”HÐ=Ñ>Ô>Ý-ñ/ô /ˆ	àˆ	�&Ô˜#˜tÑ$Ô$Ð$åŒv�d‰|Œ|ˆÝ
Œ
×Ò  TÑ*Ô*Ð*Ð*Ð*r   c                 ó˜  — t           j                             ddd¬¦  «                             t           j        ¦  «        }t          j        dt           j        ¦  «        dz  } t          j        d¦  «        t          ¦  «        } |d         ||¦  «         t          j	        |¦  «        }t           j
                             ||¦  «         d S rM  )r  r  r'  r(  r   rc  r   r*  Úatomic_min_double_oneindexrT  rc  rd  rJ  s        r   Útest_atomic_min_double_oneindexz/TestCudaAtomics.test_atomic_min_double_oneindex®  s§   € ÝŒy× Ò   C¨bÐ Ñ1Ô1×8Ò8½¼ÑDÔDˆÝŒg�a�œÑ$Ô$ sÑ*ˆØ<•D”HÐ;Ñ<Ô<Ý&ñ(ô (ˆ	àˆ	�%Ô˜˜dÑ#Ô#Ð#åŒv�d‰|Œ|ˆÝ
Œ
×Ò  TÑ*Ô*Ð*Ð*Ð*r   c                 ó�  —  t          j        d¦  «        |¦  «        }t          j                             ddd¬¦  «                             t          j        ¦  «        }t          j        dt          j        ¦  «        t          j        z   } |d         ||¦  «         t          j	         
                    |t          j        g¦  «         d S )NrH  r   rN  ©r3   r3   r#  r3   )r   r*  r  r  r'  r(  r   r+  Únanrc  rd  )r  rý   ro  r/  r  s        r   Ú _test_atomic_minmax_nan_locationz0TestCudaAtomics._test_atomic_minmax_nan_locationÃ  sš   € à>•D”HÐ=Ñ>Ô>¸tÑDÔDˆ	åŒy× Ò   C¨eÐ Ñ4Ô4×;Ò;½B¼JÑGÔGˆÝŒh�q�"œ*Ñ%Ô%­¬Ñ.ˆØˆ	�$Œ˜˜TÑ"Ô"Ð"Ý
Œ
×Ò ¥b¤f XÑ.Ô.Ð.Ð.Ð.r   c                 ó¢  —  t          j        d¦  «        |¦  «        }t          j                             ddd¬¦  «                             t          j        ¦  «        }|                     ¦   «         }t          j        dt          j        ¦  «        t          j	        z   } |d         ||¦  «         t          j
                             ||¦  «         d S )NrH  r   rN  r3   r#  rj  )r   r*  r  r  r'  r(  r   r)  r+  rk  rc  rd  )r  rý   ro  r  r3  r/  s         r   Ú_test_atomic_minmax_nan_valz+TestCudaAtomics._test_atomic_minmax_nan_valÌ  s£   € Ø>•D”HÐ=Ñ>Ô>¸tÑDÔDˆ	åŒi×Ò  3¨QÐÑ/Ô/×6Ò6µr´zÑBÔBˆØ�xŠx‰zŒzˆÝŒx˜¥¤
Ñ+Ô+­b¬fÑ4ˆØˆ	�$Œ˜˜TÑ"Ô"Ð"å
Œ
×Ò  TÑ*Ô*Ð*Ð*Ð*r   c                 ó:   — |                       t          ¦  «         d S r   )rl  rS  r6  s    r   Útest_atomic_min_nan_locationz,TestCudaAtomics.test_atomic_min_nan_locationÖ  ó   € Ø×-Ò-­jÑ9Ô9Ð9Ð9Ð9r   c                 ó:   — |                       t          ¦  «         d S r   )rl  r*  r6  s    r   Útest_atomic_max_nan_locationz,TestCudaAtomics.test_atomic_max_nan_locationÙ  rq  r   c                 ó:   — |                       t          ¦  «         d S r   )rn  rS  r6  s    r   Útest_atomic_min_nan_valz'TestCudaAtomics.test_atomic_min_nan_valÜ  ó   € Ø×(Ò(­Ñ4Ô4Ð4Ð4Ð4r   c                 ó:   — |                       t          ¦  «         d S r   )rn  r*  r6  s    r   Útest_atomic_max_nan_valz'TestCudaAtomics.test_atomic_max_nan_valß  rv  r   c                 ó–  — t           j                             ddd¬¦  «                             t           j        ¦  «        }t          j        dt           j        ¦  «        }d} t          j        |¦  «        t          ¦  «        } |d         ||¦  «         t          j	        |¦  «        }t           j
                             ||¦  «         d S ©Nr   rA   r#  r3   rO  r$  )r  r  r'  r(  r   r+  r   r*  Úatomic_max_double_sharedr+  rc  rd  ©r  r/  r  rs  ro  r3  s         r   Útest_atomic_max_double_sharedz-TestCudaAtomics.test_atomic_max_double_sharedâ  s    € ÝŒy× Ò   B¨RÐ Ñ0Ô0×7Ò7½¼
ÑCÔCˆÝŒh�q�"œ*Ñ%Ô%ˆØ,ˆØ!•D”H˜S‘M”MÕ":Ñ;Ô;ˆ	Øˆ	�%Ô˜˜dÑ#Ô#Ð#åŒv�d‰|Œ|ˆÝ
Œ
×Ò  TÑ*Ô*Ð*Ð*Ð*r   c                 óœ  — t           j                             ddd¬¦  «                             t           j        ¦  «        }t          j        dt           j        ¦  «        dz  }d} t          j        |¦  «        t          ¦  «        } |d         ||¦  «         t          j	        |¦  «        }t           j
                             ||¦  «         d S rz  )r  r  r'  r(  r   rc  r   r*  Úatomic_min_double_sharedrT  rc  rd  r|  s         r   Útest_atomic_min_double_sharedz-TestCudaAtomics.test_atomic_min_double_sharedì  s¥   € ÝŒy× Ò   B¨RÐ Ñ0Ô0×7Ò7½¼
ÑCÔCˆÝŒg�a�œÑ$Ô$ rÑ)ˆØ,ˆØ!•D”H˜S‘M”MÕ":Ñ;Ô;ˆ	Øˆ	�%Ô˜˜dÑ#Ô#Ð#åŒv�d‰|Œ|ˆÝ
Œ
×Ò  TÑ*Ô*Ð*Ð*Ð*r   r3   c                 óä  — |g|dz  z  |g|dz  z  z   }t           j                             |¦  «         t          j        ||¬¦  «        }|dk    rd|_        t          j        |¦  «        }t           j                             dd|j        ¬¦  «                             |j        ¦  «        }	||k    }
||k    }t          j        |¦  «        }|	|
         ||
<   |||<   | 	                    ¦   «         }t          j        |¦  «        }|dk    r |d         |||	|¦  «         n |d         |||	|¦  «         t           j                             ||¦  «         t           j                             ||¦  «         d S )	Nr  r%  )é
   éÿÿÿÿr3   r‚  r#  ©r‚  r‚  )r„  r„  )r  r  ÚshuffleÚasarrayr;   Ú
zeros_liker'  r(  r&  r)  r   r*  rc  Úassert_array_equal)r  ÚnÚfillÚunfillr&  Úcas_funcÚndimr  Úoutr"   Ú	fill_maskÚunfill_maskÚ
expect_resÚ
expect_outro  s                  r   Ú	check_caszTestCudaAtomics.check_casö  sg  € Øˆf˜˜Q™Ñ 6 (¨a°1©fÑ"5Ñ5ˆÝ
Œ	×Ò˜#ÑÔÐÝŒj˜ EÐ*Ñ*Ô*ˆØ�1Š9ˆ9Ø ˆCŒIÝŒm˜CÑ Ô ˆÝŒi×Ò  2¨C¬IÐÑ6Ô6×=Ò=¸c¼iÑHÔHˆà˜4’Kˆ	Ø˜V’mˆå”] 3Ñ'Ô'ˆ
Ø # I¤ˆ
�9ÑØ"(ˆ
�;Ñà—X’X‘Z”Zˆ
å”H˜XÑ&Ô&ˆ	Ø�1Š9ˆ9ØˆI�fÔ˜c 3¨¨TÑ2Ô2Ð2Ð2à)ˆIÐ(Ô)¨#¨s°C¸Ñ>Ô>Ð>å
Œ
×%Ò% j°#Ñ6Ô6Ð6Ý
Œ
×%Ò% j°#Ñ6Ô6Ð6Ð6Ð6r   c                 óX   — |                       dddt          j        t          ¬¦  «         d S ©Nr  é�ÿÿÿrƒ  ©r‰  rŠ  r‹  r&  rŒ  )r“  r  r°  r  r6  s    r   Útest_atomic_compare_and_swapz,TestCudaAtomics.test_atomic_compare_and_swap  ó4   € Ø�Š˜ 3¨r½¼Ý 7ð 	ñ 	9ô 	9ð 	9ð 	9ð 	9r   c                 óX   — |                       dddt          j        t          ¬¦  «         d S ©Nr  éÓÿÿÿrƒ  r—  )r“  r  rb  r  r6  s    r   Útest_atomic_compare_and_swap2z-TestCudaAtomics.test_atomic_compare_and_swap2  r™  r   c                 ó  — t           j                             ddt           j        ¬¦  «        }t           j                             ddt           j        ¬¦  «        }|                      d||t           j        t
          ¬¦  «         d S ©Nr  rž  r%  r3   é   r  r—  )r  r  r'  r   r“  r  ©r  ÚrfillÚrunfills      r   Útest_atomic_compare_and_swap3z-TestCudaAtomics.test_atomic_compare_and_swap3  óp   € Ý”	×!Ò! " cµ´Ð!Ñ;Ô;ˆÝ”)×#Ò# A rµ´Ð#Ñ;Ô;ˆØ�Š˜ 5°ÅÄ	Ý 7ð 	ñ 	9ô 	9ð 	9ð 	9ð 	9r   c                 ó  — t           j                             ddt           j        ¬¦  «        }t           j                             ddt           j        ¬¦  «        }|                      d||t           j        t
          ¬¦  «         d S rŸ  )r  r  r'  r   r“  r  r¡  s      r   Útest_atomic_compare_and_swap4z-TestCudaAtomics.test_atomic_compare_and_swap4  r¥  r   c                 óX   — |                       dddt          j        t          ¬¦  «         d S r•  )r“  r  r°  r  r6  s    r   Útest_atomic_cas_1dimz$TestCudaAtomics.test_atomic_cas_1dim%  ó4   € Ø�Š˜ 3¨r½¼Ý /ð 	ñ 	1ô 	1ð 	1ð 	1ð 	1r   c                 óZ   — |                       dddt          j        t          d¬¦  «         d S )Nr  r–  rƒ  r  ©r‰  rŠ  r‹  r&  rŒ  r�  )r“  r  r°  r  r6  s    r   Útest_atomic_cas_2dimz$TestCudaAtomics.test_atomic_cas_2dim)  ó6   € Ø�Š˜ 3¨r½¼Ý /°að 	ñ 	9ô 	9ð 	9ð 	9ð 	9r   c                 óX   — |                       dddt          j        t          ¬¦  «         d S r›  )r“  r  rb  r  r6  s    r   Útest_atomic_cas2_1dimz%TestCudaAtomics.test_atomic_cas2_1dim-  rª  r   c                 óZ   — |                       dddt          j        t          d¬¦  «         d S )Nr  rœ  rƒ  r  r¬  )r“  r  rb  r  r6  s    r   Útest_atomic_cas2_2dimz%TestCudaAtomics.test_atomic_cas2_2dim1  r®  r   c                 ó  — t           j                             ddt           j        ¬¦  «        }t           j                             ddt           j        ¬¦  «        }|                      d||t           j        t
          ¬¦  «         d S rŸ  )r  r  r'  r   r“  r  r¡  s      r   Útest_atomic_cas3_1dimz%TestCudaAtomics.test_atomic_cas3_1dim5  óp   € Ý”	×!Ò! " cµ´Ð!Ñ;Ô;ˆÝ”)×#Ò# A rµ´Ð#Ñ;Ô;ˆØ�Š˜ 5°ÅÄ	Ý /ð 	ñ 	1ô 	1ð 	1ð 	1ð 	1r   c                 ó
  — t           j                             ddt           j        ¬¦  «        }t           j                             ddt           j        ¬¦  «        }|                      d||t           j        t
          d¬¦  «         d S ©	Nr  rž  r%  r3   r   r  r  r¬  )r  r  r'  r   r“  r  r¡  s      r   Útest_atomic_cas3_2dimz%TestCudaAtomics.test_atomic_cas3_2dim;  ór   € Ý”	×!Ò! " cµ´Ð!Ñ;Ô;ˆÝ”)×#Ò# A rµ´Ð#Ñ;Ô;ˆØ�Š˜ 5°ÅÄ	Ý /°að 	ñ 	9ô 	9ð 	9ð 	9ð 	9r   c                 ó  — t           j                             ddt           j        ¬¦  «        }t           j                             ddt           j        ¬¦  «        }|                      d||t           j        t
          ¬¦  «         d S rŸ  )r  r  r'  r   r“  r  r¡  s      r   Útest_atomic_cas4_1dimz%TestCudaAtomics.test_atomic_cas4_1dimA  rµ  r   c                 ó
  — t           j                             ddt           j        ¬¦  «        }t           j                             ddt           j        ¬¦  «        }|                      d||t           j        t
          d¬¦  «         d S r·  )r  r  r'  r   r“  r  r¡  s      r   Útest_atomic_cas4_2dimz%TestCudaAtomics.test_atomic_cas4_2dimG  r¹  r   c                 ó0  — t          j        dt           j        ¬¦  «        }||d<    |d         |¦  «         t          j        |¦  «        r/|                      t          j        |d         ¦  «        ¦  «         d S |                      |d         |¦  «         d S )Nr  r%  r   rj  r3   )r  r+  r   Úisnanr-  ÚassertEqual©r  rZ  Úinitialr   s       r   Ú_test_atomic_returns_oldz(TestCudaAtomics._test_atomic_returns_oldR  sˆ   € ÝŒH�Q�bœjÐ)Ñ)Ô)ˆØˆˆ!‰ØˆˆtŒ�Q‰ŒˆÝŒ8�GÑÔð 	,Ø�OŠO�BœH Q q¤T™NœNÑ+Ô+Ð+Ð+Ð+à×Ò˜Q˜qœT 7Ñ+Ô+Ð+Ð+Ð+r   c                 ó\   — t           j        d„ ¦   «         }|                      |d¦  «         d S )Nc                 óN   — t           j                             | dd¦  «        | d<   d S r2   )r   rC   rD   ©r   s    r   rZ  z;TestCudaAtomics.test_atomic_add_returns_old.<locals>.kernel\  ó!   € å”;—?’? 1 a¨Ñ+Ô+ˆAˆa‰DˆDˆDr   r‚  ©r   r*  rÃ  ©r  rZ  s     r   Útest_atomic_add_returns_oldz+TestCudaAtomics.test_atomic_add_returns_old[  ó;   € Ý	Œð	,ð 	,ñ 
Œð	,ð 	×%Ò% f¨bÑ1Ô1Ð1Ð1Ð1r   c                 ó\   — t           j        d„ ¦   «         }|                      |d¦  «         d S )Nc                 óN   — t           j                             | dd¦  «        | d<   d S r2   ©r   rC   r+  rÆ  s    r   rZ  zBTestCudaAtomics.test_atomic_max_returns_no_replace.<locals>.kernelc  rÇ  r   r‚  rÈ  rÉ  s     r   Ú"test_atomic_max_returns_no_replacez2TestCudaAtomics.test_atomic_max_returns_no_replaceb  rË  r   c                 ó\   — t           j        d„ ¦   «         }|                      |d¦  «         d S )Nc                 óN   — t           j                             | dd¦  «        | d<   d S ©Nr   r‚  r3   rÎ  rÆ  s    r   rZ  zCTestCudaAtomics.test_atomic_max_returns_old_replace.<locals>.kernelj  ó!   € å”;—?’? 1 a¨Ñ,Ô,ˆAˆa‰DˆDˆDr   r3   rÈ  rÉ  s     r   Ú#test_atomic_max_returns_old_replacez3TestCudaAtomics.test_atomic_max_returns_old_replacei  s;   € Ý	Œð	-ð 	-ñ 
Œð	-ð 	×%Ò% f¨aÑ0Ô0Ð0Ð0Ð0r   c                 óp   — t           j        d„ ¦   «         }|                      |t          j        ¦  «         d S )Nc                 óN   — t           j                             | dd¦  «        | d<   d S r2   rÎ  rÆ  s    r   rZ  zHTestCudaAtomics.test_atomic_max_returns_old_nan_in_array.<locals>.kernelq  rÇ  r   ©r   r*  rÃ  r  rk  rÉ  s     r   Ú(test_atomic_max_returns_old_nan_in_arrayz8TestCudaAtomics.test_atomic_max_returns_old_nan_in_arrayp  s=   € Ý	Œð	,ð 	,ñ 
Œð	,ð 	×%Ò% f­b¬fÑ5Ô5Ð5Ð5Ð5r   c                 ó\   — t           j        d„ ¦   «         }|                      |d¦  «         d S )Nc                 ób   — t           j                             | dt          j        ¦  «        | d<   d S r2   )r   rC   r+  r  rk  rÆ  s    r   rZ  zCTestCudaAtomics.test_atomic_max_returns_old_nan_val.<locals>.kernelx  ó#   € å”;—?’? 1 a­¬Ñ0Ô0ˆAˆa‰DˆDˆDr   r‚  rÈ  rÉ  s     r   Ú#test_atomic_max_returns_old_nan_valz3TestCudaAtomics.test_atomic_max_returns_old_nan_valw  ó;   € Ý	Œð	1ð 	1ñ 
Œð	1ð 	×%Ò% f¨bÑ1Ô1Ð1Ð1Ð1r   c                 ó\   — t           j        d„ ¦   «         }|                      |d¦  «         d S )Nc                 óN   — t           j                             | dd¦  «        | d<   d S ©Nr   é   r3   ©r   rC   rT  rÆ  s    r   rZ  zFTestCudaAtomics.test_atomic_min_returns_old_no_replace.<locals>.kernel  rÓ  r   r‚  rÈ  rÉ  s     r   Ú&test_atomic_min_returns_old_no_replacez6TestCudaAtomics.test_atomic_min_returns_old_no_replace~  ó;   € Ý	Œð	-ð 	-ñ 
Œð	-ð 	×%Ò% f¨bÑ1Ô1Ð1Ð1Ð1r   c                 ó\   — t           j        d„ ¦   «         }|                      |d¦  «         d S )Nc                 óN   — t           j                             | dd¦  «        | d<   d S rÒ  râ  rÆ  s    r   rZ  zCTestCudaAtomics.test_atomic_min_returns_old_replace.<locals>.kernel†  rÓ  r   rá  rÈ  rÉ  s     r   Ú#test_atomic_min_returns_old_replacez3TestCudaAtomics.test_atomic_min_returns_old_replace…  rä  r   c                 óp   — t           j        d„ ¦   «         }|                      |t          j        ¦  «         d S )Nc                 óN   — t           j                             | dd¦  «        | d<   d S rà  râ  rÆ  s    r   rZ  zHTestCudaAtomics.test_atomic_min_returns_old_nan_in_array.<locals>.kernel�  rÓ  r   r×  rÉ  s     r   Ú(test_atomic_min_returns_old_nan_in_arrayz8TestCudaAtomics.test_atomic_min_returns_old_nan_in_arrayŒ  s=   € Ý	Œð	-ð 	-ñ 
Œð	-ð 	×%Ò% f­b¬fÑ5Ô5Ð5Ð5Ð5r   c                 ó\   — t           j        d„ ¦   «         }|                      |d¦  «         d S )Nc                 ób   — t           j                             | dt          j        ¦  «        | d<   d S r2   )r   rC   rT  r  rk  rÆ  s    r   rZ  zCTestCudaAtomics.test_atomic_min_returns_old_nan_val.<locals>.kernel”  rÛ  r   rá  rÈ  rÉ  s     r   Ú#test_atomic_min_returns_old_nan_valz3TestCudaAtomics.test_atomic_min_returns_old_nan_val“  rÝ  r   c                 ót  — t           j                             ||d¬¦  «                             |¦  «        }||dd d…<   t          j        d|j        ¬¦  «        }t          j        t          ¦  «        } |d         ||¦  «         t          j	        |¦  «        }t           j
                             ||¦  «         d S )Nr)  r#  r3   r  r%  )r  r  r'  r(  r+  r&  r   r*  Úatomic_nanmaxÚnanmaxrc  rd  ©	r  r&  r-  r.  Úinit_valr/  r  ro  r3  s	            r   Úcheck_atomic_nanmaxz#TestCudaAtomics.check_atomic_nanmax�  s¢   € ÝŒy× Ò   R¨hÐ Ñ7Ô7×>Ò>¸uÑEÔEˆØˆˆQˆT�ˆT‰
ÝŒh�q ¤
Ð+Ñ+Ô+ˆÝ”H�]Ñ+Ô+ˆ	Øˆ	�&Ô˜#˜tÑ$Ô$Ð$ÝŒy˜‰ŒˆÝ
Œ
×Ò  TÑ*Ô*Ð*Ð*Ð*r   c                 óL   — |                       t          j        ddd¬¦  «         d S ©Nr3  r4  r   ©r&  r-  r.  rò  )ró  r  r°  r6  s    r   Útest_atomic_nanmax_int32z(TestCudaAtomics.test_atomic_nanmax_int32¦  ó4   € Ø× Ò ¥r¤x°F¸uØ*+ð 	!ñ 	-ô 	-ð 	-ð 	-ð 	-r   c                 óL   — |                       t          j        ddd¬¦  «         d S ©Nr   r4  rö  )ró  r  r   r6  s    r   Útest_atomic_nanmax_uint32z)TestCudaAtomics.test_atomic_nanmax_uint32ª  ó4   € Ø× Ò ¥r¤y°Q¸5Ø*+ð 	!ñ 	-ô 	-ð 	-ð 	-ð 	-r   c                 óL   — |                       t          j        ddd¬¦  «         d S rõ  )ró  r  rb  r6  s    r   Útest_atomic_nanmax_int64z(TestCudaAtomics.test_atomic_nanmax_int64®  rø  r   c                 óL   — |                       t          j        ddd¬¦  «         d S rú  )ró  r  r   r6  s    r   Útest_atomic_nanmax_uint64z)TestCudaAtomics.test_atomic_nanmax_uint64²  rü  r   c                 ó`   — |                       t          j        ddt          j        ¬¦  «         d S ©Nr3  r4  rö  )ró  r  r   rk  r6  s    r   Útest_atomic_nanmax_float32z*TestCudaAtomics.test_atomic_nanmax_float32¶  ó6   € Ø× Ò ¥r¤z°fÀÝ*,¬&ð 	!ñ 	2ô 	2ð 	2ð 	2ð 	2r   c                 ó`   — |                       t          j        ddt          j        ¬¦  «         d S r  )ró  r  r   rk  r6  s    r   Útest_atomic_nanmax_doublez)TestCudaAtomics.test_atomic_nanmax_doubleº  r  r   c                 ó´  — t           j                             ddd¬¦  «                             t           j        ¦  «        }t           j        |dd d…<   t          j        dg|j        ¬¦  «        }d} t          j	        |¦  «        t          ¦  «        } |d         ||¦  «         t          j        |¦  «        }t           j                             ||¦  «         d S ©	Nr   rA   r#  r3   r  r%  rO  r$  )r  r  r'  r(  r   rk  r    r&  r   r*  Úatomic_nanmax_double_sharedrð  rc  rd  r|  s         r   Ú test_atomic_nanmax_double_sharedz0TestCudaAtomics.test_atomic_nanmax_double_shared¾  s¶   € ÝŒy× Ò   B¨RÐ Ñ0Ô0×7Ò7½¼
ÑCÔCˆÝ”VˆˆQˆT�ˆT‰
ÝŒh˜�s $¤*Ð-Ñ-Ô-ˆØ,ˆØ!•D”H˜S‘M”MÕ"=Ñ>Ô>ˆ	Øˆ	�%Ô˜˜dÑ#Ô#Ð#åŒy˜‰ŒˆÝ
Œ
×Ò  TÑ*Ô*Ð*Ð*Ð*r   c                 ó¶  — t           j                             ddd¬¦  «                             t           j        ¦  «        }t           j        |dd d…<   t          j        dt           j        ¦  «        } t          j        d¦  «        t          ¦  «        } |d         ||¦  «         t          j
        |¦  «        }t           j                             ||¦  «         d S ©	Nr   rN  rA   r#  r3   r  rO  r$  )r  r  r'  r(  r   rk  r+  r   r*  rP  rð  rc  rd  rJ  s        r   Ú"test_atomic_nanmax_double_oneindexz2TestCudaAtomics.test_atomic_nanmax_double_oneindexÉ  s³   € ÝŒy× Ò   C¨bÐ Ñ1Ô1×8Ò8½¼ÑDÔDˆÝ”VˆˆQˆT�ˆT‰
ÝŒh�q�"œ*Ñ%Ô%ˆØ<•D”HÐ;Ñ<Ô<Ý&ñ(ô (ˆ	àˆ	�%Ô˜˜dÑ#Ô#Ð#åŒy˜‰ŒˆÝ
Œ
×Ò  TÑ*Ô*Ð*Ð*Ð*r   c                 óv  — t           j                             ||d¬¦  «                             |¦  «        }||dd d…<   t          j        dg|j        ¬¦  «        }t          j        t          ¦  «        } |d         ||¦  «         t          j	        |¦  «        }t           j
                             ||¦  «         d S )Nr)  r#  r3   r  r4  r%  )r  r  r'  r(  r    r&  r   r*  Úatomic_nanminÚnanminrc  rd  rñ  s	            r   Úcheck_atomic_nanminz#TestCudaAtomics.check_atomic_nanminÕ  s¤   € ÝŒy× Ò   R¨hÐ Ñ7Ô7×>Ò>¸uÑEÔEˆØˆˆQˆT�ˆT‰
ÝŒh˜�w d¤jÐ1Ñ1Ô1ˆÝ”H�]Ñ+Ô+ˆ	Øˆ	�&Ô˜#˜tÑ$Ô$Ð$åŒy˜‰ŒˆÝ
Œ
×Ò  TÑ*Ô*Ð*Ð*Ð*r   c                 óL   — |                       t          j        ddd¬¦  «         d S rõ  )r  r  r°  r6  s    r   Útest_atomic_nanmin_int32z(TestCudaAtomics.test_atomic_nanmin_int32ß  rø  r   c                 óL   — |                       t          j        ddd¬¦  «         d S rú  )r  r  r   r6  s    r   Útest_atomic_nanmin_uint32z)TestCudaAtomics.test_atomic_nanmin_uint32ã  rü  r   c                 óL   — |                       t          j        ddd¬¦  «         d S rõ  )r  r  rb  r6  s    r   Útest_atomic_nanmin_int64z(TestCudaAtomics.test_atomic_nanmin_int64ç  rø  r   c                 óL   — |                       t          j        ddd¬¦  «         d S rú  )r  r  r   r6  s    r   Útest_atomic_nanmin_uint64z)TestCudaAtomics.test_atomic_nanmin_uint64ë  rü  r   c                 ó`   — |                       t          j        ddt          j        ¬¦  «         d S r  )r  r  r   rk  r6  s    r   Útest_atomic_nanmin_floatz(TestCudaAtomics.test_atomic_nanmin_floatï  r  r   c                 ó`   — |                       t          j        ddt          j        ¬¦  «         d S r  )r  r  r   rk  r6  s    r   Útest_atomic_nanmin_doublez)TestCudaAtomics.test_atomic_nanmin_doubleó  r  r   c                 ó´  — t           j                             ddd¬¦  «                             t           j        ¦  «        }t           j        |dd d…<   t          j        dg|j        ¬¦  «        }d} t          j	        |¦  «        t          ¦  «        } |d         ||¦  «         t          j        |¦  «        }t           j                             ||¦  «         d S r  )r  r  r'  r(  r   rk  r    r&  r   r*  Úatomic_nanmin_double_sharedr  rc  rd  r|  s         r   Ú test_atomic_nanmin_double_sharedz0TestCudaAtomics.test_atomic_nanmin_double_shared÷  s¶   € ÝŒy× Ò   B¨RÐ Ñ0Ô0×7Ò7½¼
ÑCÔCˆÝ”VˆˆQˆT�ˆT‰
ÝŒh˜�t 4¤:Ð.Ñ.Ô.ˆØ,ˆØ!•D”H˜S‘M”MÕ"=Ñ>Ô>ˆ	Øˆ	�%Ô˜˜dÑ#Ô#Ð#åŒy˜‰ŒˆÝ
Œ
×Ò  TÑ*Ô*Ð*Ð*Ð*r   c                 ó¸  — t           j                             ddd¬¦  «                             t           j        ¦  «        }t           j        |dd d…<   t          j        dgt           j        ¦  «        } t          j        d¦  «        t          ¦  «        } |d         ||¦  «         t          j
        |¦  «        }t           j                             ||¦  «         d S r  )r  r  r'  r(  r   rk  r    r   r*  rg  r  rc  rd  rJ  s        r   Ú"test_atomic_nanmin_double_oneindexz2TestCudaAtomics.test_atomic_nanmin_double_oneindex  sµ   € ÝŒy× Ò   C¨bÐ Ñ1Ô1×8Ò8½¼ÑDÔDˆÝ”VˆˆQˆT�ˆT‰
ÝŒh˜�u�bœjÑ)Ô)ˆØ<•D”HÐ;Ñ<Ô<Ý&ñ(ô (ˆ	àˆ	�%Ô˜˜dÑ#Ô#Ð#åŒy˜‰ŒˆÝ
Œ
×Ò  TÑ*Ô*Ð*Ð*Ð*r   c                 ó¨  — t          j        dt           j        ¬¦  «        }||d<   t           j        |d<    |d         |¦  «         t          j        |¦  «        r\|                      t          j        |d         ¦  «        ¦  «         |                      t          j        |d         ¦  «        ¦  «         d S |                      |d         |¦  «         d S )Nr  r%  r   r3   rj  )r  r+  r   rk  r¿  ÚassertFalser-  rÀ  rÁ  s       r   Ú_test_atomic_nan_returns_oldz,TestCudaAtomics._test_atomic_nan_returns_old  s³   € ÝŒH�Q�bœjÐ)Ñ)Ô)ˆØˆˆ!‰ÝŒvˆˆ!‰ØˆˆtŒ�Q‰ŒˆÝŒ8�GÑÔð 	,Ø×Ò�RœX a¨¤d™^œ^Ñ,Ô,Ð,Ø�OŠO�BœH Q q¤T™NœNÑ+Ô+Ð+Ð+Ð+à×Ò˜Q˜qœT 7Ñ+Ô+Ð+Ð+Ð+r   c                 ó\   — t           j        d„ ¦   «         }|                      |d¦  «         d S )Nc                 óN   — t           j                             | dd¦  «        | d<   d S r2   ©r   rC   rð  rÆ  s    r   rZ  zITestCudaAtomics.test_atomic_nanmax_returns_old_no_replace.<locals>.kernel  ó#   € å”;×%Ò% a¨¨AÑ.Ô.ˆAˆa‰DˆDˆDr   r‚  ©r   r*  r%  rÉ  s     r   Ú)test_atomic_nanmax_returns_old_no_replacez9TestCudaAtomics.test_atomic_nanmax_returns_old_no_replace  s;   € Ý	Œð	/ð 	/ñ 
Œð	/ð 	×)Ò)¨&°"Ñ5Ô5Ð5Ð5Ð5r   c                 ó\   — t           j        d„ ¦   «         }|                      |d¦  «         d S )Nc                 óN   — t           j                             | dd¦  «        | d<   d S rÒ  r(  rÆ  s    r   rZ  zFTestCudaAtomics.test_atomic_nanmax_returns_old_replace.<locals>.kernel"  ó#   € å”;×%Ò% a¨¨BÑ/Ô/ˆAˆa‰DˆDˆDr   r3   r*  rÉ  s     r   Ú&test_atomic_nanmax_returns_old_replacez6TestCudaAtomics.test_atomic_nanmax_returns_old_replace!  s;   € Ý	Œð	0ð 	0ñ 
Œð	0ð 	×)Ò)¨&°!Ñ4Ô4Ð4Ð4Ð4r   c                 óp   — t           j        d„ ¦   «         }|                      |t          j        ¦  «         d S )Nc                 óN   — t           j                             | dd¦  «        | d<   d S r2   r(  rÆ  s    r   rZ  zKTestCudaAtomics.test_atomic_nanmax_returns_old_nan_in_array.<locals>.kernel)  r)  r   ©r   r*  r%  r  rk  rÉ  s     r   Ú+test_atomic_nanmax_returns_old_nan_in_arrayz;TestCudaAtomics.test_atomic_nanmax_returns_old_nan_in_array(  s=   € Ý	Œð	/ð 	/ñ 
Œð	/ð 	×)Ò)¨&µ"´&Ñ9Ô9Ð9Ð9Ð9r   c                 ó\   — t           j        d„ ¦   «         }|                      |d¦  «         d S )Nc                 ób   — t           j                             | dt          j        ¦  «        | d<   d S r2   )r   rC   rð  r  rk  rÆ  s    r   rZ  zFTestCudaAtomics.test_atomic_nanmax_returns_old_nan_val.<locals>.kernel0  ó%   € å”;×%Ò% a¨­B¬FÑ3Ô3ˆAˆa‰DˆDˆDr   r‚  r*  rÉ  s     r   Ú&test_atomic_nanmax_returns_old_nan_valz6TestCudaAtomics.test_atomic_nanmax_returns_old_nan_val/  ó;   € Ý	Œð	4ð 	4ñ 
Œð	4ð 	×)Ò)¨&°"Ñ5Ô5Ð5Ð5Ð5r   c                 ó\   — t           j        d„ ¦   «         }|                      |d¦  «         d S )Nc                 óN   — t           j                             | dd¦  «        | d<   d S rà  ©r   rC   r  rÆ  s    r   rZ  zITestCudaAtomics.test_atomic_nanmin_returns_old_no_replace.<locals>.kernel7  r.  r   r‚  r*  rÉ  s     r   Ú)test_atomic_nanmin_returns_old_no_replacez9TestCudaAtomics.test_atomic_nanmin_returns_old_no_replace6  ó;   € Ý	Œð	0ð 	0ñ 
Œð	0ð 	×)Ò)¨&°"Ñ5Ô5Ð5Ð5Ð5r   c                 ó\   — t           j        d„ ¦   «         }|                      |d¦  «         d S )Nc                 óN   — t           j                             | dd¦  «        | d<   d S rÒ  r;  rÆ  s    r   rZ  zFTestCudaAtomics.test_atomic_nanmin_returns_old_replace.<locals>.kernel>  r.  r   rá  r*  rÉ  s     r   Ú&test_atomic_nanmin_returns_old_replacez6TestCudaAtomics.test_atomic_nanmin_returns_old_replace=  r=  r   c                 óp   — t           j        d„ ¦   «         }|                      |t          j        ¦  «         d S )Nc                 óN   — t           j                             | dd¦  «        | d<   d S rà  r;  rÆ  s    r   rZ  zKTestCudaAtomics.test_atomic_nanmin_returns_old_nan_in_array.<locals>.kernelE  r.  r   r2  rÉ  s     r   Ú+test_atomic_nanmin_returns_old_nan_in_arrayz;TestCudaAtomics.test_atomic_nanmin_returns_old_nan_in_arrayD  s=   € Ý	Œð	0ð 	0ñ 
Œð	0ð 	×)Ò)¨&µ"´&Ñ9Ô9Ð9Ð9Ð9r   c                 ó\   — t           j        d„ ¦   «         }|                      |d¦  «         d S )Nc                 ób   — t           j                             | dt          j        ¦  «        | d<   d S r2   )r   rC   r  r  rk  rÆ  s    r   rZ  zFTestCudaAtomics.test_atomic_nanmin_returns_old_nan_val.<locals>.kernelL  r6  r   rá  r*  rÉ  s     r   Ú&test_atomic_nanmin_returns_old_nan_valz6TestCudaAtomics.test_atomic_nanmin_returns_old_nan_valK  r8  r   )T)r3   )—Ú__name__Ú
__module__Ú__qualname__r  r5  r<  r@  rF  rJ  rL  r]  rg  rl  rp  ru  rx  rz  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  r  r  r  r  r  r  r  r"  r%  r'  r0  r7  r;  r>  r@  rB  rE  rK  rQ  rU  rW  rY  r[  r]  r_  ra  re  rh  rl  rn  rp  rs  ru  rx  r}  r€  r“  r˜  r�  r¤  r§  r©  r­  r°  r²  r´  r¸  r»  r½  rÃ  rÊ  rÏ  rÔ  rØ  rÜ  rã  rç  rê  rí  ró  r÷  rû  rþ  r   r  r  r
  r  r  r  r  r  r  r  r  r   r"  r%  r+  r/  r3  r7  r<  r@  rC  rF  Ú__classcell__)r   s   @r   r  r  Õ  s-	  ø€ € € € € ðð ð ð ð ð2ð 2ð 2ð$6ð 6ð 6ð1ð 1ð 1ð2ð 2ð 2ð$6ð 6ð 6ð1ð 1ð 1ð3ð 3ð 3ð 3ð22ð 2ð 2ð(7ð 7ð 7ð 4ð 4ð 4ðGð Gð Gð*Gð Gð Gð"Bð Bð Bð
-ð 
-ð 
-ð1ð 1ð 1ð1ð 1ð 1ð
-ð 
-ð 
-ð1ð 1ð 1ð1ð 1ð 1ð
+ð 
+ð 
+ð/ð /ð /ð/ð /ð /ð+ð +ð +ð/ð /ð /ð/ð /ð /ð-ð -ð -ð:ð :ð :ð:ð :ð :ð+ð +ð +ð8ð 8ð 8ð-ð -ð -ð:ð :ð :ð:ð :ð :ð+ð +ð +ð8ð 8ð 8ð-ð -ð -ð:ð :ð :ð:ð :ð :ð+ð +ð +ð8ð 8ð 8ð%ð %ð %ðð ð ð
Lð Lð LðLð Lð LðLð Lð LðMð Mð Mð
Mð Mð Mð
Gð Gð Gð
Gð Gð Gð
Dð Dð Dð
1ð 1ð 1ð1ð 1ð 1ðLð Lð Lð
Lð Lð Lð
Cð Cð CðCð Cð CðCð Cð CðMð Mð Mð
Mð Mð Mð
Gð Gð Gð
Gð Gð Gð
Dð Dð Dð
1ð 1ð 1ð1ð 1ð 1ðLð Lð Lð
Lð Lð Lð
1ð 1ð 1ð1ð 1ð 1ð1ð 1ð 1ð1ð 1ð 1ð+ð +ð +ðCð Cð Cð?ð ?ð ?ðCð Cð Cð?ð ?ð ?ðEð Eð EðEð Eð Eð+ð +ð +ð+ð +ð +ð+ð +ð +ðCð Cð Cð?ð ?ð ?ðCð Cð Cð?ð ?ð ?ðEð Eð EðEð Eð Eð+ð +ð +ð+ð +ð +ð*/ð /ð /ð+ð +ð +ð:ð :ð :ð:ð :ð :ð5ð 5ð 5ð5ð 5ð 5ð+ð +ð +ð+ð +ð +ð7ð 7ð 7ð 7ð69ð 9ð 9ð9ð 9ð 9ð9ð 9ð 9ð9ð 9ð 9ð1ð 1ð 1ð9ð 9ð 9ð1ð 1ð 1ð9ð 9ð 9ð1ð 1ð 1ð9ð 9ð 9ð1ð 1ð 1ð9ð 9ð 9ð,ð ,ð ,ð2ð 2ð 2ð2ð 2ð 2ð1ð 1ð 1ð6ð 6ð 6ð2ð 2ð 2ð2ð 2ð 2ð2ð 2ð 2ð6ð 6ð 6ð2ð 2ð 2ð+ð +ð +ð-ð -ð -ð-ð -ð -ð-ð -ð -ð-ð -ð -ð2ð 2ð 2ð2ð 2ð 2ð	+ð 	+ð 	+ð	+ð 	+ð 	+ð+ð +ð +ð-ð -ð -ð-ð -ð -ð-ð -ð -ð-ð -ð -ð2ð 2ð 2ð2ð 2ð 2ð	+ð 	+ð 	+ð	+ð 	+ð 	+ð	,ð 	,ð 	,ð6ð 6ð 6ð5ð 5ð 5ð:ð :ð :ð6ð 6ð 6ð6ð 6ð 6ð6ð 6ð 6ð:ð :ð :ð6ð 6ð 6ð 6ð 6ð 6ð 6r   r  Ú__main__)qÚnumpyr  Útextwrapr   Únumbar   r   r   r   r   Únumba.cuda.testingr	   r
   r   Ú
numba.corer   r*  r   r   r   r.   r0   r9   r<   r>   rF   rI   rP   rS   rU   r\   r`   rd   rh   rk   rq   rt   rx   rz   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ò   rõ   r÷   rù   rû   r  r*  rI  rP  r{  rS  rd  rg  r  rï  Ú$atomic_nanmax_double_normalizedindexÚatomic_nanmax_double_oneindexr	  r  Ú$atomic_nanmin_double_normalizedindexÚatomic_nanmin_double_oneindexr  r  r  r  r  rG  Úmainr   r   r   ú<module>rV     sv  ðØ Ð Ð Ð Ø Ð Ð Ð Ð Ð à 8Ð 8Ð 8Ð 8Ð 8Ð 8Ð 8Ð 8Ð 8Ð 8Ð 8Ð 8Ð 8Ð 8Ø DÐ DÐ DÐ DÐ DÐ DÐ DÐ DÐ DÐ DØ Ð Ð Ð Ð Ð ð 
€„�ÐÑÔðð ñ Ôðð 
€„�ÐÑÔðð ñ Ôðð 
€„�ÐÑÔðð ñ Ôðð 
€„�ÐÑÔðð ñ Ôðð 
€„�ÐÑÔð	ð 	ñ Ôð	ð 
€„�ÐÑÔðð ñ Ôðð 
€„�ÐÑÔðð ñ Ôðð 
€„�ÐÑÔðð ñ ÔððKð Kð Kð
Jð Jð Jð
Hð Hð Hð
Gð Gð Gð
Mð Mð Mð
Oð Oð Oð
Nð Nð Nð
Hð Hð Hð
Gð Gð Gð
Mð Mð Mð
Ið Ið IðHð Hð HðPð Pð PðOð Oð Oð%ð %ð %ð
Mð Mð Mð
Lð Lð Lð
Hð Hð Hð
Gð Gð Gð
Mð Mð Mð
Kð Kð Kð
Hð Hð Hð
Mð Mð Mð
Oð Oð Oð
Hð Hð Hð
Mð Mð Mð
Mð Mð Mð
Hð Hð Hð
Mð Mð Mð
Ið Ið Ið%ð %ð %ð
Mð Mð Mð
Lð Lð Lð
Ið Ið Ið
Nð Nð Nð
Jð Jð Jð7ð 7ð 7ð
Kð Kð Kð
Hð Hð Hð
Mð Mð Mð
Ið Ið Ið7ð 7ð 7ð
Kð Kð Kð
Hð Hð Hð
Mð Mð Mð
Ið Ið Ið7ð 7ð 7ð
Bð Bð Bð
Dð Dð Dð
Hð Hð Hð
Hð Hð Hð
Mð Mð Mð
Ið Ið Ið7ð 7ð 7ð
Bð Bð Bð
Dð Dð Dð
Hð Hð Hð
Hð Hð Hð
Mð Mð Mð
Ið Ið Ið7ð 7ð 7ð
Cð Cð Cð
Ið Ið Ið
Ið Ið Ið
Jð Jð Jð!Fð !Fð !FðJ 6Ð5Ð6GÑHÔHñ€Ð.Ð0JØà5Ð5Ð6GÑHÔHñ€Ð.Ð0JØð ÐÐ1Ñ2Ô2ñ=€Ð4ØÐ ;ð ÐÐ1Ñ2Ô2ñ=€Ð4ØÐ ;ðOð Oð OðAð Að AðAð Að Að{6ð {6ð {6ð {6ð {6�lñ {6ô {6ð {6ð|# ˆzÒÐØ€H„M�O„O€O€O€Oð Ðr   