U
    hâËd)  ã                   @   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   ú[/home/sam/Atlas/atlas_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é
   ©ÚdtypeéP   )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