
    xKg                         d dl Zd dlmZmZmZ d dlmZmZ d dl	m
Z
 e
j                  rdZndZeefZ G d de      Zed	k(  r ej                           yy)
    N)cudafloat64void)unittestCUDATestCase)config      c                       e Zd Zd Zy)TestCudaLaplacec           
         t        j                  t        t        t              dd      d        t        j                  t        t        d d d d f   t        d d d d f   t        d d d d f               fd       }t        j
                  rd\  }}d}nd\  }}d}t        j                  ||ft        j                  	      }t        j                  ||ft        j                  	      }|}d
}d}	t        |      D ]  }
d||
df<   d||
df<    d}t        t        f}||d   z  ||d   z  f}t        j                  |      }t        j                         }t        j                  ||      }t        j                  ||      }t        j                  ||      }|	|kD  r||k  r| j                  |j                  t        j                  k(          ||||f   |||       |j                  ||       |j                          t        j                   |      j#                         }	|}|}|}|dz  }|	|kD  r||k  ry y y y )NT)deviceinlinec                     | |kD  r| S |S )N )abs     h/home/alanp/www/video.onchill/myenv/lib/python3.12/site-packages/numba/cuda/tests/cudapy/test_laplace.pyget_maxz3TestCudaLaplace.test_laplace_small.<locals>.get_max   s    1u    c                    t         j                  j                  t        t              }t         j
                  j                  }t         j
                  j                  }t         j                  j                  }t         j                  j                  }| j                  d   }| j                  d   }	t        j                  d      \  }
}d|||f<   |dk\  r\||dz
  k  rT|
dk\  rO|
|	dz
  k  rGd| ||
dz   f   | ||
dz
  f   z   | |dz
  |
f   z   | |dz   |
f   z   z  |||
f<   |||
f   | ||
f   z
  |||f<   t        j                          t        dz  }|dkD  r?||k  r |||f   |||z   |f         |||f<   |dz  }t        j                          |dkD  r?t        dz  }|dkD  rD||k  r |dk(  r |||f   ||||z   f         |||f<   |dz  }t        j                          |dkD  rD|dk(  r|dk(  r|d   |||f<   y y y )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_corez=TestCudaLaplace.test_laplace_small.<locals>.jocabi_relax_core   s#   [[&&wg&>F!!B!!BBB
A
A99Q<DAqF2r6NAv!a!e)Q1q1u9!a1q5kAaQhK&?()!a%('467Aqk'B CQT
!%ada1g!5r2v qAa%6%,VBF^VBFBJ=O%PF2r6Na  "	 a% qAa%6bAg%,VBF^VBQJ=O%PF2r6Na  "	 a% Qw27 &tb"f #wr   )r	   r	      )   r6   i  r   gư>g      ?r   r   )stream)r   jitr   r   r   ENABLE_CUDASIMnpzerosranger&   r7   	to_device
assertTruer   copy_to_hostsynchronizeabsmax)selfr4   NNNMiter_maxr'   r(   r/   tolr)   r2   iterblockdimgriddim
error_gridr7   dAdAnewderror_gridtmpr   s                       @r   test_laplace_smallz"TestCudaLaplace.test_laplace_small   s   	''7+D	F	 
G	 
$wq!t}gadmWQT]C	D&	- 
E&	-P   FBHFBHHHb"XRZZ0xxR

3qAAadGDAJ  :!$bHQK&78XXg&
^^Av&tV,nnZ8ckdXoOOJ,,

:;8gx78UKP$$Z$?  FF:&**,E CBEAID# ckdXokokr   N)__name__
__module____qualname__rP   r   r   r   r   r      s    cr   r   __main__)numpyr:   numbar   r   r   numba.cuda.testingr   r   
numba.corer   r9   r&   r   r   rQ   mainr   r   r   <module>rZ      s[     % % 5  

C
C
s(dl dN zHMMO r   