o
    8ήc*                     @   s   d dl Zd dlm  mZ d dlmZ d dlm	Z	m
Z
mZ d dlmZ d dlmZ d dlmZmZ d dlZd dlZd dlmZ d dlmZ e
dfd	d
ZedG dd deZedkrbe  dS dS )    N)
namedtuple)voidfloat32float64)guvectorize)cuda)skip_on_cudasimCUDATestCase)NumbaPerformanceWarning)override_configc                 C   sb   t t| d d d d f | d d d d f | d d d d f gddddd }|}|r/||_|S )Nz(m,n),(n,p)->(m,p)r   targetc           	   
   S   sv   | j \}}|j \}}t|D ]*}t|D ]#}d|||f< t|D ]}|||f  | ||f |||f  7  < q qqd S Nr   shaperange)	ABCmnpijk r   J/tmp/pip-target-vg8gfxp4/lib/python/numba/cuda/tests/cudapy/test_gufunc.py
matmulcore   s   

*z*_get_matmulcore_gufunc.<locals>.matmulcore)r   r   max_blocksize)dtyper   r   gufuncr   r   r   _get_matmulcore_gufunc   s   >
	r!   z&ufunc API unsupported in the simulatorc                   @   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d'S )(TestCUDAGufuncc                 C   s~   t dd}d}tj|d d tjd|dd}tj|d d tjd|dd}|||}t||}| t|| d S N   r         r      	r!   nparanger   reshapeutmatrix_multiply
assertTrueallcloseselfr    	matrix_ctr   r   r   Goldr   r   r   test_gufunc_small%      

z TestCUDAGufunc.test_gufunc_smallc                 C   s   t dd}d}tj|d d tjd|dd}tj|d d tjd|dd}t|}||| }t	||}| 
t|| d S r#   )r!   r+   r,   r   r-   r   	to_devicecopy_to_hostr.   r/   r0   r1   )r3   r    r4   r   r   dBr   r5   r   r   r   test_gufunc_auto_transfer3   s   

z(TestCUDAGufunc.test_gufunc_auto_transferc                 C   s~   t dd}d}tj|d d tjd|dd}tj|d d tjd|dd}|||}t||}| t|| d S )Nr$   r%     r&   r'   r(   r)   r*   r2   r   r   r   test_gufuncC   r7   zTestCUDAGufunc.test_gufuncc                 C   s   t dd}d}tj|d d tjddddd}tj|d d tjddddd}|||}t||}| t|| d S )	Nr$   r%   d   r&   r'   r(      r)   r*   r2   r   r   r   test_gufunc_hidimQ   s   
$$
z TestCUDAGufunc.test_gufunc_hidimc                 C   sp   t td}tjddd}tjdd}t||}|||}tj|| ||t	|d}tj|| d S )Nr(   
      )rA      rC   )
r!   r   r+   randomrandnr.   r/   testingassert_allclosetile)r3   r    XYgoldres1res2r   r   r   test_gufunc_new_axis]   s   

z#TestCUDAGufunc.test_gufunc_new_axisc                 C   s   t dd}d}tj|d d tjd|dd}tj|d d tjd|dd}d|_|||}t||}| t	|| d S )	Nr$   r%   r<   r&   r'   r(   r)       )
r!   r+   r,   r   r-   r   r.   r/   r0   r1   r2   r   r   r   test_gufunc_adjust_blocksizel   s   

z+TestCUDAGufunc.test_gufunc_adjust_blocksizec                 C   s   t dd}d}tj|d d tjd|dd}tj|d d tjd|dd}t }t||}t||}tjd|j	|d	}|||||d
}|j
|d}	|  t||}
| t|	|
 d S )Nr$   r%   r<   r&   r'   r(   r)   )r<   r&   r)   )r   r   stream)outrQ   )rQ   )r!   r+   r,   r   r-   r   rQ   r8   device_arrayr   r9   synchronizer.   r/   r0   r1   )r3   r    r4   r   r   rQ   dAr:   dCr   r5   r   r   r   test_gufunc_stream{   s"   
