o
    Y۷i                     @   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 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   [/home/ubuntu/vllm_env/lib/python3.10/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_asmassertIsNotNoneassertIsNoner
   searchr   )	r   fnsigexpectllvmptxassertfnr   matchr   r   r   _check   s,   


	zTestCudaLineInfo._checkc                 C   2   t jdddd }| j|td d  fdd d S )NFlineinfoc                 S      d| d< d S N   r   r   xr   r   r   fooK      z5TestCudaLineInfo.test_no_lineinfo_in_asm.<locals>.foor   r   r   jitr   r   r   r'   r   r   r   test_no_lineinfo_in_asmJ      

z(TestCudaLineInfo.test_no_lineinfo_in_asmc                 C   r   )NTr    c                 S   r"   r#   r   r%   r   r   r   r'   R   r(   z2TestCudaLineInfo.test_lineinfo_in_asm.<locals>.foor)   r*   r,   r   r   r   test_lineinfo_in_asmQ   r.   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   assertNotIn)r   r   r1   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      | d  d7  < d S Nr   r$   r   r%   r   r   r   calleei      zDTestCudaLineInfo.test_no_lineinfo_in_device_function.<locals>.calleec                       d| d<  |  d S r#   r   r%   r6   r   r   callerm      zDTestCudaLineInfo.test_no_lineinfo_in_device_function.<locals>.callerFr)   )r   r+   r   r   )r   r:   r   r   r9   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 urF| 	d	|  q5| 
 }d
}|D ]}||d ur`d|v r`d} nqO|sk| 	d|  ||}	d}
|	 D ]
}d|v r|
d7 }
qvd}| |
|d| d|
  d S )NTr    c                 S   r4   r5   r   r%   r   r   r   r6   y   r7   zATestCudaLineInfo.test_lineinfo_in_device_function.<locals>.calleec                    r8   r#   r   r%   r9   r   r   r:   }   r;   zATestCudaLineInfo.test_lineinfo_in_device_function.<locals>.callerr)   z^\.weak\s+\.funczFound device function in PTX:

F
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   failr   r   r   assertEqual)r   r:   r   r   ptxlinesdevfn_startlineloc_directivefoundr   subprogramsexpected_subprogramsr   r9   r    test_lineinfo_in_device_functionu   sJ   





z1TestCudaLineInfo.test_lineinfo_in_device_functionc                 C   s   t jdd}t  tjdddddd }W d    n1 s w   Y  | 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+   rA   lencategoryr   assertInstrmessage)r   wrM   r   r   r   test_debug_and_lineinfo_warning   s   	z0TestCudaLineInfo.test_debug_and_lineinfo_warningN)__name__
__module____qualname__r   r   r-   r/   r3   r<   rI   rV   r   r   r   r   r	   
   s    3Ar	   __main__)numbar   r   r   numba.core.errorsr   numba.cuda.testingr   r   numba.tests.supportr   r
   unittestrN   r	   rW   mainr   r   r   r   <module>   s     <