
    `i                     0   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                     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                 L    t           j                                         d| d<   d S Ng      ?r   )r   cg	this_gridAs    /home/jaya/work/projects/VOICE-AGENT/VIET/agent-env/lib/python3.11/site-packages/numba/cuda/tests/cudapy/test_cooperative_groups.pyr   r      s#    GAaDDD    c                 t    t           j                                        }|                                 d| d<   d S r   )r   r   r   sync)r   gs     r   
sync_groupr      s/    AFFHHHAaDDDr   c                 4    t          j        d          | d<   d S N   r   )r   gridr   s    r   no_syncr      s    9Q<<AaDDDr   c                 &   t          j        d          }t           j                                        }| j        d         }| j        d         }t          d|          D ]3}||z
  dz
  }| |dz
  |f         dz   | ||f<   |                                 4d S r   )r   r   r   r   shaperanger   )Mcolr   rowscolsrowopposites          r   sequential_rowsr'      s     )A,,CA71:D71: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d
S )TestCudaCooperativeGroupsc                     t          j        dt           j                  }t          d         |           |                     t          j        |d                   d           d S 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   sU    GA"&)))$ 	!A$)<=====r   zFSimulator doesn't differentiate between normal and cooperative kernelsc                     t          j        dt           j                  }t          d         |           t          j                                        D ]\  }}|                     |j                    d S 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<   sr     GA"&)))$ '06688 	2 	2MCOOH01111	2 	2r   c                     t          j        dt           j                  }t          d         |           |                     t          j        |d                   d           d S r+   )r/   r0   r1   r   r2   r3   r4   s     r   test_sync_groupz)TestCudaCooperativeGroups.test_sync_groupG   sV    GA"&)))4 	!A$)<=====r   c                     t          j        dt           j                  }t          d         |           t          j                                        D ]\  }}|                     |j                    d S r8   )r/   r0   r1   r   r9   r:   r;   r<   r=   s       r   test_sync_group_is_cooperativez8TestCudaCooperativeGroups.test_sync_group_is_cooperativeO   sq     GA"&)))4'17799 	2 	2MCOOH01111	2 	2r   z$Simulator does not implement linkingc                 6   t          j        dt           j                  }t          d         |           t          j                                        D ]D\  }}|                     |j                   |j        j	        D ]}| 
                    d|           Ed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/   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     GA"&)))a$.4466 	4 	4MCX1222 -< 4 4  d33334	4 	4r   c                 L   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	        |         }|
                    |          }||k    rt          j        d            |||f         |           t          j        t          j        |d                   |d         df          j        }	t          j                            ||	           d S )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     	! EE EHU"(+++'!*(QQQsssU|o)DHSMM/::$.s311(;;R<<MMNNN,'8+,Q///GBIeAh//%(A??A	

9-----r   c                 d   t           d d d d df         f} t          j        |          t                    }|j        |         }|                    d          }|                    d          }|                    d          }|                     ||           |                     ||           d S )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     QQQsssU|o)DHSMM/::$.s377<<77AA77
C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 2, , 2 _;<<4 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      st   % % % % % %     % % % % % % % % % %5 5 5 5 5 5 5 5 5 5 5 5 5 5 5 5
   

   
   
  ( 011\- \- \- \- \- \- \- 21 \-~ zHMOOOOO r   