U
    9%e)                     @   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
rLe	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 )    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xZdxhx r   a/var/www/html/Darija-Ai-API/env/lib/python3.8/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   Z
null_shapeZshape1Z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 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
assertTrueZallcloselistrange)r   r   r   yzr   r   r   test_gpu_array_strided    s    
z$TestCudaArray.test_gpu_array_stridedc                 C   sb   t ddd }tjdtjd}|d dd }zt j| W n tk
rT   Y n
X tdd S )	Nzvoid(double[:], double[:])c                 S   s,   t d}|| jd k r(|| |< |||< d S r   r   )r   r.   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   r1   r   r.   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 )Nr3   )r   r4   r5   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:   Z
array_liker   r   r   _test_array_like_sameM   s    z#TestCudaArray._test_array_like_samec              
   C   s@   t jddd}tD ](}| j|d | || W 5 Q R X qd S )Nr    CorderrB   r   r   ARRAY_LIKE_FUNCTIONSsubTestrC   r   d_arB   r   r   r   test_array_like_1d[   s    z TestCudaArray.test_array_like_1dc              
   C   s@   t jddd}tD ](}| j|d | || W 5 Q R X qd S Nr       rD   rE   rG   rH   rK   r   r   r   test_array_like_2da   s    z TestCudaArray.test_array_like_2dc              
   C   s@   t jddd}tD ](}| j|d | || W 5 Q R X qd S rN   rH   rK   r   r   r   test_array_like_2d_transposeg   s    z*TestCudaArray.test_array_like_2d_transposec              
   C   s@   t jddd}tD ](}| j|d | || W 5 Q R X qd S )Nr    rP      rD   rE   rG   rH   rK   r   r   r   test_array_like_3dm   s    z TestCudaArray.test_array_like_3dc              
   C   s@   t jddd}tD ](}| j|d | || W 5 Q R X qd S )Nr    FrE   rG   rH   rK   r   r   r   test_array_like_1d_fs   s    z"TestCudaArray.test_array_like_1d_fc              
   C   s@   t jddd}tD ](}| j|d | || W 5 Q R X qd S NrO   rV   rE   rG   rH   rK   r   r   r   test_array_like_2d_fy   s    z"TestCudaArray.test_array_like_2d_fc              
   C   s@   t jddd}tD ](}| j|d | || W 5 Q R X qd S rX   rH   rK   r   r   r   test_array_like_2d_f_transpose   s    z,TestCudaArray.test_array_like_2d_f_transposec              
   C   s@   t jddd}tD ](}| j|d | || W 5 Q R X qd S )NrS   rV   rE   rG   rH   rK   r   r   r   test_array_like_3d_f   s    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.
        r>   r?   N)r   r   r"   r	   
zeros_liker@   rA   )r   rB   viewd_viewnb_likenp_liker   r   r   _test_array_like_view   s    
z#TestCudaArray._test_array_like_viewc              
   C   s`   d}t |d d d }t|d d d }tD ]*}| j|d | ||| W 5 Q R X q0d S )Nr    r3   rG   r	   zerosr   r   rI   rJ   ra   r   r   r]   r^   rB   r   r   r   test_array_like_1d_view   s    z%TestCudaArray.test_array_like_1d_viewc              
   C   sh   d}t j|ddd d d }tj|ddd d d }tD ]*}| j|d | ||| W 5 Q R X q8d S )Nr    rV   rE   r3   rG   rb   rd   r   r   r   test_array_like_1d_view_f   s    z'TestCudaArray.test_array_like_1d_view_fc              
   C   st   d}t |d d dd d df }t|d d dd d df }tD ]*}| j|d | ||| W 5 Q R X qDd S )NrO   r3   rG   rb   rd   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 5 Q R X qLd S NrO   rV   rE   r3   rG   rb   rd   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 ]r}| j|d\ ||}| |j|j | |j|j | d|j | 	|j
d  | |j
d  W 5 Q R X q(d S )NrO   r3   rG   )(      r>   r?   )r   r   TrI   rJ   r   r   r"   r@   r+   rA   ZassertFalse)r   r   r^   rB   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 ]}| j|dx t |}||}| |j	|j	 | |j
|j
 | |j|j | |jd |jd  | |jd |jd  W 5 Q R X qHd S )NrO   r3   rG   r>   r?   )r	   rc   rl   r   r   rI   rJ   r\   r   r   r"   r@   rA   )r   r   r]   r^   rB   r`   r_   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 5 Q R X qPd S rh   )r	   rc   rl   r   r   rI   rJ   ra   rd   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   r3   )r   r   )Aoutr   r   r   r   func   s    
z+TestCudaArray.test_issue_4628.<locals>.func   )r   rt   r   )	r   r(   r	   Zonesr   rc   r   lenZ	overloads)r   rs   narL   resultr   r   r   test_issue_4628   s    	

zTestCudaArray.test_issue_4628N)__name__
__module____qualname__r   r   r0   r8   r=   rC   rM   rQ   rR   rU   rW   rY   rZ   r[   ra   re   rf   rg   ri   r   rn   r   ro   rp   ry   r   r   r   r   r      s4   	

r   __main__)numpyr	   Znumba.cuda.testingr   r   r   r   Znumbar   r   ZENABLE_CUDASIMr   Zpinned_array_likerI   Zmapped_array_liker   rz   mainr   r   r   r   <module>   s    u