
    xKg}                     |    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e      Z
edk(  r ej                          yy)    N)cuda)unittestxfail_unless_cudasimCUDATestCase)configc                   R     e Zd Z fdZd Zd Zd Zd Zd Zd Z	e
d        Z xZS )	TestExceptionc                 D    t         |           | j                  d       y )Nz!Exceptions not supported with LTO)supersetUpskip_if_lto)self	__class__s    j/home/alanp/www/video.onchill/myenv/lib/python3.12/site-packages/numba/cuda/tests/cudapy/test_exception.pyr   zTestException.setUp	   s     	<=    c                    d }t        j                  |      } t        j                  dd      |      }t        j                  s  |d   t	        j
                  ddg             | j                  t              5 } |d   t	        j
                  ddg             d d d        | j                  dt        j                               y # 1 sw Y   /xY w)	Nc                 d    t         j                  j                  }|dk(  r| j                  |     y y )N   )r   	threadIdxxshape)aryr   s     r   fooz)TestException.test_exception.<locals>.foo   s,      AAv 		1"	 r   TF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cms        r   test_exceptionzTestException.test_exception   s    	 XXc]
2488$E237$$ JtRXXq!f-.z*bHTN288QF+, +0#bll2CD +*s   ;!C

Cc                     t        j                  dd      d        } |d   d       | j                  t              5   |d   d       d d d        y # 1 sw Y   y xY w)NTFr   c                     | rt         y N
ValueError)do_raises    r   r   z*TestException.test_user_raise.<locals>.foo%   s       r   r   r   )r   r   r#   r0   )r   r   s     r   test_user_raisezTestException.test_user_raise$   sV    	%	(	! 
)	! 	D	%z*CIdO +**s   AAc                    | }t        j                  ||      d        }t         j                  d        }d}dt        j                  |      dz   z  }dt        j                  |      dz   z  } |d|f   ||       dt        j                  |      dz   z  }dt        j                  |      dz   z  }	 |d|f   ||	       t        j                  j                  ||       t        j                  j                  |	|       y)	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                 B   t         j                  j                  }t         j                  j                  }|dkD  r&t	        |      D ]  }||xx   | |   ||   z  z  cc<    t        j
                          |dk  r't	        |      D ]  }| |xx   | |   ||   z  z  cc<    y y )N      r   r   r   blockDimrangesyncthreadsr   ytidntidis        r   problematiczBTestException.case_raise_causing_warp_diverge.<locals>.problematic7   s    ..""C==??DRxtAaDAaDAaDL(D % RxtAaDAaDAaDL(D % r   c                 f   t         j                  j                  }t         j                  j                  }|dkD  r/t	        |      D ]!  }||   dk7  s||xx   | |   ||   z  z  cc<   # t        j
                          |dk  r0t	        |      D ]!  }||   dk7  s| |xx   | |   ||   z  z  cc<   # y y )Nr6   r   r7   r8   r<   s        r   oraclez=TestException.case_raise_causing_warp_diverge.<locals>.oracleE   s    ..""C==??DRxtAtqy!!!, % RxtAtqy!!!, % r       g      ?g{Gz?r   N)r   r   r!   arangetestingassert_almost_equal)
r   with_debug_modewith_opt_moderA   rC   ngot_xgot_yexpect_xexpect_ys
             r   case_raise_causing_warp_divergez-TestException.case_raise_causing_warp_diverge.   s     ,+	]	;	) 
<	) 
	- 
	- biilT)*biilT)*AqD%'1,-1,-q!tXx(


&&x7


&&x7r   c                 (    | j                  d       y)z#Test case for issue #2655.
        F)rH   N)rO   )r   s    r   test_raise_causing_warp_divergez-TestException.test_raise_causing_warp_divergea   s     	,,U,Cr   c                 J   t         j                  d        }t        j                  d      }t        j                  d      }t        j                  d      } |d   |||       | j                  t        j                  |d         d       | j                  |d   |d   d       y )Nc                 0    |d   |d   z  | d<   |d   | d<   y Nr   r    rr   r=   s      r   fz4TestException.test_no_zero_division_error.<locals>.fm   $    Q4!A$;AaDQ4AaDr   r   r   r2   r   zExpected inf from div by zerozExpected execution to continue)r   r   r!   zerosones
assertTrueisinfassertEqual)r   rX   rW   r   r=   s        r   test_no_zero_division_errorz)TestException.test_no_zero_division_errori   s     
	 
	 HHQKHHQKGGAJ$1a1(GH1qt%EFr   c                    t        j                  dd      d        }t        j                  d      }t        j                  d      }t        j                  d      }t
        j                  rt        }nt        }| j                  |      5   |d   |||       d d d        | j                  |d   dd	       | j                  |d   dd
       y # 1 sw Y   6xY w)NTFr   c                 0    |d   |d   z  | d<   |d   | d<   y rT   rU   rV   s      r   rX   z:TestException.test_zero_division_error_in_debug.<locals>.f   rY   r   r   r   r2   r   z Expected result to be left unsetzExpected execution to stop)r   r   r!   rZ   r[   r   r    FloatingPointErrorZeroDivisionErrorr#   r^   )r   rX   rW   r   r=   excs         r   !test_zero_division_error_in_debugz/TestException.test_zero_division_error_in_debug{   s     
%	(	 
)	 HHQKHHQKGGAJ   $C#Cs#AdGAq! $ 	1q"DE1q">?	 $#s   CCc                 6   dt        j                  d      fd       t        j                  d      fd       }| j                  t              5 } |d           d d d        | j	                  t        j                               y # 1 sw Y   /xY w)NzDevice Function ErrorT)devicec                      t               r.   r/   )msgs   r   rX   z6TestException.test_raise_in_device_function.<locals>.f   s    S/!r   )r   c                                y r.   rU   )rX   s   r   kernelz;TestException.test_raise_in_device_function.<locals>.kernel   s    Cr   r2   )r   r   r#   r0   r%   r&   r'   )r   rk   raisesrX   ri   s      @@r   test_raise_in_device_functionz+TestException.test_raise_in_device_function   s    
 &			" 
	" 
		 
	 z*fF4LN + 	c3v//01 +*s   BB)__name__
__module____qualname__r   r+   r3   rO   rQ   r_   re   r   rm   __classcell__)r   s   @r   r	   r	      s?    >E*18fDG$@: 2 2r   r	   __main__)numpyr!   numbar   numba.cuda.testingr   r   r   
numba.corer   r	   rn   mainrU   r   r   <module>rx      s=      K K b2L b2J zHMMO r   