o
    8ήc'(                     @   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rPe  dS 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)compileinspect_asm)selffnsig r   M/tmp/pip-target-vg8gfxp4/lib/python/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 )Nz$debuginfo not generated for NVVM 3.4)r   z\.section\s+\.debug_info\s+{)msg)	r   	is_nvvm70skipTestr   rer   searchassertIsNotNoneassertIsNone)r   r   r   expectasmre_section_dbginfomatch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      d| d< d S N   r   r   xr   r   r   foo"      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   r"   r#   r   r%   r   r   r   r'   )   r(   z4TestCudaDebugInfo.test_debuginfo_in_asm.<locals>.foor)   r*   r-   r   r   r   test_debuginfo_in_asm(   s   
 z'TestCudaDebugInfo.test_debuginfo_in_asmc                 C   s   t dd9 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 d    d S 1 sAw   Y  d S )NCUDA_DEBUGINFO_DEFAULTr$   F)r0   c                 S   r"   r#   r   r%   r   r   r   r'   2   r(   z8TestCudaDebugInfo.test_environment_override.<locals>.fooTr)   r    c                 S   r"   r#   r   r%   r   r   r   bar9   r(   z8TestCudaDebugInfo.test_environment_override.<locals>.bar)r   r   r+   r   r   r,   )r   r'   r3   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   r(   z,TestCudaDebugInfo.test_issue_5835.<locals>.fr   r+   r   r,   r   r5   r   r   r   test_issue_5835?   s   z!TestCudaDebugInfo.test_issue_5835c                 C   s   t jd d d f}tj|ddddd }||}t jr;dd | D }| t	|d |d }| 
d	| d S d
d | D }| t	|d |d }| 
d| d S )Nr$   Tr   r/   c                 S   r"   r#   r   r%   r   r   r   r5   J   r(   z7TestCudaDebugInfo.test_wrapper_has_debuginfo.<locals>.fc                 S      g | ]}d |v r|qS )zdefine void @"_ZN6cudapyr   .0liner   r   r   
<listcomp>S       
z@TestCudaDebugInfo.test_wrapper_has_debuginfo.<locals>.<listcomp>znoinline !dbgc                 S   r9   )786478r   r:   r   r   r   r=   ^   r>   
_ZN6cudapy)r   r,   r   r+   inspect_llvmr   r   
splitlinesassertEquallenassertIn)r   r   r5   llvm_irdefineswrapper_definedisubprogramswrapper_disubprogramr   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v rdnd|d< d S )Nr   )      r$   rM   r   )inpoutpr   r   r   r5   w   s   zDTestCudaDebugInfo.test_debug_function_calls_internal_impl.<locals>.fr6   r7   r   r   r   'test_debug_function_calls_internal_impli   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!   r0   c                   S   s   t jjt jj t jj S r
   )r   blockDimr&   blockIdx	threadIdxr   r   r   r   threadid   s   zMTestCudaDebugInfo.test_debug_function_calls_device_function.<locals>.threadidr/   c                    s(   t d}|t| k r  | |< d S d S Nr$   )r   gridrD   )arrirV   r   r   kernel   s   
