§
    ƒÛg  ã                   óR  — d dl mZ d dlmZmZmZmZmZmZ d dl	m
Z
mZ d dlmZmZmZ  ed¦  «         G d„ dej        ¦  «        ¦   «         Z ed¦  «         G d„ d	e¦  «        ¦   «         Z ed¦  «         G d
„ dej        ¦  «        ¦   «         Zedk    r ej        ¦   «          dS dS )é    ©Úsqrt)ÚcudaÚfloat32Úint16Úint32Úuint32Úvoid)Úcompile_ptxÚcompile_ptx_for_current_device)Úskip_on_cudasimÚunittestÚCUDATestCasez(Compilation unsupported in the simulatorc                   óJ   — 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S )ÚTestCompileToPTXc                 óD  — d„ }t           d d …         t           d d …         t           d d …         f}t          ||¦  «        \  }}|                      d|¦  «         |                      d|¦  «         |                      d|¦  «         |                      |t
          ¦  «         d S )Nc                 ó€   — t          j        d¦  «        }|t          | ¦  «        k     r||         ||         z   | |<   d S d S )Né   )r   ÚgridÚlen)ÚrÚxÚyÚis       ú[/granary/edev/quantum/lib/python3.11/site-packages/numba/cuda/tests/cudapy/test_compiler.pyÚfz.TestCompileToPTX.test_global_kernel.<locals>.f   s>   € Ý”	˜!‘”ˆAØ•3�q‘6”6ŠzˆzØ˜”t˜a œd‘{��!‘��ð ˆzó    Úfunc_retvalú.visible .funcú.visible .entry)r   r   ÚassertNotInÚassertInÚassertEqualr
   ©Úselfr   ÚargsÚptxÚrestys        r   Útest_global_kernelz#TestCompileToPTX.test_global_kernel
   sª   € ð	#ð 	#ð 	#õ
 ˜˜˜”
�G A A AœJ­°°°¬
Ð3ˆÝ   DÑ)Ô)‰
ˆˆUð 	×Ò˜¨Ñ,Ô,Ð,à×ÒÐ)¨3Ñ/Ô/Ð/à�ŠÐ'¨Ñ-Ô-Ð-à×Ò˜¥Ñ%Ô%Ð%Ð%Ð%r   c                 ó˜  — d„ }t           t           f}t          ||d¬¦  «        \  }}|                      d|¦  «         |                      d|¦  «         |                      d|¦  «         |                      |t           ¦  «         t          t
          t
          ¦  «        }t          ||d¬¦  «        \  }}|                      |t
          ¦  «         t          t          t          ¦  «        }t          ||d¬¦  «        \  }}|                      |t          ¦  «         d}t          ||d¬¦  «        \  }}|                      |t          ¦  «         d S )Nc                 ó   — | |z   S ©N© ©r   r   s     r   Úaddz2TestCompileToPTX.test_device_function.<locals>.add   ó   € Ø�q‘5ˆLr   T©Údevicer   r   r    zuint32(uint32, uint32))r   r   r"   r!   r#   r   r   r	   )r%   r/   r&   r'   r(   Ú	sig_int32Ú	sig_int16Ú
