o
    8ήc7                     @   s   d dl mZ d dlZd dlmZ d dlmZ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G dd deZedkrJe  dS dS )    )productN)cuda)unittestCUDATestCaseskip_on_cudasim)patchc                   @   $   e Zd Zdd Zdd Zdd ZdS )CudaArrayIndexingc                 C   s   t d}t|}|j\}t| |D ]}| || ||  q| t || d   W d    n1 s8w   Y  | t ||  W d    d S 1 sRw   Y  d S )N
      )	nparanger   	to_deviceshaperangeassertEqualassertRaises
IndexError)selfarrdarrxi r   W/tmp/pip-target-vg8gfxp4/lib/python/numba/cuda/tests/cudadrv/test_cuda_array_slicing.pytest_index_1d   s   


"zCudaArrayIndexing.test_index_1dc                 C   sP  t ddd}t|}|j\}}t| |D ]}t| |D ]}| |||f |||f  q!q| t	 || d df  W d    n1 sLw   Y  | t	 ||df  W d    n1 sgw   Y  | t	 |d| d f  W d    n1 sw   Y  | t	 |d|f  W d    d S 1 sw   Y  d S )N         r   r   
r   r   reshaper   r   r   r   r   r   r   )r   r   r   r   yr   jr   r   r   test_index_2d   s&   

"zCudaArrayIndexing.test_index_2dc           	   
   C   s  t dddd}t|}|j\}}}t| |D ]$}t| |D ]}t| |D ]}| ||||f ||||f  q+q#q| t	 || d ddf  W d    n1 sZw   Y  | t	 ||ddf  W d    n1 svw   Y  | t	 |d| d df  W d    n1 sw   Y  | t	 |d|df  W d    n1 sw   Y  | t	 |dd| d f  W d    n1 sw   Y  | t	 |dd|f  W d    d S 1 sw   Y  d S )N<   r   r      r   r   r   )	r   r   r   r   r!   zr   r"   kr   r   r   test_index_3d&   s6   
""zCudaArrayIndexing.test_index_3dN)__name__
__module____qualname__r   r#   r(   r   r   r   r   r	   
   s    r	   c                   @   r   )CudaArrayStridedSlicec                 C   sN   t d}t|}t|jD ]}t j||d d ||d d   qd S )Nr
      )	r   r   r   r   r   sizetestingassert_equalcopy_to_hostr   r   r   r   r   r   r   test_strided_index_1d>   s
   

(z+CudaArrayStridedSlice.test_strided_index_1dc              
   C   s   t ddd}t|}t|jd D ])}t|jd D ]}t j||d d|d df ||d d|d df 	  qqd S )N*         r   r   r-   
r   r   r    r   r   r   r   r/   r0   r1   r   r   r   r   r"   r   r   r   test_strided_index_2dD   s   
z+CudaArrayStridedSlice.test_strided_index_2dc                 C   s   t dddd}t|}t|jd D ];}t|jd D ]1}t|jd D ]'}t j||d d|d d|d df ||d d|d d|d df 	  q(qqd S )NiP  r5   r6      r   r   r-   r7   )r   r   r   r   r"   r'   r   r   r   test_strided_index_3dM   s   
"z+CudaArrayStridedSlice.test_strided_index_3dN)r)   r*   r+   r3   r9   r;   r   r   r   r   r,   <   s    	r,   c                   @   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S )CudaArraySlicingc                 C   sX   t d}t|}t|jD ]}||d  }||d   }| t ||k qd S )Nr%   )	r   r   r   r   r   r.   r1   
assertTrueall)r   r   r   r   expectgotr   r   r   test_prefix_1dZ   s   

