U
    hâËdÄ%  ã                   @   s    d dl mZ d dlmZ d dlmZ d dlmZ d dlm	Z	 d dl
mZ d dlmZ d dlZd dlZd dlZd dlZed	ƒG d
d„ deƒƒZedkrœe ¡  dS )é    )Úoverride_config)Úskip_on_cudasim)Úcuda)ÚNVVM)Útypes)ÚNumbaInvalidConfigWarning)ÚCUDATestCaseNz&Simulator does not produce debug dumpsc                   @   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dd„ Zdd„ Zdd„ Zdd„ Zd d!„ Zd"S )#ÚTestCudaDebugInfozH
    These tests only checks the compiled PTX for debuginfo section
    c                 C   s   |  |¡ | |¡S ©N)ÚcompileZinspect_asm)ÚselfÚfnÚsig© r   ú_/home/sam/Atlas/atlas_env/lib/python3.8/site-packages/numba/cuda/tests/cudapy/test_debuginfo.pyÚ_getasm   s    
zTestCudaDebugInfo._getasmc                 C   sT   t ƒ js|  d¡ | j||d�}t d¡}| |¡}|r>| jn| j}|||d� d S )Nú$debuginfo not generated for NVVM 3.4)r   z\.section\s+\.debug_info\s+{)Úmsg)	r   Ú	is_nvvm70ÚskipTestr   Úrer   ÚsearchZassertIsNotNoneZassertIsNone)r   r   r   ÚexpectÚasmZre_section_dbginfoÚmatchZassertfnr   r   r   Ú_check   s    


zTestCudaDebugInfo._checkc                 C   s4   t jdd�dd„ ƒ}| j|tjd d … fdd� d S )NF©Údebugc                 S   s   d| d< d S ©Né   r   r   ©Úxr   r   r   Úfoo"   s    z7TestCudaDebugInfo.test_no_debuginfo_in_asm.<locals>.foo©r   r   ©r   Újitr   r   Úint32©r   r"   r   r   r   Útest_no_debuginfo_in_asm!   s    

z*TestCudaDebugInfo.test_no_debuginfo_in_asmc                 C   s6   t jddd�dd„ ƒ}| j|tjd d … fdd� d S )NTF©r   Úoptc                 S   s   d| d< d S r   r   r    r   r   r   r"   )   s    z4TestCudaDebugInfo.test_debuginfo_in_asm.<locals>.foor#   r$   r'   r   r   r   Útest_debuginfo_in_asm(   s    
z'TestCudaDebugInfo.test_debuginfo_in_asmc              	   C   sz   t ddƒ�f tjdd�dd„ ƒ}| j|tjd d … fdd� tjdd	�d
d„ ƒ}| j|tjd d … fdd� W 5 Q R X d S )NZCUDA_DEBUGINFO_DEFAULTr   F)r*   c                 S   s   d| d< d S r   r   r    r   r   r   r"   2   s    z8TestCudaDebugInfo.test_environment_override.<locals>.fooTr#   r   c                 S   s   d| d< d S r   r   r    r   r   r   Úbar9   s    z8TestCudaDebugInfo.test_environment_override.<locals>.bar)r   r   r%   r   r   r&   )r   r"   r,   r   r   r   Útest_environment_override/   s    



z+TestCudaDebugInfo.test_environment_overridec                 C   s*   t jtjd d d… fddd�dd„ ƒ}d S )Nr   TFr)   c                 S   s   d| d< d S )Nr   r   r    r   r   r   ÚfC   s    z,TestCudaDebugInfo.test_issue_5835.<locals>.f©r   r%   r   r&   ©r   r.   r   r   r   Útest_issue_5835?   s    z!TestCudaDebugInfo.test_issue_5835c                 C   s€   t ƒ js|  d¡ tjd d d… f}tj|ddd�dd„ ƒ}| |¡}dd	„ | ¡ D ƒ}|  	t
|ƒd¡ |d }|  d
|¡ d S )Nr   r   Tr   r)   c                 S   s   d| d< d S r   r   r    r   r   r   r.   M   s    z7TestCudaDebugInfo.test_wrapper_has_debuginfo.<locals>.fc                 S   s   g | ]}d |kr|‘qS )zdefine void @"_ZN6cudapyr   )Ú.0Úliner   r   r   Ú
<listcomp>S   s    ÿz@TestCudaDebugInfo.test_wrapper_has_debuginfo.<locals>.<listcomp>z!dbg)r   r   r   r   r&   r   r%   Zinspect_llvmÚ
splitlinesÚassertEqualÚlenÚassertIn)r   r   r.   Zllvm_irZdefinesZwrapper_definer   r   r   Útest_wrapper_has_debuginfoG   s    


