
    ܙd                         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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)NVVM)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     Elib/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>.fooL       AaDDDr   r   r   )r   jitr#   r   r   r-   s     r   test_no_lineinfo_in_asmz(TestCudaLineInfo.test_no_lineinfo_in_asmK   sS    	5	!	!	!	 	 
"	!	 	CeAAAh[77777r   c                     t                      j        s|                     d           t          j        d          d             }|                     |t          d d          fd           d S )N#lineinfo not generated for NVVM 3.4Tr%   c                     d| d<   d S r(   r*   r+   s    r   r-   z2TestCudaLineInfo.test_lineinfo_in_asm.<locals>.fooV   r.   r   r/   )r	   	is_nvvm70skipTestr   r0   r#   r   r1   s     r   test_lineinfo_in_asmz%TestCudaLineInfo.test_lineinfo_in_asmR   sv    vv 	AMM?@@@	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   r<   r   s       r   #test_lineinfo_maintains_error_modelz4TestCudaLineInfo.test_lineinfo_maintains_error_model\   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>.calleem       aDDDAIDDDDDr   c                 (    d| d<    |            d S r(   r*   r,   rB   s    r   callerzDTestCudaLineInfo.test_no_lineinfo_in_device_function.<locals>.callerq       AaDF1IIIIIr   Fr/   )r   r0   r   r#   )r   rF   r   rB   s      @r   #test_no_lineinfo_in_device_functionz4TestCudaLineInfo.test_no_lineinfo_in_device_functionk   st    		 	 
	 
	 	 	 	 
	 QQQxkFE22222r   c                 R   t                      j        s|                     d           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 )Nr4   Tr%   c                 &    | dxx         dz  cc<   d S rA   r*   r+   s    r   rB   zATestCudaLineInfo.test_lineinfo_in_device_function.<locals>.callee   rC   r   c                 (    d| d<    |            d S r(   r*   rE   s    r   rF   zATestCudaLineInfo.test_lineinfo_in_device_function.<locals>.caller   rG   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	   r6   r7   r   r0   r   r#   r   
splitlinesr   r   r"   failr   r   r   assertEqual)r   rF   r   r    ptxlinesdevfn_startlineloc_directivefoundr   subprogramsexpected_subprogramsrB   s               @r    test_lineinfo_in_device_functionz1TestCudaLineInfo.test_lineinfo_in_device_functiony   s:   vv 	AMM?@@@
 
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   rP   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   r8   r>   rH   rX   rg   r*   r   r   r   r      s        	 	 	1* 1* 1*f8 8 87 7 7, , ,3 3 3B/ B/ B/H) ) ) ) )r   r   __main__)numbar   r   r   numba.core.errorsr   numba.cuda.testingr   r   numba.cuda.cudadrv.nvvmr	   numba.tests.supportr
   r   unittestr_   r   rh   mainr*   r   r   <module>rs      s    & & & & & & & & & & 7 7 7 7 7 7 < < < < < < < < ( ( ( ( ( ( 8 8 8 8 8 8 				   677~) ~) ~) ~) ~)| ~) ~) 87~)B zHMOOOOO r   