zCudaArraySlicing.test_prefix_1dc                 C   s   t ddd}t|}t|jd D ]>}t|jd D ]4}||d |d f }||d |d f }| |j|j | |j|j |	 }| 
t ||k qqd S )N	   r   r   r   r   r   r    r   r   r   r   r   stridesr1   r=   r>   )r   r   r   r   r"   r?   slicedr@   r   r   r   test_prefix_2db   s   
zCudaArraySlicing.test_prefix_2dc                 C   s   t dddd}t|}t|jd D ](}|| }|| }| |j|j | |j|j |	 }| 
t ||k qt|jd D ]6}t|jd D ],}|||f }|||f }| |j|j | |j|j |	 }| 
t ||k qOqFd S )Nr$   r   r   r%   r   r   rC   )r   r   r   r   r?   rE   r@   r"   r   r   r   test_select_3d_first_two_dimn   s&   
z-CudaArraySlicing.test_select_3d_first_two_dimc                 C      t djddddd}t|}t|jd D ]H}t|jd D ]}| t |||d d f 	 |||d d f  q!t|jd	 D ]}| t ||d d |f 	 ||d d |f  qDqt|jd D ]%}t|jd	 D ]}| t |d d ||f 	 |d d ||f  qqqhd S )
N   r%   r5   r6   Forderr   r   r-   
r   r   r    r   r   r   r   r=   array_equalr1   r   adar   r"   r   r   r   test_select_f   &   
zCudaArraySlicing.test_select_fc                 C   rH   )
NrI   r%   r5   r6   CrK   r   r   r-   rM   rO   r   r   r   test_select_c   rS   zCudaArraySlicing.test_select_cc                 C   sT   t djdddd}t|}| t |d ddf  |d ddf k d S )N#   r%   r6   rJ   rK   r   )r   r   r    r   r   r=   r>   r1   r   r   r   r   r   r   test_prefix_select   s   
4z#CudaArraySlicing.test_prefix_selectc                 C   sV   t d}t|}ttddddD ]\}}t j||| |||   qd S )Nr
   ir-   repeat)	r   r   r   r   r   r   r/   assert_array_equalr1   r8   r   r   r   test_negative_slicing_1d   s   

z)CudaArraySlicing.test_negative_slicing_1dc                 C   sr   t ddd}t|}ttddddD ]\}}}}t j|||||f |||||f 	  qd S )Nr   r   r   rY   )
r   r   r    r   r   r   r   r/   r[   r1   )r   r   r   r   r!   wsr   r   r   test_negative_slicing_2d   s   
z)CudaArraySlicing.test_negative_slicing_2dc                 C   s   t d}t|}t|jd D ]}t j|||  |||  q| 	|d d d d   t j|d d d d  |d d d d  t j|d d dd   |d d dd   d S )Nr%   r   r   )
r   r   r   r   r   r   r/   r[   r1   assertFalser2   r   r   r   test_empty_slice_1d   s   

$z$CudaArraySlicing.test_empty_slice_1dc                 C   s   t ddd}t|}t j|d d  |d d  t j|dd df  |dd df  | |d d d d   t j|d d d d  |d d d d  t j|d d dd   |d d dd   d S )NrV   r%   r6   r   r   r   ra   )	r   r   r    r   r   r/   r[   r1   rb   rW   r   r   r   test_empty_slice_2d   s   
"*2z$CudaArraySlicing.test_empty_slice_2dN)r)   r*   r+   rA   rF   rG   rR   rU   rX   r\   r`   rc   rd   r   r   r   r   r<   Y   s    r<   c                   @   s   e Zd 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eddd Zeddd Zeddd Zedd d! Zd"S )#CudaArraySettingzp
    Most of the slicing logic is tested in the cases above, so these
    tests focus on the setting logic.
    c                 C   B   t ddd}t|}d|d< d|d< t j| | d S )NrV   r%   r6     )r-   r-   r   r   r    r   r   r/   r[   r1   rW   r   r   r   test_scalar   
   
zCudaArraySetting.test_scalarc                 C   rf   NrV   r%   r6   rg   r-   rh   rW   r   r   r   	test_rank   rj   zCudaArraySetting.test_rankc                 C   sR   t ddd}t|}d|d d df< d|d d df< t j| | d S rk   rh   rW   r   r   r   test_broadcast   s
   
