a
    
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   o/var/www/html/stable-diffusion-webui/venv/lib/python3.9/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   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   s   t ddp 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    n1 s0    Y  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 |v 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v rdnd|d< d S )Nr   )      r   r:   r   )inpZoutpr   r   r   r-   j   s    zDTestCudaDebugInfo.test_debug_function_calls_internal_impl.<locals>.fr.   r/   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   Zgridr6   )Zarrir?   r   r   kernelw   s    
zKTestCudaDebugInfo.test_debug_function_calls_device_function.<locals>.kernelr.   )r   rC   r   rB   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rE   r   r   f1   s    z;TestCudaDebugInfo._test_chained_device_function.<locals>.f1r(   c                    s    | | d S r
   r   rF   rI   r   r   rC      s    z?TestCudaDebugInfo._test_chained_device_function.<locals>.kernelr   r   r   r9   r.   r   kernel_debugf1_debugf2_debugrC   r   rI   rE   r   _test_chained_device_function}   s    
z/TestCudaDebugInfo._test_chained_device_functionc              	   C   sb   t jdgd  }|D ]H\}}}| j|||d | ||| W d    q1 sR0    Y  qd S NTFr:   rM   rN   rO   )	itertoolsproductsubTestrQ   r   
debug_optsrM   rN   rO   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   rE      s    zETestCudaDebugInfo._test_chained_device_function_two_calls.<locals>.f2c                    s   |  | S r
   r   rF   rH   r   r   rI      s    zETestCudaDebugInfo._test_chained_device_function_two_calls.<locals>.f1r(   c                    s    | | |  d S r
   r   rF   rP   r   r   rC      s    
zITestCudaDebugInfo._test_chained_device_function_two_calls.<locals>.kernelrK   r   r9   r   r$   rL   r   rP   r   '_test_chained_device_function_two_calls   s    
z9TestCudaDebugInfo._test_chained_device_function_two_callsc              	   C   sb   t jdgd  }|D ]H\}}}| j|||d | ||| W d    q1 sR0    Y  qd S rR   )rU   rV   rW   r\   rX   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   r5   r6   ZassertGreaterZassertIscategoryr   r7   strmessage)r   warningswarning_expectedwarningr   r   r   check_warnings   s    z TestCudaDebugInfo.check_warningsc              
   C   s   t jdgd  }|D ]\}}}| j|||db tjdd}| ||| W d    n1 s`0    Y  |pt|pt|}| || W d    q1 s0    Y  qd S )NrS   r:   rT   Trecord)rU   rV   rW   ra   catch_warningsr\   rd   )r   rY   rM   rN   rO   wrb   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 s60    Y  | |d t jdd}|ddd W d    n1 sz0    Y  | |d t jdd}|ddd W d    n1 s0    Y  | |d t jdd}|ddd W d    n1 s0    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 )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   )rj   r   r   rE      s    z[TestCudaDebugInfo.test_chained_device_three_functions.<locals>.three_device_fns.<locals>.f2c                    s   |  | S r
   r   rF   rH   r   r   rI      s    z[TestCudaDebugInfo.test_chained_device_three_functions.<locals>.three_device_fns.<locals>.f1r(   c                    s    | | d S r
   r   rF   rJ   r   r   rC      s    z_TestCudaDebugInfo.test_chained_device_three_functions.<locals>.three_device_fns.<locals>.kernelrK   r   r9   r[   )rM   
leaf_debugrC   r   )rI   rE   rj   r   three_device_fns   s    


zOTestCudaDebugInfo.test_chained_device_three_functions.<locals>.three_device_fnsTre   )rM   rk   r9   Fr   r   )ra   rg   rd   )r   rl   rh   r   r   r   #test_chained_device_three_functions   s    ***,z5TestCudaDebugInfo.test_chained_device_three_functionsN)__name__
__module____qualname____doc__r   r   r'   r*   r,   r0   r8   r<   rD   rQ   rZ   r\   r]   rd   ri   rm   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   rU   r   Zunittestra   r	   rn   mainr   r   r   r   <module>   s    w