o
    gµ„i³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rDe ¡  dS 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   Úverify_cuda_ndarray_interfaceÚnpÚemptyÚ	to_deviceÚasarrayÚassertEqualÚndim)ÚselfÚdaryÚary© r   ú]/home/kuhnn/.local/lib/python3.10/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ÚflagsÚ	writeableÚassertFalser   r   Úcopy_to_hostÚtestingÚassert_array_equal)r   r   r   Ú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   Úf4)r	   r   )r   r
   r   r   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   sj   t  ¡ }| ¡ # t jdtj|d}|  | |¡j|¡ |  |j|¡ W d   ƒ d S 1 s.w   Y  d S )N)r+   r+   )r   Ústream)r   r6   Úauto_synchronizer
   r   Úfloat64r   Úbind)r   r6   Ú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   ó   

zTestCudaNDArray.test_len_1dc                 C   ó,   t  d¡}t d¡}|  t|ƒt|ƒ¡ d S )N)r+   é   r<   r/   r   r   r   Útest_len_2dG   r?   zTestCudaNDArray.test_len_2dc                 C   r@   )N)r+   rA   é   r<   r/   r   r   r   Útest_len_3dL   r?   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   r2   r&   r   r   ÚsplitÚ
assertTrueÚallr   )r   ÚNr'   r3   r4   Ú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*   )Útor1   )r   rH   r'   r3   r4   r   r   r   Útest_devicearray_replacea   s   

z(TestCudaNDArray.test_devicearray_replacezThis works in the simulatorc                 C   sj   t  t t d¡¡ ddd¡¡}|  t¡}t |¡ W d   ƒ n1 s%w   Y  |  	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   r4   Ú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*   ©Úaxes)
r   r'   r   r.   rQ   r   r   r   rF   rG   ©r   r3   r'   r   r   r   Ú#test_devicearray_transpose_identityu   s   ÿÿz3TestCudaNDArray.test_devicearray_transpose_identityc                 C   sr   t  t t d¡¡ dd¡¡}|  t¡}tj|dd W d   ƒ n1 s&w   Y  | j	t
|jƒddgd d S )	NrN   r+   r,   )r   r   rW   zinvalid axes list (0, 0)zrepeated axis in transpose©Ú	container©r   r   r   r'   r   r.   rO   Ú
ValueErrorrQ   ÚassertInrR   rS   rT   r   r   r   Ú)test_devicearray_transpose_duplicatedaxis|   s   ÿþ
þz9TestCudaNDArray.test_devicearray_transpose_duplicatedaxisc                 C   sr   t  t t d¡¡ dd¡¡}|  t¡}tj|dd W d   ƒ n1 s&w   Y  | j	t
|jƒg d¢d d S )NrN   r+   r,   )r   r*   rW   )zinvalid axes list (0, 2)zinvalid axis for this arrayz0axis 2 is out of bounds for array of dimension 2r[   r]   rT   r   r   r   Ú$test_devicearray_transpose_wrongaxis‰   s   ÿ
þz4TestCudaNDArray.test_devicearray_transpose_wrongaxisc              	   C   s~   t jt  d¡dd dd¡}t |¡}dD ]'}| j|d t j | 	|¡ 
¡ | 	|¡¡ W d   ƒ n1 s7w   Y  qd S )NrN   Úi2r   r+   r,   )Úi4Úu4Úi8Úf8)r   r'   r   r.   r   r   ÚsubTestr    r!   Úviewr   )r   r3   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é    rb   r   r,   é   r*   Úu2)
r   r'   r   r.   r   r   r    r!   rh   r   rY   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 d   ƒ n1 s1w   Y  t	|j
ƒ}|  d|¡ d	|v }d
|v }|  |pN|d¡ d S )Nrj   rb   r   r,   rk   r*   rc   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   rO   r^   rh   rR   rS   r_   rF   )r   r3   r'   rU   ÚmsgÚcontiguous_pre_np123Úcontiguous_post_np123r   r   r   Ú&test_devicearray_view_bad_not_c_contigª   s   ÿ

ÿz6TestCudaNDArray.test_devicearray_view_bad_not_c_contigc                 C   sp   t jt  d¡dd dd¡}t |¡}|  t¡}| d¡ W d   ƒ n1 s(w   Y  |  	dt
|jƒ¡ d S )NrN   rb   r   r,   r+   rc   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   rO   r^   rh   r   rR   rS   )r   r3   r'   rU   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 ©NrN   r+   r,   )r   r'   r   r.   rQ   r   r   r   rF   rG   ÚTrY   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 rs   )
r   r'   r   r.   r   r   rt   r   rF   rG   rY   r   r   r   Útest_devicearray_transpose_TÈ   s   z,TestCudaNDArray.test_devicearray_transpose_Tc                 C   s†   t  d¡jdddd}t jddd}t |¡}||d< |  t¡}|d  |¡ W d   ƒ n1 s2w   Y  |  	t
jt|jƒ¡ d S )Né   rA   ÚF©Úorder)rA   )Ú
fill_valuer	   r*   )r   r   r.   Úfullr   r   rO   r^   Úcopy_to_devicer   r   Úerrmsg_contiguous_bufferrR   rS   )r   ÚaÚsÚdrU   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 ]<\}}d|jjr'dnd|jjr.dndf }t |¡}| |¡ | j t | 	¡ |k¡|d | j t | 	¡ |k¡|d qdS )z-
        Checks host->device memcpys
        z%s => %sÚCrx   )rn   N)
