
    ܙdW                     ^   d dl mZ d dlmZmZmZmZmZmZ d dl	m
Z
mZ d dlmZ d dlmZmZm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uint32void)compile_ptxcompile_ptx_for_current_device)NVVM)skip_on_cudasimunittestCUDATestCasez(Compilation unsupported in the simulatorc                   J    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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)rxyis       Elib/python3.11/site-packages/numba/cuda/tests/cudapy/test_compiler.pyfz.TestCompileToPTX.test_global_kernel.<locals>.f   s>    	!A3q66zztad{! z    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 N 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&   r0   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>.f<   s    Qa(((r   Tr2   z
fma.rn.f32z
div.rn.f32zsqrt.rn.f32)r3   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                     t                      j        s|                     d           |                     |d           |                     |d           d S )Nz$debuginfo not generated for NVVM 3.4z\.section\s+\.debug_info\.file.*test_compiler.py")r   	is_nvvm70skipTestassertRegexr&   r(   s     r   check_debug_infoz!TestCompileToPTX.check_debug_infoO   s]    vv 	BMM@AAA 	;<<< 	:;;;;;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>.fb       Dr   r.   T)r3   debugr   rD   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>.fj   rG   r   r.   T)rH   rI   rJ   s       r   test_kernel_with_debugz'TestCompileToPTX.test_kernel_with_debugh   sE    	 	 	 !Bd333
Uc"""""r   c                 2    |                      |d           d S )Nr?   )rB   rC   s     r   check_line_infoz TestCompileToPTX.check_line_infop   s!     	:;;;;;r   c                     t                      j        s|                     d           d }t          |ddd          \  }}|                     |           d S )N#lineinfo not generated for NVVM 3.4c                      d S r-   r.   r.   r   r   r   z?TestCompileToPTX.test_device_function_with_line_info.<locals>.fz   rG   r   r.   T)r3   lineinfor   r@   rA   r   rP   rJ   s       r   #test_device_function_with_line_infoz4TestCompileToPTX.test_device_function_with_line_infov   sj    vv 	AMM?@@@	 	 	 !BtdCCC
US!!!!!r   c                     t                      j        s|                     d           d }t          |dd          \  }}|                     |           d S )NrR   c                      d S r-   r.   r.   r   r   r   z6TestCompileToPTX.test_kernel_with_line_info.<locals>.f   rG   r   r.   T)rT   rU   rJ   s       r   test_kernel_with_line_infoz+TestCompileToPTX.test_kernel_with_line_info   sh    vv 	AMM?@@@	 	 	 !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N)__name__
__module____qualname__r*   r7   r=   rD   rK   rN   rP   rV   rY   r^   r.   r   r   r   r   	   s        & & &$( ( (<2 2 2(
< 
< 
<# # ## # #< < <" " "" " "7 7 7 7 7r   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   r0   zQTestCompileToPTXForCurrentDevice.test_compile_ptx_for_current_device.<locals>.add   r1   r   Tr2   z.target sm_r   r   )	r   r   r   get_current_devicecompute_capabilitycudadrvnvvmfind_closest_archr#   )r&   r0   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)r_   r`   ra   rn   r.   r   r   rc   rc      s#        # # # # #r   rc   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   )rl   r   
znanosleep.u32r      zGot z" nanosleep instructions, expected )r   r	   splitr$   )r&   ru   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)r_   r`   ra   __doc__r}   r.   r   r   rp   rp      s-        > >3 3 3 3 3r   rp   __main__N)mathr   numbar   r   r   r   r	   r
   
numba.cudar   r   numba.cuda.cudadrv.nvvmr   numba.cuda.testingr   r   r   TestCaser   rc   rp   r_   mainr.   r   r   <module>r      s         ; ; ; ; ; ; ; ; ; ; ; ; ; ; ; ; B B B B B B B B ( ( ( ( ( ( F F F F F F F F F F ;<<E7 E7 E7 E7 E7x( E7 E7 =<E7P ;<<# # # # #| # # =<#  ;<<3 3 3 3 38, 3 3 =<30 zHMOOOOO r   