
    sh                         S SK r S SKJrJr  S SKJr  \" S5       " S S\5      5       r\S:X  a  \ R                  " 5         gg)    N)CUDATestCaseskip_on_cudasim)captured_stdoutz4cudasim doesn't support cuda import at non-top-levelc                   >   ^  \ rS rSrSrU 4S jrU 4S jrS rSrU =r	$ )TestReduction   z
Test shared memory reduction
c                 t   > [        5       U l        U R                  R                  5         [        TU ]  5         g N)r   _captured_stdout	__enter__supersetUpself	__class__s    ڀ/Users/tiagomarins/Projetos/claudeai/copy_bank/venv/lib/python3.13/site-packages/numba/cuda/tests/doc_examples/test_reduction.pyr   TestReduction.setUp   s)     / 1'')    c                 \   > U R                   R                  S S S 5        [        TU ]  5         g r
   )r   __exit__r   tearDownr   s    r   r   TestReduction.tearDown   s&    &&tT48r   c                   ^^^ SS K nSSKJm  SSKJm  TR                  UR                  S5      5      n[        U5      mTR                  UUU4S j5       nUST4   " U5        [        US   5        [        [        UR                  S5      5      5        UR                  R                  US   [        UR                  S5      5      5        g )Nr   )cuda)int32i   c                   > TR                   R                  n[        U 5      nX:  a  TR                  S5      nTR                  R                  TT5      nX   XA'   TR                  5         SnUTR                  R                  :  aN  USU-  -  S:X  a  XA==   XAU-      -  ss'   US-  nTR                  5         UTR                  R                  :  a  MN  US:X  a  XA   X'   g g g )N      r   )	threadIdxxlengridsharedarraysyncthreadsblockDim)	datatidsizeishrsr   r   nelems	         r   	array_sum2TestReduction.test_ex_reduction.<locals>.array_sum&   s    ..""Ct9DzIIaL kk''u57   "$--//)a!e})CaL0FA$$& $--//) !8 #DI ) r   r   )numpynumbar   numba.typesr   	to_devicearanger!   jitprintsumtestingassert_equal)r   npar.   r   r   r-   s       @@@r   test_ex_reductionTestReduction.test_ex_reduction   s    %
 NN299T?+A 
	) 
	)8 	!U(Aadc"))D/"# 	

!c"))D/&:;r   )r   )
__name__
__module____qualname____firstlineno____doc__r   r   r<   __static_attributes____classcell__)r   s   @r   r   r      s    
0< 0<r   r   __main__)	unittestnumba.cuda.testingr   r   numba.tests.supportr   r   r>   main r   r   <module>rK      sL     < / GH@<L @< I@<F zMMO r   