U
    9%e                     @   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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)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   d/var/www/html/Darija-Ai-API/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   fooK   s    z5TestCudaLineInfo.test_no_lineinfo_in_asm.<locals>.foor   r   r   jitr   r   r   r"   r   r   r   test_no_lineinfo_in_asmJ   s    

z(TestCudaLineInfo.test_no_lineinfo_in_asmc                 C   s2   t jdddd }| j|td d  fdd d S )NTr   c                 S   s   d| d< d S r   r   r    r   r   r   r"   R   s    z2TestCudaLineInfo.test_lineinfo_in_asm.<locals>.foor#   r$   r&   r   r   r   test_lineinfo_in_asmQ   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_modelX   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   calleei   s    zDTestCudaLineInfo.test_no_lineinfo_in_device_function.<locals>.calleec                    s   d| d<  |  d S r   r   r    r-   r   r   callerm   s    zDTestCudaLineInfo.test_no_lineinfo_in_device_function.<locals>.callerFr#   )r   r%   r   r   )r   r/   r   r   r.   r   #test_no_lineinfo_in_device_functiong   s    
z4TestCudaLineInfo.test_no_lineinfo_in_device_functionc                    s$  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	rj| 	d	|  qj| 
 }d
}|D ]"}||d k	rd|krd} qq|s| 	d|  ||}	d}
|	 D ]}d|kr|
d7 }
qd}| |
|d| d|
  d S )NTr   c                 S   s   | d  d7  < d S r,   r   r    r   r   r   r-   y   s    zATestCudaLineInfo.test_lineinfo_in_device_function.<locals>.calleec                    s   d| d<  |  d S r   r   r    r.   r   r   r/   }   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   
splitlinesr
   r   r   Zfailr   r   r   assertEqual)r   r/   r   r   ZptxlinesZdevfn_startlineZloc_directivefoundr   ZsubprogramsZexpected_subprogramsr   r.   r    test_lineinfo_in_device_functionu   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%   r3   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+   r0   r6   rB   r   r   r   r   r	   
   s   3Ar	   __main__)Znumbar   r   r   Znumba.core.errorsr   Znumba.cuda.testingr   r   Znumba.tests.supportr   r
   Zunittestr;   r	   rC   mainr   r   r   r   <module>   s    <