§
    1@Äj¬  ã                   óà   — d dl Zd dlmZmZmZ d dlmZmZm	Z	 d dl
mZ d„ Zd„ Zd„ Zd„ Zd	„ Zd
„ Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Z G d„ de	¦  «        Zedk    r ej        ¦   «          dS dS )é    N)ÚcudaÚint32Úfloat32)Úskip_on_cudasimÚunittestÚCUDATestCase)ÚENABLE_CUDASIMc                 ó^   — t          j        d¦  «        }t          j        ¦   «          || |<   d S ©Né   )r   ÚgridÚsyncthreads©ÚaryÚis     úRC:\Projects\Historical\venv\Lib\site-packages\numba/cuda/tests/cudapy/test_sync.pyÚuseless_syncthreadsr      s+   € ÝŒ	�!‰Œ€AÝÔÑÔÐØ€Cˆ�F€F€Fó    c                 ó^   — t          j        d¦  «        }t          j        ¦   «          || |<   d S r   ©r   r   Úsyncwarpr   s     r   Úuseless_syncwarpr      s'   € ÝŒ	�!‰Œ€AÝ„M�O„O€OØ€Cˆ�F€F€Fr   c                 ó`   — t          j        d¦  «        }t          j        d¦  «         || |<   d S )Nr   éÿÿ  r   r   s     r   Úuseless_syncwarp_with_maskr      s,   € ÝŒ	�!‰Œ€AÝ„M�&ÑÔÐØ€Cˆ�F€F€Fr   c                 óh  — t           j                             dt          ¦  «        }t          j        d¦  «        }|||<   t          j        ¦   «          |dk     r+||         ||dz            z   ||<   t          j        d¦  «         |dk     r+||         ||dz            z   ||<   t          j        d¦  «         |dk     r+||         ||dz            z   ||<   t          j        d¦  «         |d	k     r+||         ||d	z            z   ||<   t          j        d
¦  «         |dk    r|d         |d         z   | d<   d S d S )Né    r   é   r   é   éÿ   é   é   é   é   r   )r   ÚsharedÚarrayr   r   r   )ÚresÚsmr   s      r   Úcoop_syncwarpr)      s/  € Ý	Œ×	Ò	˜2�uÑ	%Ô	%€BÝŒ	�!‰Œ€Aà€B€q�EÝ„M�O„O€Oàˆ2‚v€vØ�1”˜˜1˜r™6œ
Ñ"ˆˆ1‰ÝŒ�fÑÔÐàˆ1‚u€uØ�1”˜˜1˜q™5œ	Ñ!ˆˆ1‰ÝŒ�dÑÔÐàˆ1‚u€uØ�1”˜˜1˜q™5œ	Ñ!ˆˆ1‰ÝŒ�cÑÔÐàˆ1‚u€uØ�1”˜˜1˜q™5œ	Ñ!ˆˆ1‰ÝŒ�cÑÔÐàˆA‚v€vØ�A”˜˜Aœ‘ˆˆA‰ˆˆð €vr   c                 óò   — d}t           j                             |t          ¦  «        }t          j        d¦  «        }|dk    rt          |¦  «        D ]}|||<   Œt          j        ¦   «          ||         | |<   d S )Néd   r   r   )r   r%   r&   r   r   Úranger   )r   ÚNr(   r   Újs        r   Úsimple_smemr/   4   st   € Ø€AÝ	Œ×	Ò	˜1�eÑ	$Ô	$€BÝŒ	�!‰Œ€AØˆA‚v€vÝ�q‘”ð 	ð 	ˆAØˆBˆq‰EˆEÝÔÑÔÐØ�ŒU€Cˆ�F€F€Fr   c                 óâ   — t          j        d¦  «        \  }}t           j                             dt          ¦  «        }|dz   |dz   z  |||f<   t          j        ¦   «          |||f         | ||f<   d S )Nr#   ©é
   é   r   ©r   r   r%   r&   r   r   )r   r   r.   r(   s       r   Úcoop_smem2dr5   ?   sl   € ÝŒ9�Q‰<Œ<�D€A€qÝ	Œ×	Ò	˜8¥WÑ	-Ô	-€BØ�A‘˜!˜a™%Ñ €B€qˆ!€t�HÝÔÑÔÐØ�1�a�4”€Cˆˆ1ˆ�I€I€Ir   c                 óÄ   — t          j        d¦  «        }t           j                             dt          ¦  «        }|dz  ||<   t          j        ¦   «          ||         | |<   d S )Nr   r   r#   r4   )r   r   r(   s      r   Údyn_shared_memoryr7   G   sS   € ÝŒ	�!‰Œ€AÝ	Œ×	Ò	˜1�gÑ	&Ô	&€BØ�‰E€B€q�EÝÔÑÔÐØ�ŒU€Cˆ�F€F€Fr   c                 ól   — | dxx         dz  cc<   t          j        ¦   «          | dxx         dz  cc<   d S ©Nr   é{   iA  )r   Úthreadfence©r   s    r   Úuse_threadfencer=   O   s?   € Øˆ€F€F„Fˆc�M€F€F�FÝÔÑÔÐØˆ€F€F„Fˆc�M€F€F�F€F€Fr   c                 ól   — | dxx         dz  cc<   t          j        ¦   «          | dxx         dz  cc<   d S r9   )r   Úthreadfence_blockr<   s    r   Úuse_threadfence_blockr@   U   s?   € Øˆ€F€F„Fˆc�M€F€F�FÝÔÑÔÐØˆ€F€F„Fˆc�M€F€F�F€F€Fr   c                 ól   — | dxx         dz  cc<   t          j        ¦   «          | dxx         dz  cc<   d S r9   )r   Úthreadfence_systemr<   s    r   Úuse_threadfence_systemrC   [   s?   € Øˆ€F€F„Fˆc�M€F€F�FÝÔÑÔÐØˆ€F€F„Fˆc�M€F€F�F€F€Fr   c                 óh   — t          j        d¦  «        }t          j        | |         ¦  «        ||<   d S r   )r   r   Úsyncthreads_count©Úary_inÚary_outr   s      r   Úuse_syncthreads_countrI   a   s+   € ÝŒ	�!‰Œ€AÝÔ'¨¨q¬	Ñ2Ô2€GˆA�J€J€Jr   c                 óh   — t          j        d¦  «        }t          j        | |         ¦  «        ||<   d S r   )r   r   Úsyncthreads_andrF   s      r   Úuse_syncthreads_andrL   f   s+   € ÝŒ	�!‰Œ€AÝÔ% f¨Q¤iÑ0Ô0€GˆA�J€J€Jr   c                 óh   — t          j        d¦  «        }t          j        | |         ¦  «        ||<   d S r   )r   r   Úsyncthreads_orrF   s      r   Úuse_syncthreads_orrO   k   s+   € ÝŒ	�!‰Œ€AÝÔ$ V¨A¤YÑ/Ô/€GˆA�J€J€Jr   c                 óL   — t           rdS t          j        ¦   «         j        | k    S )NT)r	   r   Úget_current_deviceÚcompute_capability)Úccs    r   Ú_safe_cc_checkrT   p   s'   € Ýð BØˆtåÔ&Ñ(Ô(Ô;¸rÒAÐAr   c                   ó~  — e Zd Zd„ Zd„ Z ed¦  «        d„ ¦   «         Z ed¦  «         ej         e	d¦  «        d¦  «        d„ ¦   «         ¦   «         Z
 ed¦  «         ej         e	d¦  «        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S )ÚTestCudaSyncc                 ó*  —  t          j        d¦  «        |¦  «        }d}t          j        |t          j        ¬¦  «        }t          j        |t          j        ¬¦  «        } |d|f         |¦  «         t          j                             ||¦  «         d S )Núvoid(int32[::1])r2   ©Údtyper   )r   ÚjitÚnpÚemptyr   ÚarangeÚtestingÚassert_equal)ÚselfÚkernelÚcompiledÚnelemr   Úexps         r   Ú_test_uselesszTestCudaSync._test_uselessx   s„   € Ø/•4”8Ð.Ñ/Ô/°Ñ7Ô7ˆØˆÝŒh�u¥B¤HÐ-Ñ-Ô-ˆÝŒi˜¥R¤XÐ.Ñ.Ô.ˆØˆ��E�Ô˜3ÑÔÐÝ