zKTestCudaDebugInfo.test_debug_function_calls_device_function.<locals>.kernelr6   )r   r\   r   r[   r   )test_debug_function_calls_device_function{   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 )NTFrQ   c                 S      | d S rW   r   r%   r   r   r   f2      z;TestCudaDebugInfo._test_chained_device_function.<locals>.f2c                       |  | S r
   r   r&   yr_   r   r   f1   r(   z;TestCudaDebugInfo._test_chained_device_function.<locals>.f1r/   c                        | | d S r
   r   rb   re   r   r   r\         z?TestCudaDebugInfo._test_chained_device_function.<locals>.kernelr$   r$   r$   rL   r6   r   kernel_debugf1_debugf2_debugr\   r   re   r_   r   _test_chained_device_function   s   
z/TestCudaDebugInfo._test_chained_device_functionc              	   C   b   t jdgd  }|D ]$\}}}| j|||d | ||| W d    n1 s)w   Y  q
d S NTFrM   rk   rl   rm   )	itertoolsproductsubTestro   r   
debug_optsrk   rl   rm   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 )NTFrQ   c                 S   r^   rW   r   r%   r   r   r   r_      r`   zETestCudaDebugInfo._test_chained_device_function_two_calls.<locals>.f2c                    ra   r
   r   rb   rd   r   r   re      r(   zETestCudaDebugInfo._test_chained_device_function_two_calls.<locals>.f1r/   c                    s    | | |  d S r
   r   rb   rn   r   r   r\      s   
zITestCudaDebugInfo._test_chained_device_function_two_calls.<locals>.kernelri   r$   rL   r   r+   rj   r   rn   r   '_test_chained_device_function_two_calls   s   
z9TestCudaDebugInfo._test_chained_device_function_two_callsc              	   C   rp   rq   )rt   ru   rv   r{   rw   r   r   r   &test_chained_device_function_two_calls   s   z8TestCudaDebugInfo.test_chained_device_function_two_callsc                 C   sZ   t  jr| t|d d S | t|| |D ]}| |jt | dt|j	 qd S )Nr   zdebuginfo is not generated)
r   r   rC   rD   assertIscategoryr   rE   strmessage)r   warnings
warn_countwarningr   r   r   check_warnings   s   z TestCudaDebugInfo.check_warningsc              
   C   s   t jdgd  }|D ]F\}}}| j|||d1 tjdd}| ||| W d    n1 s0w   Y  || | }| || W d    n1 sKw   Y  q
d S )Nrr   rM   rs   Trecord)rt   ru   rv   r   catch_warningsr{   r   )r   rx   rk   rl   rm   wr   r   r   r   test_debug_warning   s"   z$TestCudaDebugInfo.test_debug_warningc                 C   s  dd }t jdd}|ddd W d    n1 sw   Y  | |d t jdd}|ddd W d    n1 s=w   Y  | |d t jdd}|ddd W d    n1 s_w   Y  | |d t jdd}|ddd W d    n1 sw   Y  | |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 )NTFrQ   c                 S   s   | |  S r
   r   r%   r   r   r   f3   r`   z[TestCudaDebugInfo.test_chained_device_three_functions.<locals>.three_device_fns.<locals>.f3)rR   c                    s    | d S rW   r   r%   )r   r   r   r_      r(   z[TestCudaDebugInfo.test_chained_device_three_functions.<locals>.three_device_fns.<locals>.f2c                    ra   r
   r   rb   rd   r   r   re      r(   z[TestCudaDebugInfo.test_chained_device_three_functions.<locals>.three_device_fns.<locals>.f1r/   c                    rf   r
   r   rb   rg   r   r   r\      rh   z_TestCudaDebugInfo.test_chained_device_three_functions.<locals>.three_device_fns.<locals>.kernelri   r$   rL   rz   )rk   
leaf_debugr\   r   )re   r_   r   r   three_device_fns   s   


zOTestCudaDebugInfo.test_chained_device_three_functions.<locals>.three_device_fnsTr   )rk   r   rL   Fr$   r   )r   r   r   )r   r   r   r   r   r   #test_chained_device_three_functions   s"   z5TestCudaDebugInfo.test_chained_device_three_functionsN)__name__
__module____qualname____doc__r   r   r.   r1   r4   r8   rK   rP   r]   ro   ry   r{   r|   r   r   r   r   r   r   r   r	      s$    
"r	   __main__)numba.tests.supportr   numba.cuda.testingr   numbar   numba.cuda.cudadrv.nvvmr   
numba.corer   numba.core.errorsr   r   rt   r   unittestr   r	   r   mainr   r   r   r   <module>   s$      