o
    8ήcH                     @   s   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 ejdd Zejdd Zejd	d
 Zdd ZeG dd de
ZedkrOe	  dS dS )    )print_functionN)configcudaint32)unittestCUDATestCaseskip_on_cudasimskip_unless_cc_60skip_if_cudadevrt_missingc                 C   s   t j  d| d< d S Ng      ?r   )r   cg	this_gridA r   V/tmp/pip-target-vg8gfxp4/lib/python/numba/cuda/tests/cudapy/test_cooperative_groups.pyr   
   s   
r   c                 C   s   t j }|  d| d< d S r   )r   r   r   sync)r   gr   r   r   
sync_group   s   
r   c                 C   s   t d| d< d S N   r   )r   gridr   r   r   r   no_sync   s   r   c                 C   sl   t d}t j }| jd }| jd }td|D ]}|| d }| |d |f d | ||f< |  qd S r   )r   r   r   r   shaperanger   )Mcolr   rowscolsrowoppositer   r   r   sequential_rows   s   




r!   c                   @   st   e Zd Zedd Zeeddd Zedd Zeeddd	 Zed
dd Z	edd Z
edd ZdS )TestCudaCooperativeGroupsc                 C   6   t jdt jd}td | | t |d d d S Nr   
fill_valuer   r   r   zValue was not set)npfullnanr   assertFalseisnanselfr   r   r   r   test_this_grid2      z(TestCudaCooperativeGroups.test_this_gridzFSimulator doesn't differentiate between normal and cooperative kernelsc                 C   @   t jdt jd}td | tj D ]
\}}| |j qd S Nr   r%   r'   )r(   r)   r*   r   	overloadsitems
assertTruecooperativer.   r   keyoverloadr   r   r   test_this_grid_is_cooperative:   s
   z7TestCudaCooperativeGroups.test_this_grid_is_cooperativec                 C   r#   r$   )r(   r)   r*   r   r+   r,   r-   r   r   r   test_sync_groupE   r0   z)TestCudaCooperativeGroups.test_sync_groupc                 C   r1   r2   )r(   r)   r*   r   r3   r4   r5   r6   r7   r   r   r   test_sync_group_is_cooperativeM   s
   z8TestCudaCooperativeGroups.test_sync_group_is_cooperativez$Simulator does not implement linkingc                 C   sZ   t jdt jd}td | tj D ]\}}| |j |jj	D ]}| 
d| q!qdS )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(   r)   r*   r   r3   r4   r+   r6   _codelibrary_linking_filesassertNotIn)r.   r   r8   r9   linkr   r   r   ,test_false_cooperative_doesnt_link_cudadevrtW   s   zFTestCudaCooperativeGroups.test_false_cooperative_doesnt_link_cudadevrtc           
      C   s   t jrd}nd}tj|tjd}d}|jd | }td d d d df f}t|t}|j	| }|
|}||kr?td |||f | tt|d |d dfj}	tj||	 d S )N)    rC   )   rD   )dtyperC   r   z1GPU cannot support enough cooperative grid blocksr   )r   ENABLE_CUDASIMr(   zerosr   r   r   jitr!   r3   max_cooperative_grid_blocksr   skiptilearangeTtestingassert_equal)
r.   r   r   blockdimgriddimsigc_sequential_rowsr9   mb	referencer   r   r   test_sync_at_matrix_rowf   s   


 z1TestCudaCooperativeGroups.test_sync_at_matrix_rowc                 C   sj   t d d d d df f}t|t}|j| }|d}|d}|d}| || | || d S )Nr      )   rX   )rX      rY   )r   r   rH   r!   r3   rI   assertEqual)r.   rR   rS   r9   blocks1dblocks2dblocks3dr   r   r    test_max_cooperative_grid_blocks   s   



z:TestCudaCooperativeGroups.test_max_cooperative_grid_blocksN)__name__
__module____qualname__r	   r/   r   r:   r;   r<   rB   rV   r^   r   r   r   r   r"   0   s"    



r"   __main__)
__future__r   numpyr(   numbar   r   r   numba.cuda.testingr   r   r   r	   r
   rH   r   r   r   r!   r"   r_   mainr   r   r   r   <module>   s     


_