Œ
×Ò  SÑ)Ô)Ð)Ð)Ð)r   c                 ó:   — |                       t          ¦  «         d S ©N)rf   r   ©ra   s    r   Útest_useless_syncthreadsz%TestCudaSync.test_useless_syncthreads€   s   € Ø×ÒÕ.Ñ/Ô/Ð/Ð/Ð/r   z#syncwarp not implemented on cudasimc                 ó:   — |                       t          ¦  «         d S rh   )rf   r   ri   s    r   Útest_useless_syncwarpz"TestCudaSync.test_useless_syncwarpƒ   s   € à×ÒÕ+Ñ,Ô,Ð,Ð,Ð,r   )é   r   z'Partial masks require CC 7.0 or greaterc                 ó:   — |                       t          ¦  «         d S rh   )rf   r   ri   s    r   Útest_useless_syncwarp_with_maskz,TestCudaSync.test_useless_syncwarp_with_mask‡   s   € ð 	×ÒÕ5Ñ6Ô6Ð6Ð6Ð6r   c                 ó  — d}d}d} t          j        d¦  «        t          ¦  «        }t          j        dt          j        ¬¦  «        } |||f         |¦  «         t          j                             ||d         ¦  «         d S )Nið  r   r   rX   rY   r   )r   r[   r)   r\   Úzerosr   r_   r`   )ra   ÚexpectedÚnthreadsÚnblocksrc   r'   s         r   Útest_coop_syncwarpzTestCudaSync.test_coop_syncwarp�   s   € ð ˆØˆØˆà/•4”8Ð.Ñ/Ô/µÑ>Ô>ˆÝŒh�q¥¤Ð)Ñ)Ô)ˆØ#ˆ�˜(Ð"Ô# CÑ(Ô(Ð(Ý
Œ
×Ò ¨#¨a¬&Ñ1Ô1Ð1Ð1Ð1r   c           	      óF  —  t          j        d¦  «        t          ¦  «        }d}t          j        |t          j        ¬¦  «        } |d|f         |¦  «         |                      t          j        |t          j        |t          j        ¬¦  «        k    ¦  «        ¦  «         d S )NrX   r+   rY   r   )	r   r[   r/   r\   r]   r   Ú
assertTrueÚallr^   )ra   rc   rd   r   s       r   Útest_simple_smemzTestCudaSync.test_simple_smemœ   sˆ   € Ø/•4”8Ð.Ñ/Ô/µÑ<Ô<ˆØˆÝŒh�u¥B¤HÐ-Ñ-Ô-ˆØˆ��E�Ô˜3ÑÔÐØ�Š�œ˜s¥b¤i°½R¼XÐ&FÑ&FÔ&FÒFÑGÔGÑHÔHÐHÐHÐHr   c                 óÀ  —  t          j        d¦  «        t          ¦  «        }d}t          j        |t          j        ¬¦  «        } |d|f         |¦  «         t          j        |¦  «        }t          |j        d         ¦  «        D ]/}t          |j        d         ¦  «        D ]}|dz   |dz   z  |||f<   ŒŒ0|  	                    t          j
        ||¦  «        ¦  «         d S )Nzvoid(float32[:,::1])r1   rY   r   r   )r   r[   r5   r\   r]   r   Ú
