
    J/Ph                         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e                      Zed	k    r ej                     dS dS )
    )cudafloat32int32)NumbaInvalidConfigWarning)CUDATestCaseskip_on_cudasim)ignore_internal_warningsNz#Simulator does not produce lineinfoc                   >    e Zd Zd Zd Zd Zd Zd Zd Zd Z	d Z
d	S )
TestCudaLineInfoc                 .    d}t          j        |          S )Nz \.loc\s+[0-9]+\s+[0-9]+\s+[0-9]+)recompile)selfpats     e/var/www/html/test/jupyter/venv/lib/python3.11/site-packages/numba/cuda/tests/cudapy/test_lineinfo.py_loc_directive_regexz%TestCudaLineInfo._loc_directive_regex   s     	 z#    c                    |                     |           |                    |          }|                    |          }|r| j        n| j        }d}t          j         |                              |          } |||           d}t          j         |                              |          }|                     ||           d}t          j         |                              |          } |||           |                                                     |            |||           d}t          j         |                              |          }|                     ||           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   matchs	            r   _checkzTestCudaLineInfo._check   s}   


3s##nnS!!+1H4''t7H
# 	 
3&&t,,C     	
 
3&&t,,%S)))
$ 	
 
3&&s++C     	!!##**3///C    
 	 
3&&s++%S)))))r   c                     t          j        d          d             }|                     |t          d d          fd           d S )NFlineinfoc                     d| d<   d S N   r    xs    r   fooz5TestCudaLineInfo.test_no_lineinfo_in_asm.<locals>.fooK       AaDDDr   r   r   r   jitr"   r   r   r,   s     r   test_no_lineinfo_in_asmz(TestCudaLineInfo.test_no_lineinfo_in_asmJ   sS    	5	!	!	!	 	 
"	!	 	CeAAAh[77777r   c                     t          j        d          d             }|                     |t          d d          fd           d S )NTr$   c                     d| d<   d S r'   r)   r*   s    r   r,   z2TestCudaLineInfo.test_lineinfo_in_asm.<locals>.fooR   r-   r   r.   r/   r1   s     r   test_lineinfo_in_asmz%TestCudaLineInfo.test_lineinfo_in_asmQ   sS    	4	 	 	 	 	 
!	 	 	CeAAAh[66666r   c                     t           d d d         t           d d d         f}t          j        |d          d             }|                    |          }|                     d|           d S )Nr(   Tr$   c                 2    | dxx         |d         z  cc<   d S )Nr   r)   )r+   ys     r   divide_kernelzKTestCudaLineInfo.test_lineinfo_maintains_error_model.<locals>.divide_kernel[   s    aDDDAaDLDDDDDr   z	ret i32 1)r   r   r0   r   assertNotIn)r   r   r9   r   s       r   #test_lineinfo_maintains_error_modelz4TestCudaLineInfo.test_lineinfo_maintains_error_modelX   s|    sss|WSSqS\*	#	%	%	%	 	 
&	%	 ))#.. 	d+++++r   c                     t           j        d             t           j        fd            }t          d d          f}|                     ||d           d S )Nc                 &    | dxx         dz  cc<   d S Nr   r(   r)   r*   s    r   calleezDTestCudaLineInfo.test_no_lineinfo_in_device_function.<locals>.calleei       aDDDAIDDDDDr   c                 (    d| d<    |            d S r'   r)   r+   r?   s    r   callerzDTestCudaLineInfo.test_no_lineinfo_in_device_function.<locals>.callerm       AaDF1IIIIIr   Fr.   )r   r0   r   r"   )r   rC   r   r?   s      @r   #test_no_lineinfo_in_device_functionz4TestCudaLineInfo.test_no_lineinfo_in_device_functiong   st    		 	 
	 
	 	 	 	 
	 QQQxkFE22222r   c                    t          j        d          d             t          j        d          fd            }t          d d          f}|                     ||d           |                    |          }|                                }t          j        d          }|D ]/}|                    |          | 	                    d|            0| 
                                }d}|D ]}|                    |          d	|v rd} n |s| 	                    d
|            |                    |          }	d}
|	                                D ]}d|v r|
dz  }
d}|                     |
|d| d|
            d S )NTr$   c                 &    | dxx         dz  cc<   d S r>   r)   r*   s    r   r?   zATestCudaLineInfo.test_lineinfo_in_device_function.<locals>.calleey   r@   r   c                 (    d| d<    |            d S r'   r)   rB   s    r   rC   zATestCudaLineInfo.test_lineinfo_in_device_function.<locals>.caller}   rD   r   r.   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   r0   r   r"   r   
splitlinesr   r   r!   failr   r   r   assertEqual)r   rC   r   r   ptxlinesdevfn_startlineloc_directivefoundr   subprogramsexpected_subprogramsr?   s               @r    test_lineinfo_in_device_functionz1TestCudaLineInfo.test_lineinfo_in_device_functionu   s    
4	 	 	 	 	 
!	 	 
4	 	 	 	 	 	 	 
!	 	 QQQxkFD111   %%>>##
 j!455 	E 	ED  &&2		CcCCDDD 1133 	 	D##D))54'' EE 	'II & #& & ' ' ' ""3''OO%% 	! 	!D'4//q   !&:.&: . . +. .	/ 	/ 	/ 	/ 	/r   c                    t          j        d          5 }t                       t          j        ddd          d             }d d d            n# 1 swxY w Y   |                     t          |          d           |                     |d         j        t                     | 	                    dt          |d         j                             d S )	NT)recordF)debugr%   optc                      d S )Nr)   r)   r   r   fz;TestCudaLineInfo.test_debug_and_lineinfo_warning.<locals>.f   s    r   r(   r   z)debug and lineinfo are mutually exclusive)warningscatch_warningsr	   r   r0   rM   lencategoryr   assertInstrmessage)r   wr[   s      r   test_debug_and_lineinfo_warningz0TestCudaLineInfo.test_debug_and_lineinfo_warning   s   $D111 	Q$&&& XD4U;;;  <;	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	Q###1(ABBBA!A$,''	) 	) 	) 	) 	)s   /AAAN)__name__
__module____qualname__r   r"   r2   r5   r;   rE   rU   rd   r)   r   r   r   r   
   s        	 	 	1* 1* 1*f8 8 87 7 7, , ,3 3 3?/ ?/ ?/B) ) ) ) )r   r   __main__)numbar   r   r   numba.core.errorsr   numba.cuda.testingr   r   numba.tests.supportr	   r   unittestr\   r   re   mainr)   r   r   <module>ro      s    & & & & & & & & & & 7 7 7 7 7 7 < < < < < < < < 8 8 8 8 8 8 				   677x) x) x) x) x)| x) x) 87x)v zHMOOOOO r   