§
    ƒÛgÇ  ã                   óÆ   — 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     ú[/granary/edev/quantum/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Ñ#Ô#ˆØ�nŠn˜SÑ!Ô!ˆØ+1ÐH�4Ô'Ð'°tÔ7Hˆð
#ð 	õ ”
˜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   ó   € àˆAˆa‰DˆDˆDr   ©r   r   ©r   Újitr"   r   ©r   r,   s     r   Útest_no_lineinfo_in_asmz(TestCudaLineInfo.test_no_lineinfo_in_asmJ   sS   € Ý	Œ˜5Ð	!Ñ	!Ô	!ð	ð 	ñ 
"Ô	!ð	ð 	�Š�C�e A A Aœh˜[°ˆÑ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Ð	 Ñ	 Ô	 ð	ð 	ñ 
!Ô	 ð	ð 	�Š�C�e A A Aœh˜[°ˆÑ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   € àˆaˆDˆDŒD�A�a”D‰LˆDˆD‰DˆDˆDr   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˜�sŒ|�W S S q Sœ\Ð*ˆå	Œ�# Ð	%Ñ	%Ô	%ð	ð 	ñ 
&Ô	%ð	ð ×)Ò)¨#Ñ.Ô.ˆð 	×Ò˜ 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   ó   € àˆaˆDˆDŒD�A‰IˆDˆD‰DˆDˆDr   c                 ó(   •— d| d<    ‰| ¦  «         d S r'   r)   ©r+   r?   s    €r   ÚcallerzDTestCudaLineInfo.test_no_lineinfo_in_device_function.<locals>.callerm   ó   ø€ àˆAˆa‰DØˆF�1‰IŒIˆIˆIˆIr   Fr.   )r   r0   r   r"   )r   rC   r   r?   s      @r   Ú#test_no_lineinfo_in_device_functionz4TestCudaLineInfo.test_no_lineinfo_in_device_functiong   st   ø€ å	Œð	ð 	ñ 
Œð	õ 
Œð	ð 	ð 	ð 	ñ 
Œð	õ �Q�Q�QŒxˆkˆØ�Š�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Ð	 Ñ	 Ô	 ð	ð 	ð 	ð 	ñ 
!Ô	 ð	õ �Q�Q�QŒxˆkˆØ�Š�F ¨DˆÑ1Ô1Ð1ð × Ò  Ñ%Ô%ˆØ—>’>Ñ#Ô#ˆõ
 ”jÐ!4Ñ5Ô5ˆàð 	Eð 	EˆDØ× Ò  Ñ&Ô&Ð2Ø—	’	ÐC¸cÐCÐCÑDÔDÐDøð ×1Ò1Ñ3Ô3ˆØˆàð 	ð 	ˆDØ×#Ò# DÑ)Ô)Ð5Ø 4Ð'Ð'Ø �EØ�Eøàð 	'Ø�IŠIð &Ø #ð&ð &ñ 'ô 'ð 'ð ×"Ò" 3Ñ'Ô'ˆØˆØ—O’OÑ%Ô%ð 	!ð 	!ˆ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Ý$Ñ&Ô&Ð&õ ŒX˜D¨4°UÐ;Ñ;Ô;ðð ñ <Ô;ðð	ð 	ð 	ñ 	ô 	ð 	ð 	ð 	ð 	ð 	ð 	øøøð 	ð 	ð 	ð 	ð 	×Ò�˜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ÒÐØ€H„M�O„O€O€O€Oð Ðr   