
    xKg                         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                          yy)    )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                 D    t         |           | j                  d       y )Nz!Exceptions not supported with LTO)supersetUpskip_if_lto)self	__class__s    j/home/alanp/www/video.onchill/myenv/lib/python3.12/site-packages/numba/cuda/tests/cudapy/test_debuginfo.pyr   zTestCudaDebugInfo.setUp   s     	<=    c                 F    |j                  |       |j                  |      S N)compileinspect_asm)r   fnsigs      r   _getasmzTestCudaDebugInfo._getasm   s    


3~~c""r   c                     | j                  ||      }t        j                  d      }|j                  |      }|r| j                  n| j
                  } |||       y )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   sS    ll23l'ZZ(GH"))#.+14''t7H7HC r   c                     t        j                  d      d        }| j                  |t        j                  d d  fd       y )NFdebugc                     d| d<   y N   r    xs    r   fooz7TestCudaDebugInfo.test_no_debuginfo_in_asm.<locals>.foo&       AaDr   r   r   r   jitr$   r   int32r   r.   s     r   test_no_debuginfo_in_asmz*TestCudaDebugInfo.test_no_debuginfo_in_asm%   s<    			 
	 	Cekk!n.u=r   c                     t        j                  dd      d        }| j                  |t        j                  d d  fd       y )NTFr'   optc                     d| d<   y 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,   s>    	%	(	 
)	 	Cekk!n.t<r   c                 N   t        dd      5  t        j                  d      d        }| j                  |t        j
                  d d  fd       t        j                  d      d	        }| j                  |t        j
                  d d  fd       d d d        y # 1 sw Y   y xY w)
NCUDA_DEBUGINFO_DEFAULTr*   F)r8   c                     d| d<   y r)   r+   r,   s    r   r.   z8TestCudaDebugInfo.test_environment_override.<locals>.foo6       !r   Tr0   r&   c                     d| d<   y 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XX%  ! KK%++a.!24K@ XXE" # KK%++a.!25KA :99s   BBB$c                 j    t        j                  t        j                  d d d   fdd      d        }y )Nr*   TFr7   c                     d| d<   y )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   s3     
5;;ss#%Tu	=	 
>	r   c                 @   t         j                  d d d   f}t        j                  |dd      d        }|j	                  |      }|j                         D cg c]  }d|v r|
 }}| j                  t        |      d       |d   }| j                  d|       y c c}w )Nr*   Tr   r7   c                     d| d<   y r)   r+   r,   s    r   rD   z7TestCudaDebugInfo.test_wrapper_has_debuginfo.<locals>.fN   r/   r   zdefine void @"_ZN6cudapyz!dbg)	r   r3   r   r2   inspect_llvm
splitlinesassertEquallenassertIn)r   r   rD   llvm_irlinedefineswrapper_defines          r   test_wrapper_has_debuginfoz,TestCudaDebugInfo.test_wrapper_has_debuginfoK   s    {{3Q3!	#Tq	)	 
*	 ..%$+$6$6$8 :$8D0D8 $8 : 	Wq) fn-:s   Bc                     t        j                  t        j                  d d  t        j                  d d  fdd      d        }y )NTFr7   c                 (    | d   dv rd|d<   y d|d<   y )Nr   )      r*   rW   r+   )inpoutps     r   rD   zDTestCudaDebugInfo.test_debug_function_calls_internal_impl.<locals>.fk   s    q6V+aDGDGr   rE   rF   s     r   'test_debug_function_calls_internal_implz9TestCudaDebugInfo.test_debug_function_calls_internal_impl]   s9     
5;;q>5;;q>2$E	J	3 
K	3r   c                     t        j                  ddd      d        t        j                  t        j                  d d  fdd      fd       }y )NTr   devicer'   r8   c                      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~~7G7GGGr   r7   c                 b    t        j                  d      }|t        |       k  r        | |<   y y Nr*   )r   gridrM   )arrirb   s     r   kernelzKTestCudaDebugInfo.test_debug_function_calls_device_function.<locals>.kernelx   s+    		!A3s8|!A r   rE   )r   rh   rb   s     @r   )test_debug_function_calls_device_functionz;TestCudaDebugInfo.test_debug_function_calls_device_functiono   sR    
 
