o
    i)                     @   s   d dl Zd dlmZmZ d dlmZmZ d dlmZm	Z	 ej
r&e	je	jfZne	je	je	jfZG dd deZedkr@e  dS dS )    N)unittestCUDATestCase)skip_on_cudasimskip_unless_cudasim)configcudac                   @   s   e Zd Zdd Zdd Zdd Zdd Zd	d
 Zdd Zdd Z	dd Z
dd Zdd Zdd Zdd Zdd Zdd Zdd Zdd  Zd!d" Zd#d$ Zd%d& Zed'd(d) Zed'd*d+ Zd,d- Zed.d/d0 Zd1S )2TestCudaArrayc                 C   s`   t d}t|}| }| |j|j | |j|j | |j|j | |j|j d S )Nr   )nparanger   	to_devicecopy_to_hostassertEqualshapesize)selfxdxhx r   _/home/ubuntu/veenaModal/venv/lib/python3.10/site-packages/numba/cuda/tests/cudapy/test_array.pytest_gpu_array_zero_length   s   

z(TestCudaArray.test_gpu_array_zero_lengthc                 C   s>   d}t dj}t tdj}| || | || d S )Nr   )r   device_arrayr   device_array_liker	   ndarrayr   )r   
null_shapeshape1shape2r   r   r   test_null_shape   s
   zTestCudaArray.test_null_shapec                 C   sx   t ddd }tjdtjd}tjd|tjd}tjd|d	d
 tjd}|d | | t|t	t
d d S )Nzvoid(double[:])c                 S   s(   t d}|| jd k r|| |< d S d S N   r   r   gridr   )r   ir   r   r   kernel"   s   
z4TestCudaArray.test_gpu_array_strided.<locals>.kernel
   dtypeP   )r   bufferr&   	      )r(   r&   )r$   r$   )r   jitr	   r
   doubler   byte
assertTrueallcloselistrange)r   r#   r   yzr   r   r   test_gpu_array_strided    s   
z$TestCudaArray.test_gpu_array_stridedc                 C   s\   t ddd }tjdtjd}|d dd }zt j| W td ty-   Y d S w )	Nzvoid(double[:], double[:])c                 S   s0   t d}|| jd k r|| |< |||< d S d S r   r    )r   r3   r"   r   r   r   
copykernel0   s
   
z<TestCudaArray.test_gpu_array_interleaved.<locals>.copykernelr$   r%      zDShould raise exception complaining the contiguous-ness of the array.)	r   r,   r	   r
   r-   devicearrayauto_device
ValueErrorAssertionError)r   r6   r   r3   r   r   r   test_gpu_array_interleaved.   s   
z(TestCudaArray.test_gpu_array_interleavedc                 C   s2   t jd\}}| t| tdk d S )Nr8   )r   r9   r:   r/   r	   allr   array)r   d_r   r   r   test_auto_device_constI   s   "z$TestCudaArray.test_auto_device_constc                 C   sl   ||}|  |j|j |  |j|j |  |j|j |  |jd |jd  |  |jd |jd  dS )zk
        Tests of *_array_like where shape, strides, dtype, and flags should
        all be equal.
        C_CONTIGUOUSF_CONTIGUOUSN)r   r   stridesr&   flags)r   	like_funcr?   
array_liker   r   r   _test_array_like_sameM   s   z#TestCudaArray._test_array_like_samec              	   C   T   t jddd}tD ]}| j|d | || W d    n1 s"w   Y  q	d S )Nr$   CorderrG   r   r   ARRAY_LIKE_FUNCTIONSsubTestrI   r   d_arG   r   r   r   test_array_like_1d[      z TestCudaArray.test_array_like_1dc              	   C   rJ   Nr$      rK   rL   rN   rO   rR   r   r   r   test_array_like_2da   rU   z TestCudaArray.test_array_like_2dc              	   C   rJ   rV   rO   rR   r   r   r   test_array_like_2d_transposeg   rU   z*TestCudaArray.test_array_like_2d_transposec              	   C   rJ   )Nr$   rX      rK   rL   rN   rO   rR   r   r   r   test_array_like_3dm   rU   z TestCudaArray.test_array_like_3dc              	   C   rJ   )Nr$   FrL   rN   rO   rR   r   r   r   test_array_like_1d_fs   rU   z"TestCudaArray.test_array_like_1d_fc              	   C   rJ   NrW   r^   rL   rN   rO   rR   r   r   r   test_array_like_2d_fy   rU   z"TestCudaArray.test_array_like_2d_fc              	   C   rJ   r`   rO   rR   r   r   r   test_array_like_2d_f_transpose   rU   z,TestCudaArray.test_array_like_2d_f_transposec              	   C   rJ   )Nr[   r^   rL   rN   rO   rR   r   r   r   test_array_like_3d_f   rU   z"TestCudaArray.test_array_like_3d_fc                 C   sv   ||}|  |j|j |  |j|j t|}|  |j|j |  |jd |jd  |  |jd |jd  dS )z
        Tests of device_array_like where the original array is a view - the
        strides should not be equal because a contiguous array is expected.
        rC   rD   N)r   r   r&   r	   
