
    sh                         S SK Jr  S SKJr  S SKJr  S SKJr  S SKJr  S SK	r	S SK
r
S SKr\" S5       " S S	\5      5       r\S
:X  a  \R                  " 5         gg)    )override_config)skip_on_cudasim)cuda)types)CUDATestCaseNz&Simulator does not produce debug dumpsc                      ^  \ rS rSrSrU 4S jrS rS rS rS r	S r
S	 rS
 rS rS rS rS rS rS rS rSrU =r$ )TestCudaDebugInfo   z@
These tests only checks the compiled PTX for debuginfo section
c                 D   > [         TU ]  5         U R                  S5        g )Nz!Exceptions not supported with LTO)supersetUpskip_if_lto)self	__class__s    z/Users/tiagomarins/Projetos/claudeai/copy_bank/venv/lib/python3.13/site-packages/numba/cuda/tests/cudapy/test_debuginfo.pyr   TestCudaDebugInfo.setUp   s     	<=    c                 F    UR                  U5        UR                  U5      $ N)compileinspect_asm)r   fnsigs      r   _getasmTestCudaDebugInfo._getasm   s    


3~~c""r   c                     U R                  XS9n[        R                  " S5      nUR                  U5      nU(       a  U R                  OU R
                  nU" XdS9  g )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   _checkTestCudaDebugInfo._check   sO    ll2l'ZZ(GH"))#.+14''t7H7H r   c                 ~    [         R                  " SS9S 5       nU R                  U[        R                  S S  4SS9  g )NFdebugc                     SU S'   g N   r    xs    r   foo7TestCudaDebugInfo.test_no_debuginfo_in_asm.<locals>.foo&       AaDr   r   r"   r   jitr'   r   int32r   r2   s     r   test_no_debuginfo_in_asm*TestCudaDebugInfo.test_no_debuginfo_in_asm%   s<    			 
	 	Cekk!n.u=r   c                     [         R                  " SSS9S 5       nU R                  U[        R                  S S  4SS9  g )NTFr+   optc                     SU S'   g r-   r/   r0   s    r   r2   4TestCudaDebugInfo.test_debuginfo_in_asm.<locals>.foo-   r4   r   r5   r6   r9   s     r   test_debuginfo_in_asm'TestCudaDebugInfo.test_debuginfo_in_asm,   s>    	%	(	 
)	 	Cekk!n.t<r   c                 D   [        SS5         [        R                  " SS9S 5       nU R                  U[        R
                  S S  4SS9  [        R                  " SS9S	 5       nU R                  U[        R
                  S S  4SS9  S S S 5        g ! , (       d  f       g = f)
NCUDA_DEBUGINFO_DEFAULTr.   F)r>   c                     SU S'   g r-   r/   r0   s    r   r2   8TestCudaDebugInfo.test_environment_override.<locals>.foo6       !r   Tr5   r*   c                     SU S'   g r-   r/   r0   s    r   bar8TestCudaDebugInfo.test_environment_override.<locals>.bar=   rG   r   )r   r   r7   r'   r   r8   )r   r2   rI   s      r   test_environment_override+TestCudaDebugInfo.test_environment_override3   s    5q9XX%  ! KK%++a.!24K@ XXE" # KK%++a.!25KA :99s   A;B
Bc                 f    [         R                  " [        R                  S S S2   4SSS9S 5       ng )Nr.   TFr=   c                     SU S'   g )Nr   r/   r0   s    r   f,TestCudaDebugInfo.test_issue_5835.<locals>.fG   r4   r   r   r7   r   r8   r   rO   s     r   test_issue_5835!TestCudaDebugInfo.test_issue_5835C   s3     
5;;ss#%Tu	=	 
>	r   c                 J   [         R                  S S S2   4n[        R                  " USSS9S 5       nUR	                  U5      nUR                  5        Vs/ s H  nSU;   d  M  UPM     nnU R                  [        U5      S5        US   nU R                  SU5        g s  snf )Nr.   Tr   r=   c                     SU S'   g r-   r/   r0   s    r   rO   7TestCudaDebugInfo.test_wrapper_has_debuginfo.<locals>.fN   r4   r   zdefine void @"_ZN6cudapyz!dbg)	r   r8   r   r7   inspect_llvm
splitlinesassertEquallenassertIn)r   r   rO   llvm_irlinedefineswrapper_defines          r   test_wrapper_has_debuginfo,TestCudaDebugInfo.test_wrapper_has_debuginfoK   s    {{3Q3!	#Tq	)	 
*	 ..%$+$6$6$8 :$8D0D8 $8 : 	Wq) fn-:s   
B %B c                     [         R                  " [        R                  S S  [        R                  S S  4SSS9S 5       ng )NTFr=   c                 $    U S   S;   a  SOSUS'   g )Nr   )      r.   rf   r/   )inpoutps     r   rO   DTestCudaDebugInfo.test_debug_function_calls_internal_impl.<locals>.fk   s    q6V+aDGr   rQ   rR   s     r   'test_debug_function_calls_internal_impl9TestCudaDebugInfo.test_debug_function_calls_internal_impl]   s9     
5;;q>5;;q>2$E	J	3 
K	3r   c                    ^ [         R                  " SSSS9S 5       m[         R                  " [        R                  S S  4SSS9U4S j5       ng )NTr   devicer+   r>   c                      [         R                  R                  [         R                  R                  -  [         R                  R                  -   $ r   )r   blockDimr1   blockIdx	threadIdxr/   r   r   threadidMTestCudaDebugInfo.test_debug_function_calls_device_function.<locals>.threadidt   s,    ==??T]]__4t~~7G7GGGr   r=   c                 d   > [         R                  " S5      nU[        U 5      :  a
  T" 5       X'   g g Nr.   )r   gridr[   )arrirs   s     r   kernelKTestCudaDebugInfo.test_debug_function_calls_device_function.<locals>.kernelx   s)    		!A3s8|! r   rQ   )r   rz   rs   s     @r   )test_debug_function_calls_device_function;TestCudaDebugInfo.test_debug_function_calls_device_functiono   sR    
 
Tq	1	H 
2	H 
5;;q>#4Q	7	$ 
8	$r   c                   ^^ [         R                  " SUSS9S 5       m[         R                  " SUSS9U4S j5       m[         R                  " [        R                  [        R                  4USS9U4S j5       nUS   " S	S
5        g )NTFrm   c                     U S-   $ rv   r/   r0   s    r   f2;TestCudaDebugInfo._test_chained_device_function.<locals>.f2       q5Lr   c                    > U T" U5      -
  $ r   r/   r1   yr   s     r   f1;TestCudaDebugInfo._test_chained_device_function.<locals>.f1       r!u9r   r=   c                    > T" X5        g r   r/   r1   r   r   s     r   rz   ?TestCudaDebugInfo._test_chained_device_function.<locals>.kernel   s
    qHr   r.   r.   r.   re   rQ   r   kernel_debugf1_debugf2_debugrz   r   r   s        @@r   _test_chained_device_function/TestCudaDebugInfo._test_chained_device_function~   s    	X5	9	 
:	 
X5	9	 
:	 
5;;,Le	L	 
M	 	tQr   c                     [         R                  " S/S-  6 nU H3  u  p#nU R                  UUUS9   U R                  UUU5        S S S 5        M5     g ! , (       d  f       MG  = fN)TFrf   )r   r   r   )	itertoolsproductsubTestr   r   
debug_optsr   r   r   s        r   test_chained_device_function.TestCudaDebugInfo.test_chained_device_function   sp    
 &&!(;<
0:,LH<'/'/  1 22<3;3;=1 1 1;1 1   A
A!	c                    ^^ [         R                  " SUSS9S 5       m[         R                  " SUSS9U4S j5       m[         R                  " USS9UU4S j5       nUS   " S	S
5        g )NTFrm   c                     U S-   $ rv   r/   r0   s    r   r   ETestCudaDebugInfo._test_chained_device_function_two_calls.<locals>.f2   r   r   c                    > U T" U5      -
  $ r   r/   r   s     r   r   ETestCudaDebugInfo._test_chained_device_function_two_calls.<locals>.f1   r   r   r=   c                 &   > T" X5        T" U 5        g r   r/   )r1   r   r   r   s     r   rz   ITestCudaDebugInfo._test_chained_device_function_two_calls.<locals>.kernel   s    qHqEr   r   r.   re   r   r7   r   s        @@r   '_test_chained_device_function_two_calls9TestCudaDebugInfo._test_chained_device_function_two_calls   st     
X5	9	 
:	 
X5	9	 
:	 
%	0	 
1	 	tQr   c                     [         R                  " S/S-  6 nU H3  u  p#nU R                  UUUS9   U R                  UUU5        S S S 5        M5     g ! , (       d  f       MG  = fr   )r   r   r   r   r   s        r   &test_chained_device_function_two_calls8TestCudaDebugInfo.test_chained_device_function_two_calls   sq     &&!(;<
0:,LH<'/'/  1 <<\=E=EG1 1 1;1 1r   c                 B    S nU" SSS9  U" SSS9  U" SSS9  U" SSS9  g )Nc                 
  ^^^ [         R                  " SUSS9S 5       m[         R                  " SS9U4S j5       m[         R                  " SS9U4S j5       m[         R                  " U SS9U4S	 j5       nUS
   " SS5        g )NTFrm   c                 
    X -  $ r   r/   r0   s    r   f3[TestCudaDebugInfo.test_chained_device_three_functions.<locals>.three_device_fns.<locals>.f3   s	    ur   )rn   c                    > T" U 5      S-   $ rv   r/   )r1   r   s    r   r   [TestCudaDebugInfo.test_chained_device_three_functions.<locals>.three_device_fns.<locals>.f2   s    !uqy r   c                    > U T" U5      -
  $ r   r/   r   s     r   r   [TestCudaDebugInfo.test_chained_device_three_functions.<locals>.three_device_fns.<locals>.f1   s    2a5y r   r=   c                    > T" X5        g r   r/   r   s     r   rz   _TestCudaDebugInfo.test_chained_device_three_functions.<locals>.three_device_fns.<locals>.kernel   s
    1r   r   r.   re   r   )r   
leaf_debugrz   r   r   r   s      @@@r   three_device_fnsOTestCudaDebugInfo.test_chained_device_three_functions.<locals>.three_device_fns   s    XXT? @ XXT"! #! XXT"! #! XXLe4 5 4LAr   T)r   r   Fr/   )r   r   s     r   #test_chained_device_three_functions5TestCudaDebugInfo.test_chained_device_three_functions   s1    	( 	dt<du=e=e>r   r/   )__name__
__module____qualname____firstlineno____doc__r   r   r'   r:   rA   rK   rS   ra   rj   r|   r   r   r   r   r   __static_attributes____classcell__)r   s   @r   r	   r	      sZ    >#!>=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      sZ    / .   +  	  9:M? M? ;M?` zMMO r   