empty_liker,   Úshaperw   Úallclose)ra   rc   r|   r   re   r   r.   s          r   Útest_coop_smem2dzTestCudaSync.test_coop_smem2d£   sÞ   € Ø3•4”8Ð2Ñ3Ô3µKÑ@Ô@ˆØˆÝŒh�u¥B¤JÐ/Ñ/Ô/ˆØˆ��E�Ô˜3ÑÔÐÝŒm˜CÑ Ô ˆÝ�s”y ”|Ñ$Ô$ð 	.ð 	.ˆAÝ˜3œ9 Qœ<Ñ(Ô(ð .ð .�Ø ™U q¨1¡uÑ-��A�q�D‘	�	ð.à�Š�œ C¨Ñ-Ô-Ñ.Ô.Ð.Ð.Ð.r   c           
      ój  —  t          j        d¦  «        t          ¦  «        }d}t          j        |t          j        ¬¦  «        } |d|d|j        dz  f         |¦  «         |                      t          j        |dt          j	        |j        t          j
        ¬¦  «        z  k    ¦  «        ¦  «         d S )Nzvoid(float32[::1])é2   rY   r   r   r!   r#   )r   r[   r7   r\   r]   r   Úsizerw   rx   r^   r   )ra   rc   r|   r   s       r   Útest_dyn_shared_memoryz#TestCudaSync.test_dyn_shared_memory®   s›   € Ø1•4”8Ð0Ñ1Ô1Õ2CÑDÔDˆØˆÝŒh�u¥B¤JÐ/Ñ/Ô/ˆØ+ˆ��E˜1˜cœh¨™lÐ*Ô+¨CÑ0Ô0Ð0Ø�Š�œ˜s a­"¬)°C´HÅBÄHÐ*MÑ*MÔ*MÑ&MÒMÑNÔNÑOÔOÐOÐOÐOr   c                 óh  — t           d d …         f} t          j        |¦  «        t          ¦  «        }t	          j        dt          j         ¬¦  «        } |d         |¦  «         |                      d|d         ¦  «         t          s+|                      d| 	                    |¦  «        ¦  «         d S d S )Nr2   rY   ©r   r   é¼  r   z
