
    J/Ph                         d dl mZ d dlmZ d dlmZ d dlmZ d dlmZ d dl	Z	d dl
Z
d dlZ ed           G d d	e                      Zed
k    r ej                     dS dS )    )override_config)skip_on_cudasim)cuda)types)CUDATestCaseNz&Simulator does not produce debug dumpsc                   v     e Zd ZdZ f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 xZS )TestCudaDebugInfozH
    These tests only checks the compiled PTX for debuginfo section
    c                 r    t                                                       |                     d           d S )Nz!Exceptions not supported with LTO)supersetUpskip_if_lto)self	__class__s    f/var/www/html/test/jupyter/venv/lib/python3.11/site-packages/numba/cuda/tests/cudapy/test_debuginfo.pyr   zTestCudaDebugInfo.setUp   s1     	<=====    c                 V    |                     |           |                    |          S N)compileinspect_asm)r   fnsigs      r   _getasmzTestCudaDebugInfo._getasm   s#    


3~~c"""r   c                     |                      ||          }t          j        d          }|                    |          }|r| j        n| j        } |||           d S )N)r   z\.section\s+\.debug_info\s+{)msg)r   rer   searchassertIsNotNoneassertIsNone)r   r   r   expectasmre_section_dbginfomatchassertfns           r   _checkzTestCudaDebugInfo._check   sk    ll23l''Z(GHH"))#..+1H4''t7HC      r   c                     t          j        d          d             }|                     |t          j        d d          fd           d S )NFdebugc                     d| d<   d S N   r    xs    r   fooz7TestCudaDebugInfo.test_no_debuginfo_in_asm.<locals>.foo&       AaDDDr   r   r   r   jitr$   r   int32r   r.   s     r   test_no_debuginfo_in_asmz*TestCudaDebugInfo.test_no_debuginfo_in_asm%   sV    					 	 
		 	Cek!!!n.u=====r   c                     t          j        dd          d             }|                     |t          j        d d          fd           d S )NTFr'   optc                     d| d<   d S r)   r+   r,   s    r   r.   z4TestCudaDebugInfo.test_debuginfo_in_asm.<locals>.foo-   r/   r   r0   r1   r4   s     r   test_debuginfo_in_asmz'TestCudaDebugInfo.test_debuginfo_in_asm,   sX    	%	(	(	(	 	 
)	(	 	Cek!!!n.t<<<<<r   c                 |   t          dd          5  t          j        d          d             }|                     |t          j        d d          fd           t          j        d          d	             }|                     |t          j        d d          fd           d d d            d S # 1 swxY w Y   d S )
NCUDA_DEBUGINFO_DEFAULTr*   F)r8   c                     d| d<   d S r)   r+   r,   s    r   r.   z8TestCudaDebugInfo.test_environment_override.<locals>.foo6       !r   Tr0   r&   c                     d| d<   d S r)   r+   r,   s    r   barz8TestCudaDebugInfo.test_environment_override.<locals>.bar=   r>   r   )r   r   r2   r$   r   r3   )r   r.   r@   s      r   test_environment_overridez+TestCudaDebugInfo.test_environment_override3   s)   5q99 	B 	BX%     !  KK%+aaa.!24K@@@ XE"""  #" KK%+aaa.!25KAAA	B 	B 	B 	B 	B 	B 	B 	B 	B 	B 	B 	B 	B 	B 	B 	B 	B 	Bs   BB11B58B5c                 n    t          j        t          j        d d d         fdd          d             }d S )Nr*   TFr7   c                     d| d<   d S )Nr   r+   r,   s    r   fz,TestCudaDebugInfo.test_issue_5835.<locals>.fG   r/   r   r   r2   r   r3   r   rD   s     r   test_issue_5835z!TestCudaDebugInfo.test_issue_5835C   sI     
5;sss#%Tu	=	=	=	 	 
>	=	 	 	r   c                 Z   t           j        d d d         f}t          j        |dd          d             }|                    |          }d |                                D             }|                     t          |          d           |d         }|                     d|           d S )Nr*   Tr   r7   c                     d| d<   d S r)   r+   r,   s    r   rD   z7TestCudaDebugInfo.test_wrapper_has_debuginfo.<locals>.fN   r/   r   c                     g | ]}d |v |	S )zdefine void @"_ZN6cudapyr+   ).0lines     r   
<listcomp>z@TestCudaDebugInfo.test_wrapper_has_debuginfo.<locals>.<listcomp>T   s,     : : :D0D88 888r   z!dbg)	r   r3   r   r2   inspect_llvm
splitlinesassertEquallenassertIn)r   r   rD   llvm_irdefineswrapper_defines         r   test_wrapper_has_debuginfoz,TestCudaDebugInfo.test_wrapper_has_debuginfoK   s    {33Q3!	#Tq	)	)	)	 	 
*	)	 ..%%: :G$6$6$8$8 : : : 	Wq))) fn-----r   c                     t          j        t          j        d d          t          j        d d          fdd          d             }d S )NTFr7   c                 (    | d         dv rdnd|d<   d S )Nr   )      r*   rZ   r+   )inpoutps     r   rD   zDTestCudaDebugInfo.test_debug_function_calls_internal_impl.<locals>.fk   s!    q6V++aaDGGGr   rE   rF   s     r   'test_debug_function_calls_internal_implz9TestCudaDebugInfo.test_debug_function_calls_internal_impl]   sS     
5;qqq>5;qqq>2$E	J	J	J	3 	3 
K	J	3 	3 	3r   c                     t          j        ddd          d             t          j        t          j        d d          fdd          fd            }d S )NTr   devicer'   r8   c                  l    t           j        j        t           j        j        z  t           j        j        z   S r   )r   blockDimr-   blockIdx	threadIdxr+   r   r   threadidzMTestCudaDebugInfo.test_debug_function_calls_device_function.<locals>.threadidt   s    =?T]_4t~7GGGr   r7   c                 t    t          j        d          }|t          |           k     r             | |<   d S d S Nr*   )r   gridrQ   )arrire   s     r   kernelzKTestCudaDebugInfo.test_debug_function_calls_device_function.<locals>.kernelx   s9    	!A3s88||!A |r   rE   )r   rk   re   s     @r   )test_debug_function_calls_device_functionz;TestCudaDebugInfo.test_debug_function_calls_device_functiono   s    
 
