U
    ?h¶O  ã                   @   sˆ   d dl Z d dlZd dlmZ d dlmZ d dlmZm	Z	 d dlm
Z
 G dd„ de	ƒZG dd	„ d	e	ƒZG d
d„ de	ƒZedkr„e ¡  dS )é    N)Údevicearray)Úcuda)ÚunittestÚCUDATestCase)Úskip_on_cudasimc                   @   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edƒ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d0d1„ Zd2d3„ Zd4d5„ Zd6d7„ Zd8d9„ Zd:d;„ Z d<d=„ Z!ed>ƒd?d@„ ƒZ"dAdB„ Z#edCƒdDdE„ ƒZ$edCƒdFdG„ ƒZ%edCƒdHdI„ ƒZ&edCƒdJdK„ ƒZ'edCƒdLdM„ ƒZ(edCƒdNdO„ ƒZ)edCƒdPdQ„ ƒZ*edCƒdRdS„ ƒZ+edCƒdTdU„ ƒZ,dVdW„ Z-ed>ƒdXdY„ ƒZ.dZS )[ÚTestCudaNDArrayc                 C   sd   t jdd}t |¡ t d¡}t  |¡}t |¡ t d¡}t  |¡}|  |j	d¡ t |¡ d S )Néd   )ÚshapegX9´Èv¾ó?r   )
r   Údevice_arrayr   Zverify_cuda_ndarray_interfaceÚnpÚemptyÚ	to_deviceÚasarrayÚassertEqualsÚndim)ÚselfÚdaryÚary© r   ú\/var/www/html/venv/lib/python3.8/site-packages/numba/cuda/tests/cudadrv/test_cuda_ndarray.pyÚtest_device_array_interface
   s    





z+TestCudaNDArray.test_device_array_interfacec                 C   sJ   t jdt jd}d|j_|  |jj¡ t |¡}| ¡ }t j	 
||¡ d S )Nr   ©ÚdtypeF)r   ÚarangeÚfloat32ÚflagsZ	writeableZassertFalser   r   Úcopy_to_hostÚtestingÚassert_array_equal)r   r   r   Zretrr   r   r   Útest_device_array_from_readonly   s    
z/TestCudaNDArray.test_device_array_from_readonlyc                 C   s&   t jddd}|  |jt d¡¡ d S )N)r   Zf4)r	   r   )r   r
   ÚassertEqualr   r   )r   r   r   r   r   Útest_devicearray_dtype!   s    z&TestCudaNDArray.test_devicearray_dtypec                 C   s"   t jdt jd}tj|dd d S )Nr   r   F)Úcopy)r   r   r   r   r   )r   Úarrayr   r   r   Útest_devicearray_no_copy%   s    z(TestCudaNDArray.test_devicearray_no_copyc                 C   sR   t  d¡ ddd¡}t |¡}|  |j|j¡ |  |jdd … |jdd … ¡ d S )Né   é   é   é   é   )r   r   Úreshaper   r   r   r	   ©r   r   r   r   r   r   Útest_devicearray_shape)   s    
z&TestCudaNDArray.test_devicearray_shapec                 C   sJ   t jdt jd}| ¡ }t |¡}d|d d …< | |¡ t j ||¡ d S )Nr   r   r   ©	r   r   Úint32r"   r   r   r   r   r   )r   r#   ÚoriginalÚgpumemr   r   r   Útest_devicearray/   s    

z TestCudaNDArray.test_devicearrayc              	   C   sT   t  ¡ }| ¡ : t jdtj|d}|  | |¡j|¡ |  |j|¡ W 5 Q R X d S )N)r'   r'   )r   Ústream)r   r2   Zauto_synchronizer
   r   Úfloat64r    Úbind)r   r2   Úarrr   r   r   Útest_stream_bind8   s    
ýz TestCudaNDArray.test_stream_bindc                 C   s,   t  d¡}t d¡}|  t|ƒt|ƒ¡ d S )N)r'   r'   ©r   r   r   r
   r    Úlenr+   r   r   r   Útest_len_1dB   s    

