
    wg                     "   d dl mZ d dlZd dlmZmZmZ d dlm	Z	m
Z
mZmZmZmZ ej                  d        Zej                  d        Zej                  d        Zd Ze ed	       G d
 de
                    Zedk(  r e	j,                          yy)    )print_functionN)configcudaint32)unittestCUDATestCaseskip_on_cudasimskip_unless_cc_60skip_if_cudadevrt_missingskip_if_mvc_enabledc                 J    t         j                  j                          d| d<   y Ng      ?r   )r   cg	this_gridAs    t/home/mcse/projects/flask/flask-venv/lib/python3.12/site-packages/numba/cuda/tests/cudapy/test_cooperative_groups.pyr   r      s    GGAaD    c                 j    t         j                  j                         }|j                          d| d<   y r   )r   r   r   sync)r   gs     r   
sync_groupr      s&    AFFHAaDr   c                 4    t        j                  d      | d<   y N   r   )r   gridr   s    r   no_syncr      s    99Q<AaDr   c                    t        j                  d      }t         j                  j                         }| j                  d   }| j                  d   }t        d|      D ],  }||z
  dz
  }| |dz
  |f   dz   | ||f<   |j                          . y 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                       e Zd Zed        Ze ed      d               Zed        Ze ed      d               Z ed      d        Z	ed        Z
ed	        Zy
)TestCudaCooperativeGroupsc                     t        j                  dt         j                        }t        d   |       | j	                  t        j
                  |d         d       y Nr   
fill_valuer   r   r   zValue was not set)npfullnanr   assertFalseisnanselfr   s     r   test_this_gridz(TestCudaCooperativeGroups.test_this_grid4   sA    GGA"&&)$ 	!A$)<=r   zFSimulator doesn't differentiate between normal and cooperative kernelsc                     t        j                  dt         j                        }t        d   |       t        j                  j                         D ]   \  }}| j                  |j                         " y Nr   r,   r.   )r/   r0   r1   r   	overloadsitems
assertTruecooperativer5   r   keyoverloads       r   test_this_grid_is_cooperativez7TestCudaCooperativeGroups.test_this_grid_is_cooperative<   sZ     GGA"&&)$ '00668 	2MCOOH001	2r   c                     t        j                  dt         j                        }t        d   |       | j	                  t        j
                  |d         d       y r+   )r/   r0   r1   r   r2   r3   r4   s     r   test_sync_groupz)TestCudaCooperativeGroups.test_sync_groupG   sB    GGA"&&)4 	!A$)<=r   c                     t        j                  dt         j                        }t        d   |       t        j                  j                         D ]   \  }}| j                  |j                         " y r8   )r/   r0   r1   r   r9   r:   r;   r<   r=   s       r   test_sync_group_is_cooperativez8TestCudaCooperativeGroups.test_sync_group_is_cooperativeO   sY     GGA"&&)4'11779 	2MCOOH001	2r   z$Simulator does not implement linkingc                 F   t        j                  dt         j                        }t        d   |       t        j                  j                         D ]M  \  }}| j                  |j                         |j                  j                  D ]  }| j                  d|        O y)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)r/   r0   r1   r   r9   r:   r2   r<   _codelibrary_linking_filesassertNotIn)r5   r   r>   r?   links        r   ,test_false_cooperative_doesnt_link_cudadevrtzFTestCudaCooperativeGroups.test_false_cooperative_doesnt_link_cudadevrtY   s     GGA"&&)a$..446 	4MCX112 --<< 4  d34	4r   c                 D   t         j                  rd}nd}t        j                  |t        j                        }d}|j
                  d   |z  }t        d d d d df   f} t        j                  |      t              }|j                  |   }|j                  |      }||kD  rt        j                  d        |||f   |       t        j                  t        j                  |d         |d   df      j                  }	t        j                   j#                  ||	       y )N)    rM   )   rN   )dtyperM   r   z1GPU cannot support enough cooperative grid blocksr   )r   ENABLE_CUDASIMr/   zerosr   r   r   jitr'   r9   max_cooperative_grid_blocksr   skiptilearangeTtestingassert_equal)
r5   r   r   blockdimgriddimsigc_sequential_rowsr?   mb	references
             r   test_sync_at_matrix_rowz1TestCudaCooperativeGroups.test_sync_at_matrix_rowh   s       E EHHU"((+''!*(QssU|o)DHHSM/:$..s311(;R<MMMN,'8+,Q/GGBIIeAh/%(A?AA	


9-r   c                 0   t         d d d d df   f} t        j                  |      t              }|j                  |   }|j                  d      }|j                  d      }|j                  d      }| j                  ||       | j                  ||       y )Nr      )   rc   )rc      rd   )r   r   rR   r'   r9   rS   assertEqual)r5   r\   r]   r?   blocks1dblocks2dblocks3ds          r    test_max_cooperative_grid_blocksz:TestCudaCooperativeGroups.test_max_cooperative_grid_blocks   s     QssU|o)DHHSM/:$..s377<77A77
C8,8,r   N)__name__
__module____qualname__r
   r6   r	   r@   rB   rD   rK   r`   ri    r   r   r)   r)   1   s     > >  + ,2, 2 > >  + ,2, 2 ;<4 =4 . .0 - -r   r)   __main__)
__future__r   numpyr/   numbar   r   r   numba.cuda.testingr   r   r	   r
   r   r   rR   r   r   r   r'   r)   rj   mainrm   r   r   <module>rt      s    %  % %5 5
  

  
  
( 01\- \- 2 \-~ zHMMO r   