
    sh                         S SK rS SKJrJrJr  S SKJrJr  S SK	J
r
  \
R                  (       a  SrOSr\\4r " S S\5      r\S	:X  a  \R                   " 5         gg)
    N)cudafloat64void)unittestCUDATestCase)config      c                       \ rS rSrS rSrg)TestCudaLaplace   c           
      t  ^ [         R                  " [        [        [        5      SSS9S 5       m[         R                  " [        [        S S 2S S 24   [        S S 2S S 24   [        S S 2S S 24   5      5      U4S j5       n[        R
                  (       a  Su  p#SnOSu  p#Sn[        R                  " X#4[        R                  S	9n[        R                  " X#4[        R                  S	9nUnS
nSn	[        U5       H  n
SXZS4'   SXjS4'   M     Sn[        [        4nX,S   -  X<S   -  4n[        R                  " U5      n[         R                  " 5       n[         R                  " X_5      n[         R                  " Xo5      n[         R                  " X5      nX:  a  X:  a  U R                  UR                  [        R                  :H  5        XX4   " UUU5        UR                  XS9  UR                  5         [        R                   " U5      R#                  5       n	UnUnUnUS-  nX:  a  X:  a  M  g g g g )NT)deviceinlinec                     X:  a  U $ U$ )N )abs     x/Users/tiagomarins/Projetos/claudeai/copy_bank/venv/lib/python3.13/site-packages/numba/cuda/tests/cudapy/test_laplace.pyget_max3TestCudaLaplace.test_laplace_small.<locals>.get_max   s    u    c                   > [         R                  R                  [        [        S9n[         R
                  R                  n[         R
                  R                  n[         R                  R                  n[         R                  R                  nU R                  S   nU R                  S   n	[         R                  " S5      u  pSX4U4'   US:  aU  XS-
  :  aM  U
S:  aG  XS-
  :  a?  SXU
S-   4   XU
S-
  4   -   XS-
  U
4   -   XS-   U
4   -   -  XU
4'   XU
4   XU
4   -
  X4U4'   [         R                  " 5         [        S-  nUS:  a?  XL:  a  T" X4U4   X4U-   U4   5      X4U4'   US-  n[         R                  " 5         US:  a  M?  [        S-  nUS:  aD  X\:  a  US:X  a  T" X4U4   X4X\-   4   5      X4U4'   US-  n[         R                  " 5         US:  a  MD  US:X  a  US:X  a
  US   X'U4'   g g g )Ndtyper         g      ?)r   r   )r   sharedarraySM_SIZEr   	threadIdxxyblockIdxshapegridsyncthreadstpb)AAnewerrorerr_smtytxbxbynmijtr   s                r   jocabi_relax_core=TestCudaLaplace.test_laplace_small.<locals>.jocabi_relax_core   s   [[&&wg&>F!!B!!BBB
A
A99Q<DAFr6NAv!!e)Q11u9!a1q5kAQhK&?()a%('467Aqk'B CT
!%da1g!52v qAa%6%,VF^VFBJ=O%PFr6Na  "	 a% qAa%6bAg%,VF^VJ=O%PFr6Na  "	 a% Qw27 &t"f #wr   )r	   r	      )   r9   i  r   gư>g      ?r   r   )stream)r   jitr   r   r   ENABLE_CUDASIMnpzerosranger(   r:   	to_device
assertTruer   copy_to_hostsynchronizeabsmax)selfr6   NNNMiter_maxr)   r*   r1   tolr+   r4   iterblockdimgriddim
error_gridr:   dAdAnewderror_gridtmpr   s                       @r   test_laplace_small"TestCudaLaplace.test_laplace_small   s   	''7+D	F	 
G	 
$wq!t}gadmWQT]C	D&	- 
E&	-P   FBHFBHHHbXRZZ0xx

3qAAdGDAJ  :!$bQK&78XXg&
^^A&t,nnZ8kdoOOJ,,

:;x78UKP$$Z$?  FF:&**,E CBEAID# kdokokr   r   N)__name__
__module____qualname____firstlineno__rS   __static_attributes__r   r   r   r   r      s    cr   r   __main__)numpyr=   numbar   r   r   numba.cuda.testingr   r   
numba.corer   r<   r(   r    r   rU   mainr   r   r   <module>r`      sY     % % 5  

C
C
s(dl dN zMMO r   