
    ,i#                     \   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 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_ptxcompile_ptx_for_current_device)skip_on_cudasimunittestCUDATestCasec                     | |z   S N xys     c/var/www/html/speakWrite/venv/lib/python3.11/site-packages/numba/cuda/tests/cudapy/test_compiler.pyf_moduler   
   s    q5L    z(Compilation unsupported in the simulatorc                   h    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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)rr   r   is       r   fz.TestCompileToPTX.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#TestCompileToPTX.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   addz2TestCompileToPTX.test_device_function.<locals>.add#       q5Lr   Tdevicer$   r%   r&   zuint32(uint32, uint32))r   r   r(   r'   r)   r   r   r
   )r+   r2   r,   r-   r.   	sig_int32	sig_int16
sig_strings           r   test_device_functionz%TestCompileToPTX.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)TestCompileToPTX.test_fastmath.<locals>.fA   s    Qa(((r   Tr4   z
fma.rn.f32z
div.rn.f32zsqrt.rn.f32)r5   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 	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+   r-   s     r   check_debug_infoz!TestCompileToPTX.check_debug_infoT   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>.fd       Dr   r   T)r5   debugr   rE   r+   r#   r-   r.   s       r   test_device_function_with_debugz0TestCompileToPTX.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#   z2TestCompileToPTX.test_kernel_with_debug.<locals>.fl   rH   r   r   T)rI   rJ   rK   s       r   test_kernel_with_debugz'TestCompileToPTX.test_kernel_with_debugj   sE    	 	 	 !Bd333
Uc"""""r   c                 2    |                      |d           d S )NrA   rB   rD   s     r   check_line_infoz TestCompileToPTX.check_line_infor   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>.fy   rH   r   r   T)r5   lineinfor   rQ   rK   s       r   #test_device_function_with_line_infoz4TestCompileToPTX.test_device_function_with_line_infox   sG    	 	 	 !Btd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>.f   rH   r   r   T)rT   rU   rK   s       r   test_kernel_with_line_infoz+TestCompileToPTX.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#   z5TestCompileToPTX.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*TestCompileToPTX.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#   z<TestCompileToPTX.test_c_abi_disallowed_for_kernel.<locals>.f   r3   r   z&The C ABI is not supported for kernelscabir\   NotImplementedErrorr   r   r^   s     r    test_c_abi_disallowed_for_kernelz1TestCompileToPTX.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#   z0TestCompileToPTX.test_unsupported_abi.<locals>.f   r3   r   zUnsupported ABI: fastcallfastcallrc   re   r^   s     r   test_unsupported_abiz%TestCompileToPTX.test_unsupported_abi   s    	 	 	 ##$7$?A A 	; 	;E5>z::::	; 	; 	; 	; 	; 	; 	; 	; 	; 	; 	; 	; 	; 	; 	; 	; 	; 	;rh   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#   z6TestCompileToPTX.test_c_abi_device_function.<locals>.f   r3   r   Trb   r5   rd   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'   rC   r	   rK   s       r   test_c_abi_device_functionz+TestCompileToPTX.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 )NTrb   ro   zD\.visible\s+\.func\s+\(\.param\s+\.b32\s+func_retval0\)\s+f_module\(r   r   r   rC   )r+   r-   r.   s      r   'test_c_abi_device_function_module_scopez8TestCompileToPTX.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	_Z4funciiTrb   )r5   rd   abi_infozE\.visible\s+\.func\s+\(\.param\s+\.b32\s+func_retval0\)\s+_Z4funcii\(rs   )r+   rx   r-   r.   s       r   test_c_abi_with_abi_namez)TestCompileToPTX.test_c_abi_with_abi_name   sd    , 5+>+>t%(8= = =
U
 	 > 	? 	? 	? 	? 	?r   N)__name__
__module____qualname__r/   r9   r?   rE   rL   rO   rQ   rV   rY   r_   rg   rl   rq   rt   ry   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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   r2   zQTestCompileToPTXForCurrentDevice.test_compile_ptx_for_current_device.<locals>.add   r3   r   Tr4   z.target sm_r   r   )	r   r   r   get_current_devicecompute_capabilitycudadrvnvvmfind_closest_archr(   )r+   r2   r,   r-   r.   	device_cccctargets           r   #test_compile_ptx_for_current_devicezDTestCompileToPTXForCurrentDevice.test_compile_ptx_for_current_device   s    	 	 	 !3CdKKK
U +--@	\00;;-r!u-be--fc"""""r   N)rz   r{   r|   r   r   r   r   r~   r~      s#        # # # # #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   r-   r.   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)rz   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   numba.cuda.testingr   r   r   r   TestCaser   r~   r   rz   mainr   r   r   <module>r      s         B B B B B B B B B B B B B B B B B B B B B B B B B B F F F F F F F F F F
   ;<<r? r? r? r? r?x( r? r? =<r?j ;<<# # # # #| # # =<#  ;<<3 3 3 3 38, 3 3 =<30 zHMOOOOO r   