zeros_likerE   rF   )r   rG   viewd_viewnb_likenp_liker   r   r   _test_array_like_view   s   
z#TestCudaArray._test_array_like_viewc              	   C   st   d}t |d d d }t|d d d }tD ]}| j|d | ||| W d    n1 s2w   Y  qd S )Nr$   r8   rN   r	   zerosr   r   rP   rQ   ri   r   r   re   rf   rG   r   r   r   test_array_like_1d_view   s   z%TestCudaArray.test_array_like_1d_viewc              	   C   s|   d}t j|ddd d d }tj|ddd d d }tD ]}| j|d | ||| W d    n1 s6w   Y  qd S )Nr$   r^   rL   r8   rN   rj   rl   r   r   r   test_array_like_1d_view_f   s   z'TestCudaArray.test_array_like_1d_view_fc              	   C   s   d}t |d d dd d df }t|d d dd d df }tD ]}| j|d | ||| W d    n1 s<w   Y  q"d S )NrW   r8   rN   rj   rl   r   r   r   test_array_like_2d_view   s   z%TestCudaArray.test_array_like_2d_viewc              	   C   s   d}t j|ddd d dd d df }tj|ddd d dd d df }tD ]}| j|d | ||| W d    n1 s@w   Y  q&d S NrW   r^   rL   r8   rN   rj   rl   r   r   r   test_array_like_2d_view_f   s   ""z'TestCudaArray.test_array_like_2d_view_fz5Numba and NumPy stride semantics differ for transposec              	   C   s   d}t |d d dd d df j}tD ]C}| j|d3 ||}| |j|j | |j|j | d|j | 	|j
d  | |j
d  W d    n1 sRw   Y  qd S )NrW   r8   rN   )(      rC   rD   )r   r   TrP   rQ   r   r   r&   rE   r/   rF   assertFalse)r   r   rf   rG   liker   r   r   (test_array_like_2d_view_transpose_device   s    z6TestCudaArray.test_array_like_2d_view_transpose_devicec              	   C   s   d}t |d d dd d df j}t|d d dd d df j}tD ]Q}| j|dA t |}||}| |j	|j	 | |j
|j
 | |j|j | |jd |jd  | |jd |jd  W d    n1 spw   Y  q$d S )NrW   r8   rN   rC   rD   )r	   rk   rt   r   r   rP   rQ   rd   r   r   r&   rE   rF   )r   r   re   rf   rG   rh   rg   r   r   r   +test_array_like_2d_view_transpose_simulator   s&     
z9TestCudaArray.test_array_like_2d_view_transpose_simulatorc              	   C   s   d}t j|ddd d dd d df j}tj|ddd d dd d df j}tD ]}| j|d | ||| W d    n1 sBw   Y  q(d S rp   )r	   rk   rt   r   r   rP   rQ   ri   rl   r   r   r   #test_array_like_2d_view_f_transpose   s   $$z1TestCudaArray.test_array_like_2d_view_f_transposez-Kernel overloads not created in the simulatorc                 C   sf   t jdd }d}t|f}t |}t|f}|d || |d || | dt|j d S )Nc                 S   s   t d}| | d ||< d S )Nr   r8   )r   r!   )Aoutr"   r   r   r   func   s   
z+TestCudaArray.test_issue_4628.<locals>.func   )r   r}   r   )	r   r,   r	   onesr   rk   r   len	overloads)r   r|   narS   resultr   r   r   test_issue_4628   s   	

zTestCudaArray.test_issue_4628N)__name__
__module____qualname__r   r   r5   r=   rB   rI   rT   rY   rZ   r]   r_   ra   rb   rc   ri   rm   rn   ro   rq   r   rw   r   rx   ry   r   r   r   r   r   r      s6    	

r   __main__)numpyr	   numba.cuda.testingr   r   r   r   numbar   r   ENABLE_CUDASIMr   pinned_array_likerP   mapped_array_liker   r   mainr   r   r   r   <module>   s     u