§
    ƒÛg%  ã                   óÆ   — d dl mZ d dlmZ d dl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 )é    )Úoverride_config)Úskip_on_cudasim)Úcuda)Útypes)ÚCUDATestCaseNz&Simulator does not produce debug dumpsc                   óf   — e Zd ZdZd„ Zd„ Zd„ Zd„ Zd„ Zd„ Z	d„ Z
d	„ Zd
„ Zd„ Zd„ Zd„ Zd„ Zd„ ZdS )ÚTestCudaDebugInfozH
    These tests only checks the compiled PTX for debuginfo section
    c                 óV   — |                      |¦  «         |                     |¦  «        S ©N)ÚcompileÚinspect_asm)ÚselfÚfnÚsigs      ú\/granary/edev/quantum/lib/python3.11/site-packages/numba/cuda/tests/cudapy/test_debuginfo.pyÚ_getasmzTestCudaDebugInfo._getasm   s#   € Ø
�
Š
�3‰ŒˆØ�~Š~˜cÑ"Ô"Ð"ó    c                 óÀ   — |                       ||¬¦  «        }t          j        d¦  «        }|                     |¦  «        }|r| j        n| j        } |||¬¦  «         d S )N)r   z\.section\s+\.debug_info\s+{)Úmsg)r   Úrer   ÚsearchÚassertIsNotNoneÚassertIsNone)r   r   r   ÚexpectÚasmÚre_section_dbginfoÚmatchÚassertfns           r   Ú_checkzTestCudaDebugInfo._check   sk   € Ø�lŠl˜2 3ˆlÑ'Ô'ˆÝœZÐ(GÑHÔHÐØ"×)Ò)¨#Ñ.Ô.ˆØ+1ÐH�4Ô'Ð'°tÔ7HˆØˆ�˜CÐ Ñ Ô Ð Ð Ð r   c                 ó˜   — t          j        d¬¦  «        d„ ¦   «         }|                      |t          j        d d …         fd¬¦  «         d S )NF©Údebugc                 ó   — d| d<   d S ©Né   r   © ©Úxs    r   Úfooz7TestCudaDebugInfo.test_no_debuginfo_in_asm.<locals>.foo   ó   € àˆAˆa‰DˆDˆDr   ©r   r   ©r   Újitr   r   Úint32©r   r)   s     r   Útest_no_debuginfo_in_asmz*TestCudaDebugInfo.test_no_debuginfo_in_asm   sV   € Ý	Œ˜Ð	Ñ	Ô	ð	ð 	ñ 
Ô	ð	ð 	�Š�C�eœk¨!¨!¨!œnÐ.°uˆÑ=Ô=Ð=Ð=Ð=r   c                 óš   — t          j        dd¬¦  «        d„ ¦   «         }|                      |t          j        d d …         fd¬¦  «         d S )NTF©r"   Úoptc                 ó   — d| d<   d S r$   r&   r'   s    r   r)   z4TestCudaDebugInfo.test_debuginfo_in_asm.<locals>.foo#   r*   r   r+   r,   r/   s     r   Útest_debuginfo_in_asmz'TestCudaDebugInfo.test_debuginfo_in_asm"   sX   € Ý	Œ˜ %Ð	(Ñ	(Ô	(ð	ð 	ñ 
)Ô	(ð	ð 	�Š�C�eœk¨!¨!¨!œnÐ.°tˆÑ<Ô<Ð<Ð<Ð<r   c                 ó|  — t          dd¦  «        5  t          j        d¬¦  «        d„ ¦   «         }|                      |t          j        d d …         fd¬¦  «         t          j        d¬¦  «        d	„ ¦   «         }|                      |t          j        d d …         fd¬¦  «         d d d ¦  «         d S # 1 swxY w Y   d S )
NÚCUDA_DEBUGINFO_DEFAULTr%   F)r3   c                 ó   — d| d<   d S r$   r&   r'   s    r   r)   z8TestCudaDebugInfo.test_environment_override.<locals>.foo,   ó   € à��!‘��r   Tr+   r!   c                 ó   — d| d<   d S r$   r&   r'   s    r   Úbarz8TestCudaDebugInfo.test_environment_override.<locals>.bar3   r9   r   )r   r   r-   r   r   r.   )r   r)   r;   s      r   Útest_environment_overridez+TestCudaDebugInfo.test_environment_override)   s)  € ÝÐ5°qÑ9Ô9ð 	Bð 	BåŒX˜%Ð Ñ Ô ðð ñ !Ô ðð �KŠK˜¥%¤+¨a¨a¨a¤.Ð!2¸4ˆKÑ@Ô@Ð@õ ŒX˜EÐ"Ñ"Ô"ðð ñ #Ô"ðð �KŠK˜¥%¤+¨a¨a¨a¤.Ð!2¸5ˆKÑAÔAÐAð	Bð 	Bð 	Bñ 	Bô 	Bð 	Bð 	Bð 	Bð 	Bð 	Bð 	Bð 	Bøøøð 	Bð 	Bð 	Bð 	Bð 	Bð 	Bs   ‘BB1Â1B5Â8B5c                 ón   — t          j        t          j        d d d…         fdd¬¦  «        d„ ¦   «         }d S )Nr%   TFr2   c                 ó   — d| d<   d S )Nr   r&   r'   s    r   Úfz,TestCudaDebugInfo.test_issue_5835.<locals>.f=   r*   r   ©r   r-   r   r.   ©r   r?   s     r   Útest_issue_5835z!TestCudaDebugInfo.test_issue_58359   sI   € õ 
Œ•5”;˜s˜s ˜sÔ#Ð%¨T°uÐ	=Ñ	=Ô	=ð	ð 	ñ 
>Ô	=ð	ð 	ð 	r   c                 óZ  — t           j        d d d…         f}t          j        |dd¬¦  «        d„ ¦   «         }|                     |¦  «        }d„ |                     ¦   «         D ¦   «         }|                      t          |¦  «        d¦  «         |d         }|                      d|¦  «         d S )Nr%   Tr   r2   c                 ó   — d| d<   d S r$   r&   r'   s    r   r?   z7TestCudaDebugInfo.test_wrapper_has_debuginfo.<locals>.fD   r*   r   c                 ó   — g | ]}d |v ¯|‘Œ	S )zdefine void @"_ZN6cudapyr&   )Ú.0Úlines     r   ú
<listcomp>z@TestCudaDebugInfo.test_wrapper_has_debuginfo.<locals>.<listcomp>J   s,   € ð :ð :ð :˜DØ0°DÐ8Ð8ð Ø8Ð8Ð8r   z!dbg)	r   r.   r   r-   Úinspect_llvmÚ
splitlinesÚassertEqualÚlenÚassertIn)r   r   r?   Úllvm_irÚdefinesÚwrapper_defines         r   Útest_wrapper_has_debuginfoz,TestCudaDebugInfo.test_wrapper_has_debuginfoA   s¿   € ÝŒ{˜3˜3˜Q˜3ÔÐ!ˆå	Œ�#˜T qÐ	)Ñ	)Ô	)ð	ð 	ñ 
*Ô	)ð	ð —.’. Ñ%Ô%ˆð:ð : G×$6Ò$6Ñ$8Ô$8ð :ñ :ô :ˆð 	×Ò�˜W™œ qÑ)Ô)Ð)à  œˆØ�Š�f˜nÑ-Ô-Ð-Ð-Ð-r   c                 ó’   — t          j        t          j        d d …         t          j        d d …         fdd¬¦  «        d„ ¦   «         }d S )NTFr2   c                 ó(   — | d         dv rdnd|d<   d S )Nr   )é   é   r%   rU   r&   )ÚinpÚoutps     r   r?   zDTestCudaDebugInfo.test_debug_function_calls_internal_impl.<locals>.fa   s!   € à˜qœ6 VÐ+Ð+�a�a°ˆD�‰GˆGˆGr   r@   rA   s     r   Ú'test_debug_function_calls_internal_implz9TestCudaDebugInfo.test_debug_function_calls_internal_implS   sS   € õ 
Œ•5”;˜q˜q˜q”>¥5¤;¨q¨q¨q¤>Ð2¸$ÀEÐ	JÑ	JÔ	Jð	3ð 	3ñ 
KÔ	Jð	3ð 	3ð 	3r   c                 ó²   ‡— t          j        ddd¬¦  «        d„ ¦   «         Št          j        t          j        d d …         fdd¬¦  «        ˆfd„¦   «         }d S )NTr   ©Údevicer"   r3   c                  ól   — t           j        j        t           j        j        z  t           j        j        z   S r   )r   ÚblockDimr(   ÚblockIdxÚ	threadIdxr&   r   r   ÚthreadidzMTestCudaDebugInfo.test_debug_function_calls_device_function.<locals>.threadidj   s   € å”=”?¥T¤]¤_Ñ4µt´~Ô7GÑGÐGr   r2   c                 ót   •— t          j        d¦  «        }|t          | ¦  «        k     r ‰¦   «         | |<   d S d S ©Nr%   )r   ÚgridrL   )ÚarrÚir`   s     €r   ÚkernelzKTestCudaDebugInfo.test_debug_function_calls_device_function.<locals>.kerneln   s9   ø€ å”	˜!‘”ˆAØ•3�s‘8”8Š|ˆ|Ø!˜™œ��A‘��ð ˆ|r   r@   )r   rf   r`   s     @r   Ú)test_debug_function_calls_device_functionz;TestCudaDebugInfo.test_debug_function_calls_device_functione   s�   ø€ õ
 
