o
    8ήc
                     @   sX   d dl Z d dlmZmZmZmZ eeedG dd deZedkr*e   dS dS )    N)CUDATestCaseskip_on_cudasimskip_if_cudadevrt_missingskip_unless_cc_60z4cudasim doesn't support cuda import at non-top-levelc                   @   s   e Zd Zdd ZdS )TestCooperativeGroupsc           
         s   ddl m m} dd l}|d d d d df f} | fdd}|jd|jd}d}|jd | }|j| |}||k rE| 	d	 |||f | |
|d
dj}	|j||	 d S )Nr   )cudaint32   c                    sl     d} j }| jd }| jd }td|D ]}|| d }| |d |f d | ||f< |  qd S )Nr	   r   )gridcg	this_gridshaperangesync)Mcolgrowscolsrowoppositer    L/tmp/pip-target-vg8gfxp4/lib/python/numba/cuda/tests/doc_examples/test_cg.pysequential_rows   s   




z@TestCooperativeGroups.test_ex_grid_sync.<locals>.sequential_rows)   r   )dtype    z0Device does not support a large enough coop gridr   )r   r	   )numbar   r   numpyjitzerosr   	overloadsmax_cooperative_grid_blocksskipTesttilearangeTtestingassert_equal)
selfr   npsigr   Ablockdimgriddimmb	referencer   r   r   test_ex_grid_sync   s   
z'TestCooperativeGroups.test_ex_grid_syncN)__name__
__module____qualname__r2   r   r   r   r   r   	   s    r   __main__)	unittestnumba.cuda.testingr   r   r   r   r   r3   mainr   r   r   r   <module>   s   >