§
    ƒÛgƒ  ã                   ó„   — 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        ¦   «          dS dS )é    N)Úcuda)ÚunittestÚxfail_unless_cudasimÚCUDATestCase)Úconfigc                   óH   — e Zd Zd„ Zd„ Zd„ Zd„ Zd„ Zd„ Ze	d„ ¦   «         Z
dS )	ÚTestExceptionc                 óÒ  — d„ }t          j        |¦  «        } t          j        dd¬¦  «        |¦  «        }t          j        s% |d         t	          j        ddg¦  «        ¦  «         |                      t          ¦  «        5 } |d         t	          j        ddg¦  «        ¦  «         d d d ¦  «         n# 1 swxY w Y   |                      dt          |j
        ¦  «        ¦  «         d S )	Nc                 óT   — t           j        j        }|dk    r| j        |           d S d S )Né   )r   Ú	threadIdxÚxÚshape)Úaryr   s     ú\/granary/edev/quantum/lib/python3.11/site-packages/numba/cuda/tests/cudapy/test_exception.pyÚfooz)TestException.test_exception.<locals>.foo
   s/   € Ý”Ô ˆAØ�AŠvˆvð ”	˜1˜"”���ð	 ˆvó    TF©ÚdebugÚopt)é   é   r   r   ztuple index out of range)r   Újitr   ÚENABLE_CUDASIMÚnpÚarrayÚassertRaisesÚ
IndexErrorÚassertInÚstrÚ	exception)Úselfr   Ú
unsafe_fooÚsafe_fooÚcms        r   Útest_exceptionzTestException.test_exception	   s  € ð	ð 	ð 	õ ”X˜c‘]”]ˆ
Ø2•4”8 $¨EÐ2Ñ2Ô2°3Ñ7Ô7ˆåÔ$ð 	/ð ˆJ�tÔ�RœX q¨! fÑ-Ô-Ñ.Ô.Ð.à×Ò�zÑ*Ô*ð 	-¨bØˆH�TŒN�2œ8 Q¨ FÑ+Ô+Ñ,Ô,Ð,ð	-ð 	-ð 	-ñ 	-ô 	-ð 	-ð 	-ð 	-ð 	-ð 	-ð 	-øøøð 	-ð 	-ð 	-ð 	-à�ŠÐ0µ#°b´lÑ2CÔ2CÑDÔDÐDÐDÐDs   Â&B4Â4B8Â;B8c                 óî   — t          j        dd¬¦  «        d„ ¦   «         } |d         d¦  «         |                      t          ¦  «        5   |d         d¦  «         d d d ¦  «         d S # 1 swxY w Y   d S )NTFr   c                 ó   — | rt           ‚d S ©N©Ú
ValueError)Údo_raises    r   r   z*TestException.test_user_raise.<locals>.foo   s   € àð !Ý Ð ð!ð !r   ©r   r   )r   r   r   r+   )r"   r   s     r   Útest_user_raisezTestException.test_user_raise   sÊ   € Ý	Œ˜ %Ð	(Ñ	(Ô	(ð	!ð 	!ñ 
)Ô	(ð	!ð 	ˆˆDŒ	�%ÑÔÐØ×Ò�zÑ*Ô*ð 	ð 	ØˆC�ŒI�d‰OŒOˆOð	ð 	ð 	ñ 	ô 	ð 	ð 	ð 	ð 	ð 	ð 	ð 	øøøð 	ð 	ð 	ð 	ð 	ð 	s   ÁA*Á*A.Á1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                             ||¦  «         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                 óP  — t           j        j        }t           j        j        }|dk    r1t	          |¦  «        D ]!}||xx         | |         ||         z  z  cc<   Œ"t          j        ¦   «          |dk     r1t	          |¦  «        D ]#}| |xx         | |         ||         z  z  cc<   Œ"d S d S )Né   é   ©r   r   r   ÚblockDimÚrangeÚsyncthreads©r   ÚyÚtidÚntidÚis        r   ÚproblematiczBTestException.case_raise_causing_warp_diverge.<locals>.problematic1   s¹   € å”.Ô"ˆCÝ”=”?ˆDà�RŠxˆxÝ˜t™œð )ð )�AØ�a�D�D”D˜A˜aœD A a¤D™LÑ(�D�D‘D�DåÔÑÔÐØ�RŠxˆxÝ˜t™œð )ð )�AØ�a�D�D”D˜A˜aœD A a¤D™LÑ(�D�D‘D�Dð ˆxð)ð )r   c                 ó€  — t           j        j        }t           j        j        }|dk    r=t	          |¦  «        D ]-}||         dk    r||xx         | |         ||         z  z  cc<   Œ.t          j        ¦   «          |dk     r=t	          |¦  «        D ]/}||         dk    r| |xx         | |         ||         z  z  cc<   Œ.d S d S )Nr1   r   r2   r3   r7   s        r   Úoraclez=TestException.case_raise_causing_warp_diverge.<locals>.oracle?   sÑ   € å”.Ô"ˆCÝ”=”?ˆDà�RŠxˆxÝ˜t™œð -ð -�AØ˜”t˜q’y�yØ˜!˜˜œ  !¤¨¨!¬¡Ñ,˜˜™øåÔÑÔÐØ�RŠxˆxÝ˜t™œð -ð -�AØ˜”t˜q’y�yØ˜!˜˜œ  !¤¨¨!¬¡Ñ,˜˜™øð ˆxð-ð -r   é    g      ð?g{®Gáz„?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_ys
             r   Úcase_raise_causing_warp_divergez-TestException.case_raise_causing_warp_diverge(   s   € ð ,Ð+ˆå	Œ˜¨]Ð	;Ñ	;Ô	;ð	)ð 	)ñ 
