
    sh                         S SK rS SKJrJrJr  S SKJrJrJ	r	  S SK
Jr  S rS rS rS rS	 rS
 rS rS rS rS rS rS rS rS r " S S\	5      r\S:X  a  \R8                  " 5         gg)    N)cudaint32float32)skip_on_cudasimunittestCUDATestCase)ENABLE_CUDASIMc                 b    [         R                  " S5      n[         R                  " 5         XU'   g N   )r   gridsyncthreadsaryis     u/Users/tiagomarins/Projetos/claudeai/copy_bank/venv/lib/python3.13/site-packages/numba/cuda/tests/cudapy/test_sync.pyuseless_syncthreadsr      s"    		!AF    c                 b    [         R                  " S5      n[         R                  " 5         XU'   g r   r   r   syncwarpr   s     r   useless_syncwarpr      s    		!AMMOFr   c                 d    [         R                  " S5      n[         R                  " S5        XU'   g )Nr     r   r   s     r   useless_syncwarp_with_maskr      s"    		!AMM&Fr   c                 $   [         R                  R                  S[        5      n[         R                  " S5      nX!U'   [         R
                  " 5         US:  a$  X   XS-      -   X'   [         R
                  " S5        US:  a$  X   XS-      -   X'   [         R
                  " S5        US:  a$  X   XS-      -   X'   [         R
                  " S5        US	:  a$  X   XS	-      -   X'   [         R
                  " S
5        US:X  a  US   US   -   U S'   g g )N    r      r                     r   )r   sharedarrayr   r   r   )ressmr   s      r   coop_syncwarpr)      s    			2u	%B		!AqEMMO2vr6
"f1uq5	!d1uq5	!c1uq5	!cAvAAA r   c                     Sn[         R                  R                  U[        5      n[         R                  " S5      nUS:X  a  [        U5       H  nXBU'   M	     [         R                  " 5         X#   X'   g )Nd   r   r   )r   r%   r&   r   r   ranger   )r   Nr(   r   js        r   simple_smemr/   4   s[    A			1e	$B		!AAvqAqE UCFr   c                     [         R                  " S5      u  p[         R                  R                  S[        5      nUS-   US-   -  X1U4'   [         R
                  " 5         X1U4   XU4'   g )Nr#   
      r   r   r   r%   r&   r   r   )r   r   r.   r(   s       r   coop_smem2dr5   ?   s\    99Q<DA			8W	-BA!a% B!tHa4C1Ir   c                     [         R                  " S5      n[         R                  R                  S[        5      nUS-  X!'   [         R
                  " 5         X!   X'   g )Nr   r   r#   r4   )r   r   r(   s      r   dyn_shared_memoryr7   G   sE    		!A			1g	&BEBEUCFr   c                 b    U S==   S-  ss'   [         R                  " 5         U S==   S-  ss'   g Nr   {   iA  )r   threadfencer   s    r   use_threadfencer=   O   s(    FcMFFcMFr   c                 b    U S==   S-  ss'   [         R                  " 5         U S==   S-  ss'   g r9   )r   threadfence_blockr<   s    r   use_threadfence_blockr@   U   s(    FcMFFcMFr   c                 b    U S==   S-  ss'   [         R                  " 5         U S==   S-  ss'   g r9   )r   threadfence_systemr<   s    r   use_threadfence_systemrC   [   s(    FcMFFcMFr   c                 d    [         R                  " S5      n[         R                  " X   5      X'   g r   )r   r   syncthreads_countary_inary_outr   s      r   use_syncthreads_countrI   a   s#    		!A''	2GJr   c                 d    [         R                  " S5      n[         R                  " X   5      X'   g r   )r   r   syncthreads_andrF   s      r   use_syncthreads_andrL   f   s#    		!A%%fi0GJr   c                 d    [         R                  " S5      n[         R                  " X   5      X'   g r   )r   r   syncthreads_orrF   s      r   use_syncthreads_orrO   k   s#    		!A$$VY/GJr   c                 ^    [         (       a  g[        R                  " 5       R                  U :  $ )NT)r	   r   get_current_devicecompute_capability)ccs    r   _safe_cc_checkrT   p   s$    ~&&(;;rAAr   c                   P   \ rS rSrS rS r\" S5      S 5       r\" S5      \R                  " \
" S5      S5      S 5       5       r\" S5      \R                  " \
" S5      S5      S	 5       5       rS
 rS rS rS rS rS rS rS rS rS rS rS rS rS rS rS rS rS rSrg)TestCudaSyncw   c                 $   [         R                  " S5      " U5      nSn[        R                  " U[        R                  S9n[        R
                  " U[        R                  S9nUSU4   " U5        [        R                  R                  XE5        g )Nvoid(int32[::1])r2   dtyper   )r   jitnpemptyr   arangetestingassert_equal)selfkernelcompilednelemr   exps         r   _test_uselessTestCudaSync._test_uselessx   sf    88./7hhuBHH-iiRXX.E3


)r   c                 .    U R                  [        5        g N)rg   r   rb   s    r   test_useless_syncthreads%TestCudaSync.test_useless_syncthreads   s    ./r   z#syncwarp not implemented on cudasimc                 .    U R                  [        5        g rj   )rg   r   rk   s    r   test_useless_syncwarp"TestCudaSync.test_useless_syncwarp   s    +,r   )   r   z'Partial masks require CC 7.0 or greaterc                 .    U R                  [        5        g rj   )rg   r   rk   s    r   test_useless_syncwarp_with_mask,TestCudaSync.test_useless_syncwarp_with_mask   s     	56r   c                     SnSnSn[         R                  " S5      " [        5      n[        R                  " S[        R
                  S9nXCU4   " U5        [        R                  R                  XS   5        g )Ni  r   r   rY   rZ   r   )r   r\   r)   r]   zerosr   r`   ra   )rb   expectednthreadsnblocksrd   r'   s         r   test_coop_syncwarpTestCudaSync.test_coop_syncwarp   sa     88./>hhq)("#C(


a&1r   c           
      :   [         R                  " S5      " [        5      nSn[        R                  " U[        R
                  S9nUSU4   " U5        U R                  [        R                  " U[        R                  " U[        R
                  S9:H  5      5        g )NrY   r+   rZ   r   )	r   r\   r/   r]   r^   r   
assertTrueallr_   )rb   rd   re   r   s       r   test_simple_smemTestCudaSync.test_simple_smem   si    88./<hhuBHH-E3sbiiRXX&FFGHr   c                    [         R                  " S5      " [        5      nSn[        R                  " U[        R
                  S9nUSU4   " U5        [        R                  " U5      n[        UR                  S   5       H1  n[        UR                  S   5       H  nUS-   US-   -  XEU4'   M     M3     U R                  [        R                  " X45      5        g )Nzvoid(float32[:,::1])r1   rZ   r   r   )r   r\   r5   r]   r^   r   
empty_liker,   shaper}   allclose)rb   rd   r   r   rf   r   r.   s          r   test_coop_smem2dTestCudaSync.test_coop_smem2d   s    8823K@hhuBJJ/E3mmC syy|$A399Q<(Uq1u-qD	 ) % 	C-.r   c                 r   [         R                  " S5      " [        5      nSn[        R                  " U[        R
                  S9nUSUSUR                  S-  4   " U5        U R                  [        R                  " US[        R                  " UR                  [        R                  S9-  :H  5      5        g )Nzvoid(float32[::1])2   rZ   r   r   r!   r#   )r   r\   r7   r]   r^   r   sizer}   r~   r_   r   )rb   rd   r   r   s       r   test_dyn_shared_memory#TestCudaSync.test_dyn_shared_memory   s    88012CDhhuBJJ/E1chhl*+C0sa"))CHHBHH*M&MMNOr   c                 8   [         S S  4n[        R                  " U5      " [        5      n[        R
                  " S[        R                   S9nUS   " U5        U R                  SUS   5        [        (       d"  U R                  SUR                  U5      5        g g )Nr2   rZ   r   r     r   z
membar.gl;)
r   r   r\   r=   r]   rv   assertEqualr	   assertIninspect_asmrb   sigrd   r   s       r   test_threadfence_codegen%TestCudaSync.test_threadfence_codegen   sv    Qxk88C=1hhr*sCF+~MM,(<(<S(AB r   c                 8   [         S S  4n[        R                  " U5      " [        5      n[        R
                  " S[        R                   S9nUS   " U5        U R                  SUS   5        [        (       d"  U R                  SUR                  U5      5        g g )Nr2   rZ   r   r   r   zmembar.cta;)
r   r   r\   r@   r]   rv   r   r	   r   r   r   s       r   test_threadfence_block_codegen+TestCudaSync.test_threadfence_block_codegen   sw    Qxk88C=!67hhr*sCF+~MM-)=)=c)BC r   c                 8   [         S S  4n[        R                  " U5      " [        5      n[        R
                  " S[        R                   S9nUS   " U5        U R                  SUS   5        [        (       d"  U R                  SUR                  U5      5        g g )Nr2   rZ   r   r   r   zmembar.sys;)
r   r   r\   rC   r]   rv   r   r	   r   r   r   s       r   test_threadfence_system_codegen,TestCudaSync.test_threadfence_system_codegen   sw    Qxk88C=!78hhr*sCF+~MM-)=)=c)BC r   c                 "   [         R                  " [        5      n[        R                  " SUS9n[        R
                  " S[        R                  S9nSUS'   SUS'   US   " X45        U R                  [        R                  " US:H  5      5        g )NH   rZ   r      *   )r   r   F   )	r   r\   rI   r]   onesrv   r   r}   r~   )rb   in_dtyperd   rG   rH   s        r   _test_syncthreads_count$TestCudaSync._test_syncthreads_count   sl    88128,((2RXX.r
r
(w"}-.r   c                 B    U R                  [        R                  5        g rj   )r   r]   r   rk   s    r   test_syncthreads_count#TestCudaSync.test_syncthreads_count       $$RXX.r   c                 B    U R                  [        R                  5        g rj   )r   r]   int16rk   s    r   test_syncthreads_count_upcast*TestCudaSync.test_syncthreads_count_upcast   r   r   c                 B    U R                  [        R                  5        g rj   )r   r]   int64rk   s    r   test_syncthreads_count_downcast,TestCudaSync.test_syncthreads_count_downcast   r   r   c                    [         R                  " [        5      nSn[        R                  " X1S9n[        R
                  " U[        R                  S9nUSU4   " XE5        U R                  [        R                  " US:H  5      5        SUS'   USU4   " XE5        U R                  [        R                  " US:H  5      5        g Nr+   rZ   r   r   r   )	r   r\   rL   r]   r   rv   r   r}   r~   rb   r   rd   re   rG   rH   s         r   _test_syncthreads_and"TestCudaSync._test_syncthreads_and   s    88/0/((51E6+w!|,-r
E6+w!|,-r   c                 B    U R                  [        R                  5        g rj   )r   r]   r   rk   s    r   test_syncthreads_and!TestCudaSync.test_syncthreads_and       ""288,r   c                 B    U R                  [        R                  5        g rj   )r   r]   r   rk   s    r   test_syncthreads_and_upcast(TestCudaSync.test_syncthreads_and_upcast   r   r   c                 B    U R                  [        R                  5        g rj   )r   r]   r   rk   s    r   test_syncthreads_and_downcast*TestCudaSync.test_syncthreads_and_downcast   r   r   c                    [         R                  " [        5      nSn[        R                  " X1S9n[        R                  " U[        R
                  S9nUSU4   " XE5        U R                  [        R                  " US:H  5      5        SUS'   USU4   " XE5        U R                  [        R                  " US:H  5      5        g r   )r   r\   rO   r]   rv   r   r}   r~   r   s         r   _test_syncthreads_or!TestCudaSync._test_syncthreads_or   s    88./%0((51E6+w!|,-r
E6+w!|,-r   c                 B    U R                  [        R                  5        g rj   )r   r]   r   rk   s    r   test_syncthreads_or TestCudaSync.test_syncthreads_or      !!"((+r   c                 B    U R                  [        R                  5        g rj   )r   r]   r   rk   s    r   test_syncthreads_or_upcast'TestCudaSync.test_syncthreads_or_upcast  r   r   c                 B    U R                  [        R                  5        g rj   )r   r]   r   rk   s    r   test_syncthreads_or_downcast)TestCudaSync.test_syncthreads_or_downcast
  r   r    N) __name__
__module____qualname____firstlineno__rg   rl   r   ro   r   
skipUnlessrT   rs   rz   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   __static_attributes__r   r   r   rV   rV   w   s    *0 :;- <- :;/BD7D <7 :;/BD
2D <
2I	/PC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 ,63
1
0
BT,< T,n zMMO r   