U
    hâËd¶  ã                   @   sŒ   d dl mZmZmZ d dlmZ d dlmZ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 )é    )ÚcudaÚfloat32Úint32)ÚNumbaInvalidConfigWarning)ÚCUDATestCaseÚskip_on_cudasim)ÚNVVM)Úignore_internal_warningsNz#Simulator does not produce lineinfoc                   @   sL   e 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S )ÚTestCudaLineInfoc                 C   s   d}t  |¡S )Nz \.loc\s+[0-9]+\s+[0-9]+\s+[0-9]+)ÚreÚcompile)ÚselfÚpat© r   ú^/home/sam/Atlas/atlas_env/lib/python3.8/site-packages/numba/cuda/tests/cudapy/test_lineinfo.pyÚ_loc_directive_regex   s    ÿz%TestCudaLineInfo._loc_directive_regexc           	      C   sÐ   |  |¡ | |¡}| |¡}|r(| jn| j}d}t  |¡ |¡}|||d� d}t  |¡ |¡}| j||d� d}t  |¡ |¡}|||d� |  ¡  |¡ |||d� d}t  |¡ |¡}| j||d� d S )Nz5!DICompileUnit\(.*emissionKind:\s+DebugDirectivesOnly)Úmsgz+!DICompileUnit\(.*emissionKind:\s+FullDebugz&\.file\s+[0-9]+\s+".*test_lineinfo.py"z\.section\s+\.debug_info)r   Úinspect_llvmÚinspect_asmZassertIsNotNoneZassertIsNoner   Úsearchr   )	r   ÚfnÚsigÚexpectÚllvmÚptxZassertfnr   Úmatchr   r   r   Ú_check   s,    


ÿ	ÿÿÿzTestCudaLineInfo._checkc                 C   s2   t jdd�dd„ ƒ}| j|td d … fdd� d S )NF©Úlineinfoc                 S   s   d| d< d S ©Né   r   r   ©Úxr   r   r   ÚfooL   s    z5TestCudaLineInfo.test_no_lineinfo_in_asm.<locals>.foo©r   r   )r   Újitr   r   ©r   r#   r   r   r   Útest_no_lineinfo_in_asmK   s    

z(TestCudaLineInfo.test_no_lineinfo_in_asmc                 C   sD   t ƒ js|  d¡ tjdd�dd„ ƒ}| j|td d … fdd� d S )Nú#lineinfo not generated for NVVM 3.4Tr   c                 S   s   d| d< d S r   r   r!   r   r   r   r#   V   s    z2TestCudaLineInfo.test_lineinfo_in_asm.<locals>.foor$   )r   Ú	is_nvvm70ÚskipTestr   r%   r   r   r&   r   r   r   Útest_lineinfo_in_asmR   s
    


z%TestCudaLineInfo.test_lineinfo_in_asmc                 C   sL   t d d d… t d d d… f}tj|dd�dd„ ƒ}| |¡}|  d|¡ d S )Nr    Tr   c                 S   s   | d  |d   < d S )Nr   r   )r"   Úyr   r   r   Údivide_kernel_   s    zKTestCudaLineInfo.test_lineinfo_maintains_error_model.<locals>.divide_kernelz	ret i32 1)r   r   r%   r   ZassertNotIn)r   r   r-   r   r   r   r   Ú#test_lineinfo_maintains_error_model\   s
    

z4TestCudaLineInfo.test_lineinfo_maintains_error_modelc                    sB   t jdd„ ƒ‰ t j‡ fdd„ƒ}td d … f}| j||dd� d S )Nc                 S   s   | d  d7  < d S ©Nr   r    r   r!   r   r   r   Úcalleem   s    zDTestCudaLineInfo.test_no_lineinfo_in_device_function.<locals>.calleec                    s   d| d< ˆ | ƒ d S r   r   r!   ©r0   r   r   Úcallerq   s    zDTestCudaLineInfo.test_no_lineinfo_in_device_function.<locals>.callerFr$   )r   r%   r   r   )r   r2   r   r   r1   r   Ú#test_no_lineinfo_in_device_functionk   s    
z4TestCudaLineInfo.test_no_lineinfo_in_device_functionc                    s6  t ƒ js|  d¡ tjdd�dd„ ƒ‰ tjdd�‡ fdd„ƒ}td d … f}| j||dd� | |¡}| ¡ }t	 
d	¡}|D ]"}| |¡d k	r||  d
|› �¡ q||  ¡ }d}|D ]"}| |¡d k	r°d|kr°d} qÔq°|sè|  d|› �¡ | |¡}	d}
|	 ¡ D ]}d|krþ|
d7 }
qþd}|  |
|d|› d|
› �¡ d S )Nr(   Tr   c                 S   s   | d  d7  < d S r/   r   r!   r   r   r   r0   €   s    zATestCudaLineInfo.test_lineinfo_in_device_function.<locals>.calleec                    s   d| d< ˆ | ƒ d S r   r   r!   r1   r   r   r2   „   s    zATestCudaLineInfo.test_lineinfo_in_device_function.<locals>.callerr$   z^\.weak\s+\.funczFound device function in PTX:

FZ
inlined_atz1No .loc directive with inlined_at info foundin:

r   zdistinct !DISubprogramr    é   z
"Expected z DISubprograms; got )r   r)   r*   r   r%   r   r   r   Ú
splitlinesr   r   r   Zfailr   r   r   ÚassertEqual)r   r2   r   r   ZptxlinesZdevfn_startÚlineZloc_directiveÚfoundr   ZsubprogramsZexpected_subprogramsr   r1   r   Ú test_lineinfo_in_device_functiony   s@    







ÿz1TestCudaLineInfo.test_lineinfo_in_device_functionc              	   C   sr   t jdd��$}tƒ  tjdddd�dd„ ƒ}W 5 Q R X |  t|ƒd¡ |  |d jt¡ |  	d	t
|d jƒ¡ d S )
NT)ÚrecordF)Údebugr   Úoptc                   S   s   d S )Nr   r   r   r   r   ÚfÃ   s    z;TestCudaLineInfo.test_debug_and_lineinfo_warning.<locals>.fr    r   z)debug and lineinfo are mutually exclusive)ÚwarningsÚcatch_warningsr	   r   r%   r6   ÚlenÚcategoryr   ZassertInÚstrÚmessage)r   Úwr=   r   r   r   Útest_debug_and_lineinfo_warning½   s    ÿz0TestCudaLineInfo.test_debug_and_lineinfo_warningN)Ú__name__Ú
__module__Ú__qualname__r   r   r'   r+   r.   r3   r9   rE   r   r   r   r   r
      s   3
Dr
   Ú__main__)Znumbar   r   r   Znumba.core.errorsr   Znumba.cuda.testingr   r   Znumba.cuda.cudadrv.nvvmr   Znumba.tests.supportr	   r   Zunittestr>   r
   rF   Úmainr   r   r   r   Ú<module>   s    B