<Ô	;ð	)õ 
Œð	-ð 	-ñ 
Œð	-ð ˆØ•b”i ‘l”l TÑ)Ñ*ˆØ•b”i ‘l”l TÑ)Ñ*ˆØˆ�A�q�DÔ˜% Ñ'Ô'Ð'à�œ 1™œ¨Ñ,Ñ-ˆØ�œ 1™œ¨Ñ,Ñ-ˆØˆˆq�!ˆtŒ�X˜xÑ(Ô(Ð(å
Œ
×&Ò& x°Ñ7Ô7Ð7Ý
Œ
×&Ò& x°Ñ7Ô7Ð7Ð7Ð7r   c                 ó2   — |                       d¬¦  «         dS )z#Test case for issue #2655.
        F)rC   N)rJ   )r"   s    r   Útest_raise_causing_warp_divergez-TestException.test_raise_causing_warp_diverge[   s!   € ð 	×,Ò,¸UÐ,ÑCÔCÐCÐCÐCr   c                 óp  — t           j        d„ ¦   «         }t          j        d¦  «        }t          j        d¦  «        }t          j        d¦  «        } |d         |||¦  «         |                      t          j        |d         ¦  «        d¦  «         |                      |d         |d         d¦  «         d S )Nc                 óD   — |d         |d         z  | d<   |d         | d<   d S ©Nr   r   © ©Úrr   r8   s      r   Úfz4TestException.test_no_zero_division_error.<locals>.fg   ó(   € à�Q”4˜!˜Aœ$‘;ˆAˆa‰DØ�Q”4ˆAˆa‰DˆDˆDr   r   r   r-   r   zExpected inf from div by zerozExpected execution to continue)r   r   r   ÚzerosÚonesÚ
assertTrueÚisinfÚassertEqual)r"   rS   rR   r   r8   s        r   Útest_no_zero_division_errorz)TestException.test_no_zero_division_errorc   s¨   € õ 
Œð	ð 	ñ 
Œð	õ ŒH�Q‰KŒKˆÝŒH�Q‰KŒKˆÝŒG�A‰JŒJˆàˆˆ$Œ��1�aÑÔÐà�Š�œ  1¤™œÐ(GÑHÔHÐHØ×Ò˜˜1œ˜q œtÐ%EÑFÔFÐFÐFÐFr   c                 óæ  — t          j        dd¬¦  «        d„ ¦   «         }t          j        d¦  «        }t          j        d¦  «        }t          j        d¦  «        }t
          j        rt          }nt          }|  	                    |¦  «        5   |d         |||¦  «         d d d ¦  «         n# 1 swxY w Y   |  
                    |d         dd	¦  «         |  
                    |d         dd