zCudaArraySetting.test_broadcastc                 C   sP   t ddd}t|}t jddd}||d< ||d< t j| | d S NrV   r%   r6     r   
fill_valuer-   	r   r   r    r   r   fullr/   r[   r1   r   r   r   _400r   r   r   test_array_assign_column   s   
z)CudaArraySetting.test_array_assign_columnc                 C   s`   t ddd}t|}t jddd}||d d df< ||d d df< t j| | d S rn   rr   rt   r   r   r   test_array_assign_row   s   
z&CudaArraySetting.test_array_assign_rowc                 C   sR   t dddd}t|}t jddd}||d< ||d< t j| | d S )	NrI   r%   r5   r6   )r5   r6   ro   rp   r-   rr   rt   r   r   r   test_array_assign_subarray   s   
z+CudaArraySetting.test_array_assign_subarrayc                 C   sp   t ddddd}t|}t jddd}||d d d d d	f< ||d d d d d	f< t j| | d S )
Ni  r%   r5   r6   r:   )r%   r5   r:   ro   rp   r-   rr   rt   r   r   r   test_array_assign_deep_subarray  s   
z0CudaArraySetting.test_array_assign_deep_subarrayc                 C   sX   t ddd}t|}t jddd}||d d < ||d d < t j| | d S )NrV   r%   r6   )r%   r6   ro   rp   rr   rt   r   r   r   test_array_assign_all  s   
z&CudaArraySetting.test_array_assign_allc                 C   sF   t d}t|}d|d d d< d|d d d< t j| | d S )N   rg   r-   )r   onesr   r   r/   r[   r1   rW   r   r   r   test_strides  s
   

zCudaArraySetting.test_stridesc                 C   sj   t td}| t}tjdd|d d < W d    n1 s"w   Y  | jt|j	g dd d S )NrV   )r   r-   r   )r   )z"Can't assign 3-D array to 1-D selfz@could not broadcast input array from shape (2,3) into shape (35)zAcould not broadcast input array from shape (2,3) into shape (35,)member	container)
r   r   r   r   r   
ValueErrorr|   assertInstr	exceptionr   r   er   r   r   test_incompatible_highdim  s   
z*CudaArraySetting.test_incompatible_highdimc                 C   sf   t td}| t}ddg|d d < W d    n1 s w   Y  | jt|jg dd d S )Nr%   r   r   )z@Can't copy sequence with size 2 to array axis 0 with dimension 5z?cannot copy sequence with size 2 to array axis with dimension 5z?could not broadcast input array from shape (2,) into shape (5,)r~   )	r   r   r   r   r   r   r   r   r   r   r   r   r   test_incompatible_shape+  s   
z(CudaArraySetting.test_incompatible_shapez7cudasim does not use streams and operates synchronouslyc                 C   sZ   t td}tjt jjjdd d}d|d< W d    n1 s"w   Y  |	  d S )Nr%   synchronizereturn_valuer
   r   )
r   r   r   r   r   objectcudadrvdriverStreamassert_called_once)r   r   	mock_syncr   r   r   	test_sync<  s   
zCudaArraySetting.test_syncc              	   C   s   t  t  t  t  f}|D ]/}t jtd|d}tj	t j
jjdd d}d|d< W d    n1 s6w   Y  |  qd S )Nr%   streamr   r   r
   r   )r   r   default_streamlegacy_default_streamper_thread_default_streamr   r   r   r   r   r   r   r   assert_not_calledr   streamsr   r   r   r   r   r   test_no_sync_default_streamG  s   

z,CudaArraySetting.test_no_sync_default_streamc              	   C   s   t  t  t  t  f}|D ]1}t td}tj	t j
jjdd d}|jdd|d W d    n1 s8w   Y  |  qd S )Nr%   r   r   r   r
   r   )r   r   r   r   r   r   r   r   r   r   r   r   r   setitemr   r   r   r   r   test_no_sync_supplied_streamY  s   
z-CudaArraySetting.test_no_sync_supplied_streamzRequires PR #6367c                 C   sJ   t jdtjd}d|d d < |d}d|d< d|d< | t|d d S )Nr-   )dtyper   u1r   r%   i   )r   mapped_arrayr   int32viewr   sum)r   aryary_vr   r   r   test_issue_6505k  s   
z CudaArraySetting.test_issue_6505N)r)   r*   r+   __doc__ri   rl   rm   rv   rw   rx   ry   rz   r}   r   r   r   r   r   r   r   skipr   r   r   r   r   re      s*    



re   __main__)	itertoolsr   numpyr   numbar   numba.cuda.testingr   r   r   unittest.mockr   r	   r,   r<   re   r)   mainr   r   r   r   <module>   s    2x *