zTestCudaNDArray.test_len_1dc                 C   s,   t  d¡}t d¡}|  t|ƒt|ƒ¡ d S )N)r'   é   r7   r+   r   r   r   Útest_len_2dG   s    

zTestCudaNDArray.test_len_2dc                 C   s,   t  d¡}t d¡}|  t|ƒt|ƒ¡ d S )N)r'   r:   é   r7   r+   r   r   r   Útest_len_3dL   s    

zTestCudaNDArray.test_len_3dc                 C   sœ   d}t j|t jd}| ¡ }t |¡}| |d ¡\}}d|d d …< |  t  |dk¡¡ | 	||d d … ¡ | 	|d |d … ¡ |  t  ||k¡¡ d S )Nr   r   r&   r   )
r   r   r.   r"   r   r   ÚsplitÚ
assertTrueÚallr   )r   ÚNr#   r/   r0   ÚleftÚrightr   r   r   Útest_devicearray_partitionQ   s    
z*TestCudaNDArray.test_devicearray_partitionc                 C   sX   d}t j|t jd}| ¡ }t |¡}tj|d |d | |¡ t j ||d ¡ d S )Nr   r   r&   )Útor-   )r   rA   r#   r/   r0   r   r   r   Útest_devicearray_replacea   s    

z(TestCudaNDArray.test_devicearray_replacezThis works in the simulatorc              	   C   sV   t  t t d¡¡ ddd¡¡}|  t¡}t |¡ W 5 Q R X |  	dt
|jƒ¡ d S )Né   r'   r(   r)   z2transposing a non-2D DeviceNDArray isn't supported)r   r   r   r#   r   r*   ÚassertRaisesÚNotImplementedErrorÚ	transposer    ÚstrÚ	exception©r   r0   Úer   r   r   Ú#test_devicearray_transpose_wrongdimj   s     þz3TestCudaNDArray.test_devicearray_transpose_wrongdimc                 C   sJ   t  t  d¡¡ ddd¡}t jt |¡dd ¡ }|  t  	||k¡¡ d S )Nr%   r'   r(   r&   )r   r)   r&   ©Zaxes)
r   r#   r   r*   rJ   r   r   r   r?   r@   ©r   r/   r#   r   r   r   Ú#test_devicearray_transpose_identityu   s
    ÿ
z3TestCudaNDArray.test_devicearray_transpose_identityc              	   C   s^   t  t t d¡¡ dd¡¡}|  t¡}tj|dd W 5 Q R X | j	t
|jƒddgd d S )	NrG   r'   r(   )r   r   rP   zinvalid axes list (0, 0)zrepeated axis in transpose©Ú	container©r   r   r   r#   r   r*   rH   Ú
ValueErrorrJ   ÚassertInrK   rL   rM   r   r   r   Ú)test_devicearray_transpose_duplicatedaxis|   s    þþz9TestCudaNDArray.test_devicearray_transpose_duplicatedaxisc              	   C   s`   t  t t d¡¡ dd¡¡}|  t¡}tj|dd W 5 Q R X | j	t
|jƒdddgd	 d S )
NrG   r'   r(   )r   r&   rP   zinvalid axes list (0, 2)zinvalid axis for this arrayz0axis 2 is out of bounds for array of dimension 2rS   rU   rM   r   r   r   Ú$test_devicearray_transpose_wrongaxis‰   s    ýþz4TestCudaNDArray.test_devicearray_transpose_wrongaxisc              
   C   sj   t jt  d¡dd dd¡}t |¡}dD ]:}| j|d$ t j | 	|¡ 
¡ | 	|¡¡ W 5 Q R X q*d S )NrG   Úi2r   r'   r(   )Úi4Zu4Úi8Zf8)r   r#   r   r*   r   r   ZsubTestr   r   Úviewr   )r   r/   r#   r   r   r   r   Útest_devicearray_view_ok—   s    
þz(TestCudaNDArray.test_devicearray_view_okc                 C   sp   t jt  d¡dd dd¡}t |¡d d …d d d…f }|d d …d d d…f }t j | d¡ 	¡ | d¡¡ d S )Né    rZ   r   r(   é   r&   Úu2)
r   r#   r   r*   r   r   r   r   r]   r   rQ   r   r   r   Ú%test_devicearray_view_ok_not_c_contig¡   s    þz5TestCudaNDArray.test_devicearray_view_ok_not_c_contigc              	   C   s’   t jt  d¡dd dd¡}t |¡d d …d d d…f }|  t¡}| d¡ W 5 Q R X t	|j
ƒ}|  d|¡ d	|k}d
|k}|  |pˆ|d¡ d S )Nr_   rZ   r   r(   r`   r&   r[   z)To change to a dtype of a different size,zthe array must be C-contiguousz the last axis must be contiguousz&Expected message to mention contiguity)r   r#   r   r*   r   r   rH   rV   r]   rK   rL   rW   r?   )r   r/   r#   rN   ÚmsgZcontiguous_pre_np123Zcontiguous_post_np123r   r   r   Ú&test_devicearray_view_bad_not_c_contigª   s    

