
    shR#                         S SK rS SKJrJr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5       " S S\
5      5       r\S:X  a  \	R<                  " 5         gg)    N)cudaint32int64float32float64)unittestCUDATestCaseskip_on_cudasim)configc                     [         R                  " S5      nUS:X  a  SU S'   [         R                  " S5        U S   X'   g )N   r   *       )r   gridsyncwarp)aryis     y/Users/tiagomarins/Projetos/claudeai/copy_bank/venv/lib/python3.13/site-packages/numba/cuda/tests/cudapy/test_warp_ops.pyuseful_syncwarpr      s6    		!AAvAMM*VCF    c                 f    [         R                  " S5      n[         R                  " SX!5      nX0U'   g Nr   r   r   r   	shfl_sync)r   idxr   vals       r   use_shfl_sync_idxr      s&    		!A
..Q
,CFr   c                 f    [         R                  " S5      n[         R                  " SX!5      nX0U'   g r   )r   r   shfl_up_syncr   deltar   r   s       r   use_shfl_sync_upr"      s(    		!A


J
1CFr   c                 f    [         R                  " S5      n[         R                  " SX!5      nX0U'   g r   )r   r   shfl_down_syncr    s       r   use_shfl_sync_downr%      s(    		!A


j!
3CFr   c                 f    [         R                  " S5      n[         R                  " SX!5      nX0U'   g r   )r   r   shfl_xor_sync)r   xorr   r   s       r   use_shfl_sync_xorr)   !   s(    		!A


Z
0CFr   c                 h    [         R                  " S5      n[         R                  " SUS5      nX0U'   g Nr   r   r   r   )r   intor   r   s       r   use_shfl_sync_with_valr-   '   s(    		!A
..T1
-CFr   c                 j    [         R                  " S5      n[         R                  " SX   5      nX1U'   g r   )r   r   all_syncary_inary_outr   preds       r   use_vote_sync_allr4   -   (    		!A==VY/DAJr   c                 j    [         R                  " S5      n[         R                  " SX   5      nX1U'   g r   )r   r   any_syncr0   s       r   use_vote_sync_anyr8   3   r5   r   c                 j    [         R                  " S5      n[         R                  " SX   5      nX1U'   g r   )r   r   eq_syncr0   s       r   use_vote_sync_eqr;   9   s(    		!A<<
FI.DAJr   c                 n    [         R                  R                  n[         R                  " SS5      nX U'   g )Nr   Tr   	threadIdxxballot_sync)r   r   ballots      r   use_vote_sync_ballotrB   ?   s*    Aj$/FFr   c                 j    [         R                  " S5      n[         R                  " SX   5      nX1U'   g r   )r   r   match_any_sync)r1   r2   r   rA   s       r   use_match_any_syncrE   E   s*    		!A  VY7FAJr   c                     [         R                  " S5      n[         R                  " SX   5      u  p4U(       a  UOSX'   g r+   )r   r   match_all_sync)r1   r2   r   rA   r3   s        r   use_match_all_syncrH   K   s0    		!A&&z69=LFQGJr   c                 F   [         R                  R                  nUS-  S:X  a  [         R                  " SS5      nObUS-  S:X  a  [         R                  " SS5      nOAUS-  S:X  a  [         R                  " SS5      nO US-  S	:X  a  [         R                  " S
S5      nWX'   g )N   r   Tr   """"   DDDD       r=   )arrr   rA   s      r   use_independent_schedulingrR   Q   s    A1uz!!*d3	
Q!!!*d3	
Q!!!*d3	
Q!!!*d3CFr   c                 r    [         R                  (       a  g[        R                  " 5       R                  U :  $ )NT)r   ENABLE_CUDASIMr   get_current_devicecompute_capability)ccs    r   _safe_cc_checkrX   ^   s(    &&(;;rAAr   z2Warp Operations are not yet implemented on cudasimc                      \ 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\R                  " \" S5      S5      S 5       r\R                  " \" S5      S5      S 5       r\R                  " \" S5      S5      S 5       rS rS rSrg)TestCudaWarpOperationse   c                     [         R                  " S5      " [        5      nSn[        R                  " U[        R
                  S9nUSU4   " U5        U R                  [        R                  " US:H  5      5        g )Nzvoid(int32[:])    dtyper   r   )r   jitr   npemptyr   
assertTrueallselfcompilednelemr   s       r   test_useful_syncwarp+TestCudaWarpOperations.test_useful_syncwarpg   sY    88,-o>hhuBHH-E3sby)*r   c                     [         R                  " S5      " [        5      nSnSn[        R                  " U[        R
                  S9nUSU4   " XC5        U R                  [        R                  " XC:H  5      5        g Nvoid(int32[:], int32)r]   rJ   r^   r   )r   r`   r   ra   rb   r   rc   rd   )rf   rg   rh   r   r   s        r   test_shfl_sync_idx)TestCudaWarpOperations.test_shfl_sync_idxn   s]    88345FGhhuBHH-E3$sz*+r   c                 Z   [         R                  " S5      " [        5      nSnSn[        R                  " U[        R
                  S9n[        R                  " U[        R
                  S9nXSS === U-  sss& USU4   " XC5        U R                  [        R                  " XE:H  5      5        g rl   )	r   r`   r"   ra   rb   r   arangerc   rd   rf   rg   rh   r!   r   exps         r   test_shfl_sync_up(TestCudaWarpOperations.test_shfl_sync_upv   s    88345EFhhuBHH-iiRXX.FuE3&sz*+r   c                 ^   [         R                  " S5      " [        5      nSnSn[        R                  " U[        R
                  S9n[        R                  " U[        R
                  S9nUS U* === U-  sss& USU4   " XC5        U R                  [        R                  " XE:H  5      5        g rl   )	r   r`   r%   ra   rb   r   rq   rc   rd   rr   s         r   test_shfl_sync_down*TestCudaWarpOperations.test_shfl_sync_down   s    88345GHhhuBHH-iiRXX.GeVE3&sz*+r   c                 F   [         R                  " S5      " [        5      nSnSn[        R                  " U[        R
                  S9n[        R                  " U[        R
                  S9U-  nUSU4   " XC5        U R                  [        R                  " XE:H  5      5        g )Nrm   r]      r^   r   )	r   r`   r)   ra   rb   r   rq   rc   rd   )rf   rg   rh   r(   r   rs   s         r   test_shfl_sync_xor)TestCudaWarpOperations.test_shfl_sync_xor   sv    88345FGhhuBHH-iiRXX.4E3$sz*+r   c                 .   [         [        [        [        4n[        R                   " S5      [        R                  " S5      [        R                  " [        R
                  5      [        R                  " [        R
                  5      4n[        X5       H~  u  p4[        R                  " US S  U45      " [        5      nSn[        R                  " XdR                  S9nUSU4   " Xt5        U R                  [        R                  " Xt:H  5      5        M     g )Nl        r]   r^   r   )r   r   r   r   ra   pizipr   r`   r-   rb   r_   rc   rd   )rf   typesvaluestypr   rg   rh   r   s           r   test_shfl_sync_types+TestCudaWarpOperations.test_shfl_sync_types   s    ugw.((2, 1**RUU#RZZ%68E*HCxxQ./EFHE((5		2CQXs(OOBFF3:./ +r   c                    [         R                  " S5      " [        5      nSn[        R                  " U[        R
                  S9n[        R                  " U[        R
                  S9nUSU4   " X45        U R                  [        R                  " US:H  5      5        SUS'   USU4   " X45        U R                  [        R                  " US:H  5      5        g )Nvoid(int32[:], int32[:])r]   r^   r   r   r~   )	r   r`   r4   ra   onesr   rb   rc   rd   rf   rg   rh   r1   r2   s        r   test_vote_sync_all)TestCudaWarpOperations.test_vote_sync_all   s    88678IJbhh/((51E6+w!|,-r
E6+w!|,-r   c                    [         R                  " S5      " [        5      nSn[        R                  " U[        R
                  S9n[        R                  " U[        R
                  S9nUSU4   " X45        U R                  [        R                  " US:H  5      5        SUS'   SUS'   USU4   " X45        U R                  [        R                  " US:H  5      5        g )Nr   r]   r^   r   r   rM      )	r   r`   r8   ra   zerosr   rb   rc   rd   r   s        r   test_vote_sync_any)TestCudaWarpOperations.test_vote_sync_any   s    88678IJ%rxx0((51E6+w!|,-q	q	E6+w!|,-r   c                 &   [         R                  " S5      " [        5      nSn[        R                  " U[        R
                  S9n[        R                  " U[        R
                  S9nUSU4   " X45        U R                  [        R                  " US:H  5      5        SUS'   USU4   " X45        U R                  [        R                  " US:H  5      5        SUS S & USU4   " X45        U R                  [        R                  " US:H  5      5        g )Nr   r]   r^   r   r   )	r   r`   r;   ra   r   r   rb   rc   rd   r   s        r   test_vote_sync_eq(TestCudaWarpOperations.test_vote_sync_eq   s    88678HI%rxx0((51E6+w!|,-q	E6+w!|,-q	E6+w!|,-r   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
                  " S5      :H  5      5        g )Nvoid(uint32[:])r]   r^   r   r   )r   r`   rB   ra   rb   uint32rc   rd   re   s       r   test_vote_sync_ballot,TestCudaWarpOperations.test_vote_sync_ballot   sd    88-./CDhhuBII.E3sbii
&;;<=r   )   r   z-Matching requires at least Volta Architecturec                 p   [         R                  " S5      " [        5      nSn[        R                  " U[        R
                  S9S-  n[        R                  " U[        R
                  S9n[        R                  " SS5      nUSU4   " X45        U R                  [        R                  " XE:H  5      5        g )Nr   
   r^   rM   )iU  i  r   r   )
r   r`   rE   ra   rq   r   rb   tilerc   rd   )rf   rg   rh   r1   r2   rs   s         r   test_match_any_sync*TestCudaWarpOperations.test_match_any_sync   s     88678JK51A5((51gg2A6E6+w~./r   c                    [         R                  " S5      " [        5      nSn[        R                  " U[        R
                  S9n[        R                  " U[        R
                  S9nUSU4   " X45        U R                  [        R                  " US:H  5      5        SUS'   USU4   " X45        U R                  [        R                  " US:H  5      5        g )Nr   r   r^   r   i  rJ   r   )	r   r`   rH   ra   r   r   rb   rc   rd   r   s        r   test_match_all_sync*TestCudaWarpOperations.test_match_all_sync   s     88678JK%rxx0((51E6+w,678q	E6+w!|,-r   z;Independent scheduling requires at least Volta Architecturec                    [         R                  " S5      " [        5      n[        R                  " S[        R
                  S9n[        R                  " SS5      nUS   " U5        U R                  [        R                  " X#:H  5      5        g )Nr   r]   r^   )rK   rL   rN   rP      r   r]   )	r   r`   rR   ra   rb   r   r   rc   rd   )rf   rg   rQ   rs   s       r   test_independent_scheduling2TestCudaWarpOperations.test_independent_scheduling   sa     88-./IJhhr+ggFJsz*+r   c                     [         R                  S 5       n[        R                  " S[        R                  S9nUS   " U5        [        R
                  " SS5      n[        R                  R                  X25        g )Nc                     [         R                  " S5      nUS-  S:X  a  [         R                  " 5       X'   g [         R                  " 5       X'   g )Nr   rM   r   )r   r   
activemaskr?   r   s     r   use_activemask>TestCudaWarpOperations.test_activemask.<locals>.use_activemask   s:    		!AA!| ( (r   r]   r^   r   )iUUUUl   *UU rz   )r   r`   ra   r   r   r   testingassert_equal)rf   r   outexpecteds       r   test_activemask&TestCudaWarpOperations.test_activemask   s`    			) 
		) hhr+uc" 773R8


.r   c                 Z   [         R                  S 5       n[        R                  " S[        R                  S9nUS   " U5        [        R
                  " [        S5       Vs/ s H  nSU-  S-
  PM     sn[        R                  S9n[        R                  R                  XB5        g s  snf )Nc                 ^    [         R                  " S5      n[         R                  " 5       X'   g )Nr   )r   r   lanemask_ltr   s     r   use_lanemask_lt@TestCudaWarpOperations.test_lanemask_lt.<locals>.use_lanemask_lt  s    		!A##%ADr   r]   r^   r   rM   r   )	r   r`   ra   r   r   asarrayranger   r   )rf   r   r   r   r   s        r   test_lanemask_lt'TestCudaWarpOperations.test_lanemask_lt  s    		& 
	& hhr+s# ::U2Y?YQ!|Y?$&II/


. @s   "B( N)__name__
__module____qualname____firstlineno__ri   rn   rt   rw   r{   r   r   r   r   r   r   
skipUnlessrX   r   r   r   r   r   __static_attributes__r   r   r   rZ   rZ   e   s    +,,,,	0	.
..> /HJ0J0 /HJ	.J	. /(),),/*/r   rZ   __main__)numpyra   numbar   r   r   r   r   numba.cuda.testingr   r	   r
   
numba.corer   r   r   r"   r%   r)   r-   r4   r8   r;   rB   rE   rH   rR   rX   rZ   r   mainr   r   r   <module>r      s     6 6 F F '
B EFj/\ j/ Gj/Z zMMO r   