o
    Y۷iL                     @   s   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 deZed	kr@e  dS dS )
    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                   @   sp   e Zd Zedd Zdd Zdd Zdd Zed	d
d Z	ed	dd Z
eddd Zeddd ZdS )TestInspectc                 C   s   t  jjS N)r   current_contextdevicecompute_capability)self r   Z/home/ubuntu/vllm_env/lib/python3.10/site-packages/numba/cuda/tests/cudapy/test_inspect.pycc      zTestInspect.ccc                 C   s   t tf}t|dd }t }|j|d | }| d| | d| |  |	|}| d| | d| | d| |
|}| d| | d| d S )	Nc                 S      d S r   r   xyr   r   r   foo      z'TestInspect.test_monotyped.<locals>.foo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asmr   r   r   test_monotyped   s    


zTestInspect.test_monotypedc                 C   s~  t jdd }|d dd |d dd t }|j|d | }|  | dt| | d	| |	 }| 
d
t| | ttf| | ttf| | d|ttf  | d|ttf  | d|ttf  | d|ttf  | d|ttf  | d|ttf  | }| 
d
t| | ttf| | ttf| | d|ttf  | d|ttf  d S )Nc                 S   r   r   r   r   r   r   r   r   1   r   z'TestInspect.test_polytyped.<locals>.foo)   r-   r-   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asmdictr   r   r   test_polytyped0   s4   
zTestInspect.test_polytypedc                 C   sd   d}|  D ]}d|v r||v rd}q| | | |d | d| | d| | d| d S )NFz.textTz*//## File ".*/test_inspect.py", line [0-9]S2RBRAEXIT)split
assertTrueassertRegexr$   )r   kernelnamesassseen_functionliner   r   r   _test_inspect_sass]   s   
zTestInspect._test_inspect_sassz"nvdisasm needed for inspect_sass()c                 C   sJ   t d d d td d d f}tj|dddd }| |d|| d S )Nr-   Tlineinfoc                 S   2   t d}|t| k r| |  || 7  < d S d S Nr-   r   gridr1   r   r   ir   r   r   addq      
z0TestInspect.test_inspect_sass_eager.<locals>.addrI   )r   r   r   r!   r@   inspect_sassr   r(   rI   r   r   r   test_inspect_sass_eagerm   s   
z#TestInspect.test_inspect_sass_eagerc                 C   sz   t jdddd }tdtj}tdtj}|d || td d d td d d f}| |d|| d S )	NTrA   c                 S   rC   rD   rE   rG   r   r   r   rI   {   rJ   z/TestInspect.test_inspect_sass_lazy.<locals>.add
   )r-   rN   r-   rI   )	r   r!   nparangeastyper   r   r@   rK   )r   rI   r   r   	signaturer   r   r   test_inspect_sass_lazyy   s   

z"TestInspect.test_inspect_sass_lazyz@Missing nvdisasm exception only generated when it is not presentc                 C   sf   t td d d fdd }| t}|  W d    n1 s#w   Y  | dt|j d S )Nr-   c                 S   s   d| d< d S )Nr   r   )r   r   r   r   f   r   z9TestInspect.test_inspect_sass_nvdisasm_missing.<locals>.fznvdisasm has not been found)	r   r!   r   assertRaisesRuntimeErrorrK   r$   str	exception)r   rT   raisesr   r   r   "test_inspect_sass_nvdisasm_missing   s   

z.TestInspect.test_inspect_sass_nvdisasm_missingz&nvdisasm needed for inspect_sass_cfg()c                 C   sF   t d d d td d d f}t|dd }| |j|dd d S )Nr-   c                 S   rC   rD   rE   rG   r   r   r   rI      rJ   z.TestInspect.test_inspect_sass_cfg.<locals>.add)rR   zdigraph\s*\w\s*{(.|\n)*\n})r   r   r   r!   r:   inspect_sass_cfgrL   r   r   r   test_inspect_sass_cfg   s   

z!TestInspect.test_inspect_sass_cfgN)__name__
__module____qualname__propertyr   r,   r4   r@   r   rM   rS   r   rZ   r\   r   r   r   r   r   
   s    
 -



r   __main__)numpyrO   ior   numbar   r   r   r   r   numba.cuda.testingr   r	   r
   r   r   r   r]   mainr   r   r   r   <module>   s     