U
    hd;                     @   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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fd	d
ZedG dd deZedG dd deZedkre  dS )    N)
namedtuple)voidint32float32float64)guvectorize)cuda)skip_on_cudasimCUDATestCase)NumbaPerformanceWarning)override_configc                 C   sT   t t| d d d d f | d d d d f | d d d d f gddddd }|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   \/home/sam/Atlas/atlas_env/lib/python3.8/site-packages/numba/cuda/tests/cudapy/test_gufunc.py
matmulcore   s    

z*_get_matmulcore_gufunc.<locals>.matmulcore)r   r   )dtyper   r   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'd( Zd)d* Zd+d, Zd-S ).TestCUDAGufuncc                 C   sz   t  }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    nparanger   reshapeutmatrix_multiply
assertTrueallcloseselfgufunc	matrix_ctr   r   r   Goldr   r   r   test_gufunc_small"   s    
z TestCUDAGufunc.test_gufunc_smallc                 C   s   t  }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,   r-   r.   )r0   r1   r2   r   r   dBr   r3   r   r   r   test_gufunc_auto_transfer0   s    
z(TestCUDAGufunc.test_gufunc_auto_transferc                 C   sz   t  }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%   r&   r'   r/   r   r   r   test_gufunc@   s    
zTestCUDAGufunc.test_gufuncc                 C   s~   t  }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 )Nd   r#   r$   r%      r&   r'   r/   r   r   r   test_gufunc_hidimN   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>      r@   )
r    r   r(   randomZrandnr+   r,   testingassert_allcloseZtile)r0   r1   XYZgoldZres1Zres2r   r   r   test_gufunc_new_axisZ   s    

z#TestCUDAGufunc.test_gufunc_new_axisc                 C   s   t  }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 )
Nr9   r#   r$   r%   r&   )r9   r#   r&   )r   r   stream)outrG   )rG   )r    r(   r)   r   r*   r   rG   r5   Zdevice_arrayr   r6   Zsynchronizer+   r,   r-   r.   )r0   r1   r2   r   r   rG   ZdAr7   ZdCr   r3   r   r   r   test_gufunc_streami   s"    z!TestCUDAGufunc.test_gufunc_streamc                 C   sj   t ttd d  td d  gddddd }tjdtjdd }t|}|||d	 tj|| 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%   r@   rH   r   r   r   r(   r)   
zeros_likerB   rC   r0   rO   r   r   r   r   r   	test_copy   s    

zTestCUDAGufunc.test_copyc                 C   sl   t td d  td d  fgddddd }tjdtjdd }t|}|||d	 | t|| d S )
NrJ   r   r   c                 S   s    t |jD ]}| | ||< q
d S rK   rL   rN   r   r   r   rO      s    z9TestCUDAGufunc.test_copy_unspecified_return.<locals>.copyr>   r%   r@   rP   )r   r   r(   r)   rR   r-   r.   rS   r   r   r   test_copy_unspecified_return   s    

