§
    1@Äj   ã                   óÆ   — 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                   óv   ‡ — e Zd ZdZˆ f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ˆ xZS )ÚTestCudaDebugInfozH
    These tests only checks the compiled PTX for debuginfo section
    c                 ór   •— t          ¦   «                              ¦   «          |                      d¦  «         d S )Nz!Exceptions not supported with LTO)ÚsuperÚsetUpÚskip_if_lto)ÚselfÚ	__class__s    €úWC:\Projects\Historical\venv\Lib\site-packages\numba/cuda/tests/cudapy/test_debuginfo.pyr   zTestCudaDebugInfo.setUp   s1   ø€ Ý‰Œ�Š‰Œˆð 	×ÒÐ<Ñ=Ô=Ð=Ð=Ð=ó    c                 óV   — |                      |¦  «         |                     |¦  «        S ©N)ÚcompileÚinspect_asm)r   ÚfnÚsigs      r   Ú_getasmzTestCudaDebugInfo._getasm   s#   € Ø
�
Š
�3‰ŒˆØ�~Š~˜cÑ"Ô"Ð"r   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   r0   r1   r4   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)r8   c                 ó   — d| d<   d S r)   r+   r,   s    r   r.   z8TestCudaDebugInfo.test_environment_override.<locals>.foo6   ó   € à��!‘��r   Tr0   r&   c                 ó   — d| d<   d S r)   r+   r,   s    r   Úbarz8TestCudaDebugInfo.test_environment_override.<locals>.bar=   r>   r   )r   r   r2   r$   r   r3   )r   r.   r@   s      r   Útest_environment_overridez+TestCudaDebugInfo.test_environment_override3   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*   TFr7   c                 ó   — d| d<   d S )Nr   r+   r,   s    r   Úfz,TestCudaDebugInfo.test_issue_5835.<locals>.fG   r/   r   ©r   r2   r   r3   ©r   rD   s     r   Útest_issue_5835z!TestCudaDebugInfo.test_issue_5835C   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   r7   c                 ó   — d| d<   d S r)   r+   r,   s    r   rD   z7TestCudaDebugInfo.test_wrapper_has_debuginfo.<locals>.fN   r/   r   c                 ó   — g | ]}d |v ¯|‘Œ	S )zdefine void @"_ZN6cudapyr+   )Ú.0Úlines     r   ú
<listcomp>z@TestCudaDebugInfo.test_wrapper_has_debuginfo.<locals>.<listcomp>T   s,   € ð :ð :ð :˜DØ0°DÐ8Ð8ð Ø8Ð8Ð8r   z!dbg)	r   r3   r   r2   Úinspect_llvmÚ
splitlinesÚassertEqualÚlenÚassertIn)r   r   rD   Úllvm_irÚdefinesÚwrapper_defines         r   Útest_wrapper_has_debuginfoz,TestCudaDebugInfo.test_wrapper_has_debuginfoK   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 )NTFr7   c                 ó(   — | d         dv rdnd|d<   d S )Nr   )é   é   r*   rZ   r+   )ÚinpÚoutps     r   rD   zDTestCudaDebugInfo.test_debug_function_calls_internal_impl.<locals>.fk   s!   € à˜qœ6 VÐ+Ð+�a�a°ˆD�‰GˆGˆGr   rE   rF   s     r   Ú'test_debug_function_calls_internal_implz9TestCudaDebugInfo.test_debug_function_calls_internal_impl]   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'   r8   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>.threadidt   s   € å”=”?¥T¤]¤_Ñ4µt´~Ô7GÑGÐGr   r7   c                 ót   •— t          j        d¦  «        }|t          | ¦  «        k     r ‰¦   «         | |<   d S d S ©Nr*   )r   ÚgridrQ   )ÚarrÚire   s     €r   ÚkernelzKTestCudaDebugInfo.test_debug_function_calls_device_function.<locals>.kernelx   s9   ø€ å”	˜!‘”ˆAØ•3�s‘8”8Š|ˆ|Ø!˜™œ��A‘��ð ˆ|r   rE   )r   rk   re   s     @r   Ú)test_debug_function_calls_device_functionz;TestCudaDebugInfo.test_debug_function_calls_device_functiono   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 )NTFr_   c                 ó   — | dz   S rg   r+   r,   s    r   Úf2z;TestCudaDebugInfo._test_chained_device_function.<locals>.f2   ó   € à�q‘5ˆLr   c                 ó    •— |  ‰|¦  «        z
  S r   r+   ©r-   Úyro   s     €r   Úf1z;TestCudaDebugInfo._test_chained_device_function.<locals>.f1ƒ   ó   ø€ à�r�r˜!‘u”u‘9Ðr   r7   c                 ó    •—  ‰| |¦  «         d S r   r+   ©r-   rs   rt   s     €r   rk   z?TestCudaDebugInfo._test_chained_device_function.<locals>.kernel‡   s   ø€ àˆBˆq�!‰HŒHˆHˆHˆHr   ©r*   r*   r*   rY   rE   ©r   Úkernel_debugÚf1_debugÚf2_debugrk   rt   ro   s        @@r   Ú_test_chained_device_functionz/TestCudaDebugInfo._test_chained_device_function~   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)TFrZ   )rz   r{   r|   )Ú	itertoolsÚproductÚsubTestr}   ©r   Ú