Œ˜ T¨qÐ	1Ñ	1Ô	1ð	Hð 	Hñ 
2Ô	1ð	Hõ 
Œ•5”;˜q˜q˜q”>Ð#¨4°QÐ	7Ñ	7Ô	7ð	$ð 	$ð 	$ð 	$ñ 
8Ô	7ð	$ð 	$ð 	$r   c                 ó"  ‡‡— t          j        d|d¬¦  «        d„ ¦   «         Št          j        d|d¬¦  «        ˆfd„¦   «         Št          j        t          j        t          j        f|d¬¦  «        ˆfd„¦   «         } |d         d	d
¦  «         d S )NTFrZ   c                 ó   — | dz   S rb   r&   r'   s    r   Úf2z;TestCudaDebugInfo._test_chained_device_function.<locals>.f2u   ó   € à�q‘5ˆLr   c                 ó    •— |  ‰|¦  «        z
  S r   r&   ©r(   Úyrj   s     €r   Úf1z;TestCudaDebugInfo._test_chained_device_function.<locals>.f1y   ó   ø€ à�r�r˜!‘u”u‘9Ðr   r2   c                 ó    •—  ‰| |¦  «         d S r   r&   ©r(   rn   ro   s     €r   rf   z?TestCudaDebugInfo._test_chained_device_function.<locals>.kernel}   s   ø€ àˆBˆq�!‰HŒHˆHˆHˆHr   ©r%   r%   r%   rT   r@   ©r   Úkernel_debugÚf1_debugÚf2_debugrf   ro   rj   s        @@r   Ú_test_chained_device_functionz/TestCudaDebugInfo._test_chained_device_functiont   sÅ   øø€ Ý	Œ˜ X°5Ð	9Ñ	9Ô	9ð	ð 	ñ 
:Ô	9ð	õ 
Œ˜ X°5Ð	9Ñ	9Ô	9ð	ð 	ð 	ð 	ñ 
:Ô	9ð	õ 
Œ•5”;¥¤Ð,°LÀeÐ	LÑ	LÔ	Lð	ð 	ð 	ð 	ñ 
MÔ	Lð	ð 	ˆˆtŒ�Q˜ÑÔÐÐÐr   c                 óÊ   — t          j        dgdz  Ž }|D ]M\  }}}|                      |||¬¦  «        5  |                      |||¦  «         d d d ¦  «         n# 1 swxY w Y   ŒNd S ©N)TFrU   )ru   rv   rw   )Ú	itertoolsÚproductÚsubTestrx   ©r   Ú
debug_optsru   rv   rw   s        r   Útest_chained_device_functionz.TestCudaDebugInfo.test_chained_device_functionƒ   sÑ   € õ
 Ô&¨¨¸!Ñ(;Ð<ˆ
à0:ð 	=ð 	=Ñ,ˆL˜( HØ—’¨<Ø'/Ø'/ð ñ 1ô 1ð =ð =ð ×2Ò2°<Ø3;Ø3;ñ=ô =ð =ð=ð =ð =ñ =ô =ð =ð =ð =ð =ð =ð =øøøð =ð =ð =ð =øð	=ð 	=ó   ³AÁA	ÁA	c                 óö   ‡‡— t          j        d|d¬¦  «        d„ ¦   «         Št          j        d|d¬¦  «        ˆfd„¦   «         Št          j        |d¬¦  «        ˆˆfd„¦   «         } |d         d	d
¦  «         d S )NTFrZ   c                 ó   — | dz   S rb   r&   r'   s    r   rj   zETestCudaDebugInfo._test_chained_device_function_two_calls.<locals>.f2•   rk   r   c                 ó    •— |  ‰|¦  «        z
  S r   r&   rm   s     €r   ro   zETestCudaDebugInfo._test_chained_device_function_two_calls.<locals>.f1™   rp   r   r2   c                 ó6   •—  ‰| |¦  «          ‰| ¦  «         d S r   r&   )r(   rn   ro   rj   s     €€r   rf   zITestCudaDebugInfo._test_chained_device_function_two_calls.<locals>.kernel�   s#   ø€ àˆBˆq�!‰HŒHˆHØˆBˆq‰EŒEˆEˆEˆEr   rs   r%   rT   ©r   r-   rt   s        @@r   Ú'_test_chained_device_function_two_callsz9TestCudaDebugInfo._test_chained_device_function_two_calls’   sÀ   øø€ õ 
Œ˜ X°5Ð	9Ñ	9Ô	9ð	ð 	ñ 
:Ô	9ð	õ 
Œ˜ X°5Ð	9Ñ	9Ô	9ð	ð 	ð 	ð 	ñ 
:Ô	9ð	õ 
Œ˜¨%Ð	0Ñ	0Ô	0ð	ð 	ð 	ð 	ð 	ñ 
1Ô	0ð	ð 	ˆˆtŒ�Q˜ÑÔÐÐÐr   c                 óÊ   — t          j        dgdz  Ž }|D ]M\  }}}|                      |||¬¦  «        5  |                      |||¦  «         d d d ¦  «         n# 1 swxY w Y   ŒNd S rz   )r{   r|   r}   r‡   r~   s        r   Ú&test_chained_device_function_two_callsz8TestCudaDebugInfo.test_chained_device_function_two_calls¤   sé   € õ Ô&¨¨¸!Ñ(;Ð<ˆ
à0:ð 	Gð 	GÑ,ˆL˜( HØ—’¨<Ø'/Ø'/ð ñ 1ô 1ð Gð Gð ×<Ò<¸\Ø=EØ=EñGô Gð GðGð Gð Gñ Gô Gð Gð Gð Gð Gð Gð Gøøøð Gð Gð Gð Gøð	Gð 	Gr�   c                 ót   — d„ } |dd¬¦  «          |dd¬¦  «          |dd¬¦  «          |dd¬¦  «         d S )Nc                 ó2  ‡‡‡— t          j        d|d¬¦  «        d„ ¦   «         Št          j        d¬¦  «        ˆfd„¦   «         Št          j        d¬¦  «        ˆfd„¦   «         Št          j        | d¬¦  «        ˆfd	„¦   «         } |d
         dd¦  «         d S )NTFrZ   c                 ó   — | | z  S r   r&   r'   s    r   Úf3z[TestCudaDebugInfo.test_chained_device_three_functions.<locals>.three_device_fns.<locals>.f3¹   s   € à˜1‘u�r   )r[   c                 ó    •—  ‰| ¦  «        dz   S rb   r&   )r(   r�   s    €r   rj   z[TestCudaDebugInfo.test_chained_device_three_functions.<locals>.three_device_fns.<locals>.f2½   s   ø€ à�r˜!‘u”u˜q‘yÐ r   c                 ó    •— |  ‰|¦  «        z
  S r   r&   rm   s     €r   ro   z[TestCudaDebugInfo.test_chained_device_three_functions.<locals>.three_device_fns.<locals>.f1Á   s   ø€ à˜2˜2˜a™5œ5‘yÐ r   r2   c                 ó    •—  ‰| |¦  «         d S r   r&   rr   s     €r   rf   z_TestCudaDebugInfo.test_chained_device_three_functions.<locals>.three_device_fns.<locals>.kernelÅ   s   ø€ à��1�a‘”���r   rs   r%   rT   r†   )ru   Ú