membar.gl;)
r   r   r[   r=   r\   rq   ÚassertEqualr	   ÚassertInÚinspect_asm©ra   Úsigrc   r   s       r   Útest_threadfence_codegenz%TestCudaSync.test_threadfence_codegenµ   s©   € å�Q�Q�QŒxˆkˆØ •4”8˜C‘=”=¥Ñ1Ô1ˆÝŒh�r¥¤Ð*Ñ*Ô*ˆØˆ�Œ�sÑÔÐØ×Ò˜ C¨¤FÑ+Ô+Ð+Ýð 	CØ�MŠM˜,¨×(<Ò(<¸SÑ(AÔ(AÑBÔBÐBÐBÐBð	Cð 	Cr   c                 óh  — t           d d …         f} t          j        |¦  «        t          ¦  «        }t	          j        dt          j         ¬¦  «        } |d         |¦  «         |                      d|d         ¦  «         t          s+|                      d| 	                    |¦  «        ¦  «         d S d S )Nr2   rY   r„   r…   r   zmembar.cta;)
r   r   r[   r@   r\   rq   r†   r	   r‡   rˆ   r‰   s       r   Útest_threadfence_block_codegenz+TestCudaSync.test_threadfence_block_codegen¿   sª   € å�Q�Q�QŒxˆkˆØ •4”8˜C‘=”=Õ!6Ñ7Ô7ˆÝŒh�r¥¤Ð*Ñ*Ô*ˆØˆ�Œ�sÑÔÐØ×Ò˜ C¨¤FÑ+Ô+Ð+Ýð 	DØ�MŠM˜-¨×)=Ò)=¸cÑ)BÔ)BÑCÔCÐCÐCÐCð	Dð 	Dr   c                 óh  — t           d d …         f} t          j        |¦  «        t          ¦  «        }t	          j        dt          j         ¬¦  «        } |d         |¦  «         |                      d|d         ¦  «         t          s+|                      d| 	                    |¦  «        ¦  «         d S d S )Nr2   rY   r„   r…   r   zmembar.sys;)
r   r   r[   rC   r\   rq   r†   r	   r‡   rˆ   r‰   s       r   Útest_threadfence_system_codegenz,TestCudaSync.test_threadfence_system_codegenÉ   sª   € å�Q�Q�QŒxˆkˆØ •4”8˜C‘=”=Õ!7Ñ8Ô8ˆÝŒh�r¥¤Ð*Ñ*Ô*ˆØˆ�Œ�sÑÔÐØ×Ò˜ C¨¤FÑ+Ô+Ð+Ýð 	DØ�MŠM˜-¨×)=Ò)=¸cÑ)BÔ)BÑCÔCÐCÐCÐCð	Dð 	Dr   c                 ó2  — t          j        t          ¦  «        }t          j        d|¬¦  «        }t          j        dt          j        ¬¦  «        }d|d<   d|d<    |d         ||¦  «         |                      t          j        |dk    ¦  «        ¦  «         d S )NéH   rY   r   é   é*   )r   r‘   éF   )	r   r[   rI   r\   Úonesrq   r   rw   rx   )ra   Úin_dtyperc   rG   rH   s        r   Ú_test_syncthreads_countz$TestCudaSync._test_syncthreads_countÓ   sŠ   € Ý”8Õ1Ñ2Ô2ˆÝ”˜ 8Ð,Ñ,Ô,ˆÝ”(˜2¥R¤XÐ.Ñ.Ô.ˆØˆˆr‰
Øˆˆr‰
Øˆ�Œ˜ Ñ(Ô(Ð(Ø�Š�œ˜w¨"š}Ñ-Ô-Ñ.Ô.Ð.Ð.Ð.r   c                 óD   — |                       t          j        ¦  «         d S rh   )r—   r\   r   ri   s    r   Útest_syncthreads_countz#TestCudaSync.test_syncthreads_countÜ   ó   € Ø×$Ò$¥R¤XÑ.Ô.Ð.Ð.Ð.r   c                 óD   — |                       t          j        ¦  «         d S rh   )r—   r\   Úint16ri   s    r   Útest_syncthreads_count_upcastz*TestCudaSync.test_syncthreads_count_upcastß   rš   r   c                 óD   — |                       t          j        ¦  «         d S rh   )r—   r\   Úint64ri   s    r   Útest_syncthreads_count_downcastz,TestCudaSync.test_syncthreads_count_downcastâ   rš   r   c                 ó®  — t          j        t          ¦  «        }d}t          j        ||¬¦  «        }t          j        |t          j        ¬¦  «        } |d|f         ||¦  «         |                      t          j        |dk    ¦  «        ¦  «         d|d<    |d|f         ||¦  «         |                      t          j        |dk    ¦  «        ¦  «         d S ©Nr+   rY   r   r   r’   )	r   r[   rL   r\   r•   rq   r   rw   rx   ©ra   r–   rc   rd   rG   rH   s         r   Ú_test_syncthreads_andz"TestCudaSync._test_syncthreads_andå   sÆ   € Ý”8Õ/Ñ0Ô0ˆØˆÝ”˜ hÐ/Ñ/Ô/ˆÝ”(˜5­¬Ð1Ñ1Ô1ˆØˆ��E�Ô˜6 7Ñ+Ô+Ð+Ø�Š�œ˜w¨!š|Ñ,Ô,Ñ-Ô-Ð-Øˆˆr‰
Øˆ��E�Ô˜6 7Ñ+Ô+Ð+Ø�Š�œ˜w¨!š|Ñ,Ô,Ñ-Ô-Ð-Ð-Ð-r   c                 óD   — |                       t          j        ¦  «         d S rh   )r¤   r\   r   ri   s    r   Útest_syncthreads_andz!TestCudaSync.test_syncthreads_andð   ó   € Ø×"Ò"¥2¤8Ñ,Ô,Ð,Ð,Ð,r   c                 óD   — |                       t          j        ¦  «         d S rh   )r¤   r\   rœ   ri   s    r   Útest_syncthreads_and_upcastz(TestCudaSync.test_syncthreads_and_upcastó   r§   r   c                 óD   — |                       t          j        ¦  «         d S rh   )r¤   r\   rŸ   ri   s    r   Útest_syncthreads_and_downcastz*TestCudaSync.test_syncthreads_and_downcastö   r§   r   c                 ó®  — t          j        t          ¦  «        }d}t          j        ||¬¦  «        }t          j        |t          j        ¬¦  «        } |d|f         ||¦  «         |                      t          j        |dk    ¦  «        ¦  «         d|d<    |d|f         ||¦  «         |                      t          j        |dk    ¦  «        ¦  «         d S r¢   )r   r[   rO   r\   rq   r   rw   rx   r£   s         r   Ú_test_syncthreads_orz!TestCudaSync._test_syncthreads_orù   sÆ   € Ý”8Õ.Ñ/Ô/ˆØˆÝ”˜% xÐ0Ñ0Ô0ˆÝ”(˜5­¬Ð1Ñ1Ô1ˆØˆ��E�Ô˜6 7Ñ+Ô+Ð+Ø�Š�œ˜w¨!š|Ñ,Ô,Ñ-Ô-Ð-Øˆˆr‰
Øˆ��E�Ô˜6 7Ñ+Ô+Ð+Ø�Š�œ˜w¨!š|Ñ,Ô,Ñ-Ô-Ð-Ð-Ð-r   c                 óD   — |                       t          j        ¦  «         d S rh   )r­   r\   r   ri   s    r   Útest_syncthreads_orz TestCudaSync.test_syncthreads_or  ó   € Ø×!Ò!¥"¤(Ñ+Ô+Ð+Ð+Ð+r   c                 óD   — |                       t          j        ¦  «         d S rh   )r­   r\   rœ   ri   s    r   Útest_syncthreads_or_upcastz'TestCudaSync.test_syncthreads_or_upcast  r°   r   c                 óD   — |                       t          j        ¦  «         d S rh   )r­   r\   rŸ   ri   s    r   Útest_syncthreads_or_downcastz)TestCudaSync.test_syncthreads_or_downcast
  r°   r   N)Ú__name__Ú