sig_strings           r   Útest_device_functionz%TestCompileToPTX.test_device_function   sB  € ð	ð 	ð 	õ �Ð!ˆÝ   d°4Ð8Ñ8Ô8‰
ˆˆUð 	�Š�m SÑ)Ô)Ð)à�ŠÐ&¨Ñ,Ô,Ð,à×ÒÐ*¨CÑ0Ô0Ð0à×Ò˜¥Ñ(Ô(Ð(õ �%¥Ñ'Ô'ˆ	Ý   i¸Ð=Ñ=Ô=‰
ˆˆUØ×Ò˜¥Ñ&Ô&Ð&å�%¥Ñ'Ô'ˆ	Ý   i¸Ð=Ñ=Ô=‰
ˆˆUØ×Ò˜¥Ñ&Ô&Ð&à-ˆ
Ý   j¸Ð>Ñ>Ô>‰
ˆˆUØ×Ò˜¥Ñ'Ô'Ð'Ð'Ð'r   c                 óž  — d„ }t           t           t           t           f}t          ||d¬¦  «        \  }}|                      d|¦  «         |                      d|¦  «         |                      d|¦  «         t          ||dd¬¦  «        \  }}|                      d|¦  «         |                      d	|¦  «         |                      d
|¦  «         d S )Nc                 ó2   — t          | |z  |z   |z  ¦  «        S r,   r   )r   r   ÚzÚds       r   r   z)TestCompileToPTX.test_fastmath.<locals>.f;   s   € Ý˜˜Q™ ™ a™Ñ(Ô(Ð(r   Tr1   z
fma.rn.f32z
div.rn.f32zsqrt.rn.f32)r2   Úfastmathzfma.rn.ftz.f32zdiv.approx.ftz.f32zsqrt.approx.ftz.f32)r   r   r"   r$   s        r   Útest_fastmathzTestCompileToPTX.test_fastmath:   sØ   € ð	)ð 	)ð 	)õ �¥'­7Ð3ˆÝ   D°Ð6Ñ6Ô6‰
ˆˆUð 	�Š�l CÑ(Ô(Ð(Ø�Š�l CÑ(Ô(Ð(Ø�Š�m SÑ)Ô)Ð)å   D°ÀÐEÑEÔE‰
ˆˆUð 	�ŠÐ&¨Ñ,Ô,Ð,Ø�ŠÐ*¨CÑ0Ô0Ð0Ø�ŠÐ+¨SÑ1Ô1Ð1Ð1Ð1r   c                 ó^   — |                       |d¦  «         |                       |d¦  «         d S )Nz\.section\s+\.debug_infoú\.file.*test_compiler.py"©ÚassertRegex©r%   r'   s     r   Úcheck_debug_infoz!TestCompileToPTX.check_debug_infoN   s:   € ð 	×Ò˜Ð;Ñ<Ô<Ð<ð 	×Ò˜Ð:Ñ;Ô;Ð;Ð;Ð;r   c                 ób   — d„ }t          |ddd¬¦  «        \  }}|                      |¦  «         d S )Nc                  ó   — d S r,   r-   r-   r   r   r   z;TestCompileToPTX.test_device_function_with_debug.<locals>.f^   ó   € ØˆDr   r-   T)r2   Údebug©r   rB   ©r%   r   r'   r(   s       r   Útest_device_function_with_debugz0TestCompileToPTX.test_device_function_with_debugW   sG   € ð	ð 	ð 	õ !  B¨t¸4Ð@Ñ@Ô@‰
ˆˆUØ×Ò˜cÑ"Ô"Ð"Ð"Ð"r   c                 ó`   — d„ }t          |dd¬¦  «        \  }}|                      |¦  «         d S )Nc                  ó   — d S r,   r-   r-   r   r   r   z2TestCompileToPTX.test_kernel_with_debug.<locals>.ff   rE   r   r-   T)rF   rG   rH   s       r   Útest_kernel_with_debugz'TestCompileToPTX.test_kernel_with_debugd   sE   € ð	ð 	ð 	õ !  B¨dÐ3Ñ3Ô3‰
ˆˆUØ×Ò˜cÑ"Ô"Ð"Ð"Ð"r   c                 ó2   — |                       |d¦  «         d S )Nr>   r?   rA   s     r   Úcheck_line_infoz TestCompileToPTX.check_line_infol   s!   € ð 	×Ò˜Ð:Ñ;Ô;Ð;Ð;Ð;r   c                 ób   — d„ }t          |ddd¬¦  «        \  }}|                      |¦  «         d S )Nc                  ó   — d S r,   r-   r-   r   r   r   z?TestCompileToPTX.test_device_function_with_line_info.<locals>.fs   rE   r   r-   T)r2   Úlineinfo©r   rN   rH   s       r   Ú#test_device_function_with_line_infoz4TestCompileToPTX.test_device_function_with_line_infor   sG   € ð	ð 	ð 	õ !  B¨t¸dÐCÑCÔC‰
ˆˆUØ×Ò˜SÑ!Ô!Ð!Ð!Ð!r   c                 ó`   — d„ }t          |dd¬¦  «        \  }}|                      |¦  «         d S )Nc                  ó   — d S r,   r-   r-   r   r   r   z6TestCompileToPTX.test_kernel_with_line_info.<locals>.fz   rE   r   r-   T)rQ   rR   rH   s       r   Útest_kernel_with_line_infoz+TestCompileToPTX.test_kernel_with_line_infoy   sE   € ð	ð 	ð 	õ !  B°Ð6Ñ6Ô6‰
ˆˆUØ×Ò˜SÑ!Ô!Ð!Ð!Ð!r   c           	      óÐ   — d„ }|                       t          d¦  «        5  t          |t          d d d…         t          d d d…         f¦  «         d d d ¦  «         d S # 1 swxY w Y   d S )Nc                 ó$   — | d         |d         z   S )Nr   r-   r.   s     r   r   z5TestCompileToPTX.test_non_void_return_type.<locals>.f�   s   € Ø�Q”4˜!˜Aœ$‘;Ðr   zmust have void return typer   )ÚassertRaisesRegexÚ	TypeErrorr   r	   )r%   r   s     r   Útest_non_void_return_typez*TestCompileToPTX.test_non_void_return_type€   s·   € ð	ð 	ð 	ð ×#Ò#¥IÐ/KÑLÔLð 	7ð 	7Ý˜�F 3 3 Q 3œK­°°°!°¬Ð5Ñ6Ô6Ð6ð	7ð 	7ð 	7ñ 	7ô 	7ð 	7ð 	7ð 	7ð 	7ð 	7ð 	7ð 	7øøøð 	7ð 	7ð 	7ð 	7ð 	7ð 	7s   Ÿ/AÁAÁ"AN)Ú__name__Ú
__module__Ú__qualname__r)   r6   r<   rB   rI   rL   rN   rS   rV   r[   r-   r   r   r   r      sª   € € € € € ð&ð &ð &ð$(ð (ð (ð<2ð 2ð 2ð(<ð <ð <ð#ð #ð #ð#ð #ð #ð<ð <ð <ð"ð "ð "ð"ð "ð "ð7ð 7ð 7ð 7ð 7r   r   c                   ó   — e Zd Zd„ ZdS )Ú TestCompileToPTXForCurrentDevicec                 ó  — d„ }t           t           f}t          ||d¬¦  «        \  }}t          j        ¦   «         j        }t          j        j                             |¦  «        }d|d         › |d         › �}|                      ||¦  «         d S )Nc                 ó   — | |z   S r,   r-   r.   s     r   r/   zQTestCompileToPTXForCurrentDevice.test_compile_ptx_for_current_device.<locals>.add‹   r0   r   Tr1   z.target sm_r   r   )	r   r   r   Úget_current_deviceÚcompute_capabilityÚcudadrvÚnvvmÚfind_closest_archr"   )r%   r/   r&   r'   r(   Ú	device_ccÚccÚtargets           r   Ú#test_compile_ptx_for_current_devicezDTestCompileToPTXForCurrentDevice.test_compile_ptx_for_current_deviceŠ   s”   € ð	ð 	ð 	õ �Ð!ˆÝ3°C¸ÀdÐKÑKÔK‰
ˆˆUõ Ô+Ñ-Ô-Ô@ˆ	ÝŒ\Ô×0Ò0°Ñ;Ô;ˆØ-˜r !œuÐ- b¨¤eÐ-Ð-ˆØ�Š�f˜cÑ"Ô"Ð"Ð"Ð"r   N)r\   r]   r^   rk   r-   r   r   r`   r`   ˆ   s#   € € € € € ð#ð #ð #ð #ð #r   r`   c                   ó   — e Zd ZdZd„ ZdS )ÚTestCompileOnlyTestsz‚For tests where we can only check correctness by examining the compiler
    output rather than observing the effects of execution.c                 óÆ   — d„ }t          |t          fd¬¦  «        \  }}d}|                     d¦  «        D ]}d|v r|dz  }Œd}|                      ||d	|› d
