+
    :i                         ^ RI t^ RIHtHtHt ^ RIHtHtH	t	 ^ RI
Ht R tR tR tR tR	 tR
 tR tR tR tR tR tR tR tR t ! R R]	4      t]R8X  d   ]P8                  ! 4        R# R# )    N)cudaint32float32)skip_on_cudasimunittestCUDATestCase)ENABLE_CUDASIMc                 d    \         P                  ! ^4      p\         P                  ! 4        WV&   R#    N)r   gridsyncthreadsaryis   & _/var/www/html/photoedit/myenv/lib/python3.14/site-packages/numba/cuda/tests/cudapy/test_sync.pyuseless_syncthreadsr      s"    		!AF    c                 d    \         P                  ! ^4      p\         P                  ! 4        WV&   R# r   r   r   syncwarpr   s   & r   useless_syncwarpr      s    		!AMMOFr   c                 f    \         P                  ! ^4      p\         P                  ! R4       WV&   R# )r     Nr   r   s   & r   useless_syncwarp_with_maskr      s"    		!AMM&Fr   c                    \         P                  P                  ^ \        4      p\         P                  ! ^4      pW!V&   \         P
                  ! 4        V^8  d5   W,          W^,           ,          ,           W&   \         P
                  ! R4       V^8  d5   W,          W^,           ,          ,           W&   \         P
                  ! ^4       V^8  d5   W,          W^,           ,          ,           W&   \         P
                  ! ^4       V^8  d5   W,          W^,           ,          ,           W&   \         P
                  ! ^4       V^ 8X  d   V^ ,          V^,          ,           V ^ &   R# R# )    r   N)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                     ^dp\         P                  P                  V\        4      p\         P                  ! ^4      pV^ 8X  d   \        V4       F  pWBV&   K	  	  \         P                  ! 4        W#,          W&   R# )d   N)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                     \         P                  ! ^4      w  r\         P                  P                  R\        4      pV^,           V^,           ,          W1V3&   \         P
                  ! 4        W1V3,          WV3&   R# )   N
      r   r   r   r   r   r   )r   r   r'   r!   s   &   r   coop_smem2dr/   ?   s\    99Q<DA			8W	-BA!a% B!tHa4C1Ir   c                     \         P                  ! ^4      p\         P                  P                  ^ \        4      pV^,          W!&   \         P
                  ! 4        W!,          W&   R# r   r.   )r   r   r!   s   &  r   dyn_shared_memoryr1   G   sE    		!A			1g	&BEBEUCFr   c                     V ^ ;;,          ^{,          uu&   \         P                  ! 4        V ^ ;;,          R,          uu&   R# r   iA  N)r   threadfencer   s   &r   use_threadfencer6   O   s(    FcMFFcMFr   c                     V ^ ;;,          ^{,          uu&   \         P                  ! 4        V ^ ;;,          R,          uu&   R# r3   )r   threadfence_blockr5   s   &r   use_threadfence_blockr9   U   s(    FcMFFcMFr   c                     V ^ ;;,          ^{,          uu&   \         P                  ! 4        V ^ ;;,          R,          uu&   R# r3   )r   threadfence_systemr5   s   &r   use_threadfence_systemr<   [   s(    FcMFFcMFr   c                 n    \         P                  ! ^4      p\         P                  ! W,          4      W&   R# r   )r   r   syncthreads_countary_inary_outr   s   && r   use_syncthreads_countrB   a   s#    		!A''	2GJr   c                 n    \         P                  ! ^4      p\         P                  ! W,          4      W&   R# r   )r   r   syncthreads_andr?   s   && r   use_syncthreads_andrE   f   s#    		!A%%fi0GJr   c                 n    \         P                  ! ^4      p\         P                  ! W,          4      W&   R# r   )r   r   syncthreads_orr?   s   && r   use_syncthreads_orrH   k   s#    		!A$$VY/GJr   c                 b    \         '       d   R # \        P                  ! 4       P                  V 8  # )T)r	   r   get_current_devicecompute_capability)ccs   &r   _safe_cc_checkrM   p   s$    ~&&(;;rAAr   c                   \  a  ] tR t^wt o R tR t]! R4      R 4       t]! R4      ]P                  ! ]