rF   r   Úc_contiguousÚf_contiguousr   r   r}   r   rG   r   )r   Úa_cÚa_fr3   r&   rn   r   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é}   rA   rx   ry   )r   r   r.   r'   rˆ   ©r   r†   r‡   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 )NrA   rx   ry   )r   r   r'   rˆ   rŠ   r   r   r   Ú(test_devicearray_contiguous_copy_host_1dý   s   
z8TestCudaNDArray.test_devicearray_contiguous_copy_host_1dc                 C   sV  t  d¡ ddd¡}t j|dd}|  |jj¡ |  |jj¡ t 	|¡}|  
t¡}| t 	|¡¡ W d   ƒ n1 s<w   Y  |  d |j|j¡t|jƒ¡ | t 	|¡¡ |  t  | ¡ |k¡¡ t 	|¡}|  
t¡}| t 	|¡¡ W d   ƒ n1 sw   Y  |  d |j|j¡t|jƒ¡ | t 	|¡¡ |  t  | ¡ |k¡¡ d S )Nr‰   rA   rx   ry   zincompatible strides: {} vs. {})r   r   r.   r'   rF   r   r„   r…   r   r   rO   r^   r}   r   ÚformatÚstridesrR   rS   rG   r   )r   r†   r‡   r   rU   r   r   r   Ú'test_devicearray_contiguous_copy_device  s0   
ÿþ
ÿþz7TestCudaNDArray.test_devicearray_contiguous_copy_devicec                 C   s  d}d}t  |¡}t  |¡j|dd}t  |¡j|dd}tt|ƒƒD ]d}td ƒf| t jf }|d |… |f ||d …  }t  || |¡}	t  || |¡}
t	 
|	¡}t	 
|
¡}t j | ¡ |	¡ t j | ¡ |
¡ | |
¡ | |	¡ t j | ¡ |
¡ t j | ¡ |	¡ q#d S )Nr,   )r*   r+   rƒ   ry   rx   )r   Úprodr   r.   Úranger=   ÚsliceÚnewaxisÚbroadcast_tor   r   r    r!   r   r}   )r   Ú	broadsizeÚ	coreshapeÚcoresizeÚcore_cÚcore_fÚdimÚnewindexÚ
broadshapeÚbroad_cÚbroad_fÚdbroad_cÚ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   r}   r    r!   r   )r   r†   r   r:   r   r   r   Ú(test_devicearray_contiguous_host_strided3  s
   


z8TestCudaNDArray.test_devicearray_contiguous_host_stridedc                 C   sv   t  t d¡¡}t d¡}|  t¡}| t  |¡d d d… ¡ W d   ƒ n1 s*w   Y  |  tj	t
|jƒ¡ d S )Nr£   r*   )r   r   r   r   rO   r^   r}   r   r   r~   rR   rS   )r   r   r:   rU   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   rk   ÚC_CONTIGUOUSÚF_CONTIGUOUS)r   ÚDeviceNDArrayr   r8   rF   r   )r   r:   r   r   r   Ú test_devicearray_relaxed_stridesD  s   z0TestCudaNDArray.test_devicearray_relaxed_stridesc                 C   sj   d}d}t  ||¡D ](\}}tj||d}t |¡}|  |jd |jd ¡ |  |jd |jd ¡ q
d S )N))r-   r,   )r,   r-   )rƒ   rx   ry   r¦   r§   )Ú	itertoolsÚproductr   Úndarrayr   r   r   r   )r   ÚshapesÚordersr	   rz   r:   Ú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   ó,   t jddd}t |¡}|  |jjd¡ d S ©Nr¢   rƒ   ry   ©r   Úzerosr   r   r   Ú_numba_type_Úlayout©r   r   r   r   r   r   Ú&test_devicearray_typing_order_simple_ca  ó   
z6TestCudaNDArray.test_devicearray_typing_order_simple_cc                 C   s,   t jddd}t |¡}|  |jjd¡ d S )Nr¢   rx   ry   rƒ   r³   r·   r   r   r   Ú&test_devicearray_typing_order_simple_fh  r¹   z6TestCudaNDArray.test_devicearray_typing_order_simple_fc                 C   r±   )N©r*   r¢   rƒ   ry   r³   r·   r   r   r   Ú"test_devicearray_typing_order_2d_co  r¹   z2TestCudaNDArray.test_devicearray_typing_order_2d_cc                 C   r±   )Nr»   rx   ry   r³   r·   r   r   r   Ú"test_devicearray_typing_order_2d_fv  r¹   z2TestCudaNDArray.test_devicearray_typing_order_2d_fc                 C   ó8   t jddd}t |¡d d …df }|  |jjd¡ d S )N©rA   rA   rƒ   ry   r*   ÚAr³   r·   r   r   r   Ú/test_devicearray_typing_order_noncontig_slice_c}  ó   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¿   rx   ry   r*   rÀ   r³   r·   r   r   r   Ú/test_devicearray_typing_order_noncontig_slice_f„  rÂ   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¿   rƒ   ry   r*   r³   r·   r   r   r   Ú,test_devicearray_typing_order_contig_slice_c‹  rÂ   z<TestCudaNDArray.test_devicearray_typing_order_contig_slice_cc                 C   r¾   )Nr¿   rx   ry   r*   rƒ   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   Ú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   )rk   )r	   rŽ   r   )r   r¨   r   Úint8r¬   r   rŽ   r   r}   r   r   r	   )r   Ú	dev_arrayÚ
host_arrayÚdev_array_from_hostr   r   r   Útest_issue_8477§  s   	ÿ