z!TestCUDAGufunc.test_gufunc_streamc                 C   n   t ttd d  td d  gddddd }tjdtjdd }t|}|||d	 | t|| d S )
N(x)->(x)r   r   c                 S       t |jD ]}| | ||< qd S Nr   sizer   r   r   r   r   r   copy      z&TestCUDAGufunc.test_copy.<locals>.copyrA   r(   rC   rR   r   r   r   r+   r,   
zeros_liker0   r1   r3   r_   r   r   r   r   r   	test_copy      

zTestCUDAGufunc.test_copyc                 C   rX   )
NrY   r   r   c                 S   rZ   r[   r\   r^   r   r   r   r_      r`   z*TestCUDAGufunc.test_copy_odd.<locals>.copy   r(   rC   ra   rb   rd   r   r   r   test_copy_odd   rf   zTestCUDAGufunc.test_copy_oddc                 C   s   t ttd d d d f td d d d f gddddd }tjdtjddd	d
 }t|}|||d | t|| d S )Nz(x, y)->(x, y)r   r   c                 S   s@   t |jd D ]}t |jd D ]}| ||f |||f< qqd S )Nr   rC   )r   r   )r   r   xyr   r   r   copy2d   s
   z*TestCUDAGufunc.test_copy2d.<locals>.copy2d   r(   r)      rC   ra   )	r   r   r   r+   r,   r-   rc   r0   r1   )r3   rk   r   r   r   r   r   test_copy2d   s   ,

zTestCUDAGufunc.test_copy2dc              	   C   s  t dgddddd }tjdd}tjdd}t|jd	 d}td
dL tj	dd-}|||| | 
|d	 jt | dt|d	 j | dt|d	 j W d    n1 sew   Y  W d    d S W d    d S 1 s}w   Y  d S )N(void(float32[:], float32[:], float32[:])(n),(n)->(n)r   r   c                 S   0   | j d }t|D ]}| | ||  ||< q	d S r   r   abdistlenr   r   r   r   numba_dist_cuda      
zMTestCUDAGufunc.test_inefficient_launch_configuration.<locals>.numba_dist_cudai   r   r   CUDA_LOW_OCCUPANCY_WARNINGSrC   Trecordz	Grid sizezlow occupancy)r   r+   rD   randastypezerosr   r   warningscatch_warningsassertEqualcategoryr
   assertInstrmessage)r3   rw   rs   rt   ru   wr   r   r   %test_inefficient_launch_configuration   s"   
"z4TestCUDAGufunc.test_inefficient_launch_configurationc              	   C   s   t dgdddddd }tjdd	d
}tjdd	d
}t|}tdd5 tj	dd}|||| | 
t|d W d    n1 sOw   Y  W d    d S W d    d S 1 sgw   Y  d S )Nro   rp   Tr   )nopythonr   c                 S   rq   r   r   rr   r   r   r   numba_dist_cuda2   rx   zLTestCUDAGufunc.test_efficient_launch_configuration.<locals>.numba_dist_cuda2i   r   )i   r&   ry   rC   rz   r   )r   r+   rD   r|   r}   r-   rc   r   r   r   r   rv   )r3   r   rs   rt   ru   r   r   r   r   #test_efficient_launch_configuration   s&   

"z2TestCUDAGufunc.test_efficient_launch_configurationc                 C   s   dd }t ttd d  td d  gdddd| | t}t ttd d  td d  gdddd| W d    n1 s@w   Y  | dt|j d S )	Nc                 S      d S r[   r   r   r   r   r   r   foo      z.TestCUDAGufunc.test_nopython_flag.<locals>.foorY   r   T)r   r   Fznopython flag must be True)r   r   r   assertRaises	TypeErrorr   r   	exception)r3   r   raisesr   r   r   test_nopython_flag   s    z!TestCUDAGufunc.test_nopython_flagc                 C   s   dd }|  t}tttd d  td d  gddddd| W d    n1 s+w   Y  d}t|j}| |d t| | |t|d  	 
d	}d
d |D }| tddgt| d S )Nc                 S   r   r[   r   r   r   r   r   r      r   z.TestCUDAGufunc.test_invalid_flags.<locals>.foorY   r   TF)r   what1ever2z/The following target options are not supported:,c                 S   s   g | ]}| d qS )z'" )strip).0r   r   r   r   
<listcomp>  s    z5TestCUDAGufunc.test_invalid_flags.<locals>.<listcomp>r   r   )r   r   r   r   r   r   r   r   rv   r   splitset)r3   r   r   headmsgitemsr   r   r   test_invalid_flags   s   
z!TestCUDAGufunc.test_invalid_flagsc                 C   s   t ttd d  td d  gddddd }tjdtjd }}| t}||||d W d    n1 s8w   Y  d	}| t|j	| d S )
NrY   r   r   c                 S   r   r[   r   )inprR   r   r   r   r     s   z2TestCUDAGufunc.test_duplicated_output.<locals>.foorA   r(   ra   z>cannot specify 'out' as both a positional and keyword argument)
r   r   r   r+   r~   r   
ValueErrorr   r   r   )r3   r   r   rR   r   r   r   r   r   test_duplicated_output  s   $
z%TestCUDAGufunc.test_duplicated_outputc                 C   sp   t td d  td d  td d  fgddddd }|||}tjt|t| dd}tj|| d S )Nz(n),(n)->()r   r   c                 S   s6   d}t t| D ]}|| | ||  7 }q||d< d S r   )r   rv   )ri   rj   rsr   r   r   r   	gu_reduce  s   z1TestCUDAGufunc.check_tuple_arg.<locals>.gu_reducerC   )axis)r   r   r+   sumasarrayrF   assert_equal)r3   rs   rt   r   r   expectedr   r   r   check_tuple_arg  s   &

zTestCUDAGufunc.check_tuple_argc                 C   s   d}d}|  || d S )N)      ?       @      @      @      @      @)      ?      @      @      @      @      @)r   r3   rs   rt   r   r   r   test_tuple_of_tuple_arg  s   z&TestCUDAGufunc.test_tuple_of_tuple_argc                 C   sR   t dd}|dddd|ddddf}|d	d
dd|ddddf}| || d S )NPoint)ri   rj   zr   r   r   r   r   r   r   r   r   r   r   r   )r   r   )r3   r   rs   rt   r   r   r   test_tuple_of_namedtuple_arg$  s   
z+TestCUDAGufunc.test_tuple_of_namedtuple_argc                 C   s8   t dt df}t dt df}| || d S )Nr   r   r   r   )r+   r   r   r   r   r   r   test_tuple_of_array_arg,  s   z&TestCUDAGufunc.test_tuple_of_array_argN)__name__
__module____qualname__r6   r;   r=   r@   rN   rP   rW   re   rh   rn   r   r   r   r   r   r   r   r   r   r   r   r   r   r"   "   s(    r"   __main__)numpyr+   numpy.core.umath_testscoreumath_testsr.   collectionsr   numbar   r   r   r   r   numba.cuda.testingr   r	   unittestr   numba.core.errorsr
   numba.tests.supportr   r!   r"   r   mainr   r   r   r   <module>   s&      