ÿz6TestCudaNDArray.test_devicearray_view_bad_not_c_contigc              	   C   s\   t jt  d¡dd dd¡}t |¡}|  t¡}| d¡ W 5 Q R X |  	dt
|jƒ¡ d S )NrG   rZ   r   r(   r'   r[   zuWhen changing to a larger dtype, its size must be a divisor of the total size in bytes of the last axis of the array.)r   r#   r   r*   r   r   rH   rV   r]   r    rK   rL   )r   r/   r#   rN   r   r   r   Ú"test_devicearray_view_bad_itemsize¸   s    
üz2TestCudaNDArray.test_devicearray_view_bad_itemsizec                 C   sF   t  t  d¡¡ dd¡}t  t |¡¡ ¡ }|  t  	||j
k¡¡ d S ©NrG   r'   r(   )r   r#   r   r*   rJ   r   r   r   r?   r@   ÚTrQ   r   r   r   Útest_devicearray_transpose_okÃ   s    z-TestCudaNDArray.test_devicearray_transpose_okc                 C   sB   t  t  d¡¡ dd¡}t |¡j ¡ }|  t  	||jk¡¡ d S rf   )
r   r#   r   r*   r   r   rg   r   r?   r@   rQ   r   r   r   Útest_devicearray_transpose_TÈ   s    z,TestCudaNDArray.test_devicearray_transpose_Tc              	   C   sr   t  d¡jdddd}t jddd}t |¡}||d< |  t¡}|d  |¡ W 5 Q R X |  	t
jt|jƒ¡ d S )Né   r:   ÚF©Úorder)r:   )Z
fill_valuer	   r&   )r   r   r*   Úfullr   r   rH   rV   Úcopy_to_devicer    r   Úerrmsg_contiguous_bufferrK   rL   )r   ÚaÚsÚdrN   r   r   r   Ú!test_devicearray_contiguous_sliceÍ   s    
þz1TestCudaNDArray.test_devicearray_contiguous_slicec                 C   s¶   |   |jj¡ |   |jj¡ ||f||f||f||ffD ]x\}}d|jjrNdnd|jjr\dndf }t |¡}| |¡ | j t | 	¡ |k¡|d | j t | 	¡ |k¡|d q8dS )z-
        Checks host->device memcpys
        z%s => %sÚCrk   )rc   N)
r?   r   Úc_contiguousÚf_contiguousr   r   ro   r   r@   r   )r   Úa_cÚa_fr/   r"   rc   rs   r   r   r   Ú&_test_devicearray_contiguous_host_copyá   s    üþ

