
    sh                     $   S SK Jr  S SKrS SKJrJrJr  S SKJ	r	J
r
JrJrJrJr  \R                  S 5       r\R                  S 5       r\R                  S 5       rS r\\" S	5       " S
 S\
5      5       5       r\S:X  a  \	R,                  " 5         gg)    )print_functionN)configcudaint32)unittestCUDATestCaseskip_on_cudasimskip_unless_cc_60skip_if_cudadevrt_missingskip_if_mvc_enabledc                 J    [         R                  R                  5         SU S'   g Ng      ?r   )r   cg	this_gridAs    ڃ/Users/tiagomarins/Projetos/claudeai/copy_bank/venv/lib/python3.13/site-packages/numba/cuda/tests/cudapy/test_cooperative_groups.pyr   r      s    GGAaD    c                 j    [         R                  R                  5       nUR                  5         SU S'   g r   )r   r   r   sync)r   gs     r   
sync_groupr      s&    AFFHAaDr   c                 6    [         R                  " S5      U S'   g N   r   )r   gridr   s    r   no_syncr      s    99Q<AaDr   c                    [         R                  " S5      n[         R                  R                  5       nU R                  S   nU R                  S   n[        SU5       H*  nXA-
  S-
  nXS-
  U4   S-   XU4'   UR                  5         M,     g r   )r   r   r   r   shaperanger   )Mcolr   rowscolsrowopposites          r   sequential_rowsr'      s}     ))A,CA771:D771:DQ~:>a)*Q.s(	 r   zCG not supported with MVCc                       \ rS rSr\S 5       r\\" S5      S 5       5       r\S 5       r\\" S5      S 5       5       r	\" S5      S 5       r
\S	 5       r\S
 5       rSrg)TestCudaCooperativeGroups1   c                     [         R                  " S[         R                  S9n[        S   " U5        U R	                  [         R
                  " US   5      S5        g Nr   
fill_valuer   r   r   zValue was not set)npfullnanr   assertFalseisnanselfr   s     r   test_this_grid(TestCudaCooperativeGroups.test_this_grid4   sA    GGA"&&)$ 	!A$)<=r   zFSimulator doesn't differentiate between normal and cooperative kernelsc                     [         R                  " S[         R                  S9n[        S   " U5        [        R                  R                  5        H   u  p#U R                  UR                  5        M"     g Nr   r-   r/   )r0   r1   r2   r   	overloadsitems
assertTruecooperativer6   r   keyoverloads       r   test_this_grid_is_cooperative7TestCudaCooperativeGroups.test_this_grid_is_cooperative<   sS     GGA"&&)$ '00668MCOOH001 9r   c                     [         R                  " S[         R                  S9n[        S   " U5        U R	                  [         R
                  " US   5      S5        g r,   )r0   r1   r2   r   r3   r4   r5   s     r   test_sync_group)TestCudaCooperativeGroups.test_sync_groupG   sB    GGA"&&)4 	!A$)<=r   c                     [         R                  " S[         R                  S9n[        S   " U5        [        R                  R                  5        H   u  p#U R                  UR                  5        M"     g r:   )r0   r1   r2   r   r;   r<   r=   r>   r?   s       r   test_sync_group_is_cooperative8TestCudaCooperativeGroups.test_sync_group_is_cooperativeO   sR     GGA"&&)4'11779MCOOH001 :r   z$Simulator does not implement linkingc                 J   [         R                  " S[         R                  S9n[        S   " U5        [        R                  R                  5        HO  u  p#U R                  UR                  5        UR                  R                   H  nU R                  SU5        M     MQ     g)z
We should only mark a kernel as cooperative and link cudadevrt if the
kernel uses grid sync. Here we ensure that one that doesn't use grid
synsync isn't marked as such.
r   r-   r/   	cudadevrtN)r0   r1   r2   r   r;   r<   r3   r>   _codelibrary_linking_filesassertNotIn)r6   r   r@   rA   links        r   ,test_false_cooperative_doesnt_link_cudadevrtFTestCudaCooperativeGroups.test_false_cooperative_doesnt_link_cudadevrtY   sy     GGA"&&)a$..446MCX112 --<<  d3 = 7r   c                 N   [         R                  (       a  SnOSn[        R                  " U[        R                  S9nSnUR
                  S   U-  n[        S S 2S S S24   4n[        R                  " U5      " [        5      nUR                  U   nUR                  U5      nXH:  a  [        R                  " S5        XdU4   " U5        [        R                  " [        R                  " US   5      US   S45      R                  n	[        R                   R#                  X)5        g )N)    rS   )   rT   )dtyperS   r   z1GPU cannot support enough cooperative grid blocksr   )r   ENABLE_CUDASIMr0   zerosr   r   r   jitr'   r;   max_cooperative_grid_blocksr   skiptilearangeTtestingassert_equal)
r6   r   r   blockdimgriddimsigc_sequential_rowsrA   mb	references
             r   test_sync_at_matrix_row1TestCudaCooperativeGroups.test_sync_at_matrix_rowh   s       E EHHU"((+''!*(QssU|o HHSM/:$..s311(;<MMMN8+,Q/GGBIIeAh/%(A?AA	


-r   c                 .   [         S S 2S S S24   4n[        R                  " U5      " [        5      nUR                  U   nUR                  S5      nUR                  S5      nUR                  S5      nU R                  XE5        U R                  XF5        g )Nr      )   rj   )rj      rk   )r   r   rX   r'   r;   rY   assertEqual)r6   rb   rc   rA   blocks1dblocks2dblocks3ds          r    test_max_cooperative_grid_blocks:TestCudaCooperativeGroups.test_max_cooperative_grid_blocks   s     QssU|o HHSM/:$..s377<77A77
C,,r    N)__name__
__module____qualname____firstlineno__r
   r7   r	   rB   rE   rH   rP   rf   rp   __static_attributes__rr   r   r   r)   r)   1   s     > >  + ,2, 2 > >  + ,2, 2 ;<4 =4 . .0 - -r   r)   __main__)
__future__r   numpyr0   numbar   r   r   numba.cuda.testingr   r   r	   r
   r   r   rX   r   r   r   r'   r)   rs   mainrr   r   r   <module>r~      s    %  % %5 5
  

  
  
( 01\- \- 2 \-~ zMMO r   