|› �¦  «         d S )Nc                 óV   — t          j        d¦  «         t          j        | ¦  «         d S )Né    )r   Ú	nanosleep)r   s    r   Úuse_nanosleepz:TestCompileOnlyTests.test_nanosleep.<locals>.use_nanosleepŸ   s(   € åŒN˜2ÑÔÐåŒN˜1ÑÔÐÐÐr   )é   r   )ri   r   ú
znanosleep.u32r   é   zGot z" nanosleep instructions, expected )r   r	   Úsplitr#   )r%   rr   r'   r(   Únanosleep_countÚlineÚexpecteds          r   Útest_nanosleepz#TestCompileOnlyTests.test_nanosleepž   sª   € ð	ð 	ð 	õ ! µ°	¸fÐEÑEÔE‰
ˆˆUàˆØ—I’I˜d‘O”Oð 	%ð 	%ˆDØ $Ð&Ð&Ø 1Ñ$�øàˆØ×Ò˜ ?ð1 ð 1ð 1Ø&.ð1ð 1ñ	3ô 	3ð 	3ð 	3ð 	3r   N)r\   r]   r^   Ú__doc__rz   r-   r   r   rm   rm   ™   s-   € € € € € ð>ð >ð3ð 3ð 3ð 3ð 3r   rm   Ú__main__N)Úmathr   Únumbar   r   r   r   r	   r
   Ú
numba.cudar   r   Únumba.cuda.testingr   r   r   ÚTestCaser   r`   rm   r\   Úmainr-   r   r   ú<module>rƒ      sx  ðØ Ð Ð Ð Ð Ð Ø ;Ð ;Ð ;Ð ;Ð ;Ð ;Ð ;Ð ;Ð ;Ð ;Ð ;Ð ;Ð ;Ð ;Ð ;Ð ;Ø BÐ BÐ BÐ BÐ BÐ BÐ BÐ Bà FÐ FÐ FÐ FÐ FÐ FÐ FÐ FÐ FÐ Fð €Ð;Ñ<Ô<ð|7ð |7ð |7ð |7ð |7�xÔ(ñ |7ô |7ñ =Ô<ð|7ð~ €Ð;Ñ<Ô<ð#ð #ð #ð #ð # |ñ #ô #ñ =Ô<ð#ð  €Ð;Ñ<Ô<ð3ð 3ð 3ð 3ð 3˜8Ô,ñ 3ô 3ñ =Ô<ð3ð0 ˆzÒÐØ€H„M�O„O€O€O€Oð Ðr   