o
    Û¾iâ  ã                   @   sT   d dl Z d dlmZmZ d dlmZ edƒG dd„ deƒƒZedkr(e  ¡  dS dS )é    N)ÚCUDATestCaseÚskip_on_cudasim)Úcaptured_stdoutz4cudasim doesn't support cuda import at non-top-levelc                       s4   e Zd ZdZ‡ fdd„Z‡ fdd„Zdd„ Z‡  ZS )ÚTestReductionz&
    Test shared memory reduction
    c                    s    t ƒ | _| j ¡  tƒ  ¡  d S ©N)r   Ú_captured_stdoutÚ	__enter__ÚsuperÚsetUp©Úself©Ú	__class__© ú`/home/ubuntu/.local/lib/python3.10/site-packages/numba/cuda/tests/doc_examples/test_reduction.pyr
      s   
zTestReduction.setUpc                    s   | j  d d d ¡ tƒ  ¡  d S r   )r   Ú__exit__r	   ÚtearDownr   r   r   r   r      s   zTestReduction.tearDownc                    sœ   dd l }ddlm‰  ddlm‰ ˆ  | d¡¡}t|ƒ‰ˆ j‡ ‡‡fdd„ƒ}|dˆf |ƒ t	|d ƒ t	t
| d¡ƒƒ |j |d t
| d¡ƒ¡ d S )Nr   )Úcuda)Úint32i   c                    s¸   ˆ j j}t| ƒ}||k rXˆ  d¡}ˆ j ˆˆ¡}| | ||< ˆ  ¡  d}|ˆ jjk rL|d|  dkr>||  |||  7  < |d9 }ˆ  ¡  |ˆ jjk s*|dkrZ|| | |< d S d S d S )Né   é   r   )Ú	threadIdxÚxÚlenÚgridÚsharedÚarrayÚsyncthreadsÚblockDim)ÚdataÚtidÚsizeÚiÚshrÚs©r   r   Únelemr   r   Ú	array_sum&   s$   
ûëz2TestReduction.test_ex_reduction.<locals>.array_sumr   )ÚnumpyÚnumbar   Únumba.typesr   Ú	to_deviceÚaranger   ÚjitÚprintÚsumÚtestingÚassert_equal)r   ÚnpÚar'   r   r%   r   Útest_ex_reduction   s    zTestReduction.test_ex_reduction)Ú__name__Ú
__module__Ú__qualname__Ú__doc__r
   r   r4   Ú__classcell__r   r   r   r   r      s
    r   Ú__main__)	ÚunittestÚnumba.cuda.testingr   r   Únumba.tests.supportr   r   r5   Úmainr   r   r   r   Ú<module>   s    Cÿ