U
    -e%                     @   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 )    )override_config)skip_on_cudasim)cuda)types)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S )TestCudaDebugInfozH
    These tests only checks the compiled PTX for debuginfo section
    c                 C   s   | | ||S N)compileZinspect_asm)selffnsig r   g/var/www/html/Darija-Ai-Train/env/lib/python3.8/site-packages/numba/cuda/tests/cudapy/test_debuginfo.py_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	   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   ,   s    z8TestCudaDebugInfo.test_environment_override.<locals>.fooTr   r   c                 S   s   d| d< d S r   r   r   r   r   r   bar3   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   f=   s    z,TestCudaDebugInfo.test_issue_5835.<locals>.fr   r    r   r!   r
   r)   r   r   r   test_issue_58359   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   r$   c                 S   s   d| d< d S r   r   r   r   r   r   r)   D   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>J   s    z@TestCudaDebugInfo.test_wrapper_has_debuginfo.<locals>.<listcomp>z!dbg)	r   r!   r   r    Zinspect_llvm
splitlinesassertEquallenZassertIn)r
   r   r)   Zllvm_irZdefinesZwrapper_definer   r   r   test_wrapper_has_debuginfoA   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   r5   r   )ZinpZoutpr   r   r   r)   a   s    zDTestCudaDebugInfo.test_debug_function_calls_internal_impl.<locals>.fr*   r+   r   r   r   'test_debug_function_calls_internal_implS   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   threadidj   s    zMTestCudaDebugInfo.test_debug_function_calls_device_function.<locals>.threadidr$   c                    s$   t d}|t| k r   | |< d S Nr   )r   gridr2   )Zarrir9   r   r   kerneln   s    
zKTestCudaDebugInfo.test_debug_function_calls_device_function.<locals>.kernelr*   )r
   r>   r   r=   r   )test_debug_function_calls_device_functione   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 )NTFr7   c                 S   s   | d S r:   r   r   r   r   r   f2u   s    z;TestCudaDebugInfo._test_chained_device_function.<locals>.f2c                    s   |  | S r   r   r   yr@   r   r   f1y   s    z;TestCudaDebugInfo._test_chained_device_function.<locals>.f1r$   c                    s    | | d S r   r   rA   rD   r   r   r>   }   s    z?TestCudaDebugInfo._test_chained_device_function.<locals>.kernelr   r   r   r4   r*   r
   kernel_debugf1_debugf2_debugr>   r   rD   r@   r   _test_chained_device_functiont   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)TFr5   )rH   rI   rJ   )	itertoolsproductsubTestrL   r
   Z
debug_optsrH   rI   rJ   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 )NTFr7   c                 S   s   | d S r:   r   r   r   r   r   r@      s    zETestCudaDebugInfo._test_chained_device_function_two_calls.<locals>.f2c                    s   |  | S r   r   rA   rC   r   r   rD      s    zETestCudaDebugInfo._test_chained_device_function_two_calls.<locals>.f1r$   c                    s    | | |  d S r   r   rA   rK   r   r   r>      s    
zITestCudaDebugInfo._test_chained_device_function_two_calls.<locals>.kernelrF   r   r4   r   r    rG   r   rK   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 rM   )rN   rO   rP   rT   rQ   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 )NTFr7   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)r8   c                    s    | d S r:   r   r   )rV   r   r   r@      s    z[TestCudaDebugInfo.test_chained_device_three_functions.<locals>.three_device_fns.<locals>.f2c                    s   |  | S r   r   rA   rC   r   r   rD      s    z[TestCudaDebugInfo.test_chained_device_three_functions.<locals>.three_device_fns.<locals>.f1r$   c                    s    | | d S r   r   rA   rE   r   r   r>      s    z_TestCudaDebugInfo.test_chained_device_three_functions.<locals>.three_device_fns.<locals>.kernelrF   r   r4   rS   )rH   
leaf_debugr>   r   )rD   r@   rV   r   three_device_fns   s    


zOTestCudaDebugInfo.test_chained_device_three_functions.<locals>.three_device_fnsT)rH   rW   Fr   )r
   rX   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,   r3   r6   r?   rL   rR   rT   rU   rY   r   r   r   r   r      s   r   __main__)Znumba.tests.supportr   Znumba.cuda.testingr   Znumbar   Z
numba.corer   r   rN   r   Zunittestr   rZ   mainr   r   r   r   <module>   s    G