
    xKgL                         d dl Zd dlmZ d dlmZmZmZmZm	Z	 d dl
mZmZ d dl
mZmZmZ  ed       G d de             Zed	k(  r ej$                          yy)
    N)StringIO)cudafloat32float64int32intp)unittestCUDATestCase)skip_on_cudasimskip_with_nvdisasmskip_without_nvdisasmz0Simulator does not generate code to be inspectedc                       e Zd Zed        Zd Zd Zd Z ed      d        Z	 ed      d        Z
 ed      d	        Z ed
      d        Zy)TestInspectc                 R    t        j                         j                  j                  S N)r   current_contextdevicecompute_capability)selfs    h/home/alanp/www/video.onchill/myenv/lib/python3.12/site-packages/numba/cuda/tests/cudapy/test_inspect.pycczTestInspect.cc   s    ##%,,???    c                 
   t         t        f}t        j                  |      d        }t	               }|j                  |       |j                         }| j                  d|       | j                  d|       |j                          |j                  |      }| j                  d|       | j                  d|       | j                  d|       |j                  |      }| j                  d|       | j                  d|       y )Nc                      y r    xys     r   fooz'TestInspect.test_monotyped.<locals>.foo       r   filer   z(float32, int32)cuda.kernel.wrapperdefine linkonce_odr i32z!Generated by NVIDIA NVVM Compiler)r   r   r   jitr   inspect_typesgetvalueassertIncloseinspect_llvminspect_asm)r   sigr   r"   typeannollvmasms          r   test_monotypedzTestInspect.test_monotyped   s    	#	 
	 zt$==?eX&((3

$eT" 	+T2 	/6ooc" 	eS!93?r   c                    t         j                  d        } |d   dd        |d   dd       t               }|j                  |       |j	                         }|j                          | j                  dj                  t              |       | j                  d|       |j                         }| j                  d	t        |             | j                  t        t        f|       | j                  t        t        f|       | j                  d
|t        t        f          | j                  d
|t        t        f          | j                  d|t        t        f          | j                  d|t        t        f          | j                  d|t        t        f          | j                  d|t        t        f          |j                         }| j                  d	t        |             | j                  t        t        f|       | j                  t        t        f|       | j                  d
|t        t        f          | j                  d
|t        t        f          y )Nc                      y r   r   r   s     r   r   z'TestInspect.test_polytyped.<locals>.foo1   r    r   )   r3   r3   g333333?g333333@r!   z
({0}, {0})z(float64, float64)   r   r#   r$   )r   r%   r   r&   r'   r)   r(   formatr   r*   assertEquallenr   r+   )r   r   r"   r-   llvmirsasmdicts         r   test_polytypedzTestInspect.test_polytyped0   s   		 
	 	D	!QD	#szt$==?

l))$/:*H5 ""$CL+tTlG,w('2 	eWT4Z01eWWg%567 	+WT4Z-@A+WWg5E-FG 	/t1DE/'9I1JK//# 	CL+tTlG,w('2 	eWT4Z01eWWg%567r   c                     d}|j                         D ]  }d|v s||v sd} | j                  |       | j                  |d       | j                  d|       | j                  d|       | j                  d|       y )NFz.textTz*//## File ".*/test_inspect.py", line [0-9]S2RBRAEXIT)split
assertTrueassertRegexr(   )r   kernelnamesassseen_functionlines         r   _test_inspect_sasszTestInspect._test_inspect_sass]   st    JJLD$44< $ ! 	&LM 	eT"eT"fd#r   z"nvdisasm needed for inspect_sass()c                     t         d d d   t        d d d   f}t        j                  |d      d        }| j	                  |d|j                  |             y )Nr3   Tlineinfoc                 l    t        j                  d      }|t        |       k  r| |xx   ||   z  cc<   y y Nr3   r   gridr7   r   r   is      r   addz0TestInspect.test_inspect_sass_eager.<locals>.addq   0    		!A3q6z!! r   rQ   )r   r   r   r%   rG   inspect_sassr   r,   rQ   s      r   test_inspect_sass_eagerz#TestInspect.test_inspect_sass_eagerm   sX    ss|U3Q3Z(	#	%	 
&	
 	UC,<,<S,ABr   c                    t        j                  d      d        }t        j                  d      j	                  t        j
                        }t        j                  d      j	                  t        j                        } |d   ||       t
        d d d   t        d d d   f}| j                  |d|j                  |             y )NTrI   c                 l    t        j                  d      }|t        |       k  r| |xx   ||   z  cc<   y y rL   rM   rO   s      r   rQ   z/TestInspect.test_inspect_sass_lazy.<locals>.add{   rR   r   
   )r3   rX   r3   rQ   )	r   r%   nparangeastyper   r   rG   rS   )r   rQ   r   r   	signatures        r   test_inspect_sass_lazyz"TestInspect.test_inspect_sass_lazyy   s    	4	 	 
!	
 IIbM  *IIbM  ,E
1a3Q3Z1.	UC,<,<Y,GHr   z@Missing nvdisasm exception only generated when it is not presentc                    t        j                  t        d d d   f      d        }| j                  t              5 }|j                          d d d        | j                  dt        j                               y # 1 sw Y   /xY w)Nr3   c                     d| d<   y )Nr   r   )r   s    r   fz9TestInspect.test_inspect_sass_nvdisasm_missing.<locals>.f   s    AaDr   znvdisasm has not been found)	r   r%   r   assertRaisesRuntimeErrorrS   r(   str	exception)r   r`   raisess      r   "test_inspect_sass_nvdisasm_missingz.TestInspect.test_inspect_sass_nvdisasm_missing   so     
73Q3</	"	 
#	 |,NN - 	3S9I9I5JK -,s   A;;Bz&nvdisasm needed for inspect_sass_cfg()c                     t         d d d   t        d d d   f}t        j                  |      d        }| j	                  |j                  |      d       y )Nr3   c                 l    t        j                  d      }|t        |       k  r| |xx   ||   z  cc<   y y rL   rM   rO   s      r   rQ   z.TestInspect.test_inspect_sass_cfg.<locals>.add   rR   r   )r\   zdigraph\s*\w\s*{(.|\n)*\n})r   r   r   r%   rA   inspect_sass_cfgrT   s      r   test_inspect_sass_cfgz!TestInspect.test_inspect_sass_cfg   sZ    ss|U3Q3Z(	#	 
	
 	  3 /)	
r   N)__name__
__module____qualname__propertyr   r0   r:   rG   r   rU   r]   r   rf   rj   r   r   r   r   r   
   s    @ @@@+8Z$  ?@	C A	C ?@I AI  & 'L'L CD
 E
r   r   __main__)numpyrY   ior   numbar   r   r   r   r   numba.cuda.testingr	   r
   r   r   r   r   rk   mainr   r   r   <module>ru      s^      5 5 57 7 CDV
, V
 EV
r zHMMO r   