leaf_debugrf   ro   rj   r�   s      @@@r   Úthree_device_fnszOTestCudaDebugInfo.test_chained_device_three_functions.<locals>.three_device_fns¸   sê   øøø€ ÝŒX˜T¨¸Ð?Ñ?Ô?ðð ñ @Ô?ðõ ŒX˜TÐ"Ñ"Ô"ð!ð !ð !ð !ñ #Ô"ð!õ ŒX˜TÐ"Ñ"Ô"ð!ð !ð !ð !ñ #Ô"ð!õ ŒX˜L¨eÐ4Ñ4Ô4ðð ð ð ñ 5Ô4ðð ˆF�4ŒL˜˜AÑÔÐÐÐr   T)ru   r‘   Fr&   )r   r’   s     r   Ú#test_chained_device_three_functionsz5TestCudaDebugInfo.test_chained_device_three_functions´   sq   € ð	ð 	ð 	ð( 	Ð d°tÐ<Ñ<Ô<Ð<ØÐ d°uÐ=Ñ=Ô=Ð=ØÐ e¸Ð=Ñ=Ô=Ð=ØÐ e¸Ð>Ñ>Ô>Ð>Ð>Ð>r   N)Ú__name__Ú
__module__Ú__qualname__Ú__doc__r   r   r0   r5   r<   rB   rQ   rX   rg   rx   r€   r‡   r‰   r“   r&   r   r   r	   r	      sö   € € € € € ðð ð#ð #ð #ð!ð !ð !ð>ð >ð >ð=ð =ð =ðBð Bð Bð ð ð ð.ð .ð .ð$3ð 3ð 3ð$$ð $ð $ðð ð ð=ð =ð =ðð ð ð$Gð Gð Gð ?ð ?ð ?ð ?ð ?r   r	   Ú__main__)Únumba.tests.supportr   Únumba.cuda.testingr   Únumbar   Ú
numba.corer   r   r{   r   Úunittestr	   r”   Úmainr&   r   r   ú<module>rŸ      sî   ðØ /Ð /Ð /Ð /Ð /Ð /Ø .Ð .Ð .Ð .Ð .Ð .Ø Ð Ð Ð Ð Ð Ø Ð Ð Ð Ð Ð Ø +Ð +Ð +Ð +Ð +Ð +Ø Ð Ð Ð Ø 	€	€	€	Ø €€€ð €Ð9Ñ:Ô:ðC?ð C?ð C?ð C?ð C?˜ñ C?ô C?ñ ;Ô:ðC?ðL ˆzÒÐØ€H„M�O„O€O€O€Oð Ðr   