
    sh                         S SK JrJrJr  S SKJr  S SKJrJr  S SK	J
r
  S SKrS SKrS SKr\" S5       " S S\5      5       r\S	:X  a  \R                   " 5         gg)
    )cudafloat32int32)NumbaInvalidConfigWarning)CUDATestCaseskip_on_cudasim)ignore_internal_warningsNz#Simulator does not produce lineinfoc                   D    \ rS rSrS rS rS rS rS rS r	S r
S	 rS
rg)TestCudaLineInfo
   c                 2    Sn[         R                  " U5      $ )Nz \.loc\s+[0-9]+\s+[0-9]+\s+[0-9]+)recompile)selfpats     y/Users/tiagomarins/Projetos/claudeai/copy_bank/venv/lib/python3.13/site-packages/numba/cuda/tests/cudapy/test_lineinfo.py_loc_directive_regex%TestCudaLineInfo._loc_directive_regex   s     	 zz#    c                 ~   UR                  U5        UR                  U5      nUR                  U5      nU(       a  U R                  OU R                  nSn[
        R                   " U5      R                  U5      nU" XS9  Sn[
        R                   " U5      R                  U5      nU R	                  XS9  Sn[
        R                   " U5      R                  U5      nU" XS9  U R                  5       R                  U5        U" XS9  Sn[
        R                   " U5      R                  U5      nU R	                  XS9  g )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   _checkTestCudaLineInfo._check   s!   


3s#nnS!+14''t7H7H
# 	 

3&&t,  	
 

3&&t,%)
$ 	
 

3&&s+  	!!#**3/ 
 	 

3&&s+%)r   c                 j    [         R                  " SS9S 5       nU R                  U[        S S  4SS9  g )NFlineinfoc                     SU S'   g N   r    xs    r   foo5TestCudaLineInfo.test_no_lineinfo_in_asm.<locals>.fooK       AaDr   r   r   r   jitr$   r   r   r/   s     r   test_no_lineinfo_in_asm(TestCudaLineInfo.test_no_lineinfo_in_asmJ   s7    	5	!	 
"	 	CeAh[7r   c                 j    [         R                  " SS9S 5       nU R                  U[        S S  4SS9  g )NTr'   c                     SU S'   g r*   r,   r-   s    r   r/   2TestCudaLineInfo.test_lineinfo_in_asm.<locals>.fooR   r1   r   r2   r3   r5   s     r   test_lineinfo_in_asm%TestCudaLineInfo.test_lineinfo_in_asmQ   s7    	4	 	 
!	 	CeAh[6r   c                     [         S S S2   [         S S S2   4n[        R                  " USS9S 5       nUR                  U5      nU R	                  SU5        g )Nr+   Tr'   c                 $    U S==   US   -  ss'   g )Nr   r,   )r.   ys     r   divide_kernelKTestCudaLineInfo.test_lineinfo_maintains_error_model.<locals>.divide_kernel[   s    aDAaDLDr   z	ret i32 1)r   r   r4   r   assertNotIn)r   r   r@   r    s       r   #test_lineinfo_maintains_error_model4TestCudaLineInfo.test_lineinfo_maintains_error_modelX   s]    ss|WSqS\*	#	%	 
&	 ))#. 	d+r   c                    ^ [         R                  S 5       m[         R                  U4S j5       n[        S S  4nU R                  XSS9  g )Nc                     U S==   S-  ss'   g Nr   r+   r,   r-   s    r   calleeDTestCudaLineInfo.test_no_lineinfo_in_device_function.<locals>.calleei       aDAIDr   c                     > SU S'   T" U 5        g r*   r,   r.   rH   s    r   callerDTestCudaLineInfo.test_no_lineinfo_in_device_function.<locals>.callerm       AaD1Ir   Fr2   )r   r4   r   r$   )r   rM   r   rH   s      @r   #test_no_lineinfo_in_device_function4TestCudaLineInfo.test_no_lineinfo_in_device_functiong   sN    		 
	 
	 
	 QxkFE2r   c                   ^ [         R                  " SS9S 5       m[         R                  " SS9U4S j5       n[        S S  4nU R                  XSS9  UR	                  U5      nUR                  5       n[        R                  " S5      nU H+  nUR                  U5      c  M  U R                  SU 35        M-     U R                  5       nSnU H!  nUR                  U5      c  M  S	U;   d  M  Sn  O   U(       d  U R                  S
U 35        UR                  U5      n	Sn
U	R                  5        H  nSU;   d  M  U
S-  n
M     SnU R                  XSU SU
 35        g )NTr'   c                     U S==   S-  ss'   g rG   r,   r-   s    r   rH   ATestCudaLineInfo.test_lineinfo_in_device_function.<locals>.calleey   rJ   r   c                     > SU S'   T" U 5        g r*   r,   rL   s    r   rM   ATestCudaLineInfo.test_lineinfo_in_device_function.<locals>.caller}   rO   r   r2   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   r4   r   r$   r   
splitlinesr   r   r#   failr   r   r   assertEqual)r   rM   r   r!   ptxlinesdevfn_startlineloc_directivefoundr    subprogramsexpected_subprogramsrH   s               @r    test_lineinfo_in_device_function1TestCudaLineInfo.test_lineinfo_in_device_functionu   s    
4	 	 
!	 
4	 	 
!	 QxkFD1   %>>#
 jj!45D  &2		=cUCD  113D##D)54' E	  II   #u& ' ""3'OO%D'4/q  &  !%&:%; <  +}.	/r   c                 x   [         R                  " SS9 n[        5         [        R                  " SSSS9S 5       nS S S 5        U R                  [        W5      S5        U R                  US   R                  [        5        U R                  S[        US   R                  5      5        g ! , (       d  f       Nu= f)	NT)recordF)debugr(   optc                      g )Nr,   r,   r   r   f;TestCudaLineInfo.test_debug_and_lineinfo_warning.<locals>.f   s    r   r+   r   z)debug and lineinfo are mutually exclusive)warningscatch_warningsr	   r   r4   r[   lencategoryr   assertInstrmessage)r   wrj   s      r   test_debug_and_lineinfo_warning0TestCudaLineInfo.test_debug_and_lineinfo_warning   s    $$D1Q$& XXD4U; < 2 	Q#1(ABA!A$,,'	) 21s   'B++
B9r,   N)__name__
__module____qualname____firstlineno__r   r$   r6   r;   rC   rP   rc   rt   __static_attributes__r,   r   r   r   r   
   s,    	1*f87,3?/B)r   r   __main__)numbar   r   r   numba.core.errorsr   numba.cuda.testingr   r   numba.tests.supportr	   r   unittestrl   r   rv   mainr,   r   r   <module>r      sZ    & & 7 < 8 	   67x)| x) 8x)v zMMO r   