¦  «         d S )NTFr   c                 óD   — |d         |d         z  | d<   |d         | d<   d S rO   rP   rQ   s      r   rS   z:TestException.test_zero_division_error_in_debug.<locals>.fy   rT   r   r   r   r-   r   z Expected result to be left unsetzExpected execution to stop)r   r   r   rU   rV   r   r   ÚFloatingPointErrorÚZeroDivisionErrorr   rY   )r"   rS   rR   r   r8   Úexcs         r   Ú!test_zero_division_error_in_debugz/TestException.test_zero_division_error_in_debugu   s2  € õ 
Œ˜ %Ð	(Ñ	(Ô	(ð	ð 	ñ 
)Ô	(ð	õ ŒH�Q‰KŒKˆÝŒH�Q‰KŒKˆÝŒG�A‰JŒJˆõ Ô ð 	$Ý$ˆCˆCå#ˆCà×Ò˜sÑ#Ô#ð 	ð 	ØˆAˆdŒG�A�q˜!ÑÔÐð	ð 	ð 	ñ 	ô 	ð 	ð 	ð 	ð 	ð 	ð 	øøøð 	ð 	ð 	ð 	ð 	×Ò˜˜1œ˜qÐ"DÑEÔEÐEØ×Ò˜˜1œ˜qÐ">Ñ?Ô?Ð?Ð?Ð?s   ÂB,Â,B0Â3B0c                 ób  ‡‡— dŠt          j        d¬¦  «        ˆfd„¦   «         Št          j        d¬¦  «        ˆfd„¦   «         }|                      t          ¦  «        5 } |d         ¦   «          d d d ¦  «         n# 1 swxY w Y   |                      ‰t          |j        ¦  «        ¦  «         d S )NzDevice Function ErrorT)Údevicec                  ó"   •— t          ‰ ¦  «        ‚r)   r*   )Úmsgs   €r   rS   z6TestException.test_raise_in_device_function.<locals>.f™   s   ø€ å˜S‘/”/Ð!r   )r   c                  ó   •—  ‰ ¦   «          d S r)   rP   )rS   s   €r   Úkernelz;TestException.test_raise_in_device_function.<locals>.kernel�   s   ø€ àˆA‰CŒCˆCˆCˆCr   r-   )r   r   r   r+   r   r    r!   )r"   rf   ÚraisesrS   rd   s      @@r   Útest_raise_in_device_functionz+TestException.test_raise_in_device_function’   s  øø€ ð
 &ˆå	Œ˜Ð	Ñ	Ô	ð	"ð 	"ð 	"ð 	"ñ 
Ô	ð	"õ 
Œ˜Ð	Ñ	Ô	ð	ð 	ð 	ð 	ñ 
Ô	ð	ð ×Ò�zÑ*Ô*ð 	¨fØˆF�4ŒL‰NŒNˆNð	ð 	ð 	ñ 	ô 	ð 	ð 	ð 	ð 	ð 	ð 	øøøð 	ð 	ð 	ð 	ð 	�Š�c�3˜vÔ/Ñ0Ô0Ñ1Ô1Ð1Ð1Ð1s   ÁA<Á<B ÂB N)Ú__name__Ú
__module__Ú__qualname__r&   r.   rJ   rL   rZ   r`   r   rh   rP   r   r   r	   r	      s—   € € € € € ðEð Eð Eð*ð ð ð18ð 18ð 18ðfDð Dð DðGð Gð Gð$@ð @ð @ð: ð2ð 2ñ Ôð2ð 2ð 2r   r	   Ú__main__)Únumpyr   Únumbar   Únumba.cuda.testingr   r   r   Ú
numba.corer   r	   ri   ÚmainrP   r   r   ú<module>rr      s¬   ðØ Ð Ð Ð à Ð Ð Ð Ð Ð Ø KÐ KÐ KÐ KÐ KÐ KÐ KÐ KÐ KÐ KØ Ð Ð Ð Ð Ð ð\2ð \2ð \2ð \2ð \2�Lñ \2ô \2ð \2ð~ ˆzÒÐØ€H„M�O„O€O€O€Oð Ðr   