U
    hâËdƒ  ã                   @   sX   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rTe ¡  dS )é    N)Úcuda)ÚunittestÚxfail_unless_cudasimÚCUDATestCase)Úconfigc                   @   sH   e Zd Zdd„ Zdd„ Zdd„ Zdd„ Zd	d
„ Zdd„ Ze	dd„ ƒZ
dS )ÚTestExceptionc              	   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 5 Q R X |  d	t	|j
ƒ¡ d S )
Nc                 S   s    t jj}|dkr| j|   d S )Né   )r   Ú	threadIdxÚxÚshape)Zaryr
   © r   ú_/home/sam/Atlas/atlas_env/lib/python3.8/site-packages/numba/cuda/tests/cudapy/test_exception.pyÚ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)Úselfr   Z
unsafe_fooZsafe_fooÚcmr   r   r   Útest_exception	   s    
 zTestException.test_exceptionc              	   C   sH   t jddd�dd„ ƒ}|d dƒ |  t¡� |d dƒ W 5 Q R X d S )NTFr   c                 S   s   | rt ‚d S ©N©Ú
ValueError)Z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 )z¼Testing 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rBt|ƒD ] }||  | | ||  7  < q t  ¡  |dk r|t|ƒD ] }| |  | | ||  7  < qZd S )Né   é   ©r   r	   r
   ZblockDimÚrangeZsyncthreads©r
   ÚyÚtidZntidÚir   r   r   Úproblematic1   s    zBTestException.case_raise_causing_warp_diverge.<locals>.problematicc                 S   s˜   t jj}t jj}|dkrNt|ƒD ],}|| dkr ||  | | ||  7  < q t  ¡  |dk r”t|ƒD ],}|| dkrf| |  | | ||  7  < qfd S )Nr%   r   r&   r'   r)   r   r   r   Úoracle?   s    z=TestException.case_raise_causing_warp_diverge.<locals>.oracleé    g      ð?g{®Gáz„?r   N)r   r   r   ZarangeÚtestingZassert_almost_equal)
r   Úwith_debug_modeZwith_opt_moder-   r.   ÚnZgot_xZgot_yZexpect_xZ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)r1   N)r3   )r   r   r   r   Útest_raise_causing_warp_diverge[   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   s$   |d |d  | d< |d | d< d S ©Nr   r   r   ©Úrr
   r*   r   r   r   Úfg   s    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   r8   r7   r
   r*   r   r   r   Útest_no_zero_division_errorc   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 5 Q R X |  
|d	 d	d
¡ |  
|d d	d¡ d S )NTFr   c                 S   s$   |d |d  | d< |d | d< d S r5   r   r6   r   r   r   r8   y   s    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   r9   r:   r   r   ÚFloatingPointErrorÚZeroDivisionErrorr   r=   )r   r8   r7   r
   r*   Úexcr   r   r   Ú!test_zero_division_error_in_debugu   s    



z/TestException.test_zero_division_error_in_debugc              	      sj   d‰t jdd�‡fdd„ƒ‰ t jdd�‡ fdd„ƒ}|  t¡�}|d	 ƒ  W 5 Q R X |  ˆt|jƒ¡ d S )
NzDevice Function ErrorT)Zdevicec                      s   t ˆ ƒ‚d S r    r!   r   )Úmsgr   r   r8   ™   s    z6TestException.test_raise_in_device_function.<locals>.f)r   c                      s
   ˆ ƒ  d S r    r   r   )r8   r   r   Úkernel�   s    z;TestException.test_raise_in_device_function.<locals>.kernelr#   )r   r   r   r"   r   r   r   )r   rD   Zraisesr   )r8   rC   r   Útest_raise_in_device_function’   s    

z+TestException.test_raise_in_device_functionN)Ú__name__Ú
__module__Ú__qualname__r   r$   r3   r4   r>   rB   r   rE   r   r   r   r   r      s   
3r   Ú__main__)Únumpyr   Znumbar   Znumba.cuda.testingr   r   r   Z
numba.corer   r   rF   Úmainr   r   r   r   Ú<module>   s     