U
    ,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re  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   J/tmp/pip-unpacked-wheel-eu7e0c37/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_nvvm70ZskipTestr   rer   searchZassertIsNotNoneZassertIsNone)r   r   r   expectZ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d d d f}tj|ddddd }||}t jrtdd | D }| t	|d |d }| 
d	| n6d
d | D }| t	|d |d }| 
d| d S )Nr   Tr   r&   c                 S   s   d| d< d S r   r   r   r   r   r   r+   J   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noinline !dbgc                 S   s   g | ]}d |kr|qS )Z786478r   r/   r   r   r   r2   ^   s    Z
_ZN6cudapy)r   r#   r   r"   Zinspect_llvmr   r   
splitlinesassertEquallenassertIn)r   r   r+   Zllvm_irZdefinesZwrapper_defineZdisubprogramsZ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krdnd|d< d S )Nr   )      r   r9   r   )inpZoutpr   r   r   r+   w   s    zDTestCudaDebugInfo.test_debug_function_calls_internal_impl.<locals>.fr,   r-   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   r'   c                   S   s   t jjt jj t jj S r
   )r   ZblockDimr   ZblockIdxZ	threadIdxr   r   r   r   threadid   s    zMTestCudaDebugInfo.test_debug_function_calls_device_function.<locals>.threadidr&   c                    s$   t d}|t| k r   | |< d S Nr   )r   Zgridr5   )Zarrir>   r   r   kernel   s    
zKTestCudaDebugInfo.test_debug_function_calls_device_function.<locals>.kernelr,   )r   rB   r   rA   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 )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rD   r   r   f1   s    z;TestCudaDebugInfo._test_chained_device_function.<locals>.f1r&   c                    s    | | d S r
   r   rE   rH   r   r   rB      s    z?TestCudaDebugInfo._test_chained_device_function.<locals>.kernelr   r   r   r8   r,   r   kernel_debugf1_debugf2_debugrB   r   rH   rD   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TFr9   rL   rM   rN   )	itertoolsproductsubTestrP   r   
debug_optsrL   rM   rN   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   rD      s    zETestCudaDebugInfo._test_chained_device_function_two_calls.<locals>.f2c                    s   |  | S r
   r   rE   rG   r   r   rH      s    zETestCudaDebugInfo._test_chained_device_function_two_calls.<locals>.f1r&   c                    s    | | |  d S r
   r   rE   rO   r   r   rB      s    
zITestCudaDebugInfo._test_chained_device_function_two_calls.<locals>.kernelrJ   r   r8   r   r"   rK   r   rO   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 rQ   )rT   rU   rV   r[   rW   r   r   r   &test_chained_device_function_two_calls   s    z8TestCudaDebugInfo.test_chained_device_function_two_callsc                 C   sX   t  jr| t|d n:| t|| |D ]$}| |jt | dt|j	 q.d S )Nr   zdebuginfo is not generated)
r   r   r4   r5   ZassertIscategoryr   r6   strmessage)r   warnings
warn_count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 || | }| || W 5 Q R X qd S )NrR   r9   rS   Trecord)rT   rU   rV   r`   catch_warningsr[   rc   )r   rX   rL   rM   rN   wra   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   )ri   r   r   rD      s    z[TestCudaDebugInfo.test_chained_device_three_functions.<locals>.three_device_fns.<locals>.f2c                    s   |  | S r
   r   rE   rG   r   r   rH      s    z[TestCudaDebugInfo.test_chained_device_three_functions.<locals>.three_device_fns.<locals>.f1r&   c                    s    | | d S r
   r   rE   rI   r   r   rB      s    z_TestCudaDebugInfo.test_chained_device_three_functions.<locals>.three_device_fns.<locals>.kernelrJ   r   r8   rZ   )rL   
leaf_debugrB   r   )rH   rD   ri   r   three_device_fns   s    


zOTestCudaDebugInfo.test_chained_device_three_functions.<locals>.three_device_fnsTrd   )rL   rj   r8   Fr   r   )r`   rf   rc   )r   rk   rg   r   r   r   #test_chained_device_three_functions   s    z5TestCudaDebugInfo.test_chained_device_three_functionsN)__name__
__module____qualname____doc__r   r   r%   r(   r*   r.   r7   r;   rC   rP   rY   r[   r\   rc   rh   rl   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   rT   r   Zunittestr`   r	   rm   mainr   r   r   r   <module>   s      