U
    hâË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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rÚe ¡  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 )Néd   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&     ;\