Tq	1	1	1	H 	H 
2	1	H 
5;qqq>#4Q	7	7	7	$ 	$ 	$ 	$ 
8	7	$ 	$ 	$r   c                 "   t          j        d|d          d             t          j        d|d          fd            t          j        t          j        t          j        f|d          fd            } |d         d	d
           d S )NTFr_   c                     | dz   S rg   r+   r,   s    r   f2z;TestCudaDebugInfo._test_chained_device_function.<locals>.f2       q5Lr   c                      |  |          z
  S r   r+   r-   yro   s     r   f1z;TestCudaDebugInfo._test_chained_device_function.<locals>.f1       rr!uu9r   r7   c                       | |           d S r   r+   r-   rs   rt   s     r   rk   z?TestCudaDebugInfo._test_chained_device_function.<locals>.kernel   s    Bq!HHHHHr   r*   r*   r*   rY   rE   r   kernel_debugf1_debugf2_debugrk   rt   ro   s        @@r   _test_chained_device_functionz/TestCudaDebugInfo._test_chained_device_function~   s    	X5	9	9	9	 	 
:	9	 
X5	9	9	9	 	 	 	 
:	9	 
5;,Le	L	L	L	 	 	 	 
M	L	 	tQr   c                     t          j        dgdz   }|D ]M\  }}}|                     |||          5  |                     |||           d d d            n# 1 swxY w Y   Nd S N)TFrZ   )rz   r{   r|   )	itertoolsproductsubTestr}   r   
debug_optsrz   r{   r|   s        r   test_chained_device_functionz.TestCudaDebugInfo.test_chained_device_function   s    
 &!(;<
0: 	= 	=,L(H<'/'/  1 1 = = 22<3;3;= = == = = = = = = = = = = = = = =	= 	=   AA	A	c                     t          j        d|d          d             t          j        d|d          fd            t          j        |d          fd            } |d         d	d
           d S )NTFr_   c                     | dz   S rg   r+   r,   s    r   ro   zETestCudaDebugInfo._test_chained_device_function_two_calls.<locals>.f2   rp   r   c                      |  |          z
  S r   r+   rr   s     r   rt   zETestCudaDebugInfo._test_chained_device_function_two_calls.<locals>.f1   ru   r   r7   c                 6     | |            |            d S r   r+   )r-   rs   rt   ro   s     r   rk   zITestCudaDebugInfo._test_chained_device_function_two_calls.<locals>.kernel   s#    Bq!HHHBqEEEEEr   rx   r*   rY   r   r2   ry   s        @@r   '_test_chained_device_function_two_callsz9TestCudaDebugInfo._test_chained_device_function_two_calls   s     
