o
    8ήc                     @   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                   @   s`   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dS )TestInspectc                 C   s   t  jjS N)r   current_contextdevicecompute_capability)self r   K/tmp/pip-target-vg8gfxp4/lib/python/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 S )Nc                 S      d S r   r   xyr   r   r   foo      z'TestInspect.test_monotyped.<locals>.foofiler   z(float32, int32)z!Generated by NVIDIA NVVM Compiler)r   r   r   jitr   inspect_typesgetvalueassertIncloseinspect_llvminspect_asm)r   sigr   r   typeanno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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   *   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"   formatr   r$   assertEquallenr   r%   )r   r   r   r'   llvmirsasmdictr   r   r   test_polytyped)   s,   
zTestInspect.test_polytypedc                 C   sX   d}|  D ]}d|v r||v rd}q| | | d| | d| | d| d S )NFz.textTS2RBRAEXIT)split
assertTruer"   )r   kernelnamesassseen_functionliner   r   r   _test_inspect_sassN   s   
zTestInspect._test_inspect_sassz"nvdisasm needed for inspect_sass()c                 C   sF   t d d d td d d f}t|dd }| |d|| d S )Nr*   c                 S   2   t d}|t| k r| |  || 7  < d S d S Nr*   r   gridr.   r   r   ir   r   r   add`      
z0TestInspect.test_inspect_sass_eager.<locals>.addrC   )r   r   r   r   r<   inspect_sass)r   r&   rC   r   r   r   test_inspect_sass_eager\   s   
z#TestInspect.test_inspect_sass_eagerc                 C   st   t jdd }tdtj}tdtj}|d || td d d td d d f}| |d|| d S )Nc                 S   r=   r>   r?   rA   r   r   r   rC   j   rD   z/TestInspect.test_inspect_sass_lazy.<locals>.add
   )r*   rG   r*   rC   )	r   r   nparangeastyper   r   r<   rE   )r   rC   r   r   	signaturer   r   r   test_inspect_sass_lazyh   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   fz   r   z9TestInspect.test_inspect_sass_nvdisasm_missing.<locals>.fznvdisasm is required)	r   r   r   assertRaisesRuntimeErrorrE   r"   str	exception)r   rM   raisesr   r   r   "test_inspect_sass_nvdisasm_missingw   s   

z.TestInspect.test_inspect_sass_nvdisasm_missingN)__name__
__module____qualname__propertyr   r)   r1   r<   r   rF   rL   r   rS   r   r   r   r   r   
   s    
%

r   __main__)numpyrH   ior   numbar   r   r   r   r   numba.cuda.testingr   r	   r
   r   r   r   rT   mainr   r   r   r   <module>   s    y