Tq	1	H 
2	H 
5;;q>#4Q	7	$ 
8	$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
       y )NTFr\   c                     | dz   S rd   r+   r,   s    r   f2z;TestCudaDebugInfo._test_chained_device_function.<locals>.f2       q5Lr   c                     |  |      z
  S r   r+   r-   yrl   s     r   f1z;TestCudaDebugInfo._test_chained_device_function.<locals>.f1       r!u9r   r7   c                      | |       y r   r+   r-   rp   rq   s     r   rh   z?TestCudaDebugInfo._test_chained_device_function.<locals>.kernel   s    q!Hr   r*   r*   r*   rV   rE   r   kernel_debugf1_debugf2_debugrh   rq   rl   s        @@r   _test_chained_device_functionz/TestCudaDebugInfo._test_chained_device_function~   s    	X5	9	 
:	 
X5	9	 
:	 
5;;,Le	L	 
M	 	tQr   c                     t        j                  dgdz   }|D ]6  \  }}}| j                  |||      5  | j                  |||       d d d        8 y # 1 sw Y   CxY wN)TFrW   )rw   rx   ry   )	itertoolsproductsubTestrz   r   
debug_optsrw   rx   ry   s        r   test_chained_device_functionz.TestCudaDebugInfo.test_chained_device_function   sr    
 &&!(;<
0:,L(H<'/'/  1 22<3;3;=1 1 1;1 1   AA	c                     t        j                  d|d      d        t        j                  d|d      fd       t        j                  |d      fd       } |d   d	d
       y )NTFr\   c                     | dz   S rd   r+   r,   s    r   rl   zETestCudaDebugInfo._test_chained_device_function_two_calls.<locals>.f2   rm   r   c                     |  |      z
  S r   r+   ro   s     r   rq   zETestCudaDebugInfo._test_chained_device_function_two_calls.<locals>.f1   rr   r   r7   c                 (     | |        |        y r   r+   )r-   rp   rq   rl   s     r   rh   zITestCudaDebugInfo._test_chained_device_function_two_calls.<locals>.kernel   s    q!HqEr   ru   r*   rV   r   r2   rv   s        @@r   '_test_chained_device_function_two_callsz9TestCudaDebugInfo._test_chained_device_function_two_calls   sv     
X5	9	 
:	 
X5	9	 
:	 
%	0	 
1	 	tQr   c                     t        j                  dgdz   }|D ]6  \  }}}| j                  |||      5  | j                  |||       d d d        8 y # 1 sw Y   CxY wr|   )r}   r~   r   r   r   s        r   &test_chained_device_function_two_callsz8TestCudaDebugInfo.test_chained_device_function_two_calls   ss     &&!(;<
0:,L(H<'/'/  1 <<\=E=EG1 1 1;1 1r   c                 Z    d } |dd        |dd        |dd        |dd       y )Nc                    t        j                  d|d      d        t        j                  d      fd       t        j                  d      fd       t        j                  | d      fd	       } |d
   dd       y )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 rd   r+   )r-   r   s    r   rl   z[TestCudaDebugInfo.test_chained_device_three_functions.<locals>.three_device_fns.<locals>.f2   s    !uqy r   c                     |  |      z
  S r   r+   ro   s     r   rq   z[TestCudaDebugInfo.test_chained_device_three_functions.<locals>.three_device_fns.<locals>.f1   s    2a5y r   r7   c                      | |       y r   r+   rt   s     r   rh   z_TestCudaDebugInfo.test_chained_device_three_functions.<locals>.three_device_fns.<locals>.kernel   s    1ar   ru   r*   rV   r   )rw   
leaf_debugrh   rq   rl   r   s      @@@r   three_device_fnszOTestCudaDebugInfo.test_chained_device_three_functions.<locals>.three_device_fns   s    XXT? @ XXT"! #! XXT"! #! XXLe4 5 F4LAr   T)rw   r   Fr+   )r   r   s     r   #test_chained_device_three_functionsz5TestCudaDebugInfo.test_chained_device_three_functions   s1    	( 	dt<du=e=e>r   )__name__
__module____qualname____doc__r   r   r$   r5   r:   rA   rG   rS   rZ   ri   rz   r   r   r   r   __classcell__)r   s   @r   r	   r	      sU    >#!>=B .$3$$=$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?` zHMMO r   