
    xKg                         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                           yy)
    )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
y	)
TestCudaLineInfoc                 0    d}t        j                  |      S )Nz \.loc\s+[0-9]+\s+[0-9]+\s+[0-9]+)recompile)selfpats     i/home/alanp/www/video.onchill/myenv/lib/python3.12/site-packages/numba/cuda/tests/cudapy/test_lineinfo.py_loc_directive_regexz%TestCudaLineInfo._loc_directive_regex   s     	 zz#    c                    |j                  |       |j                  |      }|j                  |      }|r| j                  n| j                  }d}t        j                   |      j                  |      } |||       d}t        j                   |      j                  |      }| j	                  ||       d}t        j                   |      j                  |      } |||       | j                         j                  |        |||       d}t        j                   |      j                  |      }| j	                  ||       y )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!+14''t7H7H
# 	 

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

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

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

3&&s+%S)r   c                 t    t        j                  d      d        }| j                  |t        d d  fd       y )NFlineinfoc                     d| d<   y N   r    xs    r   fooz5TestCudaLineInfo.test_no_lineinfo_in_asm.<locals>.fooK       AaDr   r   r   r   jitr"   r   r   r,   s     r   test_no_lineinfo_in_asmz(TestCudaLineInfo.test_no_lineinfo_in_asmJ   s7    	5	!	 
"	 	CeAh[7r   c                 t    t        j                  d      d        }| j                  |t        d d  fd       y )NTr$   c                     d| d<   y 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   s7    	4	 	 
!	 	CeAh[6r   c                     t         d d d   t         d d d   f}t        j                  |d      d        }|j                  |      }| j	                  d|       y )Nr(   Tr$   c                 $    | dxx   |d   z  cc<   y )Nr   r)   )r+   ys     r   divide_kernelzKTestCudaLineInfo.test_lineinfo_maintains_error_model.<locals>.divide_kernel[   s    aDAaDLDr   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]    ss|WSqS\*	#	%	 
&	 ))#. 	d+r   c                     t         j                  d        t         j                  fd       }t        d d  f}| j                  ||d       y )Nc                     | dxx   dz  cc<   y Nr   r(   r)   r*   s    r   calleezDTestCudaLineInfo.test_no_lineinfo_in_device_function.<locals>.calleei       aDAIDr   c                      d| d<    |        y r'   r)   r+   r?   s    r   callerzDTestCudaLineInfo.test_no_lineinfo_in_device_function.<locals>.callerm       AaD1Ir   Fr.   )r   r0   r   r"   )r   rC   r   r?   s      @r   #test_no_lineinfo_in_device_functionz4TestCudaLineInfo.test_no_lineinfo_in_device_functiong   sP    		 
	 
	 
	 QxkFE2r   c                    t        j                  d      d        t        j                  d      fd       }t        d d  f}| j                  ||d       |j	                  |      }|j                         }t        j                  d      }|D ](  }|j                  |      | j                  d|        * | j                         }d}|D ]  }|j                  |      d	|v sd} n |s| j                  d
|        |j                  |      }	d}
|	j                         D ]  }d|v s|
dz  }
 d}| j                  |
|d| d|
        y )NTr$   c                     | dxx   dz  cc<   y r>   r)   r*   s    r   r?   zATestCudaLineInfo.test_lineinfo_in_device_function.<locals>.calleey   r@   r   c                      d| d<    |        y 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	 	 
!	 QxkFD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                 v   t        j                  d      5 }t                t        j                  ddd      d        }d d d        | j                  t              d       | j                  |d   j                  t               | j                  dt        |d   j                               y # 1 sw Y   pxY w)	NT)recordF)debugr%   optc                       y )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1Q$& XXD4U; < 2 	Q#1(ABA!A$,,'	) 21s   )B//B8N)__name__
__module____qualname__r   r"   r2   r5   r;   rE   rU   rd   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   unittestr\   r   re   mainr)   r   r   <module>ro      s\    & & 7 < 8 	   67x)| x) 8x)v zHMMO r   