§
    1@ÄjY+  ã                   óp  — d dl mZ d dlmZmZmZmZmZmZm	Z	 d dl
mZmZmZmZ d dlmZ d dlmZmZmZ d„ 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Úint64Úuint32Úvoid)ÚcompileÚcompile_for_current_deviceÚcompile_ptxÚcompile_ptx_for_current_device)Úruntime)Úskip_on_cudasimÚunittestÚCUDATestCasec                 ó   — | |z   S ©N© ©ÚxÚys     úVC:\Projects\Historical\venv\Lib\site-packages\numba/cuda/tests/cudapy/test_compiler.pyÚf_moduler      s   € Øˆq‰5€Ló    z(Compilation unsupported in the simulatorc                   óz   — 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„ Zd„ Zd„ ZdS )ÚTestCompilec                 ó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)Úrr   r   Úis       r   Úfz)TestCompile.test_global_kernel.<locals>.f   s>   € Ý”	˜!‘”ˆAØ•3�q‘6”6ŠzˆzØ˜”t˜a œd‘{��!‘��ð ˆzr   Úfunc_retvalú.visible .funcú.visible .entry)r   r   ÚassertNotInÚassertInÚassertEqualr   ©Úselfr&   ÚargsÚptxÚrestys        r   Útest_global_kernelzTestCompile.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 r   r   r   s     r   Úaddz-TestCompile.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.   r5   r/   r0   r1   Ú	sig_int32Ú	sig_int16Ú
sig_strings           r   Útest_device_functionz TestCompile.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$TestCompile.test_fastmath.<locals>.fB   s   € Ý˜˜Q™ ™ a™Ñ(Ô(Ð(r   Tr7   z
fma.rn.f32z
div.rn.f32zsqrt.rn.f32)r8   Úfastmathzfma.rn.ftz.f32zdiv.approx.ftz.f32zsqrt.approx.ftz.f32)r   r   r+   r-   s        r   Útest_fastmathzTestCompile.test_fastmathA   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.   r0   s     r   Úcheck_debug_infozTestCompile.check_debug_infoU   s:   € ð 	×Ò˜Ð;Ñ<Ô<Ð<ð 	×Ò˜Ð:Ñ;Ô;Ð;Ð;Ð;r   c                 ób   — d„ }t          |ddd¬¦  «        \  }}|                      |¦  «         d S )Nc                  ó   — d S r   r   r   r   r   r&   z6TestCompile.test_device_function_with_debug.<locals>.fe   ó   € ØˆDr   r   T)r8   Údebug©r   rH   ©r.   r&   r0   r1   s       r   Útest_device_function_with_debugz+TestCompile.test_device_function_with_debug^   sG   € ð	ð 	ð 	õ !  B¨t¸4Ð@Ñ@Ô@‰
ˆˆUØ×Ò˜cÑ"Ô"Ð"Ð"Ð"r   c                 ó`   — d„ }t          |dd¬¦  «        \  }}|                      |¦  «         d S )Nc                  ó   — d S r   r   r   r   r   r&   z-TestCompile.test_kernel_with_debug.<locals>.fm   rK   r   r   T)rL   rM   rN   s       r   Útest_kernel_with_debugz"TestCompile.test_kernel_with_debugk   sE   € ð	ð 	ð 	õ !  B¨dÐ3Ñ3Ô3‰
ˆˆUØ×Ò˜cÑ"Ô"Ð"Ð"Ð"r   c                 ó2   — |                       |d¦  «         d S )NrD   rE   rG   s     r   Úcheck_line_infozTestCompile.check_line_infos   s!   € ð 	×Ò˜Ð:Ñ;Ô;Ð;Ð;Ð;r   c                 ób   — d„ }t          |ddd¬¦  «        \  }}|                      |¦  «         d S )Nc                  ó   — d S r   r   r   r   r   r&   z:TestCompile.test_device_function_with_line_info.<locals>.fz   rK   r   r   T)r8   Úlineinfo©r   rT   rN   s       r   Ú#test_device_function_with_line_infoz/TestCompile.test_device_function_with_line_infoy   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&   z1TestCompile.test_kernel_with_line_info.<locals>.f�   rK   r   r   T)rW   rX   rN   s       r   Útest_kernel_with_line_infoz&TestCompile.test_kernel_with_line_info€   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&   z0TestCompile.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%TestCompile.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c                 ó°   — d„ }|                       t          d¦  «        5  t          |t          t          fd¬¦  «         d d d ¦  «         d S # 1 swxY w Y   d S )Nc                 ó   — | |z   S r   r   r   s     r   r&   z7TestCompile.test_c_abi_disallowed_for_kernel.<locals>.f�   r6   r   z&The C ABI is not supported for kernelsÚc©Úabi©r_   ÚNotImplementedErrorr   r   ra   s     r   Ú test_c_abi_disallowed_for_kernelz,TestCompile.test_c_abi_disallowed_for_kernelŽ   s®   € ð	ð 	ð 	ð ×#Ò#Õ$7Ø$LñNô Nð 	4ð 	4å˜�E¥5˜>¨sÐ3Ñ3Ô3Ð3ð	4ð 	4ð 	4ñ 	4ô 	4ð 	4ð 	4ð 	4ð 	4ð 	4ð 	4ð 	4øøøð 	4ð 	4ð 	4ð 	4ð 	4ð 	4ó   ŸAÁAÁAc                 ó°   — d„ }|                       t          d¦  «        5  t          |t          t          fd¬¦  «         d d d ¦  «         d S # 1 swxY w Y   d S )Nc                 ó   — | |z   S r   r   r   s     r   r&   z+TestCompile.test_unsupported_abi.<locals>.f—   r6   r   zUnsupported ABI: fastcallÚfastcallrf   rh   ra   s     r   Útest_unsupported_abiz TestCompile.test_unsupported_abi–   s®   € ð	ð 	ð 	ð ×#Ò#Õ$7Ø$?ñAô Að 	;ð 	;å˜�E¥5˜>¨zÐ:Ñ:Ô:Ð:ð	;ð 	;ð 	;ñ 	;ô 	;ð 	;ð 	;ð 	;ð 	;ð 	;ð 	;ð 	;øøøð 	;ð 	;ð 	;ð 	;ð 	;ð 	;rk   c                 óH  — d„ }t          |t          t          t          ¦  «        dd¬¦  «        \  }}|                      |d¦  «         |                      |d¦  «         t          |t	          t          t          ¦  «        dd¬¦  «        \  }}|                      |d¦  «         d S )Nc                 ó   — | |z   S r   r   r   s     r   r&   z1TestCompile.test_c_abi_device_function.<locals>.fŸ   r6   r   Tre   ©r8   rg   Úparam_2z=\.visible\s+\.func\s+\(\.param\s+\.b32\s+func_retval0\)\s+f\(z&\.visible\s+\.func\s+\(\.param\s+\.b64)r   r   r*   rF   r	   rN   s       r   Útest_c_abi_device_functionz&TestCompile.test_c_abi_device_functionž   s±   € ð	ð 	ð 	õ ! ¥E­%µÑ$7Ô$7ÀÈ#ÐNÑNÔN‰
ˆˆUà×Ò˜˜iÑ(Ô(Ð(ð
 	×Ò˜ð 6ñ 	7ô 	7ð 	7õ
 ! ¥E­%µÑ$7Ô$7ÀÈ#ÐNÑNÔN‰
ˆˆUØ×Ò˜ÐGÑHÔHÐHÐHÐHr   c                 ó˜   — t          t          t          t          t          ¦  «        dd¬¦  «        \  }}|                      |d¦  «         d S )NTre   rr   úD\.visible\s+\.func\s+\(\.param\s+\.b32\s+func_retval0\)\s+f_module\(©r   r   r   rF   ©r.   r0   r1   s      r   Ú'test_c_abi_device_function_module_scopez3TestCompile.test_c_abi_device_function_module_scope±   sX   € Ý ¥­5µ½Ñ+>Ô+>ÀtØ%(ð*ñ *ô *‰
ˆˆUð
 	×Ò˜ð =ñ 	>ô 	>ð 	>ð 	>ð 	>r   c                 ó¢   — ddi}t          t          t          t          t          ¦  «        dd|¬¦  «        \  }}|                      |d¦  «         d S )NÚabi_nameÚ	_Z4funciiTre   )r8   rg   Úabi_infozE\.visible\s+\.func\s+\(\.param\s+\.b32\s+func_retval0\)\s+_Z4funcii\(rw   )r.   r}   r0   r1   s       r   Útest_c_abi_with_abi_namez$TestCompile.test_c_abi_with_abi_nameº   sd   € Ø Ð,ˆÝ ¥­5µ½Ñ+>Ô+>ÀtØ%(°8ð=ñ =ô =‰
ˆˆUð
 	×Ò˜ð >ñ 	?ô 	?ð 	?ð 	?ð 	?r   c                 ó–   — t          t          t          t          t          ¦  «        d¬¦  «        \  }}|                      |d¦  «         d S )NTr7   rv   )r   r   r   rF   rx   s      r   Útest_compile_defaults_to_c_abiz*TestCompile.test_compile_defaults_to_c_abiÄ   sO   € Ý�X¥u­UµEÑ':Ô':À4ÐHÑHÔH‰
