
    xKgE*                     L   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j6                          yy)    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     i/home/alanp/www/video.onchill/myenv/lib/python3.12/site-packages/numba/cuda/tests/cudapy/test_compiler.pyf_moduler      s    q5L    z(Compilation unsupported in the simulatorc                   x    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y)TestCompilec                     d }t         d d  t         d d  t         d d  f}t        ||      \  }}| j                  d|       | j                  d|       | j                  d|       | j	                  |t
               y )Nc                 h    t        j                  d      }|t        |       k  r||   ||   z   | |<   y y )N   )r   gridlen)rr   r   is       r   fz)TestCompile.test_global_kernel.<locals>.f   s4    		!A3q6ztad{! r   func_retval.visible .func.visible .entry)r   r   assertNotInassertInassertEqualr   selfr&   argsptxrestys        r   test_global_kernelzTestCompile.test_global_kernel   ss    	#
 
GAJ
3 D)
U 	,)3/'-%r   c                    d }t         t         f}t        ||d      \  }}| j                  d|       | j                  d|       | j                  d|       | j	                  |t                t        t
        t
              }t        ||d      \  }}| j	                  |t
               t        t        t              }t        ||d      \  }}| j	                  |t               d}t        ||d      \  }}| j	                  |t               y )Nc                     | |z   S r   r   r   s     r   addz-TestCompile.test_device_function.<locals>.add$       q5Lr   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#   s    	 ! d48
U 	mS)&,*C0( %'	 i=
U&%'	 i=
U&-
 j>
U'r   c                 T   d }t         t         t         t         f}t        ||d      \  }}| j                  d|       | j                  d|       | j                  d|       t        ||dd      \  }}| j                  d|       | j                  d	|       | j                  d
|       y )Nc                 *    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
U 	lC(lC(mS) DE
U 	&,*C0+S1r   c                 L    | j                  |d       | j                  |d       y )Nz\.section\s+\.debug_info\.file.*test_compiler.py"assertRegexr.   r0   s     r   check_debug_infozTestCompile.check_debug_infoU   s(     	;< 	:;r   c                 P    d }t        |ddd      \  }}| j                  |       y )Nc                       y r   r   r   r   r   r&   z6TestCompile.test_device_function_with_debug.<locals>.fe       r   r   T)r8   debugr   rH   r.   r&   r0   r1   s       r   test_device_function_with_debugz+TestCompile.test_device_function_with_debug^   s+    	 !Bt4@
Uc"r   c                 N    d }t        |dd      \  }}| j                  |       y )Nc                       y 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   s)    	 !Bd3
Uc"r   c                 (    | j                  |d       y )NrD   rE   rG   s     r   check_line_infozTestCompile.check_line_infos   s     	:;r   c                 P    d }t        |ddd      \  }}| j                  |       y )Nc                       y 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   s+    	 !BtdC
US!r   c                 N    d }t        |dd      \  }}| j                  |       y )Nc                       y 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   s)    	 !B6
US!r   c           	          d }| j                  t        d      5  t        |t        d d d   t        d d d   f       d d d        y # 1 sw Y   y xY w)Nc                     | d   |d   z   S )Nr   r   r   s     r   r&   z0TestCompile.test_non_void_return_type.<locals>.f   s    Q4!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   sF    	 ##I/KLF3Q3K!56 MLLs   #AAc                     d }| j                  t        d      5  t        |t        t        fd       d d d        y # 1 sw Y   y xY w)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   sA    	 ##$7$LNE5>s3N N N	   <Ac                     d }| j                  t        d      5  t        |t        t        fd       d d d        y # 1 sw Y   y xY w)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   sA    	 ##$7$?AE5>z:A A Ark   c                    d }t        |t        t        t              dd      \  }}| j                  |d       | j                  |d       t        |t	        t        t              dd      \  }}| j                  |d       y )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   sy    	 !E%$7#N
Ui(
 	 6 	7
 !E%$7#N
UGHr   c                 x    t        t        t        t        t              dd      \  }}| j                  |d       y )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   s7     5+>t%(*
U
 	 = 	>r   c                     ddi}t        t        t        t        t              dd|      \  }}| j                  |d       y )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   sC    , 5+>t%(8=
U
 	 > 	?r   c                 v    t        t        t        t        t              d      \  }}| j                  |d       y )NTr7   rv   )r   r   r   rF   rx   s      r   test_compile_defaults_to_c_abiz*TestCompile.test_compile_defaults_to_c_abi   s2    XuUE':4H
U 	 = 	>r   c                 ,   t        j                         dk  r| j                  d       t        t        t        t
        t
              dd      \  }}d}t        j                  |d d d	      }| j                  ||       | j                  |t
               y )
N)      z,-gen-lto unavailable in this toolkit versionTltoirr8   outputiCN   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*MMHIxue)<T&-/u !bq	X>-&r   c                     d}d| }| j                  t        |      5  t        t        t	        t        t              d|       d d d        y # 1 sw Y   y xY w)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   sI    ").)9:##$7=HeE51$)+ >==s   &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(<##<""74;I&>?>'+r   r   c                       e Zd Zd Zd Zd Zy)TestCompileForCurrentDevicec                 
   d }t         t         f} |||d      \  }}t        j                         j                  }t        j                  j
                  j                  |      }d|d    |d    }| j                  ||       y )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{    	 !%c4=
U ++-@@	\\00;r!ugbeW-fc"r   c                 .    | j                  t               y r   )r   r   r.   s    r   #test_compile_ptx_for_current_devicez?TestCompileForCurrentDevice.test_compile_ptx_for_current_device   s    **+IJr   c                 .    | j                  t               y r   )r   r   r   s    r   test_compile_for_current_devicez;TestCompileForCurrentDevice.test_compile_for_current_device   s    **+EFr   N)r   r   r   r   r   r   r   r   r   r   r      s    #KGr   r   c                       e Zd ZdZd Zy)TestCompileOnlyTestszFor 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}|j                  d      D ]  }d|v s|dz  } d}| j                  ||d	| d
|        y )Nc                 X    t        j                  d       t        j                  |        y )N    )r   	nanosleep)r   s    r   use_nanosleepz: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,   )r.   r   r0   r1   nanosleep_countlineexpecteds          r   test_nanosleepz#TestCompileOnlyTests.test_nanosleep   sv    	 !	fE
UIIdOD$&1$ $ ?  1 2&&.Z1	3r   N)r   r   r   __doc__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HMMO r   