o
    Y۷i}                     @   s\   d dl Zd dlmZ d dlmZmZmZ d dlm	Z	 G dd deZ
edkr,e  dS dS )    N)cuda)unittestxfail_unless_cudasimCUDATestCase)configc                       sX   e Zd Z fddZdd Zdd Zdd Zd	d
 Zdd Zdd Z	e
dd Z  ZS )TestExceptionc                    s   t    | d d S )Nz!Exceptions not supported with LTO)supersetUpskip_if_ltoself	__class__ \/home/ubuntu/vllm_env/lib/python3.10/site-packages/numba/cuda/tests/cudapy/test_exception.pyr	   	   s   
zTestException.setUpc                 C   s   dd }t |}t jddd|}tjs |d tddg | t}|d tddg W d    n1 s;w   Y  | d	t	|j
 d S )
Nc                 S   s$   t jj}|dkr| j|   d S d S )N   )r   	threadIdxxshape)aryr   r   r   r   foo   s   z)TestException.test_exception.<locals>.fooTFdebugopt)      r   r   ztuple index out of range)r   jitr   ENABLE_CUDASIMnparrayassertRaises
IndexErrorassertInstr	exception)r   r   
unsafe_foosafe_foocmr   r   r   test_exception   s   
zTestException.test_exceptionc                 C   s^   t jddddd }|d d | t |d d W d    d S 1 s(w   Y  d S )NTFr   c                 S   s   | rt d S N
ValueError)do_raiser   r   r   r   %   s   z*TestException.test_user_raise.<locals>.foor   r   )r   r   r    r+   )r   r   r   r   r   test_user_raise$   s   
"zTestException.test_user_raisec           
      C   s   | }t j||ddd }t jdd }d}dt|d  }dt|d  }|d	|f || dt|d  }dt|d  }	|d	|f ||	 tj|| tj|	| d
S )zTesting issue #2655.

        Exception raising code can cause the compiler to miss location
        of unifying branch target and resulting in unexpected warp
        divergence.
        r   c                 S   s   t jj}t jj}|dkr!t|D ]}||  | | ||  7  < qt   |dk r>t|D ]}| |  | | ||  7  < q-d S d S )N      r   r   r   blockDimrangesyncthreadsr   ytidntidir   r   r   problematic7   s   zBTestException.case_raise_causing_warp_diverge.<locals>.problematicc                 S   s   t jj}t jj}|dkr't|D ]}|| dkr&||  | | ||  7  < qt   |dk rJt|D ]}|| dkrI| |  | | ||  7  < q3d S d S )Nr/   r   r0   r1   r5   r   r   r   oracleE   s   z=TestException.case_raise_causing_warp_diverge.<locals>.oracle    g      ?g{Gz?r   N)r   r   r   arangetestingassert_almost_equal)
r   with_debug_modewith_opt_moder:   r;   ngot_xgot_yexpect_xexpect_yr   r   r   case_raise_causing_warp_diverge.   s   

z-TestException.case_raise_causing_warp_divergec                 C   s   | j dd dS )z#Test case for issue #2655.
        F)r@   N)rG   r   r   r   r   test_raise_causing_warp_divergea   s   z-TestException.test_raise_causing_warp_divergec                 C   sl   t jdd }td}td}td}|d ||| | t|d d | |d |d d d S )	Nc                 S   $   |d |d  | d< |d | d< d S Nr   r   r   rr   r6   r   r   r   fm      z4TestException.test_no_zero_division_error.<locals>.fr   r   r-   r   zExpected inf from div by zerozExpected execution to continue)r   r   r   zerosones
assertTrueisinfassertEqual)r   rM   rL   r   r6   r   r   r   test_no_zero_division_errori   s   



z)TestException.test_no_zero_division_errorc                 C   s   t jddddd }td}td}td}tjr t}nt}| 	| |d ||| W d    n1 s:w   Y  | 
|d	 d	d
 | 
|d d	d d S )NTFr   c                 S   rI   rJ   r   rK   r   r   r   rM      rN   z:TestException.test_zero_division_error_in_debug.<locals>.fr   r   r-   r   z Expected result to be left unsetzExpected execution to stop)r   r   r   rO   rP   r   r   FloatingPointErrorZeroDivisionErrorr    rS   )r   rM   rL   r   r6   excr   r   r   !test_zero_division_error_in_debug{   s   



z/TestException.test_zero_division_error_in_debugc                    s~   dt jddfdd t jdd fdd}| t}|d	   W d    n1 s/w   Y  | t|j d S )
NzDevice Function ErrorT)devicec                      s   t  r)   r*   r   )msgr   r   rM      s   z6TestException.test_raise_in_device_function.<locals>.f)r   c                      s
      d S r)   r   r   )rM   r   r   kernel   s   
z;TestException.test_raise_in_device_function.<locals>.kernelr-   )r   r   r    r+   r"   r#   r$   )r   r[   raisesr   )rM   rZ   r   test_raise_in_device_function   s   

z+TestException.test_raise_in_device_function)__name__
__module____qualname__r	   r(   r.   rG   rH   rT   rX   r   r]   __classcell__r   r   r   r   r      s    
3r   __main__)numpyr   numbar   numba.cuda.testingr   r   r   
numba.corer   r   r^   mainr   r   r   r   <module>   s     &