U
    hd                     @   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re  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	rd|krd} qq|s| d|  ||}	d}
|	 D ]}d|kr|
d7 }
qd}| |
|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