U
    hd%                     @   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   _/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