
    J/PhE*                     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     e/var/www/html/test/jupyter/venv/lib/python3.11/site-packages/numba/cuda/tests/cudapy/test_compiler.pyf_moduler      s    q5L    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3q66zzta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    	# 	# 	#
 
GAAAJ
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$       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#   sB   	 	 	 ! d4888
U 	mS)))&,,,*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 	lC(((lC(((mS))) 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    	 	 	 !Bt4@@@
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    	 	 	 !Bd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    	 	 	 !Btd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    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   s    	 	 	 ##I/KLL 	7 	7F33Q3K!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E5>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 	; 	;E5>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    XuUE':':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   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IIxue)<)<T&-/ / /u !bqb	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== 	+ 	+HeE5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    	 	 	 !%%c4===
U +--@	\00;;-r!u-be--fc"""""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 )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}|                    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(    N2N1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IIdOO 	% 	%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 38, 3 3 =<30 zHMOOOOO r   