U
    ,d*                     @   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re  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}t|D ]F}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   G/tmp/pip-unpacked-wheel-eu7e0c37/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%   s    

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   s    

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(   
      )r@      rB   )
r!   r   r+   randomZrandnr.   r/   testingZassert_allcloseZtile)r3   r    XYZgoldZres1Z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)outrJ   )rJ   )r!   r+   r,   r   r-   r   rJ   r7   Zdevice_arrayr   r8   Zsynchronizer.   r/   r0   r1   )r3   r    r4   r   r   rJ   ZdAr9   ZdCr   r5   r   r   r   test_gufunc_stream{   s"    
z!TestCUDAGufunc.test_gufunc_streamc                 C   sn   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   s    t |jD ]}| | ||< q
d S Nr   sizer   r   r   r   r   r   copy   s    z&TestCUDAGufunc.test_copy.<locals>.copyr@   r(   rB   rK   r   r   r   r+   r,   
zeros_liker0   r1   r3   rR   r   r   r   r   r   	test_copy   s    

zTestCUDAGufunc.test_copyc                 C   sn   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 )
NrM   r   r   c                 S   s    t |jD ]}| | ||< q
d S rN   rO   rQ   r   r   r   rR      s    z*TestCUDAGufunc.test_copy_odd.<locals>.copy   r(   rB   rS   rT   rV   r   r   r   test_copy_odd   s    

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   rB   )r   r   )r   r   xyr   r   r   copy2d   s    z*TestCUDAGufunc.test_copy2d.<locals>.copy2d   r(   r)      rB   rS   )	r   r   r   r+   r,   r-   rU   r0   r1   )r3   r\   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h tj	ddP}|||| | 
|d	 jt | dt|d	 j | dt|d	 j W 5 Q R X W 5 Q R X d S )N(void(float32[:], float32[:], float32[:])(n),(n)->(n)r   r   c                 S   s0   | j d }t|D ]}| | ||  ||< qd S r   r   abdistlenr   r   r   r   numba_dist_cuda   s    
zMTestCUDAGufunc.test_inefficient_launch_configuration.<locals>.numba_dist_cudai   r   r   CUDA_LOW_OCCUPANCY_WARNINGSrB   Trecordz	Grid sizezlow occupancy)r   r+   rC   randastypezerosr   r   warningscatch_warningsassertEqualcategoryr
   ZassertInstrmessage)r3   rg   rc   rd   re   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: tj	dd"}|||| | 
t|d W 5 Q R X W 5 Q R X d S )Nr`   ra   Tr   )nopythonr   c                 S   s0   | j d }t|D ]}| | ||  ||< qd S r   r   rb   r   r   r   numba_dist_cuda2   s    
zLTestCUDAGufunc.test_efficient_launch_configuration.<locals>.numba_dist_cuda2i   r   )i   r&   rh   rB   ri   r   )r   r+   rC   rk   rl   r-   rU   r   rn   ro   rp   rf   )r3   rw   rc   rd   re   rt   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2}t ttd d  td d  gdddd| W 5 Q R X | dt|j d S )	Nc                 S   s   d S rN   r   r   r   r   r   r   foo   s    z.TestCUDAGufunc.test_nopython_flag.<locals>.foorM   r   T)r   rv   Fznopython flag must be True)r   r   r   assertRaises	TypeErrorrp   rr   	exception)r3   rz   raisesr   r   r   test_nopython_flag   s      z!TestCUDAGufunc.test_nopython_flagc              	   C   s   dd }|  t4}tttd d  td d  gddddd| W 5 Q R X d}t|j}| |d t| | |t|d  	 
d	}d
d |D }| tddgt| d S )Nc                 S   s   d S rN   r   ry   r   r   r   rz      s    z.TestCUDAGufunc.test_invalid_flags.<locals>.foorM   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   rr   r}   rp   rf   r   splitset)r3   rz   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 5 Q R X d	}| t|j	| d S )
NrM   r   r   c                 S   s   d S rN   r   )inprK   r   r   r   rz     s    z2TestCUDAGufunc.test_duplicated_output.<locals>.foor@   r(   rS   z>cannot specify 'out' as both a positional and keyword argument)
r   r   r   r+   rm   r{   
ValueErrorrp   rr   r}   )r3   rz   r   rK   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   rf   )rZ   r[   rsr   r   r   r   	gu_reduce  s    z1TestCUDAGufunc.check_tuple_arg.<locals>.gu_reducerB   )Zaxis)r   r   r+   sumasarrayrD   Zassert_equal)r3   rc   rd   r   r   expectedr   r   r   check_tuple_arg  s    &

zTestCUDAGufunc.check_tuple_argc                 C   s   d}d}|  || d S )N)      ?       @      @      @      @      @)      ?      @      @      @      @      @)r   r3   rc   rd   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)rZ   r[   zr   r   r   r   r   r   r   r   r   r   r   r   )r   r   )r3   r   rc   rd   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?   rG   rI   rL   rW   rY   r_   ru   rx   r   r   r   r   r   r   r   r   r   r   r   r"   "   s&   r"   __main__)Znumpyr+   Znumpy.core.umath_testscoreZumath_testsr.   collectionsr   Znumbar   r   r   r   r   Znumba.cuda.testingr   r	   Zunittestrn   Znumba.core.errorsr
   Znumba.tests.supportr   r!   r"   r   mainr   r   r   r   <module>   s"     