! R4      R4      R 4       4       t]! R4      ]P                  ! ]
! R4      R4      R 4       4       tR tR	 tR
 tR tR tR tR tR tR tR tR tR tR tR tR tR tR tR tRtV t R# )TestCudaSyncc                :   \         P                  ! R 4      ! V4      p^
p\        P                  ! V\        P                  R7      p\        P
                  ! V\        P                  R7      pV^V3,          ! V4       \        P                  P                  WE4       R# void(int32[::1])dtypeN)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                0    V P                  \        4       R # N)r`   r   r[   s   &r   test_useless_syncthreads%TestCudaSync.test_useless_syncthreads   s    ./r   z#syncwarp not implemented on cudasimc                0    V P                  \        4       R # rc   )r`   r   rd   s   &r   test_useless_syncwarp"TestCudaSync.test_useless_syncwarp   s    +,r   z'Partial masks require CC 7.0 or greaterc                0    V P                  \        4       R # rc   )r`   r   rd   s   &r   test_useless_syncwarp_with_mask,TestCudaSync.test_useless_syncwarp_with_mask   s     	56r   c                
   R p^ p^p\         P                  ! R4      ! \        4      p\        P                  ! ^\        P
                  R7      pWCV3,          ! V4       \        P                  P                  W^ ,          4       R# )i  rR   rS   N)r   rU   r"   rV   zerosr   rY   rZ   )r[   expectednthreadsnblocksr]   r    s   &     r   test_coop_syncwarpTestCudaSync.test_coop_syncwarp   sa     88./>hhq)("#C(


a&1r   c           
     P   \         P                  ! R 4      ! \        4      p^dp\        P                  ! V\        P
                  R7      pV^V3,          ! V4       V P                  \        P                  ! V\        P                  ! V\        P
                  R7      8H  4      4       R# rQ   )	r   rU   r(   rV   rW   r   
assertTrueallrX   )r[   r]   r^   r   s   &   r   test_simple_smemTestCudaSync.test_simple_smem   si    88./<hhuBHH-E3sbiiRXX&FFGHr   c                   \         P                  ! R 4      ! \        4      pRp\        P                  ! V\        P
                  R7      pV^V3,          ! V4       \        P                  ! V4      p\        VP                  ^ ,          4       FA  p\        VP                  ^,          4       F  pV^,           V^,           ,          WEV3&   K   	  KC  	  V P                  \        P                  ! W44      4       R# )zvoid(float32[:,::1])rS   Nr+   )r   rU   r/   rV   rW   r   
