o
    Y۷i                     @   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	Z	d dl
Z
d dlZedG dd	 d	eZed
kr@e  dS dS )    )override_config)skip_on_cudasim)cuda)types)CUDATestCaseNz&Simulator does not produce debug dumpsc                       s   e Zd ZdZ f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  ZS ) TestCudaDebugInfozH
    These tests only checks the compiled PTX for debuginfo section
    c                    s   t    | d d S )Nz!Exceptions not supported with LTO)supersetUpskip_if_lto)self	__class__ \/home/ubuntu/vllm_env/lib/python3.10/site-packages/numba/cuda/tests/cudapy/test_debuginfo.pyr	      s   
zTestCudaDebugInfo.setUpc                 C   s   | | ||S N)compileinspect_asm)r   fnsigr   r   r   _getasm   s   

zTestCudaDebugInfo._getasmc                 C   sB   | j ||d}td}||}|r| jn| j}|||d d S )N)r   z\.section\s+\.debug_info\s+{)msg)r   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)r1   c                 S   r#   r$   r   r&   r   r   r   r(   6   r)   z8TestCudaDebugInfo.test_environment_override.<locals>.fooTr*   r!   c                 S   r#   r$   r   r&   r   r   r   bar=   r)   z8TestCudaDebugInfo.test_environment_override.<locals>.bar)r   r   r,   r    r   r-   )r   r(   r4   r   r   r   test_environment_override3   s   



"z+TestCudaDebugInfo.test_environment_overridec                 C   s*   t jtjd d d fddddd }d S )Nr%   TFr0   c                 S   s   d| d< d S )Nr   r   r&   r   r   r   fG   r)   z,TestCudaDebugInfo.test_issue_5835.<locals>.fr   r,   r   r-   r   r6   r   r   r   test_issue_5835C   s   z!TestCudaDebugInfo.test_issue_5835c                 C   sn   t jd d d f}tj|ddddd }||}dd | D }| t|d |d }| d	| d S )
Nr%   Tr   r0   c                 S   r#   r$   r   r&   r   r   r   r6   N   r)   z7TestCudaDebugInfo.test_wrapper_has_debuginfo.<locals>.fc                 S   s   g | ]}d |v r|qS )zdefine void @"_ZN6cudapyr   ).0liner   r   r   
<listcomp>T   s    z@TestCudaDebugInfo.test_wrapper_has_debuginfo.<locals>.<listcomp>z!dbg)	r   r-   r   r,   inspect_llvm
splitlinesassertEquallenassertIn)r   r   r6   llvm_irdefineswrapper_definer   r   r   test_wrapper_has_debuginfoK   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 )NTFr0   c                 S   s   | d dv rdnd|d< d S )Nr   )      r%   rG   r   )inpoutpr   r   r   r6   k   s   zDTestCudaDebugInfo.test_debug_function_calls_internal_impl.<locals>.fr7   r8   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"   r1   c                   S   s   t jjt jj t jj S r   )r   blockDimr'   blockIdx	threadIdxr   r   r   r   threadidt   s   zMTestCudaDebugInfo.test_debug_function_calls_device_function.<locals>.threadidr0   c                    s(   t d}|t| k r  | |< d S d S Nr%   )r   gridr@   )arrirP   r   r   kernelx   s   
zKTestCudaDebugInfo.test_debug_function_calls_device_function.<locals>.kernelr7   )r   rV   r   rU   r   )test_debug_function_calls_device_functiono   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 )NTFrK   c                 S      | d S rQ   r   r&   r   r   r   f2      z;TestCudaDebugInfo._test_chained_device_function.<locals>.f2c                       |  | S r   r   r'   yrY   r   r   f1   r)   z;TestCudaDebugInfo._test_chained_device_function.<locals>.f1r0   c                        | | d S r   r   r\   r_   r   r   rV         z?TestCudaDebugInfo._test_chained_device_function.<locals>.kernelr%   r%   r%   rF   r7   r   kernel_debugf1_debugf2_debugrV   r   r_   rY   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)TFrG   )re   rf   rg   )	itertoolsproductsubTestri   r   
debug_optsre   rf   rg   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 )NTFrK   c                 S   rX   rQ   r   r&   r   r   r   rY      rZ   zETestCudaDebugInfo._test_chained_device_function_two_calls.<locals>.f2c                    r[   r   r   r\   r^   r   r   r_      r)   zETestCudaDebugInfo._test_chained_device_function_two_calls.<locals>.f1r0   c                    s    | | |  d S r   r   r\   rh   r   r   rV      s   
zITestCudaDebugInfo._test_chained_device_function_two_calls.<locals>.kernelrc   r%   rF   r   r,   rd   r   rh   r   '_test_chained_device_function_two_calls   s   
z9TestCudaDebugInfo._test_chained_device_function_two_callsc              	   C   rj   rk   )rl   rm   rn   rs   ro   r   r   r   &test_chained_device_function_two_calls   s   z8TestCudaDebugInfo.test_chained_device_function_two_callsc                 C   s<   dd }|ddd |ddd |ddd |dd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 )NTFrK   c                 S   s   | |  S r   r   r&   r   r   r   f3   rZ   z[TestCudaDebugInfo.test_chained_device_three_functions.<locals>.three_device_fns.<locals>.f3)rL   c                    s    | d S rQ   r   r&   )ru   r   r   rY      r)   z[TestCudaDebugInfo.test_chained_device_three_functions.<locals>.three_device_fns.<locals>.f2c                    r[   r   r   r\   r^   r   r   r_      r)   z[TestCudaDebugInfo.test_chained_device_three_functions.<locals>.three_device_fns.<locals>.f1r0   c                    r`   r   r   r\   ra   r   r   rV      rb   z_TestCudaDebugInfo.test_chained_device_three_functions.<locals>.three_device_fns.<locals>.kernelrc   r%   rF   rr   )re   
leaf_debugrV   r   )r_   rY   ru   r   three_device_fns   s   


zOTestCudaDebugInfo.test_chained_device_three_functions.<locals>.three_device_fnsT)re   rv   Fr   )r   rw   r   r   r   #test_chained_device_three_functions   s
   z5TestCudaDebugInfo.test_chained_device_three_functions)__name__
__module____qualname____doc__r	   r   r    r/   r2   r5   r9   rE   rJ   rW   ri   rq   rs   rt   rx   __classcell__r   r   r   r   r      s"    	r   __main__)numba.tests.supportr   numba.cuda.testingr   numbar   
numba.corer   r   rl   r   unittestr   ry   mainr   r   r   r   <module>   s     Q