o
    ۾i                     @   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mZ ejdd Zejdd Zejd	d
 Zdd ZeedG dd de
ZedkrUe	  dS dS )    )print_functionN)configcudaint32)unittestCUDATestCaseskip_on_cudasimskip_unless_cc_60skip_if_cudadevrt_missingskip_if_mvc_enabledc                 C   s   t j  d| d< d S Ng      ?r   )r   cg	this_gridA r   c/home/ubuntu/.local/lib/python3.10/site-packages/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"   zCG not supported with MVCc                   @   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_grid4      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_groupG   r1   z)TestCudaCooperativeGroups.test_sync_groupc                 C   r2   r3   )r)   r*   r+   r   r4   r5   r6   r7   r8   r   r   r   test_sync_group_is_cooperativeO   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   r4   r5   r,   r7   _codelibrary_linking_filesassertNotIn)r/   r   r9   r:   linkr   r   r   ,test_false_cooperative_doesnt_link_cudadevrtY   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)    rD   )   rE   )dtyperD   r   z1GPU cannot support enough cooperative grid blocksr   )r   ENABLE_CUDASIMr)   zerosr   r   r   jitr"   r4   max_cooperative_grid_blocksr   skiptilearangeTtestingassert_equal)
r/   r   r   blockdimgriddimsigc_sequential_rowsr:   mb	referencer   r   r   test_sync_at_matrix_rowh   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      )   rY   )rY      rZ   )r   r   rI   r"   r4   rJ   assertEqual)r/   rS   rT   r:   blocks1dblocks2dblocks3dr   r   r    test_max_cooperative_grid_blocks   s   



z:TestCudaCooperativeGroups.test_max_cooperative_grid_blocksN)__name__
__module____qualname__r	   r0   r   r;   r<   r=   rC   rW   r_   r   r   r   r   r#   1   s"    



r#   __main__)
__future__r   numpyr)   numbar   r   r   numba.cuda.testingr   r   r   r	   r
   r   rI   r   r   r   r"   r#   r`   mainr   r   r   r   <module>   s"     


_