ˆˆUð 	×Ò˜ð =ñ 	>ô 	>ð 	>ð 	>ð 	>r   c                 ór  — t          j        ¦   «         dk     r|                      d¦  «         t          t          t          t
          t
          ¦  «        dd¬¦  «        \  }}d}t                               |d d…         d¬	¦  «        }|                      ||¦  «         |                      |t
          ¦  «         d S )
N)é   é   z,-gen-lto unavailable in this toolkit versionTÚltoir©r8   ÚoutputiíCNé   Úlittle)Ú	byteorder)	r   Úget_versionÚskipTestr   r   r   ÚintÚ
from_bytesr,   )r.   r„   r1   ÚLTOIR_MAGICÚheaders        r   Útest_compile_to_ltoirz!TestCompile.test_compile_to_ltoirÌ   sª   € ÝÔÑ Ô  7Ò*Ð*Ø�MŠMÐHÑIÔIÐIå�x­­uµeÑ)<Ô)<ÀTØ&-ð/ñ /ô /‰ˆˆuð !ˆÝ—’  b q b¤	°X�Ñ>Ô>ˆØ×Ò˜ Ñ-Ô-Ð-Ø×Ò˜¥Ñ&Ô&Ð&Ð&Ð&r   c                 óÜ   — d}d|› �}|                       t          |¦  «        5  t          t          t	          t          t          ¦  «        d|¬¦  «         d d d ¦  «         d S # 1 swxY w Y   d S )NÚillegalzUnsupported output type: Tr…   )r_   ri   r   r   r   )r.   Úillegal_outputÚmsgs      r   Útest_compile_to_invalid_errorz)TestCompile.test_compile_to_invalid_errorÛ   s¹   € Ø"ˆØ:¨.Ð:Ð:ˆØ×#Ò#Õ$7¸Ñ=Ô=ð 	+ð 	+Ý•H�e¥E­5Ñ1Ô1¸$Ø)ð+ñ +ô +ð +ð	+ð 	+ð 	+ñ 	+ô 	+ð 	+ð 	+ð 	+ð 	+ð 	+ð 	+ð 	+øøøð 	+ð 	+ð 	+ð 	+ð 	+ð 	+s   £1A!Á!A%Á(A%N)Ú__name__Ú
__module__Ú__qualname__r2   r<   rB   rH   rO   rR   rT   rY   r\   rb   rj   ro   rt   ry   r~   r€   r�   r•   r   r   r   r   r      s%  € € € € € ð&ð &ð &ð$(ð (ð (ð<2ð 2ð 2ð(<ð <ð <ð#ð #ð #ð#ð #ð #ð<ð <ð <ð"ð "ð "ð"ð "ð "ð7ð 7ð 7ð4ð 4ð 4ð;ð ;ð ;ðIð Ið Ið&>ð >ð >ð?ð ?ð ?ð>ð >ð >ð'ð 'ð 'ð+ð +ð +ð +ð +r   r   c                   ó    — e Zd Zd„ Zd„ Zd„ ZdS )ÚTestCompileForCurrentDevicec                 ó  — d„ }t           t           f} |||d¬¦  «        \  }}t          j        ¦   «         j        }t          j        j                             |¦  «        }d|d         › |d         › �}|                      ||¦  «         d S )Nc                 ó   — | |z   S r   r   r   s     r   r5   zFTestCompileForCurrentDevice._check_ptx_for_current_device.<locals>.addæ   r6   r   Tr7   z.target sm_r   r!   )r   r   Úget_current_deviceÚcompute_capabilityÚcudadrvÚnvvmÚfind_closest_archr+   )	r.   Úcompile_functionr5   r/   r0   r1   Ú	device_ccÚccÚtargets	            r   Ú_check_ptx_for_current_devicez9TestCompileForCurrentDevice._check_ptx_for_current_deviceå   s—   € ð	ð 	ð 	õ �Ð!ˆØ%Ð% c¨4¸Ð=Ñ=Ô=‰
ˆˆUõ Ô+Ñ-Ô-Ô@ˆ	ÝŒ\Ô×0Ò0°Ñ;Ô;ˆØ-˜r !œuÐ- b¨¤eÐ-Ð-ˆØ�Š�f˜cÑ"Ô"Ð"Ð"Ð"r   c                 ó:   — |                       t          ¦  «         d S r   )r¦   r   ©r.   s    r   Ú#test_compile_ptx_for_current_devicez?TestCompileForCurrentDevice.test_compile_ptx_for_current_deviceó   s   € Ø×*Ò*Õ+IÑJÔJÐJÐJÐJr   c                 ó:   — |                       t          ¦  «         d S r   )r¦   r   r¨   s    r   Útest_compile_for_current_devicez;TestCompileForCurrentDevice.test_compile_for_current_deviceö   s   € Ø×*Ò*Õ+EÑFÔFÐFÐFÐFr   N)r–   r—   r˜   r¦   r©   r«   r   r   r   rš   rš   ã   sI   € € € € € ð#ð #ð #ðKð Kð KðGð Gð Gð Gð G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   )r¤   r   ú
znanosleep.u32r!   é   zGot z" nanosleep instructions, expected )r   r
   Úsplitr,   )r.   r²   r0   r1   Ú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__rº   r   r   r   r­   r­   ú   s-   € € € € € ð>ð >ð3ð 3ð 3ð 3ð 3r   r­   Ú__main__N)Úmathr   Únumbar   r   r   r   r	   r
   r   Ú
numba.cudar   r   r   r   Únumba.cuda.cudadrvr   Únumba.cuda.testingr   r   r   r   ÚTestCaser   rš   r­   r–   Úmainr   r   r   ú<module>rÄ      sË  ðØ Ð Ð Ð Ð Ð Ø BÐ BÐ BÐ BÐ BÐ BÐ BÐ BÐ BÐ BÐ BÐ BÐ BÐ BÐ BÐ BÐ BÐ Bð8ð 8ð 8ð 8ð 8ð 8ð 8ð 8ð 8ð 8ð 8ð 8à &Ð &Ð &Ð &Ð &Ð &Ø FÐ FÐ FÐ FÐ FÐ FÐ FÐ FÐ FÐ Fð
ð ð ð €Ð;Ñ<Ô<ðP+ð P+ð P+ð P+ð P+�(Ô#ñ P+ô P+ñ =Ô<ðP+ðf €Ð;Ñ<Ô<ðGð Gð Gð Gð G ,ñ Gô Gñ =Ô<ðGð, €Ð;Ñ<Ô<ð3ð 3ð 3ð 3ð 3˜8Ô,ñ 3ô 3ñ =Ô<ð3ð0 ˆzÒÐØ€H„M�O„O€O€O€Oð Ðr   