empty_liker%   shaperu   allclose)r[   r]   r{   r   r_   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                   \         P                  ! R 4      ! \        4      p^2p\        P                  ! V\        P
                  R7      pV^V^ VP                  ^,          3,          ! V4       V P                  \        P                  ! V^\        P                  ! VP                  \        P                  R7      ,          8H  4      4       R# )zvoid(float32[::1])rS   N)r   rU   r1   rV   rW   r   sizeru   rv   rX   r   )r[   r]   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                \   \         R ,          3p\        P                  ! V4      ! \        4      p\        P
                  ! ^
\        P                   R7      pVR,          ! V4       V P                  RV^ ,          4       \        '       g$   V P                  RVP                  V4      4       R# R# )NNNrS   z
membar.gl;Nr   r     )
r   r   rU   r6   rV   rn   assertEqualr	   assertIninspect_asmr[   sigr]   r   s   &   r   test_threadfence_codegen%TestCudaSync.test_threadfence_codegen   sv    Qxk88C=1hhr*sCF+~MM,(<(<S(AB r   c                \   \         R ,          3p\        P                  ! V4      ! \        4      p\        P
                  ! ^
\        P                   R7      pVR,          ! V4       V P                  RV^ ,          4       \        '       g$   V P                  RVP                  V4      4       R# R# )r   rS   zmembar.cta;Nr   r   )
r   r   rU   r9   rV   rn   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                \   \         R ,          3p\        P                  ! V4      ! \        4      p\        P
                  ! ^
\        P                   R7      pVR,          ! V4       V P                  RV^ ,          4       \        '       g$   V P                  RVP                  V4      4       R# R# )r   rS   zmembar.sys;Nr   r   )
r   r   rU   r<   rV   rn   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                8   \         P                  ! \        4      p\        P                  ! ^HVR7      p\        P
                  ! ^H\        P                  R7      p^ V^&   ^ V^*&   VR,          ! W44       V P                  \        P                  ! V^F8H  4      4       R# )H   rS   N)r   r   )	r   rU   rB   rV   onesrn   r   ru   rv   )r[   in_dtyper]   r@   rA   s   &&   r   _test_syncthreads_count$TestCudaSync._test_syncthreads_count   sl    88128,((2RXX.r
r
(w"}-.r   c                D    V P                  \        P                  4       R # rc   )r   rV   r   rd   s   &r   test_syncthreads_count#TestCudaSync.test_syncthreads_count       $$RXX.r   c                D    V P                  \        P                  4       R # rc   )r   rV   int16rd   s   &r   test_syncthreads_count_upcast*TestCudaSync.test_syncthreads_count_upcast   r   r   c                D    V P                  \        P                  4       R # rc   )r   rV   int64rd   s   &r   test_syncthreads_count_downcast,TestCudaSync.test_syncthreads_count_downcast   r   r   c                   \         P                  ! \        4      p^dp\        P                  ! W1R7      p\        P
                  ! V\        P                  R7      pV^V3,          ! WE4       V P                  \        P                  ! V^8H  4      4       ^ V^&   V^V3,          ! WE4       V P                  \        P                  ! V^ 8H  4      4       R# r$   rS   N)	r   rU   rE   rV   r   rn   r   ru   rv   r[   r   r]   r^   r@   rA   s   &&    r   _test_syncthreads_and"TestCudaSync._test_syncthreads_and   s    88/0/((51E6+w!|,-r
E6+w!|,-r   c                D    V P                  \        P                  4       R # rc   )r   rV   r   rd   s   &r   test_syncthreads_and!TestCudaSync.test_syncthreads_and       ""288,r   c                D    V P                  \        P                  4       R # rc   )r   rV   r   rd   s   &r   test_syncthreads_and_upcast(TestCudaSync.test_syncthreads_and_upcast   r   r   c                D    V P                  \        P                  4       R # rc   )r   rV   r   rd   s   &r   test_syncthreads_and_downcast*TestCudaSync.test_syncthreads_and_downcast   r   r   c                   \         P                  ! \        4      p^dp\        P                  ! W1R7      p\        P                  ! V\        P
                  R7      pV^V3,          ! WE4       V P                  \        P                  ! V^ 8H  4      4       ^V^&   V^V3,          ! WE4       V P                  \        P                  ! V^8H  4      4       R# r   )r   rU   rH   rV   rn   r   ru   rv   r   s   &&    r   _test_syncthreads_or!TestCudaSync._test_syncthreads_or   s    88./%0((51E6+w!|,-r
E6+w!|,-r   c                D    V P                  \        P                  4       R # rc   )r   rV   r   rd   s   &r   test_syncthreads_or TestCudaSync.test_syncthreads_or      !!"((+r   c                D    V P                  \        P                  4       R # rc   )r   rV   r   rd   s   &r   test_syncthreads_or_upcast'TestCudaSync.test_syncthreads_or_upcast  r   r   c                D    V P                  \        P                  4       R # rc   )r   rV   r   rd   s   &r   test_syncthreads_or_downcast)TestCudaSync.test_syncthreads_or_downcast
  r   r    N)   r   )!__name__
__module____qualname____firstlineno__r`   re   r   rh   r   
skipUnlessrM   rk   rr   rw   r}   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   __static_attributes____classdictcell__)__classdict__s   @r   rO   rO   w   s     *0 :;- <- :;/BD7D <7 :;/BD
2D <
2I	/PCDD////	.---	.,,, ,r   rO   __main__)numpyrV   numbar   r   r   numba.cuda.testingr   r   r   numba.core.configr	   r   r   r   r"   r(   r/   r1   r6   r9   r<   rB   rE   rH   rM   rO   r   mainr   r   r   <module>r      s     & & F F ,63
1
0
BT,< T,n zMMO r   