
    ,i%                         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                   f    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 )TestCudaDebugInfozH
    These tests only checks the compiled PTX for debuginfo section
    c                 V    |                     |           |                    |          S N)compileinspect_asm)selffnsigs      d/var/www/html/speakWrite/venv/lib/python3.11/site-packages/numba/cuda/tests/cudapy/test_debuginfo.py_getasmzTestCudaDebugInfo._getasm   s#    


3~~c"""    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   r+   r,   r/   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)r3   c                     d| d<   d S r$   r&   r'   s    r   r)   z8TestCudaDebugInfo.test_environment_override.<locals>.foo,       !r   Tr+   r!   c                     d| d<   d S r$   r&   r'   s    r   barz8TestCudaDebugInfo.test_environment_override.<locals>.bar3   r9   r   )r   r   r-   r   r   r.   )r   r)   r;   s      r   test_environment_overridez+TestCudaDebugInfo.test_environment_override)   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%   TFr2   c                     d| d<   d S )Nr   r&   r'   s    r   fz,TestCudaDebugInfo.test_issue_5835.<locals>.f=   r*   r   r   r-   r   r.   r   r?   s     r   test_issue_5835z!TestCudaDebugInfo.test_issue_58359   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   r2   c                     d| d<   d S r$   r&   r'   s    r   r?   z7TestCudaDebugInfo.test_wrapper_has_debuginfo.<locals>.fD   r*   r   c                     g | ]}d |v |	S )zdefine void @"_ZN6cudapyr&   ).0lines     r   
<listcomp>z@TestCudaDebugInfo.test_wrapper_has_debuginfo.<locals>.<listcomp>J   s,     : : :D0D88 888r   z!dbg)	r   r.   r   r-   inspect_llvm
splitlinesassertEquallenassertIn)r   r   r?   llvm_irdefineswrapper_defines         r   test_wrapper_has_debuginfoz,TestCudaDebugInfo.test_wrapper_has_debuginfoA   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 )NTFr2   c                 (    | d         dv rdnd|d<   d S )Nr   )      r%   rU   r&   )inpoutps     r   r?   zDTestCudaDebugInfo.test_debug_function_calls_internal_impl.<locals>.fa   s!    q6V++aaDGGGr   r@   rA   s     r   'test_debug_function_calls_internal_implz9TestCudaDebugInfo.test_debug_function_calls_internal_implS   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"   r3   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>.threadidj   s    =?T]_4t~7GGGr   r2   c                 t    t          j        d          }|t          |           k     r             | |<   d S d S Nr%   )r   gridrL   )arrir`   s     r   kernelzKTestCudaDebugInfo.test_debug_function_calls_device_function.<locals>.kerneln   s9    	!A3s88||!A |r   r@   )r   rf   r`   s     @r   )test_debug_function_calls_device_functionz;TestCudaDebugInfo.test_debug_function_calls_device_functione   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 )NTFrZ   c                     | dz   S rb   r&   r'   s    r   f2z;TestCudaDebugInfo._test_chained_device_function.<locals>.f2u       q5Lr   c                      |  |          z
  S r   r&   r(   yrj   s     r   f1z;TestCudaDebugInfo._test_chained_device_function.<locals>.f1y       rr!uu9r   r2   c                       | |           d S r   r&   r(   rn   ro   s     r   rf   z?TestCudaDebugInfo._test_chained_device_function.<locals>.kernel}   s    Bq!HHHHHr   r%   r%   r%   rT   r@   r   kernel_debugf1_debugf2_debugrf   ro   rj   s        @@r   _test_chained_device_functionz/TestCudaDebugInfo._test_chained_device_functiont   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)TFrU   )ru   rv   rw   )	itertoolsproductsubTestrx   r   
debug_optsru   rv   rw   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 )NTFrZ   c                     | dz   S rb   r&   r'   s    r   rj   zETestCudaDebugInfo._test_chained_device_function_two_calls.<locals>.f2   rk   r   c                      |  |          z
  S r   r&   rm   s     r   ro   zETestCudaDebugInfo._test_chained_device_function_two_calls.<locals>.f1   rp   r   r2   c                 6     | |            |            d S r   r&   )r(   rn   ro   rj   s     r   rf   zITestCudaDebugInfo._test_chained_device_function_two_calls.<locals>.kernel   s#    Bq!HHHBqEEEEEr   rs   r%   rT   r   r-   rt   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 rz   )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 )NTFrZ   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 rb   r&   )r(   r   s    r   rj   z[TestCudaDebugInfo.test_chained_device_three_functions.<locals>.three_device_fns.<locals>.f2   s    r!uuqy r   c                      |  |          z
  S r   r&   rm   s     r   ro   z[TestCudaDebugInfo.test_chained_device_three_functions.<locals>.three_device_fns.<locals>.f1   s    22a55y r   r2   c                       | |           d S r   r&   rr   s     r   rf   z_TestCudaDebugInfo.test_chained_device_three_functions.<locals>.three_device_fns.<locals>.kernel   s    1ar   rs   r%   rT   r   )ru   
leaf_debugrf   ro   rj   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)ru   r   Fr&   )r   r   s     r   #test_chained_device_three_functionsz5TestCudaDebugInfo.test_chained_device_three_functions   sq    	 	 	( 	dt<<<<du====e====e>>>>>>r   N)__name__
__module____qualname____doc__r   r   r0   r5   r<   rB   rQ   rX   rg   rx   r   r   r   r   r&   r   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::C? C? C? C? C? C? C? ;:C?L zHMOOOOO r   