U
    ?h                     @   sP   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rLe   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__ ^/var/www/html/venv/lib/python3.8/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 r d} j}| | ||<    d}| jjk r|d|  dkr|||  |||  7  < |d9 }   qH|dkr|| | |< d S )N      r   )Z	threadIdxxlengridZsharedarrayZsyncthreadsZblockDim)datatidsizeiZshrsr   r   Znelemr   r   	array_sum&   s    

z2TestReduction.test_ex_reduction.<locals>.array_sumr   )numpyZnumbar   Znumba.typesr   Z	to_deviceZaranger   ZjitprintsumZtestingZassert_equal)r   npar!   r   r    r   test_ex_reduction   s    zTestReduction.test_ex_reduction)__name__
__module____qualname____doc__r
   r   r'   __classcell__r   r   r   r   r      s   r   __main__)	ZunittestZnumba.cuda.testingr   r   Znumba.tests.supportr   r   r(   mainr   r   r   r   <module>   s   C