
    iE*                     N   S SK Jr  S SKJrJrJrJrJrJrJ	r	  S SK
JrJrJrJr  S SKJr  S SKJrJrJr  S r\" S5       " S S	\R,                  5      5       r\" S5       " S
 S\5      5       r\" S5       " S S\R,                  5      5       r\S:X  a  \R6                  " 5         gg)    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                 
    X-   $ N xys     `/var/www/html/trading/venv/lib/python3.13/site-packages/numba/cuda/tests/cudapy/test_compiler.pyf_moduler      s	    5L    z(Compilation unsupported in the simulatorc                       \ rS rSrS rS rS rS rS rS r	S r
S	 rS
 rS rS rS rS rS rS rS rS rS rSrg)TestCompile   c                     S n[         S S  [         S S  [         S S  4n[        X5      u  p4U R                  SU5        U R                  SU5        U R                  SU5        U R	                  U[
        5        g )Nc                 f    [         R                  " S5      nU[        U 5      :  a  X   X#   -   X'   g g )N   )r   gridlen)rr   r   is       r   f)TestCompile.test_global_kernel.<locals>.f   s.    		!A3q6ztad{ r   func_retval.visible .func.visible .entry)r   r   assertNotInassertInassertEqualr   selfr'   argsptxrestys        r   test_global_kernelTestCompile.test_global_kernel   so    	#
 
GAJ
3 )
 	,)3/'-%r   c                    S n[         [         4n[        XSS9u  p4U R                  SU5        U R                  SU5        U R                  SU5        U R	                  U[         5        [        [
        [
        5      n[        XSS9u  p4U R	                  U[
        5        [        [        [        5      n[        XSS9u  p4U R	                  U[        5        Sn[        XSS9u  p4U R	                  U[        5        g )Nc                 
    X-   $ r   r   r   s     r   add-TestCompile.test_device_function.<locals>.add$   	    5Lr   Tdevicer)   r*   r+   zuint32(uint32, uint32))r   r   r-   r,   r.   r   r   r
   )r0   r8   r1   r2   r3   	sig_int32	sig_int16
sig_strings           r   test_device_function TestCompile.test_device_function#   s    	 ! 48
 	mS)&,*C0( %'	 =
&%'	 =
&-
 >
'r   c                 @   S n[         [         [         [         4n[        XSS9u  p4U R                  SU5        U R                  SU5        U R                  SU5        [        XSSS9u  p4U R                  SU5        U R                  S	U5        U R                  S
U5        g )Nc                 (    [        X-  U-   U-  5      $ r   r   )r   r   zds       r   r'   $TestCompile.test_fastmath.<locals>.fB   s    a((r   Tr;   z
fma.rn.f32z
div.rn.f32zsqrt.rn.f32)r<   fastmathzfma.rn.ftz.f32zdiv.approx.ftz.f32zsqrt.approx.ftz.f32)r   r   r-   r/   s        r   test_fastmathTestCompile.test_fastmathA   s    	) '73 6
 	lC(lC(mS) E
 	&,*C0+S1r   c                 L    U R                  US5        U R                  US5        g )Nz\.section\s+\.debug_info\.file.*test_compiler.py"assertRegexr0   r2   s     r   check_debug_infoTestCompile.check_debug_infoU   s(     	;< 	:;r   c                 H    S n[        USSSS9u  p#U R                  U5        g )Nc                      g r   r   r   r   r   r'   6TestCompile.test_device_function_with_debug.<locals>.fe       r   r   T)r<   debugr   rO   r0   r'   r2   r3   s       r   test_device_function_with_debug+TestCompile.test_device_function_with_debug^   s)    	 !Bt4@
c"r   c                 F    S n[        USSS9u  p#U R                  U5        g )Nc                      g r   r   r   r   r   r'   -TestCompile.test_kernel_with_debug.<locals>.fm   rT   r   r   T)rU   rV   rW   s       r   test_kernel_with_debug"TestCompile.test_kernel_with_debugk   s'    	 !Bd3
c"r   c                 (    U R                  US5        g )NrK   rL   rN   s     r   check_line_infoTestCompile.check_line_infos   s     	:;r   c                 H    S n[        USSSS9u  p#U R                  U5        g )Nc                      g r   r   r   r   r   r'   :TestCompile.test_device_function_with_line_info.<locals>.fz   rT   r   r   T)r<   lineinfor   r`   rW   s       r   #test_device_function_with_line_info/TestCompile.test_device_function_with_line_infoy   s)    	 !BtdC
S!r   c                 F    S n[        USSS9u  p#U R                  U5        g )Nc                      g r   r   r   r   r   r'   1TestCompile.test_kernel_with_line_info.<locals>.f   rT   r   r   T)re   rf   rW   s       r   test_kernel_with_line_info&TestCompile.test_kernel_with_line_info   s'    	 !B6
S!r   c           	          S nU R                  [        S5         [        U[        S S S2   [        S S S2   45        S S S 5        g ! , (       d  f       g = f)Nc                     U S   US   -   $ )Nr   r   r   s     r   r'   0TestCompile.test_non_void_return_type.<locals>.f   s    Q4!A$;r   zmust have void return typer"   )assertRaisesRegex	TypeErrorr   r
   r0   r'   s     r   test_non_void_return_type%TestCompile.test_non_void_return_type   sF    	 ##I/KLF3Q3K!56 MLLs   #A
Ac                     S nU R                  [        S5         [        U[        [        4SS9  S S S 5        g ! , (       d  f       g = f)Nc                 
    X-   $ r   r   r   s     r   r'   7TestCompile.test_c_abi_disallowed_for_kernel.<locals>.f   r:   r   z&The C ABI is not supported for kernelscabirq   NotImplementedErrorr   r   rs   s     r    test_c_abi_disallowed_for_kernel,TestCompile.test_c_abi_disallowed_for_kernel   sA    	 ##$7$LNE5>s3N N N	   9
Ac                     S nU R                  [        S5         [        U[        [        4SS9  S S S 5        g ! , (       d  f       g = f)Nc                 
    X-   $ r   r   r   s     r   r'   +TestCompile.test_unsupported_abi.<locals>.f   r:   r   zUnsupported ABI: fastcallfastcallrz   r|   rs   s     r   test_unsupported_abi TestCompile.test_unsupported_abi   sA    	 ##$7$?AE5>z:A A Ar   c                     S n[        U[        [        [        5      SSS9u  p#U R                  US5        U R                  US5        [        U[	        [        [        5      SSS9u  p#U R                  US5        g )Nc                 
    X-   $ r   r   r   s     r   r'   1TestCompile.test_c_abi_device_function.<locals>.f   r:   r   Try   r<   r{   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,   rM   r	   rW   s       r   test_c_abi_device_function&TestCompile.test_c_abi_device_function   su    	 !E%$7#N
i(
 	 6 	7
 !E%$7#N
GHr   c                 p    [        [        [        [        [        5      SSS9u  pU R                  US5        g )NTry   r   D\.visible\s+\.func\s+\(\.param\s+\.b32\s+func_retval0\)\s+f_module\(r   r   r   rM   r0   r2   r3   s      r   'test_c_abi_device_function_module_scope3TestCompile.test_c_abi_device_function_module_scope   s5     5+>t%(*

 	 = 	>r   c                 z    SS0n[        [        [        [        [        5      SSUS9u  p#U R                  US5        g )Nabi_name	_Z4funciiTry   )r<   r{   abi_infozE\.visible\s+\.func\s+\(\.param\s+\.b32\s+func_retval0\)\s+_Z4funcii\(r   )r0   r   r2   r3   s       r   test_c_abi_with_abi_name$TestCompile.test_c_abi_with_abi_name   sA    , 5+>t%(8=

 	 > 	?r   c                 n    [        [        [        [        [        5      SS9u  pU R                  US5        g )NTr;   r   )r   r   r   rM   r   s      r   test_compile_defaults_to_c_abi*TestCompile.test_compile_defaults_to_c_abi   s0    XuUE':4H
 	 = 	>r   c                     [         R                  " 5       S:  a  U R                  S5        [        [        [        [
        [
        5      SSS9u  pSn[        R                  US S SS	9nU R                  XC5        U R                  U[
        5        g )
N)      z,-gen-lto unavailable in this toolkit versionTltoirr<   outputiCN   little)	byteorder)	r   get_versionskipTestr   r   r   int
from_bytesr.   )r0   r   r3   LTOIR_MAGICheaders        r   test_compile_to_ltoir!TestCompile.test_compile_to_ltoir   sy     7*MMHIxue)<T&-/ !bq	X>-&r   c                     SnSU 3nU R                  [        U5         [        [        [	        [        [        5      SUS9  S S S 5        g ! , (       d  f       g = f)NillegalzUnsupported output type: Tr   )rq   r}   r   r   r   )r0   illegal_outputmsgs      r   test_compile_to_invalid_error)TestCompile.test_compile_to_invalid_error   sI    ").)9:##$7=HeE51$)+ >==s   #A


Ar   N)__name__
__module____qualname____firstlineno__r4   r@   rH   rO   rX   r]   r`   rg   rl   rt   r~   r   r   r   r   r   r   r   __static_attributes__r   r   r   r   r      s]    &$(<2(<##<""74;I&>?>'+r   r   c                   &    \ rS rSrS rS rS rSrg)TestCompileForCurrentDevice   c                     S n[         [         4nU" X#SS9u  pE[        R                  " 5       R                  n[        R                  R
                  R                  U5      nSUS    US    3nU R                  X5        g )Nc                 
    X-   $ r   r   r   s     r   r8   FTestCompileForCurrentDevice._check_ptx_for_current_device.<locals>.add   r:   r   Tr;   z.target sm_r   r"   )r   r   get_current_devicecompute_capabilitycudadrvnvvmfind_closest_archr-   )	r0   compile_functionr8   r1   r2   r3   	device_cccctargets	            r   _check_ptx_for_current_device9TestCompileForCurrentDevice._check_ptx_for_current_device   su    	 !%c=
 ++-@@	\\00;r!ugbeW-f"r   c                 .    U R                  [        5        g r   )r   r   r0   s    r   #test_compile_ptx_for_current_device?TestCompileForCurrentDevice.test_compile_ptx_for_current_device   s    **+IJr   c                 .    U R                  [        5        g r   )r   r   r   s    r   test_compile_for_current_device;TestCompileForCurrentDevice.test_compile_for_current_device   s    **+EFr   r   N)r   r   r   r   r   r   r   r   r   r   r   r   r      s    #KGr   r   c                       \ rS rSrSrS rSrg)TestCompileOnlyTests   z~For tests where we can only check correctness by examining the compiler
output rather than observing the effects of execution.c                     S n[        U[        4SS9u  p#SnUR                  S5       H  nSU;   d  M  US-  nM     SnU R                  XdS	U S
U 35        g )Nc                 \    [         R                  " S5        [         R                  " U 5        g )N    )r   	nanosleep)r   s    r   use_nanosleep:TestCompileOnlyTests.test_nanosleep.<locals>.use_nanosleep   s    NN2NN1r   )   r   )r   r   
znanosleep.u32r"      zGot z" nanosleep instructions, expected )r   r
   splitr.   )r0   r   r2   r3   nanosleep_countlineexpecteds          r   test_nanosleep#TestCompileOnlyTests.test_nanosleep   sr    	 !	fE
IIdOD$&1$ $   1 2&&.Z1	3r   r   N)r   r   r   r   __doc__r   r   r   r   r   r   r      s    >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8 8 & F F
 ;<P+(## P+ =P+f ;<G, G =G, ;<38,, 3 =30 zMMO r   