z6TestCudaNDArray._test_devicearray_contiguous_host_copyc                 C   s2   t  d¡ ddd¡}t j|dd}|  ||¡ d S )Né}   r:   rk   rl   )r   r   r*   r#   rz   ©r   rx   ry   r   r   r   Ú(test_devicearray_contiguous_copy_host_3dø   s    z8TestCudaNDArray.test_devicearray_contiguous_copy_host_3dc                 C   s(   t  d¡}t j|dd}|  ||¡ d S )Nr:   rk   rl   )r   r   r#   rz   r|   r   r   r   Ú(test_devicearray_contiguous_copy_host_1dý   s    
z8TestCudaNDArray.test_devicearray_contiguous_copy_host_1dc              	   C   s.  t  d¡ ddd¡}t j|dd}|  |jj¡ |  |jj¡ t 	|¡}|  
t¡}| t 	|¡¡ W 5 Q R X |  d |j|j¡t|jƒ¡ | t 	|¡¡ |  t  | ¡ |k¡¡ t 	|¡}|  
t¡}| t 	|¡¡ W 5 Q R X |  d |j|j¡t|jƒ¡ | t 	|¡¡ |  t  | ¡ |k¡¡ d S )Nr{   r:   rk   rl   zincompatible strides: {} vs. {})r   r   r*   r#   r?   r   rv   rw   r   r   rH   rV   ro   r    ÚformatÚstridesrK   rL   r@   r   )r   rx   ry   rs   rN   r   r   r   Ú'test_devicearray_contiguous_copy_device  s,    
þ
þz7TestCudaNDArray.test_devicearray_contiguous_copy_devicec                 C   s  d}d}t  |¡}t  |¡j|dd}t  |¡j|dd}tt|ƒƒD ]È}td ƒf| t jf }|d |… |f ||d …  }t  || |¡}	t  || |¡}
t	 
|	¡}t	 
|
¡}t j | ¡ |	¡ t j | ¡ |
¡ | |
¡ | |	¡ t j | ¡ |
¡ t j | ¡ |	¡ qFd S )Nr(   )r&   r'   ru   rl   rk   )r   Úprodr   r*   Úranger8   ÚsliceZnewaxisÚbroadcast_tor   r   r   r   r   ro   )r   Z	broadsizeZ	coreshapeZcoresizeZcore_cZcore_fÚdimZnewindexZ
broadshapeZbroad_cZbroad_fZdbroad_cZdbroad_fr   r   r   Ú$test_devicearray_broadcast_host_copy  s$    