X5	9	9	9	 	 
:	9	 
X5	9	9	9	 	 	 	 
:	9	 
%	0	0	0	 	 	 	 	 
1	0	 	tQr   c                     t          j        dgdz   }|D ]M\  }}}|                     |||          5  |                     |||           d d d            n# 1 swxY w Y   Nd S r   )r   r   r   r   r   s        r   &test_chained_device_function_two_callsz8TestCudaDebugInfo.test_chained_device_function_two_calls   s     &!(;<
0: 	G 	G,L(H<'/'/  1 1 G G <<\=E=EG G GG G G G G G G G G G G G G G G	G 	Gr   c                 t    d } |dd            |dd            |dd            |dd           d S )Nc                 2   t          j        d|d          d             t          j        d          fd            t          j        d          fd            t          j        | d          fd	            } |d
         dd           d S )NTFr_   c                     | | z  S r   r+   r,   s    r   f3z[TestCudaDebugInfo.test_chained_device_three_functions.<locals>.three_device_fns.<locals>.f3   s    1ur   )r`   c                       |           dz   S rg   r+   )r-   r   s    r   ro   z[TestCudaDebugInfo.test_chained_device_three_functions.<locals>.three_device_fns.<locals>.f2   s    r!uuqy r   c                      |  |          z
  S r   r+   rr   s     r   rt   z[TestCudaDebugInfo.test_chained_device_three_functions.<locals>.three_device_fns.<locals>.f1   s    22a55y r   r7   c                       | |           d S r   r+   rw   s     r   rk   z_TestCudaDebugInfo.test_chained_device_three_functions.<locals>.three_device_fns.<locals>.kernel   s    1ar   rx   r*   rY   r   )rz   
leaf_debugrk   rt   ro   r   s      @@@r   three_device_fnszOTestCudaDebugInfo.test_chained_device_three_functions.<locals>.three_device_fns   s    XT???  @? XT"""! ! ! ! #"! XT"""! ! ! ! #"! XLe444    54 F4LAr   T)rz   r   Fr+   )r   r   s     r   #test_chained_device_three_functionsz5TestCudaDebugInfo.test_chained_device_three_functions   sq    	 	 	( 	dt<<<<du====e====e>>>>>>r   )__name__
__module____qualname____doc__r   r   r$   r5   r:   rA   rG   rV   r]   rl   r}   r   r   r   r   __classcell__)r   s   @r   r	   r	      s        > > > > ># # #! ! !> > >= = =B B B   . . .$3 3 3$$ $ $  = = =  $G G G ? ? ? ? ? ? ?r   r	   __main__)numba.tests.supportr   numba.cuda.testingr   numbar   
numba.corer   r   r   r   unittestr	   r   mainr+   r   r   <module>r      s    / / / / / / . . . . . .             + + + + + +     				  9::M? M? M? M? M? M? M? ;:M?` zHMOOOOO r   