debug_optsrz   r{   r|   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 )NTFr_   c                 ó   — | dz   S rg   r+   r,   s    r   ro   zETestCudaDebugInfo._test_chained_device_function_two_calls.<locals>.f2Ÿ   rp   r   c                 ó    •— |  ‰|¦  «        z
  S r   r+   rr   s     €r   rt   zETestCudaDebugInfo._test_chained_device_function_two_calls.<locals>.f1£   ru   r   r7   c                 ó6   •—  ‰| |¦  «          ‰| ¦  «         d S r   r+   )r-   rs   rt   ro   s     €€r   rk   zITestCudaDebugInfo._test_chained_device_function_two_calls.<locals>.kernel§   s#   ø€ àˆBˆq�!‰HŒHˆHØˆBˆq‰EŒEˆEˆEˆEr   rx   r*   rY   ©r   r2   ry   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 r   )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 )NTFr_   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 rg   r+   )r-   r’   s    €r   ro   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+   rr   s     €r   rt   z[TestCudaDebugInfo.test_chained_device_three_functions.<locals>.three_device_fns.<locals>.f1Ë   s   ø€ à˜2˜2˜a™5œ5‘yÐ r   r7   c                 ó    •—  ‰| |¦  «         d S r   r+   rw   s     €r   rk   z_TestCudaDebugInfo.test_chained_device_three_functions.<locals>.three_device_fns.<locals>.kernelÏ   s   ø€ à��1�a‘”���r   rx   r*   rY   r‹   )rz   Ú
leaf_debugrk   rt   ro   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)rz   r–   Fr+   )r   r—   s     r   Ú#test_chained_device_three_functionsz5TestCudaDebugInfo.test_chained_device_three_functions¾   sq   € ð	ð 	ð 	ð( 	Ð d°tÐ<Ñ<Ô<Ð<ØÐ d°uÐ=Ñ=Ô=Ð=ØÐ e¸Ð=Ñ=Ô=Ð=ØÐ e¸Ð>Ñ>Ô>Ð>Ð>Ð>r   )Ú__name__Ú
__module__Ú__qualname__Ú__doc__r   r   r$   r5   r:   rA   rG   rV   r]   rl   r}   r…   rŒ   rŽ   r˜   Ú__classcell__)r   s   @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Ñ:Ô:ðM?ð M?ð M?ð M?ð M?˜ñ M?ô M?ñ ;Ô:ðM?ð` ˆzÒÐØ€H„M�O„O€O€O€Oð Ðr   