z4TestCudaNDArray.test_devicearray_broadcast_host_copyc                 C   sH   t  d¡}t |¡}t  d¡d d d… }| |¡ t j | ¡ |¡ d S )Né
   é   r&   )r   r   r   r   ro   r   r   r   )r   rx   rs   r5   r   r   r   Ú(test_devicearray_contiguous_host_strided3  s
    


z8TestCudaNDArray.test_devicearray_contiguous_host_stridedc              	   C   sb   t  t d¡¡}t d¡}|  t¡ }| t  |¡d d d… ¡ W 5 Q R X |  tj	t
|jƒ¡ d S )Nr‰   r&   )r   r   r   r   rH   rV   ro   r    r   rp   rK   rL   )r   rs   r5   rN   r   r   r   Ú*test_devicearray_contiguous_device_strided:  s    
$þz:TestCudaNDArray.test_devicearray_contiguous_device_stridedz,DeviceNDArray class not present in simulatorc                 C   s4   t  ddtj¡}|  |jd ¡ |  |jd ¡ d S )N)r)   rˆ   )i   r`   ÚC_CONTIGUOUSÚF_CONTIGUOUS)r   ÚDeviceNDArrayr   r3   r?   r   )r   r5   r   r   r   Ú test_devicearray_relaxed_stridesD  s    z0TestCudaNDArray.test_devicearray_relaxed_stridesc                 C   sj   d}d}t  ||¡D ]P\}}tj||d}t |¡}|  |jd |jd ¡ |  |jd |jd ¡ qd S )N))r)   r(   )r(   r)   )ru   rk   rl   rŒ   r   )Ú	itertoolsÚproductr   Úndarrayr   r   r    r   )r   ZshapesZordersr	   rm   r5   Zd_arrr   r   r   Ú!test_c_f_contiguity_matches_numpyS  s    
ÿÿz1TestCudaNDArray.test_c_f_contiguity_matches_numpyz Typing not done in the simulatorc                 C   s,   t jddd}t |¡}|  |jjd¡ d S ©Nrˆ   ru   rl   ©r   Úzerosr   r   r    Ú_numba_type_Úlayout©r   rq   rs   r   r   r   Ú&test_devicearray_typing_order_simple_ca  s    
z6TestCudaNDArray.test_devicearray_typing_order_simple_cc                 C   s,   t jddd}t |¡}|  |jjd¡ d S )Nrˆ   rk   rl   ru   r•   r™   r   r   r   Ú&test_devicearray_typing_order_simple_fh  s    
z6TestCudaNDArray.test_devicearray_typing_order_simple_fc                 C   s,   t jddd}t |¡}|  |jjd¡ d S )N©r&   rˆ   ru   rl   r•   r™   r   r   r   Ú"test_devicearray_typing_order_2d_co  s    
z2TestCudaNDArray.test_devicearray_typing_order_2d_cc                 C   s,   t jddd}t |¡}|  |jjd¡ d S )Nrœ   rk   rl   r•   r™   r   r   r   Ú"test_devicearray_typing_order_2d_fv  s    
z2TestCudaNDArray.test_devicearray_typing_order_2d_fc                 C   s8   t jddd}t |¡d d …df }|  |jjd¡ d S )N©r:   r:   ru   rl   r&   ÚAr•   r™   r   r   r   Ú/test_devicearray_typing_order_noncontig_slice_c}  s    z?TestCudaNDArray.test_devicearray_typing_order_noncontig_slice_cc                 C   s8   t jddd}t |¡dd d …f }|  |jjd¡ d S )NrŸ   rk   rl   r&   r    r•   r™   r   r   r   Ú/test_devicearray_typing_order_noncontig_slice_f„  s    z?TestCudaNDArray.test_devicearray_typing_order_noncontig_slice_fc                 C   s8   t jddd}t |¡dd d …f }|  |jjd¡ d S )NrŸ   ru   rl   r&   r•   r™   r   r   r   Ú,test_devicearray_typing_order_contig_slice_c‹  s    z<TestCudaNDArray.test_devicearray_typing_order_contig_slice_cc                 C   s8   t jddd}t |¡d d …df }|  |jjd¡ d S )NrŸ   rk   rl   r&   ru   r•   r™   r   r   r   Ú,test_devicearray_typing_order_contig_slice_f’  s    z<TestCudaNDArray.test_devicearray_typing_order_contig_slice_fc                 C   s2   t  t  dg¡d¡}t |¡}|  |jjd¡ d S )Nr)   )rˆ   r    )r   r…   r#   r   r   r    r—   r˜   r™   r   r   r   Ú)test_devicearray_typing_order_broadcastedš  s    
z9TestCudaNDArray.test_devicearray_typing_order_broadcastedc                 C   s8   t jdt jd}t |¡}t  |¡}|  |j|j¡ d S )Nrˆ   r   )r   r   Zint16r   r   r   r    r   )r   r   r   Úgotr   r   r   Útest_bug6697¡  s    

zTestCudaNDArray.test_bug6697c                 C   s„   t jddtjd}tjddtjd}|  |jd¡ | |¡ | |¡ t	 
|¡}|  |jd¡ |  |jd¡ | |¡ | |¡ d S )N)r   )r`   )r	   r€   r   )r   rŽ   r   Zint8r’   r    r€   r   ro   r   r   r	   )r   Z	dev_arrayZ
host_arrayZdev_array_from_hostr   r   r   Útest_issue_8477§  s    	ÿ



zTestCudaNDArray.test_issue_8477N)/Ú__name__Ú
__module__Ú__qualname__r   r   r!   r$   r,   r1   r6   r9   r;   r=   rD   rF   r   rO   rR   rX   rY   r^   rb   rd   re   rh   ri   rt   rz   r}   r~   r   r‡   rŠ   r‹   r   r“   rš   r›   r   rž   r¡   r¢   r£   r¤   r¥   r§   r¨   r   r   r   r   r   	   sn   
	
	


	










r   c                   @   s   e Zd Zdd„ ZdS )ÚTestRecarrayc                 C   s®   t jddt jfdt jfgd}t j|jt jd|_t j|jt jdd |_|j}|j}dd„ }t  |¡}t  |¡}t	 
|¡d|jf |||ƒ t j ||¡ t j ||¡ d S )	N)é   Úvalue1Úvalue2r   r   c                 S   s4   t  d¡}|| jk r0| j| ||< | j| ||< d S )Nr)   )r   ÚgridÚsizer®   r¯   )ÚxZout1Zout2Úir   r   r   ÚtestØ  s    

