o
    ۾iO                     @   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 )Nd   )shapegX9v?r   )
r   device_arrayr   verify_cuda_ndarray_interfacenpempty	to_deviceasarrayassertEqualndim)selfdaryary r   ^/home/ubuntu/.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<