z,TestCudaDebugInfo.test_wrapper_has_debuginfoc                 C   s4   t jtjd d … tjd d … fddd�dd„ ƒ}d S )NTFr)   c                 S   s   | d dkrdnd|d< d S )Nr   )é   é   r   r;   r   )ZinpZoutpr   r   r   r.   j   s    zDTestCudaDebugInfo.test_debug_function_calls_internal_impl.<locals>.fr/   r0   r   r   r   Ú'test_debug_function_calls_internal_impl\   s    &z9TestCudaDebugInfo.test_debug_function_calls_internal_implc                    sD   t jdddd�dd„ ƒ‰ t jtjd d … fddd�‡ fdd„ƒ}d S )	NTr   ©Údevicer   r*   c                   S   s   t jjt jj t jj S r
   )r   ZblockDimr!   ZblockIdxZ	threadIdxr   r   r   r   Úthreadids   s    zMTestCudaDebugInfo.test_debug_function_calls_device_function.<locals>.threadidr)   c                    s$   t  d¡}|t| ƒk r ˆ ƒ | |< d S ©Nr   )r   Úgridr7   )ZarrÚi©r?   r   r   Úkernelw   s    
zKTestCudaDebugInfo.test_debug_function_calls_device_function.<locals>.kernelr/   )r   rD   r   rC   r   Ú)test_debug_function_calls_device_functionn   s    
z;TestCudaDebugInfo.test_debug_function_calls_device_functionc                    sj   t jd|dd�dd„ ƒ‰t jd|dd�‡fdd„ƒ‰ t jtjtjf|dd�‡ fd	d
„ƒ}|d ddƒ d S )NTFr=   c                 S   s   | d S r@   r   r    r   r   r   Úf2~   s    z;TestCudaDebugInfo._test_chained_device_function.<locals>.f2c                    s   | ˆ |ƒ S r
   r   ©r!   Úy©rF   r   r   Úf1‚   s    z;TestCudaDebugInfo._test_chained_device_function.<locals>.f1r)   c                    s   ˆ | |ƒ d S r
   r   rG   ©rJ   r   r   rD   †   s    z?TestCudaDebugInfo._test_chained_device_function.<locals>.kernel©r   r   r   r:   r/   ©r   Úkernel_debugÚf1_debugÚf2_debugrD   r   ©rJ   rF   r   Ú_test_chained_device_function}   s    
z/TestCudaDebugInfo._test_chained_device_functionc              
   C   sN   t jdgd Ž }|D ]4\}}}| j|||d�� |  |||¡ W 5 Q R X qd S ©N©TFr;   ©rN   rO   rP   )Ú	itertoolsÚproductÚsubTestrR   ©r   Ú
debug_optsrN   rO   rP   r   r   r   Útest_chained_device_functionŒ   s    þþz.TestCudaDebugInfo.test_chained_device_functionc                    sb   t jd|dd�dd„ ƒ‰t jd|dd�‡fdd„ƒ‰ t j|dd�‡ ‡fd	d
„ƒ}|d ddƒ d S )NTFr=   c                 S   s   | d S r@   r   r    r   r   r   rF   ž   s    zETestCudaDebugInfo._test_chained_device_function_two_calls.<locals>.f2c                    s   | ˆ |ƒ S r
   r   rG   rI   r   r   rJ   ¢   s    zETestCudaDebugInfo._test_chained_device_function_two_calls.<locals>.f1r)   c                    s   ˆ | |ƒ ˆ| ƒ d S r
   r   rG   rQ   r   r   rD   ¦   s    
