U
    -e                     @   st   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rDdZndZeefZG dd deZed	krpe  dS )
    N)cudafloat64void)unittestCUDATestCase)config      c                   @   s   e Zd Zdd ZdS )TestCudaLaplacec              	      s  t jtttddddd  t ttd d d d f td d d d f td d d d f  fdd}tjrd\}}d}nd	\}}d
}tj||ftjd}tj||ftjd}|}d}d}	t|D ]}
d||
df< d||
df< qd}t	t	f}||d  ||d  f}t|}t 
 }t ||}t ||}t ||}|	|kr||k r| |jtjk ||||f ||| |j||d |  t| }	|}|}|}|d7 }q@d S )NT)Zdeviceinlinec                 S   s   | |kr| S |S d S )N )abr   r   e/var/www/html/Darija-Ai-Train/env/lib/python3.8/site-packages/numba/cuda/tests/cudapy/test_laplace.pyget_max   s    z3TestCudaLaplace.test_laplace_small.<locals>.get_maxc                    s  t jjttd}t jj}t jj}t jj}t jj}| j	d }| j	d }	t 
d\}
}d|||f< |dkr||d k r|
dkr|
|	d k rd| ||
d f | ||
d f  | |d |
f  | |d |
f   |||
f< |||
f | ||
f  |||f< t   td }|dkrP||k r< |||f ||| |f |||f< |d }t   qtd }|dkr||k r|dkr |||f |||| f |||f< |d }t   qX|dkr|dkr|d |||f< d S )Ndtyper         g      ?)r   r   )r   ZsharedarraySM_SIZEr   Z	threadIdxxyZblockIdxshapegridZsyncthreadstpb)AAnewerrorZerr_smtyZtxZbxZbynmijtr   r   r   jocabi_relax_core   s>    

(  

&
&z=TestCudaLaplace.test_laplace_small.<locals>.jocabi_relax_core)r   r      )   r(   i  r   gư>g      ?r   r   )stream)r   Zjitr   r   r   ENABLE_CUDASIMnpZzerosranger   r)   Z	to_device
assertTruer   Zcopy_to_hostZsynchronizeabsmax)selfr&   ZNNZNMZiter_maxr   r   r    Ztolr   r#   iterZblockdimZgriddimZ
error_gridr)   ZdAZdAnewZderror_gridtmpr   r%   r   test_laplace_small   sF    
@(
z"TestCudaLaplace.test_laplace_smallN)__name__
__module____qualname__r3   r   r   r   r   r
      s   r
   __main__)numpyr+   Znumbar   r   r   Znumba.cuda.testingr   r   Z
numba.corer   r*   r   r   r
   r4   mainr   r   r   r   <module>   s   g