zTestCudaNDArray.test_issue_8477N)/Ú__name__Ú
__module__Ú__qualname__r   r#   r%   r(   r0   r5   r;   r>   rB   rD   rK   rM   r   rV   rZ   r`   ra   ri   rm   rq   rr   ru   rv   r‚   rˆ   r‹   rŒ   r   r¡   r¤   r¥   r©   r°   r¸   rº   r¼   r½   rÁ   rÃ   rÄ   rÅ   rÆ   rÉ   rÎ   r   r   r   r   r   	   sp    
	
	


	










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   s8   t  d¡}|| jk r| j| ||< | j| ||< d S d S )Nr-   )r   ÚgridÚsizerÔ   rÕ   )ÚxÚout1Úout2Úir   r   r   ÚtestØ  s
   

þz(TestRecarray.test_recarray.<locals>.testr-   )r   ÚrecarrayÚint64r8   r   r×   rÔ   rÕ   Ú
zeros_liker   Újitr    r!   )r   r   Úexpect1Úexpect2rÜ   Úgot1Ú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   Úis_contiguousÚ
array_corer   )r   rh   r   r   r   Ú_test_against_array_coreç  s   þz+TestCoreContiguous._test_against_array_corec                 C   ó   t jddd}|  |¡ d S r²   ©r   r
   ré   ©r   Úd_ar   r   r   Útest_device_array_like_1dí  ó   z,TestCoreContiguous.test_device_array_like_1dc                 C   rê   ©N©r¢   rN   rƒ   ry   rë   rì   r   r   r   Útest_device_array_like_2dñ  rï   z,TestCoreContiguous.test_device_array_like_2dc                 C   ó   t jddd}|  |j¡ d S rð   ©r   r
   ré   rt   rì   r   r   r   Ú#test_device_array_like_2d_transposeõ  ó   z6TestCoreContiguous.test_device_array_like_2d_transposec                 C   rê   )N©r¢   rN   é   rƒ   ry   rë   rì   r   r   r   Útest_device_array_like_3dù  rï   z,TestCoreContiguous.test_device_array_like_3dc                 C   rê   )Nr¢   rx   ry   rë   rì   r   r   r   Útest_device_array_like_1d_fý  rï   z.TestCoreContiguous.test_device_array_like_1d_fc                 C   rê   ©Nrñ   rx   ry   rë   rì   r   r   r   Útest_device_array_like_2d_f  rï   z.TestCoreContiguous.test_device_array_like_2d_fc                 C   ró   rû   rô   rì   r   r   r   Ú%test_device_array_like_2d_f_transpose  rö   z8TestCoreContiguous.test_device_array_like_2d_f_transposec                 C   rê   )Nr÷   rx   ry   rë   rì   r   r   r   Útest_device_array_like_3d_f	  rï   z.TestCoreContiguous.test_device_array_like_3d_fc                 C   s&   d}t  |¡d d d… }|  |¡ d S )Nr¢   r*   ©r   r´   ré   ©r   r	   rh   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¢   rx   ry   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ñ   rx   ry   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   Únumba.cuda.cudadrvr   Únumbar   Únumba.cuda.testingr   r   r   r   rÒ   ræ   rÏ   Úmainr   r   r   r   Ú<module>   s       E<ÿ