__module__Ú__qualname__rf   rj   r   rl   r   Ú
skipUnlessrT   ro   ru   ry   r~   r‚   r‹   r�   r�   r—   r™   r�   r    r¤   r¦   r©   r«   r­   r¯   r²   r´   © r   r   rV   rV   w   s"  € € € € € ð*ð *ð *ð0ð 0ð 0ð €_Ð:Ñ;Ô;ð-ð -ñ <Ô;ð-ð €_Ð:Ñ;Ô;Ø€XÔ˜˜¨Ñ/Ô/ØBñDô Dð7ð 7ñDô Dñ <Ô;ð7ð €_Ð:Ñ;Ô;Ø€XÔ˜˜¨Ñ/Ô/ØBñDô Dð
2ð 
2ñDô Dñ <Ô;ð
2ðIð Ið Ið	/ð 	/ð 	/ðPð Pð PðCð Cð CðDð Dð DðDð Dð Dð/ð /ð /ð/ð /ð /ð/ð /ð /ð/ð /ð /ð	.ð 	.ð 	.ð-ð -ð -ð-ð -ð -ð-ð -ð -ð	.ð 	.ð 	.ð,ð ,ð ,ð,ð ,ð ,ð,ð ,ð ,ð ,ð ,r   rV   Ú__main__)Únumpyr\   Únumbar   r   r   Únumba.cuda.testingr   r   r   Únumba.core.configr	   r   r   r   r)   r/   r5   r7   r=   r@   rC   rI   rL   rO   rT   rV   rµ   Úmainr¹   r   r   ú<module>rÀ      s�  ðØ Ð Ð Ð Ø &Ð &Ð &Ð &Ð &Ð &Ð &Ð &Ð &Ð &Ø FÐ FÐ FÐ FÐ FÐ FÐ FÐ FÐ FÐ FØ ,Ð ,Ð ,Ð ,Ð ,Ð ,ðð ð ðð ð ðð ð ðð ð ð6ð ð ðð ð ðð ð ðð ð ðð ð ðð ð ð3ð 3ð 3ð
1ð 1ð 1ð
0ð 0ð 0ð
Bð Bð BðT,ð T,ð T,ð T,ð T,�<ñ T,ô T,ð T,ðn ˆzÒÐØ€H„M�O„O€O€O€Oð Ðr   