z(TestRecarray.test_recarray.<locals>.testr)   )r   ZrecarrayZint64r3   r   r±   r®   r¯   Z
zeros_liker   Zjitr   r   )r   rq   Zexpect1Zexpect2r´   Zgot1Zgot2r   r   r   Útest_recarrayÌ  s    þ

zTestRecarray.test_recarrayN)r©   rª   r«   rµ   r   r   r   r   r¬   Ë  s   r¬   c                   @   st   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S )ÚTestCoreContiguousc                 C   s"   |   t |¡t |¡jd ¡ d S )NrŒ   )r    r   Zis_contiguousZ
array_corer   )r   r]   r   r   r   Ú_test_against_array_coreç  s    þz+TestCoreContiguous._test_against_array_corec                 C   s   t jddd}|  |¡ d S r”   ©r   r
   r·   ©r   Zd_ar   r   r   Útest_device_array_like_1dí  s    z,TestCoreContiguous.test_device_array_like_1dc                 C   s   t jddd}|  |¡ d S ©N©rˆ   rG   ru   rl   r¸   r¹   r   r   r   Útest_device_array_like_2dñ  s    z,TestCoreContiguous.test_device_array_like_2dc                 C   s   t jddd}|  |j¡ d S r»   ©r   r
   r·   rg   r¹   r   r   r   Ú#test_device_array_like_2d_transposeõ  s    z6TestCoreContiguous.test_device_array_like_2d_transposec                 C   s   t jddd}|  |¡ d S )N©rˆ   rG   é   ru   rl   r¸   r¹   r   r   r   Útest_device_array_like_3dù  s    z,TestCoreContiguous.test_device_array_like_3dc                 C   s   t jddd}|  |¡ d S )Nrˆ   rk   rl   r¸   r¹   r   r   r   Útest_device_array_like_1d_fý  s    z.TestCoreContiguous.test_device_array_like_1d_fc                 C   s   t jddd}|  |¡ d S ©Nr¼   rk   rl   r¸   r¹   r   r   r   Útest_device_array_like_2d_f  s    z.TestCoreContiguous.test_device_array_like_2d_fc                 C   s   t jddd}|  |j¡ d S rÄ   r¾   r¹   r   r   r   Ú%test_device_array_like_2d_f_transpose  s    z8TestCoreContiguous.test_device_array_like_2d_f_transposec                 C   s   t jddd}|  |¡ d S )NrÀ   rk   rl   r¸   r¹   r   r   r   Útest_device_array_like_3d_f	  s    z.TestCoreContiguous.test_device_array_like_3d_fc                 C   s&   d}t  |¡d d d… }|  |¡ d S )Nrˆ   r&   ©r   r–   r·   ©r   r	   r]   r   r   r   Útest_1d_view  s    zTestCoreContiguous.test_1d_viewc                 C   s*   d}t j|ddd d d… }|  |¡ d S )Nrˆ   rk   rl   r&   rÈ   rÉ   r   r   r   Útest_1d_view_f  s    z!TestCoreContiguous.test_1d_view_fc                 C   s0   d}t  |¡d d d…d d d…f }|  |¡ d S )Nr¼   r&   rÈ   rÉ   r   r   r   Útest_2d_view  s    zTestCoreContiguous.test_2d_viewc                 C   s4   d}t j|ddd d d…d d d…f }|  |¡ d S )Nr¼   rk   rl   r&   rÈ   rÉ   r   r   r   Útest_2d_view_f  s    "z!TestCoreContiguous.test_2d_view_fN)r©   rª   r«   r·   rº   r½   r¿   rÂ   rÃ   rÅ   rÆ   rÇ   rÊ   rË   rÌ   rÍ   r   r   r   r   r¶   æ  s   r¶   Ú__main__)r   Únumpyr   Znumba.cuda.cudadrvr   Znumbar   Znumba.cuda.testingr   r   r   r   r¬   r¶   r©   Úmainr   r   r   r   Ú<module>   s      E<