zITestCudaDebugInfo._test_chained_device_function_two_calls.<locals>.kernelrL   r   r:   ©r   r%   rM   r   rQ   r   Ú'_test_chained_device_function_two_calls›   s    
z9TestCudaDebugInfo._test_chained_device_function_two_callsc              
   C   sN   t jdgd Ž }|D ]4\}}}| j|||d�� |  |||¡ W 5 Q R X qd S rS   )rV   rW   rX   r]   rY   r   r   r   Ú&test_chained_device_function_two_calls­   s    þþz8TestCudaDebugInfo.test_chained_device_function_two_callsc                 C   s\   t ƒ jr|  t|ƒd¡ n>|r.|  t|ƒd¡ |D ]$}|  |jt¡ |  dt	|j
ƒ¡ q2d S )Nr   zdebuginfo is not generated)r   r   r6   r7   ZassertGreaterZassertIsÚcategoryr   r8   ÚstrÚmessage)r   ÚwarningsÚwarning_expectedÚwarningr   r   r   Úcheck_warnings½   s    ÿz TestCudaDebugInfo.check_warningsc                 C   s~   t jdgd Ž }|D ]d\}}}| j|||d��D tjdd��}|  |||¡ W 5 Q R X |p`|p`|}|  ||¡ W 5 Q R X qd S )NrT   r;   rU   T©Úrecord)rV   rW   rX   rb   Úcatch_warningsr]   re   )r   rZ   rN   rO   rP   Úwrc   r   r   r   Útest_debug_warningÊ   s    þþz$TestCudaDebugInfo.test_debug_warningc              	   C   sÌ   dd„ }t jdd��}|ddd� W 5 Q R X |  |d¡ t jdd��}|ddd� W 5 Q R X |  |d¡ t jdd��}|ddd� W 5 Q R X |  |d¡ t jdd��}|ddd� W 5 Q R X |  |d	¡ d S )
Nc                    st   t jd|dd�dd„ ƒ‰t jdd�‡fdd„ƒ‰t jdd�‡fd	d
„ƒ‰ t j| dd�‡ fdd„ƒ}|d ddƒ d S )NTFr=   c                 S   s   | |  S r
   r   r    r   r   r   Úf3à   s    z[TestCudaDebugInfo.test_chained_device_three_functions.<locals>.three_device_fns.<locals>.f3)r>   c                    s   ˆ | ƒd S r@   r   r    )rk   r   r   rF   ä   s    z[TestCudaDebugInfo.test_chained_device_three_functions.<locals>.three_device_fns.<locals>.f2c                    s   | ˆ |ƒ S r
   r   rG   rI   r   r   rJ   è   s    z[TestCudaDebugInfo.test_chained_device_three_functions.<locals>.three_device_fns.<locals>.f1r)   c                    s   ˆ | |ƒ d S r
   r   rG   rK   r   r   rD   ì   s    z_TestCudaDebugInfo.test_chained_device_three_functions.<locals>.three_device_fns.<locals>.kernelrL   r   r:   r\   )rN   Ú
leaf_debugrD   r   )rJ   rF   rk   r   Úthree_device_fnsß   s    


zOTestCudaDebugInfo.test_chained_device_three_functions.<locals>.three_device_fnsTrf   )rN   rl   r:   Fr   r   )rb   rh   re   )r   rm   ri   r   r   r   Ú#test_chained_device_three_functionsÛ   s    z5TestCudaDebugInfo.test_chained_device_three_functionsN)Ú__name__Ú
__module__Ú__qualname__Ú__doc__r   r   r(   r+   r-   r1   r9   r<   rE   rR   r[   r]   r^   re   rj   rn   r   r   r   r   r	      s"   
r	   Ú__main__)Znumba.tests.supportr   Znumba.cuda.testingr   Znumbar   Znumba.cuda.cudadrv.nvvmr   Z
numba.corer   Znumba.core.errorsr   r   rV   r   Zunittestrb   r	   ro   Úmainr   r   r   r   Ú<module>   s    w