z+TestCUDAGufunc.test_copy_unspecified_returnc                 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 )
NrJ   r   r   c                 S   s    t |jD ]}| | ||< q
d S rK   rL   rN   r   r   r   rO      s    z*TestCUDAGufunc.test_copy_odd.<locals>.copy   r%   r@   rP   )r   r   r   r(   r)   rR   r-   r.   rS   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   r@   )r   r   )r   r   xyr   r   r   copy2d   s    z*TestCUDAGufunc.test_copy2d.<locals>.copy2d   r%   r&      r@   rP   )	r   r   r   r(   r)   r*   rR   r-   r.   )r0   rZ   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_WARNINGSr@   Trecordz	Grid sizezlow occupancy)r   r(   rA   randastypezerosr   r   warningscatch_warningsassertEqualcategoryr   assertInstrmessage)r0   re   ra   rb   rc   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^   r_   Tr   )nopythonr   c                 S   s0   | j d }t|D ]}| | ||  ||< qd S r   r   r`   r   r   r   numba_dist_cuda2   s    
zLTestCUDAGufunc.test_efficient_launch_configuration.<locals>.numba_dist_cuda2i   r   )i   r#   rf   r@   rg   r   )r   r(   rA   ri   rj   r*   rR   r   rl   rm   rn   rd   )r0   rv   ra   rb   rc   rs   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 rK   r   r   r   r   r   r   foo   s    z.TestCUDAGufunc.test_nopython_flag.<locals>.foorJ   r   T)r   ru   Fznopython flag must be True)r   r   r   assertRaises	TypeErrorrn   rq   	exception)r0   ry   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 rK   r   rx   r   r   r   ry      s    z.TestCUDAGufunc.test_invalid_flags.<locals>.foorJ   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   )rz   r{   r   r   r   rq   r|   rn   rd   r   splitset)r0   ry   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 )
NrJ   r   r   c                 S   s   d S rK   r   )inprH   r   r   r   ry     s    z2TestCUDAGufunc.test_duplicated_output.<locals>.foor>   r%   rP   z<cannot specify argument 'out' as both positional and keyword)
r   r   r   r(   rk   rz   
ValueErrorrn   rq   r|   )r0   ry   r   rH   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   rd   )rX   rY   rsr   r   r   r   	gu_reduce  s    z1TestCUDAGufunc.check_tuple_arg.<locals>.gu_reducer@   )Zaxis)r   r   r(   sumasarrayrB   Zassert_equal)r0   ra   rb   r   r   expectedr   r   r   check_tuple_arg  s    &

zTestCUDAGufunc.check_tuple_argc                 C   s   d}d}|  || d S )N)      ?       @      @      @      @      @)      ?      @      @      @      @      @)r   r0   ra   rb   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rX   rY   zr   r   r   r   r   r   r   r   r   r   r   r   )r   r   )r0   r   ra   rb   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_argc                 C   s   t  }| |jd d S )Nr   )r    rn   __name__)r0   r1   r   r   r   test_gufunc_name0  s    zTestCUDAGufunc.test_gufunc_namec              	   C   sj   |  t4}tttd d  td d  gddddd }W 5 Q R X t|j}| d| | d| d S )Nz(m)->(m)r   r   c                 S   s   d S rK   r   )rX   rY   r   r   r   f6  s    z.TestCUDAGufunc.test_bad_return_type.<locals>.fz+guvectorized functions cannot return valueszspecifies int32 return type)rz   r{   r   r   rq   r|   rp   )r0   ter   r   r   r   r   test_bad_return_type4  s    $
z#TestCUDAGufunc.test_bad_return_typec              	   C   s   t td d  td d  td d  fgddddd }td}| t}|| W 5 Q R X t|j}| d| | d| | d	| | t}||||| W 5 Q R X t|j}| d| | d| | d
| d S )Nz(m),(m)->(m)r   r   c                 S   s   d S rK   r   r   r   r   r   r   ?  s    z;TestCUDAGufunc.test_incorrect_number_of_pos_args.<locals>.fr&   %gufunc accepts 2 positional argumentszor 3 positional argumentsGot 1 positional argument.zGot 4 positional arguments.	r   r   r(   r)   rz   r{   rq   r|   rp   r0   r   Zarrr   r   r   r   r   !test_incorrect_number_of_pos_args>  s$    $ 



z0TestCUDAGufunc.test_incorrect_number_of_pos_argsN)r   
__module____qualname__r4   r8   r:   r=   rF   rI   rT   rU   rW   r]   rt   rw   r~   r   r   r   r   r   r   r   r   r   r   r   r   r   r!      s,   
r!   c                   @   s4   e Zd Zdd Zdd Zdd Zdd Zd	d
 ZdS )TestMultipleOutputsc                 C   s   t ttd d  td d  td d  gddddd }tjdtjdd }t|}t|}|||| tj|| tj|| d S )	N(x)->(x),(x)r   r   c                 S   s,   t |jD ]}| | ||< | | ||< q
d S rK   rL   r   r   r   r   r   r   r   rO   \  s    zKTestMultipleOutputs.test_multiple_outputs_same_type_passed_in.<locals>.copyr>   r%   r@   rQ   )r0   rO   r   r   r   r   r   r   )test_multiple_outputs_same_type_passed_in[  s    &


z=TestMultipleOutputs.test_multiple_outputs_same_type_passed_inc                 C   s   t ttd d  td d  td d  gddddd }tjdtjdd }t|}t|}|||| tj|| tj|d	 | d S )
Nr   r   r   c                 S   s0   t |jD ] }| | ||< | | d ||< q
d S Nr#   rL   r   r   r   r   copy_and_doublem  s    zRTestMultipleOutputs.test_multiple_outputs_distinct_values.<locals>.copy_and_doubler>   r%   r@   r#   rQ   r0   r   r   r   r   r   r   r   %test_multiple_outputs_distinct_valuesk  s    &


z9TestMultipleOutputs.test_multiple_outputs_distinct_valuesc                 C   s|   t ttd d  td d  td d  gddddd }tjdtjdd }||\}}tj|| tj|d	 | d S )
Nr   r   r   c                 S   s0   t |jD ] }| | ||< | | d ||< q
d S r   rL   r   r   r   r   r   }  s    zLTestMultipleOutputs.test_multiple_output_allocation.<locals>.copy_and_doubler>   r%   r@   r#   )r   r   r   r(   r)   rB   rC   r   r   r   r   test_multiple_output_allocation|  s    &
z3TestMultipleOutputs.test_multiple_output_allocationc                 C   s   t ttd d  td d  td d  gddddd }tjdtjdd }t|}tj|tjd}|||| tj|| tj|td	 | d S )
Nr   r   r   c                 S   s0   t |jD ] }| | ||< | | d ||< q
d S )Nr   rL   r   r   r   r   copy_and_multiply  s    zJTestMultipleOutputs.test_multiple_output_dtypes.<locals>.copy_and_multiplyr>   r%   r@   r   )	r   r   r   r   r(   r)   rR   rB   rC   )r0   r   r   r   r   r   r   r   test_multiple_output_dtypes  s    &

z/TestMultipleOutputs.test_multiple_output_dtypesc              	   C   s   t td d  td d  td d  td d  fgddddd }td}| t}|| W 5 Q R X t|j}| d| | d| | d	| | t}|||||| W 5 Q R X t|j}| d| | d| | d
| d S )Nz(m),(m)->(m),(m)r   r   c                 S   s   d S rK   r   )rX   rY   r   rs   r   r   r   r     s    z@TestMultipleOutputs.test_incorrect_number_of_pos_args.<locals>.fr&   r   zor 4 positional argumentsr   zGot 5 positional arguments.r   r   r   r   r   r     s$    . 



z5TestMultipleOutputs.test_incorrect_number_of_pos_argsN)r   r   r   r   r   r   r   r   r   r   r   r   r   Y  s
   r   __main__)numpyr(   Znumpy.core.umath_testscoreZumath_testsr+   collectionsr   Znumbar   r   r   r   r   r   Znumba.cuda.testingr	   r
   Zunittestrl   Znumba.core.errorsr   Znumba.tests.supportr   r    r!   r   r   mainr   r   r   r   <module>   s&     ;\